UCT/ROCM: Enable dmabuf by default with peermem fallback by ddebonis-amd · Pull Request #11934 · openucx/ucx

ddebonis-amd · 2026-09-09T21:55:29Z

What?

Use dma-buf for GPU-direct when the kernel supports it, and keep peermem as the fallback on older kernels. Also report allocation-base offsets for interior ROCm addresses.

Why?

The rocm components in UCX have support for the dma-buf mechanism to enable direct data transfer between an RDMA NIC and a GPU. However, the user explicitly needs to request that support by setting UCX_ROCM_COPY_DMABUF=y, dma-buf support is not turned on by default.

How?

Starting from linux kernel > 6.0, the peermem mechanism that was used previously is not useable in the vast majority of deployments, necessitating using dma-buf instead. The goal of this story is to enable dma-buf by default in the rocm components of UCX, and ensure that if dma-buf support is not available (e.g. linux kernel 5.15 in most cases) we can fall back to the peermem mechanism if that is available.

svc-nvidia-pr-review · 2026-09-09T21:55:45Z

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

svc-ucx · 2026-09-09T21:58:47Z

🤖 CI Triage Agent — UCX PR (Codestyle AUTHORS file update check) · commit a5336a72

TL;DR: This is not a code/test bug — the "AUTHORS file update check" failed because PR #11934's commit author, David DeBonis ddebonis@amd.com, is a new contributor who is not listed in the AUTHORS file. The fix is for the PR to add that line to AUTHORS (easiest by running ./contrib/authors_update.sh locally and committing the result).

Full analysis

Summary: The AUTHORS file check bash step in the Codestyle job exited with code 1 because git diff --exit-code detected an uncommitted change to AUTHORS after contrib/authors_update.sh ran.

Root cause: The pipeline step computes the PR commit range (8d1fb96..a5336a72) and runs ./contrib/authors_update.sh on it, then asserts the working tree is clean. The script reported a new, unknown author:

Names:
++ David DeBonis <ddebonis@amd.com>

Since neither the email ddebonis@amd.com nor the name David DeBonis matches any existing entry in AUTHORS (the file jumps from Daniel Pressler at line 26 straight to David Wootton at line 27), the script appended the entry, git diff --exit-code saw the modification and returned 1, which the Bash task surfaced as ##[error]Bash exited with code '1'. This is the check working as designed — it enforces that new contributors add themselves to AUTHORS in the same PR. There is no hang, timeout, infrastructure fault, or regression involved; the whole job ran in about 1 second.

Implicated commit: [REDACTED:Hex High Entropy String] (David DeBonis ddebonis@amd.com) — the PR head commit, which introduces a new author without an AUTHORS entry. The check script itself is unchanged and correct.

File: AUTHORS:26 (insertion point, between Daniel Pressler and David Wootton); check logic at contrib/authors_update.sh:36

Suggested fix: In PR #11934, add the author to AUTHORS and amend/append a commit:

# from the repo root, on the PR branch
./contrib/authors_update.sh $(git merge-base HEAD origin/develop)..HEAD
git add AUTHORS && git commit -m "AUTHORS: add David DeBonis"

Or simply insert David DeBonis <ddebonis@amd.com> manually after line 26 (Daniel Pressler <danielpr@nvidia.com>), preserving LC_COLLATE=C sort order as the script does. No pipeline or script change is warranted. One caveat worth noting to the contributor: the script's dedup check at contrib/authors_update.sh:36 uses a substring grep -iq "$name", so if the same person previously contributed under a different name/email spelling, the correct action is to normalize via .mailmap rather than adding a duplicate entry.

Related: PR #11934 (the failing PR). No pre-existing issue tracks this failure mode; recent AUTHORS commits (e.g. c2567cbe, 6d75cec7) show the same pattern of new contributors adding themselves alongside their change.

🛡️ This comment had 1 potential secret(s) redacted (Hex High Entropy String). See request_id 305cfe92-def8-4a3b-b19a-4e065a82fa8b in the triage console for the audit trail.

svc-nvidia-pr-review · 2026-09-09T22:01:29Z

Residual coverage gap: the ROCm CI containers (ubuntu2004_rocm_5_4_0, ubuntu2204_rocm_6_0_0) are build-only, so neither mem_query_dmabuf nor the new disable test is exercised by the PR pipeline; both need a run on real AMD hardware to validate the default change.

svc-nvidia-pr-review · 2026-09-09T22:22:03Z

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

svc-ucx · 2026-09-09T22:23:42Z

🤖 CI Triage Agent — UCX PR (Codestyle commit title) · commit 6ce012cd

TL;DR: The "commit title" codestyle job failed because one commit in PR #11934 is titled added author and reduced warning to debug, which doesn't match UCX's required MODULE: Description prefix format — the author needs to reword that commit (e.g. UCT/ROCM: Add author and reduce warning to debug) and force-push.

Full analysis

Summary: Azure Pipelines job Codestyle / commit title (build 135339) exited with code 1 after codestyle_check_commit_title rejected a commit title in PR #11934.

Root cause: Not a product or infrastructure bug. The log shows:

Bad commit title: 'added author and reduced warning to debug'
Good commit title: 'UCT/ROCM: Enable dmabuf by default with peermem fallback ...'

codestyle_check_commit_title() requires every non-merge commit subject in base..head to match the regex ^[0-9A-Z/_-]+:\ [[:alnum:]_] — i.e. an uppercase module prefix followed by ": " and an alphanumeric first character (e.g. UCT/ROCM: Add ...). The commit added author and reduced warning to debug has no module prefix and starts lowercase, so err=1 was returned and the job failed.

Secondary (cosmetic) CI issue visible in the same log: the step runs with set -eE, so the non-zero return from codestyle_check_commit_title aborts the script immediately. The if [[ $? -ne 0 ]] block in buildlib/pr/codestyle.yml:25-30 is therefore dead code — the helpful Bad commit title(s), see <wiki url> message never printed, leaving only a bare Bash exited with code '1', which makes the failure harder to self-diagnose.

Implicated commit: [REDACTED:Hex High Entropy String] (PR #11934 head, merged as 63e9a2a for CI) — the offending commit is the one titled added author and reduced warning to debug. Author not shown in this job's log.

File: buildlib/tools/codestyle.sh:14 (the enforcing regex); buildlib/pr/codestyle.yml:17-30 (the failing step)

Suggested fix:

Contributor action (fixes this build): reword the commit, e.g.
```
git rebase -i <base>   # or: git commit --amend  if it's the tip commit
# retitle to: "UCT/ROCM: Add author and reduce warning to debug"
git push --force-with-lease
```
Also consider trimming the other commit's subject — its body text appears to be on the subject line, making a very long title (it passes the regex but violates the ≤50-char subject convention).

Optional CI hardening so future failures are self-explanatory: in buildlib/pr/codestyle.yml, stop the set -e from swallowing the message, e.g.

if ! codestyle_check_commit_title "$base_revision..$head_revision"; then
    url="https://github.com/openucx/ucx/wiki/Guidance-for-contributors#general-guidelines"
    echo "##vso[task.logissue type=error]Bad commit title(s), see $url for more info."
    echo "##vso[task.complete result=Failed;]"
fi

Related: PR #11934; contributor guidance: 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 e90fcf24-79fe-4ec1-ae9d-8aedfd0649f2 in the triage console for the audit trail.

svc-ucx · 2026-09-09T22:26:46Z

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

TL;DR: The "commit title" codestyle job failed because two commits in PR #11934 have titles that don't match UCX's required MODULE: Description format (Update src/uct/rocm/base/rocm_base.c and added author and reduced warning to debug). Fix by rebasing/squashing the PR into commits with properly prefixed titles, e.g. UCT/ROCM: ....

Full analysis

Summary: Azure job UCX PR (Codestyle commit title) build 135343 failed with Bash exited with code '1' after codestyle_check_commit_title rejected two commit titles in PR #11934.

Root cause: Not an infrastructure or product bug. buildlib/pr/codestyle.yml runs codestyle_check_commit_title over base..head of the PR, which requires every non-merge commit title to match the regex ^[0-9A-Z/_-]+:\ [[:alnum:]_] (an uppercase module prefix, a colon, a space, then alphanumeric text). The log shows:

Bad commit title: 'Update src/uct/rocm/base/rocm_base.c' — GitHub web-editor default message, no module prefix.
Bad commit title: 'added author and reduced warning to debug' — lowercase, no module prefix.
Good commit title: 'UCT/ROCM: Enable dmabuf by default with peermem fallback ...' — passes the regex, but note the entire commit body appears to be crammed onto the subject line, which is itself a style problem.

Both bad titles set err=1, the script returns non-zero under set -eE, and the task fails immediately (total job runtime ~1s — no hang or timeout involved).

Implicated commit: [REDACTED:Hex High Entropy String] (PR #11934 head, merged as 2f6f623 for testing) — the offending commits are authored by the PR submitter, not by a UCX maintainer change.

File: buildlib/tools/codestyle.sh:14 (check logic) — enforced via buildlib/pr/codestyle.yml:24

Suggested fix: Rewrite the PR history so every commit title has the required form. Simplest is to squash the three commits into one:

git rebase -i <base>        # or: git reset --soft <base> && git commit
# subject line, <= ~50-72 chars, module prefix required:
UCT/ROCM: Enable dmabuf by default with peermem fallback

# body (blank line above), free-form:
Use dma-buf for GPU-direct when the kernel supports it, and keep
peermem as the fallback on older kernels. Also report allocation-base
offsets for interior ROCm addresses.
git push --force-with-lease

If the commits must stay separate, amend each title to e.g. UCT/ROCM: Reduce warning to debug level and AUTHORS: Add contributor. Also move the long description out of the subject into the commit body — the current "good" commit only passes because the regex checks the prefix, not the length.

Related: PR #11934; contributor guidance referenced by the job: 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 3e1de382-fe07-40ac-aac3-0b9ed929ca15 in the triage console for the audit trail.

svc-nvidia-pr-review · 2026-09-09T22:29:29Z

Residual gaps: both mem_query_dmabuf and the new disable test need an AMD GPU run to validate the offset math and the default flip; the disable test also gives no coverage that a failed hsa_amd_portable_export_dmabuf() with the new try default degrades gracefully (fd -1 → IB non-dmabuf path), which is the main behavioral risk of this PR.

svc-nvidia-pr-review · 2026-09-09T22:30:03Z

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

svc-nvidia-pr-review · 2026-09-09T22:35:14Z

Residual coverage gap: the ROCm CI jobs are build-only, so neither the compile break nor the offset math would be caught by the PR pipeline; an AMD GPU run of test_md* is needed.

svc-nvidia-pr-review · 2026-09-09T22:37:24Z

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

svc-ucx · 2026-09-09T22:39:57Z

🤖 CI Triage Agent — UCX PR (Codestyle commit title) · commit 384eac45

TL;DR: This is not an infrastructure failure — the "commit title" codestyle job correctly rejected PR #11934 because two of its commits have titles that lack the required MODULE: prefix; the author needs to reword those commits (and the CI script should be fixed so it prints the helpful error instead of a bare exit code 1).

Full analysis

Summary: Azure Pipelines job UCX PR (Codestyle / commit title) build 135347 failed with ##[error]Bash exited with code '1' after codestyle_check_commit_title flagged two commit titles in PR #11934.

Root cause: The job runs codestyle_check_commit_title "$base..$head", which requires every non-merge commit subject to match ^[0-9A-Z/_-]+:\ [[:alnum:]_] (i.e. an uppercase module prefix like UCT/ROCM: ). The log shows the checker's verdict per commit:

Good: UCX/ROCM: Update src/uct/rocm/base/rocm_base.c
Bad: Update src/uct/rocm/base/rocm_base.c — no module prefix
Bad: added author and reduced warning to debug — no module prefix, lowercase
Good: UCT/ROCM: Enable dmabuf by default with peermem fallback Use dma-buf for GPU-direct when the kernel supports it, and keep peermem as the fallback on older kernels. Also report allocation-base offsets for interior ROCm addresses.

Two bad titles set err=1, so the function returned 1. Note a secondary CI-script bug: the step begins with set -eE, so the non-zero return aborts the script immediately and the if [[ $? -ne 0 ]] block (which prints the wiki guidance link) is dead code — hence the log shows only a raw Bash exited with code '1' with no explanation. Also visible: the last commit's subject line has swallowed its whole body (no blank line after the subject), which passes the regex but violates the 50-char-subject convention.

Implicated commit: No product regression. The offending commits are the intermediate commits of PR #11934 (merge commit 813df758b, PR head [REDACTED:Hex High Entropy String]); the CI-script defect predates this build in buildlib/pr/codestyle.yml.

File: buildlib/tools/codestyle.sh:14 (the title regex) and buildlib/pr/codestyle.yml:18-30 (set -eE making the error-reporting branch unreachable)

Suggested fix:

PR author: interactively rebase and reword/squash the commits so every title has a module prefix and a short subject, e.g.
- Update src/uct/rocm/base/rocm_base.c → UCT/ROCM: Update rocm_base dmabuf detection
- added author and reduced warning to debug → UCT/ROCM: Add author entry, lower warning to debug
- Split the long fourth title: keep UCT/ROCM: Enable dmabuf by default with peermem fallback as the subject and move the rest into the body after a blank line.
  Simplest path: git rebase -i and squash all four into one properly titled commit, then force-push.
CI maintainers: make the failure self-explanatory by replacing the set -eE + $? pattern with if ! codestyle_check_commit_title "$base_revision..$head_revision"; then ... fi (or set +e around the call) in buildlib/pr/codestyle.yml, so the wiki guidance message is actually emitted.

Related: PR #11934; guidelines referenced by the job: 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 5f81c1af-1070-40f7-b044-a9ed75a3dceb in the triage console for the audit trail.

svc-nvidia-pr-review · 2026-09-10T13:43:39Z

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

svc-nvidia-pr-review · 2026-09-10T13:52:53Z

Residual coverage gap: ROCm CI jobs are build-only, so neither mem_query_dmabuf (now newly active for rocm_cpy because of the default flip) nor the new disable test is exercised; both need an AMD GPU run.

Residual coverage gap: mem_query_dmabuf still asserts EXPECT_EQ(mem_buf.ptr(), mem_attr.base_address) and EXPECT_EQ(offset, mem_attr.dmabuf_offset), which additionally assumes HSA returns offset 0 for the allocation base and that HIP does not sub-allocate; only a hardware run can confirm.

Residual coverage gap: nothing covers a failing hsa_amd_portable_export_dmabuf() under the new try default (fd -1 → IB falls back to the non-dmabuf path).

svc-nvidia-pr-review · 2026-09-10T13:53:34Z

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

svc-nvidia-pr-review · 2026-09-10T14:04:15Z


 #include <ucs/sys/string.h>
 #include <ucs/sys/module.h>
+#include <ucs/sys/ptr_arith.h>


minor: ptr_arith.h isn't needed here, UCS_PTR_BYTE_DIFF comes from ucs/sys/compiler_def.h (ze_copy_md.c uses it for the same dmabuf offset math without this include).

svc-nvidia-pr-review · 2026-09-10T14:04:15Z

+                                    UCT_MD_MEM_ATTR_V2_FIELD_DMABUF_OFFSET);
+    }
+
+    status                 = uct_rocm_base_mem_query(uct_md, addr, length,


Redundant state in the new wrapper: can we pass md->have_dmabuf down to uct_rocm_base_mem_query() instead of clearing and restoring the caller's field_mask? It has only this one caller, so that would drop both the mask juggling and the fd/offset fixup below.

svc-nvidia-pr-review · 2026-09-10T14:04:17Z

Residual gap: mem_type_detect_mds now asserts EXPECT_GE(alloc_length, buffer_size) for ROCm but keeps EXPECT_EQ(address, base_address). If HSA reports a larger sizeInBytes because HIP sub-allocated from a slab (rather than page-rounding), base_address would also differ and the strict check would fail — only an AMD GPU run distinguishes the two cases.

svc-nvidia-pr-review · 2026-09-10T14:04:57Z

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

svc-nvidia-pr-review · 2026-09-10T14:12:02Z

+        (mem_attr_p->field_mask & UCT_MD_MEM_ATTR_V2_FIELD_DMABUF_OFFSET)) {
+        uct_rocm_base_dmabuf_export(export_addr, export_len, mem_type,
+                                    &dmabuf_fd, &dmabuf_offset);
+        if (mem_attr_p->field_mask & UCT_MD_MEM_ATTR_V2_FIELD_DMABUF_FD) {


blocker: duplicated if (DMABUF_FD) block — returns a closed fd and leaks the fd.

if (mem_attr_p->field_mask & ..._DMABUF_FD) { mem_attr_p->dmabuf_fd = dmabuf_fd; if (mem_attr_p->field_mask & ..._DMABUF_FD) { mem_attr_p->dmabuf_fd = dmabuf_fd; } else { ucs_close_fd(&dmabuf_fd); } ucs_close_fd(&dmabuf_fd); }

Brace matching makes this compile but it is semantically wrong in both directions:

When DMABUF_FD is requested, the fd is stored into mem_attr_p->dmabuf_fd and then unconditionally closed by the trailing ucs_close_fd(). The caller (ucp_memh_register_internal() → ibv_reg_dmabuf_mr(), and ucp_memh_register_internal()'s own ucs_close_fd(&reg_params.dmabuf_fd)) receives a stale fd number: registration fails with EBADF, or worse, the number is reused by another thread and a wrong file gets registered / double-closed.

When only DMABUF_OFFSET is requested, the whole block is skipped, so the exported fd is never closed — fd leak on every offset-only query (gtest check_fd_leaks() will also flag it).

Suggested change

if (mem_attr_p->field_mask & UCT_MD_MEM_ATTR_V2_FIELD_DMABUF_FD) {

if (mem_attr_p->field_mask & UCT_MD_MEM_ATTR_V2_FIELD_DMABUF_FD) {

mem_attr_p->dmabuf_fd = dmabuf_fd;

} else {

ucs_close_fd(&dmabuf_fd);

}

svc-nvidia-pr-review · 2026-09-10T14:12:02Z

    }
 }

+UCS_TEST_P(test_md_dmabuf, rocm_dmabuf_disable, "ROCM_COPY_DMABUF~=no")


minor (test placement/coverage), both non-blocking:

The test asks for UCT_MD_MEM_ATTR_FIELD_DMABUF_OFFSET but never checks it; EXPECT_EQ(0, mem_attr.dmabuf_offset) would also cover the disabled path of the new offset math.

It lives in the generic test_md_dmabuf fixture and relies on ROCM_COPY_DMABUF~=no to skip all other MDs. Existing precedent for a component-specific test is a dedicated fixture plus _UCT_MD_INSTANTIATE_TEST_CASE(test_rocm, rocm_cpy) (see test_gdr_copy), which makes the intent explicit instead of depending on the config-skip trick.

svc-nvidia-pr-review · 2026-09-10T14:12:04Z

Residual gaps: the two mangled hunks mean the current head almost certainly does not build with ROCm enabled; ROCm CI jobs are build-only and not part of the default pipeline, so neither the compile break nor the fd lifetime bug would be caught. An AMD GPU run of gtest.uct.test_md* is required.

Residual gap: nothing exercises a failing hsa_amd_portable_export_dmabuf() under the new try default (fd -1 → IB must fall back to the non-dmabuf path).

svc-nvidia-pr-review · 2026-09-10T14:12:46Z

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

svc-nvidia-pr-review · 2026-09-10T14:22:16Z

     UCS_CONFIG_TYPE_TABLE(ucs_config_rcache_table)},

-    {"DMABUF", "no",
+    {"DMABUF", "try",


With try as default, md_attr.dmabuf_mem_types now advertises ROCm dmabuf based only on the kernel-config/kallsyms heuristic in uct_rocm_base_is_dmabuf_supported(). UCP then adds every dmabuf-capable MD to reg_md_map[ROCM] (ucp_context.c, "we can register memory of this type with any md that supports dmabuf"), so ROCm buffers are pushed to IB MDs even when hsa_amd_portable_export_dmabuf() fails at runtime and peermem is not loaded — that turns "not registered on IB" into a registration error. cuda_copy handles this by clearing UCS_MEM_FLAG_REGISTRABLE when the export fails (uct_cuda_copy_md_is_registrable()), while uct_rocm_base_mem_query() always reports UCS_MEM_FLAG_REGISTRABLE in ucs_memtype_cache_update(). Can we detect the export failure the same way, or is the peermem fallback guaranteed to cover it?

svc-nvidia-pr-review · 2026-09-10T14:22:16Z

            ASSERT_UCS_OK(status);
            EXPECT_EQ(alloc_mem_type, mem_attr.mem_type);
            if ((alloc_mem_type == UCS_MEMORY_TYPE_CUDA) ||
+                (alloc_mem_type == UCS_MEMORY_TYPE_ROCM) ||


Minor: the outer condition now lists UCS_MEMORY_TYPE_ROCM and then re-tests it inside; pls give ROCm its own branch so the expectations read directly.

svc-nvidia-pr-review · 2026-09-10T14:22:18Z

Residual coverage gap: ROCm CI is build-only, so neither the compile break nor the offset math is exercised by the PR pipeline; an AMD GPU run of gtest.uct.test_md* is still required.

svc-nvidia-pr-review · 2026-09-10T14:27:48Z

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

svc-nvidia-pr-review · 2026-09-10T14:34:56Z

-                     length);
+            ucs_debug("failed to export dmabuf handle for addr %p / %zu: 0x%x",
+                      addr, length, status);
+            ucs_debug("failed to export dmabuf handle for addr %p / %zu: 0x%x",


the log line is duplicated

Suggested change

ucs_debug("failed to export dmabuf handle for addr %p / %zu: 0x%x",

ucs_debug("failed to export dmabuf handle for addr %p / %zu: 0x%x",

addr, length, status);

svc-nvidia-pr-review · 2026-09-10T14:34:56Z

-        mem_attr_p->dmabuf_fd = dmabuf_fd;
-        is_exported           = 1;
-    }
+        if (mem_attr_p->field_mask & UCT_MD_MEM_ATTR_V2_FIELD_DMABUF_FD) {


minor: ucs_close_fd() already ignores UCT_DMABUF_FD_INVALID, so the extra check can go

Suggested change

if (mem_attr_p->field_mask & UCT_MD_MEM_ATTR_V2_FIELD_DMABUF_FD) {

if (mem_attr_p->field_mask & UCT_MD_MEM_ATTR_V2_FIELD_DMABUF_FD) {

mem_attr_p->dmabuf_fd = dmabuf_fd;

} else {

ucs_close_fd(&dmabuf_fd);

}

svc-nvidia-pr-review · 2026-09-10T14:34:56Z

    return UCS_OK;
 }

+static ucs_status_t


can we pass md->have_dmabuf down to uct_rocm_base_mem_query() instead of clearing and restoring mem_attr_p->field_mask? it avoids temporarily modifying the caller's input struct, and keeps the dmabuf decision in one place like cuda_copy does.

svc-nvidia-pr-review · 2026-09-10T14:34:56Z

    }
 }

+UCS_TEST_P(test_md_dmabuf, rocm_dmabuf_disable, "ROCM_COPY_DMABUF~=no")


minor: can we put this in a rocm-only fixture with _UCT_MD_INSTANTIATE_TEST_CASE(test_rocm, rocm_cpy) (like test_gdr_copy) instead of relying on the ~= modifier to skip it for every other MD?

svc-nvidia-pr-review · 2026-09-10T14:34:56Z

@@ -542,10 +542,16 @@ UCS_TEST_P(test_md, mem_type_detect_mds) {
            ASSERT_UCS_OK(status);
            EXPECT_EQ(alloc_mem_type, mem_attr.mem_type);
            if ((alloc_mem_type == UCS_MEMORY_TYPE_CUDA) ||


minor: ROCM is tested twice in nested conditions; a separate if (alloc_mem_type == UCS_MEMORY_TYPE_ROCM) branch before the CUDA/ZE one would read better.

svc-nvidia-pr-review · 2026-09-10T14:34:58Z

Residual gap: ROCm CI jobs are build-only, so neither the offset math (mem_query_dmabuf, now active for rocm_cpy because of the try default) nor the new disable test runs in this pipeline; an AMD GPU run of gtest.uct.test_md* is still needed.

Residual gap: EXPECT_GE(alloc_length, buffer_size) combined with EXPECT_EQ(address, base_address) in mem_type_detect_mds only holds if HSA page-rounds sizeInBytes; if HIP sub-allocates from a slab, base_address differs too and the strict check fails. Only hardware distinguishes the two.

svc-nvidia-pr-review · 2026-09-10T19:57:28Z

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

svc-nvidia-pr-review · 2026-09-10T20:07:46Z

+                        uct_md_mem_attr_v2_t *mem_attr_p)
+{
+    uct_rocm_copy_md_t *md = ucs_derived_of(uct_md, uct_rocm_copy_md_t);
+    uint64_t orig_mask     = mem_attr_p->field_mask;


why not pass md->have_dmabuf into uct_rocm_base_mem_query() and gate uct_rocm_base_dmabuf_export() there? it has no other caller, so that drops the temporary mutation of the caller's mem_attr_p->field_mask and the manual re-init of the two fields.

Rationale: mem_attr_p is a caller-owned in/out struct; clearing bits in it and restoring them after the callee returns is fragile (it also makes the base run with an empty mask purely for its side effects, where an HSA failure is still propagated as a raw hsa_status_t). Passing the flag down keeps the base's behavior identical to before and removes ~25 lines.

svc-nvidia-pr-review · 2026-09-10T20:07:46Z

    }
 }

+UCS_TEST_P(test_md_dmabuf, rocm_dmabuf_disable, "ROCM_COPY_DMABUF~=no")


minor: this test only applies to rocm_cpy, so a dedicated fixture with _UCT_MD_INSTANTIATE_TEST_CASE(test_rocm_copy, rocm_cpy) (as done for test_gdr_copy) is clearer than relying on ~= to skip the other 12 MD instances.

svc-nvidia-pr-review · 2026-09-11T15:46:41Z

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

svc-nvidia-pr-review · 2026-09-11T15:55:35Z

    ucs_sys_device_t sys_dev   = UCS_SYS_DEVICE_ID_UNKNOWN;
-    int dmabuf_fd;
+    int dmabuf_fd              = UCT_DMABUF_FD_INVALID;
    hsa_status_t status;


the early exit below returns status (an hsa_status_t) as ucs_status_t, so a uct_rocm_base_get_ptr_info() failure propagates e.g. 0x1002 up to ucp_memh_register_internal(). this path becomes reachable in gtest now that test_md_dmabuf runs for rocm_cpy and queries host memory too. can we return UCS_ERR_INVALID_ADDR instead?

svc-nvidia-pr-review · 2026-09-11T15:55:35Z

+        if (have_dmabuf) {
+            uct_rocm_base_dmabuf_export(base_addr, base_size, mem_type,
+                                        &dmabuf_fd, &dmabuf_offset);
+        }


minor: pls add a blank line after the block, per docs/CodeStyle.md ("Use blank line after if/else blocks").

Suggested change

}

}

svc-nvidia-pr-review · 2026-09-11T16:02:03Z

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

svc-nvidia-pr-review · 2026-09-11T16:11:23Z

@@ -447,7 +455,7 @@ ucs_status_t uct_rocm_base_mem_query(uct_md_h md, const void *addr,
                                        &base_size, &hsa_mem_type, &agent,


minor: ucs_close_fd() already returns early for -1, so the extra check can go.

Suggested change

&base_size, &hsa_mem_type, &agent,

if (mem_attr_p->field_mask & UCT_MD_MEM_ATTR_V2_FIELD_DMABUF_FD) {

mem_attr_p->dmabuf_fd = dmabuf_fd;

} else {

ucs_close_fd(&dmabuf_fd);

}

svc-nvidia-pr-review · 2026-09-11T16:39:23Z

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

svc-nvidia-pr-review · 2026-09-11T16:47:53Z

-                                        &dmabuf_offset);
+        if (mem_attr_p->field_mask & UCT_MD_MEM_ATTR_V2_FIELD_DMABUF_OFFSET) {
+            mem_attr_p->dmabuf_offset = dmabuf_offset +
+                                        UCS_PTR_BYTE_DIFF(base_addr, addr);


dmabuf_offset is reported even when no dmabuf was exported. When have_dmabuf == 0 (e.g. ROCM_COPY_DMABUF=no), or when hsa_amd_portable_export_dmabuf() fails, dmabuf_fd stays UCT_DMABUF_FD_INVALID but we still report dmabuf_offset = addr - base_addr. Per the dmabuf_offset doc in uct.h (src/uct/api/uct.h:1686-1690) the offset is only defined relative to the dmabuf region identified by dmabuf_fd, so a non-zero value with an invalid fd has no meaning — please leave it 0 when nothing was exported. That also removes the hidden dependency behind EXPECT_EQ(0ul, mem_attr.dmabuf_offset) in the new dmabuf_disable test, which currently only passes because HSA happens to report agentBaseAddress == hipMalloc() pointer for that allocation. Note the fix cannot simply test dmabuf_fd == UCT_DMABUF_FD_INVALID at that point, because ucs_close_fd() resets it to -1 on the offset-only path; a separate "exported" flag or reordering is needed.

svc-nvidia-pr-review · 2026-09-11T16:47:53Z

 ## Current
 ### Features:
+#### ROCM
+ * Enabled dmabuf by default for GPU-direct transfers, falling back to peermem when unavailable


minor: user-facing doc not updated for the ROCm default flip. Can we also update the dmabuf section in docs/source/faq.md ("What is needed for dmabuf support?")? It currently lists only Cuda requirements, and since this PR makes dmabuf the default for ROCm, the ROCm prerequisites (kernel with CONFIG_DMABUF_MOVE_NOTIFY/CONFIG_PCI_P2PDMA, hsa_amd_portable_export_dmabuf availability) belong there too.

svc-nvidia-pr-review · 2026-09-11T16:47:54Z

Residual coverage gap: ROCm CI containers are build-only, so neither mem_query_dmabuf (newly active for rocm_cpy after the default flip) nor the new dmabuf_disable test runs in this pipeline; an AMD GPU run of gtest.uct.test_md* is needed.

Residual coverage gap: nothing exercises a failing hsa_amd_portable_export_dmabuf() under the new try default (fd -1 → uct_ib_reg_mr() falls back to ibv_reg_mr/peermem), which remains the main behavioral risk of the default change.

svc-nvidia-pr-review · 2026-09-11T16:58:34Z

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

svc-nvidia-pr-review · 2026-09-11T17:05:27Z

+        if (mem_attr_p->field_mask & UCT_MD_MEM_ATTR_V2_FIELD_DMABUF_FD) {
+            mem_attr_p->dmabuf_fd = dmabuf_fd;
+        } else if (dmabuf_fd != UCT_DMABUF_FD_INVALID) {
+            ucs_close_fd(&dmabuf_fd);


minor: ucs_close_fd() already returns early for an invalid fd, so the extra check and the second sentence of the comment above can be dropped.

Suggested change

ucs_close_fd(&dmabuf_fd);

} else {

ucs_close_fd(&dmabuf_fd);

}

svc-nvidia-pr-review · 2026-09-11T17:05:27Z

-                EXPECT_EQ(buffer_size, mem_attr.alloc_length);
+                /* ROCm HSA may round sizeInBytes up to a page */
+                if (alloc_mem_type == UCS_MEMORY_TYPE_ROCM) {
+                    EXPECT_GE(mem_attr.alloc_length, buffer_size);


minor: can we test ROCm in its own branch instead of nesting a second alloc_mem_type == UCS_MEMORY_TYPE_ROCM check inside the branch that just matched it?

svc-nvidia-pr-review · 2026-09-11T17:05:28Z

Residual coverage gap: The ROCm CI containers (ubuntu2004_rocm_5_4_0, ubuntu2204_rocm_6_0_0) only run build_rocm, so neither mem_query_dmabuf (newly active for rocm_cpy because of the try default) nor the new dmabuf_disable test executes in this pipeline; an AMD GPU run of gtest.uct.test_md* is needed.

Residual coverage gap: mem_type_detect_mds keeps EXPECT_EQ(address, mem_attr.base_address) while relaxing the length check. That only holds if HSA page-rounds sizeInBytes; if HIP sub-allocates from a slab, base_address differs too and the strict check fails. Only hardware distinguishes the two.

Residual coverage gap: Nothing exercises a failing hsa_amd_portable_export_dmabuf() under the new try default, nor the DMABUF=yes + unsupported md_open() error path whose leak this PR fixes.

svc-nvidia-pr-review · 2026-09-11T17:15:37Z

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

svc-nvidia-pr-review · 2026-09-14T13:00:52Z

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

svc-nvidia-pr-review · 2026-09-14T13:11:26Z

-        if (!is_exported) {
-            uct_rocm_base_dmabuf_export(addr, length, mem_type, &dmabuf_fd,
-                                        &dmabuf_offset);
+        /* Offset is defined only relative to a valid dmabuf_fd. Check the fd


minor: the second sentence is hard to parse ("Check the fd before close, which sets it to -1..."), which obscures the real reason the offset block must precede the fd/close block.

Suggested change

/* Offset is defined only relative to a valid dmabuf_fd. Check the fd

/* The offset is meaningful only with a valid dmabuf_fd, so evaluate it

* before the fd is closed below on the offset-only query path. */

svc-nvidia-pr-review · 2026-09-14T13:11:26Z

                (alloc_mem_type == UCS_MEMORY_TYPE_ZE_DEVICE) ||
                (alloc_mem_type == UCS_MEMORY_TYPE_ZE_MANAGED)) {
-                EXPECT_EQ(buffer_size, mem_attr.alloc_length);
+                /* ROCm HSA may round sizeInBytes up to a page */


minor: the comment about ROCm page rounding sits above a check that covers CUDA/ZE too, so it reads as if it applies to all of them. Pls move the comment into the ROCm branch, it does not apply to the CUDA/ZE check.

Suggested change

/* ROCm HSA may round sizeInBytes up to a page */

if (alloc_mem_type == UCS_MEMORY_TYPE_ROCM) {

/* ROCm HSA may round the allocation size up to a page */

EXPECT_GE(mem_attr.alloc_length, buffer_size);

} else {

EXPECT_EQ(buffer_size, mem_attr.alloc_length);

}

svc-ucx · 2026-09-14T14:47:32Z

🤖 CI Triage Agent — UCX PR (AddressSanitizer roce on worker 1) · commit dcda3360

TL;DR: The "AddressSanitizer roce on worker 1" job was cancelled by Azure mid-step (not a test failure, no ASAN report) while still in the very first ASAN phase (run_malloc_hook_gtest); PR #11934 only touches src/uct/rocm/*, which this lane doesn't even build, so this is a CI wall-clock/cancellation issue — re-run the lane.

Full analysis

Summary: Run ./contrib/test_jenkins.sh in AddressSanitizer → roce → worker 1 ended with ##[error]The Operation will be canceled / The operation was canceled, right after starting the "Running cuda hooks with far jump" gtest step.

Root cause: From the log, no test failed and no AddressSanitizer report was produced. The last work shown is run_malloc_hook_gtest (contrib/test_jenkins.sh:1224-1251): the cuda_hooks.* run completed cleanly (Running 0 tests from 0 test suites — the roce host has no GPU: [ INFO ] srd,cuda_copy,rocm_copy is not available), then the far-jump variant started at 14:44:02 and was cut at 14:44:04. Timestamps in the whole retrievable window are continuous (max gap ≈1.2 s), so nothing was hung at cancellation time — the job was killed from outside while making progress. Notably the cancellation hit several long-running "roce on worker N" jobs, including one from a different stage (Tests roce as well as AddressSanitizer roce), within the same second, which points at a build-level cancellation/timeout of the still-in-flight roce lanes rather than a per-test fault.

Contributing structural issue: worker 1 had spent the whole job budget without even reaching the main suite (run_gtest_make, contrib/test_jenkins.sh:1315). Each gtest phase in run_malloc_hook_gtest is wrapped in TIMEOUT="timeout 200m" (contrib/test_jenkins.sh:45) while the Azure job cap is timeoutInMinutes: 360 (buildlib/pr/tests.yml:17). Two slow ASAN phases therefore exceed the job budget, and the inner timeout never gets to fail loudly — you get exactly this uninformative "canceled mid-step" signature instead of an actionable error.

Implicated commit: Not the PR commit. dcda3360 / 29cb04cb (David DeBonis, "UCT/ROCM: Enable dmabuf by default with peermem fallback") only touches src/uct/rocm/base/rocm_base.c and friends; the roce lane runs the ubuntu2204_ib container with no ROCm, so that code is not exercised.

File: contrib/test_jenkins.sh:1246 (step that was cut); contrib/test_jenkins.sh:45 and buildlib/pr/tests.yml:17 (timeout budgets)

Suggested fix:

Re-run the AddressSanitizer roce lane — this build carries no evidence of a code defect from PR UCT/ROCM: Enable dmabuf by default with peermem fallback #11934.
If it recurs, make the failure diagnosable instead of a bare cancellation: shrink the per-phase budget so it fits inside the job cap (e.g. TIMEOUT="timeout 60m" for ASAN_CHECK=yes, since 2×200 m > 360 m), and/or wrap each $TIMEOUT ... make -C test/gtest test so a 124 exit prints which GTEST_FILTER phase blew up.
Add elapsed-time echoes around each phase of run_malloc_hook_gtest / run_specific_tests so the log shows which ASAN phase consumed the ~6 h before the cut (the current tail-only log cannot attribute it).
Check the roce agent host for over-subscription — four ASAN gtest workers plus the non-ASAN Tests roce matrix on the same pool is a plausible reason these lanes are the only ones still running at cancel time.

Related: PR #11934 (UCT/ROCM: Enable dmabuf by default with peermem fallback) — unrelated to this lane's failure; no existing issue found for the roce ASAN cancellation signature.

…for GPU-direct when the kernel supports it, and keep peermem as the fallback on older kernels. Also report allocation-base offsets for interior ROCm addresses.

svc-nvidia-pr-review · 2026-09-15T13:16:35Z

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

svc-nvidia-pr-review · 2026-09-15T13:26:01Z

🤖 Review complete — no issues found.

svc-nvidia-pr-review reviewed Sep 9, 2026

View reviewed changes

Comment thread src/uct/rocm/base/rocm_base.c Outdated

Comment thread src/uct/rocm/copy/rocm_copy_md.c

Comment thread src/uct/rocm/base/rocm_base.c Outdated

svc-nvidia-pr-review reviewed Sep 9, 2026

View reviewed changes

Comment thread src/uct/rocm/base/rocm_base.c Outdated

Comment thread src/uct/rocm/base/rocm_base.c Outdated

Comment thread src/uct/rocm/base/rocm_base.c

Comment thread src/uct/rocm/copy/rocm_copy_md.c

svc-nvidia-pr-review reviewed Sep 9, 2026

View reviewed changes

Comment thread src/uct/rocm/base/rocm_base.c

Comment thread src/uct/rocm/base/rocm_base.c Outdated

Comment thread test/gtest/uct/test_md.cc Outdated

svc-nvidia-pr-review reviewed Sep 9, 2026

View reviewed changes

Comment thread src/uct/rocm/base/rocm_base.c

Comment thread src/uct/rocm/base/rocm_base.c Outdated

Comment thread test/gtest/uct/test_md.cc Outdated

svc-nvidia-pr-review reviewed Sep 10, 2026

View reviewed changes

Comment thread src/uct/rocm/base/rocm_base.c Outdated

Comment thread src/uct/rocm/base/rocm_base.c Outdated

Comment thread src/uct/rocm/base/rocm_base.c Outdated

svc-nvidia-pr-review reviewed Sep 10, 2026

View reviewed changes

ddebonis-amd force-pushed the develop branch from 7e3dddc to 4a9f32a Compare September 10, 2026 19:57

svc-nvidia-pr-review reviewed Sep 10, 2026

View reviewed changes

ddebonis-amd force-pushed the develop branch from 4a9f32a to ba5cbb6 Compare September 11, 2026 15:46

svc-nvidia-pr-review reviewed Sep 11, 2026

View reviewed changes

ddebonis-amd force-pushed the develop branch from ba5cbb6 to cb780f7 Compare September 11, 2026 16:01

svc-nvidia-pr-review reviewed Sep 11, 2026

View reviewed changes

ddebonis-amd force-pushed the develop branch from cb780f7 to 2e8426a Compare September 11, 2026 16:39

svc-nvidia-pr-review reviewed Sep 11, 2026

View reviewed changes

ddebonis-amd force-pushed the develop branch from 2e8426a to 6b6bb7b Compare September 11, 2026 16:58

svc-nvidia-pr-review reviewed Sep 11, 2026

View reviewed changes

ddebonis-amd force-pushed the develop branch from 6b6bb7b to 29cb04c Compare September 11, 2026 17:15

svc-nvidia-pr-review reviewed Sep 11, 2026

View reviewed changes

Comment thread test/gtest/uct/test_md.cc

svc-nvidia-pr-review reviewed Sep 14, 2026

View reviewed changes

ddebonis-amd force-pushed the develop branch from dcda336 to 29cb04c Compare September 15, 2026 13:10

UCT/ROCM: Enable dmabuf by default with peermem fallback Use dma-buf …

57c450e

…for GPU-direct when the kernel supports it, and keep peermem as the fallback on older kernels. Also report allocation-base offsets for interior ROCm addresses.

ddebonis-amd force-pushed the develop branch from 29cb04c to 57c450e Compare September 15, 2026 13:16

edgargabriel approved these changes Sep 16, 2026

View reviewed changes

edgargabriel merged commit 9c85cdf into openucx:master Sep 16, 2026
162 checks passed

	ucs_debug("failed to export dmabuf handle for addr %p / %zu: 0x%x",
	ucs_debug("failed to export dmabuf handle for addr %p / %zu: 0x%x",
	addr, length, status);

		@@ -447,7 +455,7 @@ ucs_status_t uct_rocm_base_mem_query(uct_md_h md, const void *addr,
		&base_size, &hsa_mem_type, &agent,

	/* Offset is defined only relative to a valid dmabuf_fd. Check the fd
	/* The offset is meaningful only with a valid dmabuf_fd, so evaluate it
	* before the fd is closed below on the offset-only query path. */

Conversation

ddebonis-amd commented Sep 9, 2026

What?

Why?

How?

Uh oh!

svc-nvidia-pr-review commented Sep 9, 2026

Uh oh!

svc-ucx commented Sep 9, 2026

Uh oh!

Uh oh!

Uh oh!

Uh oh!

svc-nvidia-pr-review commented Sep 9, 2026

Uh oh!

svc-nvidia-pr-review commented Sep 9, 2026

Uh oh!

svc-ucx commented Sep 9, 2026

Uh oh!

svc-ucx commented Sep 9, 2026

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

svc-nvidia-pr-review commented Sep 9, 2026

Uh oh!

svc-nvidia-pr-review commented Sep 9, 2026

Uh oh!

Uh oh!

Uh oh!

Uh oh!

svc-nvidia-pr-review commented Sep 9, 2026

Uh oh!

svc-nvidia-pr-review commented Sep 9, 2026

Uh oh!

svc-ucx commented Sep 9, 2026

Uh oh!

Uh oh!

Uh oh!

Uh oh!

svc-nvidia-pr-review commented Sep 10, 2026

Uh oh!

Uh oh!

Uh oh!

Uh oh!

svc-nvidia-pr-review commented Sep 10, 2026

Uh oh!

svc-nvidia-pr-review commented Sep 10, 2026

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

svc-nvidia-pr-review commented Sep 10, 2026

Uh oh!

svc-nvidia-pr-review commented Sep 10, 2026

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

svc-nvidia-pr-review commented Sep 10, 2026

Uh oh!

svc-nvidia-pr-review commented Sep 10, 2026

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

svc-nvidia-pr-review commented Sep 10, 2026

Uh oh!

svc-nvidia-pr-review commented Sep 10, 2026

Uh oh!

Choose a reason for hiding this comment