From 5293b1a3a514f93c36f381a282d6faf43df2e73e Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Wed, 1 May 2024 18:05:24 +0100 Subject: [PATCH] SWDEV-450109 - add sanity tests for cg functions Change-Id: If54be97c07502beaaf068fdc249758931dfa5abf --- catch/unit/cooperativeGrps/CMakeLists.txt | 10 + .../unit/cooperativeGrps/binary_partition.cc | 105 +++ catch/unit/cooperativeGrps/cg_any_all.cc | 671 ++++++++++++++++++ catch/unit/cooperativeGrps/cg_ballot.cc | 162 +++++ 4 files changed, 948 insertions(+) create mode 100644 catch/unit/cooperativeGrps/binary_partition.cc create mode 100644 catch/unit/cooperativeGrps/cg_any_all.cc create mode 100644 catch/unit/cooperativeGrps/cg_ballot.cc diff --git a/catch/unit/cooperativeGrps/CMakeLists.txt b/catch/unit/cooperativeGrps/CMakeLists.txt index 9732e58f23..9ace14cccb 100644 --- a/catch/unit/cooperativeGrps/CMakeLists.txt +++ b/catch/unit/cooperativeGrps/CMakeLists.txt @@ -18,16 +18,26 @@ set(TEST_SRC coalesced_group.cc grid_group.cc coalesced_tiled_groups_metagrp.cc + binary_partition.cc + cg_ballot.cc + cg_any_all.cc ) + if(HIP_PLATFORM STREQUAL "nvidia") set_source_files_properties(hipCGMultiGridGroupType_old.cc PROPERTIES COMPILE_FLAGS "-D_CG_ABI_EXPERIMENTAL -rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80") set_source_files_properties(hipLaunchCooperativeKernelMultiDevice_old.cc PROPERTIES COMPILE_FLAGS "-D_CG_ABI_EXPERIMENTAL -rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80") set_source_files_properties(multi_grid_group.cc PROPERTIES COMPILE_FLAGS "-D_CG_ABI_EXPERIMENTAL -rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80") + set_source_files_properties(cg_ballot.cc PROPERTIES COMPILE_FLAGS "-arch=sm_70") + set_source_files_properties(cg_any_all.cc PROPERTIES COMPILE_FLAGS "-arch=sm_70") + set_source_files_properties(binary_partition.cc PROPERTIES COMPILE_FLAGS "-arch=sm_70") hip_add_exe_to_target(NAME coopGrpTest TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests LINKER_LIBS "-rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80, -gencode arch=compute_86,code=sm_86, -gencode=arch=compute_86,code=compute_86") else() + set_source_files_properties(cg_ballot.cc PROPERTIES COMPILE_FLAGS "-DHIP_ENABLE_WARP_SYNC_BUILTINS") + set_source_files_properties(cg_any_all.cc PROPERTIES COMPILE_FLAGS "-DHIP_ENABLE_WARP_SYNC_BUILTINS") + set_source_files_properties(binary_partition.cc PROPERTIES COMPILE_FLAGS "-DHIP_ENABLE_WARP_SYNC_BUILTINS") hip_add_exe_to_target(NAME coopGrpTest TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests) diff --git a/catch/unit/cooperativeGrps/binary_partition.cc b/catch/unit/cooperativeGrps/binary_partition.cc new file mode 100644 index 0000000000..0460e9f7a0 --- /dev/null +++ b/catch/unit/cooperativeGrps/binary_partition.cc @@ -0,0 +1,105 @@ +/* +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. +*/ + +#include +#include + +#include + +static __global__ void binary_part_odd_even_val(int* data, int* res) { + auto block = cooperative_groups::this_thread_block(); + auto tile = cooperative_groups::coalesced_threads(); + + int val = data[block.thread_rank()]; + auto part = cooperative_groups::binary_partition(tile, (val & 1)); + res[block.thread_rank()] = part.any(val % 2); +} + +static __global__ void binary_part_tiled_odd_even_val(int* data, int* odd_res, int* even_res) { + auto block = cooperative_groups::this_thread_block(); + auto tile16 = cooperative_groups::tiled_partition<16>(block); + + int val = data[block.thread_rank()]; + auto part = cooperative_groups::binary_partition(tile16, (val & 1)); + if (val & 1) { + *odd_res = part.all((val % 2) != 0); + } else { + *even_res = part.all((val % 2) == 0); + } +} + +TEST_CASE("Unit_cg_binary_part") { + const size_t warp_size = getWarpSize(); + + int *data, *odd_res, *even_res; + + HIP_CHECK(hipMalloc(&data, sizeof(int) * warp_size)); + HIP_CHECK(hipMalloc(&odd_res, sizeof(int) * warp_size)); + HIP_CHECK(hipMalloc(&even_res, sizeof(int) * warp_size)); + + HIP_CHECK(hipMemset(data, 0, sizeof(int) * warp_size)); + + SECTION("odd_even part") { + std::vector input; + std::vector output(warp_size, -1); + + auto& res = odd_res; + input.reserve(warp_size); + for (size_t i = 0; i < warp_size; i++) { + if (i < 16) + input.push_back(10); + else + input.push_back(11); + } + + HIP_CHECK(hipMemcpy(data, input.data(), sizeof(int) * input.size(), hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(res, output.data(), sizeof(int) * output.size(), hipMemcpyHostToDevice)); + + binary_part_odd_even_val<<<1, warp_size>>>(data, res); + HIP_CHECK(hipMemcpy(output.data(), res, sizeof(int) * output.size(), hipMemcpyDeviceToHost)); + for (size_t i = 0; i < warp_size; i++) { + if (i < 16) { + INFO("Output <16, index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 0); + } else { + INFO("Output >=16, index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 1); + } + } + } + + SECTION("tiled odd/even part") { + std::vector input; + input.reserve(warp_size); + for (size_t i = 0; i < warp_size; i++) { + input.push_back(i + 1); + } + + HIP_CHECK(hipMemcpy(data, input.data(), sizeof(int) * input.size(), hipMemcpyHostToDevice)); + binary_part_tiled_odd_even_val<<<1, warp_size>>>(data, odd_res, even_res); + int odd_output, even_output; + HIP_CHECK(hipMemcpy(&odd_output, odd_res, sizeof(int), hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(&even_output, even_res, sizeof(int), hipMemcpyDeviceToHost)); + REQUIRE(odd_output == 1); + REQUIRE(even_output == 1); + } + + HIP_CHECK(hipFree(odd_res)); + HIP_CHECK(hipFree(even_res)); +} \ No newline at end of file diff --git a/catch/unit/cooperativeGrps/cg_any_all.cc b/catch/unit/cooperativeGrps/cg_any_all.cc new file mode 100644 index 0000000000..b9b2a092d7 --- /dev/null +++ b/catch/unit/cooperativeGrps/cg_any_all.cc @@ -0,0 +1,671 @@ +/* +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. +*/ + +#include +#include + +#include + +// Wavefront sized launch +__global__ void coop_any_coal(unsigned int* data, unsigned int val, unsigned int* res) { + auto cg = cooperative_groups::coalesced_threads(); + unsigned int tmp = data[threadIdx.x]; + res[threadIdx.x] = cg.any(tmp == val); +} + +__global__ void coop_any_coal_odd_even(unsigned int* data, unsigned int even_val, + unsigned int odd_val, unsigned int* res) { + if ((threadIdx.x % 2) == 0) { + auto cg = cooperative_groups::coalesced_threads(); + unsigned int tmp = data[threadIdx.x]; + res[threadIdx.x] = cg.any(tmp == even_val); + } else { + auto cg = cooperative_groups::coalesced_threads(); + unsigned int tmp = data[threadIdx.x]; + res[threadIdx.x] = cg.any(tmp == odd_val); + } +} + +TEST_CASE("Unit_coopgroups_any") { + const size_t warp_size = getWarpSize(); + + unsigned int *data, *res; + + HIP_CHECK(hipMalloc(&data, sizeof(int) * warp_size)); + HIP_CHECK(hipMalloc(&res, sizeof(int) * warp_size)); + + HIP_CHECK(hipMemset(data, 0, sizeof(int) * warp_size)); + + SECTION("all set") { + constexpr unsigned int any_val = 10; + std::vector input(warp_size, any_val); + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_any_coal<<<1, warp_size>>>(data, any_val, res); + std::vector output(warp_size, 0); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 1); + } + } + + SECTION("all set - odd/even coal") { + constexpr unsigned int odd_val = 10, even_val = 20; + std::vector input; + input.reserve(warp_size); + for (size_t i = 0; i < warp_size; i++) { + if ((i % 2) == 0) { + input.push_back(even_val); + } else { + input.push_back(odd_val); + } + } + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_any_coal_odd_even<<<1, warp_size>>>(data, even_val, odd_val, res); + std::vector output(warp_size, 0); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 1); + } + } + + SECTION("None set") { + constexpr unsigned int any_val = 20; + std::vector input(warp_size, any_val); + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_any_coal<<<1, warp_size>>>(data, 10, res); + std::vector output(warp_size, 0); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 0); + } + } + + SECTION("None set - odd/even coal") { + constexpr unsigned int odd_val = 10, even_val = 20; + std::vector input; + input.reserve(warp_size); + for (size_t i = 0; i < warp_size; i++) { + if ((i % 2) == 0) { + input.push_back(even_val); + } else { + input.push_back(odd_val); + } + } + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_any_coal_odd_even<<<1, warp_size>>>(data, odd_val, even_val, res); + std::vector output(warp_size, 0); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 0); + } + } + + SECTION("Alternate set") { + constexpr unsigned int any_val = 20; + std::vector input; + for (size_t i = 0; i < warp_size; i++) { + if ((i % 2) == 0) { + input.push_back(any_val + 1); + } else { + input.push_back(any_val); + } + } + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_any_coal<<<1, warp_size>>>(data, any_val, res); + std::vector output(warp_size, 0); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 1); + } + } + + SECTION("first half set") { + constexpr unsigned int any_val = 20; + std::vector input; + for (size_t i = 0; i < warp_size; i++) { + if (i < (warp_size / 2)) { + input.push_back(any_val); + } else { + input.push_back(any_val + 1); + } + } + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_any_coal<<<1, warp_size>>>(data, any_val, res); + std::vector output(warp_size, 0); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 1); + } + } + + SECTION("last half set") { + constexpr unsigned int any_val = 20; + std::vector input; + for (size_t i = 0; i < warp_size; i++) { + if (i < (warp_size / 2)) { + input.push_back(any_val + 1); + } else { + input.push_back(any_val); + } + } + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_any_coal<<<1, warp_size>>>(data, any_val, res); + std::vector output(warp_size, 0); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 1); + } + } + + SECTION("First set") { + constexpr unsigned int any_val = 20; + std::vector input(warp_size, any_val + 1); + input[0] = any_val; + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_any_coal<<<1, warp_size>>>(data, any_val, res); + std::vector output(warp_size, 0); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 1); + } + } + + SECTION("First set - odd/even coal") { + constexpr unsigned int odd_val = 10, even_val = 20; + std::vector input(warp_size, 0); + input[0] = even_val; + input[1] = odd_val; + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_any_coal_odd_even<<<1, warp_size>>>(data, even_val, odd_val, res); + std::vector output(warp_size, 0); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 1); + } + } + + SECTION("Last set") { + constexpr unsigned int any_val = 20; + std::vector input(warp_size, any_val + 1); + input[warp_size - 1] = any_val; + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_any_coal<<<1, warp_size>>>(data, any_val, res); + std::vector output(warp_size, 0); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 1); + } + } + + SECTION("Last set - odd/even coal") { + constexpr unsigned int odd_val = 10, even_val = 20; + std::vector input(warp_size, 0); + input[warp_size - 2] = even_val; + input[warp_size - 1] = odd_val; + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_any_coal_odd_even<<<1, warp_size>>>(data, even_val, odd_val, res); + std::vector output(warp_size, 0); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 1); + } + } + + HIP_CHECK(hipFree(data)); + HIP_CHECK(hipFree(res)); +} + +__global__ void coop_all_coal(unsigned int* data, unsigned int val, unsigned int* res) { + auto cg = cooperative_groups::coalesced_threads(); + unsigned int tmp = data[threadIdx.x]; + res[threadIdx.x] = cg.all(tmp == val); +} + +__global__ void coop_all_coal_odd_even(unsigned int* data, unsigned int even_val, + unsigned int odd_val, unsigned int* res) { + if ((threadIdx.x % 2) == 0) { + auto cg = cooperative_groups::coalesced_threads(); + unsigned int tmp = data[threadIdx.x]; + res[threadIdx.x] = cg.all(tmp == even_val); + } else { + auto cg = cooperative_groups::coalesced_threads(); + unsigned int tmp = data[threadIdx.x]; + res[threadIdx.x] = cg.all(tmp == odd_val); + } +} + +TEST_CASE("Unit_coopgroups_coal_all") { + const size_t warp_size = getWarpSize(); + + unsigned int *data, *res; + + HIP_CHECK(hipMalloc(&data, sizeof(unsigned int) * warp_size)); + HIP_CHECK(hipMalloc(&res, sizeof(unsigned int) * warp_size)); + + HIP_CHECK(hipMemset(data, 0, sizeof(unsigned int) * warp_size)); + + SECTION("All set - all in sync") { + constexpr unsigned int any_val = 10; + std::vector input(warp_size, any_val); + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_all_coal<<<1, warp_size>>>(data, any_val, res); + std::vector output(warp_size, 0); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 1); + } + } + + SECTION("All set - odd/even in sync") { + constexpr unsigned int even_val = 10, odd_val = 20; + std::vector input; + input.reserve(warp_size); + for (size_t i = 0; i < warp_size; i++) { + if ((i % 2) == 0) { + input.push_back(even_val); + } else { + input.push_back(odd_val); + } + } + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_all_coal_odd_even<<<1, warp_size>>>(data, even_val, odd_val, res); + std::vector output(warp_size, 0); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 1); + } + } + + SECTION("None set") { + constexpr unsigned int any_val = 20; + std::vector input(warp_size, any_val); + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_all_coal<<<1, warp_size>>>(data, 10, res); + std::vector output(warp_size, 0); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 0); + } + } + + SECTION("None set - odd/even sync") { + constexpr unsigned int even_val = 10, odd_val = 20; + std::vector input; + input.reserve(warp_size); + for (size_t i = 0; i < warp_size; i++) { + if ((i % 2) == 0) { + input.push_back(odd_val); + } else { + input.push_back(even_val); + } + } + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_all_coal_odd_even<<<1, warp_size>>>(data, even_val, odd_val, res); + std::vector output(warp_size, 1); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 0); + } + } + + SECTION("Alternate set") { + constexpr unsigned int any_val = 20; + std::vector input; + for (size_t i = 0; i < warp_size; i++) { + if ((i % 2) == 0) { + input.push_back(any_val + 1); + } else { + input.push_back(any_val); + } + } + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_all_coal<<<1, warp_size>>>(data, any_val, res); + std::vector output(warp_size, 0); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 0); + } + } + + SECTION("first half set") { + constexpr unsigned int any_val = 20; + std::vector input; + for (size_t i = 0; i < warp_size; i++) { + if (i < (warp_size / 2)) { + input.push_back(any_val); + } else { + input.push_back(any_val + 1); + } + } + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_all_coal<<<1, warp_size>>>(data, any_val, res); + std::vector output(warp_size, 0); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 0); + } + } + + SECTION("last half set") { + constexpr unsigned int any_val = 20; + std::vector input; + for (size_t i = 0; i < warp_size; i++) { + if (i < (warp_size / 2)) { + input.push_back(any_val + 1); + } else { + input.push_back(any_val); + } + } + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_all_coal<<<1, warp_size>>>(data, any_val, res); + std::vector output(warp_size, 0); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 0); + } + } + + SECTION("First set") { + constexpr unsigned int any_val = 20; + std::vector input(warp_size, any_val + 1); + input[0] = any_val; + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_all_coal<<<1, warp_size>>>(data, any_val, res); + std::vector output(warp_size, 0); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 0); + } + } + + SECTION("Last set") { + constexpr unsigned int any_val = 20; + std::vector input(warp_size, any_val + 1); + input[warp_size - 1] = any_val; + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_all_coal<<<1, warp_size>>>(data, any_val, res); + std::vector output(warp_size, 0); + HIP_CHECK( + hipMemcpy(output.data(), res, sizeof(unsigned int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == 0); + } + } + + HIP_CHECK(hipFree(data)); + HIP_CHECK(hipFree(res)); +} + +__global__ void coop_match_any_coal(unsigned int* data, unsigned long long* res) { + auto cg = cooperative_groups::coalesced_threads(); + unsigned int tmp = data[threadIdx.x]; + res[threadIdx.x] = cg.match_any(tmp); +} + +__global__ void coop_match_any_coal_odd_even(unsigned int* data, unsigned long long* res) { + unsigned int tmp = data[threadIdx.x]; + if ((threadIdx.x % 2) == 0) { + auto cg = cooperative_groups::coalesced_threads(); + res[threadIdx.x] = cg.match_any(tmp); + } else { + auto cg = cooperative_groups::coalesced_threads(); + res[threadIdx.x] = cg.match_any(tmp); + } +} + +TEST_CASE("Unit_coopgroups_match_any_coal") { + const size_t warp_size = getWarpSize(); + + unsigned int* data; + unsigned long long* res; + + HIP_CHECK(hipMalloc(&data, sizeof(unsigned int) * warp_size)); + HIP_CHECK(hipMalloc(&res, sizeof(unsigned long long) * warp_size)); + + HIP_CHECK(hipMemset(data, 0, sizeof(unsigned int) * warp_size)); + HIP_CHECK(hipMemset(res, 0, sizeof(unsigned long long) * warp_size)); + + SECTION("all set") { + constexpr unsigned int any_val = 10; + std::vector input(warp_size, any_val); + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_match_any_coal<<<1, warp_size>>>(data, res); + std::vector output(warp_size, 0); + HIP_CHECK(hipMemcpy(output.data(), res, sizeof(unsigned long long) * warp_size, + hipMemcpyDeviceToHost)); + auto all_set_mask = warp_size == 32 ? 0xFFFF'FFFFull : 0xFFFF'FFFF'FFFF'FFFFull; + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == all_set_mask); + } + } + + SECTION("all set - odd/even") { + constexpr unsigned int even_val = 10, odd_val = 20; + std::vector input; + for (size_t i = 0; i < warp_size; i += 2) { + input.push_back(even_val); + input.push_back(odd_val); + } + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_match_any_coal_odd_even<<<1, warp_size>>>(data, res); + std::vector output(warp_size, 0); + HIP_CHECK(hipMemcpy(output.data(), res, sizeof(unsigned long long) * warp_size, + hipMemcpyDeviceToHost)); + auto all_set_mask = warp_size == 32 ? 0xFFFFull : 0xFFFF'FFFFull; + for (size_t i = 0; i < output.size(); i++) { + INFO("Checking for index: " << i << " output: " << output[i]); + REQUIRE(output[i] == all_set_mask); + } + } + + HIP_CHECK(hipFree(data)); + HIP_CHECK(hipFree(res)); +} + +__global__ void coop_match_all_coal(unsigned int* data, unsigned long long* res, int* pred_res) { + auto cg = cooperative_groups::coalesced_threads(); + unsigned int tmp = data[threadIdx.x]; + res[threadIdx.x] = cg.match_all(tmp, pred_res[threadIdx.x]); +} + +__global__ void coop_match_all_coal_odd_even(unsigned int* data, unsigned long long* res, + int* pred_res) { + unsigned int tmp = data[threadIdx.x]; + if ((threadIdx.x % 2) == 0) { + auto cg = cooperative_groups::coalesced_threads(); + res[threadIdx.x] = cg.match_all(tmp, pred_res[threadIdx.x]); + } else { + auto cg = cooperative_groups::coalesced_threads(); + res[threadIdx.x] = cg.match_all(tmp, pred_res[threadIdx.x]); + } +} + +TEST_CASE("Unit_coopgroups_match_all_coal") { + const size_t warp_size = getWarpSize(); + + unsigned int* data; + unsigned long long* res; + int* pred_res; + + HIP_CHECK(hipMalloc(&data, sizeof(unsigned int) * warp_size)); + HIP_CHECK(hipMalloc(&res, sizeof(unsigned long long) * warp_size)); + HIP_CHECK(hipMalloc(&pred_res, sizeof(int) * warp_size)); + + HIP_CHECK(hipMemset(data, 0, sizeof(unsigned int) * warp_size)); + HIP_CHECK(hipMemset(res, 0, sizeof(unsigned long long) * warp_size)); + + SECTION("all set") { + constexpr unsigned int any_val = 10; + std::vector input(warp_size, any_val); + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_match_all_coal<<<1, warp_size>>>(data, res, pred_res); + std::vector mask_output(warp_size, 0); + std::vector pred_output(warp_size, 0); + HIP_CHECK(hipMemcpy(mask_output.data(), res, sizeof(unsigned long long) * warp_size, + hipMemcpyDeviceToHost)); + HIP_CHECK( + hipMemcpy(pred_output.data(), pred_res, sizeof(int) * warp_size, hipMemcpyDeviceToHost)); + auto all_set_mask = warp_size == 32 ? 0xFFFF'FFFFull : 0xFFFF'FFFF'FFFF'FFFFull; + for (size_t i = 0; i < mask_output.size(); i++) { + INFO("Checking for index: " << i << " output: " << mask_output[i] + << " pred: " << pred_output[i]); + REQUIRE(mask_output[i] == all_set_mask); + REQUIRE(pred_output[i] == 1); + } + } + + SECTION("one mismatch") { + constexpr unsigned int any_val = 10; + std::vector input(warp_size, any_val); + input[warp_size - 1]--; + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_match_all_coal<<<1, warp_size>>>(data, res, pred_res); + std::vector mask_output(warp_size, 0); + std::vector pred_output(warp_size, 0); + HIP_CHECK(hipMemcpy(mask_output.data(), res, sizeof(unsigned long long) * warp_size, + hipMemcpyDeviceToHost)); + HIP_CHECK( + hipMemcpy(pred_output.data(), pred_res, sizeof(int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < mask_output.size(); i++) { + INFO("Checking for index: " << i << " output: " << mask_output[i] + << " pred: " << pred_output[i]); + REQUIRE(mask_output[i] == 0); + REQUIRE(pred_output[i] == 0); + } + } + + SECTION("all set - odd/even") { + constexpr unsigned int even_val = 10, odd_val = 20; + std::vector input; + for (size_t i = 0; i < warp_size; i += 2) { + input.push_back(even_val); + input.push_back(odd_val); + } + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_match_all_coal_odd_even<<<1, warp_size>>>(data, res, pred_res); + std::vector mask_output(warp_size, 0); + std::vector pred_output(warp_size, 0); + HIP_CHECK(hipMemcpy(mask_output.data(), res, sizeof(unsigned long long) * warp_size, + hipMemcpyDeviceToHost)); + HIP_CHECK( + hipMemcpy(pred_output.data(), pred_res, sizeof(int) * warp_size, hipMemcpyDeviceToHost)); + auto all_set_mask = warp_size == 32 ? 0xFFFFull : 0xFFFF'FFFFull; + for (size_t i = 0; i < mask_output.size(); i++) { + INFO("Checking for index: " << i << " output: " << mask_output[i] + << " pred: " << pred_output[i]); + REQUIRE(mask_output[i] == all_set_mask); + REQUIRE(pred_output[i] == 1); + } + } + + SECTION("one mismatch - odd/even") { + constexpr unsigned int even_val = 10, odd_val = 20; + std::vector input; + for (size_t i = 0; i < warp_size; i += 2) { + input.push_back(even_val); + input.push_back(odd_val); + } + input[warp_size - 1]--; + input[warp_size - 2]--; + HIP_CHECK( + hipMemcpy(data, input.data(), sizeof(unsigned int) * warp_size, hipMemcpyHostToDevice)); + coop_match_all_coal_odd_even<<<1, warp_size>>>(data, res, pred_res); + std::vector mask_output(warp_size, 0); + std::vector pred_output(warp_size, 0); + HIP_CHECK(hipMemcpy(mask_output.data(), res, sizeof(unsigned long long) * warp_size, + hipMemcpyDeviceToHost)); + HIP_CHECK( + hipMemcpy(pred_output.data(), pred_res, sizeof(int) * warp_size, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < mask_output.size(); i++) { + INFO("Checking for index: " << i << " output: " << mask_output[i] + << " pred: " << pred_output[i]); + REQUIRE(mask_output[i] == 0); + REQUIRE(pred_output[i] == 0); + } + } + + HIP_CHECK(hipFree(data)); + HIP_CHECK(hipFree(res)); + HIP_CHECK(hipFree(pred_res)); +} diff --git a/catch/unit/cooperativeGrps/cg_ballot.cc b/catch/unit/cooperativeGrps/cg_ballot.cc new file mode 100644 index 0000000000..4cd0e62e52 --- /dev/null +++ b/catch/unit/cooperativeGrps/cg_ballot.cc @@ -0,0 +1,162 @@ +/* +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. +*/ + +#include +#include + +#include + +__global__ void coop_ballot_coal_alternate(int* data, unsigned long long* d_even_mask, + unsigned long long* d_odd_mask) { + int tid = threadIdx.x; + auto val = data[tid]; + + if ((tid % 2) == 0) { + auto cg = cooperative_groups::coalesced_threads(); + auto cg_ballot_res = cg.ballot(val == 10); + if (tid == 0) { + *d_even_mask = cg_ballot_res; + } + } else { + auto cg = cooperative_groups::coalesced_threads(); + auto cg_ballot_res = cg.ballot(val == 10); + if (tid == 1) { + *d_odd_mask = cg_ballot_res; + } + } +} + +TEST_CASE("Unit_coopgroups_ballot") { + const size_t warp_size = getWarpSize(); + std::vector input; + input.reserve(warp_size); + // Push 10 - 10 - 20 - 20 in vector + for (size_t i = 0; i < warp_size; i += 4) { + input.push_back(10); + input.push_back(10); + input.push_back(20); + input.push_back(20); + } + + int* data; + HIP_CHECK(hipMalloc(&data, sizeof(int) * input.size())); + HIP_CHECK(hipMemcpy(data, input.data(), sizeof(int) * input.size(), hipMemcpyHostToDevice)); + + unsigned long long *d_even_mask, *d_odd_mask, even_mask = 0, odd_mask = 0l; + HIP_CHECK(hipMalloc(&d_even_mask, sizeof(unsigned long long))); + HIP_CHECK(hipMalloc(&d_odd_mask, sizeof(unsigned long long))); + + coop_ballot_coal_alternate<<<1, warp_size>>>(data, d_even_mask, d_odd_mask); + HIP_CHECK(hipMemcpy(&even_mask, d_even_mask, sizeof(unsigned long long), hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(&odd_mask, d_odd_mask, sizeof(unsigned long long), hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + + HIP_CHECK(hipFree(data)); + HIP_CHECK(hipFree(d_even_mask)); + HIP_CHECK(hipFree(d_odd_mask)); + + if (warp_size == 32) { + REQUIRE(even_mask == 0x5555); + REQUIRE(odd_mask == 0x5555); + } else { + REQUIRE(even_mask == 0x55555555); + REQUIRE(odd_mask == 0x55555555); + } +} + +static __global__ void coop_ballot_binary_part(int* data, unsigned long long* res) { + int tid = threadIdx.x; + auto val = data[tid]; + auto part = + cooperative_groups::binary_partition(cooperative_groups::coalesced_threads(), (tid % 5) == 0); + // Only tid % 5 lane will vote yes + res[tid] = part.ballot((val % 5) == 0); +} + +static __global__ void coop_ballot_binary_part_3_5_7(int* data, unsigned long long* res) { + int tid = threadIdx.x; + auto val = data[tid]; + auto part = cooperative_groups::binary_partition(cooperative_groups::coalesced_threads(), + (tid % 3) == 0 || (tid % 5) == 0); + res[tid] = part.ballot((val % 7) == 0); +} + +TEST_CASE("Unit_binary_part_ballot") { + const size_t warp_size = getWarpSize(); + std::vector input; + input.reserve(warp_size); + for (size_t i = 0; i < warp_size; i++) { + input.push_back(i); + } + + int* data; + HIP_CHECK(hipMalloc(&data, sizeof(int) * input.size())); + HIP_CHECK(hipMemcpy(data, input.data(), sizeof(int) * input.size(), hipMemcpyHostToDevice)); + + unsigned long long* d_res; + HIP_CHECK(hipMalloc(&d_res, sizeof(unsigned long long) * warp_size)); + + SECTION("Basic partition of 5th lane") { + std::vector res(warp_size, 1); + coop_ballot_binary_part<<<1, warp_size>>>(data, d_res); + HIP_CHECK(hipMemcpy(res.data(), d_res, sizeof(unsigned long long) * warp_size, + hipMemcpyDeviceToHost)); + unsigned long long expected_result = warp_size == 32 ? 0x7F : 0x1FFF; + for (size_t i = 0; i < res.size(); i++) { + INFO("Index: " << i << " result: " << res[i]); + REQUIRE((i % 5 == 0 ? res[i] == expected_result : res[i] == 0)); + } + } + + SECTION("Lanes index multiple of 2 and 3 voting on value % 5") { + std::vector res(warp_size, 1); + + std::vector expected; + expected.reserve(warp_size); + /* This will be fun to explain + * So lane which are multiple of 3 or 5 will call ballot on values which are multiple of 7 + * lane: 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 + * active:x x x x x x x x x x x x x x x + * res1: 1 0 0 0 0 0 0 0 0 0 1 0 0 0 0 + * res2: 0 0 0 1 0 0 0 1 0 0 0 0 0 0 1 0 0 + * Active lane here is how the warp will be partitioned into i.e. tid % 3 || tid % 5, so active + * in one partition and ~active in other. res1 is what active part will report as result and + * res2 will be reported by ~active partition. + */ + const unsigned long long res1 = warp_size == 32 ? 1025 : 537986049; + const unsigned long long res2 = warp_size == 32 ? 16520 : 603996296; + for (size_t i = 0; i < warp_size; i++) { + if ((i % 3) == 0 || (i % 5) == 0) { + expected.push_back(res1); + } else { + expected.push_back(res2); + } + } + coop_ballot_binary_part_3_5_7<<<1, warp_size>>>(data, d_res); + HIP_CHECK(hipMemcpy(res.data(), d_res, sizeof(unsigned long long) * warp_size, + hipMemcpyDeviceToHost)); + for (size_t i = 0; i < res.size(); i++) { + INFO("index: " << i << "expected: " << expected[i] << " got: " << res[i]); + CHECK(expected[i] == res[i]); + } + } + + HIP_CHECK(hipFree(d_res)); + HIP_CHECK(hipFree(data)); +} \ No newline at end of file