diff --git a/catch/unit/warp/CMakeLists.txt b/catch/unit/warp/CMakeLists.txt index 7b205f915b..292da243c4 100644 --- a/catch/unit/warp/CMakeLists.txt +++ b/catch/unit/warp/CMakeLists.txt @@ -5,6 +5,7 @@ set(TEST_SRC warp_shfl_up.cc warp_shfl_down.cc warp_ballot.cc + warp_any.cc ) hip_add_exe_to_target(NAME WarpTest diff --git a/catch/unit/warp/warp_any.cc b/catch/unit/warp/warp_any.cc new file mode 100644 index 0000000000..91dadbb12c --- /dev/null +++ b/catch/unit/warp/warp_any.cc @@ -0,0 +1,119 @@ +/* +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_vote_common.hh" + +#include + +/** + * @addtogroup any any + * @{ + * @ingroup DeviceLanguageTest + * `int __any(int predicate)` - + * Contains unit test for warp vote any function + */ + +namespace cg = cooperative_groups; + + +__global__ void kernel_any(int* const out, const uint64_t* const active_masks, uint64_t predicate) { + if (deactivate_thread(active_masks)) { + return; + } + + const auto grid = cg::this_grid(); + const auto warp = cg::tiled_partition(cg::this_thread_block(), warpSize); + + out[grid.thread_rank()] = __any((predicate & (static_cast(1) << warp.thread_rank()))); +} + +class WarpAny : public WarpVoteTest { + public: + void launch_kernel(int* const arr_dev, const uint64_t* const active_masks) { + auto test_case = GENERATE(range(0, 5)); + predicate_mask_ = get_predicate_mask(test_case, this->warp_size_); + INFO("Predicate mask: " << predicate_mask_); + kernel_any<<grid_.grid_dim_, this->grid_.block_dim_>>>(arr_dev, active_masks, + predicate_mask_); + } + + void validate(const int* const arr) { + ArrayAllOf(arr, this->grid_.thread_count_, [this](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 warp_idx = this->warps_in_block_ * (i / this->grid_.threads_in_block_count_) + + rank_in_block / this->warp_size_; + const auto block_rank = warp_idx / this->warps_in_block_; + const std::bitset active_mask(this->active_masks_[warp_idx]); + + auto partition_size = this->warp_size_; + // If the number of threads in a block is not a multiple of warp size, the + // last warp will have inactive threads and partition size must be recalculated + if (warp_idx == this->warps_in_block_ * (block_rank + 1) - 1) { + partition_size = + this->grid_.threads_in_block_count_ - (this->warps_in_block_ - 1) * this->warp_size_; + } + + if (!active_mask.test(rank_in_warp)) + return std::nullopt; + else { + // Active predicate mask must be calculated as partition can be smaller than warp_size + auto active_predicate = get_active_predicate(predicate_mask_, partition_size); + return ((active_predicate & this->active_masks_[warp_idx]) != 0); + } + }); + } + + private: + uint64_t predicate_mask_; +}; + +/** + * Test Description + * ------------------------ + * - Validates the warp vote any function behavior. Threads are deactivated based on the passed + * active mask. The predicate for each thread is determined according to the generated predicate + * mask. + * Test source + * ------------------------ + * - unit/warp/warp_any.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + * - Device supports warp vote + */ +TEST_CASE("Unit_Warp_Vote_Any_Positive_Basic") { + int device; + hipDeviceProp_t device_properties; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&device_properties, device)); + + if (!device_properties.arch.hasWarpVote) { + HipTest::HIP_SKIP_TEST("Device doesn't support Warp Vote!"); + return; + } + + SECTION("Warp Vote Any with specified active mask") { + WarpAny().run(false); + } + + SECTION("Warp Vote Any with random active mask") { + WarpAny().run(true); + } +} \ No newline at end of file