Add Logic for Using TDM and Async Loads and Stores by alex-breslow-amd · Pull Request #322 · ROCm/TransferBench

alex-breslow-amd · 2026-06-11T06:41:01Z

Motivation

Run fast on gfx1250

Technical Details

Adds simple copy kernels using TDM and async/loads and stores

Test Plan

Code has correctness testing, but I need to add more corner case testing

Test Result

Tests pass on a hardware simulator for both variants:
./TransferBench cmdline 1024K "1 4 G0 T0 G0"
./TransferBench cmdline 1024K "1 4 G0 L0 G0"

Submission Checklist

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

…control TDM kernel execution

Co-authored-by: Cursor <cursoragent@cursor.com> # Conflicts: # src/header/TransferBench.hpp

alex-breslow-amd · 2026-06-11T06:42:10Z

 ROCM_PATH ?= /opt/rocm
 CUDA_PATH ?= /usr/local/cuda
 MPI_PATH  ?= /usr/local/openmpi
+# pip ROCm wheels ship bin/amdclang++ but often omit bin/amdllvm (which that stub execs).


Logic for running building and running this in a venv on WSL

alex-breslow-amd · 2026-06-12T18:52:01Z

Discussed with @AtlantaPepsi and @gilbertlee-amd. The plan is to merge as is. There's some cleanup that needs to be performed however before the candidate branch can be released:

Add proper gating for gfx1250
Handle case where we are on an NVIDIA platform
Add tunability for cache hints to TDM/async load store implementation (Let's sync on this @AtlantaPepsi)

Regarding async loads/stores, the current code is just proof of concept. The waitcnt placement and unrolling isn't optimized. Feel free to remove if you think it's best. TBH, I think this code could be rolled into the existing TransferBench kernels. I was thinking of working on that (actually what I started on before moving to TDM). TDM feels sufficiently different where this likely isn't warranted. I concur with Gilbert on that one.

alex-breslow-amd · 2026-06-12T18:56:32Z

@@ -0,0 +1,355 @@
+/*


This file is a standard header but was missing from my installation of ROCm, not sure why. It can probably be removed. This is standard in CLR:

Copilot

Pull request overview

This PR adds two new GPU executor variants intended to improve copy performance on gfx1250 by using (1) TDM-based tensor load/store operations and (2) async global load/store-to-LDS memory ops, along with CLI/help/env-var wiring and a small Makefile improvement for ROCm wheel layouts.

Changes:

Add new executor types EXE_GPU_ASYNC_TENSOR and EXE_GPU_ASYNC_MEMOPS with corresponding dispatch, validation, topology naming, and client stringification.
Introduce TDM configuration (TdmOptions, TDM_* env vars) and implement new HIP kernels for the TDM tensor path and async load/store path.
Improve HIP compiler discovery in the Makefile for ROCm installations where bin/amdllvm is absent.

Reviewed changes

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

Show a summary per file

File	Description
`src/header/TransferBench.hpp`	Adds new executor types, TDM options, and implements/dispatches async GPU kernels (TDM + async LD/ST).
`src/header/tdm.h`	Adds gfx1250 TDM register layout/bitfield definitions used by the TDM kernel path.
`src/client/Utilities.hpp`	Extends executor-to-string mapping for the new executor types.
`src/client/Presets/Help.hpp`	Updates help text to document new executor letters and examples.
`src/client/EnvVars.hpp`	Adds env vars for TDM kernel tuning and wires them into `ConfigOptions`.
`Makefile`	Adjusts HIP compiler selection to support ROCm wheel layouts.

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

AtlantaPepsi · 2026-06-12T20:55:35Z

+    // Copy from global memory to LDS
+    for (size_t i = 0; i < elementsToProcess; i += elements_per_dwordx4_load_store) {
+      asyncLoadX4(srcPtr + i, shmemPtr + i);
+      __builtin_amdgcn_s_wait_asynccnt(0);
+    }
+
+    // Write back from LDS to global
+    for (size_t i = 0; i < elementsToProcess; i += elements_per_dwordx4_load_store){
+      asyncStoreX4(shmemPtr + i, dstPtr + i);
+      __builtin_amdgcn_s_wait_asynccnt(0);
+    }


async l/s will undergo further changes

AtlantaPepsi · 2026-06-12T20:54:27Z

+      if (rss.numBytes % sizeof(float) != 0)
+        return {ERR_FATAL, "Async tensor executor (TDM): numBytes (%zu) must be a multiple of %zu", rss.numBytes, sizeof(float)};


this will be elevated into TransfersHaveErrors guard

AtlantaPepsi · 2026-06-12T21:45:49Z

+      int maxShmem = 0;
+      ERR_CHECK(hipDeviceGetAttribute(&maxShmem, hipDeviceAttributeMaxSharedMemoryPerBlock, exeIndex));
+      maxShmem = std::min<int>(maxShmem, cfg.tdm.maxLDSBytes);
+      int kWavesPerWorkgroup = threads / warpSize;
+      int pipelineDepth = usePipelinedTensorOps ? 2 : 1;
+      int numFloatsPerTile = std::max<int>(1, maxShmem / (sizeof(float) *  kWavesPerWorkgroup * pipelineDepth));
+      unsigned int shmemBytes = numFloatsPerTile * kWavesPerWorkgroup * sizeof(float) * pipelineDepth;


to be further validated in next PR

Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>

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

alex-breslow-amd added 12 commits May 15, 2026 15:47

Put in stubs

4023f2e

Add kernel and fix up

f061148

Finalize tunability of threads per block and number of blocks.

14e8f62

Implement pipelining

37fe5c5

Add debugging logic

f4411a6

Fix segfault

120cf52

Disable debug prints

e97fa5a

Add missing timing information

cc4c22b

Add TDM_BLOCK_SIZE, TDM_MAX_LDS_BYTES, and TDM_PIPELINED env vars to …

014dce1

…control TDM kernel execution

Not functioning but proof of concept async load/store implementation

037d966

Fix async load store kernel MVP

c72fd7f

Merge remote-tracking branch 'upstream/candidate-1.68' into asyncOps

d06901c

Co-authored-by: Cursor <cursoragent@cursor.com> # Conflicts: # src/header/TransferBench.hpp

alex-breslow-amd commented Jun 11, 2026

View reviewed changes

alex-breslow-amd marked this pull request as ready for review June 12, 2026 18:48

alex-breslow-amd requested a review from a team as a code owner June 12, 2026 18:48

alex-breslow-amd requested review from AtlantaPepsi and Copilot June 12, 2026 18:55

Copilot started reviewing on behalf of alex-breslow-amd June 12, 2026 18:56 View session

alex-breslow-amd commented Jun 12, 2026

View reviewed changes

Copilot AI reviewed Jun 12, 2026

View reviewed changes

Potential fix for pull request finding

a25a01c

Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>

AtlantaPepsi merged commit a9b5eb2 into ROCm:candidate-1.68 Jun 12, 2026
1 check passed

AtlantaPepsi mentioned this pull request Jul 8, 2026

1.69 feature merge into candidate #331

Closed

1 task

Navigation Menu

Uh oh!

Add Logic for Using TDM and Async Loads and Stores - #322

Add Logic for Using TDM and Async Loads and Stores#322
AtlantaPepsi merged 13 commits into
ROCm:candidate-1.68from
alex-breslow-amd:asyncOps

alex-breslow-amd commented Jun 11, 2026

Uh oh!

alex-breslow-amd Jun 11, 2026

Uh oh!

alex-breslow-amd commented Jun 12, 2026 •

edited

Loading

Uh oh!

alex-breslow-amd Jun 12, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

AtlantaPepsi Jun 12, 2026

Uh oh!

AtlantaPepsi Jun 12, 2026

Uh oh!

AtlantaPepsi Jun 12, 2026

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

		if (rss.numBytes % sizeof(float) != 0)
		return {ERR_FATAL, "Async tensor executor (TDM): numBytes (%zu) must be a multiple of %zu", rss.numBytes, sizeof(float)};

Uh oh!

Conversation

alex-breslow-amd commented Jun 11, 2026

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Uh oh!

alex-breslow-amd Jun 11, 2026

Choose a reason for hiding this comment

Uh oh!

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

Uh oh!

Uh oh!

alex-breslow-amd Jun 12, 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

Reviewed changes

Uh oh!

AtlantaPepsi Jun 12, 2026

Choose a reason for hiding this comment

Uh oh!

AtlantaPepsi Jun 12, 2026

Choose a reason for hiding this comment

Uh oh!

AtlantaPepsi Jun 12, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

alex-breslow-amd commented Jun 12, 2026 •

edited

Loading