TransferBench v1.71.00 by AtlantaPepsi · Pull Request #359 · ROCm/TransferBench

AtlantaPepsi · 2026-09-23T15:27:20Z

Motivation

Pingpong latency testing integration and presets
NIC executor check against max message size
fixing logical CU ID report
inclusion of gfx1250-strict target

Technical Details

Test Plan

Test Result

To be filled

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

Copilot

Copilot review overview

🟡 Changes recommended

Unresolved critical and moderate correctness, portability, parsing, and build-target issues remain.

Get a fresh assessment by requesting another Copilot review.

Review effort: Lite
Findings: 4 · 4

Open (8)

Restrict GetXccId CDNA5 instruction to CDNA5 targets · New
Include pong subindex when enabling GFX subindices · New
Emit advanced transfer fields before the pong token · New
Align TDM device guard with supported architectures · New
Avoid indexing stale or empty transfer results after failure · New
Validate NIC message size for the selected port · New
Fix pingpong timing scaling by subiterations · New
Reject multi-device pingpong memory halves · New

What changed in this PR

Adds pingpong latency testing, NIC message-size validation, logical CU reporting fixes, and gfx1250-strict support.

Changes:

Adds pingpong parsing, execution, timing, and latency presets.
Adds NIC max_msg_sz validation and reporting.
Updates GPU architecture and TDM handling.

File	Review summary
`src/header/TransferBench.hpp`	Critical and moderate issues in XCC handling, pingpong validation, NIC limits, subindex selection, dump parsing, and timing scaling.
`src/header/tdmCopy.h`	Critical architecture guard mismatch enables unsupported TDM targets.
`src/client/Utilities.hpp`	Reviewed result and topology utilities.
`src/client/Topology.hpp`	Reviewed NIC message-size reporting.
`src/client/Presets/Presets.hpp`	Reviewed latency preset registration.
`src/client/Presets/Latency.hpp`	Moderate issues with failed-run result handling and diagonal pair measurement.
`src/client/EnvVars.hpp`	Reviewed pingpong configuration variables.
`src/client/Client.cpp`	Reviewed pingpong transfer display updates.
`docs/install/build_from_source.rst`	Reviewed strict GPU target documentation.
`CMakeLists.txt`	Moderate issue: package build targets omit `gfx1250-strict`.

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

@@ -788,7 +855,7 @@ __device__ __forceinline__ uint32_t GetXccId()
  uint32_t xccReg = 0;
  asm volatile("s_getreg_b32 %0, hwreg(HW_REG_XCC_ID)" : "=s"(xccReg));
  return xccReg & 0xF;
-#elif defined(__gfx1250__)
+#elif defined(__GFX12__)


+        exeInfo.totalSubExecs += t.numSubExecs;
+      } else {
+        exeInfo.totalPingpong ++;
+      }
      exeInfo.useSubIndices |= (t.exeSubIndex != -1 || (t.exeDevice.exeType == EXE_GPU_GFX && !cfg.gfx.prefXccTable.empty()));


+        fprintf(dumpCfgFile, " +");
+        if (t.numLaps != 1)
+          fprintf(dumpCfgFile, "%d", t.numLaps);
+        fprintf(dumpCfgFile, " ");
+        printMem(t.srcs[1]);
+        fprintf(dumpCfgFile, "->");
+        printExe(t.exeDevicePong, t.exeSubIndexPong, t.exeSubSlotPong);
+        fprintf(dumpCfgFile, "->");
+        printMem(t.dsts[1]);
+        fprintf(dumpCfgFile, " %d %lu)", t.numSubExecs, t.numBytes);


+#  if defined(__GFX12__) && \
      __has_builtin(__builtin_amdgcn_tensor_load_to_lds) && \
      TDM_TOOLCHAIN_AVAILABLE


+    if (!TransferBench::RunTransfers(cfg, transfers, results))
+      Utils::PrintErrors(results.errResults);
+
+    for (size_t k = 0; k < pairs.size(); k++)
+      latencyUs[pairs[k].first][pairs[k].second] = results.tfrResults[k].avgDurationMsec * 1000.0;


+                  // Any active port may be selected for QP setup via cfg.nic.ibPort, which is
+                  // not visible here, so keep the most restrictive limit across all of them.
+                  if (portAttr.max_msg_sz &&
+                      (!ibvDevice.maxMsgSize || portAttr.max_msg_sz < ibvDevice.maxMsgSize))
+                    ibvDevice.maxMsgSize = portAttr.max_msg_sz;


+      if (numCopyTasks == 0) {
+        exeInfo.totalDurationMsec += hasPingpong
+          ? std::chrono::duration_cast<std::chrono::duration<double>>(cpuDelta).count() * scale
+          : 0.0;


+            auto singleMemOrNull = [](vector<MemDevice> const& mems) {
+              if (mems.empty()) return MemDevice{MEM_NULL, 0, 0};
+              return mems[0];
+            };


Copilot

Copilot review overview

🟡 Changes recommended

Unresolved correctness issues remain in ping-pong execution, validation, parsing, and target gating.

Get a fresh assessment by requesting another Copilot review.

Review effort: Lite
Findings: 5 · 4

Open (9)

Run unpinned pingpong parameters on only one XCC · New
Align TDM device guard with supported architectures
Emit advanced transfer fields before the pong token
Include pong subindex when enabling GFX subindices
Restrict GetXccId CDNA5 instruction to CDNA5 targets
Reject multi-device pingpong memory halves
Fix pingpong timing scaling by subiterations
Validate NIC message size for the selected port
Avoid indexing stale or empty transfer results after failure

+    dim3 const gridSize(xccDim, numPingpong, 1);
+    dim3 const blockSize(1);


Copilot

Copilot review overview

🟡 Changes recommended

Unresolved critical and moderate issues remain in pingpong execution, validation, result handling, and architecture gating.

Get a fresh assessment by requesting another Copilot review.

Review effort: Lite
Findings: 7 · 5

Open (12)

Reject pingpong subindices below -1 · New
Prevent persistent pingpong dependency cycles from hanging · New
Run unpinned pingpong parameters on only one XCC
Align TDM device guard with supported architectures
Emit advanced transfer fields before the pong token
Include pong subindex when enabling GFX subindices
Restrict GetXccId CDNA5 instruction to CDNA5 targets
Allow default numSubExecs for pingpong transfers · New
Reject multi-device pingpong memory halves
Fix pingpong timing scaling by subiterations
Validate NIC message size for the selected port
Avoid indexing stale or empty transfer results after failure

+#else
+            useSubIndexCount[exe]++;
+            int numSubIndices = GetNumExecutorSubIndices(exe);
+            if (subIndex >= numSubIndices) {


+    dim3 const gridSize(xccDim, numPingpong, 1);
+    dim3 const blockSize(1);


      if (t.numSubExecs <= 0)
        errors.push_back({ERR_FATAL, "Transfer %d: # of subexecutors must be positive", i});
-      else
+      else if (isPingpong) {
+        if (t.numSubExecs != 1)
+          errors.push_back({ERR_WARN,


AtlantaPepsi added 8 commits September 22, 2026 09:39

ibv port maxMsgSize check

3d321ae

fixing logical CU ID report

98a43f5

gfx1250strict

bed30a5

pingpong latency test

85d92c1

adjusting output format

6f775df

relaxed atomics for pingpong kernel

93227a0

refactor GFX executor timing

6954cdd

cuda fix and cross pod fix

81865fa

Copilot AI lite review requested due to automatic review settings September 23, 2026 15:27

AtlantaPepsi requested review from a team as code owners September 23, 2026 15:27

Copilot started reviewing on behalf of AtlantaPepsi September 23, 2026 15:28 View session

Copilot AI reviewed Sep 23, 2026

View reviewed changes

Copilot AI review requested due to automatic review settings September 23, 2026 15:36

Copilot started reviewing on behalf of AtlantaPepsi September 23, 2026 15:37 View session

Copilot AI reviewed Sep 23, 2026

View reviewed changes

Comment thread src/header/TransferBench.hpp

Comment on lines +6660 to +6661

dim3 const gridSize(xccDim, numPingpong, 1);

dim3 const blockSize(1);

AtlantaPepsi changed the title ~~Candidate 1.71~~ TransferBench v1.71.00 Sep 23, 2026

AtlantaPepsi added 2 commits September 23, 2026 16:05

Latency Presets

e5166b6

TransferBench v1.71.00

07e9f67

Copilot AI review requested due to automatic review settings September 23, 2026 21:06

AtlantaPepsi force-pushed the candidate-1.71 branch from e37e7de to 07e9f67 Compare September 23, 2026 21:06

Copilot started reviewing on behalf of AtlantaPepsi September 23, 2026 21:07 View session

Copilot AI reviewed Sep 23, 2026

View reviewed changes

Navigation Menu

TransferBench v1.71.00 - #359

TransferBench v1.71.00#359
AtlantaPepsi wants to merge 10 commits into
developfrom
candidate-1.71

AtlantaPepsi commented Sep 23, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

Copilot AI left a comment

Uh oh!

Copilot AI left a comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

		dim3 const gridSize(xccDim, numPingpong, 1);
		dim3 const blockSize(1);

Conversation

AtlantaPepsi commented Sep 23, 2026

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Copilot review overview

🟡 Changes recommended

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Copilot review overview

🟡 Changes recommended

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Copilot review overview

🟡 Changes recommended

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants