From 001affc6ae5194b5bdd5f5218855c1be69538d66 Mon Sep 17 00:00:00 2001 From: "Hernandez, Gerardo" Date: Wed, 14 May 2025 11:57:05 +0100 Subject: [PATCH] SWDEV-420237 - Add tests for reduce sync operations (#102) --- catch/hipTestMain/main.cc | 9 +- catch/include/cmd_options.hh | 3 +- catch/include/resource_guards.hh | 7 +- catch/include/warp_common.hh | 519 ++++++++++++++++++++++ catch/performance/CMakeLists.txt | 1 + catch/performance/warpSync/CMakeLists.txt | 31 ++ catch/performance/warpSync/warpSync.cc | 395 ++++++++++++++++ catch/unit/rtc/CMakeLists.txt | 1 + catch/unit/rtc/rtc_reduce.cc | 215 +++++++++ catch/unit/warp/CMakeLists.txt | 1 + catch/unit/warp/warp_common.hh | 178 -------- catch/unit/warp/warp_reduce.cc | 230 ++++++++++ 12 files changed, 1407 insertions(+), 183 deletions(-) create mode 100644 catch/include/warp_common.hh create mode 100644 catch/performance/warpSync/CMakeLists.txt create mode 100644 catch/performance/warpSync/warpSync.cc create mode 100644 catch/unit/rtc/rtc_reduce.cc delete mode 100644 catch/unit/warp/warp_common.hh create mode 100644 catch/unit/warp/warp_reduce.cc diff --git a/catch/hipTestMain/main.cc b/catch/hipTestMain/main.cc index e869695572..03591dc71c 100644 --- a/catch/hipTestMain/main.cc +++ b/catch/hipTestMain/main.cc @@ -30,9 +30,6 @@ int main(int argc, char** argv) { | Opt(cmd_options.progress) ["-P"]["--progress"] ("Show progress bar when running performance tests") - | Opt(cmd_options.cg_extended_run, "cg_extened_run") - ["-E"]["--cg-extended-run"] - ("TODO: Description goes here") | Opt(cmd_options.cg_iterations, "cg_iterations") ["-C"]["--cg-iterations"] ("Number of iterations used for cooperative groups sync tests (default: 5)") @@ -42,6 +39,12 @@ int main(int argc, char** argv) { | Opt(cmd_options.accuracy_max_memory, "accuracy_max_memory") ["-M"]["--accuracy-max-memory"] ("Percentage of global device memory allowed for math accuracy tests (default: 80%)") + | Opt(cmd_options.reduce_iterations, "reduce_iterations") + ["-R"]["--reduce-iterations"] + ("Number of iterations for fuzzing reduce operations (default: 1)") + | Opt(cmd_options.reduce_input_size, "reduce_input_size") + ["-Z"]["--reduce-input-size"] + ("Size of the input for the reduce sync operations performance test (megabytes) (default: 50)") ; // clang-format on diff --git a/catch/include/cmd_options.hh b/catch/include/cmd_options.hh index 666f34ea82..b9307f302d 100644 --- a/catch/include/cmd_options.hh +++ b/catch/include/cmd_options.hh @@ -28,11 +28,12 @@ THE SOFTWARE. struct CmdOptions { int iterations = 10; int warmups = 100; - int cg_extended_run = 5; int cg_iterations = 5; bool no_display = false; bool progress = false; uint64_t accuracy_iterations = std::numeric_limits::max() + 1ull; + uint64_t reduce_iterations = 1; + uint64_t reduce_input_size = 50; int accuracy_max_memory = 80; }; diff --git a/catch/include/resource_guards.hh b/catch/include/resource_guards.hh index c545fad074..c2f0864d1b 100644 --- a/catch/include/resource_guards.hh +++ b/catch/include/resource_guards.hh @@ -55,7 +55,8 @@ template class LinearAllocGuard { LinearAllocGuard(const LinearAllocs allocation_type, const size_t size, const unsigned int flags = 0u) - : allocation_type_{allocation_type} { + : allocation_type_{allocation_type}, + size_{size} { switch (allocation_type_) { case LinearAllocs::malloc: ptr_ = host_ptr_ = reinterpret_cast(malloc(size)); @@ -92,10 +93,12 @@ template class LinearAllocGuard { allocation_type_ = o.allocation_type_; ptr_ = o.ptr_; host_ptr_ = o.host_ptr_; + size_ = o.size_; o.allocation_type_ = LinearAllocs::noAlloc; o.ptr_ = nullptr; o.host_ptr_ = nullptr; + o.size_ = 0; } return *this; @@ -105,11 +108,13 @@ template class LinearAllocGuard { T* ptr() const { return ptr_; }; T* host_ptr() const { return host_ptr_; } + size_t size_bytes() const { return size_; } private: LinearAllocs allocation_type_ = LinearAllocs::noAlloc; T* ptr_ = nullptr; T* host_ptr_ = nullptr; + size_t size_ = 0; void dealloc() { if (ptr_ == nullptr) { diff --git a/catch/include/warp_common.hh b/catch/include/warp_common.hh new file mode 100644 index 0000000000..26c66edb0f --- /dev/null +++ b/catch/include/warp_common.hh @@ -0,0 +1,519 @@ +/* +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. +*/ + +#pragma once + +#define HIP_ENABLE_WARP_SYNC_BUILTINS +#define HIP_ENABLE_EXTRA_WARP_SYNC_TYPES + +#include +#include +#include +#include +#include +#include +#include +#include + +#define MASK_SHIFT(x, n) \ + (x & (static_cast(1) << n)) >> n + +const unsigned long long Every5thBit = 0x1084210842108421; +const unsigned long long Every9thBit = 0x8040201008040201; +const unsigned long long Every5thBut9th = Every5thBit & ~Every9thBit; +const unsigned long long AllThreads = ~0; +static constexpr int kNumReduces = 5000; + +inline __device__ bool deactivate_thread(const uint64_t* const active_masks) { + const auto warp = + cooperative_groups::tiled_partition(cooperative_groups::this_thread_block(), warpSize); + const auto block = cooperative_groups::this_thread_block(); + const auto warps_per_block = (block.size() + warpSize - 1) / warpSize; + const auto block_rank = (blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x; + const auto idx = block_rank * warps_per_block + block.thread_rank() / warpSize; + return !(active_masks[idx] & (static_cast(1) << warp.thread_rank())); +} + +inline std::mt19937& GetRandomGenerator() { + static std::mt19937 mt(std::random_device{}()); + return mt; +} + +template inline T GenerateRandomInteger(const T min, const T max) { + std::uniform_int_distribution dist(min, max); + return dist(GetRandomGenerator()); +} + +template inline T GenerateRandomReal(const T min, const T max) { + std::uniform_real_distribution dist(min, max); + return dist(GetRandomGenerator()); +} + +inline int generate_width(int warp_size) { + int exponent = 0; + while (warp_size >>= 1) { + ++exponent; + } + + return GENERATE_COPY(map([](int e) { return 1 << e; }, range(1, exponent + 1))); +} + +inline uint64_t get_active_mask(unsigned int warp_id, unsigned int warp_size) { + uint64_t active_mask = 0; + switch (warp_id % 5) { + case 0: // even threads in the warp + active_mask = 0xAAAAAAAAAAAAAAAA; + break; + case 1: // odd threads in the warp + active_mask = 0x5555555555555555; + break; + case 2: // first half of the warp + for (int i = 0; i < warp_size / 2; i++) { + active_mask = active_mask | (static_cast(1) << i); + } + break; + case 3: // second half of the warp + for (int i = warp_size / 2; i < warp_size; i++) { + active_mask = active_mask | (static_cast(1) << i); + } + break; + case 4: // all threads + active_mask = 0xFFFFFFFFFFFFFFFF; + break; + } + return active_mask; +} + +template ::value, bool> = true> +inline T expandPrecision(int X) { return X; } + +template ::value, bool> = true> +inline T expandPrecision(int X) { + return X * 3.141592653589793115997963468544185161590576171875; +} + +template ::value, bool> = true> +inline __half expandPrecision(int X) { + return (__half)expandPrecision(X); +} + +template ::value, bool> = true> +inline __half2 expandPrecision(int X) { + __half H = expandPrecision(X); + return {H, H}; +} + +template ::value, bool> = true> +inline void expandPrecision(T* Array, int size) { + (void)Array; + (void)size; +} + +template ::value, bool> = true> +inline void expandPrecision(T *Array, int size) { + for (int i = 0; i != size; ++i) { + Array[i] *= 3.141592653589793115997963468544185161590576171875; + } +} + +template +inline void initializeInput(T *Input, int size) { + int Values[] = {0, -1, 2, 3, 4, 5, -6, 7, + 8, -9, 10, 11, 12, 13, -14, 15, + 16, 17, -18, 19, 20, -21, 22, 23, + 24, 25, 26, -27, 28, 29, 30, 31, + -32, 33, 34, 35, -36, 37, 38, -39, + 40, 41, 42, 43, -44, -45, 46, 47, + 48, 49, 50, -51, 52, 53, -54, 55, + 56, 57, -58, 59, 60, 61, 62, -63}; + + for (int i = 0; i != size; ++i) { + Input[i] = expandPrecision(Values[i]); + } +} + +template +inline void initializeExpected(T *Expected, int *Values, int size) { + for (int i = 0; i != size; ++i) { + Expected[i] = expandPrecision(Values[i]); + } +} + +template +inline bool compareEqual(T X, T Y) { return X == Y; } + +template <> +inline bool compareEqual(__half X, __half Y) { + return __half2float(X) == __half2float(Y); +} + +template <> +inline bool compareEqual(__half2 X, __half2 Y) { + return compareEqual(X.x, Y.x) && compareEqual(X.y, Y.y); +} + +inline bool compareMaskEqual(unsigned long long *Actual, unsigned long long *Expected, + int i, int warpSize) { + if (warpSize == 32) + return (unsigned)Actual[i] == (unsigned)Expected[i]; + return Actual[i] == Expected[i]; +} + +template +inline T alignUp(T num, size_t n) { + if (num % n == 0) { + return num; + } + + return ((num + n - 1) / n) * n; +} + +template +struct DistributionType { + using type = std::uniform_int_distribution; +}; + +// there is no std::uniform_real_distribution for 'half' type, so we cast from +// unsigned short, avoiding Nan and Infinity +template <> +struct DistributionType<__half> { + using type = std::uniform_int_distribution; +}; + +template <> +struct DistributionType { + using type = std::uniform_real_distribution; +}; + +template <> +struct DistributionType { + using type = std::uniform_real_distribution; +}; + + +template +struct MinOp { + T operator()(const T& lhs, const T& rhs) const + { + return std::min(lhs, rhs); + } +}; + +template +struct MaxOp { + T operator()(const T& lhs, const T& rhs) const + { + return std::max(lhs, rhs); + } +}; + +template +struct XorOp { + __host__ __device__ T operator()(const T& lhs, const T& rhs) + { + return (!lhs) != (!rhs) == 1; + } +}; + +// typeid(T).name() does seem to return a very descriptive name for primitive types, +// at least on clang, so we roll out an equivalent +template +const char* typeToString() +{ + if (std::is_same::value) + return "int"; + if (std::is_same::value) + return "unsigned int"; + if (std::is_same::value) + return "long long"; + if (std::is_same::value) + return "unsigned long long"; + if (std::is_same::value) + return "half"; + if (std::is_same::value) + return "float"; + if (std::is_same::value) + return "double"; + + assert(false && "Missing conversion to string for type"); + return ""; +} + +template class Op> +const char* opToString() +{ + if constexpr (std::is_same, std::plus>::value) + return "add"; + else if constexpr (std::is_same, MinOp>::value) + return "min"; + else if constexpr (std::is_same, MaxOp>::value) + return "max"; + else if constexpr (std::is_same, std::logical_and>::value) + return "logical_and"; + else if constexpr (std::is_same, std::logical_or>::value) + return "logical_or"; + else if constexpr (std::is_same, XorOp>::value) + return "logical_xor"; + else { + static_assert(std::is_void::value, "Unsupported operator"); + return ""; + } +} + +template +void genRandomMasks(LinearAllocGuard& d_buf, + LinearAllocGuard& buf, + Gen& gen, + int numItems) +{ + // masks must be != 0, hence passing 1 as the 'a' distribution parameter + std::uniform_int_distribution dist(1); + int numBytes = numItems * sizeof(T); + LinearAllocGuard tmp(LinearAllocs::malloc, numBytes); + LinearAllocGuard d_tmp(LinearAllocs::hipMalloc, numBytes); + + buf = std::move(tmp); + d_buf = std::move(d_tmp); + + for (int i = 0; i < numItems; i++) { + T mask = dist(gen); + + if (getWarpSize() == 32) + mask &= 0xFFFFFFFF; + + buf.ptr()[i] = mask; + } + + HIP_CHECK(hipMemcpy(d_buf.ptr(), buf.ptr(), numBytes, hipMemcpyHostToDevice)); +} + +// generates a random __half (instead of using uniform_real_distribution casting to __half +// which is problematic) +// @expDist needs to be between [0-2^5-2] +template +__half genRandomHalf(std::uniform_int_distribution& dist, + Gen& gen) +{ + __half_raw tmp; + + tmp.x = dist(gen); + // rewrite the exponent to force the number to be (-8 +void genRandomBuffers(LinearAllocGuard& d_buf, + LinearAllocGuard& buf, + Dist& dist, + Gen& gen, + int numItems) +{ + int numBytes = numItems * sizeof(T); + LinearAllocGuard tmp(LinearAllocs::malloc, numBytes); + LinearAllocGuard d_tmp(LinearAllocs::hipMalloc, numBytes); + + buf = std::move(tmp); + d_buf = std::move(d_tmp); + + for (int i = 0; i < numItems; i++) + if constexpr (std::is_same::value) + buf.ptr()[i] = genRandomHalf(dist, gen); + else + buf.ptr()[i] = dist(gen); + + HIP_CHECK(hipMemcpy(d_buf.ptr(), buf.ptr(), numBytes, hipMemcpyHostToDevice)); +} + +// given an operation produces the expected result of the reduction +// @mask indicates the lanes that will participate in the computation +template +T calculateExpected(const T* input, Op op, unsigned long long mask) +{ + T result; + int wavefrontSize = getWarpSize(); + + if (std::is_same>::value) { + T tmp[64] = { 0 }; + + for (int i = 0; i < wavefrontSize; i++) { + if (mask & (1ul << i)) { + tmp[i] = input[i]; + } + } + + for (int modulo = 2; modulo <= wavefrontSize; modulo *= 2) { + for (int i = 0; i < wavefrontSize; i += modulo) { + int j = i + modulo / 2; + + if (j < wavefrontSize) + tmp[i] += tmp[j]; + } + } + result = tmp[0]; + } else { + bool initialized = false; + + for (int i = 0; i < wavefrontSize; i++) { + if (mask & (1ul << i)) { + if (initialized) + result = op(input[i], result); + else { + result = input[i]; + initialized = true; + } + } + } + } + return result; +} + +template +void printMismatch(const T& result, const T& expected, const T* input, unsigned long long mask) +{ + std::ios init(NULL); + + init.copyfmt(std::cout); + std::cout << "\nMismatch\n"; + std::cout << "Mask: 0x" << std::hex << std::setfill('0') << std::setw(16) << mask << "\n"; + std::cout.copyfmt(init); + + for (int i = 0; i < getWarpSize(); i++) { + if ((1ul << i) & mask) { + if constexpr (std::is_same::value) + std::cout << "Lane " << i << ": " << __half2float(input[i]) << "\n"; + else + std::cout << "Lane " << i << ": " << input[i] << "\n"; + } + } + + if constexpr (std::is_same::value) { + std::cout << "Result: " << __half2float(result) << "\n"; + std::cout << "Expected: " << __half2float(expected) << "\n"; + } else { + std::cout << "Result: " << result << "\n"; + std::cout << "Expected: " << expected << "\n"; + } +} + +template +void compareFloatingPoint(const T& result, const T& expected, unsigned long long mask, const T* input) +{ + using namespace Catch::Matchers; + if constexpr (std::is_same::value) { + float resultFloat = __half2float(result); + float expectedFloat = __half2float(expected); + float absDifference = fabs(resultFloat - expectedFloat); + float relativeEpsilon = 0.1 * fmax(resultFloat, expectedFloat); + float eps = 0.01f; + + REQUIRE(!__hisnan(result)); + REQUIRE(!__hisinf(result)); + + if (relativeEpsilon > eps) { + if (absDifference > 0.0001) { + if (absDifference >= eps * fabs(fmax(resultFloat, expectedFloat))) { + printMismatch(result, expected, input, mask); + std::cout << "Relative epsilon: " << relativeEpsilon << "\n"; + std::cout << "Difference: " << absDifference << "\n"; + } + } + + REQUIRE_THAT(__half2float(resultFloat), WithinRel(expectedFloat, eps)); + } + } else { + // for float or double, also lossy in terms of precision + T absDifference = fabs(result - expected); + T relativeEpsilon = 0.1 * fmax(result, expected); + T eps = 0.01; + + if (relativeEpsilon > eps) { + if (absDifference > 0.0001) { + if (absDifference >= eps * fabs(fmax(result, expected))) { + printMismatch(result, expected, input, mask); + std::cout << "Relative epsilon: " << relativeEpsilon << "\n"; + std::cout << "Difference: " << absDifference << "\n"; + } + + REQUIRE_THAT(result, WithinRel(expected, eps)); + } + } + } +} + +// @tparam Reduce a functor; abstracts away kernel dispatching +// (via hiprtc or normal execution) +template class Op> +void runTestReduce(int iteration, Reduce reduce) +{ + using namespace Catch::Matchers; + using distribution = typename DistributionType::type; + unsigned int wavefrontSize = getWarpSize(); + // one result per reduce per thread to be checked + LinearAllocGuard d_output(LinearAllocs::hipMalloc, kNumReduces * wavefrontSize * sizeof(T)); + LinearAllocGuard output(LinearAllocs::malloc, kNumReduces * wavefrontSize * sizeof(T)); + std::mt19937_64 gen(iteration); + // for float16, we generate any random unsigned short, but cap the exponent later on + // to keep it in the range (-8.0..8.0) (just to avoid overflows) + // On the rest of the types, just use a bigger reduced range of numbers to avoid overflows too + T a = std::is_same::value? std::numeric_limits::lowest() : -1023; + T b = std::is_same::value? std::numeric_limits::max() : 1023; + distribution dist(a, b); + LinearAllocGuard input, d_input; + LinearAllocGuard masks, d_masks; + Op op; + int numReduce = 0; + + genRandomBuffers(d_input, input, dist, gen, kNumReduces * wavefrontSize); + genRandomMasks(d_masks, masks, gen, kNumReduces); + reduce(d_output.ptr(), d_input.ptr(), d_masks.ptr(), kNumReduces, op); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(output.ptr(), d_output.ptr(), d_output.size_bytes(), hipMemcpyDeviceToHost)); + + while (numReduce < kNumReduces) { + T expected = calculateExpected(input.ptr(), op, masks.ptr()[numReduce]); + int lane = 0; + + while (lane < wavefrontSize) { + auto result = output.ptr()[numReduce * wavefrontSize + lane]; + unsigned long long mask = masks.ptr()[numReduce]; + + if ((1ul << lane) & mask) { + if constexpr (std::is_integral::value || std::is_same, MinOp>::value || + std::is_same, MaxOp>::value) { + // for integral types or min/max the result should match exactly + if constexpr (std::is_same::value) + REQUIRE(__half2float(result) == __half2float(expected)); + else { + if (result != expected) { + printMismatch(result, expected, input.ptr(), mask); + REQUIRE(result == expected); + } + } + } else + compareFloatingPoint(result, expected, mask, input.ptr()); + + } + lane++; + } + numReduce++; + } +} diff --git a/catch/performance/CMakeLists.txt b/catch/performance/CMakeLists.txt index 2778dab03d..72dffe2111 100644 --- a/catch/performance/CMakeLists.txt +++ b/catch/performance/CMakeLists.txt @@ -23,4 +23,5 @@ add_subdirectory(memcpy) add_subdirectory(kernelLaunch) add_subdirectory(stream) add_subdirectory(event) +add_subdirectory(warpSync) add_subdirectory(example) diff --git a/catch/performance/warpSync/CMakeLists.txt b/catch/performance/warpSync/CMakeLists.txt new file mode 100644 index 0000000000..7ad4a44a95 --- /dev/null +++ b/catch/performance/warpSync/CMakeLists.txt @@ -0,0 +1,31 @@ +# Copyright (c) 2024 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. + +# As CUDA use 32-bit sync masks and not 64-bit this has been disabled on Nvidia +if(HIP_PLATFORM MATCHES "amd") + set(TEST_SRC + warpSync.cc + ) +endif() + +hip_add_exe_to_target(NAME WarpSyncPerformance + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests + COMPILE_OPTIONS -std=c++17) diff --git a/catch/performance/warpSync/warpSync.cc b/catch/performance/warpSync/warpSync.cc new file mode 100644 index 0000000000..2e67d8c506 --- /dev/null +++ b/catch/performance/warpSync/warpSync.cc @@ -0,0 +1,395 @@ +/* +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. +*/ +#define HIP_ENABLE_WARP_SYNC_BUILTINS +#define HIP_ENABLE_EXTRA_WARP_SYNC_TYPES + +#include "warp_common.hh" +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +/** +* @addtogroup __reduce_op_sync __reduce_op_sync +* @{ +* @ingroup WarpSyncPerformance +* __reduce_op_sync(MaskT mask, T val) +* Reduces the val as per the lanes described in mask and calculates the +* aggregated result +*/ + +static constexpr int kBlockDim = 1024; + +template +struct AtomicAddOp { + __device__ T operator()(T* lhs, const T& rhs) + { + return atomicAdd(lhs, rhs); + } +}; + +template +struct AtomicMinOp { + __device__ T operator()(T* lhs, const T& rhs) + { + return atomicMin(lhs, rhs); + } +}; + +template +struct AtomicMaxOp { + __device__ T operator()(T* lhs, const T& rhs) + { + return atomicMax(lhs, rhs); + } +}; + +template +struct AtomicAndOp { + __device__ T operator()(T* lhs, const T& rhs) + { + return atomicAnd(lhs, rhs); + } +}; + +template +struct AtomicOrOp { + __device__ T operator()(T* lhs, const T& rhs) + { + return atomicOr(lhs, rhs); + } +}; + +template +struct AtomicXorOp { + __device__ T operator()(T* lhs, const T& rhs) + { + return atomicXor(lhs, rhs); + } +}; + +// uses atomics to reduce the whole warp; depending on the mask our reduce should be faster +// @output to store the result, one per warp +// @numItems must be a multiple of warpSize +template class Op> +__global__ void reduceAllAtomics(T* __restrict__ output, const T* __restrict__ input, unsigned long long mask) +{ + int idx = threadIdx.x + blockIdx.x * kBlockDim; + __shared__ T result[kBlockDim / warpSize]; // one per warp + Op op; + int numWarp = threadIdx.x / warpSize; + + // initialize result[numWarp] to the "identity" element for Op + if constexpr (std::is_same, AtomicMinOp>::value) + result[numWarp] = std::numeric_limits::max(); + else if constexpr (std::is_same, AtomicMaxOp>::value) + result[numWarp] = std::numeric_limits::lowest(); + else if constexpr (std::is_same, AtomicAndOp>::value) + result[numWarp] = 1; + else + result[numWarp] = 0; + + __syncthreads(); + + if (mask & (1ul << __ockl_lane_u32())) + op(&result[numWarp], input[idx]); + + __syncthreads(); + + if (__ockl_lane_u32() == 0) + output[idx / warpSize] = result[numWarp]; +} + +template class Op> +__global__ void reduceOpSync(T* __restrict__ output, const T* __restrict__ input, unsigned long long mask) +{ + int idx = threadIdx.x + blockIdx.x * kBlockDim; + T result; + + if (mask & (1ul << __ockl_lane_u32())) { + if constexpr (std::is_same, std::plus>::value) + result = __reduce_add_sync(mask, input[idx]); + else if constexpr (std::is_same, MinOp>::value) + result = __reduce_min_sync(mask, input[idx]); + else if constexpr (std::is_same, MaxOp>::value) + result = __reduce_max_sync(mask, input[idx]); + else if constexpr (std::is_same, std::logical_and>::value) + result = __reduce_and_sync(mask, input[idx]); + else if constexpr (std::is_same, std::logical_or>::value) + result = __reduce_or_sync(mask, input[idx]); + else if constexpr (std::is_same, XorOp>::value) + result = __reduce_xor_sync(mask, input[idx]); + else + static_assert(std::is_void::value, "Unsupported operator"); + + if (__ockl_activelane_u32() == 0) + output[idx / warpSize] = result; + } +} + +template class Op> +class AtomicBenchmark : public Benchmark> { +public: + void operator()(T* output, const T* input, int numItems, unsigned long long mask) + { + dim3 blockDim = { kBlockDim }; + dim3 gridDim = { static_cast(std::ceil(numItems / static_cast(blockDim.x))) }; + + TIMED_SECTION(kTimerTypeEvent) { + if constexpr (std::is_same, std::plus>::value) + reduceAllAtomics<<>>(output, input, mask); + else if constexpr (std::is_same, MinOp>::value) + reduceAllAtomics<<>>(output, input, mask); + else if constexpr (std::is_same, MaxOp>::value) + reduceAllAtomics<<>>(output, input, mask); + else if constexpr (std::is_same, std::logical_and>::value) + reduceAllAtomics<<>>(output, input, mask); + else if constexpr (std::is_same, std::logical_or>::value) + reduceAllAtomics<<>>(output, input, mask); + else if constexpr (std::is_same, XorOp>::value) + reduceAllAtomics<<>>(output, input, mask); + else + static_assert(std::is_void::value, "Unsupported operator"); + + HIP_CHECK(hipDeviceSynchronize()); + } + } +}; + +template class Op> +class ReduceSyncBenchmark : public Benchmark> { +public: + void operator()(T* output, T* input, int numItems, unsigned long long mask) + { + dim3 blockDim = { kBlockDim }; + dim3 gridDim = { static_cast(std::ceil(numItems / static_cast(blockDim.x))) }; + + + TIMED_SECTION(kTimerTypeEvent) { + reduceOpSync<<>>(output, input, mask); + HIP_CHECK(hipDeviceSynchronize()); + } + } +}; + +template class Op> +void checkResults(T* d_atomicsResult, T* d_reduceResult, size_t numBytes, unsigned long long mask) +{ + using namespace Catch::Matchers; + LinearAllocGuard outputAtomic(LinearAllocs::malloc, numBytes); + LinearAllocGuard outputReduce(LinearAllocs::malloc, numBytes); + bool memcmpResult = std::memcmp(outputAtomic.ptr(), outputReduce.ptr(), numBytes); + + assert(numBytes % sizeof(T) == 0 && "numBytes needs to be a multiple of sizeof(T)"); + HIP_CHECK(hipMemcpy(outputAtomic.ptr(), d_atomicsResult, numBytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(outputReduce.ptr(), d_reduceResult, numBytes, hipMemcpyDeviceToHost)); + + if (memcmpResult) { + for (int i = 0; i < numBytes / sizeof(T); i++) { + auto& atomicResult = outputAtomic.ptr()[i]; + auto& reduceResult = outputReduce.ptr()[i]; + + if constexpr (std::is_integral::value || std::is_same, MinOp>::value || + std::is_same, MaxOp>::value) + // for integral types or min/max the result should match exactly + REQUIRE(atomicResult == reduceResult); + else + // floating point types or operations which are lossy in terms of precision + REQUIRE_THAT(reduceResult, WithinRel(atomicResult)); + } + } +} + +template class Op> +struct IsLogicalOp { + static constexpr bool value = false; +}; + +template +struct IsLogicalOp { + static constexpr bool value = true; +}; + +template +struct IsLogicalOp { + static constexpr bool value = true; +}; + +template +struct IsLogicalOp { + static constexpr bool value = true; +}; + +// Neither long long or fp16 have atomic operations. In those cases +// we only benchmark reduce sync operations, we cannot compare with native atomics +template +struct HasAtomicOps { + static constexpr bool value = true; +}; + +template <> +struct HasAtomicOps { + static constexpr bool value = false; +}; + +template <> +struct HasAtomicOps { + static constexpr bool value = false; +}; + +template class Op> +struct ReduceBenchmark { + void Run() + { + static constexpr int numMasks = 6; + using distribution = typename DistributionType::type; + ReduceSyncBenchmark benchmarkReduce; + uint64_t inputSize = cmd_options.reduce_input_size * 1_MB; + int numItems = inputSize / sizeof(T); + int wavefrontSize = getWarpSize(); + int outputNumBytes = inputSize / wavefrontSize; + LinearAllocGuard input(LinearAllocs::malloc, inputSize); + LinearAllocGuard d_input(LinearAllocs::hipMalloc, inputSize); + LinearAllocGuard d_outputsAtomic[numMasks]; + LinearAllocGuard d_outputsReduce[numMasks]; + LinearAllocGuard* d_outputAtomic = &d_outputsAtomic[0]; + LinearAllocGuard* d_outputReduce = &d_outputsReduce[0]; + std::mt19937_64 gen(123); + distribution dist; + int halfWaveSize = wavefrontSize / 2; + unsigned long long halfBitsOn = (1ul << (wavefrontSize / 2)) - 1; + unsigned long long fullMask = -1ul, + halfHighBitsOn = halfBitsOn << halfWaveSize, + high16BitsOn = halfBitsOn << (wavefrontSize - 16), + high8BitsOn = halfBitsOn << (wavefrontSize - 8), + high4BitsOn = halfBitsOn << (wavefrontSize - 4), + allButOne = -1 & ~1; + const char* typeStr = typeToString(); + const char* opStr = opToString(); + std::map masks; + std::pair masksPairs[] = { { "full mask", fullMask }, + { "high order 32 bits on", halfHighBitsOn }, + { "high order 16 bits on", high16BitsOn }, + { "high order 8 bits on", high8BitsOn }, + { "high order 4 bits on", high4BitsOn }, + { "all but one", allButOne } }; + int pos = 0, numMask = 0; + + for (auto& mask : masksPairs) { + // don't use 'halfHighBitsOn' on warp size 32; it's the same as high16BitsOn + if (wavefrontSize != 32 || mask.second != halfHighBitsOn) { + masks.emplace(std::to_string(numMask) + " - " + mask.first, wavefrontSize == 64? mask.second : mask.second & 0xFFFFFFFF); + numMask++; + } + } + + // avoid generating values different than 1 or 0 for logical operators; + // otherwise the atomic version of the kernels would produce different results as + // atomicAnd/Or() are bitwise operations, not logical + if constexpr (IsLogicalOp::value) + dist = distribution(0, 1); + + for (int i = 0; i < numItems; i++) { + input.ptr()[i] = dist(gen); + } + + for (auto& buffer : d_outputsAtomic) { + buffer = LinearAllocGuard(LinearAllocs::hipMalloc, outputNumBytes); + } + + for (auto& buffer : d_outputsReduce) { + buffer = LinearAllocGuard(LinearAllocs::hipMalloc, outputNumBytes); + } + + HIP_CHECK(hipMemcpy(d_input.ptr(), input.ptr(), inputSize, hipMemcpyHostToDevice)); + + if constexpr (HasAtomicOps::value) { + AtomicBenchmark benchmarkAtomics; + + printf("\n--- atomics %s %s---\n", opStr, typeStr); + + for (auto& mask : masks) { + printf("%s %llx\n", mask.first.c_str(), mask.second); + benchmarkAtomics.Run((d_outputAtomic++)->ptr(), d_input.ptr(), numItems, mask.second); + } + } + + printf("\n--- reduce %s %s--- \n", opStr, typeStr); + + for (const auto& mask : masks) { + printf("%s %llx\n", mask.first.c_str(), mask.second); + benchmarkReduce.Run((d_outputReduce++)->ptr(), d_input.ptr(), numItems, mask.second); + } + + printf("\n"); + + if constexpr (HasAtomicOps::value) { + printf("Checking results...\n"); + + for (const auto& mask : masks) { + checkResults(d_outputsAtomic[pos].ptr(), d_outputsReduce[pos].ptr(), outputNumBytes, mask.second); + pos++; + } + } + } +}; + +TEMPLATE_TEST_CASE("Performance_Reduce_Sync_Add", "", int, unsigned int, unsigned long long, + long long, float, half, double) { + ReduceBenchmark benchmark; + + benchmark.Run(); +} + +TEMPLATE_TEST_CASE("Performance_Reduce_Sync_Min", "", int, unsigned int, unsigned long long, long long, float, half, double) { + ReduceBenchmark benchmark; + + benchmark.Run(); +} + +TEMPLATE_TEST_CASE("Performance_Reduce_Sync_Max", "", int, unsigned int, unsigned long long, long long, float, half, double) { + ReduceBenchmark benchmark; + + benchmark.Run(); +} + +TEMPLATE_TEST_CASE("Performance_Reduce_Sync_And", "", int, unsigned int, unsigned long long, long long) { + ReduceBenchmark benchmark; + + benchmark.Run(); +} + +TEMPLATE_TEST_CASE("Performance_Reduce_Sync_Or", "", int, unsigned int, unsigned long long, long long) { + ReduceBenchmark benchmark; + + benchmark.Run(); +} + +TEMPLATE_TEST_CASE("Performance_Reduce_Sync_Xor", "", int, unsigned int, unsigned long long, long long) { + ReduceBenchmark benchmark; + + benchmark.Run(); +} diff --git a/catch/unit/rtc/CMakeLists.txt b/catch/unit/rtc/CMakeLists.txt index c7d6cd022e..db606ce794 100644 --- a/catch/unit/rtc/CMakeLists.txt +++ b/catch/unit/rtc/CMakeLists.txt @@ -16,6 +16,7 @@ set(AMD_TEST_SRC linker.cc shfl.cc shfl_sync.cc + rtc_reduce.cc stdheaders.cc hiprtc_MathConstants_HeaderTst.cc hiprtc_VectorTypes_HeaderTst.cc diff --git a/catch/unit/rtc/rtc_reduce.cc b/catch/unit/rtc/rtc_reduce.cc new file mode 100644 index 0000000000..d7eec6117d --- /dev/null +++ b/catch/unit/rtc/rtc_reduce.cc @@ -0,0 +1,215 @@ +/* +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. +*/ +#define HIP_ENABLE_WARP_SYNC_BUILTINS +#define HIP_ENABLE_EXTRA_WARP_SYNC_TYPES + +#include "warp_common.hh" +#include +#include +#include +#include +#include + +#define NELEMS(array) (sizeof(array) / sizeof(array[0])) + +// compiles the program, reusing the same compiling session for all the types +// (as opposed as calling the rtc compiler for each of the types) +template