From c9922ed399034ca885789b62de52e5a28698cb38 Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Wed, 29 Jul 2026 21:38:35 +0000 Subject: [PATCH 1/7] TDM reduceCopy scaffolding --- src/client/EnvVars.hpp | 8 ++ src/header/TransferBench.hpp | 168 +++++++++++++++++++++++++++++++++-- 2 files changed, 167 insertions(+), 9 deletions(-) diff --git a/src/client/EnvVars.hpp b/src/client/EnvVars.hpp index 0d9b4cbb..7d50b390 100644 --- a/src/client/EnvVars.hpp +++ b/src/client/EnvVars.hpp @@ -129,6 +129,7 @@ class EnvVars // TDM options int tdmBlockOrder; // How threadblocks for multiple Transfers are ordered 0=sequential 1=interleaved int tdmBlockSize; // Size of each threadblock for TDM kernels (must be multiple of 32) + int tdmKernel; // TDM Kernel to use (-1=auto, 0=copy-only, 1=reduce) int tdmLdsBytes; // Size of LDS (shared memory) bytes per threadblock for TDM kernels (0 = use device max) // Developer features @@ -184,6 +185,7 @@ class EnvVars sweepMinPow2 = GetEnvVar("SWEEP_MIN_POW2" , 10); tdmBlockOrder = GetEnvVar("TDM_BLOCK_ORDER" , 0); tdmBlockSize = GetEnvVar("TDM_BLOCK_SIZE" , 256); + tdmKernel = GetEnvVar("TDM_KERNEL" , 0); tdmLdsBytes = GetEnvVar("TDM_LDS_BYTES" , 0); useHipEvents = GetEnvVar("USE_HIP_EVENTS" , 1); useHsaDma = GetEnvVar("USE_HSA_DMA" , 0); @@ -406,6 +408,7 @@ class EnvVars printf(" SWEEP_MIN_POW2 - When 0 is specified for data size, this is the starting power of two exponent\n"); printf(" TDM_BLOCK_ORDER - How blocks for TDM transfers are ordered. 0=sequential, 1=interleaved\n"); printf(" TDM_BLOCK_SIZE - # of threads per threadblock for TDM (async tensor) kernels (Must be multiple of 32)\n"); + printf(" TDM_KERNEL - -1=auto, 0=force TdmCopyKernel, 1=force TdmReduceKernel (may error if ineligible)\n"); printf(" TDM_LDS_BYTES - Amount of LDS bytes to allocate per workgroup for TDM kernels (0 = device max; K/M/G suffixes accepted)\n"); printf(" USE_HIP_EVENTS - Use HIP events for GFX executor timing\n"); printf(" USE_HIP_EVENTS - Use HIP events for GFX/DMA/TDM executor timing (0=CPU wall-clock)\n"); @@ -554,6 +557,10 @@ class EnvVars Print("TDM_BLOCK_ORDER", tdmBlockOrder, "TDM Thread block ordering: %s", tdmBlockOrder == 0 ? "Sequential" : "Interleaved"); Print("TDM_BLOCK_SIZE", tdmBlockSize, "TDM threadblock size of %d", tdmBlockSize); + Print("TDM_KERNEL", tdmKernel, + "%s", tdmKernel == -1 ? "auto" : + tdmKernel == 0 ? "force TdmCopyKernel" : + tdmKernel == 1 ? "force TdmReduceKernel" : "unknown"); Print("TDM_LDS_BYTES", tdmLdsBytes, "%s", tdmLdsBytes == 0 ? "Using device max LDS bytes per workgroup" : (std::string("Setting LDS to ") + std::to_string(tdmLdsBytes) + " bytes per workgroup").c_str()); @@ -766,6 +773,7 @@ class EnvVars cfg.tdm.blockOrder = tdmBlockOrder; cfg.tdm.blockSize = tdmBlockSize; + cfg.tdm.tdmKernel = tdmKernel; cfg.tdm.ldsBytes = tdmLdsBytes; return cfg; diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index 2f7740c8..13c0cee6 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -170,6 +170,17 @@ namespace TransferBench NUM_GFX_KERNELS = 2 ///< Number of GFX kernels currently supported }; + /** + * Enumeration of supported TDM kernels + */ + enum TdmKernelType + { + TDM_KERNEL_AUTO = -1, ///< Automatically choose a kernel + TDM_KERNEL_COPY = 0, ///< Default kernel that copies a single input to a single output + TDM_KERNEL_REDUCE = 1, ///< Kernel that supports multiple input/output buffers (sum-reduce) + NUM_TDM_KERNELS = 2 ///< Number of TDM kernels currently supported + }; + /** * A MemDevice indicates a memory type on a specific device */ @@ -284,6 +295,7 @@ namespace TransferBench int blockOrder = 0; ///< Determines how threadblocks are ordered (0=sequential, 1=interleaved, 2=random) int blockSize = 256; ///< Size of each threadblock int ldsBytes = 0; ///< Amount of __shared__ memory per threadblock to use as bounce buffer (0 = device max) + int tdmKernel = 0; ///< Kernel selector: -1=auto, 0=copy-only, 1=reduce }; /** @@ -2062,6 +2074,7 @@ namespace { System::Get().Broadcast(root, sizeof(tdm), &tdm); if (tdm.blockOrder != cfg.tdm.blockOrder) ADD_ERROR("cfg.tdm.blockOrder"); if (tdm.blockSize != cfg.tdm.blockSize) ADD_ERROR("cfg.tdm.blockSize"); + if (tdm.tdmKernel != cfg.tdm.tdmKernel) ADD_ERROR("cfg.tdm.tdmKernel"); if (tdm.ldsBytes != cfg.tdm.ldsBytes) ADD_ERROR("cfg.tdm.ldsBytes"); } #undef ADD_ERROR @@ -2169,6 +2182,10 @@ namespace { "[tdm.blockSize] must be a positive multiple of 32 less than or equal to %d", MAX_BLOCKSIZE}); + if (cfg.tdm.tdmKernel < -1 || cfg.tdm.tdmKernel >= NUM_TDM_KERNELS) + errors.push_back( + {ERR_FATAL, "[tdm.tdmKernel] must be -1 for auto, or less than %d", NUM_TDM_KERNELS}); + if (cfg.tdm.ldsBytes < 0) errors.push_back({ERR_FATAL, "[tdm.ldsBytes] must be positive or 0"}); else { @@ -2384,9 +2401,22 @@ namespace { } break; case EXE_GPU_TDM: - if (t.srcs.size() != 1 || t.dsts.size() != 1) { + // The copy TDM kernel only supports a single SRC/DST; the reduce kernel supports + // any number of inputs/outputs. When auto-selecting, allow the reduce cardinalities. + if (cfg.tdm.tdmKernel == TDM_KERNEL_COPY && (t.srcs.size() != 1 || t.dsts.size() != 1)) { errors.push_back({ERR_FATAL, - "Transfer %d: GPU TDM kernel currently requires exactly 1 SRC and 1 DST", i}); + "Transfer %d: GPU TDM copy kernel currently requires exactly 1 SRC and 1 DST", i}); + hasFatalError = true; + break; + } + // The multi-SRC/DST sum-reduce kernel is only implemented on the AMD TDM + // backend; disable it on the NVIDIA platform. This covers both explicitly + // requesting the reduce kernel and auto-selecting it via reduce cardinalities. + if (TDM_PLATFORM_NV && + (cfg.tdm.tdmKernel == TDM_KERNEL_REDUCE || + (cfg.tdm.tdmKernel == TDM_KERNEL_AUTO && (t.srcs.size() != 1 || t.dsts.size() != 1)))) { + errors.push_back({ERR_FATAL, + "Transfer %d: GPU TDM reduce kernel (multi-SRC/DST) is not supported on the NVIDIA platform", i}); hasFatalError = true; break; } @@ -2951,6 +2981,7 @@ namespace { // For TDM-Executors uint32_t ldsBytesActual; ///< Actual number of LDS bytes to use as buffer + int tdmKernelToUse; ///< (TDM-only) Which TDM kernel to use }; // Structure to track PCIe topology @@ -4114,6 +4145,46 @@ namespace { return ERR_NONE; } + static bool CanUseTdmKernel(int const tdmKernelIdx, + ConfigOptions const& cfg, + vector const& transfers, + ExeInfo const& exeInfo) + { + // Reduce kernel supports any number of inputs/outputs + if (tdmKernelIdx == TDM_KERNEL_REDUCE) return true; + + // Copy kernel works if all Transfers have exactly one SRC / one DST + if (tdmKernelIdx == TDM_KERNEL_COPY) { + if (exeInfo.resources.empty()) return false; + for (auto const& rss : exeInfo.resources) { + Transfer const& t = transfers[rss.transferIdx]; + if (t.srcs.size() != 1 || t.dsts.size() != 1) return false; + } + return true; + } + + return false; + } + + static ErrResult SelectTdmKernel(ConfigOptions const& cfg, vector const& transfers, ExeInfo& exeInfo) + { + // Decide on which TDM kernel to use + // Auto-select - prefer copy kernel if eligible, otherwise fall back to reduce + if (cfg.tdm.tdmKernel == TDM_KERNEL_AUTO) { + exeInfo.tdmKernelToUse = CanUseTdmKernel(TDM_KERNEL_COPY, cfg, transfers, exeInfo) + ? TDM_KERNEL_COPY : TDM_KERNEL_REDUCE; + } else { + exeInfo.tdmKernelToUse = cfg.tdm.tdmKernel; + } + + // Warn if forcing copy kernel even though incompatible, but allow kernel to continue + if (cfg.tdm.tdmKernel == TDM_KERNEL_COPY && !CanUseTdmKernel(TDM_KERNEL_COPY, cfg, transfers, exeInfo)) { + return {ERR_WARN, + "TDM copy kernel forced even though deemed incompatible for current set of Transfers / config"}; + } + return ERR_NONE; + } + // Preparation-related functions //======================================================================================== @@ -5843,7 +5914,8 @@ namespace { // TDM Executor-related functions //======================================================================================== #if TDM_SUPPORTED - __global__ void GpuTdmKernel(SubExecParam* params, + // Copy kernel: single SRC -> single DST via tensor-DMA staged through shared memory + __global__ void TdmCopyKernel(SubExecParam* params, uint32_t ldsBytes, int numSubIterations) { @@ -5876,13 +5948,60 @@ namespace { GetXccId(p.xccId); } } + + // Reduce kernel: sum-reduce any number of SRCs into any number of DSTs + __global__ void TdmReduceKernel(SubExecParam* params, + uint32_t ldsBytes, + int numSubIterations) + { + int64_t startCycle; + bool const shouldRecordTiming = (threadIdx.x == 0); + if (shouldRecordTiming) startCycle = GetTimestamp(); + + extern __shared__ __align__(128) float shmem[]; + + // Each threadblock is a subexecutor (mirrors GpuCopyKernel). + SubExecParam& p = params[blockIdx.x]; + if (p.N == 0) return; + + // TODO: TDM-accelerated sum-reduce (mirror GpuReduceKernel using tensor loads/stores staged + // through shared memory). This naive per-thread implementation is a placeholder so the + // reduce path is wired end-to-end and compiles; replace with the tensor-DMA reduce. + int32_t const numSrcs = p.numSrcs; + int32_t const numDsts = p.numDsts; + size_t const sizeBytes = p.N * sizeof(float); + + int subIterations = 0; + while (1) { + //tdm::tdmReduce(p.dst, p.src, numSrcs, numDsts, sizeBytes, shmem, ldsBytes); + __syncthreads(); // Wait for all warps to finish this subiteration + if (++subIterations == numSubIterations) break; + } + + if (shouldRecordTiming) { + __threadfence_system(); + p.stopCycle = GetTimestamp(); + p.startCycle = startCycle; + GetHwId(p.hwId); + GetXccId(p.xccId); + } + } #else // gfx1250 tensor TDM builtins unavailable for this translation: emit empty kernel stubs with // the exact launch signatures so the host-side launch path still links. They are never // dispatched on non-gfx1250 or nvidia hardware (see TransfersHaveErrors). - __global__ void GpuTdmKernel(SubExecParam*, uint32_t, int) {} + __global__ void TdmCopyKernel(SubExecParam*, uint32_t, int) {} + __global__ void TdmReduceKernel(SubExecParam*, uint32_t, int) {} #endif // TDM_SUPPORTED + // Table of all TDM kernel functions - must match ordering in TdmKernelType + typedef void (*TdmKernelFuncPtr)(SubExecParam*, uint32_t, int); + TdmKernelFuncPtr TdmKernelsTable[NUM_TDM_KERNELS] = + { + TdmCopyKernel, // TDM_KERNEL_COPY + TdmReduceKernel, // TDM_KERNEL_REDUCE + }; + static ErrResult ExecuteTdmTransfer(int const iteration, int const exeTotalSubExecs, SubExecParam* exeSubExecParam, @@ -5892,6 +6011,7 @@ namespace { ConfigOptions const& cfg, bool const subExecParamHostAccessible, uint32_t const ldsBytes, + int const tdmKernelIdx, TransferResources& rss) { // Compute kernel launch parameters @@ -5900,18 +6020,21 @@ namespace { dim3 const blockSize(cfg.tdm.blockSize); SubExecParam* params = cfg.general.useMultiStream ? rss.subExecParamGpuPtr : exeSubExecParam; + // Select which TDM kernel to launch (must match ordering in TdmKernelType) + auto tdmKernel = TdmKernelsTable[tdmKernelIdx]; + auto cpuStart = std::chrono::high_resolution_clock::now(); #if defined(__NVCC__) if (cfg.general.useHipEvents) ERR_CHECK(hipEventRecord(startEvent, stream)); - GpuTdmKernel<<>>(params, - ldsBytes, - cfg.general.numSubIterations); + tdmKernel<<>>(params, + ldsBytes, + cfg.general.numSubIterations); if (cfg.general.useHipEvents) ERR_CHECK(hipEventRecord(stopEvent, stream)); #else - hipExtLaunchKernelGGL(GpuTdmKernel, gridSize, blockSize, (int)ldsBytes, stream, + hipExtLaunchKernelGGL(tdmKernel, gridSize, blockSize, (int)ldsBytes, stream, startEvent, stopEvent, 0, params, ldsBytes, cfg.general.numSubIterations); #endif @@ -5976,6 +6099,7 @@ namespace { std::cref(cfg), exeInfo.subExecParamHostAccessible, exeInfo.ldsBytesActual, + exeInfo.tdmKernelToUse, std::ref(exeInfo.resources[i]))); } for (auto& asyncTransfer : asyncTransfers) @@ -5985,7 +6109,8 @@ namespace { ExecuteTdmTransfer(iteration, exeInfo.totalSubExecs, exeInfo.subExecParamGpu, exeInfo.streams[0], cfg.general.useHipEvents ? exeInfo.startEvents[0] : NULL, cfg.general.useHipEvents ? exeInfo.stopEvents[0] : NULL, - cfg, exeInfo.subExecParamHostAccessible, exeInfo.ldsBytesActual, exeInfo.resources[0]); + cfg, exeInfo.subExecParamHostAccessible, exeInfo.ldsBytesActual, + exeInfo.tdmKernelToUse, exeInfo.resources[0]); } if (iteration >= 0) { @@ -6222,6 +6347,31 @@ namespace { if (exeDevice.exeType == EXE_GPU_GFX) { ERR_APPEND(SelectGfxKernel(cfg, transfers, exeInfo), errResults); } + + // Select which TDM kernel to use for this executor + if (exeDevice.exeType == EXE_GPU_TDM) { + ERR_APPEND(SelectTdmKernel(cfg, transfers, exeInfo), errResults); + + // For the TDM reduce kernel, warn about any SRC/DST buffer that is not + // 128B aligned (the tensor data mover reaches peak bandwidth on + // 128B-aligned addresses; misaligned buffers still work but slower). + if (exeInfo.tdmKernelToUse == TDM_KERNEL_REDUCE) { + for (auto const& rss : exeInfo.resources) { + for (int iSrc = 0; iSrc < (int)rss.srcMem.size(); ++iSrc) { + if (reinterpret_cast(rss.srcMem[iSrc]) & 127u) + errResults.push_back({ERR_WARN, + "Transfer %d: TDM reduce SRC[%d] (%p) is not 128B aligned; performance may be reduced", + rss.transferIdx, iSrc, (void*)rss.srcMem[iSrc]}); + } + for (int iDst = 0; iDst < (int)rss.dstMem.size(); ++iDst) { + if (reinterpret_cast(rss.dstMem[iDst]) & 127u) + errResults.push_back({ERR_WARN, + "Transfer %d: TDM reduce DST[%d] (%p) is not 128B aligned; performance may be reduced", + rss.transferIdx, iDst, (void*)rss.dstMem[iDst]}); + } + } + } + } } // Prepare reference src/dst arrays - only once for largest size. From 4734f9acb474e0bbd9cbbdfdc4143b9644780f7f Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Wed, 29 Jul 2026 21:39:00 +0000 Subject: [PATCH 2/7] small addition for TdmSweep config --- src/client/Presets/TdmSweep.hpp | 97 +++++++++++++++++++-------------- 1 file changed, 55 insertions(+), 42 deletions(-) diff --git a/src/client/Presets/TdmSweep.hpp b/src/client/Presets/TdmSweep.hpp index 31b3845f..2a3ceeb9 100644 --- a/src/client/Presets/TdmSweep.hpp +++ b/src/client/Presets/TdmSweep.hpp @@ -57,6 +57,7 @@ int TdmSweepPreset(EnvVars& ev, // Collect environment variables for this preset vector blockList = EnvVars::GetEnvVarArray("BLOCKSIZES", {64,128,256,512,1024}); vector blockOrders = EnvVars::GetEnvVarArray("BLOCK_ORDERS", {0}); + vector byteOffsets = EnvVars::GetEnvVarArray("BYTE_OFFSETS", {0}); vector ldsList = EnvVars::GetEnvVarArray("LDS_BYTES", {8192,16384,32768,65536,0}); vector numSesList = EnvVars::GetEnvVarArray("NUM_SUB_EXECS", {2,4,8,16,32,64}); int numTransfers = EnvVars::GetEnvVar( "NUM_TRANSFERS", 1); @@ -71,6 +72,7 @@ int TdmSweepPreset(EnvVars& ev, Utils::Print("[TDM Sweep Related]\n"); ev.Print("BLOCKSIZES", blockList.size(), EnvVars::ToStr(blockList).c_str()); ev.Print("BLOCK_ORDERS", blockOrders.size(), "%s (0=sequential 1=interleaved 2=random)", EnvVars::ToStr(blockOrders).c_str()); + ev.Print("BYTE_OFFSETS", byteOffsets.size(), "%s (src/dst alloc offset; drives the TDM vector head/tail path)", EnvVars::ToStr(byteOffsets).c_str()); ev.Print("LDS_BYTES", ldsList.size(), "%s (0 = device max LDS per block)", EnvVars::ToStr(ldsList).c_str()); ev.Print("NUM_SUB_EXECS", numSesList.size(), EnvVars::ToStr(numSesList).c_str()); ev.Print("NUM_TRANSFERS", numTransfers, "Number of Transfers specified in TDM_TRANSFER"); @@ -109,6 +111,13 @@ int TdmSweepPreset(EnvVars& ev, return ERR_FATAL; } } + for (int off : byteOffsets) { + if (off < 0 || off % (int)sizeof(float)) { + Utils::Print("[ERROR] BYTE_OFFSETS value %d is invalid (must be a non-negative multiple of %lu)\n", + off, sizeof(float)); + return ERR_FATAL; + } + } std::vector transfers; Utils::CheckForError(ParseTransfers(std::to_string(numTransfers) + " 1 " + transferStr, transfers)); @@ -173,7 +182,7 @@ int TdmSweepPreset(EnvVars& ev, // Print header char sep = ev.outputToCsv ? ',' : ' '; - Utils::Print(" BlkO %c BlkS %c LDSBytes ", sep, sep); + Utils::Print(" BlkO %c BlkS %c LDSBytes %c ByteOff ", sep, sep, sep); for (int numSubExec : numSesList) Utils::Print("%c SE %03d", sep, numSubExec); Utils::Print("\n"); @@ -181,60 +190,63 @@ int TdmSweepPreset(EnvVars& ev, int bestSe = -1; double overallBestBw = 0; vector bestBw(numSesList.size(), 0.0); - // best[s] = {blockOrder, blockSize, ldsBytes, numSubExec} - vector> best(numSesList.size(), vector(4)); + // best[s] = {blockOrder, blockSize, ldsBytes, byteOffset, numSubExec} + vector> best(numSesList.size(), vector(5)); // Loop over all combinations for (int blockOrder : blockOrders) { cfg.tdm.blockOrder = blockOrder; for (int blockSize : blockList) { cfg.tdm.blockSize = blockSize; for (int ldsBytes : ldsList) { cfg.tdm.ldsBytes = ldsBytes; - Utils::Print(" %1d %c %4d %c %8d ", blockOrder, sep, blockSize, sep, ldsBytes); - fflush(stdout); - for (auto s = 0; s < numSesList.size(); s++) { - int numSubExec = numSesList[s]; - for (Transfer& t : transfers) t.numSubExecs = numSubExec; + for (int byteOffset : byteOffsets) { cfg.data.byteOffset = byteOffset; + Utils::Print(" %1d %c %4d %c %8d %c %7d ", + blockOrder, sep, blockSize, sep, ldsBytes, sep, byteOffset); + fflush(stdout); + for (auto s = 0; s < numSesList.size(); s++) { + int numSubExec = numSesList[s]; + for (Transfer& t : transfers) t.numSubExecs = numSubExec; - TestResults result; - // A given combination may be rejected by the library (e.g. LDS window - // larger than the device max). Treat that as a skipped cell (N/A) and - // keep sweeping instead of aborting the whole matrix. - if (RunTransfers(cfg, transfers, result)) { - double bw = 0.0; - switch (timingMode) { - case 0: bw = result.avgTotalBandwidthGbPerSec; break; - case 1: - for (auto const& e : result.exeResults) { - bw = std::max(bw, e.second.avgBandwidthGbPerSec); - } - break; - case 2: default: - for (auto const& t : result.tfrResults) { - bw = std::max(bw, t.avgBandwidthGbPerSec); + TestResults result; + // A given combination may be rejected by the library (e.g. LDS window + // larger than the device max). Treat that as a skipped cell (N/A) and + // keep sweeping instead of aborting the whole matrix. + if (RunTransfers(cfg, transfers, result)) { + double bw = 0.0; + switch (timingMode) { + case 0: bw = result.avgTotalBandwidthGbPerSec; break; + case 1: + for (auto const& e : result.exeResults) { + bw = std::max(bw, e.second.avgBandwidthGbPerSec); + } + break; + case 2: default: + for (auto const& t : result.tfrResults) { + bw = std::max(bw, t.avgBandwidthGbPerSec); + } + break; } - break; - } - if (bw > bestBw[s]) { - bestBw[s] = bw; - best[s] = {blockOrder, blockSize, ldsBytes, numSubExec}; - if (bw > overallBestBw) { - overallBestBw = bw; - bestSe = s; + if (bw > bestBw[s]) { + bestBw[s] = bw; + best[s] = {blockOrder, blockSize, ldsBytes, byteOffset, numSubExec}; + if (bw > overallBestBw) { + overallBestBw = bw; + bestSe = s; + } } + Utils::Print("%c%8.2f", sep, bw); + } else { + Utils::Print("%c%8s", sep, "N/A"); } - Utils::Print("%c%8.2f", sep, bw); - } else { - Utils::Print("%c%8s", sep, "N/A"); + fflush(stdout); } + Utils::Print("\n"); fflush(stdout); } - Utils::Print("\n"); - fflush(stdout); } } } - Utils::Print(" BlkO %c BlkS %c LDSBytes ", sep, sep); + Utils::Print(" BlkO %c BlkS %c LDSBytes %c ByteOff ", sep, sep, sep); for (auto s = 0; s < numSesList.size(); s++) { Utils::Print("%c%8.2f", sep, bestBw[s]); } @@ -254,10 +266,11 @@ int TdmSweepPreset(EnvVars& ev, Utils::Print(" BlockOrder : %7d [TDM_BLOCK_ORDER=%d]\n", best[bestSe][0], best[bestSe][0]); Utils::Print(" BlockSize : %7d [TDM_BLOCK_SIZE=%d]\n", best[bestSe][1], best[bestSe][1]); Utils::Print(" LDS Bytes : %7d [TDM_LDS_BYTES=%d]\n", best[bestSe][2], best[bestSe][2]); - Utils::Print(" NumSubExec : %7d\n", best[bestSe][3]); + Utils::Print(" Byte Offset : %7d [BYTE_OFFSET=%d]\n", best[bestSe][3], best[bestSe][3]); + Utils::Print(" NumSubExec : %7d\n", best[bestSe][4]); Utils::Print("Command to run best result:\n"); - Utils::Print("TDM_BLOCK_ORDER=%d TDM_BLOCK_SIZE=%d TDM_LDS_BYTES=%d ./TransferBench cmdline %lu \"%d %d %s\"\n", - best[bestSe][0], best[bestSe][1], best[bestSe][2], - numBytesPerTransfer, numTransfers, best[bestSe][3], transferStr.c_str()); + Utils::Print("TDM_BLOCK_ORDER=%d TDM_BLOCK_SIZE=%d TDM_LDS_BYTES=%d BYTE_OFFSET=%d ./TransferBench cmdline %lu \"%d %d %s\"\n", + best[bestSe][0], best[bestSe][1], best[bestSe][2], best[bestSe][3], + numBytesPerTransfer, numTransfers, best[bestSe][4], transferStr.c_str()); return ERR_NONE; } From 48bde52b508ebf00a40a09e2a171c4c7f7de483d Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Wed, 29 Jul 2026 21:40:20 +0000 Subject: [PATCH 3/7] scratch reduceCopy --- src/header/TransferBench.hpp | 5 +- src/header/Untitled-2.cpp | 179 +++++++++++++++++++++++++++++ src/header/tdmCopy.h | 213 ++++++++++++++++++++++++++++++++++- 3 files changed, 388 insertions(+), 9 deletions(-) create mode 100644 src/header/Untitled-2.cpp diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index 13c0cee6..52627be4 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -5964,16 +5964,13 @@ namespace { SubExecParam& p = params[blockIdx.x]; if (p.N == 0) return; - // TODO: TDM-accelerated sum-reduce (mirror GpuReduceKernel using tensor loads/stores staged - // through shared memory). This naive per-thread implementation is a placeholder so the - // reduce path is wired end-to-end and compiles; replace with the tensor-DMA reduce. int32_t const numSrcs = p.numSrcs; int32_t const numDsts = p.numDsts; size_t const sizeBytes = p.N * sizeof(float); int subIterations = 0; while (1) { - //tdm::tdmReduce(p.dst, p.src, numSrcs, numDsts, sizeBytes, shmem, ldsBytes); + tdm::tdmReduce(p.dst, p.src, numSrcs, numDsts, sizeBytes, shmem, ldsBytes); __syncthreads(); // Wait for all warps to finish this subiteration if (++subIterations == numSubIterations) break; } diff --git a/src/header/Untitled-2.cpp b/src/header/Untitled-2.cpp new file mode 100644 index 00000000..cfdf0a5e --- /dev/null +++ b/src/header/Untitled-2.cpp @@ -0,0 +1,179 @@ +// ---- vector fallback: reduce (sum) numSrcs sources, broadcast to numDsts. ---- +// Element-wise sum of all srcs -> written to every dst. numSrcs == 1 is a plain +// copy; numDsts > 1 broadcasts the same reduced result to each destination. +// `srcs`/`dsts` hold base addresses; `offset` (in bytes) reaches the sub-range +// to copy (e.g. the tail start), so callers can share one base pointer array. +__device__ inline void warpVecCopy(uint64_t* srcs, uint64_t* dsts, + uint32_t numSrcs, uint32_t numDsts, + size_t n, size_t offset, + uint32_t warpThread, uint32_t warpThreads) { + size_t nd = n >> 2; + for (size_t i = warpThread; i < nd; i += warpThreads) { + uint32_t acc = 0; + for (uint32_t s = 0; s < numSrcs; ++s) + acc += reinterpret_cast(srcs[s] + offset)[i]; + for (uint32_t d = 0; d < numDsts; ++d) + reinterpret_cast(dsts[d] + offset)[i] = acc; + } + uint32_t rem = static_cast(n & 3u); + if (rem && warpThread == 0) { + for (uint32_t b = 0; b < rem; ++b) { + uint8_t acc = 0; + for (uint32_t s = 0; s < numSrcs; ++s) + acc += reinterpret_cast(srcs[s] + offset)[nd * 4 + b]; + for (uint32_t d = 0; d < numDsts; ++d) + reinterpret_cast(dsts[d] + offset)[nd * 4 + b] = acc; + } + } +} + +// ---- issue one chunk of whole 256B rows (2D tile) through ONE LDS window. --- +// This staging window is single-buffered, so the two TDM ops form a dependency +// chain that MUST be enforced with TENSORcnt waits -- up to 3 TDM ops are +// outstanding per wave (they overlap), so "same-wave in-order issue" does NOT +// serialize their memory effects: +// * load -> wait: the store reads the LDS the load just wrote (RAW hazard). +// * store -> wait: the caller reuses this same window next iteration; the next +// load must not overwrite LDS the store is still draining (WAR hazard). +__device__ inline void issueRows(uint64_t* srcs, uint64_t* dsts, + uint32_t numSrcs, uint32_t numDsts, + uint32_t val, uint32_t tmp, uint32_t rows, uint32_t off = 0) { + gfx1250_TDM_GROUP1 g1; + g1.dataSize(DS4); + g1.tileDim0(TD0); g1.tileDim1(rows); + g1.tensorDim0(TD0); g1.tensorDim1(rows); + g1.tensorDim0Stride(TD0); // rows back-to-back (contiguous) + + if (numSrcs) { + gfx1250_TDM_GROUP0 g0l(val, srcs[0] + off); + load(g0l, g1); waitTensor0(); + for (size_t s = 1; s < numSrcs; s++) { + gfx1250_TDM_GROUP0 g0l(tmp, srcs[s] + off); + load(g0l, g1); waitTensor0(); + for (size_t u = 0; u < rows * WIDTH; u += Warp) { + val[u] += tpm[u]; + } + } + } + + for (size_t d = 0; d < numDsts; d++) { + gfx1250_TDM_GROUP0 g0s(tmp, dsts[d] + off); + store(g0s, g1); waitTensor0(); + } +} + +// ---- issue a sub-row tail (<256B) as a 1-D tile at BYTE granularity. --------- +// Same single-buffered LDS window and the same required RAW/WAR waits as above. +__device__ inline void issueRow1d(uint64_t* srcs, uint64_t* dsts, + uint32_t numSrcs, uint32_t numDsts, + uint32_t val, uint32_t tmp, uint32_t nbytes, + uint64_t off = 0) { + gfx1250_TDM_GROUP1 g1; + g1.dataSize(DS1); // 1-byte elements: exact length + g1.tileDim0(nbytes); g1.tileDim1(1); + g1.tensorDim0(nbytes); g1.tensorDim1(1); + g1.tensorDim0Stride(nbytes); + + if (numSrcs) { + gfx1250_TDM_GROUP0 g0l(val, srcs[0] + off); // unused higher dims -> zero (see load()) + load(g0l, g1); waitTensor0(); // RAW: fill LDS before store/reduce reads it + for (size_t s = 1; s < numSrcs; s++) { + gfx1250_TDM_GROUP0 g0l(tmp, srcs[s] + off); + load(g0l, g1); waitTensor0(); + for (size_t u = 0; u < nbytes; u += Warp) { + val[u] += tmp[u]; // reduce: accumulate into val + } + } + } + + for (size_t d = 0; d < numDsts; d++) { + gfx1250_TDM_GROUP0 g0s(val, dsts[d] + off); // broadcast reduced result to each dst + store(g0s, g1); waitTensor0(); // WAR: drain store before window reuse + } +} + +__device__ inline void issue(void** dsts, const void** srcs, uint32_t numSrcs, uint32_t numDsts, + size_t sizeBytes, void* ldsBuffer, size_t ldsBufferBytes, + uint32_t startWarpId, uint32_t stopWarpId) { + const uint32_t ldsBase = static_cast(reinterpret_cast(ldsBuffer)); + const uint32_t ldsBytes = static_cast(ldsBufferBytes); // LDS is small + + const uint32_t W = warpSize; + const uint32_t nThreads = blockDim.x * blockDim.y * blockDim.z; + const uint32_t tid = (threadIdx.z * blockDim.y + threadIdx.y) * blockDim.x + + threadIdx.x; + const uint32_t warpThread = tid % W; // thread index within its warp + const uint32_t warpId = tid / W; + const uint32_t nWarps = (nThreads + W - 1) / W; + + // --- team membership: this warp participates iff in [start, stop) -------- + const uint32_t teamStop = (stopWarpId > nWarps) ? nWarps : stopWarpId; + if (startWarpId >= teamStop || warpId < startWarpId || warpId >= teamStop) + return; // not on this team + const uint32_t rank = warpId - startWarpId; // rank within the team + const uint32_t teamWarps = teamStop - startWarpId; // >= 1 + + // active threads in THIS warp (handles partial final warp); stride for vector. + const uint32_t warpThreads = (nThreads - warpId * W < W) ? (nThreads - warpId * W) : W; + + // --- split the range: [256B rows ][tail] ------------------ + size_t rows = sizeBytes / WIDTH; // whole 256B rows + size_t tail = sizeBytes % WIDTH; + size_t tdmBytes = rows * WIDTH; + + // --- base src/dst addresses for this team (byte offset applied per use) -- + uint64_t mySrcs[MAX_SRCS]; + uint64_t myDsts[MAX_DSTS]; + for (uint32_t i = 0; i < numSrcs; i++) mySrcs[i] = (uint64_t)srcs[i]; + for (uint32_t i = 0; i < numDsts; i++) myDsts[i] = (uint64_t)dsts[i]; + + // --- LDS reduce buffers: val = running sum, tmp = staging for extra srcs -- + // Base offsets (rank 0); the TDM path shifts each warp by rank*window below. + uint32_t val = ldsBase; + uint32_t tmp = ldsBase + WIDTH; + + // --- edges (team's FIRST warp = rank 0): vector head, TDM tail ----------- + if (rank == 0 && tail) { + // LDS needed: val (+ tmp when reducing multiple srcs), each holding `tail` bytes + uint32_t need = (numSrcs > 1 ? (tmp - ldsBase) : (val - ldsBase)) + tail; + if (ldsBytes >= need) { // stage tail in rank 0's window + issueRow1d(mySrcs, myDsts, numSrcs, numDsts, val, tmp, tail, tdmBytes); + } else { + warpVecCopy(mySrcs, myDsts, numSrcs, numDsts, tail, tdmBytes, + warpThread, warpThreads); + } + } + + // --- 256B rows copy via TDM ------------------------------------------------ + uint32_t maxByLds = ldsBytes / RWIDTH; // #warps we can give a window, 2*WIDTH because of double buffering + if (maxByLds == 0) { // LDS < 512B: vector fallback + if (rank == 0) + warpVecCopy(mySrcs, myDsts, numSrcs, numDsts, tdmBytes, 0, + warpThread, warpThreads); + return; + } + uint32_t issuers = teamWarps < maxByLds ? teamWarps : maxByLds; + uint32_t window = (ldsBytes / issuers) & ~(RWIDTH - 1); // per-warp 512B-multiple + uint32_t rowsPerChunk = window / RWIDTH; + + if (rank >= issuers) return; // this warp doesn't issue + + // distribute `rows` across issuers by team rank (contiguous row blocks) + size_t base = rows / issuers; + size_t extra = rows % issuers; + size_t myRows = base + (rank < extra ? 1u : 0u); + // TODO:if last warp, take the remaining edges + size_t myStart = rank * base + (rank < extra ? rank : extra); + if (myRows == 0) return; + + // shift this warp's reduce buffers into its own window + val += rank * window; + tmp += rank * window; + + for (size_t r = 0; r < myRows; r += rowsPerChunk) { + uint32_t chunkRows = (myRows - r < rowsPerChunk) + ? static_cast(myRows - r) : rowsPerChunk; + uint64_t off = (uint64_t)(myStart + r) * WIDTH; // this warp's row block + issueRows(mySrcs, myDsts, numSrcs, numDsts, val, tmp, chunkRows, off); + } +} \ No newline at end of file diff --git a/src/header/tdmCopy.h b/src/header/tdmCopy.h index 2fc79a13..94297f36 100644 --- a/src/header/tdmCopy.h +++ b/src/header/tdmCopy.h @@ -345,11 +345,14 @@ namespace tdm { namespace detail { -constexpr uint32_t WIDTH = 256; // bytes per TDM row (first tile dim) -constexpr uint32_t ELT = 4; // dword -constexpr uint32_t DS4 = 2; // data_size code for 4-byte -constexpr uint32_t DS1 = 0; // data_size code for 1-byte -constexpr uint32_t TD0 = WIDTH / ELT; // 64 elements per row +constexpr uint32_t WIDTH = 256; // bytes per TDM row (first tile dim) +constexpr uint32_t ELT = 4; // dword +constexpr uint32_t DS4 = 2; // data_size code for 4-byte +constexpr uint32_t DS1 = 0; // data_size code for 1-byte +constexpr uint32_t TD0 = WIDTH / ELT; // 64 elements per row +constexpr uint32_t RWIDTH = 512; +constexpr uint32_t MAX_SRCS = 16; // max sources reduced in one call +constexpr uint32_t MAX_DSTS = 16; // max destinations broadcast to // ---- instruction emission (the only arch-specific piece) ------------------- // The tensor DMA is a single builtin taking the FULL descriptor: five register @@ -521,6 +524,198 @@ __device__ inline void issue(void* dst, const void* src, size_t sizeBytes, } } +// ############################################################################ +// # MULTI-SOURCE REDUCE / MULTI-DEST BROADCAST VARIANTS # +// ############################################################################ +// Reduce (element-wise sum) numSrcs sources and broadcast the result to numDsts +// destinations. numSrcs == 1 && numDsts == 1 degenerates to a plain copy. These +// are overloads of the single-copy helpers above; the reduce entry point routes +// here via detail::issue(dsts, srcs, numSrcs, numDsts, ...). + +// ---- vector fallback: reduce (sum) numSrcs sources, broadcast to numDsts. ---- +// Element-wise sum of all srcs -> written to every dst. numSrcs == 1 is a plain +// copy; numDsts > 1 broadcasts the same reduced result to each destination. +// `srcs`/`dsts` hold base addresses; `offset` (in bytes) reaches the sub-range +// to copy (e.g. the tail start), so callers can share one base pointer array. +__device__ inline void warpVecCopy(uint64_t* srcs, uint64_t* dsts, + uint32_t numSrcs, uint32_t numDsts, + size_t n, size_t offset, + uint32_t warpThread, uint32_t warpThreads) { + size_t nd = n >> 2; + for (size_t i = warpThread; i < nd; i += warpThreads) { + uint32_t acc = 0; + for (uint32_t s = 0; s < numSrcs; ++s) + acc += reinterpret_cast(srcs[s] + offset)[i]; + for (uint32_t d = 0; d < numDsts; ++d) + reinterpret_cast(dsts[d] + offset)[i] = acc; + } + uint32_t rem = static_cast(n & 3u); + if (rem && warpThread == 0) { + for (uint32_t b = 0; b < rem; ++b) { + uint8_t acc = 0; + for (uint32_t s = 0; s < numSrcs; ++s) + acc += reinterpret_cast(srcs[s] + offset)[nd * 4 + b]; + for (uint32_t d = 0; d < numDsts; ++d) + reinterpret_cast(dsts[d] + offset)[nd * 4 + b] = acc; + } + } +} + +// ---- issue one chunk of whole 256B rows (2D tile) through ONE LDS window. --- +// Same single-buffered LDS window / RAW+WAR waits as the single-copy issueRows(). +// The first source lands in `val`; each subsequent source lands in `tmp` and is +// accumulated into `val`; the reduced `val` is stored to every destination. +__device__ inline void issueRows(uint64_t* srcs, uint64_t* dsts, + uint32_t numSrcs, uint32_t numDsts, + uint32_t val, uint32_t tmp, uint32_t rows, uint32_t off = 0) { + gfx1250_TDM_GROUP1 g1; + g1.dataSize(DS4); + g1.tileDim0(TD0); g1.tileDim1(rows); + g1.tensorDim0(TD0); g1.tensorDim1(rows); + g1.tensorDim0Stride(TD0); // rows back-to-back (contiguous) + + if (numSrcs) { + gfx1250_TDM_GROUP0 g0l(val, srcs[0] + off); + load(g0l, g1); waitTensor0(); + // LDS byte-offsets -> typed shared pointers for the reduce accumulation + uint32_t* vp = reinterpret_cast(static_cast(val)); + const uint32_t* tp = reinterpret_cast(static_cast(tmp)); + uint32_t nElems = (rows * WIDTH) >> 2; // 4-byte (DS4) elements in this tile + for (size_t s = 1; s < numSrcs; s++) { + gfx1250_TDM_GROUP0 g0l(tmp, srcs[s] + off); + load(g0l, g1); waitTensor0(); + for (uint32_t u = 0; u < nElems; u += warpSize) { + vp[u] += tp[u]; + } + } + } + + for (size_t d = 0; d < numDsts; d++) { + gfx1250_TDM_GROUP0 g0s(val, dsts[d] + off); // broadcast reduced result to each dst + store(g0s, g1); waitTensor0(); + } +} + +// ---- issue a sub-row tail (<256B) as a 1-D tile at BYTE granularity. --------- +// Same single-buffered LDS window and the same required RAW/WAR waits as above. +__device__ inline void issueRow1d(uint64_t* srcs, uint64_t* dsts, + uint32_t numSrcs, uint32_t numDsts, + uint32_t val, uint32_t tmp, uint32_t nbytes, + uint64_t off = 0) { + gfx1250_TDM_GROUP1 g1; + g1.dataSize(DS1); // 1-byte elements: exact length + g1.tileDim0(nbytes); g1.tileDim1(1); + g1.tensorDim0(nbytes); g1.tensorDim1(1); + g1.tensorDim0Stride(nbytes); + + if (numSrcs) { + gfx1250_TDM_GROUP0 g0l(val, srcs[0] + off); // unused higher dims -> zero (see load()) + load(g0l, g1); waitTensor0(); // RAW: fill LDS before store/reduce reads it + // LDS byte-offsets -> typed shared pointers for the reduce accumulation + uint8_t* vp = reinterpret_cast(static_cast(val)); + const uint8_t* tp = reinterpret_cast(static_cast(tmp)); + for (size_t s = 1; s < numSrcs; s++) { + gfx1250_TDM_GROUP0 g0l(tmp, srcs[s] + off); + load(g0l, g1); waitTensor0(); + for (uint32_t u = 0; u < nbytes; u += warpSize) { // 1-byte (DS1) elements + vp[u] += tp[u]; // reduce: accumulate into val + } + } + } + + for (size_t d = 0; d < numDsts; d++) { + gfx1250_TDM_GROUP0 g0s(val, dsts[d] + off); // broadcast reduced result to each dst + store(g0s, g1); waitTensor0(); // WAR: drain store before window reuse + } +} + +// ---- core: partition + issue a reduce/broadcast for the team [start, stop). - +__device__ inline void issue(void** dsts, const void** srcs, uint32_t numSrcs, uint32_t numDsts, + size_t sizeBytes, void* ldsBuffer, size_t ldsBufferBytes, + uint32_t startWarpId, uint32_t stopWarpId) { + const uint32_t ldsBase = static_cast(reinterpret_cast(ldsBuffer)); + const uint32_t ldsBytes = static_cast(ldsBufferBytes); // LDS is small + + const uint32_t W = warpSize; + const uint32_t nThreads = blockDim.x * blockDim.y * blockDim.z; + const uint32_t tid = (threadIdx.z * blockDim.y + threadIdx.y) * blockDim.x + + threadIdx.x; + const uint32_t warpThread = tid % W; // thread index within its warp + const uint32_t warpId = tid / W; + const uint32_t nWarps = (nThreads + W - 1) / W; + + // --- team membership: this warp participates iff in [start, stop) -------- + const uint32_t teamStop = (stopWarpId > nWarps) ? nWarps : stopWarpId; + if (startWarpId >= teamStop || warpId < startWarpId || warpId >= teamStop) + return; // not on this team + const uint32_t rank = warpId - startWarpId; // rank within the team + const uint32_t teamWarps = teamStop - startWarpId; // >= 1 + + // active threads in THIS warp (handles partial final warp); stride for vector. + const uint32_t warpThreads = (nThreads - warpId * W < W) ? (nThreads - warpId * W) : W; + + // --- split the range: [256B rows ][tail] ------------------ + size_t rows = sizeBytes / WIDTH; // whole 256B rows + size_t tail = sizeBytes % WIDTH; + size_t tdmBytes = rows * WIDTH; + + // --- base src/dst addresses for this team (byte offset applied per use) -- + uint64_t mySrcs[MAX_SRCS]; + uint64_t myDsts[MAX_DSTS]; + for (uint32_t i = 0; i < numSrcs; i++) mySrcs[i] = (uint64_t)srcs[i]; + for (uint32_t i = 0; i < numDsts; i++) myDsts[i] = (uint64_t)dsts[i]; + + // --- LDS reduce buffers: val = running sum, tmp = staging for extra srcs -- + // Base offsets (rank 0); the TDM path shifts each warp by rank*window below. + uint32_t val = ldsBase; + uint32_t tmp = ldsBase + WIDTH; + + // --- edges (team's FIRST warp = rank 0): vector head, TDM tail ----------- + if (rank == 0 && tail) { + // LDS needed: val (+ tmp when reducing multiple srcs), each holding `tail` bytes + uint32_t need = (numSrcs > 1 ? (tmp - ldsBase) : (val - ldsBase)) + tail; + if (ldsBytes >= need) { // stage tail in rank 0's window + issueRow1d(mySrcs, myDsts, numSrcs, numDsts, val, tmp, tail, tdmBytes); + } else { + warpVecCopy(mySrcs, myDsts, numSrcs, numDsts, tail, tdmBytes, + warpThread, warpThreads); + } + } + + // --- 256B rows copy via TDM ------------------------------------------------ + uint32_t maxByLds = ldsBytes / RWIDTH; // #warps we can give a window, 2*WIDTH because of double buffering + if (maxByLds == 0) { // LDS < 512B: vector fallback + if (rank == 0) + warpVecCopy(mySrcs, myDsts, numSrcs, numDsts, tdmBytes, 0, + warpThread, warpThreads); + return; + } + uint32_t issuers = teamWarps < maxByLds ? teamWarps : maxByLds; + uint32_t window = (ldsBytes / issuers) & ~(RWIDTH - 1); // per-warp 512B-multiple + uint32_t rowsPerChunk = window / RWIDTH; + + if (rank >= issuers) return; // this warp doesn't issue + + // distribute `rows` across issuers by team rank (contiguous row blocks) + size_t base = rows / issuers; + size_t extra = rows % issuers; + size_t myRows = base + (rank < extra ? 1u : 0u); + // TODO:if last warp, take the remaining edges + size_t myStart = rank * base + (rank < extra ? rank : extra); + if (myRows == 0) return; + + // shift this warp's reduce buffers into its own window + val += rank * window; + tmp += rank * window; + + for (size_t r = 0; r < myRows; r += rowsPerChunk) { + uint32_t chunkRows = (myRows - r < rowsPerChunk) + ? static_cast(myRows - r) : rowsPerChunk; + uint64_t off = (uint64_t)(myStart + r) * WIDTH; // this warp's row block + issueRows(mySrcs, myDsts, numSrcs, numDsts, val, tmp, chunkRows, off); + } +} + } // namespace detail #elif TDM_BACKEND_NV // ----- NVIDIA cp.async.bulk / TMA (sm_90+) ----- @@ -714,6 +909,14 @@ __device__ inline void tdmCopyByTeam(void* dst, const void* src, size_t sizeByte tdmWait(); // no-op on any warp that issued nothing / is off-team } +#if TDM_BACKEND_AMD // multi-source reduce path only exists on the AMD (TDM) backend +__device__ inline void tdmReduce(void** dsts, const void** srcs, uint32_t numSrcs, uint32_t numDsts, + size_t sizeBytes, void* ldsBuffer, size_t ldsBufferBytes) { + detail::issue(dsts, srcs, numSrcs, numDsts, sizeBytes, ldsBuffer, ldsBufferBytes, /*start=*/0, /*stop=*/~0u); + tdmWait(); +} +#endif // TDM_BACKEND_AMD + #endif // TDM_SUPPORTED // On an unsupported target the entry points were declared `= delete` at the top, // so there is nothing to define here -- any call is a compile-time error. From e68486042f977482738ae8f40302614707d62c04 Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Thu, 6 Aug 2026 01:47:38 +0000 Subject: [PATCH 4/7] fix reduce offset error; datatype template; support no src/dst --- src/header/Untitled-2.cpp | 179 -------------------------------------- src/header/tdmCopy.h | 104 ++++++++++++++++------ 2 files changed, 79 insertions(+), 204 deletions(-) delete mode 100644 src/header/Untitled-2.cpp diff --git a/src/header/Untitled-2.cpp b/src/header/Untitled-2.cpp deleted file mode 100644 index cfdf0a5e..00000000 --- a/src/header/Untitled-2.cpp +++ /dev/null @@ -1,179 +0,0 @@ -// ---- vector fallback: reduce (sum) numSrcs sources, broadcast to numDsts. ---- -// Element-wise sum of all srcs -> written to every dst. numSrcs == 1 is a plain -// copy; numDsts > 1 broadcasts the same reduced result to each destination. -// `srcs`/`dsts` hold base addresses; `offset` (in bytes) reaches the sub-range -// to copy (e.g. the tail start), so callers can share one base pointer array. -__device__ inline void warpVecCopy(uint64_t* srcs, uint64_t* dsts, - uint32_t numSrcs, uint32_t numDsts, - size_t n, size_t offset, - uint32_t warpThread, uint32_t warpThreads) { - size_t nd = n >> 2; - for (size_t i = warpThread; i < nd; i += warpThreads) { - uint32_t acc = 0; - for (uint32_t s = 0; s < numSrcs; ++s) - acc += reinterpret_cast(srcs[s] + offset)[i]; - for (uint32_t d = 0; d < numDsts; ++d) - reinterpret_cast(dsts[d] + offset)[i] = acc; - } - uint32_t rem = static_cast(n & 3u); - if (rem && warpThread == 0) { - for (uint32_t b = 0; b < rem; ++b) { - uint8_t acc = 0; - for (uint32_t s = 0; s < numSrcs; ++s) - acc += reinterpret_cast(srcs[s] + offset)[nd * 4 + b]; - for (uint32_t d = 0; d < numDsts; ++d) - reinterpret_cast(dsts[d] + offset)[nd * 4 + b] = acc; - } - } -} - -// ---- issue one chunk of whole 256B rows (2D tile) through ONE LDS window. --- -// This staging window is single-buffered, so the two TDM ops form a dependency -// chain that MUST be enforced with TENSORcnt waits -- up to 3 TDM ops are -// outstanding per wave (they overlap), so "same-wave in-order issue" does NOT -// serialize their memory effects: -// * load -> wait: the store reads the LDS the load just wrote (RAW hazard). -// * store -> wait: the caller reuses this same window next iteration; the next -// load must not overwrite LDS the store is still draining (WAR hazard). -__device__ inline void issueRows(uint64_t* srcs, uint64_t* dsts, - uint32_t numSrcs, uint32_t numDsts, - uint32_t val, uint32_t tmp, uint32_t rows, uint32_t off = 0) { - gfx1250_TDM_GROUP1 g1; - g1.dataSize(DS4); - g1.tileDim0(TD0); g1.tileDim1(rows); - g1.tensorDim0(TD0); g1.tensorDim1(rows); - g1.tensorDim0Stride(TD0); // rows back-to-back (contiguous) - - if (numSrcs) { - gfx1250_TDM_GROUP0 g0l(val, srcs[0] + off); - load(g0l, g1); waitTensor0(); - for (size_t s = 1; s < numSrcs; s++) { - gfx1250_TDM_GROUP0 g0l(tmp, srcs[s] + off); - load(g0l, g1); waitTensor0(); - for (size_t u = 0; u < rows * WIDTH; u += Warp) { - val[u] += tpm[u]; - } - } - } - - for (size_t d = 0; d < numDsts; d++) { - gfx1250_TDM_GROUP0 g0s(tmp, dsts[d] + off); - store(g0s, g1); waitTensor0(); - } -} - -// ---- issue a sub-row tail (<256B) as a 1-D tile at BYTE granularity. --------- -// Same single-buffered LDS window and the same required RAW/WAR waits as above. -__device__ inline void issueRow1d(uint64_t* srcs, uint64_t* dsts, - uint32_t numSrcs, uint32_t numDsts, - uint32_t val, uint32_t tmp, uint32_t nbytes, - uint64_t off = 0) { - gfx1250_TDM_GROUP1 g1; - g1.dataSize(DS1); // 1-byte elements: exact length - g1.tileDim0(nbytes); g1.tileDim1(1); - g1.tensorDim0(nbytes); g1.tensorDim1(1); - g1.tensorDim0Stride(nbytes); - - if (numSrcs) { - gfx1250_TDM_GROUP0 g0l(val, srcs[0] + off); // unused higher dims -> zero (see load()) - load(g0l, g1); waitTensor0(); // RAW: fill LDS before store/reduce reads it - for (size_t s = 1; s < numSrcs; s++) { - gfx1250_TDM_GROUP0 g0l(tmp, srcs[s] + off); - load(g0l, g1); waitTensor0(); - for (size_t u = 0; u < nbytes; u += Warp) { - val[u] += tmp[u]; // reduce: accumulate into val - } - } - } - - for (size_t d = 0; d < numDsts; d++) { - gfx1250_TDM_GROUP0 g0s(val, dsts[d] + off); // broadcast reduced result to each dst - store(g0s, g1); waitTensor0(); // WAR: drain store before window reuse - } -} - -__device__ inline void issue(void** dsts, const void** srcs, uint32_t numSrcs, uint32_t numDsts, - size_t sizeBytes, void* ldsBuffer, size_t ldsBufferBytes, - uint32_t startWarpId, uint32_t stopWarpId) { - const uint32_t ldsBase = static_cast(reinterpret_cast(ldsBuffer)); - const uint32_t ldsBytes = static_cast(ldsBufferBytes); // LDS is small - - const uint32_t W = warpSize; - const uint32_t nThreads = blockDim.x * blockDim.y * blockDim.z; - const uint32_t tid = (threadIdx.z * blockDim.y + threadIdx.y) * blockDim.x - + threadIdx.x; - const uint32_t warpThread = tid % W; // thread index within its warp - const uint32_t warpId = tid / W; - const uint32_t nWarps = (nThreads + W - 1) / W; - - // --- team membership: this warp participates iff in [start, stop) -------- - const uint32_t teamStop = (stopWarpId > nWarps) ? nWarps : stopWarpId; - if (startWarpId >= teamStop || warpId < startWarpId || warpId >= teamStop) - return; // not on this team - const uint32_t rank = warpId - startWarpId; // rank within the team - const uint32_t teamWarps = teamStop - startWarpId; // >= 1 - - // active threads in THIS warp (handles partial final warp); stride for vector. - const uint32_t warpThreads = (nThreads - warpId * W < W) ? (nThreads - warpId * W) : W; - - // --- split the range: [256B rows ][tail] ------------------ - size_t rows = sizeBytes / WIDTH; // whole 256B rows - size_t tail = sizeBytes % WIDTH; - size_t tdmBytes = rows * WIDTH; - - // --- base src/dst addresses for this team (byte offset applied per use) -- - uint64_t mySrcs[MAX_SRCS]; - uint64_t myDsts[MAX_DSTS]; - for (uint32_t i = 0; i < numSrcs; i++) mySrcs[i] = (uint64_t)srcs[i]; - for (uint32_t i = 0; i < numDsts; i++) myDsts[i] = (uint64_t)dsts[i]; - - // --- LDS reduce buffers: val = running sum, tmp = staging for extra srcs -- - // Base offsets (rank 0); the TDM path shifts each warp by rank*window below. - uint32_t val = ldsBase; - uint32_t tmp = ldsBase + WIDTH; - - // --- edges (team's FIRST warp = rank 0): vector head, TDM tail ----------- - if (rank == 0 && tail) { - // LDS needed: val (+ tmp when reducing multiple srcs), each holding `tail` bytes - uint32_t need = (numSrcs > 1 ? (tmp - ldsBase) : (val - ldsBase)) + tail; - if (ldsBytes >= need) { // stage tail in rank 0's window - issueRow1d(mySrcs, myDsts, numSrcs, numDsts, val, tmp, tail, tdmBytes); - } else { - warpVecCopy(mySrcs, myDsts, numSrcs, numDsts, tail, tdmBytes, - warpThread, warpThreads); - } - } - - // --- 256B rows copy via TDM ------------------------------------------------ - uint32_t maxByLds = ldsBytes / RWIDTH; // #warps we can give a window, 2*WIDTH because of double buffering - if (maxByLds == 0) { // LDS < 512B: vector fallback - if (rank == 0) - warpVecCopy(mySrcs, myDsts, numSrcs, numDsts, tdmBytes, 0, - warpThread, warpThreads); - return; - } - uint32_t issuers = teamWarps < maxByLds ? teamWarps : maxByLds; - uint32_t window = (ldsBytes / issuers) & ~(RWIDTH - 1); // per-warp 512B-multiple - uint32_t rowsPerChunk = window / RWIDTH; - - if (rank >= issuers) return; // this warp doesn't issue - - // distribute `rows` across issuers by team rank (contiguous row blocks) - size_t base = rows / issuers; - size_t extra = rows % issuers; - size_t myRows = base + (rank < extra ? 1u : 0u); - // TODO:if last warp, take the remaining edges - size_t myStart = rank * base + (rank < extra ? rank : extra); - if (myRows == 0) return; - - // shift this warp's reduce buffers into its own window - val += rank * window; - tmp += rank * window; - - for (size_t r = 0; r < myRows; r += rowsPerChunk) { - uint32_t chunkRows = (myRows - r < rowsPerChunk) - ? static_cast(myRows - r) : rowsPerChunk; - uint64_t off = (uint64_t)(myStart + r) * WIDTH; // this warp's row block - issueRows(mySrcs, myDsts, numSrcs, numDsts, val, tmp, chunkRows, off); - } -} \ No newline at end of file diff --git a/src/header/tdmCopy.h b/src/header/tdmCopy.h index 94297f36..c18c118a 100644 --- a/src/header/tdmCopy.h +++ b/src/header/tdmCopy.h @@ -354,6 +354,17 @@ constexpr uint32_t RWIDTH = 512; constexpr uint32_t MAX_SRCS = 16; // max sources reduced in one call constexpr uint32_t MAX_DSTS = 16; // max destinations broadcast to +constexpr uint8_t MEMSET_CHAR = 75; +constexpr uint32_t MEMSET_WORD = 0x4B4B4B4Bu; +constexpr float MEMSET_VAL = 13323083.0f; + +// Packed-float memset value (mirrors TransferBench.hpp's MemsetVal()). +template __device__ __forceinline__ T MemsetVal(); +template <> __device__ __forceinline__ float MemsetVal() { return MEMSET_VAL; } +template <> __device__ __forceinline__ float2 MemsetVal() { return make_float2(MEMSET_VAL, MEMSET_VAL); } +template <> __device__ __forceinline__ float4 MemsetVal() { return make_float4(MEMSET_VAL, MEMSET_VAL, + MEMSET_VAL, MEMSET_VAL); } + // ---- instruction emission (the only arch-specific piece) ------------------- // The tensor DMA is a single builtin taking the FULL descriptor: five register // groups plus a constant cache policy. Per the clang reference @@ -537,26 +548,40 @@ __device__ inline void issue(void* dst, const void* src, size_t sizeBytes, // copy; numDsts > 1 broadcasts the same reduced result to each destination. // `srcs`/`dsts` hold base addresses; `offset` (in bytes) reaches the sub-range // to copy (e.g. the tail start), so callers can share one base pointer array. +template __device__ inline void warpVecCopy(uint64_t* srcs, uint64_t* dsts, uint32_t numSrcs, uint32_t numDsts, size_t n, size_t offset, uint32_t warpThread, uint32_t warpThreads) { - size_t nd = n >> 2; - for (size_t i = warpThread; i < nd; i += warpThreads) { - uint32_t acc = 0; + // Bulk: reduce whole PACKED_FLOAT elements. + size_t np = n / sizeof(PACKED_FLOAT); + for (size_t i = warpThread; i < np; i += warpThreads) { + PACKED_FLOAT acc = numSrcs ? PACKED_FLOAT{} : MemsetVal(); for (uint32_t s = 0; s < numSrcs; ++s) - acc += reinterpret_cast(srcs[s] + offset)[i]; + acc += reinterpret_cast(srcs[s] + offset)[i]; for (uint32_t d = 0; d < numDsts; ++d) - reinterpret_cast(dsts[d] + offset)[i] = acc; + reinterpret_cast(dsts[d] + offset)[i] = acc; } + // Remaining whole floats (n not a PACKED_FLOAT multiple); empty when + // PACKED_FLOAT == float. + size_t nf = n >> 2; + size_t fStart = (np * sizeof(PACKED_FLOAT)) >> 2; + for (size_t i = fStart + warpThread; i < nf; i += warpThreads) { + float acc = numSrcs ? 0.f : MemsetVal(); + for (uint32_t s = 0; s < numSrcs; ++s) + acc += reinterpret_cast(srcs[s] + offset)[i]; + for (uint32_t d = 0; d < numDsts; ++d) + reinterpret_cast(dsts[d] + offset)[i] = acc; + } + // Ragged sub-float tail (< 4 bytes). uint32_t rem = static_cast(n & 3u); if (rem && warpThread == 0) { for (uint32_t b = 0; b < rem; ++b) { - uint8_t acc = 0; + uint8_t acc = numSrcs ? 0 : MEMSET_CHAR; for (uint32_t s = 0; s < numSrcs; ++s) - acc += reinterpret_cast(srcs[s] + offset)[nd * 4 + b]; + acc += reinterpret_cast(srcs[s] + offset)[nf * 4 + b]; for (uint32_t d = 0; d < numDsts; ++d) - reinterpret_cast(dsts[d] + offset)[nd * 4 + b] = acc; + reinterpret_cast(dsts[d] + offset)[nf * 4 + b] = acc; } } } @@ -565,9 +590,11 @@ __device__ inline void warpVecCopy(uint64_t* srcs, uint64_t* dsts, // Same single-buffered LDS window / RAW+WAR waits as the single-copy issueRows(). // The first source lands in `val`; each subsequent source lands in `tmp` and is // accumulated into `val`; the reduced `val` is stored to every destination. +template __device__ inline void issueRows(uint64_t* srcs, uint64_t* dsts, uint32_t numSrcs, uint32_t numDsts, - uint32_t val, uint32_t tmp, uint32_t rows, uint32_t off = 0) { + uint32_t val, uint32_t tmp, uint32_t rows, + uint32_t warpThread, uint32_t off = 0) { gfx1250_TDM_GROUP1 g1; g1.dataSize(DS4); g1.tileDim0(TD0); g1.tileDim1(rows); @@ -577,17 +604,23 @@ __device__ inline void issueRows(uint64_t* srcs, uint64_t* dsts, if (numSrcs) { gfx1250_TDM_GROUP0 g0l(val, srcs[0] + off); load(g0l, g1); waitTensor0(); - // LDS byte-offsets -> typed shared pointers for the reduce accumulation - uint32_t* vp = reinterpret_cast(static_cast(val)); - const uint32_t* tp = reinterpret_cast(static_cast(tmp)); - uint32_t nElems = (rows * WIDTH) >> 2; // 4-byte (DS4) elements in this tile + PACKED_FLOAT* vp = reinterpret_cast(static_cast(val)); + const PACKED_FLOAT* tp = reinterpret_cast(static_cast(tmp)); + uint32_t nElems = (rows * WIDTH) / sizeof(PACKED_FLOAT); for (size_t s = 1; s < numSrcs; s++) { gfx1250_TDM_GROUP0 g0l(tmp, srcs[s] + off); load(g0l, g1); waitTensor0(); - for (uint32_t u = 0; u < nElems; u += warpSize) { + for (uint32_t u = warpThread; u < nElems; u += warpSize) { vp[u] += tp[u]; } } + } else { + // Empty source: fill this warp's reduce window with the MEMSET_CHAR byte + // pattern so the TDM store below writes a memset() result to each dst + // (mirrors GpuReduceKernel's numSrcs==0 path). No load is issued. + uint32_t* vp = reinterpret_cast(static_cast(val)); + uint32_t nWords = (rows * WIDTH) / sizeof(uint32_t); + for (uint32_t u = warpThread; u < nWords; u += warpSize) vp[u] = MEMSET_WORD; } for (size_t d = 0; d < numDsts; d++) { @@ -598,10 +631,11 @@ __device__ inline void issueRows(uint64_t* srcs, uint64_t* dsts, // ---- issue a sub-row tail (<256B) as a 1-D tile at BYTE granularity. --------- // Same single-buffered LDS window and the same required RAW/WAR waits as above. +template __device__ inline void issueRow1d(uint64_t* srcs, uint64_t* dsts, uint32_t numSrcs, uint32_t numDsts, uint32_t val, uint32_t tmp, uint32_t nbytes, - uint64_t off = 0) { + uint32_t warpThread, uint64_t off = 0) { gfx1250_TDM_GROUP1 g1; g1.dataSize(DS1); // 1-byte elements: exact length g1.tileDim0(nbytes); g1.tileDim1(1); @@ -612,14 +646,32 @@ __device__ inline void issueRow1d(uint64_t* srcs, uint64_t* dsts, gfx1250_TDM_GROUP0 g0l(val, srcs[0] + off); // unused higher dims -> zero (see load()) load(g0l, g1); waitTensor0(); // RAW: fill LDS before store/reduce reads it // LDS byte-offsets -> typed shared pointers for the reduce accumulation - uint8_t* vp = reinterpret_cast(static_cast(val)); - const uint8_t* tp = reinterpret_cast(static_cast(tmp)); + PACKED_FLOAT* vp = reinterpret_cast(static_cast(val)); + const PACKED_FLOAT* tp = reinterpret_cast(static_cast(tmp)); + uint32_t nP = nbytes / sizeof(PACKED_FLOAT); // whole PACKED_FLOAT elements + // Remaining whole floats (nbytes not a PACKED_FLOAT multiple); float data is + // 4-byte aligned so the DS1 tail is always a whole number of floats. + float* vpf = reinterpret_cast(static_cast(val)); + const float* tpf = reinterpret_cast(static_cast(tmp)); + uint32_t nF = nbytes >> 2; + uint32_t fStart = nP * (sizeof(PACKED_FLOAT) >> 2); for (size_t s = 1; s < numSrcs; s++) { gfx1250_TDM_GROUP0 g0l(tmp, srcs[s] + off); load(g0l, g1); waitTensor0(); - for (uint32_t u = 0; u < nbytes; u += warpSize) { // 1-byte (DS1) elements - vp[u] += tp[u]; // reduce: accumulate into val - } + for (uint32_t u = warpThread; u < nP; u += warpSize) vp[u] += tp[u]; // PACKED_FLOAT bulk + for (uint32_t u = fStart + warpThread; u < nF; u += warpSize) vpf[u] += tpf[u]; // float remainder + } + } else { + // Empty source: fill this warp's tail window with the MEMSET_CHAR byte + // pattern (memset semantics), then the store below broadcasts it to each + // dst. No load is issued. + uint32_t* vp = reinterpret_cast(static_cast(val)); + uint32_t nWords = nbytes / sizeof(uint32_t); + for (uint32_t u = warpThread; u < nWords; u += warpSize) vp[u] = MEMSET_WORD; + uint32_t rem = nbytes & 3u; // sub-word remainder (rare for float data) + if (rem && warpThread == 0) { + uint8_t* vb = reinterpret_cast(static_cast(val)); + for (uint32_t b = 0; b < rem; ++b) vb[nWords * 4 + b] = MEMSET_CHAR; } } @@ -630,6 +682,7 @@ __device__ inline void issueRow1d(uint64_t* srcs, uint64_t* dsts, } // ---- core: partition + issue a reduce/broadcast for the team [start, stop). - +template __device__ inline void issue(void** dsts, const void** srcs, uint32_t numSrcs, uint32_t numDsts, size_t sizeBytes, void* ldsBuffer, size_t ldsBufferBytes, uint32_t startWarpId, uint32_t stopWarpId) { @@ -675,9 +728,9 @@ __device__ inline void issue(void** dsts, const void** srcs, uint32_t numSrcs, u // LDS needed: val (+ tmp when reducing multiple srcs), each holding `tail` bytes uint32_t need = (numSrcs > 1 ? (tmp - ldsBase) : (val - ldsBase)) + tail; if (ldsBytes >= need) { // stage tail in rank 0's window - issueRow1d(mySrcs, myDsts, numSrcs, numDsts, val, tmp, tail, tdmBytes); + issueRow1d(mySrcs, myDsts, numSrcs, numDsts, val, tmp, tail, warpThread, tdmBytes); } else { - warpVecCopy(mySrcs, myDsts, numSrcs, numDsts, tail, tdmBytes, + warpVecCopy(mySrcs, myDsts, numSrcs, numDsts, tail, tdmBytes, warpThread, warpThreads); } } @@ -686,7 +739,7 @@ __device__ inline void issue(void** dsts, const void** srcs, uint32_t numSrcs, u uint32_t maxByLds = ldsBytes / RWIDTH; // #warps we can give a window, 2*WIDTH because of double buffering if (maxByLds == 0) { // LDS < 512B: vector fallback if (rank == 0) - warpVecCopy(mySrcs, myDsts, numSrcs, numDsts, tdmBytes, 0, + warpVecCopy(mySrcs, myDsts, numSrcs, numDsts, tdmBytes, 0, warpThread, warpThreads); return; } @@ -712,7 +765,7 @@ __device__ inline void issue(void** dsts, const void** srcs, uint32_t numSrcs, u uint32_t chunkRows = (myRows - r < rowsPerChunk) ? static_cast(myRows - r) : rowsPerChunk; uint64_t off = (uint64_t)(myStart + r) * WIDTH; // this warp's row block - issueRows(mySrcs, myDsts, numSrcs, numDsts, val, tmp, chunkRows, off); + issueRows(mySrcs, myDsts, numSrcs, numDsts, val, tmp, chunkRows, warpThread, off); } } @@ -910,9 +963,10 @@ __device__ inline void tdmCopyByTeam(void* dst, const void* src, size_t sizeByte } #if TDM_BACKEND_AMD // multi-source reduce path only exists on the AMD (TDM) backend +template __device__ inline void tdmReduce(void** dsts, const void** srcs, uint32_t numSrcs, uint32_t numDsts, size_t sizeBytes, void* ldsBuffer, size_t ldsBufferBytes) { - detail::issue(dsts, srcs, numSrcs, numDsts, sizeBytes, ldsBuffer, ldsBufferBytes, /*start=*/0, /*stop=*/~0u); + detail::issue(dsts, srcs, numSrcs, numDsts, sizeBytes, ldsBuffer, ldsBufferBytes, /*start=*/0, /*stop=*/~0u); tdmWait(); } #endif // TDM_BACKEND_AMD From d6e346d87d6b26f9c021cf86242f19f98bd8a055 Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Mon, 10 Aug 2026 06:48:20 +0000 Subject: [PATCH 5/7] temporary fix for LDS access in reduction loop --- src/header/TransferBench.hpp | 4 +++- src/header/tdmCopy.h | 43 ++++++++++++++++++++++++++---------- 2 files changed, 34 insertions(+), 13 deletions(-) diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index 52627be4..3e5f989d 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -5970,7 +5970,9 @@ namespace { int subIterations = 0; while (1) { - tdm::tdmReduce(p.dst, p.src, numSrcs, numDsts, sizeBytes, shmem, ldsBytes); + tdm::tdmReduce(reinterpret_cast(p.dst), + const_cast(reinterpret_cast(p.src)), + numSrcs, numDsts, sizeBytes, shmem, ldsBytes); __syncthreads(); // Wait for all warps to finish this subiteration if (++subIterations == numSubIterations) break; } diff --git a/src/header/tdmCopy.h b/src/header/tdmCopy.h index c18c118a..0ef3cbc0 100644 --- a/src/header/tdmCopy.h +++ b/src/header/tdmCopy.h @@ -586,6 +586,22 @@ __device__ inline void warpVecCopy(uint64_t* srcs, uint64_t* dsts, } } +// ---- LDS (address space 3) pointer helper -------------------------------- +// `val`/`tmp` are 32-bit LDS byte offsets (exactly what the TDM load/store +// builtins consume). To touch that staging memory with ordinary scalar +// loads/stores we must build a pointer TAGGED as LDS (address space 3). A plain +// reinterpret_cast yields a generic/flat pointer whose numeric value lands +// in the GLOBAL aperture (an LDS offset like 0x19000 is a valid global VA), so +// dereferencing it faults. Casting to an address_space(3) pointer makes the +// compiler emit ds_* (LDS) accesses against the offset instead. +template +using LdsPtr = T __attribute__((address_space(3)))*; + +template +__device__ inline LdsPtr ldsCast(uint32_t off) { + return reinterpret_cast>(off); +} + // ---- issue one chunk of whole 256B rows (2D tile) through ONE LDS window. --- // Same single-buffered LDS window / RAW+WAR waits as the single-copy issueRows(). // The first source lands in `val`; each subsequent source lands in `tmp` and is @@ -604,8 +620,8 @@ __device__ inline void issueRows(uint64_t* srcs, uint64_t* dsts, if (numSrcs) { gfx1250_TDM_GROUP0 g0l(val, srcs[0] + off); load(g0l, g1); waitTensor0(); - PACKED_FLOAT* vp = reinterpret_cast(static_cast(val)); - const PACKED_FLOAT* tp = reinterpret_cast(static_cast(tmp)); + LdsPtr vp = ldsCast(val); + LdsPtr tp = ldsCast(tmp); uint32_t nElems = (rows * WIDTH) / sizeof(PACKED_FLOAT); for (size_t s = 1; s < numSrcs; s++) { gfx1250_TDM_GROUP0 g0l(tmp, srcs[s] + off); @@ -618,7 +634,7 @@ __device__ inline void issueRows(uint64_t* srcs, uint64_t* dsts, // Empty source: fill this warp's reduce window with the MEMSET_CHAR byte // pattern so the TDM store below writes a memset() result to each dst // (mirrors GpuReduceKernel's numSrcs==0 path). No load is issued. - uint32_t* vp = reinterpret_cast(static_cast(val)); + LdsPtr vp = ldsCast(val); uint32_t nWords = (rows * WIDTH) / sizeof(uint32_t); for (uint32_t u = warpThread; u < nWords; u += warpSize) vp[u] = MEMSET_WORD; } @@ -646,13 +662,13 @@ __device__ inline void issueRow1d(uint64_t* srcs, uint64_t* dsts, gfx1250_TDM_GROUP0 g0l(val, srcs[0] + off); // unused higher dims -> zero (see load()) load(g0l, g1); waitTensor0(); // RAW: fill LDS before store/reduce reads it // LDS byte-offsets -> typed shared pointers for the reduce accumulation - PACKED_FLOAT* vp = reinterpret_cast(static_cast(val)); - const PACKED_FLOAT* tp = reinterpret_cast(static_cast(tmp)); + LdsPtr vp = ldsCast(val); + LdsPtr tp = ldsCast(tmp); uint32_t nP = nbytes / sizeof(PACKED_FLOAT); // whole PACKED_FLOAT elements // Remaining whole floats (nbytes not a PACKED_FLOAT multiple); float data is // 4-byte aligned so the DS1 tail is always a whole number of floats. - float* vpf = reinterpret_cast(static_cast(val)); - const float* tpf = reinterpret_cast(static_cast(tmp)); + LdsPtr vpf = ldsCast(val); + LdsPtr tpf = ldsCast(tmp); uint32_t nF = nbytes >> 2; uint32_t fStart = nP * (sizeof(PACKED_FLOAT) >> 2); for (size_t s = 1; s < numSrcs; s++) { @@ -665,12 +681,12 @@ __device__ inline void issueRow1d(uint64_t* srcs, uint64_t* dsts, // Empty source: fill this warp's tail window with the MEMSET_CHAR byte // pattern (memset semantics), then the store below broadcasts it to each // dst. No load is issued. - uint32_t* vp = reinterpret_cast(static_cast(val)); + LdsPtr vp = ldsCast(val); uint32_t nWords = nbytes / sizeof(uint32_t); for (uint32_t u = warpThread; u < nWords; u += warpSize) vp[u] = MEMSET_WORD; uint32_t rem = nbytes & 3u; // sub-word remainder (rare for float data) if (rem && warpThread == 0) { - uint8_t* vb = reinterpret_cast(static_cast(val)); + LdsPtr vb = ldsCast(val); for (uint32_t b = 0; b < rem; ++b) vb[nWords * 4 + b] = MEMSET_CHAR; } } @@ -757,9 +773,12 @@ __device__ inline void issue(void** dsts, const void** srcs, uint32_t numSrcs, u size_t myStart = rank * base + (rank < extra ? rank : extra); if (myRows == 0) return; - // shift this warp's reduce buffers into its own window - val += rank * window; - tmp += rank * window; + // Give this warp its own window and split it evenly into two halves: + // val = running sum (first half), tmp = staging for extra srcs (second half). + // window is a multiple of RWIDTH (= 2*WIDTH), so window/2 is a WIDTH-multiple + // large enough to hold rowsPerChunk (= window/RWIDTH) rows in each half. + val = ldsBase + rank * window; + tmp = val + window / 2; for (size_t r = 0; r < myRows; r += rowsPerChunk) { uint32_t chunkRows = (myRows - r < rowsPerChunk) From ecb3b520a74fdd37ea227bbfae605c6fc7fe0e80 Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Wed, 12 Aug 2026 16:51:10 +0000 Subject: [PATCH 6/7] fix0.5 --- src/header/tdmCopy.h | 61 +++++++++++++++++++++++++++++--------------- 1 file changed, 40 insertions(+), 21 deletions(-) diff --git a/src/header/tdmCopy.h b/src/header/tdmCopy.h index 0ef3cbc0..9b7c8933 100644 --- a/src/header/tdmCopy.h +++ b/src/header/tdmCopy.h @@ -594,12 +594,29 @@ __device__ inline void warpVecCopy(uint64_t* srcs, uint64_t* dsts, // in the GLOBAL aperture (an LDS offset like 0x19000 is a valid global VA), so // dereferencing it faults. Casting to an address_space(3) pointer makes the // compiler emit ds_* (LDS) accesses against the offset instead. +// ---- LDS pointer from the real base pointer + absolute byte offset -------- +// `val`/`tmp` are ABSOLUTE LDS byte offsets (ldsBase + per-warp shift), exactly +// what the TDM load/store builtins consume. To touch that staging memory with +// ordinary scalar loads/stores we must use a pointer that carries the correct +// LDS aperture. Building one from a bare integer offset does NOT (it resolves to +// the GLOBAL aperture and reads/writes don't alias the tensor-written LDS); +// instead we offset from `ldsMem`, the real generic pointer to this block's LDS. template -using LdsPtr = T __attribute__((address_space(3)))*; +__device__ inline T* ldsPtr(void* ldsMem, uint32_t absOff) { + uint32_t base = static_cast(reinterpret_cast(ldsMem)); + return reinterpret_cast(static_cast(ldsMem) + (absOff - base)); +} -template -__device__ inline LdsPtr ldsCast(uint32_t off) { - return reinterpret_cast>(off); +// ---- LDS visibility fence between the tensor and vector memory pipes ------- +// s_wait_tensorcnt only orders tensor-op vs tensor-op, so it makes the pure copy +// path (load->wait->store) correct but does NOT make the tensor engine's LDS +// writes visible to this wave's vector (ds) reads, nor the ds writes visible to a +// following tensor store. The reduce path interleaves ds reads/writes between the +// tensor load and store, so it needs an explicit workgroup-scope fence at those +// boundaries. (A __syncthreads() barrier is unusable here: warps with +// rank >= issuers or myRows == 0 return early and would deadlock it.) +__device__ __forceinline__ void ldsFence() { + __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "workgroup"); } // ---- issue one chunk of whole 256B rows (2D tile) through ONE LDS window. --- @@ -609,7 +626,7 @@ __device__ inline LdsPtr ldsCast(uint32_t off) { template __device__ inline void issueRows(uint64_t* srcs, uint64_t* dsts, uint32_t numSrcs, uint32_t numDsts, - uint32_t val, uint32_t tmp, uint32_t rows, + uint32_t val, uint32_t tmp, void* ldsMem, uint32_t rows, uint32_t warpThread, uint32_t off = 0) { gfx1250_TDM_GROUP1 g1; g1.dataSize(DS4); @@ -619,13 +636,13 @@ __device__ inline void issueRows(uint64_t* srcs, uint64_t* dsts, if (numSrcs) { gfx1250_TDM_GROUP0 g0l(val, srcs[0] + off); - load(g0l, g1); waitTensor0(); - LdsPtr vp = ldsCast(val); - LdsPtr tp = ldsCast(tmp); + load(g0l, g1); waitTensor0(); ldsFence(); // src0 LDS write visible to ds reads + PACKED_FLOAT* vp = ldsPtr(ldsMem, val); + const PACKED_FLOAT* tp = ldsPtr(ldsMem, tmp); uint32_t nElems = (rows * WIDTH) / sizeof(PACKED_FLOAT); for (size_t s = 1; s < numSrcs; s++) { gfx1250_TDM_GROUP0 g0l(tmp, srcs[s] + off); - load(g0l, g1); waitTensor0(); + load(g0l, g1); waitTensor0(); ldsFence(); // srcS LDS write visible to ds reads for (uint32_t u = warpThread; u < nElems; u += warpSize) { vp[u] += tp[u]; } @@ -634,11 +651,12 @@ __device__ inline void issueRows(uint64_t* srcs, uint64_t* dsts, // Empty source: fill this warp's reduce window with the MEMSET_CHAR byte // pattern so the TDM store below writes a memset() result to each dst // (mirrors GpuReduceKernel's numSrcs==0 path). No load is issued. - LdsPtr vp = ldsCast(val); + uint32_t* vp = ldsPtr(ldsMem, val); uint32_t nWords = (rows * WIDTH) / sizeof(uint32_t); for (uint32_t u = warpThread; u < nWords; u += warpSize) vp[u] = MEMSET_WORD; } + ldsFence(); // ds writes (reduced/memset result) visible to the tensor store for (size_t d = 0; d < numDsts; d++) { gfx1250_TDM_GROUP0 g0s(val, dsts[d] + off); // broadcast reduced result to each dst store(g0s, g1); waitTensor0(); @@ -650,7 +668,7 @@ __device__ inline void issueRows(uint64_t* srcs, uint64_t* dsts, template __device__ inline void issueRow1d(uint64_t* srcs, uint64_t* dsts, uint32_t numSrcs, uint32_t numDsts, - uint32_t val, uint32_t tmp, uint32_t nbytes, + uint32_t val, uint32_t tmp, void* ldsMem, uint32_t nbytes, uint32_t warpThread, uint64_t off = 0) { gfx1250_TDM_GROUP1 g1; g1.dataSize(DS1); // 1-byte elements: exact length @@ -660,20 +678,20 @@ __device__ inline void issueRow1d(uint64_t* srcs, uint64_t* dsts, if (numSrcs) { gfx1250_TDM_GROUP0 g0l(val, srcs[0] + off); // unused higher dims -> zero (see load()) - load(g0l, g1); waitTensor0(); // RAW: fill LDS before store/reduce reads it + load(g0l, g1); waitTensor0(); ldsFence(); // src0 LDS write visible to ds reads // LDS byte-offsets -> typed shared pointers for the reduce accumulation - LdsPtr vp = ldsCast(val); - LdsPtr tp = ldsCast(tmp); + PACKED_FLOAT* vp = ldsPtr(ldsMem, val); + const PACKED_FLOAT* tp = ldsPtr(ldsMem, tmp); uint32_t nP = nbytes / sizeof(PACKED_FLOAT); // whole PACKED_FLOAT elements // Remaining whole floats (nbytes not a PACKED_FLOAT multiple); float data is // 4-byte aligned so the DS1 tail is always a whole number of floats. - LdsPtr vpf = ldsCast(val); - LdsPtr tpf = ldsCast(tmp); + float* vpf = ldsPtr(ldsMem, val); + const float* tpf = ldsPtr(ldsMem, tmp); uint32_t nF = nbytes >> 2; uint32_t fStart = nP * (sizeof(PACKED_FLOAT) >> 2); for (size_t s = 1; s < numSrcs; s++) { gfx1250_TDM_GROUP0 g0l(tmp, srcs[s] + off); - load(g0l, g1); waitTensor0(); + load(g0l, g1); waitTensor0(); ldsFence(); // srcS LDS write visible to ds reads for (uint32_t u = warpThread; u < nP; u += warpSize) vp[u] += tp[u]; // PACKED_FLOAT bulk for (uint32_t u = fStart + warpThread; u < nF; u += warpSize) vpf[u] += tpf[u]; // float remainder } @@ -681,16 +699,17 @@ __device__ inline void issueRow1d(uint64_t* srcs, uint64_t* dsts, // Empty source: fill this warp's tail window with the MEMSET_CHAR byte // pattern (memset semantics), then the store below broadcasts it to each // dst. No load is issued. - LdsPtr vp = ldsCast(val); + uint32_t* vp = ldsPtr(ldsMem, val); uint32_t nWords = nbytes / sizeof(uint32_t); for (uint32_t u = warpThread; u < nWords; u += warpSize) vp[u] = MEMSET_WORD; uint32_t rem = nbytes & 3u; // sub-word remainder (rare for float data) if (rem && warpThread == 0) { - LdsPtr vb = ldsCast(val); + uint8_t* vb = ldsPtr(ldsMem, val); for (uint32_t b = 0; b < rem; ++b) vb[nWords * 4 + b] = MEMSET_CHAR; } } + ldsFence(); // ds writes (reduced/memset result) visible to the tensor store for (size_t d = 0; d < numDsts; d++) { gfx1250_TDM_GROUP0 g0s(val, dsts[d] + off); // broadcast reduced result to each dst store(g0s, g1); waitTensor0(); // WAR: drain store before window reuse @@ -744,7 +763,7 @@ __device__ inline void issue(void** dsts, const void** srcs, uint32_t numSrcs, u // LDS needed: val (+ tmp when reducing multiple srcs), each holding `tail` bytes uint32_t need = (numSrcs > 1 ? (tmp - ldsBase) : (val - ldsBase)) + tail; if (ldsBytes >= need) { // stage tail in rank 0's window - issueRow1d(mySrcs, myDsts, numSrcs, numDsts, val, tmp, tail, warpThread, tdmBytes); + issueRow1d(mySrcs, myDsts, numSrcs, numDsts, val, tmp, ldsBuffer, tail, warpThread, tdmBytes); } else { warpVecCopy(mySrcs, myDsts, numSrcs, numDsts, tail, tdmBytes, warpThread, warpThreads); @@ -784,7 +803,7 @@ __device__ inline void issue(void** dsts, const void** srcs, uint32_t numSrcs, u uint32_t chunkRows = (myRows - r < rowsPerChunk) ? static_cast(myRows - r) : rowsPerChunk; uint64_t off = (uint64_t)(myStart + r) * WIDTH; // this warp's row block - issueRows(mySrcs, myDsts, numSrcs, numDsts, val, tmp, chunkRows, warpThread, off); + issueRows(mySrcs, myDsts, numSrcs, numDsts, val, tmp, ldsBuffer, chunkRows, warpThread, off); } } From eb2af9da5444dc2a192f83fc54c8374a91ea86d7 Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Wed, 12 Aug 2026 16:51:27 +0000 Subject: [PATCH 7/7] fix --- src/header/tdmCopy.h | 138 ++++++++++++++++++------------------------- 1 file changed, 57 insertions(+), 81 deletions(-) diff --git a/src/header/tdmCopy.h b/src/header/tdmCopy.h index 9b7c8933..e23fa06c 100644 --- a/src/header/tdmCopy.h +++ b/src/header/tdmCopy.h @@ -543,13 +543,25 @@ __device__ inline void issue(void* dst, const void* src, size_t sizeBytes, // are overloads of the single-copy helpers above; the reduce entry point routes // here via detail::issue(dsts, srcs, numSrcs, numDsts, ...). +// ---- optimization guard for the multi-source reduce path -------------------- +// The gfx1250 tensor load/store builtins are memory(inaccessiblemem): the backend +// does NOT model them as touching the LDS that the vector (ds) reduce reads/writes. +// At -O2/-O3 the scheduler/inliner exploits that missing dependency and miscompiles +// the reduce (drops the tensor store, or stores stale LDS) in ways that flip with +// unrelated codegen changes (inlining, register pressure). The generated code is +// correct at -O0/-O1, so the whole reduce path is pinned to no-optimization; the +// scalar overhead here is negligible (bandwidth comes from the TDM engine / ds +// pipe), and it keeps the result deterministically correct regardless of the TU's +// optimization level. Remove once the backend models the tensor<->LDS dependency. +#define TDM_REDUCE_NOOPT __attribute__((optnone, noinline)) + // ---- vector fallback: reduce (sum) numSrcs sources, broadcast to numDsts. ---- // Element-wise sum of all srcs -> written to every dst. numSrcs == 1 is a plain // copy; numDsts > 1 broadcasts the same reduced result to each destination. // `srcs`/`dsts` hold base addresses; `offset` (in bytes) reaches the sub-range // to copy (e.g. the tail start), so callers can share one base pointer array. template -__device__ inline void warpVecCopy(uint64_t* srcs, uint64_t* dsts, +TDM_REDUCE_NOOPT __device__ void warpVecCopy(uint64_t* srcs, uint64_t* dsts, uint32_t numSrcs, uint32_t numDsts, size_t n, size_t offset, uint32_t warpThread, uint32_t warpThreads) { @@ -607,6 +619,21 @@ __device__ inline T* ldsPtr(void* ldsMem, uint32_t absOff) { return reinterpret_cast(static_cast(ldsMem) + (absOff - base)); } +// ---- explicit LDS (address_space(3)) pointer from a segment byte offset ------- +// The generic `ldsPtr` above resolves fine, but marking the result `volatile` +// (needed to stop the compiler DCE'ing the reduce RMW, since the tensor store +// intrinsic is not modeled as an LDS reader) blocks address-space inference and +// forces `flat_*` lowering. `flat_*` routes through the small FLAT LDS aperture, +// so large per-warp offsets fault (HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION). +// An explicit address_space(3) pointer always lowers to `ds_*`, which uses a +// 32-bit LDS offset and reaches the full allocated LDS. `absOff` is already the +// 0-based LDS segment offset (low 32 bits of the block's flat LDS address). +template using LdsPtrT = T __attribute__((address_space(3)))*; +template +__device__ inline LdsPtrT ldsPtr3(uint32_t absOff) { + return reinterpret_cast>(static_cast(absOff)); +} + // ---- LDS visibility fence between the tensor and vector memory pipes ------- // s_wait_tensorcnt only orders tensor-op vs tensor-op, so it makes the pure copy // path (load->wait->store) correct but does NOT make the tensor engine's LDS @@ -624,7 +651,7 @@ __device__ __forceinline__ void ldsFence() { // The first source lands in `val`; each subsequent source lands in `tmp` and is // accumulated into `val`; the reduced `val` is stored to every destination. template -__device__ inline void issueRows(uint64_t* srcs, uint64_t* dsts, +TDM_REDUCE_NOOPT __device__ void issueRows(uint64_t* srcs, uint64_t* dsts, uint32_t numSrcs, uint32_t numDsts, uint32_t val, uint32_t tmp, void* ldsMem, uint32_t rows, uint32_t warpThread, uint32_t off = 0) { @@ -634,24 +661,27 @@ __device__ inline void issueRows(uint64_t* srcs, uint64_t* dsts, g1.tensorDim0(TD0); g1.tensorDim1(rows); g1.tensorDim0Stride(TD0); // rows back-to-back (contiguous) + uint32_t nElems = (rows * WIDTH) / sizeof(PACKED_FLOAT); if (numSrcs) { + // Seed the running-sum window `val` with src0. gfx1250_TDM_GROUP0 g0l(val, srcs[0] + off); - load(g0l, g1); waitTensor0(); ldsFence(); // src0 LDS write visible to ds reads - PACKED_FLOAT* vp = ldsPtr(ldsMem, val); - const PACKED_FLOAT* tp = ldsPtr(ldsMem, tmp); - uint32_t nElems = (rows * WIDTH) / sizeof(PACKED_FLOAT); + load(g0l, g1); + waitTensor0(); ldsFence(); // RAW: src0 LDS write visible to vp reads + LdsPtrT vp = ldsPtr3(val); for (size_t s = 1; s < numSrcs; s++) { - gfx1250_TDM_GROUP0 g0l(tmp, srcs[s] + off); - load(g0l, g1); waitTensor0(); ldsFence(); // srcS LDS write visible to ds reads - for (uint32_t u = warpThread; u < nElems; u += warpSize) { - vp[u] += tp[u]; - } + gfx1250_TDM_GROUP0 g0lt(tmp, srcs[s] + off); // stage next source in `tmp` + load(g0lt, g1); + waitTensor0(); ldsFence(); // RAW: srcS LDS write visible to tp reads + LdsPtrT tp = ldsPtr3(tmp); + for (uint32_t u = warpThread; u < nElems; u += warpSize) + vp[u] = vp[u] + tp[u]; + ldsFence(); // WAR: accumulate done before next reload of tmp } } else { - // Empty source: fill this warp's reduce window with the MEMSET_CHAR byte - // pattern so the TDM store below writes a memset() result to each dst - // (mirrors GpuReduceKernel's numSrcs==0 path). No load is issued. - uint32_t* vp = ldsPtr(ldsMem, val); + // Empty source: fill this warp's reduce window with the MEMSET byte pattern + // so the TDM store writes a memset() result to each dst (mirrors + // GpuReduceKernel's numSrcs==0 path). No load is issued. + LdsPtrT vp = ldsPtr3(val); uint32_t nWords = (rows * WIDTH) / sizeof(uint32_t); for (uint32_t u = warpThread; u < nWords; u += warpSize) vp[u] = MEMSET_WORD; } @@ -659,66 +689,13 @@ __device__ inline void issueRows(uint64_t* srcs, uint64_t* dsts, ldsFence(); // ds writes (reduced/memset result) visible to the tensor store for (size_t d = 0; d < numDsts; d++) { gfx1250_TDM_GROUP0 g0s(val, dsts[d] + off); // broadcast reduced result to each dst - store(g0s, g1); waitTensor0(); - } -} - -// ---- issue a sub-row tail (<256B) as a 1-D tile at BYTE granularity. --------- -// Same single-buffered LDS window and the same required RAW/WAR waits as above. -template -__device__ inline void issueRow1d(uint64_t* srcs, uint64_t* dsts, - uint32_t numSrcs, uint32_t numDsts, - uint32_t val, uint32_t tmp, void* ldsMem, uint32_t nbytes, - uint32_t warpThread, uint64_t off = 0) { - gfx1250_TDM_GROUP1 g1; - g1.dataSize(DS1); // 1-byte elements: exact length - g1.tileDim0(nbytes); g1.tileDim1(1); - g1.tensorDim0(nbytes); g1.tensorDim1(1); - g1.tensorDim0Stride(nbytes); - - if (numSrcs) { - gfx1250_TDM_GROUP0 g0l(val, srcs[0] + off); // unused higher dims -> zero (see load()) - load(g0l, g1); waitTensor0(); ldsFence(); // src0 LDS write visible to ds reads - // LDS byte-offsets -> typed shared pointers for the reduce accumulation - PACKED_FLOAT* vp = ldsPtr(ldsMem, val); - const PACKED_FLOAT* tp = ldsPtr(ldsMem, tmp); - uint32_t nP = nbytes / sizeof(PACKED_FLOAT); // whole PACKED_FLOAT elements - // Remaining whole floats (nbytes not a PACKED_FLOAT multiple); float data is - // 4-byte aligned so the DS1 tail is always a whole number of floats. - float* vpf = ldsPtr(ldsMem, val); - const float* tpf = ldsPtr(ldsMem, tmp); - uint32_t nF = nbytes >> 2; - uint32_t fStart = nP * (sizeof(PACKED_FLOAT) >> 2); - for (size_t s = 1; s < numSrcs; s++) { - gfx1250_TDM_GROUP0 g0l(tmp, srcs[s] + off); - load(g0l, g1); waitTensor0(); ldsFence(); // srcS LDS write visible to ds reads - for (uint32_t u = warpThread; u < nP; u += warpSize) vp[u] += tp[u]; // PACKED_FLOAT bulk - for (uint32_t u = fStart + warpThread; u < nF; u += warpSize) vpf[u] += tpf[u]; // float remainder - } - } else { - // Empty source: fill this warp's tail window with the MEMSET_CHAR byte - // pattern (memset semantics), then the store below broadcasts it to each - // dst. No load is issued. - uint32_t* vp = ldsPtr(ldsMem, val); - uint32_t nWords = nbytes / sizeof(uint32_t); - for (uint32_t u = warpThread; u < nWords; u += warpSize) vp[u] = MEMSET_WORD; - uint32_t rem = nbytes & 3u; // sub-word remainder (rare for float data) - if (rem && warpThread == 0) { - uint8_t* vb = ldsPtr(ldsMem, val); - for (uint32_t b = 0; b < rem; ++b) vb[nWords * 4 + b] = MEMSET_CHAR; - } - } - - ldsFence(); // ds writes (reduced/memset result) visible to the tensor store - for (size_t d = 0; d < numDsts; d++) { - gfx1250_TDM_GROUP0 g0s(val, dsts[d] + off); // broadcast reduced result to each dst - store(g0s, g1); waitTensor0(); // WAR: drain store before window reuse + store(g0s, g1); waitTensor0(); // WAR: drain store before window reuse / next dst } } // ---- core: partition + issue a reduce/broadcast for the team [start, stop). - template -__device__ inline void issue(void** dsts, const void** srcs, uint32_t numSrcs, uint32_t numDsts, +TDM_REDUCE_NOOPT __device__ void issue(void** dsts, const void** srcs, uint32_t numSrcs, uint32_t numDsts, size_t sizeBytes, void* ldsBuffer, size_t ldsBufferBytes, uint32_t startWarpId, uint32_t stopWarpId) { const uint32_t ldsBase = static_cast(reinterpret_cast(ldsBuffer)); @@ -758,17 +735,16 @@ __device__ inline void issue(void** dsts, const void** srcs, uint32_t numSrcs, u uint32_t val = ldsBase; uint32_t tmp = ldsBase + WIDTH; - // --- edges (team's FIRST warp = rank 0): vector head, TDM tail ----------- - if (rank == 0 && tail) { - // LDS needed: val (+ tmp when reducing multiple srcs), each holding `tail` bytes - uint32_t need = (numSrcs > 1 ? (tmp - ldsBase) : (val - ldsBase)) + tail; - if (ldsBytes >= need) { // stage tail in rank 0's window - issueRow1d(mySrcs, myDsts, numSrcs, numDsts, val, tmp, ldsBuffer, tail, warpThread, tdmBytes); - } else { - warpVecCopy(mySrcs, myDsts, numSrcs, numDsts, tail, tdmBytes, - warpThread, warpThreads); - } - } + // --- edge (team's FIRST warp = rank 0): sub-256B tail via VECTOR reduce ---- + // The tail is < 256B (negligible for bandwidth), so it is reduced with plain + // global vector loads/stores rather than a 1-D TDM tile. Routing the tail + // through the tensor builtins needs a standalone (non-inlined) helper, and the + // current gfx1250 backend miscompiles the ds-reduce -> inaccessiblemem tensor + // store there (store dropped / stale LDS); inlining it instead bloats issue() + // and corrupts the bulk path. A vector tail sidesteps both. + if (rank == 0 && tail) + warpVecCopy(mySrcs, myDsts, numSrcs, numDsts, tail, tdmBytes, + warpThread, warpThreads); // --- 256B rows copy via TDM ------------------------------------------------ uint32_t maxByLds = ldsBytes / RWIDTH; // #warps we can give a window, 2*WIDTH because of double buffering