diff --git a/Makefile b/Makefile index db40ed20..46858abc 100644 --- a/Makefile +++ b/Makefile @@ -6,7 +6,15 @@ ROCM_PATH ?= /opt/rocm CUDA_PATH ?= /usr/local/cuda MPI_PATH ?= /usr/local/openmpi +# pip ROCm wheels ship bin/amdclang++ but often omit bin/amdllvm (which that stub execs). +# Default to llvm/bin when bin/amdllvm is absent so HIP builds work without extra flags. +ifeq ("$(shell test -e $(ROCM_PATH)/bin/amdllvm && echo found)", "found") HIPCC ?= $(ROCM_PATH)/bin/amdclang++ +else ifeq ("$(shell test -e $(ROCM_PATH)/llvm/bin/amdclang++ && echo found)", "found") +HIPCC ?= $(ROCM_PATH)/llvm/bin/amdclang++ +else +HIPCC ?= $(ROCM_PATH)/bin/amdclang++ +endif NVCC ?= $(CUDA_PATH)/bin/nvcc DEBUG ?= 0 diff --git a/src/client/EnvVars.hpp b/src/client/EnvVars.hpp index 34940b69..69bd23be 100644 --- a/src/client/EnvVars.hpp +++ b/src/client/EnvVars.hpp @@ -124,6 +124,11 @@ class EnvVars int nicTrafficClass; // DSCP/traffic class byte for RoCE GRH int roceVersion; // RoCE version number + // TDM (async tensor) options + int tdmBlockSize; // Size of each threadblock for TDM (async tensor) kernels (must be multiple of 32) + int tdmMaxLdsBytes; // Max LDS (shared memory) bytes per workgroup for TDM kernels (INT_MAX = use device max) + int tdmPipelined; // Use the pipelined (depth=2) TDM tensor kernel (0=non-pipelined, 1=pipelined) + // Developer features int gpuMaxHwQueues; // Tracks GPU_MAX_HW_QUEUES environment variable @@ -173,6 +178,9 @@ class EnvVars showBorders = GetEnvVar("SHOW_BORDERS" , 1); showIterations = GetEnvVar("SHOW_ITERATIONS" , 0); showPercentiles = GetEnvVarArray("SHOW_PERCENTILES", {}); + tdmBlockSize = GetEnvVar("TDM_BLOCK_SIZE" , 256); + tdmMaxLdsBytes = GetEnvVar("TDM_MAX_LDS_BYTES" , INT_MAX); + tdmPipelined = GetEnvVar("TDM_PIPELINED" , 0); useHipEvents = GetEnvVar("USE_HIP_EVENTS" , 1); useHsaDma = GetEnvVar("USE_HSA_DMA" , 0); useInteractive = GetEnvVar("USE_INTERACTIVE" , 0); @@ -394,6 +402,9 @@ class EnvVars printf(" SHOW_BORDERS - Show ASCII box-drawing characters in tables\n"); printf(" SHOW_ITERATIONS - Show per-iteration timing info\n"); printf(" SHOW_PERCENTILES - Comma-separated percentiles iteration duration\n"); + printf(" TDM_BLOCK_SIZE - # of threads per threadblock for TDM (async tensor) kernels (Must be multiple of 32)\n"); + printf(" TDM_MAX_LDS_BYTES - Max LDS bytes per workgroup for TDM kernels (defaults to device max; K/M/G suffixes accepted)\n"); + printf(" TDM_PIPELINED - 1=use pipelined (depth=2) TDM kernel, 0=use non-pipelined kernel\n"); printf(" USE_HIP_EVENTS - Use HIP events for GFX executor timing\n"); printf(" USE_HSA_DMA - Use hsa_amd_async_copy instead of hipMemcpy for non-targeted DMA execution\n"); printf(" USE_INTERACTIVE - Pause for user-input before starting transfer loop\n"); @@ -541,6 +552,14 @@ class EnvVars "%s per-iteration timing", showIterations ? "Showing" : "Hiding"); Print("SHOW_PERCENTILES", showPercentiles.empty() ? 0 : 1, "%s", showPercentiles.empty() ? "Disabled" : GetStr(showPercentiles).c_str()); + Print("TDM_BLOCK_SIZE", tdmBlockSize, + "TDM threadblock size of %d", tdmBlockSize); + Print("TDM_MAX_LDS_BYTES", tdmMaxLdsBytes, + "%s", tdmMaxLdsBytes == INT_MAX + ? "Using device max LDS bytes per workgroup" + : (std::string("Capping LDS to ") + std::to_string(tdmMaxLdsBytes) + " bytes per workgroup").c_str()); + Print("TDM_PIPELINED", tdmPipelined, + "Using %s TDM tensor kernel", tdmPipelined ? "pipelined (depth=2)" : "non-pipelined"); Print("USE_HIP_EVENTS", useHipEvents, "Using %s for GFX/DMA Executor timing", useHipEvents ? "HIP events" : "CPU wall time"); Print("USE_HSA_DMA", useHsaDma, @@ -749,6 +768,10 @@ class EnvVars cfg.nic.trafficClass = nicTrafficClass; cfg.nic.roceVersion = roceVersion; + cfg.tdm.blockSize = tdmBlockSize; + cfg.tdm.maxLDSBytes = tdmMaxLdsBytes; + cfg.tdm.pipelined = tdmPipelined; + return cfg; } }; diff --git a/src/client/Presets/Help.hpp b/src/client/Presets/Help.hpp index 26ede846..de92ee9e 100644 --- a/src/client/Presets/Help.hpp +++ b/src/client/Presets/Help.hpp @@ -37,13 +37,15 @@ int HelpPreset([[maybe_unused]] EnvVars& ev, printf("# SRC 1 -> Executor -> DST 1\n"); printf("# SRC X DST Y\n"); printf("\n"); - printf("# Five Executors are supported by TransferBench\n"); - printf("# Executor: SubExecutor:\n"); - printf("# 1) CPU CPU thread\n"); - printf("# 2) GPU GPU threadblock/Compute Unit (CU)\n"); - printf("# 3) DMA N/A. (Must have single SRC, at least one DST)\n"); - printf("# 4) NIC Queue Pair\n"); - printf("# 5) Batched-DMA Batch item (Must have single SRC, at least one DST)\n"); + printf("# Seven Executors are supported by TransferBench\n"); + printf("# Executor: SubExecutor:\n"); + printf("# 1) CPU CPU thread\n"); + printf("# 2) GPU GPU threadblock/Compute Unit (CU)\n"); + printf("# 3) DMA N/A. (Must have single SRC, at least one DST)\n"); + printf("# 4) NIC Queue Pair\n"); + printf("# 5) Batched-DMA Batch item (Must have single SRC, at least one DST)\n"); + printf("# 6) TDM GPU threadblock/Compute Unit (CU)\n"); + printf("# 7) Async Load/Store GPU threadblock/Compute Unit (CU)\n"); printf("\n"); printf("# Each single line in the configuration file defines a set of Transfers (a Test) to run in parallel\n"); printf("\n"); @@ -71,6 +73,8 @@ int HelpPreset([[maybe_unused]] EnvVars& ev, printf("# - B: Batched-DMA-executor (Indexed from 0 to # GPUs - 1)\n"); printf("# - I#.#: NIC executor (Indexed from 0 to # NICs - 1)\n"); printf("# - N#.#: Nearest NIC executor (Indexed from 0 to # GPUs - 1)\n"); + printf("# - T: GPU TDM kernel kernel (Indexed from 0 to # GPUs - 1)\n"); + printf("# - L: GPU async load/store (Indexed from 0 to # GPUs - 1)\n"); printf("# dstMemL : Destination memory locations (Where the data is to be written to)\n"); printf("# bytesL : Number of bytes to copy (0 means use command-line specified size)\n"); printf("# Must be a multiple of 4 and may be suffixed with ('K','M', or 'G')\n"); @@ -108,6 +112,12 @@ int HelpPreset([[maybe_unused]] EnvVars& ev, printf("## Single DMA executed Transfer between GPUs 0 and 1\n"); printf("1 1 (G0->D0->G1)\n"); printf("\n"); + printf("## Single GPU-executed Transfer between GPUs 0 and 1 using 1 CU for tensor-op path\n"); + printf("1 1 (G0->T0->G1)\n"); + printf("\n"); + printf("## Single GPU-executed Transfer between GPUs 0 and 1 using 1 CU for async load/store path\n"); + printf("1 1 (G0->L0->G1)\n"); + printf("\n"); printf("## Copy 1Mb from GPU0 to GPU1 with 4 CUs, and 2Mb from GPU1 to GPU0 with 8 CUs\n"); printf("-2 (G0->G0->G1 4 1M) (G1->G1->G0 8 2M)\n"); printf("\n"); diff --git a/src/client/Utilities.hpp b/src/client/Utilities.hpp index 9730ace6..77a44b3c 100644 --- a/src/client/Utilities.hpp +++ b/src/client/Utilities.hpp @@ -439,6 +439,8 @@ namespace TransferBench::Utils case EXE_NIC: return "NIC"; case EXE_NIC_NEAREST: return "NIC"; case EXE_GPU_BDMA: return "BMA"; + case EXE_GPU_ASYNC_TENSOR: return "AT"; // async tensor kernel path + case EXE_GPU_ASYNC_MEMOPS: return "AL"; // async load/store kernel path default: return "N/A"; } } diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index fa536155..d6b6475c 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -25,6 +25,10 @@ THE SOFTWARE. #include #include #include +#include +#include +#include +#include #include #include #include @@ -77,9 +81,9 @@ THE SOFTWARE. #ifdef AMD_SMI_ENABLED #include "amd_smi/amdsmi.h" #endif +#include "tdm.h" #endif /// @endcond - // Batched DMA executor is only supported with HIP >= 7.1 and CUDA 12.8 #if (defined(HIP_VERSION) && (HIP_VERSION >= 70100000)) || (defined(CUDA_VERSION) && (CUDA_VERSION >= 12080)) #define BMA_EXEC_ENABLED @@ -107,10 +111,15 @@ namespace TransferBench EXE_NIC = 3, ///< NIC RDMA executor (subExecutor = queue pair) EXE_NIC_NEAREST = 4, ///< NIC RDMA nearest executor (subExecutor = queue pair) EXE_GPU_BDMA = 5, ///< GPU Batched SDMA executor (subExecutor = batch item) + EXE_GPU_ASYNC_TENSOR = 6, ///< GPU async kernel — tensor-op path (subExecutor — reserved; stub uses 1-D launch) + EXE_GPU_ASYNC_MEMOPS = 7, ///< GPU async kernel — load/store path (subExecutor — reserved; stub uses 1-D launch) }; - char const ExeTypeStr[7] = "CGDINB"; + char const ExeTypeStr[9] = "CGDINBTL"; inline bool IsCpuExeType(ExeType e){ return e == EXE_CPU; } - inline bool IsGpuExeType(ExeType e){ return e == EXE_GPU_GFX || e == EXE_GPU_DMA || e == EXE_GPU_BDMA; } + inline bool IsGpuExeType(ExeType e){ + return e == EXE_GPU_GFX || e == EXE_GPU_DMA || e == EXE_GPU_BDMA + || e == EXE_GPU_ASYNC_TENSOR || e == EXE_GPU_ASYNC_MEMOPS; + } inline bool IsNicExeType(ExeType e){ return e == EXE_NIC || e == EXE_NIC_NEAREST; } /** @@ -275,6 +284,13 @@ namespace TransferBench int useNuma = 0; ///< Switch to closest numa thread for execution }; + struct TdmOptions + { + int blockSize = 256; ///< Size of each threadblock (must be multiple of 64) + int maxLDSBytes = INT_MAX; ///< Maximum number of bytes of __shared__ memory to use + int pipelined = 0; ///< Whether to call the pipelined TDM kernel + }; + /** * Configuration options for performing Transfers @@ -287,6 +303,7 @@ namespace TransferBench GfxOptions gfx; ///< GFX executor options DmaOptions dma; ///< DMA executor options NicOptions nic; ///< NIC executor options + TdmOptions tdm; ///< TDM executor options }; /** @@ -2242,7 +2259,8 @@ namespace { // Each subexecutor is assigned a multiple of cfg.data.blockBytes, however this may // mean that some subexecutors might not have any work assigned to them if the amount to // transfer is small - if (t.exeDevice.exeType == EXE_GPU_GFX || t.exeDevice.exeType == EXE_CPU || t.exeDevice.exeType == EXE_GPU_BDMA) { + if (t.exeDevice.exeType == EXE_GPU_GFX || t.exeDevice.exeType == EXE_CPU || t.exeDevice.exeType == EXE_GPU_BDMA || + t.exeDevice.exeType == EXE_GPU_ASYNC_TENSOR || t.exeDevice.exeType == EXE_GPU_ASYNC_MEMOPS) { size_t const N = t.numBytes / sizeof(float); int const targetMultiple = cfg.data.blockBytes / sizeof(float); int const maxSubExecToUse = std::min((size_t)(N + targetMultiple - 1) / targetMultiple, @@ -2300,6 +2318,27 @@ namespace { hasFatalError = true; } break; + case EXE_GPU_ASYNC_TENSOR: + case EXE_GPU_ASYNC_MEMOPS: + if (t.srcs.size() != 1 || t.dsts.size() != 1) { + errors.push_back({ERR_FATAL, + "Transfer %d: Async GPU kernel executor requires exactly 1 SRC and 1 DST", i}); + hasFatalError = true; + break; + } + if (t.exeDevice.exeIndex < 0 || t.exeDevice.exeIndex >= numExecutors) { + errors.push_back({ERR_FATAL, + "Transfer %d: Async GPU kernel device index must be between 0 and %d (instead of %d) for rank %d", + i, numExecutors - 1, t.exeDevice.exeIndex, t.exeDevice.exeRank}); + hasFatalError = true; + break; + } + if (t.exeSubIndex != -1) { + errors.push_back({ERR_FATAL, + "Transfer %d: Async GPU kernel executor does not support subindices yet", i}); + hasFatalError = true; + } + break; case EXE_GPU_GFX: if (t.exeDevice.exeIndex < 0 || t.exeDevice.exeIndex >= numExecutors) { errors.push_back({ERR_FATAL, @@ -2671,6 +2710,22 @@ namespace { } break; } + case EXE_GPU_ASYNC_TENSOR: + case EXE_GPU_ASYNC_MEMOPS: + { + int numGpuSubExec = GetNumSubExecutors(exeDevice); + if (totalSubExecs[exeDevice] > numGpuSubExec) + errors.push_back({ERR_WARN, + "GPU %d (async kernel) requests %d total subexecutors however only %d CUs available. " + "Serialization may occur", + exeDevice.exeIndex, totalSubExecs[exeDevice], numGpuSubExec}); + if (transferCount[exeDevice] > gpuMaxHwQueues) { + errors.push_back({ERR_WARN, + "GPU %d (async kernel) attempting %d parallel transfers, however GPU_MAX_HW_QUEUES only set to %d", + exeDevice.exeIndex, transferCount[exeDevice], gpuMaxHwQueues}); + } + break; + } case EXE_GPU_DMA: { // Check that if executor subindices are used, all Transfers specify executor subindices @@ -4327,13 +4382,16 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } // Prepare additional requirements for GPU-based executors - if ((exeDevice.exeType == EXE_GPU_GFX || exeDevice.exeType == EXE_GPU_DMA || exeDevice.exeType == EXE_GPU_BDMA) + if ((exeDevice.exeType == EXE_GPU_GFX || exeDevice.exeType == EXE_GPU_DMA || exeDevice.exeType == EXE_GPU_BDMA || + exeDevice.exeType == EXE_GPU_ASYNC_TENSOR || exeDevice.exeType == EXE_GPU_ASYNC_MEMOPS) && exeDevice.exeRank == localRank) { ERR_CHECK(hipSetDevice(exeDevice.exeIndex)); // Determine how many streams to use int const numStreamsToUse = (exeDevice.exeType == EXE_GPU_DMA || exeDevice.exeType == EXE_GPU_BDMA || - (exeDevice.exeType == EXE_GPU_GFX && cfg.gfx.useMultiStream)) + (exeDevice.exeType == EXE_GPU_GFX && cfg.gfx.useMultiStream) || + exeDevice.exeType == EXE_GPU_ASYNC_TENSOR || + exeDevice.exeType == EXE_GPU_ASYNC_MEMOPS) ? exeInfo.resources.size() : 1; exeInfo.streams.resize(numStreamsToUse); @@ -4351,7 +4409,8 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } } - if (cfg.gfx.useHipEvents || cfg.dma.useHipEvents) { + if (cfg.gfx.useHipEvents || cfg.dma.useHipEvents || exeDevice.exeType == EXE_GPU_ASYNC_TENSOR || + exeDevice.exeType == EXE_GPU_ASYNC_MEMOPS) { exeInfo.startEvents.resize(numStreamsToUse); exeInfo.stopEvents.resize(numStreamsToUse); for (int i = 0; i < numStreamsToUse; ++i) { @@ -4557,11 +4616,13 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } // Teardown additional requirements for GPU-based executors - if ((exeDevice.exeType == EXE_GPU_GFX || exeDevice.exeType == EXE_GPU_DMA || exeDevice.exeType == EXE_GPU_BDMA) + if ((exeDevice.exeType == EXE_GPU_GFX || exeDevice.exeType == EXE_GPU_DMA || exeDevice.exeType == EXE_GPU_BDMA || + exeDevice.exeType == EXE_GPU_ASYNC_TENSOR || exeDevice.exeType == EXE_GPU_ASYNC_MEMOPS) && exeDevice.exeRank == localRank) { for (auto stream : exeInfo.streams) ERR_CHECK(hipStreamDestroy(stream)); - if (cfg.gfx.useHipEvents || cfg.dma.useHipEvents) { + if (cfg.gfx.useHipEvents || cfg.dma.useHipEvents || exeDevice.exeType == EXE_GPU_ASYNC_TENSOR || + exeDevice.exeType == EXE_GPU_ASYNC_MEMOPS) { for (auto event : exeInfo.startEvents) ERR_CHECK(hipEventDestroy(event)); for (auto event : exeInfo.stopEvents) @@ -4827,80 +4888,244 @@ static bool IsConfiguredGid(union ibv_gid const& gid) #define TEMPORAL_STORE 2 #define TEMPORAL_BOTH 3 - template - __device__ __forceinline__ void Load(float const* src, float& dst) { - if (TEMPORAL_MODE & TEMPORAL_LOAD) { -#if !defined(__NVCC__) - dst = __builtin_nontemporal_load(src); + template + struct VecCast { + static_assert( + std::is_same_v || std::is_same_v || std::is_same_v, + "NonTemporalVec: T must be float, float2, or float4"); + using type = float __attribute__((ext_vector_type(sizeof(T) / sizeof(float)))); + }; + template + using VecCastType = typename VecCast::type; + template + __device__ __forceinline__ void Load(T const* src, T& dst) { + if constexpr (TEMPORAL_MODE & TEMPORAL_LOAD) { +#if !defined(__NVCC__) + using Vec = VecCastType; + *reinterpret_cast(&dst) = __builtin_nontemporal_load(reinterpret_cast(src)); #endif } else { dst = *src; } } - template - __device__ __forceinline__ void Load(float2 const* src, float2& dst) { - if (TEMPORAL_MODE & TEMPORAL_LOAD) { + template + __device__ __forceinline__ void Store(T const& src, T* dst) { + if constexpr (TEMPORAL_MODE & TEMPORAL_STORE) { #if !defined(__NVCC__) - dst.x = __builtin_nontemporal_load(&(src->x)); - dst.y = __builtin_nontemporal_load(&(src->y)); + using Vec = VecCastType; + __builtin_nontemporal_store(*reinterpret_cast(&src), reinterpret_cast(dst)); #endif } else { - dst = *src; + *dst = src; } } - template - __device__ __forceinline__ void Load(float4 const* src, float4& dst) { - if (TEMPORAL_MODE & TEMPORAL_LOAD) { -#if !defined(__NVCC__) - dst.x = __builtin_nontemporal_load(&(src->x)); - dst.y = __builtin_nontemporal_load(&(src->y)); - dst.z = __builtin_nontemporal_load(&(src->z)); - dst.w = __builtin_nontemporal_load(&(src->w)); -#endif - } else { - dst = *src; + //---------------------------------------------------------------------------- +#if !defined(__gfx1250__) + __global__ void NonAsyncFallbackKernel(float const* __restrict__ src, float* __restrict__ dst, size_t numFloats){ + size_t const gid = blockIdx.x * blockDim.x + threadIdx.x; + size_t const stride = gridDim.x * blockDim.x; + for (size_t i = gid; i < numFloats; i += stride){ + dst[i] = src[i]; } } +#else // __gfx1250__ +// The TDM API does not have a function to set the transfer size, so we need to do it manually. +__device__ void SetTransferSize(gfx1250_TDM_GROUP1& group1, int numElements){ + group1.tensorDim0(numElements); + group1.tensorDim0Stride(numElements); + group1.tileDim0(numElements); +} +namespace { + __constant__ constexpr bool verbose = false; // Turn off to disable debug prints from kernel TDM kernels. +} - template - __device__ __forceinline__ void Store(float const& src, float* dst) { - if (TEMPORAL_MODE & TEMPORAL_STORE) { -#if !defined(__NVCC__) - __builtin_nontemporal_store(src, dst); -#endif - } else { - *dst = src; - } +// numElementsPerTile is the number of elements to process per tile. Its maximum value is the shared memory size / number of waves per workgroup. +__global__ void GpuAsyncTensorOpsKernel(float const* __restrict__ src, float* __restrict__ dst, size_t numElements, int numElementsPerTile){ + extern __shared__ __align__(128) float shmem[]; + int waveId = threadIdx.x / warpSize; + int numWavesPerBlock = blockDim.x / warpSize; + size_t itemsProcessedPerGridIteration = numElementsPerTile * numWavesPerBlock * gridDim.x; + + float* shmemPtr = static_cast(shmem) + numElementsPerTile * waveId; + // Local per-wave source and destination pointers + const float* srcPtr = src + numElementsPerTile * (waveId + blockIdx.x * numWavesPerBlock); + float* dstPtr = dst + numElementsPerTile * (waveId + blockIdx.x * numWavesPerBlock); + gfx1250_TDM_GROUP0 group0; + + group0.ldsAddr((uintptr_t)shmemPtr); + + gfx1250_TDM_GROUP1 group1; + group1.dataSize(2); // Log2 of the element size in bytes, so 2 for float + SetTransferSize(group1, numElementsPerTile); + + constexpr __hip_uint32x4 empty_x4{}; + constexpr __hip_uint32x8 empty_x8{}; + size_t elementsToProcess = numElementsPerTile; + while(srcPtr < src + numElements){ + // Handle the last tile of the block, which may be less than num_elements_per_tile. + if(src + numElements - srcPtr < numElementsPerTile){ + size_t remainingElements = src + numElements - srcPtr; + SetTransferSize(group1, remainingElements); + if constexpr(verbose) elementsToProcess = remainingElements; + } + // Copy from global memory to LDS + group0.globalAddr((uintptr_t)srcPtr); + if constexpr(verbose) if (threadIdx.x % warpSize == 0) printf("Block %d, Wave %d: Loading %zu elements from %p to shared memory at %p\n", blockIdx.x, waveId, elementsToProcess, srcPtr, shmemPtr); + __builtin_amdgcn_tensor_load_to_lds(group0.m_bitfield, group1.m_bitfield, empty_x4, empty_x4, empty_x8, 0); + __builtin_amdgcn_s_wait_tensorcnt(0); + + // write back from LDS to global + group0.globalAddr((uintptr_t)dstPtr); + if constexpr(verbose) if (threadIdx.x % warpSize == 0) printf("Block %d, Wave %d: Storing %zu elements from shared memory at %p to %p\n", blockIdx.x, waveId, elementsToProcess, shmemPtr, dstPtr); + __builtin_amdgcn_tensor_store_from_lds(group0.m_bitfield, group1.m_bitfield, empty_x4, empty_x4, empty_x8, 0); + __builtin_amdgcn_s_wait_tensorcnt(0); + srcPtr += itemsProcessedPerGridIteration; + dstPtr += itemsProcessedPerGridIteration; } +} - template - __device__ __forceinline__ void Store(float2 const& src, float2* dst) { - if (TEMPORAL_MODE & TEMPORAL_STORE) { -#if !defined(__NVCC__) - __builtin_nontemporal_store(src.x, &(dst->x)); - __builtin_nontemporal_store(src.y, &(dst->y)); -#endif - } else { - *dst = src; - } +struct TileMover { + gfx1250_TDM_GROUP0 group0; + gfx1250_TDM_GROUP1 group1; + float* dstPtr{nullptr}; + float* shmemPtr{nullptr}; // TODO: Remove this + int numElementsToProcess{0}; // TODO: Remove this + + __device__ void LoadTile(float* shmemPtr, const float* srcPtr, float* dstPtr, int numElementsToProcess) { + group0.ldsAddr((uintptr_t)shmemPtr); + group0.globalAddr((uintptr_t)srcPtr); + group1.dataSize(2); + SetTransferSize(group1, numElementsToProcess); + this -> dstPtr = dstPtr; this -> shmemPtr = shmemPtr; this -> numElementsToProcess = numElementsToProcess; + if constexpr(verbose) if (threadIdx.x % warpSize == 0) printf("Warp %d Block %d: Loading %d elements from %p to %p\n", threadIdx.x / warpSize, blockIdx.x, numElementsToProcess, srcPtr, shmemPtr); + __builtin_amdgcn_tensor_load_to_lds(group0.m_bitfield, group1.m_bitfield, __hip_uint32x4{}, __hip_uint32x4{}, __hip_uint32x8{}, 0); + } + __device__ void StoreTile() { + assert(dstPtr != nullptr); + group0.globalAddr((uintptr_t)dstPtr); + if constexpr(verbose) if (threadIdx.x % warpSize == 0) printf("Warp %d Block %d: Storing %d elements from %p to %p\n", threadIdx.x / warpSize, blockIdx.x, numElementsToProcess, shmemPtr, dstPtr); + __builtin_amdgcn_tensor_store_from_lds(group0.m_bitfield, group1.m_bitfield, __hip_uint32x4{}, __hip_uint32x4{}, __hip_uint32x8{}, 0); + } +}; + +// numElementsPerTile is the number of elements to process per tile. Its maximum value is the shared memory size / number of waves per workgroup. +__global__ void GpuAsyncPipelinedTensorOpsKernel(float const* __restrict__ src, float* __restrict__ dst, size_t numElements, int numElementsPerSubtile){ + extern __shared__ __align__(128) float shmem[]; + int waveId = threadIdx.x / warpSize; + int numWavesPerBlock = blockDim.x / warpSize; // This only works because the blockDim.x is fixed. + constexpr int pipelineDepth = 2; + + int numElementsPerTile = pipelineDepth * numElementsPerSubtile; + size_t itemsProcessedPerGridIteration = numElementsPerTile * numWavesPerBlock * gridDim.x; + + float* shmemPtr = static_cast(shmem) + numElementsPerTile * waveId; + // Local per-wave source and destination pointers + const float* srcPtr = src + numElementsPerTile * (waveId + blockIdx.x * numWavesPerBlock); + float* dstPtr = dst + numElementsPerTile * (waveId + blockIdx.x * numWavesPerBlock); + + // Skip waves whose initial position is already past the end of the buffer. + if (srcPtr >= src + numElements){ + if constexpr(verbose) if (threadIdx.x % warpSize == 0) printf("Warp %d Block %d: Skipping wave %d\n", threadIdx.x / warpSize, blockIdx.x, waveId); + return; + } + + int numElementsToProcess = std::min(numElementsPerSubtile, (int)(src + numElements - srcPtr)); + TileMover tileMovers[pipelineDepth]{}; + unsigned int slot = 0; + tileMovers[slot].LoadTile(shmemPtr, srcPtr, dstPtr, numElementsToProcess); + while(srcPtr + (slot ^ 1) * numElementsPerSubtile < src + numElements){ + slot ^= 1; + // Handle the last tile of the block, which may be less than num_elements_per_tile. + if(src + numElements - (srcPtr + slot * numElementsPerSubtile) < numElementsPerSubtile){ + numElementsToProcess = src + numElements - (srcPtr + slot * numElementsPerSubtile); + if constexpr(verbose) if (threadIdx.x % warpSize == 0) printf("Small Subtile: Warp %d Block %d: Num Elements To Process: %d\n", threadIdx.x / warpSize, blockIdx.x, numElementsToProcess); + } + // Copy from global memory to LDS + tileMovers[slot].LoadTile(shmemPtr + slot * numElementsPerSubtile, srcPtr + slot * numElementsPerSubtile, dstPtr + slot * numElementsPerSubtile, numElementsToProcess); + __builtin_amdgcn_s_wait_tensorcnt(1); + + // write back from LDS to global for the other slot + tileMovers[slot^1].StoreTile(); + __builtin_amdgcn_s_wait_tensorcnt(1); + srcPtr += itemsProcessedPerGridIteration * slot; + dstPtr += itemsProcessedPerGridIteration * slot; + } + __builtin_amdgcn_s_wait_tensorcnt(0); + tileMovers[slot].StoreTile(); + __builtin_amdgcn_s_wait_tensorcnt(0); +} + + // Source for definitions: https://github.com/llvm/llvm-project/blob/4e3bac3ea2cc6fd778d53e317dd9fc27c1ddfc4f/clang/include/clang/Basic/BuiltinsAMDGPU.td#L956 + // and internal ISA documentation + using int32x4 = __attribute__((__vector_size__(4 * sizeof(int)))) int; + // Loads 16 bytes/lane from global (src) into LDS (dst). + __device__ void asyncLoadX4(float const* src, float* dst){ + __builtin_amdgcn_global_load_async_to_lds_b128( + (__attribute__((address_space(1))) int32x4*)src, + (__attribute__((address_space(3))) int32x4*)dst, 0, 0); + } + // Stores 16 bytes/lane from LDS (src) to global (dst). + __device__ void asyncStoreX4(float const* src, float* dst){ + __builtin_amdgcn_global_store_async_from_lds_b128( + (__attribute__((address_space(1))) int32x4*)dst, + (__attribute__((address_space(3))) int32x4*)src, 0, 0); } - template - __device__ __forceinline__ void Store(float4 const& src, float4* dst) { - if (TEMPORAL_MODE & TEMPORAL_STORE) { -#if !defined(__NVCC__) - __builtin_nontemporal_store(src.x, &(dst->x)); - __builtin_nontemporal_store(src.y, &(dst->y)); - __builtin_nontemporal_store(src.z, &(dst->z)); - __builtin_nontemporal_store(src.w, &(dst->w)); -#endif - } else { - *dst = src; + __device__ void asyncLoad(float const* src, float* dst){ + __builtin_amdgcn_global_load_async_to_lds_b32( + (__attribute__((address_space(1))) int32_t*)src, + (__attribute__((address_space(3))) int32_t*)dst, 0, 0); + } + __device__ void asyncStore(float const* src, float* dst){ + __builtin_amdgcn_global_store_async_from_lds_b32( + (__attribute__((address_space(1))) int32_t*)dst, + (__attribute__((address_space(3))) int32_t*)src, 0, 0); + } + + // numElementsPerTile should be a multiple of 128, given the asyncLoadX4 and asyncStoreX4 functions operate on a 32-lane warp and 4 elements per lane (512 bytes total) + __global__ void GpuAsyncLoadStoreKernel(float const* __restrict__ src, float* __restrict__ dst, size_t numElements, int numElementsPerTile) + { + extern __shared__ __align__(128) float shmem[]; + int waveId = threadIdx.x / warpSize; + int laneId = threadIdx.x % warpSize; + int numWavesPerBlock = blockDim.x / warpSize; + size_t itemsProcessedPerGridIteration = (size_t)numElementsPerTile * numWavesPerBlock * gridDim.x; + constexpr size_t stride_per_lane = sizeof(float4) / sizeof(float); // 4 elements per lane + // Number of elements a full 32-lane warp moves per b128 async op (32 * 4 = 128). + size_t elements_per_dwordx4_load_store = warpSize * stride_per_lane; + + float* shmemPtr = shmem + numElementsPerTile * waveId + laneId * stride_per_lane; + // Wave-tile base (lane-independent) used to gate the loop and compute the partial-tile size. + size_t tileBase = (size_t)numElementsPerTile * (waveId + (size_t)blockIdx.x * numWavesPerBlock); + + while(tileBase < numElements){ + // Handle the last tile, which may be less than numElementsPerTile. + size_t elementsToProcess = numElementsPerTile; + if(numElements - tileBase < (size_t)numElementsPerTile) + elementsToProcess = numElements - tileBase; + + const float* srcPtr = src + tileBase + laneId * stride_per_lane; + float* dstPtr = dst + tileBase + laneId * stride_per_lane; + + // Copy from global memory to LDS + for (size_t i = 0; i < elementsToProcess; i += elements_per_dwordx4_load_store) { + asyncLoadX4(srcPtr + i, shmemPtr + i); + __builtin_amdgcn_s_wait_asynccnt(0); } + + // Write back from LDS to global + for (size_t i = 0; i < elementsToProcess; i += elements_per_dwordx4_load_store){ + asyncStoreX4(shmemPtr + i, dstPtr + i); + __builtin_amdgcn_s_wait_asynccnt(0); + } + + tileBase += itemsProcessedPerGridIteration; } +} +#endif // __gfx1250__ // Simplified Kernel for GFX execution for copies only template @@ -5480,6 +5705,132 @@ static bool IsConfiguredGid(union ibv_gid const& gid) return ERR_NONE; } + static ErrResult ExecuteGpuAsyncKernelTransfer(int const iteration, + ConfigOptions const& cfg, + int const exeIndex, + hipStream_t const stream, + hipEvent_t const startEvent, + hipEvent_t const stopEvent, + bool const tensorFlavor, + int const numSubExecs, + TransferResources& rss) + { + ERR_CHECK(hipSetDevice(exeIndex)); + int const initOffset = cfg.data.byteOffset / sizeof(float); + size_t const numFloats = rss.numBytes / sizeof(float); + float const* __restrict__ src = rss.srcMem[0] + initOffset; + float* __restrict__ dst = rss.dstMem[0] + initOffset; + + int const threads = tensorFlavor ? cfg.tdm.blockSize : cfg.gfx.blockSize; + int const blocks = numSubExecs; + + auto cpuStart = std::chrono::high_resolution_clock::now(); + + int subIterations = 0; + do { + if (startEvent) + ERR_CHECK(hipEventRecord(startEvent, stream)); +#if !defined(__gfx1250__) // NVIDIA and AMD Instinct GPUs prior to GFX1250 + NonAsyncFallbackKernel<<>>(src, dst, numFloats); +#else + bool usePipelinedTensorOps = (cfg.tdm.pipelined != 0); + auto gpuKernel = tensorFlavor ? (usePipelinedTensorOps ? GpuAsyncPipelinedTensorOpsKernel : GpuAsyncTensorOpsKernel) : GpuAsyncLoadStoreKernel; + if (rss.numBytes % sizeof(float) != 0) + return {ERR_FATAL, "Async tensor executor (TDM): numBytes (%zu) must be a multiple of %zu", rss.numBytes, sizeof(float)}; + + + hipDeviceProp_t props{}; + ERR_CHECK(hipGetDeviceProperties(&props, exeIndex)); + int const warpSize = props.warpSize; + int maxShmem = 0; + ERR_CHECK(hipDeviceGetAttribute(&maxShmem, hipDeviceAttributeMaxSharedMemoryPerBlock, exeIndex)); + maxShmem = std::min(maxShmem, cfg.tdm.maxLDSBytes); + int kWavesPerWorkgroup = threads / warpSize; + int pipelineDepth = usePipelinedTensorOps ? 2 : 1; + int numFloatsPerTile = std::max(1, maxShmem / (sizeof(float) * kWavesPerWorkgroup * pipelineDepth)); + unsigned int shmemBytes = numFloatsPerTile * kWavesPerWorkgroup * sizeof(float) * pipelineDepth; + hipLaunchKernelGGL(gpuKernel, dim3(blocks), dim3(threads), + shmemBytes, stream, src, dst, numFloats, + numFloatsPerTile); +#endif + ERR_CHECK(hipGetLastError()); + if (stopEvent) + ERR_CHECK(hipEventRecord(stopEvent, stream)); + ERR_CHECK(hipStreamSynchronize(stream)); + } while (++subIterations != cfg.general.numSubIterations); + + if (iteration >= 0) { + double deltaMsec; + if (startEvent && stopEvent) { + float gpuDeltaMsec; + ERR_CHECK(hipEventElapsedTime(&gpuDeltaMsec, startEvent, stopEvent)); + deltaMsec = gpuDeltaMsec / cfg.general.numSubIterations; + } else { + auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart; + deltaMsec = std::chrono::duration_cast>(cpuDelta).count() * 1000.0 + / cfg.general.numSubIterations; + } + rss.totalDurationMsec += deltaMsec; + if (cfg.general.recordPerIteration) + rss.perIterMsec.push_back(deltaMsec); + } + return ERR_NONE; + } + + static ErrResult RunGpuAsyncKernelExecutor(int const iteration, + ConfigOptions const& cfg, + int const exeIndex, + bool const tensorFlavor, + ExeInfo& exeInfo) + { + auto cpuStart = std::chrono::high_resolution_clock::now(); + ERR_CHECK(hipSetDevice(exeIndex)); + + vector> asyncTransfers; + for (size_t i = 0; i < exeInfo.resources.size(); i++) { + hipEvent_t startEv = (!exeInfo.startEvents.empty()) ? exeInfo.startEvents[i] : nullptr; + hipEvent_t stopEv = (!exeInfo.stopEvents.empty()) ? exeInfo.stopEvents[i] : nullptr; + int const numSubExecs = + static_cast(exeInfo.resources[i].subExecParamCpu.size()); + asyncTransfers.emplace_back(std::async(std::launch::async, + ExecuteGpuAsyncKernelTransfer, + iteration, + std::cref(cfg), + exeIndex, + exeInfo.streams[i], + startEv, + stopEv, + tensorFlavor, + numSubExecs, + std::ref(exeInfo.resources[i]))); + } + for (auto& asyncTransfer : asyncTransfers) + ERR_CHECK(asyncTransfer.get()); + + auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart; + double deltaMsec = std::chrono::duration_cast>(cpuDelta).count() * 1000.0 + / cfg.general.numSubIterations; + if (iteration >= 0) + exeInfo.totalDurationMsec += deltaMsec; + return ERR_NONE; + } + + static ErrResult RunGpuAsyncTensorExecutor(int const iteration, + ConfigOptions const& cfg, + int const exeIndex, + ExeInfo& exeInfo) + { + return RunGpuAsyncKernelExecutor(iteration, cfg, exeIndex, true, exeInfo); + } + + static ErrResult RunGpuAsyncMemOpsExecutor(int const iteration, + ConfigOptions const& cfg, + int const exeIndex, + ExeInfo& exeInfo) + { + return RunGpuAsyncKernelExecutor(iteration, cfg, exeIndex, false, exeInfo); + } + // DMA Executor-related functions //======================================================================================== @@ -5700,6 +6051,8 @@ static bool IsConfiguredGid(union ibv_gid const& gid) switch (exeDevice.exeType) { case EXE_CPU: return RunCpuExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); case EXE_GPU_GFX: return RunGpuExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); + case EXE_GPU_ASYNC_TENSOR: return RunGpuAsyncTensorExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); + case EXE_GPU_ASYNC_MEMOPS: return RunGpuAsyncMemOpsExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); case EXE_GPU_DMA: return RunDmaExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); #ifdef NIC_EXEC_ENABLED case EXE_NIC: return RunNicExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); @@ -6400,6 +6753,12 @@ static bool IsConfiguredGid(union ibv_gid const& gid) result |= RecursiveWildcardTransferExpansion(wc, baseRankIndex, numBytes, numSubExecs, transfers); wc.exe.exeSubIndices[0] = -2; return result; + case EXE_GPU_ASYNC_TENSOR: + case EXE_GPU_ASYNC_MEMOPS: + wc.exe.exeSubIndices[0] = -1; + result |= RecursiveWildcardTransferExpansion(wc, baseRankIndex, numBytes, numSubExecs, transfers); + wc.exe.exeSubIndices[0] = -2; + return result; case EXE_GPU_GFX: case EXE_GPU_DMA: case EXE_GPU_BDMA: { // Iterate over all available subindices @@ -7287,6 +7646,8 @@ static bool IsConfiguredGid(union ibv_gid const& gid) topo.numExecutors[EXE_GPU_GFX] = numGpus; topo.numExecutors[EXE_GPU_DMA] = numGpus; topo.numExecutors[EXE_GPU_BDMA] = numGpus; + topo.numExecutors[EXE_GPU_ASYNC_TENSOR] = numGpus; + topo.numExecutors[EXE_GPU_ASYNC_MEMOPS] = numGpus; std::vector gpuArchNames(numGpus); @@ -7310,6 +7671,8 @@ static bool IsConfiguredGid(union ibv_gid const& gid) topo.executorName[{EXE_GPU_GFX, exeIndex}] = gpuName; topo.executorName[{EXE_GPU_DMA, exeIndex}] = gpuName; topo.executorName[{EXE_GPU_BDMA, exeIndex}] = gpuName; + topo.executorName[{EXE_GPU_ASYNC_TENSOR, exeIndex}] = gpuName + " [async tensor]"; + topo.executorName[{EXE_GPU_ASYNC_MEMOPS, exeIndex}] = gpuName + " [async LD/ST]"; #if !defined(__NVCC__) hsa_agent_t gpuAgent = gpuAgents[exeIndex]; @@ -7339,9 +7702,13 @@ static bool IsConfiguredGid(union ibv_gid const& gid) topo.numExecutorSubIndices[{EXE_GPU_GFX, exeIndex}] = numXccs; topo.numExecutorSubIndices[{EXE_GPU_DMA, exeIndex}] = numDmaEngines; topo.numExecutorSubIndices[{EXE_GPU_BDMA, exeIndex}] = 0; + topo.numExecutorSubIndices[{EXE_GPU_ASYNC_TENSOR, exeIndex}] = 0; + topo.numExecutorSubIndices[{EXE_GPU_ASYNC_MEMOPS, exeIndex}] = 0; topo.numSubExecutors[{EXE_GPU_GFX, exeIndex}] = numDeviceCUs; topo.numSubExecutors[{EXE_GPU_DMA, exeIndex}] = 1; topo.numSubExecutors[{EXE_GPU_BDMA, exeIndex}] = numDmaEngines; + topo.numSubExecutors[{EXE_GPU_ASYNC_TENSOR, exeIndex}] = numDeviceCUs; + topo.numSubExecutors[{EXE_GPU_ASYNC_MEMOPS, exeIndex}] = numDeviceCUs; topo.closestCpuNumaToGpu[exeIndex] = closestNuma; topo.closestNicsToGpu[exeIndex] = {}; @@ -7762,6 +8129,7 @@ static bool IsConfiguredGid(union ibv_gid const& gid) agent = cpuAgents[exeDevice.exeIndex]; break; case EXE_GPU_GFX: case EXE_GPU_DMA: case EXE_GPU_BDMA: + case EXE_GPU_ASYNC_TENSOR: case EXE_GPU_ASYNC_MEMOPS: if (exeIndex < 0 || exeIndex >= numGpus) return {ERR_FATAL, "GPU index must be between 0 and %d inclusively", numGpus - 1}; agent = gpuAgents[exeIndex]; diff --git a/src/header/tdm.h b/src/header/tdm.h new file mode 100644 index 00000000..c55a2bdd --- /dev/null +++ b/src/header/tdm.h @@ -0,0 +1,355 @@ +/* +Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef HIP_INCLUDE_HIP_AMD_GFX1250_TDM_H +#define HIP_INCLUDE_HIP_AMD_GFX1250_TDM_H + +#if !defined(__HIPCC_RTC__) +/* hip_runtime_api.h defines hipMemoryType before including driver_types.h; do not include + * driver_types.h alone or HIP 7.x fails with unknown type hipMemoryType. */ +#include +#include +#if __cplusplus +#include +#else +#include +#endif +#endif + +using __hip_uint32x4 = __NATIVE_VECTOR__(4, uint32_t); +using __hip_uint32x8 = __NATIVE_VECTOR__(8, uint32_t); + +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Wshift-count-overflow" // we're going to be using overflow on purpose + +union gfx1250_TDM_GROUP0 +{ + __device__ constexpr gfx1250_TDM_GROUP0() : m_bitfield{0x1, 0, 0, 0x80000000} + {} + + __device__ gfx1250_TDM_GROUP0( uintptr_t lds_addr_in, + uintptr_t global_addr_in, + uint32_t gather_idx_size_in = 0, + uint32_t gather_mode_in = 0 ) + { + globalAddr(global_addr_in); + m_count = 1; + m_lds_addr = lds_addr_in; + m_type = 2; + m_gather_index_size = gather_idx_size_in; + m_gather_mode = gather_mode_in; + } + + + struct + { + //reserve the first 32 bits + union + { + struct + { + uint32_t m_count : 2; + uint32_t m_is_restore : 1; + uint32_t m_is_store : 1; + uint32_t m_nv : 1; + uint32_t m_scope_trait : 2; + uint32_t m_th :3; + uint32_t m_reserved_space : 20; + uint32_t m_gather_index_size : 1; + uint32_t m_gather_mode : 1; + }; + + uint32_t m_reserved0; + }; + + uint32_t m_lds_addr; + uint32_t m_global_addr_lo; + union + { + struct + { + uint32_t m_global_addr_hi : 25; + uint32_t m_reserved2 : 5; + uint32_t m_type : 2; + }; + uint32_t m_sgpr3; + }; + }; + __hip_uint32x4 m_bitfield; + + // setters for all user configurable fields + __device__ void count(bool value) + { + m_count = (int)value; + } + + __device__ void ldsAddr(uint32_t value) + { + m_lds_addr = value; + } + + __device__ void gatherMode(int enable, int value) + { + m_gather_mode = enable; + m_gather_index_size = value; + } + + __device__ void globalAddr(uintptr_t value) + { + m_global_addr_lo = value & 0xFFFFFFFF; + m_global_addr_hi = (value >> 32); + } +}; + +union gfx1250_TDM_GROUP1 +{ + __device__ constexpr gfx1250_TDM_GROUP1() : m_bitfield{0,0,0,0,0,0,0,0} {} + + struct + { + union + { + struct + { + uint32_t m_workgroup_mask : 16; + uint32_t m_data_size : 2; + uint32_t m_atomic_barrier_enable : 1; + uint32_t m_iterate_enable : 1; + uint32_t m_pad_enable : 1; + uint32_t m_early_timeout : 1; + uint32_t m_pad_interval : 3; + uint32_t m_pad_amount : 7; + }; + uint32_t m_sgpr0; + }; + union + { + struct + { + uint32_t m_atomic_barrier_address : 16; + uint32_t m_tensor_dim0_lo : 16; + }; + uint32_t m_sgpr1; + }; + union + { + struct + { + uint32_t m_tensor_dim0_hi : 16; + uint32_t m_tensor_dim1_lo : 16; + }; + uint32_t m_sgpr2; + }; + union + { + struct + { + uint32_t m_tensor_dim1_hi : 16; + uint32_t m_tile_dim0 : 16; + }; + uint32_t m_sgpr3; + }; + union + { + struct + { + uint32_t m_tile_dim1 : 16; + uint32_t m_tile_dim2 : 16; + }; + uint32_t m_sgpr4; + }; + union + { + uint32_t m_tensor_dim0_stride_lo; + uint32_t m_sgpr5; + }; + union + { + struct + { + uint32_t m_tensor_dim0_stride_hi : 16; + uint32_t m_tensor_dim1_stride_lo : 16; + }; + uint32_t m_sgpr6; + }; + union + { + uint32_t m_tensor_dim1_stride_hi; + uint32_t m_sgpr7; + }; + }; + __hip_uint32x8 m_bitfield; + + // setters for all fields + void __device__ inline workgroupMask(uint32_t value) + { + m_workgroup_mask = value; + } + + void __device__ inline dataSize(uint32_t value) + { + m_data_size = value; + } + + void __device__ inline atomicBarrierEnable(bool value) + { + m_atomic_barrier_enable = value; + } + + void __device__ inline iterateEnable(bool value) + { + m_iterate_enable = value; + } + + void __device__ inline padEnable(bool value) + { + m_pad_enable = value; + } + + void __device__ inline earlyTimeout(bool value) + { + m_early_timeout = value; + } + + void __device__ inline padInterval(uint32_t value) + { + m_pad_interval = value; + } + + void __device__ inline padAmount(uint32_t value) + { + m_pad_amount = value; + } + + void __device__ inline atomicBarrierAddress(uint32_t value) + { + m_atomic_barrier_address = value; + } + + void __device__ inline tileDim0(uint32_t value) + { + m_tile_dim0 = value; + } + + void __device__ inline tileDim1(uint32_t value) + { + m_tile_dim1 = value; + } + void __device__ inline tileDim2(uint32_t value) + { + m_tile_dim2 = value; + } + + void __device__ inline tensorDim0(uint32_t value) + { + m_tensor_dim0_lo = value & 0xFFFF; + m_tensor_dim0_hi = (value >> 16); + } + void __device__ inline tensorDim1(uint32_t value) + { + m_tensor_dim1_lo = value & 0xFFFF; + m_tensor_dim1_hi = (value >> 16); + } + void __device__ inline tensorDim0Stride(uint64_t value) + { + m_tensor_dim0_stride_lo = value & 0xFFFFFFFF; + m_tensor_dim0_stride_hi = (value >> 32); + } + + void __device__ inline tensorDim1Stride(uint64_t value) + { + m_tensor_dim1_stride_lo = value & 0xFFFFFFFF; + m_tensor_dim1_stride_hi = (value >> 32); + } +}; + +union gfx1250_TDM_GROUP2 +{ + __device__ gfx1250_TDM_GROUP2() : m_bitfield{0,0,0,0} {} + + struct + { + uint32_t m_tensor_dim2; //sgpr0 + uint32_t m_tensor_dim3; //sgpr1 + uint32_t m_tensor_dim2_stride_lo; //sgpr2 + union + { + struct + { + uint32_t m_tensor_dim2_stride_hi : 16; + uint32_t m_tile_dim3 : 16; + }; + uint32_t m_sgpr3; + }; + }; + __hip_uint32x4 m_bitfield; + + void __device__ inline tensorDim2Stride(uint64_t value) + { + m_tensor_dim2_stride_lo = value & 0xFFFFFFFF; + m_tensor_dim2_stride_hi = value >> 32; + } +}; + +union gfx1250_TDM_GROUP3 +{ + __device__ gfx1250_TDM_GROUP3() : m_bitfield{0,0,0,0} {} + + struct + { + uint32_t m_tensor_dim3_stride_lo; //sgpr0 + union + { + struct + { + uint32_t m_tensor_dim3_stride_hi : 16; + uint32_t m_tensor_dim_4_lo : 16; + }; + uint32_t m_sgpr1; + }; + union + { + struct + { + uint32_t m_tensor_dim_4_hi : 16; + uint32_t m_tile_dim4 : 16; + }; + uint32_t m_sgpr2; + }; + uint32_t m_sgpr3_reserved; + }; + __hip_uint32x4 m_bitfield; + + void __device__ inline tensorDim3Stride(uint64_t value) + { + m_tensor_dim3_stride_lo = value & 0xFFFFFFFF; + m_tensor_dim3_stride_hi = value >> 32; + } + + void __device__ inline tensorDim4(uint32_t value) + { + m_tensor_dim_4_lo = value & 0xFFFF; + m_tensor_dim_4_hi = value >> 16; + } +}; +#pragma clang diagnostic pop +#endif // HIP_INCLUDE_HIP_AMD_GFX1250_TDM_H