Async load store alternative path for TDM copy kernel#344
Open
alex-breslow-amd wants to merge 8 commits into
Open
Async load store alternative path for TDM copy kernel#344alex-breslow-amd wants to merge 8 commits into
alex-breslow-amd wants to merge 8 commits into
Conversation
… use the async copy backend for the TDM path
Contributor
There was a problem hiding this comment.
Pull request overview
This PR adds a new GPU executor path based on gfx1250 Tensor Data Mover (TDM) semantics, with a compile-time-selectable backend that can use either true TDM builtins or an async-load/store-to-LDS fallback while keeping an API compatible surface.
Changes:
- Introduces a new
EXE_GPU_TDMexecutor, wiring it through config validation, topology, execution, and client presets. - Adds new device-side copy libraries (
tdmCopy.h,asyncCopy.h) plus a shared cache-policy encoding header (cachePolicy.h). - Adds sweep-bound env vars (
SWEEP_MIN_POW2,SWEEP_MAX_POW2) and Makefile improvements for ROCm toolchain/header discovery.
Reviewed changes
Copilot reviewed 12 out of 12 changed files in this pull request and generated 8 comments.
Show a summary per file
| File | Description |
|---|---|
| src/header/TransferBench.hpp | Adds the TDM executor implementation and integrates it into config/executor orchestration. |
| src/header/tdmCopy.h | New device-side TDM copy API/implementation for gfx1250 (plus availability gates). |
| src/header/asyncCopy.h | New async-to/from-LDS copy backend with API parity to tdmCopy.h. |
| src/header/cachePolicy.h | Shared cache-policy encoding used by both TDM/async copy implementations. |
| src/client/Utilities.hpp | Adds stringification for the new TDM executor. |
| src/client/Presets/Help.hpp | Updates help output to document the new executor token (T). |
| src/client/Presets/AllToAllSweep.hpp | Adds USE_TDM_EXEC and adapts sweep dimensions for TDM vs GFX. |
| src/client/Presets/AllToAll.hpp | Adds USE_TDM_EXEC support and ensures it’s mutually exclusive with DMA mode. |
| src/client/EnvVars.hpp | Adds new sweep-bound and TDM env vars and maps them into ConfigOptions. |
| src/client/Client.cpp | Uses new sweep-bound env vars when generating byte sweep ranges. |
| Makefile | Improves ROCm compiler detection and adds optional vendored gfx1250 descriptor include path + EXTRA_CXXFLAGS. |
| CHANGELOG.md | Documents the new version items (sweep bounds + TDM executor). |
Comments suppressed due to low confidence (2)
src/header/TransferBench.hpp:2349
- This subexecutor "data size too small" warning is now applied to all non-DMA executors, but NIC subexecutors are queue pairs (rss.qpCount is set directly from t.numSubExecs) and are not reduced based on cfg.data.blockBytes. This makes the warning incorrect/misleading for NIC transfers.
if (t.exeDevice.exeType != EXE_GPU_DMA) { // GPU DMA-Executor ignores subexecutors
size_t const N = t.numBytes / sizeof(float);
int const targetMultiple = cfg.data.blockBytes / sizeof(float);
int const maxSubExecToUse = std::min((size_t)(N + targetMultiple - 1) / targetMultiple,
(size_t)t.numSubExecs);
if (maxSubExecToUse < t.numSubExecs)
errors.push_back({ERR_WARN,
"Transfer %d data size is too small - will only use %d of %d subexecutors due to blockBytes of %d",
i, maxSubExecToUse, t.numSubExecs, cfg.data.blockBytes});
src/header/TransferBench.hpp:4524
- Threadblock ordering for EXE_GPU_TDM currently uses cfg.gfx.blockOrder, so the newly introduced cfg.tdm.blockOrder / TDM_BLOCK_ORDER setting has no effect for the TDM executor.
int transferOffset = 0;
if (cfg.general.useMultiStream || cfg.gfx.blockOrder == 0) {
// Threadblocks are ordered sequentially one transfer at a time
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
Comment on lines
+2199
to
+2201
| errors.push_back({ERR_FATAL, | ||
| "Unable to query max amount of shared memory per block on GPU %%d\n", i}); | ||
| } |
Comment on lines
+2176
to
+2182
| // Check TDM options | ||
| #if !defined(__NVCC__) | ||
| if (cfg.tdm.blockSize <= 0 || cfg.tdm.blockSize % 32 || cfg.tdm.blockSize > MAX_BLOCKSIZE) | ||
| errors.push_back({ERR_FATAL, | ||
| "[tdm.blockSize] must be a positive multiple of 32 less than or equal to %d", | ||
| MAX_BLOCKSIZE}); | ||
|
|
| } | ||
| } | ||
|
|
||
| // Determine how much shared memory to use for TDS |
Comment on lines
+344
to
+353
| __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]; | ||
| } |
Comment on lines
173
to
175
| // Run the specified numbers of bytes otherwise generate a range of values | ||
| for (size_t bytes = (1<<10); bytes <= (1<<29); bytes *= 2) { | ||
| for (size_t bytes = (1ULL<<ev.sweepMinPow2); bytes <= (1ULL<<ev.sweepMaxPow2); bytes *= 2) { | ||
| size_t deltaBytes = std::max(1UL, bytes / ev.samplingFactor); |
Comment on lines
+46
to
+47
| printf("# 5) Batched-DMA Batch item (Must have single SRC, at least one DST)\n"); | ||
| printf("# 6) TDM GPU threadblock/Compute Unit (CU) (Requires hardware support for Tensor Data Mover\n"); |
Comment on lines
+87
to
+88
| int sweepMaxPow2; // Maximum power of two to sweep up to when whnumber of bytes to transfer is set to 0 | ||
| int sweepMinPow2; // Minimum power of two to sweep up from when number of bytes to transfer is set to 0 |
Comment on lines
+6
to
+8
| ## v1.70.00 | ||
| ## Added | ||
| - Added support for SWEEP_MIN_POW2 and SWEEP_MAX_POW2 to set sweep bounds when bytes to transfer is 0 |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
Proof of concept with API compatibility with tdmCopy.h using async loads and stores
Technical Details
Test Plan
Test Result
Submission Checklist