1.69 feature merge into candidate by AtlantaPepsi · Pull Request #331 · ROCm/TransferBench

AtlantaPepsi · 2026-07-08T20:47:57Z

Motivation

Given recent fixes that became 1.68 release, we are promoting features originally planned for 1.68 to 1.69.
Rerouting changes on candidate-1.68to the candidate branch.

Technical Details

IB verbs dynamic loading, see Dynamic Loading IB verbs #253
TDM integration into TransferBench, including @alex-breslow-amd 's initial prototype Add Logic for Using TDM and Async Loads and Stores #322 plus subsequent clean ups, pending further testing and changes.

Test Plan

ibverbs dynamic loading have been tested fully, including packaging on no-RDMA platform and unpackaging+dynamic loading on IB/RoCE platforms.

TDM's initial results were tested on gfx1250 platforms. Subsequent changes only reports no issues from compilers and linters, testing TODO.

Test Result

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

Copilot

Pull request overview

This PR promotes features planned for 1.68 into the 1.69 candidate by (1) switching RDMA/NIC support to runtime dynamic loading of libibverbs, and (2) integrating the gfx1250 TDM copy path as a new GPU executor option.

Changes:

Add minimal ibverbs ABI headers + dlopen/dlsym-based loader, and rework NIC availability checks to be runtime-based.
Add TDM executor plumbing (config/env/help/topology) and a new tdmCopy.h device API for gfx1250.
Update build systems (Makefile/CMake) to drop build-host ibverbs dependencies and link against libdl instead.

Reviewed changes

Copilot reviewed 12 out of 12 changed files in this pull request and generated 6 comments.

Show a summary per file

File	Description
third-party/ibverbs/IbvHeader.hpp	Adds minimal ABI-compatible ibverbs types and inline dispatch helpers.
third-party/ibverbs/IbvDynLoad.hpp	Implements runtime dlopen/dlsym loading and capability probing for libibverbs.
src/header/TransferBench.hpp	Integrates runtime NIC gating, adds TDM/ALS executor types, implements TDM executor execution path.
src/header/tdmCopy.h	Adds device-side TDM copy API with compile-time availability gating for gfx1250.
src/client/Utilities.hpp	Adds string mappings for new executor types (TDM/ALS).
src/client/Topology.hpp	Makes NIC topology output conditional on runtime ibverbs availability.
src/client/Presets/Help.hpp	Updates user help text for new executors and syntax.
src/client/EnvVars.hpp	Adds TDM environment variables and updates NIC env var printing to be runtime-gated.
src/client/Client.cpp	Reports NIC support based on runtime ibverbs availability.
Makefile	Removes build-time ibverbs detection/linking; adds `-ldl` and includes third-party ibverbs headers.
CMakeLists.txt	Removes ibverbs/dmabuf build options and links against `${CMAKE_DL_LIBS}`; adds include path for ibverbs headers.
build_packages_local.sh	Drops CMake flags related to NIC/DMA-BUF build-time toggles.

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

alex-breslow-amd · 2026-07-08T21:38:16Z

+      tdm::tdmCopy(dst, src, sizeBytes, shmem, ldsBytes);
+      if (++subIterations == numSubIterations) break;
+    }
+


After the while due to the semantics of how waiting is done with the tensor wait API that calls __builtin_amdgcn_tensorcnt(0), it can happen that the first warp completes all of its works before the other warps do and can enter the if(shouldRecordTiming) block. However, this is invalid because not all other warps have completed their work.

What I would suggest is one of two options:

Change the tensor wait API call to support a scope (warp or block). If block, call __syncthreads after the __builtin_amdgcn_tensorcnt(0) call.

Implement it directly in the kernel.

__builtin_amdgcn_tensorcnt(0);
__syncthreads();

// Option 1) Release fence at system scope on the sender
__builtin_amdgcn_tensorcnt(0); // a)
__syncthreads(); // b)
__builtin_amdgcn_fence(__ATOMIC_RELEASE, "");

a) and b) can just be your implementation of block scope sync for TDM

After this, you need one of two things. Either __threadfence_system or you need to set the scope bits in the actual intrinsic for the TDM loads and stores to be system scope, and then you don't need this. The block level waiting is enough.

In general TDM with a cache policy of 0, is workgroup processor scope with regular temporal behavior. This means that the writes made by TDM to the remote GPU's memory are not guaranteed to make it unless you have the explicit release fence mentioned above, which causes the L2 cache writeback, which pushes the data out to the remote GPU.

However, if the writes made by TDM are marked as system scope by properly setting the cp field, then the writes will write through to the remote memory on the remote GPU.

TransferBench historically has relied on the cache flush that occurs when a kernel ends for correctness. Recent changes to TransferBench fix this once I brought that up. You'll notice that the new in kernel timestamping logic now has a __threadfence_system before it in the gfx kernel.

This isn't quite true. This was just a miss in v1.67. Versions before that had the threadfence_system.

alex-breslow-amd · 2026-07-08T21:40:22Z

+        u32x4{}, u32x4{}, u32x8{},
+        /*cpol=*/0);
+}
+__device__ inline void waitTensor0() { __builtin_amdgcn_s_wait_tensorcnt(0); }


It's fine to keep this warp level for most of the data copies because for performance we actually want to have warps get out of sync and run ahead if they can. However, at the end just before we collect timing, we need to verify that the writes made by TDM have completed. See my prior comment.

Synced internally and decided not to move forward with pipelining so this is enough.

alex-breslow-amd · 2026-07-08T21:43:24Z

+        g0.m_bitfield,               // D0  addresses
+        g1.m_bitfield,               // D1  2D shape
+        u32x4{}, u32x4{}, u32x8{},   // D2/D3/D4  higher dims: unused (zero)
+        /*cpol=*/0);


Set cache policy to system cpol field. Example here on how to compute the value: ROCm/rocm-systems#7019

Follow up with me if you have questions. Otherwise, you will need a __threadfence_system for in-kernel timing to be accurate.

alex-breslow-amd · 2026-07-08T21:48:41Z

+    bool const shouldRecordTiming = (threadIdx.x == 0);
+    if (shouldRecordTiming) startCycle = GetTimestamp();
+
+    extern __shared__ __align__(128) float shmem[];


Don't need the align(128), I added this but it turns out that it's not strictly necessary. Up to you to keep.

Copilot

Pull request overview

Copilot reviewed 12 out of 12 changed files in this pull request and generated 5 comments.

* Put in stubs * Add kernel and fix up * Finalize tunability of threads per block and number of blocks. * Implement pipelining * Add debugging logic * Fix segfault * Disable debug prints * Add missing timing information * Add TDM_BLOCK_SIZE, TDM_MAX_LDS_BYTES, and TDM_PIPELINED env vars to control TDM kernel execution * Not functioning but proof of concept async load/store implementation * Fix async load store kernel MVP * Potential fix for pull request finding Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> --------- Co-authored-by: Tim <43156029+AtlantaPepsi@users.noreply.github.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>

Copilot

Pull request overview

Copilot reviewed 25 out of 38 changed files in this pull request and generated 6 comments.

+                                      hipEvent_t    const  stopEvent,
+                                      int           const  numSubExecs,
+                                      TransferResources&   rss)
+  {


Copilot

Pull request overview

Copilot reviewed 12 out of 12 changed files in this pull request and generated 4 comments.

Copilot

Pull request overview

Copilot reviewed 12 out of 12 changed files in this pull request and generated 2 comments.

+__device__ inline void warpVecCopy(const uint8_t* s, uint8_t* d, size_t n,
+                                   uint32_t warpThread, uint32_t warpThreads) {
+    size_t nd = n >> 2;
+    const uint32_t* s32 = reinterpret_cast<const uint32_t*>(s);
+    uint32_t*       d32 = reinterpret_cast<uint32_t*>(d);
+    for (size_t i = warpThread; i < nd; i += warpThreads) d32[i] = s32[i];
+    uint32_t rem = static_cast<uint32_t>(n & 3u);
+    if (rem && warpThread == 0)
+        for (uint32_t b = 0; b < rem; ++b) d[nd * 4 + b] = s[nd * 4 + b];
+}


+#define IBV_FN(name, rettype, arglist) rettype(*name)arglist = nullptr;
+
+namespace {
+
+IBV_FN(ibv_alloc_pd, ibv_pd*, (ibv_context*))
+IBV_FN(ibv_close_device, int, (ibv_context*))
+IBV_FN(ibv_create_cq, ibv_cq*, (ibv_context*, int, void*, ibv_comp_channel*, int))
+IBV_FN(ibv_create_qp, ibv_qp*, (ibv_pd*, ibv_qp_init_attr*))
+IBV_FN(ibv_dealloc_pd, int, (ibv_pd*))
+IBV_FN(ibv_dereg_mr, int, (ibv_mr*))
+IBV_FN(ibv_destroy_cq, int, (ibv_cq*))
+IBV_FN(ibv_destroy_qp, int, (ibv_qp*))
+IBV_FN(ibv_free_device_list, void, (ibv_device**))
+IBV_FN(ibv_get_device_list, ibv_device**, (int*))
+IBV_FN(ibv_get_device_name, const char*, (ibv_device*))
+IBV_FN(ibv_modify_qp, int, (ibv_qp*, ibv_qp_attr*, int))
+IBV_FN(ibv_open_device, ibv_context*, (ibv_device*))
+IBV_FN(ibv_query_device, int, (ibv_context*, ibv_device_attr*))
+IBV_FN(ibv_query_gid, int, (ibv_context*, uint8_t, int, ibv_gid*))
+IBV_FN(ibv_query_port, int, (ibv_context*, uint8_t, ibv_port_attr*))
+// `ibv_reg_dmabuf_mr` is always declared; whether the underlying symbol
+// actually exists in the loaded libibverbs is decided at runtime by tryLoad().
+IBV_FN(ibv_reg_dmabuf_mr, ibv_mr*, (ibv_pd*, uint64_t, size_t, uint64_t, int, int))
+IBV_FN(ibv_reg_mr, ibv_mr*, (ibv_pd*, void*, size_t, int))
+}


AtlantaPepsi requested a review from a team as a code owner July 8, 2026 20:47

Copilot AI review requested due to automatic review settings July 8, 2026 20:47

Copilot started reviewing on behalf of AtlantaPepsi July 8, 2026 20:49 View session

Copilot AI reviewed Jul 8, 2026

View reviewed changes

Comment thread src/header/TransferBench.hpp Outdated

Comment thread src/client/Presets/Help.hpp Outdated

Comment thread src/client/Presets/Help.hpp

Comment thread src/header/TransferBench.hpp

Comment thread third-party/ibverbs/IbvDynLoad.hpp

Comment thread src/header/tdmCopy.h Outdated

alex-breslow-amd reviewed Jul 8, 2026

View reviewed changes

Copilot AI review requested due to automatic review settings July 13, 2026 20:45

Copilot started reviewing on behalf of AtlantaPepsi July 13, 2026 20:46 View session

Copilot AI reviewed Jul 13, 2026

View reviewed changes

Comment thread src/client/Presets/Help.hpp

Comment thread src/client/Presets/Help.hpp

Comment thread src/client/Presets/Help.hpp

Comment thread src/client/EnvVars.hpp

Comment thread third-party/ibverbs/IbvDynLoad.hpp

alex-breslow-amd and others added 4 commits July 13, 2026 18:02

Dynamic Loading IB verbs (#253)

9eede61

TDM productionization in TB

d7be156

temporary fixes

33d0bf7

AtlantaPepsi force-pushed the candidate-1.68 branch from be5f315 to 33d0bf7 Compare July 13, 2026 23:02

Copilot AI review requested due to automatic review settings July 13, 2026 23:54

AtlantaPepsi requested a review from a team as a code owner July 13, 2026 23:54

Copilot started reviewing on behalf of AtlantaPepsi July 13, 2026 23:55 View session

Copilot AI reviewed Jul 13, 2026

View reviewed changes

Copilot AI review requested due to automatic review settings July 14, 2026 00:22

AtlantaPepsi force-pushed the candidate-1.68 branch from bd22f40 to 33d0bf7 Compare July 14, 2026 00:22

Copilot started reviewing on behalf of AtlantaPepsi July 14, 2026 00:23 View session

Copilot AI reviewed Jul 14, 2026

View reviewed changes

Comment thread src/header/tdmCopy.h Outdated

Comment thread src/client/EnvVars.hpp

Comment thread src/client/Presets/Help.hpp

Comment thread src/client/Presets/Help.hpp

gate runtime check for absent TDM header

d2aff82

Copilot AI review requested due to automatic review settings July 14, 2026 02:21

Copilot started reviewing on behalf of AtlantaPepsi July 14, 2026 02:22 View session

Copilot AI reviewed Jul 14, 2026

View reviewed changes

Uh oh!

Conversation

AtlantaPepsi commented Jul 8, 2026

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

alex-breslow-amd Jul 8, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

gilbertlee-amd Jul 14, 2026

Choose a reason for hiding this comment

Uh oh!

alex-breslow-amd Jul 8, 2026

Choose a reason for hiding this comment

Uh oh!

alex-breslow-amd Jul 8, 2026

Choose a reason for hiding this comment

Uh oh!

alex-breslow-amd Jul 8, 2026

Choose a reason for hiding this comment

Uh oh!

alex-breslow-amd Jul 8, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

alex-breslow-amd Jul 8, 2026 •

edited

Loading