[ddebonis-amd]

What?

The ROCm IPC cache now drops a remote mapping when the owning buffer is freed.

Why?

ROCm IPC mappings stay cached, because opening a handle is expensive. hipFree cannot finish while any importer still has that mapping attached, so the exporter publishes one IPC signal per buffer and importers detach as soon as that signal becomes 0. The signal is not destroyed until every importer that could be blocking the free already holds its own reference.

How?

Wire format
uct_rocm_ipc_key_t carries the existing hsa_amd_ipc_memory_t plus one hsa_amd_ipc_signal_t. The packed rkey size is the whole key, so this is a same-node wire-format change and not a public UCT or UCP API change. The signal handle is created once, at pack time, and is the same value for every peer. An empty signal is rejected on import.

Exporter
uct_rocm_ipc_pack_key() creates the memory handle, then uct_rocm_ipc_export_add() creates an HSA_AMD_SIGNAL_IPC signal with value 1 and exports it into the key. The local signal is owned by a process-global export table keyed by buffer address, not by the memory handle. mem_dereg only frees the key. The signal stays until the allocation is actually released.

The first export installs a UCM MEM_TYPE_FREE handler at priority 100, ahead of the registration cache at 1000. Installation is lazy, because UCM events may not be enabled at static init. On a ROCm or ROCm-managed free, before the original free runs, the handler stores 0 to each overlapping signal and moves those exports from the active list to the retired list. It does not destroy the signal in the callback. The callback does not allocate.

Retired exports are destroyed later by uct_rocm_ipc_export_progress(), called from uct_rocm_ipc_iface_progress(). A plain retired count skips the export lock when nothing is waiting. Any progress thread may reap, because importers attach the signal before they map memory. Process cleanup unregisters the UCM handler, destroys both lists, and clears the count. The count is written only under the export lock.

Importer cache
The cache remains process-global and is keyed by owner pid, pid namespace, and device. A page table in each cache is keyed by remote address. A hit whose IPC memory bytes match returns the existing local mapping and does not arm another handler.

On a miss, still under the remote-cache write lock:

1. Attach the IPC signal.
2. If its value is already 0 or less, destroy that local signal and do not map. The allocation is already being freed.
3. Attach the IPC memory.
4. Insert a region whose initial reference is the page-table reference, and return it to be watched after the lock is dropped.

hsa_signal_destroy and the async handler are not called while the cache lock is held. The lock is not recursive, and the ROCr handler thread takes it.

Unmap
Three references cover a region: the page table, the watcher that arms the handler, and the handler. deactivate claims unmap once. The claim winner detaches the mapping immediately, unlinks the page table, and queues the region. Purge, range invalidate, and a stale IPC byte string are all claims. Detach happens in the handler rather than on iface progress, because hipFree is blocked until the mapping is gone.

Signal ownership is separate. Before arming, the watcher marks the handler as the signal owner. A claim that wins before that destroys the signal itself. A claim that wins after that only detaches memory and leaves the signal for the callback, so ROCr is not left waiting on a destroyed signal and the handler reference is not leaked. If the region was already claimed before the handler existed, the watcher does not arm.

The callback is one-shot. It claims unmap if nobody else has, destroys the signal when it still owns it, and drops its region reference. Returning false disarms it. handler_done ignores the re-entry that hsa_signal_destroy can cause.

After a successful arm, the watcher loads the signal again. If it is already 0, it claims unmap immediately so the free does not wait for the ROCr thread. The handler still owns the signal and its reference.

Paths that do not watch
CACHE_IPC_HANDLES=n still creates the export signal, because the key is packed at registration time. The uncached zcopy path attaches memory for one copy and detaches when that copy completes. It does not watch the signal. Device mem_elem_pack and rkey_ptr always go through the cache, which is the same rule as before.

Copy completion still uses the existing per-iface signal pool in rocm_signal.c. Those signals are unrelated to the IPC lifetime signal.

What stays out of scope
No IPC policy was added to the UCM ROCm hooks. UCM only delivers the free event. The cache, the export table, and the handler all stay in src/uct/rocm/ipc.

[svc-nvidia-pr-review]

🤖 Starting review — findings will be posted here when done.

[svc-ucx]

🤖 CI Triage Agent — UCX PR (Codestyle commit title) · commit 5296538d

TL;DR: The "Codestyle / commit title" job failed because the PR's commit title Add removal of remote mapping when owning buffer is freed to drop IPC cache lacks the mandatory MODULE: prefix required by codestyle_check_commit_title; reword/amend the commit title (e.g. UCM/ROCM: Remove remote mapping when owning buffer is freed) and force-push.

Full analysis

Summary: Azure job Codestyle → commit title (build 138795) exited with code 1 at ./buildlib/tools/codestyle.sh::codestyle_check_commit_title, printing Bad commit title: 'Add removal of remote mapping when owning buffer is freed to drop IPC cache'.

Root cause: Not a product/infrastructure bug — a contribution-guideline violation. The check validates each non-merge commit title in base..head against the regex ^[0-9A-Z/_-]+:\ [[:alnum:]_], i.e. the title must begin with an uppercase module tag followed by ": " (e.g. UCM/ROCM: ...). Commit [REDACTED:Hex High Entropy String] on branch ucm-rocm-ipc-cache starts with Add removal of ..., so the regex does not match, err=1 is returned, and the Bash step exits 1. The log shows the failure 0.03 s after the script started, with no other activity — purely the title check, no timeout or hang involved.

Secondary (cosmetic CI bug, not the cause of this failure): in buildlib/pr/codestyle.yml the step uses set -eE, so the non-zero return from codestyle_check_commit_title aborts the step immediately and the if [[ $? -ne 0 ]] block that emits the helpful ##vso[task.logissue] message with the contributor-guidelines URL is never reached — which is why the log only shows the raw Bad commit title line and a generic Bash exited with code '1'.

Implicated commit: [REDACTED:Hex High Entropy String] (PR #12077, branch ucm-rocm-ipc-cache) — the offending commit title. The checker itself is unchanged (buildlib/tools/codestyle.sh, last touched by Yossi Itigin in cfd9cd1 "BUILD/CI: Limit PR commit range checks").

File: buildlib/tools/codestyle.sh:14 (regex) — triggered by the commit title of 5296538; CI step wiring at buildlib/pr/codestyle.yml:17-30

Suggested fix:

1. Amend the commit title to include a module prefix, then force-push:
   git commit --amend -m "UCM/ROCM: Remove remote mapping when owning buffer is freed" (keep it ≤ ~50 chars and in imperative mood; squash if multiple commits in the PR have non-conforming titles — the check iterates over every non-merge commit in the range).
2. Optional CI hardening: in buildlib/pr/codestyle.yml, change the step to run the check without aborting, e.g. set -eE plus codestyle_check_commit_title "$range" || { echo "##vso[task.logissue type=error]..."; echo "##vso[task.complete result=Failed;]"; }, so contributors get the guidelines link instead of a bare exit code.

Related: PR #12077; guidelines referenced by the job: https://github.com/openucx/ucx/wiki/Guidance-for-contributors#general-guidelines; commit-range change cfd9cd1 (BUILD/CI: Limit PR commit range checks, #11527)

🛡️ This comment had 1 potential secret(s) redacted (Hex High Entropy String). See request_id 6f4af379-ddc3-4d26-bf3a-bc1534814804 in the triage console for the audit trail.

[svc-ucx]

🤖 CI Triage Agent — UCX PR (Codestyle commit title) · commit 5296538d

TL;DR: The commit title codestyle job failed because commit 5296538 has the title Add removal of remote mapping when owning buffer is freed to drop IPC cache, which lacks the mandatory MODULE: prefix; amend/reword the commit to e.g. UCT/ROCM: Remove remote mapping when owning buffer is freed.

Full analysis

Summary: Azure Pipelines job "UCX PR (Codestyle / commit title)" build 138799 failed in the Bash step with Bad commit title: 'Add removal of remote mapping when owning buffer is freed to drop IPC cache' → Bash exited with code '1'.

Root cause: codestyle_check_commit_title() enforces the regex ^[0-9A-Z/_-]+:\ [[:alnum:]_] on every non-merge commit in base..head. The commit under test starts with the lowercase-mixed English word "Add" and has no COMPONENT: prefix, so it fails the pattern and the function returns 1. This is a genuine (non-flaky) codestyle violation — the PR title is correctly prefixed (UCT/ROCM: ...), but the individual commit message is not, and the check inspects commits, not the PR title.

Secondary observation: the pipeline step runs set -eE before calling codestyle_check_commit_title, so the non-zero return aborts the script immediately and the if [[ $? -ne 0 ]] branch at buildlib/pr/codestyle.yml:25-30 is never reached. That's why the log shows only a bare Bash exited with code '1' instead of the intended helpful "see wiki for more info" message.

Implicated commit: [REDACTED:Hex High Entropy String] (PR #12077, branch ucm-rocm-ipc-cache) — the commit with the malformed title. The checker itself is unchanged and working as designed.

File: buildlib/tools/codestyle.sh:14 (the regex that rejected it); failing invocation at buildlib/pr/codestyle.yml:24

Suggested fix: Reword the offending commit message to include a module prefix, then force-push:

git rebase -i [REDACTED:Hex High Entropy String]   # reword 5296538
# new title, e.g.:
# UCM/ROCM: Drop IPC cache entry when owning buffer is freed
git push --force-with-lease

Keep the prefix in uppercase/slash/underscore characters only ([0-9A-Z/_-]+), immediately followed by ": " and an alphanumeric character, and keep the subject under ~50-72 chars. Matching the PR title (UCT/ROCM:) or UCM/ROCM: (the branch name suggests ucm) both satisfy the regex.

Optional CI hardening: in buildlib/pr/codestyle.yml, change line 24 to codestyle_check_commit_title "$base_revision..$head_revision" || true (or drop -e around that call) so the intended error message with the contributor-guidelines URL is actually emitted instead of the opaque Bash exited with code '1'.

Related: PR #12077; check introduced/limited by cfd9cd1f "BUILD/CI: Limit PR commit range checks (#11527)" (Yossi Itigin); guidelines: https://github.com/openucx/ucx/wiki/Guidance-for-contributors#general-guidelines

🛡️ This comment had 1 potential secret(s) redacted (Hex High Entropy String). See request_id fd53ed50-7eb3-44f0-9628-a05de13fce0b in the triage console for the audit trail.

[svc-nvidia-pr-review]

blocker: this destroys the signal while the async handler armed in uct_rocm_ipc_cache_region_watch() is still registered on it (purge, invalidate, stale-region and the value <= 0 paths all get here), so HSA keeps waiting on a destroyed signal and the handler's refcount is never dropped, leaking the region. handler_done only prevents re-entry, it does not deregister the handler.

[svc-nvidia-pr-review]

this creates one IPC signal per mem_reg and the export is freed only on buffer free (mem_dereg deliberately keeps it), and export_add dedups nothing on the active list, so repeated map/unmap of the same buffer accumulates signals until the buffer is freed. IPC signals are a limited resource (this is exactly the HSA_STATUS_ERROR_OUT_OF_RESOURCES the new code returns) - can we reuse the active export for the same address instead?

[ddebonis-amd]

Implemented export reuse per allocation.

[svc-nvidia-pr-review]

reaping by tid (UCT_ROCM_IPC_EXPORT_REAP_TID) means only the freeing thread can destroy these, and only if it calls rocm_ipc progress; if the allocating/freeing thread is not a progress thread (or the app only uses rocm_copy), the signals stay until UCS_STATIC_CLEANUP. can we reap all retired exports once the free is known to be complete, instead of keying on tid?

[svc-nvidia-pr-review]

ucs_atomic_fadd32(&count, 0) is a bus-locked RMW executed on every uct_rocm_ipc_iface_progress() call; a plain load is enough for the fast exit (see docs/OptimizationStyle.md on avoiding atomics on the data path):

Suggested change

	if (ucs_likely(ucs_atomic_fadd32(&uct_rocm_ipc_export_retired_count, 0) ==
	void uct_rocm_ipc_export_progress(void)
	{
	if (ucs_likely(uct_rocm_ipc_export_retired_count == 0)) {
	return;
	}

[svc-nvidia-pr-review]

does this break rkey compatibility? the new ipc_signal field changes rkey_packed_size and the importer now hard-fails (ucs_error + UCS_ERR_INVALID_PARAM) when the field is zero - worth confirming the intended compatibility story.

[svc-nvidia-pr-review]

minor: drain loop of atomic decrements under the lock where a plain assignment is enough - we already hold the lock here:

Suggested change

	while (ucs_atomic_fadd32(&uct_rocm_ipc_export_retired_count, 0) != 0) {
	uct_rocm_ipc_export_retired_count = 0;

[svc-nvidia-pr-review]

PR-level: pls split the unrelated reformatting in rocm_ipc_cache.c (pointer-cast restyling, comment rewording, log-line rewraps) into a separate PR - the functional change is hard to review inside a ~500 line diff.

PR-level: there is no test for the new export/attach/free flow (test/gtest/uct/rocm/test_rocm_ipc_device.hip does not cover it). What covers the new flow in CI? pls add a test where the exporter frees the buffer while a peer still has it mapped.

[svc-nvidia-pr-review]

🤖 Starting review — findings will be posted here when done.

[svc-nvidia-pr-review]

Each mem_reg of the same allocation adds another export entry plus HSA IPC signal, and mem_dereg never removes it, so a register/deregister loop (or several sub-ranges of one allocation) leaks signals and list entries until the buffer is freed. Please key the export on the allocation base address and refcount it instead of adding one per pack_key.

[svc-nvidia-pr-review]

This grows the packed rkey by sizeof(hsa_amd_ipc_signal_t) (32B), and docs/OptimizationStyle.md says to avoid enlarging the remote memory key at all costs. Can we exchange the signal once per (pid, pid_ns, dev_num) - the same granularity as the remote cache - instead of per rkey?

[svc-nvidia-pr-review]

In uct_rocm_ipc_export_progress: this puts a bus-locked xadd on every progress call, and as soon as any export is retired by another thread, every progress call also takes the global spinlock and walks the whole retired list. Please keep the pending count per thread so progress stays lock-free when this thread retired nothing (docs/OptimizationStyle.md: no locks/atomics in the data path).

[svc-nvidia-pr-review]

In uct_rocm_ipc_cache_region_finalize: please confirm hsa_signal_destroy is safe while an async handler is still armed on that signal. In the purge / UCS_STATIC_CLEANUP path the exporter usually has not freed yet, so the handler is still registered and holds a region reference; if it fires later it dereferences region->cache, already freed by uct_rocm_ipc_destroy_cache, and the already destroyed uct_rocm_ipc_remote_cache.lock. There is no API to cancel hsa_amd_signal_async_handler.

[svc-nvidia-pr-review]

In uct_rocm_ipc_cache_region_watch: does this break importing an rkey packed by a peer running an older UCX? ipc_signal is all-zeros there and the whole mapping now fails instead of behaving as before.

[svc-nvidia-pr-review]

Suggestion:

    uct_rocm_ipc_export_retired_count = 0;

All writers of the counter hold the lock, so a plain store is enough here.

[svc-nvidia-pr-review]

Minor: please drop (void)arg; (and the same at line 393) - UCX callbacks just leave unused parameters unused.

[svc-nvidia-pr-review]

Residual coverage gap beyond what was already raised: nothing exercises the async detach path (uct_rocm_ipc_cache_signal_cb) or the teardown path with an armed handler, which is where the hsa_signal_destroy/teardown and older-peer rkey findings above would show up, so both need hardware testing rather than CI inspection.

[svc-nvidia-pr-review]

🤖 Starting review — findings will be posted here when done.

[svc-nvidia-pr-review]

pls release the export entry in mem_dereg, or refcount entries per base address: pack_key adds a new entry plus a new HSA IPC signal on every registration (including repeated registrations of sub-ranges of the same base allocation, since key->address is the allocation base), and nothing removes it until the buffer is freed, so the active list and the number of IPC signals grow without bound. uct_rocm_ipc_export_mem_free_cb also walks this whole list under a global spinlock on every ROCm free.

[svc-nvidia-pr-review]

detach now runs from the HSA async handler thread, while other threads may still be using a mapped_addr that uct_rocm_ipc_cache_map_memhandle returned earlier - the cache-hit path hands out region->mapped_addr with no reference, and uct_rocm_ipc_md_mem_elem_pack bakes the offset into device memory. how is an in-flight hsa_amd_memory_async_copy protected from the VA being unmapped under it?

[svc-nvidia-pr-review]

ucs_atomic_fadd32(..., 0) is a locked read-modify-write, so every progress call dirties this shared cache line even when there is nothing to reap. pls declare the counter volatile and read it directly here - progress is a per-poll path (docs/OptimizationStyle.md: avoid bus-locked instructions on the data path).

[svc-nvidia-pr-review]

a retired export is reaped only by the thread that freed it, so if that thread never calls progress again the counter stays non-zero and every progress call on every other thread walks the retired list under the spinlock, and the signal is only freed at static cleanup. can we reap by address or by a free-generation counter instead of tid?

[svc-nvidia-pr-review]

minor: uct_rocm_ipc_export_in_free is always 0 here - reap is reachable only from uct_rocm_ipc_export_add and uct_rocm_ipc_export_progress, which the freeing thread cannot call while it is inside the free callback. pls remove the flag and the comment.

[svc-nvidia-pr-review]

Duplicate check: I did not re-raise the already-reported items (unrelated reformatting in rocm_ipc_cache.c, missing test for the export/attach/free flow, hsa_signal_destroy with an armed handler at teardown, older-peer rkey compatibility from the new ipc_signal field in uct_rocm_ipc_key_t). I was not able to fetch the PR conversation from this environment, so before posting, the full paginated comment/review history must be fetched and these five candidates re-checked against it; if any already has an unresolved thread, add a +1 reaction instead.

Residual coverage gap beyond what was already raised: nothing exercises repeated register/deregister of the same allocation, which is where the export-table growth above becomes visible; this needs a ROCm-hardware test rather than CI inspection.

[svc-nvidia-pr-review]

🤖 Starting review — findings will be posted here when done.

[svc-nvidia-pr-review]

Retired exports are only reaped by the thread that freed the memory: uct_rocm_ipc_export_progress() passes UCT_ROCM_IPC_EXPORT_REAP_TID with the calling thread's tid, and reap() matches exp->retired_tid == tid. Exports retired by a thread that never calls progress (app frees on the main thread while progress runs on a worker thread, or the freeing thread exits) are never reaped, so their HSA signals and export entries live until UCS_STATIC_CLEANUP and uct_rocm_ipc_export_retired_count stays non-zero forever — which makes every uct_rocm_ipc_iface_progress() call take the global export spinlock and walk the whole retired list. Can we reap independently of the tid?

[svc-nvidia-pr-review]

Region reference leaked when the claim was lost after arming the handler: uct_rocm_ipc_cache_region_watch() takes a region reference and arms uct_rocm_ipc_cache_signal_cb on local_signal. If region->deactivate was already set by purge/invalidate_regions/stale-drop, the handler is still armed, but this path destroys local_signal and only drops the watcher reference in out_unref. So uct_rocm_ipc_cache_signal_cb never runs, the handler's reference is never dropped and the region is leaked (and the ROCr handler thread is left waiting on a destroyed signal).

[svc-nvidia-pr-review]

Minor: uct_rocm_ipc_export_retired_count is only modified under uct_rocm_ipc_export_lock, so the atomic RMWs are not needed, and this ucs_atomic_fadd32(ptr, 0) is a bus-locked instruction used only to read the value on the progress path (docs/OptimizationStyle.md: avoid bus-locked instructions on the data path). The while loop in uct_rocm_ipc_export_cleanup (L283-L285) can then be a plain assignment of 0.

[svc-nvidia-pr-review]

uct_rocm_ipc_export_in_free looks unreachable: uct_rocm_ipc_export_reap() is only called from uct_rocm_ipc_export_add() and uct_rocm_ipc_export_progress(), and neither is reachable from the UCM free callback — when is this guard hit? If never, please drop uct_rocm_ipc_export_in_free (REVIEW.md: avoid redundant state).

[svc-ucx]

🤖 CI Triage Agent — UCX PR (AddressSanitizer gpu on worker 3) · commit 4425339b

TL;DR: The ASAN GPU gtest job failed on a single protocol‑selection test, shm_ib/test_ucp_proto_am_rndv.am_explicit_scheme_disables_force_provenance/0, which asserts that am/rndv is selected for 1 MB CUDA buffers while UCX_RNDV_SCHEME=get_zcopy is forced — a condition that only holds if a network MD can register CUDA memory; the test guards only on generic RNDV support, so it must be given a stronger skip guard (require_cuda_net_md()), and it is unrelated to this PR's ROCm IPC change.

Full analysis

Summary: make test in the gtest stage exited with Error 1 because 1 of 10336 tests failed: shm_ib/test_ucp_proto_am_rndv.am_explicit_scheme_disables_force_provenance/0, GetParam() = shm,ib,cuda_copy,rocm_copy (all other tests passed; no ASAN report, no crash, no hang — timestamps are continuous to the end at 00:47:54).

Root cause: The test is a pure protocol‑selection assertion, not a data‑path test. It runs with RNDV_PIPELINE_SHM_CUDA_STAGING_FORCE=y + RNDV_SCHEME=get_zcopy and calls check_am_rndv_remote_proto_config(0, 0), which in select_am_rndv_remote_proto_config() hard‑asserts EXPECT_STREQ("am/rndv", thresh->proto_config.proto->name) for a 1 MB UCS_MEMORY_TYPE_CUDA AM send (test/gtest/ucp/test_ucp_proto.cc:438). Because RNDV_SCHEME=get_zcopy sets rndv_mode != UCP_RNDV_MODE_AUTO, ucp_proto_rndv_cfg_thresh() reports UCS_MEMUNITS_INF for every rendezvous variant except get/zcopy (src/ucp/rndv/proto_rndv.inl:131‑150). Selection of am/rndv therefore depends on a get_zcopy‑capable network MD that can register CUDA memory being present on the worker. The test's only guard is require_cuda_rndv_support() (which just checks mem_buffer::is_mem_type_supported(CUDA) and is_rndv_supported()), whereas the sibling tests in the same file that depend on CUDA‑registering NICs correctly guard with require_cuda_net_md() (test/gtest/ucp/test_ucp_proto.cc:98‑112, 632, 693). On the ASAN CUDA11 container (where several memory/registration tests are explicitly skipped with [ SKIP ] (Address sanitizer), e.g. rc/test_ucp_mmap_export.*), that precondition is not guaranteed, so the forced get_zcopy variant is unavailable and a non‑am/rndv protocol (or none) is selected, tripping the expectation. The identical weak guard exists in test_ucp_proto_rma_rndv.rma_explicit_scheme_disables_force_provenance (line 777‑786), so the same flake can appear there. Note the whole‑build log is truncated to the last ~500 lines, so the individual EXPECT_* failure line scrolled out; the surviving evidence is the summary line plus the test body/env combination.

Implicated commit: Not this PR. The provenance/estimation test family and the UCP_RKEY_CONFIG_FLAG_PROTO_ESTIMATION plumbing came in with 0311ed79 (jeffnvidia, "UCP/PROTO/RNDV: Isolate remote system device namespace", #11685) and the force‑staging work 9cc08081 / b34ac9a1 (Nathan Bellalou, "UCP/RNDV: use tag-only signal for CUDA staging force" / "Add opt-in intra-node CUDA staging policy"). PR #12077 (ucm-rocm-ipc-cache, commit 4425339) touches UCT/ROCm IPC mapping teardown and cannot influence a shm,ib + CUDA AM protocol‑selection decision.

File: test/gtest/ucp/test_ucp_proto.cc:805‑813 (test body) and test/gtest/ucp/test_ucp_proto.cc:438 (EXPECT_STREQ("am/rndv", ...)); related gating logic in src/ucp/rndv/proto_rndv.inl:131‑150.

Suggested fix: Strengthen the precondition so the test only runs where the forced scheme can actually be selected:

1. In am_explicit_scheme_disables_force_provenance (and rma_explicit_scheme_disables_force_provenance), replace require_cuda_rndv_support() with require_cuda_net_md(), matching the other RNDV_SCHEME-sensitive tests in the file.
2. Additionally make select_am_rndv_remote_proto_config() tolerant: if the selected protocol name is not am/rndv (or the info string contains "stub protocol"), UCS_TEST_SKIP_R("CUDA rendezvous get_zcopy is not supported") instead of ADD_FAILURE() — the same pattern already used in rndv_force_cuda_send_child_uses_tag_provenance (line 700‑702).
3. Re-trigger CI on PR UCT/ROCM: Add removal of remote mapping when owning buffer is freed to drop IPC #12077; this failure should not block it. If the goal is to keep hard assertions, gate the test on the RNDV scheme actually being satisfiable rather than on generic CUDA RNDV support.

Related: PR #12077 (the PR under test, unrelated to the failure); prior diagnosis of this exact test name in Azure protocol selection: #11858, #11861, #11864, and #11685 / #11990.

[svc-ucx]

🤖 CI Triage Agent — UCX PR (Build build_x86_64 ubuntu2204_rocm) · commit 8aa2e5e6

TL;DR: The ROCm build fails to compile src/uct/rocm/ipc/rocm_ipc_cache.c because a bad merge (8aa2e5e) left a duplicated "remove stale region" block with an unbalanced brace in uct_rocm_ipc_cache_map_region() and an undeclared status in uct_rocm_ipc_cache_invalidate_regions(). Fix the merge conflict leftovers: add the missing }, delete the duplicated detach/trace lines, and declare ucs_status_t status;.

Full analysis

Summary: Build on x86_64 stage of build_x86_64 ubuntu2204_rocm failed at make[3]: *** [Makefile:828: ipc/libuct_rocm_la-rocm_ipc_cache.lo] Error 1 — a cascade of -Werror errors in rocm_ipc_cache.c.

Root cause: Two distinct source defects introduced by the merge commit on this branch:

1. uct_rocm_ipc_cache_invalidate_regions() uses status without declaring it:
   rocm_ipc_cache.c:454:9: error: 'status' undeclared (first use in this function) — the function body (lines 444–447) declares only region_list, region, tmp.

2. uct_rocm_ipc_cache_map_region() is missing a closing brace. Counting braces from line 492: if (ucs_likely(pgt_region != NULL)) { (depth+1) → if (memcmp(...) == 0) { (depth+1) → } else { at 502 → the else body ends with } at line 527, which closes only the else. The outer if (pgt_region != NULL) block is never closed, so the whole rest of the file is parsed inside uct_rocm_ipc_cache_map_region. That produces every remaining error in the log: invalid storage class for function 'uct_rocm_ipc_remote_cache_get', ... 'uct_rocm_ipc_cache_region_watch', ucs_initializer_ctor0/dtor1 invalid storage class, the long list of -Werror=shadow hits (cache, key, status, region, deferred "shadows a parameter/previous local" of map_region), ISO C90 forbids mixed declarations and code, and finally rocm_ipc_cache.c:836:1: error: expected declaration or statement at end of input.

   The unbalanced brace is clearly conflict-resolution debris: lines 503–517 (the old ucs_pgtable_remove + hsa_amd_ipc_memory_detach path, indented 12 spaces) and lines 519–526 (the new uct_rocm_ipc_cache_region_drop_locked path, indented 8 spaces) are two copies of the same "remove stale region" logic, including a duplicated identical ucs_trace("%s: rocm_ipc cache remove stale region:" ...). Beyond the compile break, keeping both paths would detach the IPC mapping twice for one region.

Implicated commit: [REDACTED:Hex High Entropy String] (David DeBonis, "Merge branch 'master' into ucm-rocm-ipc-cache"), resolving the merge of 3e322414 ("UCT/ROCM: Add removal of remote mapping when owning buffer is freed to drop IPC") against 1cf9701f ("UCT/ROCM_IPC: Add peer failure error handling").

File: src/uct/rocm/ipc/rocm_ipc_cache.c:454 and src/uct/rocm/ipc/rocm_ipc_cache.c:502-527

Suggested fix:

1. In uct_rocm_ipc_cache_invalidate_regions() (around line 444), add the missing declaration:

   ucs_status_t status;

2. Re-resolve the stale-region branch in uct_rocm_ipc_cache_map_region(): keep only the new uct_rocm_ipc_cache_region_drop_locked(region, deferred) path (the branch's intended behavior), delete the duplicated lines 503–518 (ucs_trace + ucs_pgtable_remove + hsa_amd_ipc_memory_detach) so the region isn't detached twice, and add the } that closes the if (ucs_likely(pgt_region != NULL)) { block after the else body. Resulting shape:

   if (ucs_likely(pgt_region != NULL)) {
       region = ucs_derived_of(pgt_region, uct_rocm_ipc_cache_region_t);
       if (memcmp(&key->ipc, &region->key.ipc, sizeof(key->ipc)) == 0) {
           ...
           return UCS_OK;
       } else {
           ucs_trace("%s: rocm_ipc cache remove stale region:" ...);
           if (!uct_rocm_ipc_cache_region_drop_locked(region, deferred)) {
               status = UCS_ERR_IO_ERROR;
               goto err;
           }
       }
   }   /* <-- missing brace */

3. Build locally with the same flags before re-pushing — this job uses -Wall -Werror -Wshadow -Werror=declaration-after-statement, so shadowed locals and mixed declarations are hard failures; a plain non-ROCm build will not catch this since src/uct/rocm is only compiled when --with-rocm is configured.

Related: PR #12077 (branch ucm-rocm-ipc-cache); prior commits on this file: 3e322414, 1cf9701f, 349e171d.

🛡️ This comment had 1 potential secret(s) redacted (Hex High Entropy String). See request_id 25527d88-cc63-4c64-9b56-789c251ecfee in the triage console for the audit trail.

[svc-nvidia-pr-review]

🤖 Starting review — findings will be posted here when done.

[svc-nvidia-pr-review]

The async handler detaches the mapping as soon as the free event fires — what prevents a detach while a put/get zcopy issued on this mapping is still in flight? region->refcount tracks only pgtable/handler/watcher, not outstanding operations (cuda_ipc refcounts per map/unmap for this reason).

[svc-nvidia-pr-review]

Suggested change

	ucs_fatal("failed to unmap addr:%p", mapped_addr);
	ucs_warn("failed to unmap addr:%p", mapped_addr);

A failed detach in a cleanup path should not abort the process; the current code uses ucs_warn. Same for the two unwind paths in uct_rocm_ipc_cache_map_region.

[svc-nvidia-pr-review]

deactivate, handler_done, signal_owner and in_pgtable are only touched while the remote-cache write lock is held, so can we make them plain fields and drop the atomics? ucs_atomic_cswap32(&region->deactivate, 0, 0) as a plain read is hard to follow.

[svc-nvidia-pr-review]

uct_rocm_ipc_export_retired_count is written under uct_rocm_ipc_export_lock but read here without it; pls make it volatile (or read it under the lock), otherwise tsan flags the race.

[svc-nvidia-pr-review]

Suggested change

	uct_rocm_ipc_ipc_signal_is_empty(const hsa_amd_ipc_signal_t *ipc_signal)
	uct_rocm_ipc_signal_is_empty(const hsa_amd_ipc_signal_t *ipc_signal)

minor: ipc is duplicated in the name; pls update the call site too.

[svc-nvidia-pr-review]

minor: pls drop (void)arg; here and (void)value; in uct_rocm_ipc_cache_signal_cb; UCX leaves unused callback parameters as-is.

[svc-nvidia-pr-review]

Suggested change

	static ucs_list_link_t uct_rocm_ipc_export_active;
	static ucs_list_link_t uct_rocm_ipc_export_active;
	static ucs_list_link_t uct_rocm_ipc_export_retired;
	static ucs_spinlock_t uct_rocm_ipc_export_lock;
	static uint32_t uct_rocm_ipc_export_retired_count;

minor: pls align the declarations on one column (docs/CodeStyle.md).

[svc-nvidia-pr-review]

Residual coverage gaps: beyond what was already raised, nothing exercises a free event arriving while a zcopy is in flight on the imported mapping, and nothing exercises UCX_ROCM_IPC_CACHE_IPC_HANDLES=n, where the direct hsa_amd_ipc_memory_attach path in uct_rocm_ipc_ep_zcopy never attaches the signal and therefore gets no protection from this change. Both need ROCm hardware rather than CI inspection.

[svc-nvidia-pr-review]

🤖 Starting review — findings will be posted here when done.

[svc-nvidia-pr-review]

pls keep this as ucs_warn - the detach failure escalated from ucs_warn to ucs_fatal. drop_locked runs from uct_rocm_ipc_cache_purge at static cleanup (via uct_rocm_ipc_destroy_cache), so a detach failure (e.g. the exporting peer already exited, which is a normal IPC scenario) would now abort every importer at exit instead of warning. docs/CodeStyle.md says a destructor should handle an error by printing a warning or error message.

Suggested change

	ucs_fatal("failed to unmap addr:%p", mapped_addr);
	if ((mapped_addr != NULL) &&
	(hsa_amd_ipc_memory_detach(mapped_addr) != HSA_STATUS_SUCCESS)) {
	ucs_warn("failed to unmap addr:%p", mapped_addr);
	}

Same escalation applies to the two new ucs_fatal detach calls in uct_rocm_ipc_cache_map_region (around lines 542 and 579).

[svc-nvidia-pr-review]

pls keep hsa_status in the message - the rewrapped log line lost the hsa_status:%d argument while hsa_status is still available, so the actual HSA failure cause is no longer printed. docs/LoggingStyle.md requires printing the exact cause of the error.

Suggested change

	ucs_error("%s: failed to open ipc mem handle. addr:%p len:%lu",
	ucs_error("%s: failed to open ipc mem handle. addr:%p len:%lu "
	"hsa_status:%d",
	cache->name, (void*)key->address, key->length, hsa_status);

[svc-nvidia-pr-review]

is handler_done needed? hsa_amd_signal_async_handler calls the handler once per registration and uct_rocm_ipc_cache_signal_cb returns false, so it cannot re-enter. The region already carries deactivate, signal_owner, in_pgtable, and refcount; REVIEW.md asks to avoid redundant state.

[svc-nvidia-pr-review]

minor: inconsistent alignment in the new globals block - pls align these on the column.

Suggested change

	static ucs_list_link_t uct_rocm_ipc_export_active;
	static ucs_list_link_t uct_rocm_ipc_export_active;
	static ucs_list_link_t uct_rocm_ipc_export_retired;
	static ucs_spinlock_t uct_rocm_ipc_export_lock;
	static uint32_t uct_rocm_ipc_export_retired_count;
	static ucs_init_once_t uct_rocm_ipc_export_handler_once =
	UCS_INIT_ONCE_INITIALIZER;

[svc-nvidia-pr-review]

Residual coverage gap: nothing exercises a detach failure in the teardown path, which is exactly where the new ucs_fatal turns a previously survivable warning into an abort. Validating this needs a dead-exporter scenario on ROCm hardware rather than CI inspection.

[svc-nvidia-pr-review]

🤖 Starting review — findings will be posted here when done.

[svc-nvidia-pr-review]

src/uct/rocm/ipc/rocm_ipc_cache.c

every imported region now holds an attached IPC signal plus an armed HSA async handler until the exporter frees the buffer, and this cache never evicts regions (there is no CACHE_MAX_REGIONS/CACHE_MAX_SIZE like cuda_ipc) - with many peers and many buffers, is the number of concurrently attached signals and armed handlers bounded? today a shortage shows up as a hard import failure in drop_mapping.

src/uct/rocm/ipc/rocm_ipc_cache.c

uct_rocm_ipc_cache_map_region: this adds an HSA signal read to every put/get zcopy cache hit, under the global remote-cache write lock, so it serializes all peers. since the armed handler already unlinks an invalidated region, do we need the load on every operation?

src/uct/rocm/ipc/rocm_ipc_cache.c

uct_rocm_ipc_cache_map_region: minor: the two hit branches are identical, pls merge them and drop the now unused signal_value. Suggested:

            if ((region->signal.handle == 0) ||
                (hsa_signal_load_scacquire(region->signal) > 0)) {
                ucs_trace("%s: rocm_ipc cache hit addr:%p size:%lu region:"
                          UCS_PGT_REGION_FMT,
                          cache->name, (void*)key->address, key->length,
                          UCS_PGT_REGION_ARG(&region->super));

                *region_p = region;
                return UCS_OK;
            }

Residual coverage gaps beyond what was already raised: nothing exercises a large number of simultaneously imported regions, which is where the per-region signal/handler accumulation becomes visible, and nothing measures the added per-operation signal read on the zcopy path. Both need ROCm hardware with two processes rather than CI inspection.

[svc-nvidia-pr-review]

🤖 Starting review — findings will be posted here when done.

[svc-nvidia-pr-review]

src/uct/rocm/ipc/rocm_ipc_cache.c

In uct_rocm_ipc_export_add, when the knob is n we still create and export an HSA IPC signal for every allocation, but no free handler is installed, so nothing ever sets it to 0 and it lives until static cleanup. Can we return early here and leave key->ipc_signal zeroed instead? The packed layout is the same either way, and today a peer that has the knob enabled attaches that signal and arms an async handler that can never fire.

Residual coverage gaps: nothing exercises an exporter with UCX_ROCM_IPC_SIGNAL_INVALIDATE=n talking to an importer with the default y (the asymmetric case), and nothing exercises ROCm managed memory, where ucm dispatches the free event as UCS_MEMORY_TYPE_ROCM_MANAGED and the new callback filters it out. Both need two ROCm processes on real hardware rather than CI inspection.

[svc-nvidia-pr-review]

🤖 Starting review — findings will be posted here when done.

[svc-nvidia-pr-review]

minor: super is at offset 0, so &unpacked->super is literally rkey and this path never touches the new fields - pls keep uct_rocm_ipc_key_t *key = (uct_rocm_ipc_key_t*)rkey; here and in get_zcopy to avoid growing an already large diff with a no-op cast round-trip.

[svc-nvidia-pr-review]

minor: in uct_rocm_ipc_cache_config_table, the second line describes the packed-rkey layout rather than what the user gets, and it reads as if the signal is still exported when disabled (it is zeroed by export_add). Suggested:

Suggested change

	"The IPC signal remains part of the packed rkey when this is n.",
	"Detach cached ROCm IPC mappings when the exported buffer is freed.\n"
	"When set to n, a mapping is replaced only when the peer's IPC handle\n"
	"changes. The packed rkey layout is the same in both cases.",

[svc-nvidia-pr-review]

src/uct/rocm/ipc/rocm_ipc_cache.c

In uct_rocm_ipc_export_add: pls install this handler at md_open instead of lazily from the first key pack: ucm_set_event_handler(UCM_EVENT_MEM_TYPE_FREE, ...) runs ucm_event_install(), which can bistro-patch the HSA entry points and load the ucm modules. Everywhere else in UCX that happens at init time (memtype cache, rcache), not from a registration call on an already running, multi-threaded app.

Residual coverage gaps beyond what was already raised: nothing exercises a first uct_md_mem_reg that triggers ucm rocm hook installation while another thread is allocating/freeing ROCm memory (the window opened by finding 1), and nothing calls ucp_rkey_ptr on a rocm_ipc rkey at all — the gtest test_md.rkey_ptr is skipped because the md does not report UCT_MD_FLAG_RKEY_PTR, so the new pin-on-first-call / unpin-on-release bookkeeping in uct_rocm_ipc_rkey_ptr is untested. Both need ROCm hardware or a targeted UCT test rather than CI inspection.

[svc-nvidia-pr-review]

🤖 Starting review — findings will be posted here when done.

[svc-nvidia-pr-review]

src/uct/rocm/ipc/rocm_ipc_cache.c

Config table help text: the first help line is duplicated, so it shows up twice in ucx_info -f. Suggested text:

{"SIGNAL_INVALIDATE", "y",
 "Detach cached ROCm IPC mappings when the exported buffer is freed.\n"
 "When set to n, a mapping is replaced only when the peer's IPC handle\n"
 "changes. The packed rkey layout is the same in both cases.",

src/uct/rocm/ipc/rocm_ipc_cache.c

uct_rocm_ipc_export_reap() is only called from export_add, so a retired export keeps its HSA signal alive until the next registration. An app that registers N buffers, frees them all and then stops registering holds N signals plus N uct_rocm_ipc_export_t until process exit — and signal exhaustion is a hard import failure elsewhere in this file. Can we also reap from the free callback path, or bound the retired list?

src/uct/rocm/ipc/rocm_ipc_md.c

uct_rocm_ipc_pack_key(): minor — uct_rocm_ipc_export_add already logs the HSA failure, and per docs/CodeStyle.md the caller does not have to log again when the callee returned an erroneous ucs_status_t. This also fires on every mem_reg once signals run out, so please use ucs_diag here instead of ucs_warn.

src/uct/rocm/ipc/rocm_ipc_cache.c

minor: uct_rocm_ipc_cache_map_memhandle() and uct_rocm_ipc_cache_pin_memhandle() are identical one-line wrappers around uct_rocm_ipc_cache_map. Can we keep a single exported uct_rocm_ipc_cache_map_memhandle(key, mapped_addr, region_p) where region_p == NULL means "do not pin", and have uct_rocm_ipc_ep_zcopy pass NULL?

Residual coverage gaps: nothing exercises the exporter retired-export list in a free-heavy / registration-idle pattern, which is where the signal accumulation becomes visible (needs ROCm hardware with two processes rather than CI inspection); and nothing checks the new UCX_ROCM_IPC_SIGNAL_INVALIDATE help text or the knob itself through ucx_info -f in CI.

[svc-nvidia-pr-review]

🤖 Starting review — findings will be posted here when done.

[svc-nvidia-pr-review]

🤖 Review complete — no issues found.

[edgargabriel]

I verified the logic both by reading the code and running tests, and it looks correct for both potential scenarios (using the ipc_signals to trigger a release of the ipc mapping as well as disabling the feature).

Except for the one comment about the location of the new parameter, this is PR is good as far as I can tell.

[svc-nvidia-pr-review]

🤖 Starting review — findings will be posted here when done.

[svc-nvidia-pr-review]

src/uct/rocm/ipc/rocm_ipc_md.c

In uct_rocm_ipc_rkey_ptr, in the if (unpacked->region == NULL) block: after the owner frees the buffer, the cache now fails new mappings with UCS_ERR_INVALID_ADDR, but here the pin taken on the first call is cached for the lifetime of the rkey, so every later call returns the stale address with UCS_OK. Is that intended, or can we report the invalidation to the caller instead of handing out a mapping of released memory?

src/uct/rocm/ipc/rocm_ipc_cache.c

In uct_rocm_ipc_cache_signal_cb: minor — claimed only guards uct_rocm_ipc_cache_region_finalize(&deferred), which is already a no-op when the list is empty, so please drop the variable.

Wire compatibility of the new ipc_signal field: it is inserted in the middle of uct_rocm_ipc_key_t (between ipc and address), so a mixed-version pair misreads address/length/dev_num/pid too, not just the new field. Appending at the end plus a flag-based check (as cuda_ipc does with UCT_CUDA_IPC_RKEY_FLAG_PID_NS and a size-dependent rkey_packed_size) would keep old peers working. Worse, garbage in the signal-field position makes hsa_amd_ipc_signal_attach return HSA_STATUS_ERROR_INVALID_ARGUMENT, which map_region reports as "already released by the owner" at debug level only — every transfer then fails with no diagnostic above debug. (Existing thread — reply/+1 there rather than a new comment.)

uct_rocm_ipc_export_cleanup: besides hoisting the config read out of the loop, the uct_rocm_ipc_signal_invalidate_enabled() check there is unreachable-false — uct_rocm_ipc_export_add returns early when the knob is off, so the table is always empty in that case. The condition can simply be removed. (Existing thread — reply/+1 there rather than a new comment.)

In-flight zcopy / missing reference on the uct_rocm_ipc_ep_zcopy path: uct_rocm_ipc_cache_map_region additionally calls uct_rocm_ipc_signal_invalidate_enabled() (a UCS_INIT_ONCE pthread mutex, and a full config-table parse on the first call) while holding the global remote-cache write lock, which is the same lock the HSA async handler needs. (Existing thread — reply/+1 there rather than a new comment.)

Residual coverage gaps: nothing exercises the pre-arm window (free event arriving between handler_armed = 1 and hsa_amd_signal_async_handler, where the signal is deliberately left for a handler that may not exist yet), and nothing exercises a pinned region being dropped while pin_count > 0 so that the detach is deferred to uct_rocm_ipc_cache_unpin. Both need two ROCm processes on real hardware rather than CI inspection. Note the existing test_rocm_ipc_rma.get_mem_elem_pack/device_put gtests do run exporter and importer in one process, so the free-event → async-detach path is partially exercised there if ROCm hardware is present in CI.

[svc-nvidia-pr-review]

🤖 Starting review — findings will be posted here when done.

[svc-nvidia-pr-review]

src/uct/rocm/ipc/rocm_ipc_cache.c

In uct_rocm_ipc_cache_signal_cb, on the uct_rocm_ipc_cache_region_signal_destroy(region, signal); line: here we destroy the signal from inside the async handler that is armed on that same signal — is that guaranteed safe for all supported ROCm versions? If the runtime still uses the handle after the callback returns, this is a use-after-free in the HSA thread; if it is safe, pls add a short comment saying so, since this is the first hsa_amd_signal_async_handler user in UCX.

[svc-ucx]

🤖 CI Triage Agent — UCX PR (Tests new on worker 3) · commit 63a1ccd9

TL;DR: The build failed on a single gtest, shm_ib/test_ucp_proto_rndv_force_cuda.rndv_force_cuda_send_child_uses_tag_provenance/0 (TLs shm,ib,cuda_copy,rocm_copy), a pure protocol-selection string-matching test that is unrelated to this ROCm-IPC PR and is known to fail on other PRs too — re-run the job and harden/skip the test instead of chasing the PR diff.

Full analysis

Summary: ./contrib/test_jenkins.sh → make test exited 2 because 1 of 10663 gtests failed: shm_ib/test_ucp_proto_rndv_force_cuda.rndv_force_cuda_send_child_uses_tag_provenance/0, where GetParam() = shm,ib,cuda_copy,rocm_copy (all other tests passed, no crash, no timeout — the run completed normally in 5419041 ms).

Root cause: The test forces RNDV_PIPELINE_SHM_CUDA_STAGING_FORCE=y, selects the RNDV_SEND protocol for CUDA(local)→host(remote), and then asserts on the human-readable protocol description, requiring it to contain "rndv_send(tag-rndv)", "pipeline cuda_copy" and "frag host" (test/gtest/ucp/test_ucp_proto.cc:728-730). This instantiation runs with shm,ib + UCP_TEST_GPU_COPY_TLS only — i.e. no gdr_copy and no cuda_ipc (ucp_test.h:435-437/462, instantiated at test_ucp_proto.cc:858) — so the winning protocol is decided purely by cost/perf estimation on that host. Its only escape hatch is a skip when the info string says "stub protocol" (line 724); if protocol estimation picks any other legitimate non-pipeline protocol, the EXPECTs fail. Recent protocol-selection/cost changes on master touch exactly this path: [REDACTED:Hex High Entropy String] "UCP/RNDV: Use local memory detection for RTR_REQ responder (#12041)" and af8174b "UCP/UCT/CUDA: Avoid gdr_copy for non gdr capable memory (#12039)" in src/ucp/rndv/proto_rndv.c. The PR under test (#12077, UCM/UCT ROCm IPC remote-mapping cache) does not touch UCP rndv protocol selection or CUDA, and the identical test failure is associated with several unrelated PRs (#12054, #12058, #12081, #11990), confirming it is a pre-existing environment-sensitive/flaky assertion, not a regression from this branch.

Implicated commit: Test assertion introduced by 9cc0808 "UCP/RNDV: use tag-only signal for CUDA staging force" (Nathan Bellalou); most likely behavior shift from [REDACTED:Hex High Entropy String] (#12041, Guy Ealey Morag) / af8174b (#12039, Thomas Vegas). Not commit 63a1ccd (PR #12077).

File: test/gtest/ucp/test_ucp_proto.cc:710-737 (assertions at :728-730 and :736); instantiation at test/gtest/ucp/test_ucp_proto.cc:858

Suggested fix: 1) Re-run the Azure job for PR #12077 — the failure is unrelated to the ROCm IPC-cache change. 2) Make the test robust instead of asserting on a formatted protocol string: gate it on the forced-pipeline path actually being selectable (as the sibling test does by first checking for "rendezvous pipeline" and require_cuda_net_md()), and extend the skip at line 724 to also UCS_TEST_SKIP_R when the selected protocol is not a pipeline/cuda_copy protocol, or assert on proto_config.proto->name / remote_op_flags (UCP_PROTO_SELECT_OP_FLAG_TAG_RNDV) rather than on substrings of the debug info. 3) Open a tracking issue against the owners of #12041/#12039 to confirm whether the forced CUDA-staging send child should still win on a shm,ib,cuda_copy (no gdr_copy) configuration.

Related: PR #12077 (this PR), and the same test reported failing on unrelated PRs #12054, #12058, #12081, #11990; relevant protocol-selection changes: #12041, #12039.

Note: the Azure log API returned only the last ~500 lines of the job, which contain the gtest summary but not the individual EXPECT failure message. For the exact failing expectation, grab the full raw log (.../_apis/build/builds/141273/logs/<logId>) and search for rndv_force_cuda_send_child_uses_tag_provenance.

🛡️ This comment had 1 potential secret(s) redacted (Hex High Entropy String). See request_id 10f08cd6-c3aa-41ac-b4a1-fda6bcfbb290 in the triage console for the audit trail.

[edgargabriel]

Looks good to me, and also passed in my testing.

However, because this is a massive PR, I will hold back merging the PR to happen after two smaller ones:

• #12150
• #12178

This way we can pin that UCX hash for a NIXL task with all required fixes before this PR is merged, just in case some issues pop up. I expect those two PRs to be merged within the next 1-2 days.