SWDEV-450109 - add sanity tests for cg functions
Change-Id: If54be97c07502beaaf068fdc249758931dfa5abf
This commit is contained in:
Jatin Chaudhary
committed by
Rakesh Roy
Rakesh Roy
parent
e835edcbab
commit
5293b1a3a5
@@ -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)
|
||||
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
|
||||
#include <vector>
|
||||
|
||||
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<int> input;
|
||||
std::vector<int> 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<int> 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));
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
|
||||
#include <vector>
|
||||
|
||||
// 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned long long> 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<unsigned int> 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<unsigned long long> 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<unsigned int> 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<unsigned long long> mask_output(warp_size, 0);
|
||||
std::vector<int> 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<unsigned int> 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<unsigned long long> mask_output(warp_size, 0);
|
||||
std::vector<int> 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<unsigned int> 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<unsigned long long> mask_output(warp_size, 0);
|
||||
std::vector<int> 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<unsigned int> 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<unsigned long long> mask_output(warp_size, 0);
|
||||
std::vector<int> 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));
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
|
||||
#include <vector>
|
||||
|
||||
__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<int> 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<int> 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<unsigned long long> 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<unsigned long long> res(warp_size, 1);
|
||||
|
||||
std::vector<unsigned long long> 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));
|
||||
}
|
||||
Reference in New Issue
Block a user