UCT/ROCM: initial commit for device initiated ipc put by mshanthagit · Pull Request #11299 · openucx/ucx

mshanthagit · 2026-03-25T18:42:14Z

What?

This PR introduces device initiated IPC put operation for the rocm component.

Why?

It is a new functionality to enable device initiated communication.

How?

It is optional, but for complex PRs, please provide information about the design,
architecture, approach, etc.

mshanthagit · 2026-04-08T17:32:00Z

This PR is now ready for review.

mshanthagit

@ofirfarjun7 addressed your comments.

Artemy-Mellanox · 2026-05-07T14:33:51Z

    ucp_context_h context  = select_params->ep->worker->context;
    unsigned ep_init_flags = ucp_wireup_ep_init_flags(select_params,
                                                      select_ctx);
+    static const ucs_memory_type_t device_mem_types[] = {


maybe add const UCP_MEMORY_TYPES_DEVICE
(UCS_BIT(UCS_MEMORY_TYPE_CUDA) | UCS_BIT(UCS_MEMORY_TYPE_ROCM)
and use ucs_memory_type_for_each to iterate skiping non-DEVICE memtypes, here and in ucp_device_detect_export_mem_type

addressed this comment. added a #define in ucp_mm.h and used ucs_for_each_bit(mem_type, UCP_DEVICE_MEM_TYPES) within the files.

Artemy-Mellanox · 2026-05-07T14:36:42Z

        return status;
    }

    status = ucp_device_remote_mem_list_create_handle(params, export_mem_type,


pass first ep so it won't need to look again

Artemy-Mellanox · 2026-05-07T14:36:57Z

    ucs_status_t status;
    uct_allocated_memory_t mem;
+    ucs_sys_device_t sys_dev;
+    ucp_ep_h ep = ucp_device_remote_mem_list_get_first_ep(params);


initialized vars first

Artemy-Mellanox · 2026-05-07T14:42:28Z

+}
+
+uct_rocm_ipc_component_t uct_rocm_ipc_component = {
+    .super     = {.query_md_resources = uct_rocm_base_query_md_resources,


strange code format

formatted it.

mshanthagit

@Artemy-Mellanox incorporated your suggestions in the latest commit.

mshanthagit · 2026-05-08T04:56:17Z

+}
+
+uct_rocm_ipc_component_t uct_rocm_ipc_component = {
+    .super     = {.query_md_resources = uct_rocm_base_query_md_resources,


formatted it.

mshanthagit · 2026-05-08T04:58:07Z

    ucp_context_h context  = select_params->ep->worker->context;
    unsigned ep_init_flags = ucp_wireup_ep_init_flags(select_params,
                                                      select_ctx);
+    static const ucs_memory_type_t device_mem_types[] = {


addressed this comment. added a #define in ucp_mm.h and used ucs_for_each_bit(mem_type, UCP_DEVICE_MEM_TYPES) within the files.

mshanthagit · 2026-05-08T04:58:33Z

    ucs_status_t status;
    uct_allocated_memory_t mem;
+    ucs_sys_device_t sys_dev;
+    ucp_ep_h ep = ucp_device_remote_mem_list_get_first_ep(params);


mshanthagit · 2026-05-08T04:58:40Z

        return status;
    }

    status = ucp_device_remote_mem_list_create_handle(params, export_mem_type,


mshanthagit · 2026-05-13T13:59:12Z

@Artemy-Mellanox did you get a chance to review the changes?

edgargabriel · 2026-05-15T21:52:17Z

I just wanted to add a comment here that I know that I am listed as reviewer and I am happy to sign off on the PR. @mshanthagit and me have discussed the issue and the actual ticket internally multiple times, so I am very well aware of the technical aspects of it.

mshanthagit · 2026-05-24T03:39:46Z

@ofirfarjun7 addressed your comments except the one related to "UCT_CUDA_IPC_IS_ALIGNED_POW2". Can we move it to ucs/sys/device_code.h and name it as UCT_IPC_IS_ALIGNED_POW2?

mshanthagit · 2026-06-02T16:12:59Z

@ofirfarjun7 rebased and addressed all comments.

ofirfarjun7 · 2026-06-07T15:33:08Z

 #define UCS_DEVICE_NUM_THREADS_IN_WARP 32

+/* Check if _n is aligned to a power-of-2 boundary _p */
+#define UCT_IPC_IS_ALIGNED_POW2(_n, _p) (!((_n) & ((_p) - 1)))


maybe UCS_DEVICE_IS_ALIGNED_POW2?

makes sense, my bad.

Artemy-Mellanox · 2026-06-23T14:59:09Z

+
+HIPCC ?= hipcc
+
+HIPCC_CMD = $(HIPCC) -DHAVE_CONFIG_H -DUCT_DEVICE_CODE_HIP -fPIE -I$(top_srcdir)/src -I$(top_builddir)/src $(BASE_CXXFLAGS) $(CXXFLAGS) $(HIP_CPPFLAGS) $(HIP_CXXFLAGS) $(HIPCC_EXTRA_FLAGS) -Wno-c++20-extensions -c $< -MT $@ -MF $(DEPDIR)/hip/$@.d -MMD -o $@


better $(DEFS) for -DHAVE_CONFIG_H

Artemy-Mellanox · 2026-06-23T15:01:31Z

-    ucs_status_t status;
-    uct_allocated_memory_t mem;
+    const ucp_worker_h worker = UCS_PARAM_VALUE(UCP_DEVICE_MEM_LIST_PARAMS_FIELD,
+                                          params, worker, WORKER, NULL);


Artemy-Mellanox · 2026-06-24T11:59:06Z

+            if (status != UCS_OK) {
+                ucs_error("Failed to create ROCm IPC component cache: %s",
+                        ucs_status_string(status));
+                pthread_mutex_unlock(&uct_rocm_ipc_component.lock);


where is it locked? isn't it leftover?

Yes, missed cleaning it up.

Artemy-Mellanox · 2026-06-24T12:02:27Z

+    ucs_status_t status;
+
+    UCS_INIT_ONCE(&cache_init_once) {
+        if (uct_rocm_ipc_component.ipc_cache == NULL) {


isn't UCS_INIT_ONCE already do this?

Artemy-Mellanox · 2026-06-24T12:18:49Z

+        UCS_DEVICE_IS_ALIGNED_POW2((intptr_t)d1, sizeof(vec4))) {
+        const vec4 *s4 = reinterpret_cast<const vec4*>(s1);
+        vec4 *d4       = reinterpret_cast<vec4*>(d1);
+        size_t n4      = len / sizeof(vec4);


n4 is num_lines in block, what's diff?
why warp/block should have separate implementations?

merged the implementations.

Artemy-Mellanox · 2026-06-24T12:24:11Z

+
+/* System-wide atomic increment */
+__device__ static inline void
+uct_rocm_ipc_atomic_inc(uint64_t *dst, uint64_t inc_value)


I'd call it atomic_add

Artemy-Mellanox · 2026-06-24T12:34:01Z

+
+#define UCS_DEVICE_LEVEL_EXEC_ID 1
+
+#define UCS_DEVICE_LEVEL_EXEC_SELECT(scope_ok, count, id) \


couldn't it be static __device__ function?

Artemy-Mellanox · 2026-06-24T12:35:36Z

+ {
+    unsigned int thread_id   = threadIdx.x;
+    unsigned int num_threads = blockDim.x;
+    unsigned int warp_id     = thread_id / 64;  // ROCm wavefront size


why not use UCT_ROCM_IPC_WAVEFRONT_SIZE ?

Artemy-Mellanox · 2026-06-24T12:53:43Z

 if HAVE_GNUXX11
+HIPCC_EXTRA_FLAGS = \
+	$(HIP_CPPFLAGS) \
+	-DHAVE_ROCM=1 \


why not use __HIPCC__?

mshanthagit · 2026-06-26T05:54:17Z

@ofirfarjun7 @Artemy-Mellanox if the changes are satisfactory and the PR is ready to be merged, I would first like to squash the commits with an appropriate commit message and force push it (unless it allows me to squash it online).

ofirfarjun7 · 2026-06-28T08:41:41Z

@mshanthagit please fix all issues that cause CI to fail.
Please check https://github.com/openucx/ucx/wiki/Guidance-for-contributors

mshanthagit · 2026-06-29T05:55:36Z

@mshanthagit please fix all issues that cause CI to fail. Please check https://github.com/openucx/ucx/wiki/Guidance-for-contributors

@ofirfarjun7 looks like the CI issues are related to "docker pull" failures. Please let me know if I missed anything obvious.

mshanthagit · 2026-07-21T13:11:40Z

@ofirfarjun7 @Artemy-Mellanox I have rebased and there are no conflicts. There seems to be a failure, but I suspect it is not related to this PR. Please let me know if there's anything else blocking the approval. As noted earlier, I would like to squash the PR into a single commit, shall I do that now?

ofirfarjun7 · 2026-07-28T12:53:27Z

    self->remote_pid = *(const pid_t*)params->iface_addr;
-
-    snprintf(target_name, sizeof(target_name), "dest:%d", *(pid_t*)params->iface_addr);
-    status = uct_rocm_ipc_create_cache(&self->remote_memh_cache, target_name);


Can it create issues if cache is not private for each EP and lookup criteria is only remote VA?

In addition to the remote VA, ipc handles are compared as well to avoid conflicts. Do you foresee any issues with this appoach?

But if ipc handles are not the same you remove it from the cache because it is "stale region".
But what if two remote processes has two regions with same VA?

In that case, we might see a ping-pong of effect that may impact performance, as the process has to now go though hsa_amd_ipc_memory_attach instead of using cached entry.

Maybe add another dimension to the key? You can examine cuda ipc cache.
If you think it will just impact performance and you prefer to leave it as is let me know.

edgargabriel · 2026-08-03T15:18:36Z

Could we try to get this PR over the finish line and have it merged? It has been opened on March 15, so its now nearly five months that it is pending.

ofirfarjun7 · 2026-08-04T11:10:12Z

    self->remote_pid = *(const pid_t*)params->iface_addr;
-
-    snprintf(target_name, sizeof(target_name), "dest:%d", *(pid_t*)params->iface_addr);
-    status = uct_rocm_ipc_create_cache(&self->remote_memh_cache, target_name);


But if ipc handles are not the same you remove it from the cache because it is "stale region".
But what if two remote processes has two regions with same VA?

ofirfarjun7 · 2026-08-06T06:52:37Z

Could we try to get this PR over the finish line and have it merged? It has been opened on March 15, so its now nearly five months that it is pending.

If these two are not a concern from your point of view let us know and we will finish the review
Can't disable cache for device:
https://github.com/openucx/ucx/pull/11299/changes/BASE..4dddf15e46735555405bf678be778a23358ec45f#r3466948308

Potential performance issue:
#11299 (comment)

anjoj0 · 2026-08-06T09:03:46Z

@edgargabriel @mshanthagit Cross-linking the W7900 compatibility result from ai-dynamo/nixl#2039.

I tested #11299 head 4dddf15e4 plus only UCT_IFACE_FLAG_ERRHANDLE_PEER_FAILURE. The branch built successfully; the PR's ROCm IPC tests reported 92 passed / 40 expected skipped / 0 failed. NIXL default peer mode selected rocm_ipc/rocm_ipc at 27.377/23.527 GB/s (same NUMA) and 27.391/23.522 GB/s (cross NUMA), with verified payloads. Twelve legal cross-NUMA peer-exit injections returned verified completion or NIXL_ERR_REMOTE_DISCONNECT without a hang.

No handle-cache or device-initiated PUT code was changed. Full methods and hashes:
https://github.com/anjoj0/vllm-awq4-qwen-1.0/blob/main/w7900_optimization/results/20260806_ucx_pr11299_compatibility_report.md

Would you prefer the capability flag and a multi-process peer-exit regression to be absorbed into this PR, or handled as a small dependent follow-up?

edgargabriel · 2026-08-06T19:39:26Z

Could we try to get this PR over the finish line and have it merged? It has been opened on March 15, so its now nearly five months that it is pending.

If these two are not a concern from your point of view let us know and we will finish the review Can't disable cache for device: https://github.com/openucx/ucx/pull/11299/changes/BASE..4dddf15e46735555405bf678be778a23358ec45f#r3466948308

Potential performance issue: #11299 (comment)

@ofirfarjun7 thank you, those things are a concern for us, and we discussed this with @mshanthagit internally. We think the most efficient way forward would be to merge this PR with a clear understanding of the shortcomings. The host-path is from the correctness perspective not affected (so we do not introduce a regression). There is potential for a performance regression in certain scenarios. However, because this is a massive PR we think its better to get it merged and fix the caching issue in a subsequent PR (which might come as soon as next week). The follow-up PR would only touch rocm_ipc and would have no impact on the rest of the UCX code base.

With this current PR, we are constantly rebasing to pull in updates to the repo, and had to resolve a number of conflicts over the weeks, its one of the reasons we would see benefits in merging this PR and fix the problems in a follow-up PR.

Introduces device initiated IPC put operation for the ROCm component. - Add device-side rocm_ipc.h with HIP __device__ put/atomic_add/get_ptr, including a strided int4/int2 vectorized copy at thread/warp/block levels. - Add MD mem_elem_pack/rkey_ptr that map the remote IPC handle via a shared, lazily-initialized process-wide handle cache; iface/ep now expose a device ep.

ofirfarjun7 · 2026-08-07T15:18:58Z

Could we try to get this PR over the finish line and have it merged? It has been opened on March 15, so its now nearly five months that it is pending.

If these two are not a concern from your point of view let us know and we will finish the review Can't disable cache for device: https://github.com/openucx/ucx/pull/11299/changes/BASE..4dddf15e46735555405bf678be778a23358ec45f#r3466948308
Potential performance issue: #11299 (comment)

@ofirfarjun7 thank you, those things are a concern for us, and we discussed this with @mshanthagit internally. We think the most efficient way forward would be to merge this PR with a clear understanding of the shortcomings. The host-path is from the correctness perspective not affected (so we do not introduce a regression). There is potential for a performance regression in certain scenarios. However, because this is a massive PR we think its better to get it merged and fix the caching issue in a subsequent PR (which might come as soon as next week). The follow-up PR would only touch rocm_ipc and would have no impact on the rest of the UCX code base.

With this current PR, we are constantly rebasing to pull in updates to the repo, and had to resolve a number of conflicts over the weeks, its one of the reasons we would see benefits in merging this PR and fix the problems in a follow-up PR.

I merged so you can progress.

rakhmets · 2026-08-07T16:23:08Z

 static ucs_status_t
-ucp_device_detect_local_sys_dev(ucp_context_h context,
-                                ucs_memory_type_t mem_type,
+ucp_device_detect_local_sys_dev(const ucp_context_h context,


Please remove this const. Since it's typedef struct ucp_context *ucp_context_h, it does not work as readers might expect.

rakhmets · 2026-08-07T16:32:33Z

+    if (worker == NULL) {
+        ucs_error("missing worker in local mem list params");
+        return UCS_ERR_INVALID_PARAM;
+    }


Please refactor this to avoid checking this parameter twice.

rakhmets · 2026-08-07T16:39:52Z

+        const ucs_memory_type_t mem_type, uct_allocated_memory_t *mem,
+        const ucs_sys_device_t local_sys_dev)


Please swap the last two parameters, so that all input parameters go before the output.

@rakhmets please review #11757, your comments are addressed there.

mshanthagit marked this pull request as draft March 27, 2026 15:54

mshanthagit force-pushed the pr/device-init-put branch 2 times, most recently from ee24de0 to 67b5b78 Compare March 30, 2026 18:28

mshanthagit force-pushed the pr/device-init-put branch from 67b5b78 to 0459bdd Compare April 8, 2026 06:09

mshanthagit marked this pull request as ready for review April 8, 2026 12:58

yosefe requested review from Artemy-Mellanox and ofirfarjun7 April 12, 2026 10:17

ofirfarjun7 reviewed Apr 15, 2026

View reviewed changes

Comment thread src/ucp/core/ucp_device.c Outdated

mshanthagit force-pushed the pr/device-init-put branch from 0459bdd to 3c778a2 Compare April 29, 2026 12:19

openucx deleted a comment from svc-nixl May 4, 2026

ofirfarjun7 reviewed May 5, 2026

View reviewed changes

Comment thread src/ucp/core/ucp_device.c

Comment thread src/ucp/core/ucp_device.c Outdated

ofirfarjun7 reviewed May 6, 2026

View reviewed changes

Comment thread src/ucp/core/ucp_device.c Outdated

Comment thread src/ucp/wireup/select.c Outdated

Comment thread src/uct/api/device/uct_device_impl.h Outdated

mshanthagit commented May 6, 2026

View reviewed changes

Comment thread src/uct/api/device/uct_device_impl.h Outdated

Comment thread src/ucp/wireup/select.c Outdated

Comment thread src/ucp/core/ucp_device.c Outdated

Artemy-Mellanox reviewed May 7, 2026

View reviewed changes

mshanthagit commented May 8, 2026

View reviewed changes

mshanthagit force-pushed the pr/device-init-put branch from 927735e to b8832e5 Compare May 18, 2026 02:24

ofirfarjun7 reviewed May 22, 2026

View reviewed changes

mshanthagit force-pushed the pr/device-init-put branch from bdbf299 to a5c4b47 Compare June 2, 2026 16:03

ofirfarjun7 reviewed Jun 7, 2026

View reviewed changes

Artemy-Mellanox reviewed Jun 24, 2026

View reviewed changes

ofirfarjun7 previously approved these changes Jun 24, 2026

View reviewed changes

Comment thread test/gtest/uct/rocm/test_rocm_ipc_device.hip

mshanthagit dismissed ofirfarjun7’s stale review via a4ae783 June 26, 2026 05:48

mshanthagit force-pushed the pr/device-init-put branch from bcb736f to a4ae783 Compare June 26, 2026 05:48

mshanthagit force-pushed the pr/device-init-put branch from a4ae783 to 7d70850 Compare July 20, 2026 18:33

ofirfarjun7 reviewed Jul 28, 2026

View reviewed changes

Comment thread src/uct/rocm/ipc/rocm_ipc.h

ofirfarjun7 reviewed Jul 28, 2026

View reviewed changes

Comment thread src/uct/rocm/ipc/rocm_ipc.h

Comment thread src/uct/api/device/uct_device_impl.h Outdated

Comment thread src/uct/rocm/ipc/rocm_ipc_cache.c Outdated

ofirfarjun7 reviewed Jul 28, 2026

View reviewed changes

ofirfarjun7 requested a review from rakhmets July 30, 2026 06:56

ofirfarjun7 reviewed Aug 4, 2026

View reviewed changes

edgargabriel mentioned this pull request Aug 5, 2026

RFC: Native same-host HIP IPC backend for ROCm VRAM ai-dynamo/nixl#2039

Open

mshanthagit force-pushed the pr/device-init-put branch from 4dddf15 to 57556bb Compare August 6, 2026 19:50

anjoj0 mentioned this pull request Aug 7, 2026

UCT/ROCM: Propagate asynchronous copy signal failures #11743

Open

ofirfarjun7 approved these changes Aug 7, 2026

View reviewed changes

ofirfarjun7 merged commit 0957aae into openucx:master Aug 7, 2026
161 checks passed

rakhmets reviewed Aug 7, 2026

View reviewed changes

This was referenced Aug 11, 2026

UCT/ROCM/IPC: Cache IPC handles per peer #11756

Open

UCP/CORE: Address device mem list review comments #11757

Open


		HIPCC ?= hipcc

		HIPCC_CMD = $(HIPCC) -DHAVE_CONFIG_H -DUCT_DEVICE_CODE_HIP -fPIE -I$(top_srcdir)/src -I$(top_builddir)/src $(BASE_CXXFLAGS) $(CXXFLAGS) $(HIP_CPPFLAGS) $(HIP_CXXFLAGS) $(HIPCC_EXTRA_FLAGS) -Wno-c++20-extensions -c $< -MT $@ -MF $(DEPDIR)/hip/$@.d -MMD -o $@


		#define UCS_DEVICE_LEVEL_EXEC_ID 1

		#define UCS_DEVICE_LEVEL_EXEC_SELECT(scope_ok, count, id) \

		const ucs_memory_type_t mem_type, uct_allocated_memory_t *mem,
		const ucs_sys_device_t local_sys_dev)

Uh oh!

Conversation

mshanthagit commented Mar 25, 2026 • edited by yosefe Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

What?

Why?

How?

Uh oh!

mshanthagit commented Apr 8, 2026

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

mshanthagit left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

mshanthagit left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

mshanthagit commented May 13, 2026

Uh oh!

edgargabriel commented May 15, 2026

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

mshanthagit commented May 24, 2026

Uh oh!

mshanthagit commented Jun 2, 2026

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

mshanthagit commented Mar 25, 2026 •

edited by yosefe

Loading

Artemy-Mellanox Jun 24, 2026 •

edited

Loading