From baabbf32d1faa52070dffbeb912b3cfa3e9e7f1c Mon Sep 17 00:00:00 2001 From: Nives Vukovic Date: Fri, 8 Dec 2023 08:59:41 +0000 Subject: [PATCH] EXSWHTEC-272 - Implement tests for warp shfl_up and shfl_down functions (#193) Change-Id: I3013d16f48ad5f607ee0f252b497fde24c7b9164 [ROCm/hip-tests commit: 26a52506733c1ab2466ff696b78ad854fe386b82] --- projects/hip-tests/catch/unit/CMakeLists.txt | 1 + .../hip-tests/catch/unit/warp/CMakeLists.txt | 2 + .../catch/unit/warp/warp_shfl_down.cc | 121 ++++++++++++++++++ .../hip-tests/catch/unit/warp/warp_shfl_up.cc | 120 +++++++++++++++++ 4 files changed, 244 insertions(+) create mode 100644 projects/hip-tests/catch/unit/warp/warp_shfl_down.cc create mode 100644 projects/hip-tests/catch/unit/warp/warp_shfl_up.cc diff --git a/projects/hip-tests/catch/unit/CMakeLists.txt b/projects/hip-tests/catch/unit/CMakeLists.txt index 7d44709d1d..cb317642b3 100644 --- a/projects/hip-tests/catch/unit/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/CMakeLists.txt @@ -35,6 +35,7 @@ add_subdirectory(multiThread) add_subdirectory(compiler) add_subdirectory(errorHandling) add_subdirectory(cooperativeGrps) +add_subdirectory(warp) add_subdirectory(context) add_subdirectory(warp) add_subdirectory(dynamicLoading) diff --git a/projects/hip-tests/catch/unit/warp/CMakeLists.txt b/projects/hip-tests/catch/unit/warp/CMakeLists.txt index cd2ec5226b..aff9c2747e 100644 --- a/projects/hip-tests/catch/unit/warp/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/warp/CMakeLists.txt @@ -2,6 +2,8 @@ set(TEST_SRC warp_shfl_xor.cc warp_shfl.cc + warp_shfl_up.cc + warp_shfl_down.cc ) hip_add_exe_to_target(NAME WarpTest diff --git a/projects/hip-tests/catch/unit/warp/warp_shfl_down.cc b/projects/hip-tests/catch/unit/warp/warp_shfl_down.cc new file mode 100644 index 0000000000..d42a19c7d8 --- /dev/null +++ b/projects/hip-tests/catch/unit/warp/warp_shfl_down.cc @@ -0,0 +1,121 @@ +/* +Copyright (c) 2023 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. +*/ + +#include "warp_shfl_common.hh" + +#include + +/** + * @addtogroup shfl_down shfl_down + * @{ + * @ingroup DeviceLanguageTest + * `T __shfl_down(T var, unsigned int lane_delta, int width = warpSize)` - + * Contains unit test for warp shfl_down function + */ + +namespace cg = cooperative_groups; + +template +__global__ void shfl_down(T* const out, const T* const in, const uint64_t* const active_masks, + const unsigned int* const deltas, const int width) { + if (deactivate_thread(active_masks)) { + return; + } + + const auto grid = cg::this_grid(); + const auto block = cg::this_thread_block(); + T var = in[grid.thread_rank()]; + out[grid.thread_rank()] = __shfl_down(var, deltas[block.thread_rank() % width], width); +} + +template class WarpShflDown : public WarpShflTest, T> { + public: + void launch_kernel(T* const arr_dev, T* const input_dev, const uint64_t* const active_masks) { + width_ = generate_width(this->warp_size_); + INFO("Width: " << width_); + const auto alloc_size = width_ * sizeof(unsigned int); + LinearAllocGuard deltas_dev(LinearAllocs::hipMalloc, alloc_size); + deltas_.resize(width_); + std::generate(deltas_.begin(), deltas_.end(), + [this] { return GenerateRandomInteger(0u, static_cast(width_)); }); + HIP_CHECK(hipMemcpy(deltas_dev.ptr(), deltas_.data(), alloc_size, hipMemcpyHostToDevice)); + shfl_down<<grid_.grid_dim_, this->grid_.block_dim_>>>(arr_dev, input_dev, active_masks, + deltas_dev.ptr(), width_); + } + + void validate(const T* const arr, const T* const input) { + ArrayAllOf(arr, this->grid_.thread_count_, [this, &input](unsigned int i) -> std::optional { + const int rank_in_block = this->grid_.thread_rank_in_block(i).value(); + const auto rank_in_warp = rank_in_block % this->warp_size_; + const auto rank_in_partition = rank_in_block % width_; + const auto mask_idx = this->warps_in_block_ * (i / this->grid_.threads_in_block_count_) + + rank_in_block / this->warp_size_; + const unsigned int delta = deltas_[rank_in_partition] % width_; + const std::bitset active_mask(this->active_masks_[mask_idx]); + + const int target = rank_in_block % width_ + delta; + if (!active_mask.test(rank_in_warp) || + (target < width_ && !active_mask.test(rank_in_warp + delta)) || + (target < width_ && rank_in_block + delta >= this->grid_.threads_in_block_count_)) { + return std::nullopt; + } + + return (target >= width_ ? input[i] : input[i + delta]); + }); + }; + + private: + std::vector deltas_; + int width_; +}; + +/** + * Test Description + * ------------------------ + * - Validates the warp shuffle down behavior for all valid width sizes {2, 4, 8, 16, 32, + * 64(if supported)} for generated delta values. The threads are deactivated based on the + * passed active mask. The test is run for all overloads of shfl_down. + * Test source + * ------------------------ + * - unit/warp/warp_shfl_down.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + * - Device supports warp shuffle + */ +TEMPLATE_TEST_CASE("Unit_Warp_Shfl_Down_Positive_Basic", "", int, unsigned int, long, unsigned long, + long long, unsigned long long, float, double) { + int device; + hipDeviceProp_t device_properties; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&device_properties, device)); + + if (!device_properties.arch.hasWarpShuffle) { + HipTest::HIP_SKIP_TEST("Device doesn't support Warp Shuffle!"); + return; + } + + SECTION("Shfl Down with specified active mask and input values") { + WarpShflDown().run(false); + } + + SECTION("Shfl Down with random active mask and input values") { + WarpShflDown().run(true); + } +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/warp/warp_shfl_up.cc b/projects/hip-tests/catch/unit/warp/warp_shfl_up.cc new file mode 100644 index 0000000000..201289f363 --- /dev/null +++ b/projects/hip-tests/catch/unit/warp/warp_shfl_up.cc @@ -0,0 +1,120 @@ +/* +Copyright (c) 2023 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. +*/ + +#include "warp_shfl_common.hh" + +#include + +/** + * @addtogroup shfl_up shfl_up + * @{ + * @ingroup DeviceLanguageTest + * `T __shfl_up(T var, unsigned int lane_delta, int width = warpSize)` - + * Contains unit test for warp shfl_up function + */ + +namespace cg = cooperative_groups; + +template +__global__ void shfl_up(T* const out, const T* const in, const uint64_t* const active_masks, + const unsigned int* const deltas, const int width) { + if (deactivate_thread(active_masks)) { + return; + } + + const auto grid = cg::this_grid(); + const auto block = cg::this_thread_block(); + T var = in[grid.thread_rank()]; + out[grid.thread_rank()] = __shfl_up(var, deltas[block.thread_rank() % width], width); +} + +template class WarpShflUp : public WarpShflTest, T> { + public: + void launch_kernel(T* const arr_dev, T* const input_dev, const uint64_t* const active_masks) { + width_ = generate_width(this->warp_size_); + INFO("Width: " << width_); + const auto alloc_size = width_ * sizeof(unsigned int); + LinearAllocGuard deltas_dev(LinearAllocs::hipMalloc, alloc_size); + deltas_.resize(width_); + std::generate(deltas_.begin(), deltas_.end(), + [this] { return GenerateRandomInteger(0u, static_cast(width_)); }); + HIP_CHECK(hipMemcpy(deltas_dev.ptr(), deltas_.data(), alloc_size, hipMemcpyHostToDevice)); + shfl_up<<grid_.grid_dim_, this->grid_.block_dim_>>>(arr_dev, input_dev, active_masks, + deltas_dev.ptr(), width_); + } + + void validate(const T* const arr, const T* const input) { + ArrayAllOf(arr, this->grid_.thread_count_, [this, &input](unsigned int i) -> std::optional { + const auto rank_in_block = this->grid_.thread_rank_in_block(i).value(); + const auto rank_in_warp = rank_in_block % this->warp_size_; + const auto rank_in_partition = rank_in_block % width_; + const auto mask_idx = this->warps_in_block_ * (i / this->grid_.threads_in_block_count_) + + rank_in_block / this->warp_size_; + const unsigned int delta = deltas_[rank_in_partition] % width_; + const std::bitset active_mask(this->active_masks_[mask_idx]); + + const int target = rank_in_block % width_ - delta; + if (!active_mask.test(rank_in_warp) || + (target >= 0 && !active_mask.test(rank_in_warp - delta))) { + return std::nullopt; + } + + return (target < 0 ? input[i] : input[i - delta]); + }); + }; + + private: + std::vector deltas_; + int width_; +}; + +/** + * Test Description + * ------------------------ + * - Validates the warp shuffle up behavior for all valid width sizes {2, 4, 8, 16, 32, + * 64(if supported)} for generated delta values. The threads are deactivated based on the + * passed active mask. The test is run for all overloads of shfl_up. + * Test source + * ------------------------ + * - unit/warp/warp_shfl_up.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + * - Device supports warp shuffle + */ +TEMPLATE_TEST_CASE("Unit_Warp_Shfl_Up_Positive_Basic", "", int, unsigned int, long, unsigned long, + long long, unsigned long long, float, double) { + int device; + hipDeviceProp_t device_properties; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&device_properties, device)); + + if (!device_properties.arch.hasWarpShuffle) { + HipTest::HIP_SKIP_TEST("Device doesn't support Warp Shuffle!"); + return; + } + + SECTION("Shfl Up with specified active mask and input values") { + WarpShflUp().run(false); + } + + SECTION("Shfl Down with random active mask and input values") { + WarpShflUp().run(true); + } +} \ No newline at end of file