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.

* 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

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.


+  struct TdmOptions
+  {
+    int blockSize    = 256; ///< Size of each threadblock (must be multiple of 64)


+  printf("## Single GPU-executed Transfer between GPUs 0 and 1 using 1 CU for async load/store path\n");
+  printf("1 1 (G0->L0->G1)\n");


+  printf("#   6) Batched-DMA         Batch item                           (Must have single SRC, at least one DST)\n");
+  printf("#   7) TDM                 GPU threadblock/Compute Unit (CU)\n");
+  printf("#   8) Async Load/Store    GPU threadblock/Compute Unit (CU)\n");


          TransferResult& tfrResult      = results.tfrResults[transferIdx];
          tfrResult.exeDevice            = exeDevice;
-#ifdef NIC_EXEC_ENABLED
          tfrResult.exeDstDevice         = {exeDevice.exeType, rss.dstNicIndex};
-#else
-          tfrResult.exeDstDevice         = exeDevice;
-#endif
          tfrResult.numBytes             = rss.numBytes;


+#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))


+///
+///        This is introduced via __device__ level memcpy-like API which can be
+///        either blocking or asynchronous, and utilize all warps, or a team of
+///        congtigous warps, to allow for other warps to do other tasks


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.

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.

alex-breslow-amd and others added 3 commits June 12, 2026 17:05

Dynamic Loading IB verbs (#253)

cca4258

TDM productionization in TB

d5fd67b

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

alex-breslow-amd reviewed Jul 8, 2026

View reviewed changes

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

1.69 feature merge into candidate#331

1.69 feature merge into candidate#331
AtlantaPepsi wants to merge 3 commits into
candidatefrom
candidate-1.68

AtlantaPepsi commented Jul 8, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

alex-breslow-amd Jul 8, 2026 •

edited

Loading

Uh oh!

alex-breslow-amd Jul 8, 2026

Uh oh!

alex-breslow-amd Jul 8, 2026

Uh oh!

alex-breslow-amd Jul 8, 2026

Uh oh!

alex-breslow-amd Jul 8, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

		printf("## Single GPU-executed Transfer between GPUs 0 and 1 using 1 CU for async load/store path\n");
		printf("1 1 (G0->L0->G1)\n");

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!

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!

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!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

alex-breslow-amd Jul 8, 2026 •

edited

Loading