From 2530a2f08411e3fb66d94268b6e205e073580f63 Mon Sep 17 00:00:00 2001 From: gilbertlee-amd <44450918+gilbertlee-amd@users.noreply.github.com> Date: Wed, 5 Jan 2022 16:33:25 -0700 Subject: [PATCH] [TransferBench] Updating for 2.11.4. Decoupling from RCCL kernel (#485) --- tools/TransferBench/EnvVars.hpp | 22 +++++ tools/TransferBench/Kernels.hpp | 124 ++++++++++++++++++++++++++ tools/TransferBench/Makefile | 2 +- tools/TransferBench/TransferBench.cpp | 7 +- tools/TransferBench/TransferBench.hpp | 66 -------------- 5 files changed, 151 insertions(+), 70 deletions(-) create mode 100644 tools/TransferBench/Kernels.hpp diff --git a/tools/TransferBench/EnvVars.hpp b/tools/TransferBench/EnvVars.hpp index f8aaaf41e4..ee6254bc3a 100644 --- a/tools/TransferBench/EnvVars.hpp +++ b/tools/TransferBench/EnvVars.hpp @@ -1,3 +1,25 @@ +/* +Copyright (c) 2021-2022 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 ENVVARS_HPP #define ENVVARS_HPP diff --git a/tools/TransferBench/Kernels.hpp b/tools/TransferBench/Kernels.hpp new file mode 100644 index 0000000000..40f8a713bd --- /dev/null +++ b/tools/TransferBench/Kernels.hpp @@ -0,0 +1,124 @@ +/* +Copyright (c) 2022 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. +*/ + +#pragma once + +#define WARP_SIZE 64 +#define BLOCKSIZE 256 + +// GPU copy kernel +__global__ void __launch_bounds__(BLOCKSIZE) +GpuCopyKernel(BlockParam* blockParams) +{ + #define PackedFloat_t float4 + #define FLOATS_PER_PACK (sizeof(PackedFloat_t) / sizeof(float)) + + // Collect the arguments for this threadblock + int Nrem = blockParams[blockIdx.x].N; + float const* src = blockParams[blockIdx.x].src; + float* dst = blockParams[blockIdx.x].dst; + + // Operate on wavefront granularity + int numWaves = BLOCKSIZE / WARP_SIZE; // Number of wavefronts per threadblock + int waveId = threadIdx.x / WARP_SIZE; // Wavefront number + int threadId = threadIdx.x % WARP_SIZE; // Thread index within wavefront + + #define LOOP1_UNROLL 8 + // 1st loop - each wavefront operates on LOOP1_UNROLL x FLOATS_PER_PACK per thread per iteration + // Determine the number of packed floats processed by the first loop + int const loop1Npack = (Nrem / (FLOATS_PER_PACK * LOOP1_UNROLL * WARP_SIZE)) * (LOOP1_UNROLL * WARP_SIZE); + int const loop1Nelem = loop1Npack * FLOATS_PER_PACK; + int const loop1Inc = BLOCKSIZE * LOOP1_UNROLL; + int loop1Offset = waveId * LOOP1_UNROLL * WARP_SIZE + threadId; + + PackedFloat_t const* packedSrc = (PackedFloat_t const*)(src) + loop1Offset; + PackedFloat_t* packedDst = (PackedFloat_t *)(dst) + loop1Offset; + while (loop1Offset < loop1Npack) + { + PackedFloat_t vals[LOOP1_UNROLL]; + #pragma unroll + for (int u = 0; u < LOOP1_UNROLL; ++u) + vals[u] = *(packedSrc + u * WARP_SIZE); + + #pragma unroll + for (int u = 0; u < LOOP1_UNROLL; ++u) + *(packedDst + u * WARP_SIZE) = vals[u]; + + packedSrc += loop1Inc; + packedDst += loop1Inc; + loop1Offset += loop1Inc; + } + Nrem -= loop1Nelem; + if (Nrem == 0) return; + + // 2nd loop - Each thread operates on FLOATS_PER_PACK per iteration + int const loop2Npack = Nrem / FLOATS_PER_PACK; + int const loop2Nelem = loop2Npack * FLOATS_PER_PACK; + int const loop2Inc = BLOCKSIZE; + int loop2Offset = threadIdx.x; + + packedSrc = (PackedFloat_t const*)(src + loop1Nelem); + packedDst = (PackedFloat_t *)(dst + loop1Nelem); + while (loop2Offset < loop2Npack) + { + packedDst[loop2Offset] = packedSrc[loop2Offset]; + loop2Offset += loop2Inc; + } + Nrem -= loop2Nelem; + if (Nrem == 0) return; + + // Deal with leftovers less than FLOATS_PER_PACK) + if (threadIdx.x < Nrem) + { + int offset = loop1Nelem + loop2Nelem + threadIdx.x; + dst[offset] = src[offset]; + } +} + +#define MEMSET_UNROLL 8 +__global__ void __launch_bounds__(BLOCKSIZE) +GpuMemsetKernel(BlockParam* blockParams) +{ + // Collect the arguments for this block + int N = blockParams[blockIdx.x].N; + float* __restrict__ dst = (float*)blockParams[blockIdx.x].dst; + + // Use non-zero value + #pragma unroll MEMSET_UNROLL + for (int tid = threadIdx.x; tid < N; tid += BLOCKSIZE) + { + dst[tid] = 1234.0; + } +} + +// CPU copy kernel +void CpuCopyKernel(BlockParam const& blockParams) +{ + memcpy(blockParams.dst, blockParams.src, blockParams.N * sizeof(float)); +} + +// CPU memset kernel +void CpuMemsetKernel(BlockParam const& blockParams) +{ + for (int i = 0; i < blockParams.N; i++) + blockParams.dst[i] = 1234.0; +} diff --git a/tools/TransferBench/Makefile b/tools/TransferBench/Makefile index 62d1046fb7..d12ca52887 100644 --- a/tools/TransferBench/Makefile +++ b/tools/TransferBench/Makefile @@ -6,7 +6,7 @@ endif HIPCC=$(HIP_PATH)/bin/hipcc EXE=TransferBench -CXXFLAGS = -O3 -I../../src/include -I. -lnuma -L$(HIP_PATH)/../hsa/lib -lhsa-runtime64 +CXXFLAGS = -O3 -I. -lnuma -L$(HIP_PATH)/../hsa/lib -lhsa-runtime64 all: $(EXE) diff --git a/tools/TransferBench/TransferBench.cpp b/tools/TransferBench/TransferBench.cpp index 79a2e9bbe7..187e6e4a21 100644 --- a/tools/TransferBench/TransferBench.cpp +++ b/tools/TransferBench/TransferBench.cpp @@ -22,14 +22,15 @@ THE SOFTWARE. // This program measures simultaneous copy performance across multiple GPUs // on the same node - -#include "TransferBench.hpp" -#include "GetClosestNumaNode.hpp" #include #include #include #include +#include "TransferBench.hpp" +#include "GetClosestNumaNode.hpp" +#include "Kernels.hpp" + // Simple configuration parameters size_t const DEFAULT_BYTES_PER_LINK = (1<<26); // Amount of data transferred per Link diff --git a/tools/TransferBench/TransferBench.hpp b/tools/TransferBench/TransferBench.hpp index 5328fdb74e..ee9fefb5b5 100644 --- a/tools/TransferBench/TransferBench.hpp +++ b/tools/TransferBench/TransferBench.hpp @@ -35,19 +35,6 @@ THE SOFTWARE. #include #include -// Include common_kernel.h from RCCL for copy kernel -// However define some variables to avoid extra includes / missing defines -#define NCCL_DEVICE_H_ // Avoid loading devcomm.h -#define WARP_SIZE 64 -typedef float half; // TransferBench doesn't actually operate on half-precision floats -typedef uint64_t PackType; -typedef ulong2 Pack128; -typedef struct -{ - uint16_t data; -} rccl_bfloat16; - -#include "../../src/collectives/device/common_kernel.h" #include "EnvVars.hpp" // Helper macro for catching HIP errors @@ -132,56 +119,3 @@ std::string GetLinkTypeDesc(uint32_t linkType, uint32_t hopCount); std::string GetDesc(MemType srcMemType, int srcIndex, MemType dstMemType, int dstIndex); std::string GetLinkDesc(Link const& link); - -#define BLOCKSIZE 256 -#define COPY_UNROLL 4 -#define MEMSET_UNROLL 4 - -// Dummy reduction function (not used because it's just a copy) -struct FuncNull { - __device__ float operator()(const float x, const float y) const { - return 0; - } -}; - -// GPU copy kernel -__global__ void __launch_bounds__(BLOCKSIZE) -GpuCopyKernel(BlockParam* blockParams) -{ - // Collect the arguments for this block - int N = blockParams[blockIdx.x].N; - const float* src[1] = {(float* )blockParams[blockIdx.x].src}; - float* dst[1] = {(float* )blockParams[blockIdx.x].dst}; - - ReduceOrCopyMulti( - threadIdx.x, BLOCKSIZE, 1, src, 1, dst, N); -} - -// GPU set kernel -__global__ void __launch_bounds__(BLOCKSIZE) -GpuMemsetKernel(BlockParam* blockParams) -{ - // Collect the arguments for this block - int N = blockParams[blockIdx.x].N; - float* __restrict__ dst = (float*)blockParams[blockIdx.x].dst; - - // Use non-zero value - #pragma unroll MEMSET_UNROLL - for (int tid = threadIdx.x; tid < N; tid += BLOCKSIZE) - { - dst[tid] = 1234.0; - } -} - -// CPU copy kernel -void CpuCopyKernel(BlockParam const& blockParams) -{ - memcpy(blockParams.dst, blockParams.src, blockParams.N * sizeof(float)); -} - -// CPU memset kernel -void CpuMemsetKernel(BlockParam const& blockParams) -{ - for (int i = 0; i < blockParams.N; i++) - blockParams.dst[i] = 1234.0; -}