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/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; } diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index 2f7740c8..3e5f989d 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,59 @@ 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; + + 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(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; + } + + 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 +6010,7 @@ namespace { ConfigOptions const& cfg, bool const subExecParamHostAccessible, uint32_t const ldsBytes, + int const tdmKernelIdx, TransferResources& rss) { // Compute kernel launch parameters @@ -5900,18 +6019,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 +6098,7 @@ namespace { std::cref(cfg), exeInfo.subExecParamHostAccessible, exeInfo.ldsBytesActual, + exeInfo.tdmKernelToUse, std::ref(exeInfo.resources[i]))); } for (auto& asyncTransfer : asyncTransfers) @@ -5985,7 +6108,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 +6346,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. diff --git a/src/header/tdmCopy.h b/src/header/tdmCopy.h index 2fc79a13..e23fa06c 100644 --- a/src/header/tdmCopy.h +++ b/src/header/tdmCopy.h @@ -345,11 +345,25 @@ 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 + +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 @@ -521,6 +535,254 @@ __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, ...). + +// ---- 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 +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) { + // 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]; + for (uint32_t d = 0; d < numDsts; ++d) + 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 = numSrcs ? 0 : MEMSET_CHAR; + for (uint32_t s = 0; s < numSrcs; ++s) + acc += reinterpret_cast(srcs[s] + offset)[nf * 4 + b]; + for (uint32_t d = 0; d < numDsts; ++d) + reinterpret_cast(dsts[d] + offset)[nf * 4 + b] = acc; + } + } +} + +// ---- 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. +// ---- 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 +__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)); +} + +// ---- 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 +// 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. --- +// 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 +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) { + 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) + + 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(); // RAW: src0 LDS write visible to vp reads + LdsPtrT vp = ldsPtr3(val); + for (size_t s = 1; s < numSrcs; s++) { + 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 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; + } + + 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 / next dst + } +} + +// ---- core: partition + issue a reduce/broadcast for the team [start, stop). - +template +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)); + 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; + + // --- 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 + 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; + + // 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) + ? 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, ldsBuffer, chunkRows, warpThread, off); + } +} + } // namespace detail #elif TDM_BACKEND_NV // ----- NVIDIA cp.async.bulk / TMA (sm_90+) ----- @@ -714,6 +976,15 @@ __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 +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); + 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.