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-ucx · 2026-09-10T18:53:27Z

🤖 CI Triage Agent — UCX PR (Build build_x86_64 ubuntu2004_rocm) · commit 595ff18c

TL;DR: The ubuntu2004_rocm build fails to compile src/uct/rocm/base/rocm_base.c because a bad merge/edit in commit 595ff18 left a duplicated ucs_close_fd(&dmabuf_fd); } block (lines 499–500), which closes uct_rocm_base_mem_query() prematurely and breaks brace balance. Delete those two stray lines.

Full analysis

Summary: make fails with -Werror in the ROCm module: rocm_base.c:508: error: expected identifier or '(' before 'if' plus control reaches end of non-void function in uct_rocm_base_mem_query.

Root cause: Syntax error from duplicated code, not an environment/timing issue. In uct_rocm_base_mem_query() the dmabuf-FD handling block is:

494:  if (mem_attr_p->field_mask & UCT_MD_MEM_ATTR_V2_FIELD_DMABUF_FD) {
495:      mem_attr_p->dmabuf_fd = dmabuf_fd;
496:  } else {
497:      ucs_close_fd(&dmabuf_fd);
498:  }
499:      ucs_close_fd(&dmabuf_fd);   <-- stray duplicate
500:  }                              <-- stray extra closing brace

The extra } on line 500 closes the enclosing if (... DMABUF_FD || ... DMABUF_OFFSET) block early, so the } on line 506 closes the function. Everything after (lines 508–514) then appears at file scope, producing "expected identifier or '(' before 'if'/'return'/'}'", and GCC also reports control reaches end of non-void function for uct_rocm_base_mem_query (fatal under -Werror). Only the ROCm-enabled build configuration compiles this file, which is why just build_x86_64 ubuntu2004_rocm failed.

Implicated commit: [REDACTED:Hex High Entropy String] — "UCT/ROCM: Enable dmabuf by default with peermem fallback", David DeBonis (PR #11934)

File: src/uct/rocm/base/rocm_base.c:499-500

Suggested fix: Delete the two duplicated lines so the block reads:

    if ((mem_attr_p->field_mask & UCT_MD_MEM_ATTR_V2_FIELD_DMABUF_FD) ||
        (mem_attr_p->field_mask & UCT_MD_MEM_ATTR_V2_FIELD_DMABUF_OFFSET)) {
        uct_rocm_base_dmabuf_export(base_addr, base_size, mem_type, &dmabuf_fd,
                                    &dmabuf_offset);
        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);
        }

        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);
        }
    }

Then re-push; a local ./contrib/configure-devel --with-rocm && make -j (or any ROCm-enabled build) will catch this before CI. Also worth double-checking the rest of the rebase in this commit for similar duplicated hunks.

Related: PR #11934 (the PR under test); prior ROCm dmabuf work in #11541.

🛡️ This comment had 1 potential secret(s) redacted (Hex High Entropy String). See request_id 693c1fb5-7d8d-40cb-9964-8dfa9f3190ca in the triage console for the audit trail.

svc-nvidia-pr-review · 2026-09-10T18:55:51Z

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


blocker: braces are unbalanced here (extra } after the offset block) and ucs_close_fd() is duplicated, so this does not compile and the fd would be closed twice. Additionally, the DMABUF_OFFSET block ends up outside the if that produced dmabuf_offset.

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_OFFSET) {

mem_attr_p->dmabuf_offset = dmabuf_offset +

UCS_PTR_BYTE_DIFF(base_addr, addr);

}

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);

}

}

Also, if ((mask & FD) || (mask & OFFSET)) can be if (mask & (FD | OFFSET)).

svc-nvidia-pr-review · 2026-09-10T18:55:51Z

+    ASSERT_UCS_OK(uct_md_mem_query(md(), mem_buf.ptr(), mem_buf.size(),
+                                   &mem_attr));
+    EXPECT_EQ(UCT_DMABUF_FD_INVALID, mem_attr.dmabuf_fd);
+    EXPECT_EQ(0ul, mem_attr.dmabuf_offset);


minor: duplicated line.

Suggested change

EXPECT_EQ(0ul, mem_attr.dmabuf_offset);

EXPECT_EQ(0ul, mem_attr.dmabuf_offset);

}

svc-nvidia-pr-review · 2026-09-10T18:55:51Z

    }
 }

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


The test is added to test_md_dmabuf, which is instantiated for every MD, and then skipped by comparing GetParam().md_name to "rocm_cpy". This file already has a pattern for MD-specific tests (_UCT_MD_INSTANTIATE_TEST_CASE(test_gdr_copy, gdr_copy)).

why not put this in its own class instantiated with _UCT_MD_INSTANTIATE_TEST_CASE(test_rocm, rocm_cpy), like test_gdr_copy? then the md_name check is not needed.

svc-nvidia-pr-review · 2026-09-10T18:55:51Z

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


Downgrading the export failure from ucs_warn to ucs_debug is right for the new try default, but it also silences the ROCM_COPY_DMABUF=yes case: uct_rocm_copy_md_open() only validates kernel support, so a runtime hsa_amd_portable_export_dmabuf() failure under an explicit yes now produces no output at the default log level.

with ROCM_COPY_DMABUF=yes an export failure is now completely silent - can we keep at least ucs_diag here, or pass the ternary down so yes still reports?

svc-nvidia-pr-review · 2026-09-10T18:55:51Z

@@ -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: the nested if inside the memory-type list is harder to read than splitting ROCm out.

Suggested change

if ((alloc_mem_type == UCS_MEMORY_TYPE_CUDA) ||

if (alloc_mem_type == UCS_MEMORY_TYPE_ROCM) {

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

EXPECT_GE(mem_attr.alloc_length, buffer_size);

EXPECT_EQ(address, mem_attr.base_address);

} else if ((alloc_mem_type == UCS_MEMORY_TYPE_CUDA) ||

(alloc_mem_type == UCS_MEMORY_TYPE_ZE_HOST) ||

(alloc_mem_type == UCS_MEMORY_TYPE_ZE_DEVICE) ||

(alloc_mem_type == UCS_MEMORY_TYPE_ZE_MANAGED)) {

EXPECT_EQ(buffer_size, mem_attr.alloc_length);

EXPECT_EQ(address, mem_attr.base_address);

svc-nvidia-pr-review · 2026-09-10T18:55:53Z

Residual coverage gap: ROCm CI jobs (ubuntu2004_rocm_5_4_0, ubuntu2204_rocm_6_0_0) are build-only, so the compile break in finding 1 and the offset math are not exercised by this pipeline; an AMD GPU run of gtest.uct.test_md* is required.

Residual coverage gap: mem_type_detect_mds combines EXPECT_GE(alloc_length, buffer_size) with EXPECT_EQ(address, base_address). That holds only 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 covers a failing hsa_amd_portable_export_dmabuf() under the new try default (fd -1 → IB must fall back to the non-dmabuf/peermem path), which is the main behavioral risk of the default flip.

svc-nvidia-pr-review · 2026-09-10T19:30:26Z

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

svc-nvidia-pr-review · 2026-09-10T19:41:29Z

    }
 }

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


minor: pls put this in its own fixture instantiated for rocm_cpy only, like test_gdr_copy below (_UCT_MD_INSTANTIATE_TEST_CASE(test_rocm, rocm_cpy)). Then the ~= modifier and the md_name != "rocm_cpy" skip are both unnecessary, and test_md_dmabuf stays generic.

svc-nvidia-pr-review · 2026-09-10T19:42:08Z

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

svc-nvidia-pr-review · 2026-09-10T19:51:31Z


+UCS_TEST_P(test_md_dmabuf, rocm_dmabuf_disable, "ROCM_COPY_DMABUF~=no")
+{
+    if (GetParam().md_name != "rocm_cpy") {


minor: this guard is dead code — "ROCM_COPY_DMABUF~=no" uses SKIP_IF_NOT_EXIST, so any MD whose config table has no ROCM_COPY_DMABUF field is already skipped in init() before the body runs. pls drop the md_name check.

svc-nvidia-pr-review · 2026-09-10T19:51:31Z

+    uint64_t orig_mask     = mem_attr_p->field_mask;
+    ucs_status_t status;
+
+    if (!md->have_dmabuf) {


maybe we can drop this wrapper? uct_cuda_copy_md_mem_query() keeps exporting the fd with CUDA_COPY_DMABUF=no and relies on dmabuf_mem_types == 0 to keep consumers from using it, so nothing in ucp reads these fields for rocm_cpy either. As written, rocm_cpy is the only MD that also mutates and restores the caller's field_mask, and the only thing that observes the difference is the new rocm_dmabuf_disable test.

svc-nvidia-pr-review · 2026-09-10T19:51:32Z

Residual coverage gap: mem_query_dmabuf becomes newly active for rocm_cpy because of the default flip, and it asserts EXPECT_EQ(mem_buf.ptr(), mem_attr.base_address) / EXPECT_EQ(offset, mem_attr.dmabuf_offset). That requires HSA's agentBaseAddress to equal the hipMalloc pointer (i.e. no ROCclr slab sub-allocation). Same coupling in mem_type_detect_mds, where EXPECT_GE(alloc_length, buffer_size) was relaxed but EXPECT_EQ(address, base_address) was kept. Only an AMD GPU run distinguishes page-rounding from sub-allocation.

Residual coverage gap: the ROCm CI containers (ubuntu2004_rocm_5_4_0, ubuntu2204_rocm_6_0_0) are build-only, so none of the above runs in this pipeline; a gtest.uct.test_md* run on real hardware is needed.

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

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.

…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-11T17:15:37Z

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

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 6812eeb to 7e3dddc Compare September 10, 2026 19:41

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

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

29cb04c

…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 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

	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);

	EXPECT_EQ(0ul, mem_attr.dmabuf_offset);
	EXPECT_EQ(0ul, mem_attr.dmabuf_offset);
	}

-            if ((alloc_mem_type == UCS_MEMORY_TYPE_CUDA) ||
+            if (alloc_mem_type == UCS_MEMORY_TYPE_ROCM) {
+                /* ROCm HSA may round sizeInBytes up to a page */
+                EXPECT_GE(mem_attr.alloc_length, buffer_size);
+                EXPECT_EQ(address, mem_attr.base_address);
+            } else if ((alloc_mem_type == UCS_MEMORY_TYPE_CUDA) ||
+                (alloc_mem_type == UCS_MEMORY_TYPE_ZE_HOST) ||
+                (alloc_mem_type == UCS_MEMORY_TYPE_ZE_DEVICE) ||
+                (alloc_mem_type == UCS_MEMORY_TYPE_ZE_MANAGED)) {
+                EXPECT_EQ(buffer_size, mem_attr.alloc_length);
+                EXPECT_EQ(address, mem_attr.base_address);

		@@ -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,

Uh oh!

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!