AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[SWDEV-438556][SWDEV-450826][SWDEV-449018] - Fix cooperative group catch tests

Browse Source 
Change-Id: I553afe9c684b9ae60ce84231d7a5fbd6532a20e6


[ROCm/hip-tests commit: 6df445fef6]
This commit is contained in:
Rahul Manocha
2024-03-08 23:37:45 +00:00
parent b46f7c8f6a
commit 92b0b9fbc2
4 changed files with 182 additions and 179 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/hipTestMain/config/config_amd_linux
+20 -56
View File
@@ -161,8 +161,6 @@
"Unit_Warp_Shfl_XOR_Positive_Basic - double",
"Unit_Warp_Shfl_XOR_Positive_Basic - __half",
"Unit_Warp_Shfl_XOR_Positive_Basic - __half2",
"Unit_Coalesced_Group_Sync_Positive_Basic - uint16_t",
"Unit_Coalesced_Group_Sync_Positive_Basic - uint32_t",
"=== SWDEV-434878: Below tests failed in stress test on 24/11/23 ===",
"Unit_hipGraphUpload_Negative_Parameters",
"Unit_hipModuleOccupancyMaxPotentialBlockSize_Negative_Parameters",
@@ -1313,18 +1311,6 @@
"Unit_hipFreeMipmappedArrayMultiTArray - char",
"Unit_hipFreeMipmappedArrayMultiTArray - int",
"Unit_hipIpcGetMemHandle_Positive_Unique_Handles_Reused_Memory",
"Unit_Coalesced_Group_Getters_Via_Base_Type_Positive_Basic",
"Unit_Coalesced_Group_Shfl_Positive_Basic - int",
"Unit_Coalesced_Group_Shfl_Positive_Basic - unsigned int",
"Unit_Coalesced_Group_Shfl_Positive_Basic - long",
"Unit_Coalesced_Group_Shfl_Positive_Basic - unsigned long",
"Unit_Coalesced_Group_Shfl_Positive_Basic - long long",
"Unit_Coalesced_Group_Shfl_Positive_Basic - unsigned long long",
"Unit_Coalesced_Group_Shfl_Positive_Basic - float",
"Unit_Coalesced_Group_Shfl_Positive_Basic - double",
"Unit_Coalesced_Group_Sync_Positive_Basic - uint8_t",
"Unit_Coalesced_Group_Sync_Positive_Basic - uint16_t",
"Unit_Coalesced_Group_Sync_Positive_Basic - uint32_t",
"SWDEV-445928: These tests fail in PSDB stress test on 09/02/2024",
"Unit_hipGraphAddNodeTypeMemset_Positive_Basic - uint8_t",
"Unit_hipGraphAddNodeTypeMemset_Positive_Basic - uint16_t",
@@ -1367,27 +1353,6 @@
"Unit_hipClock_Positive_Basic",
"=== Below tests failed in integrity test on 08/12/23 ===",
"=== SWDEV-438556:Below tests failed in stress test on 15/12/23 ===",
"Unit_Coalesced_Group_Getters_Positive_Basic",
"Unit_Coalesced_Group_Getters_Via_Non_Member_Functions_Positive_Basic",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - int",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - unsigned int",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - long",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - unsigned long",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - long long",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - unsigned long long",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - float",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - double",
"Unit_Coalesced_Group_Shfl_Down_Positive_Basic - int",
"Unit_Coalesced_Group_Shfl_Down_Positive_Basic - unsigned int",
"Unit_Coalesced_Group_Shfl_Down_Positive_Basic - long",
"Unit_Coalesced_Group_Shfl_Down_Positive_Basic - unsigned long",
"Unit_Coalesced_Group_Shfl_Down_Positive_Basic - long long",
"Unit_Coalesced_Group_Shfl_Down_Positive_Basic - unsigned long long",
"Unit_Coalesced_Group_Shfl_Down_Positive_Basic - float",
"Unit_Coalesced_Group_Shfl_Down_Positive_Basic - double",
"Unit_Coalesced_Group_Sync_Positive_Basic - uint8_t",
"Unit_Coalesced_Group_Sync_Positive_Basic - uint16_t",
"Unit_Coalesced_Group_Sync_Positive_Basic - uint32_t",
"=== SWDEV-443630 - Below tests failed in stress test on 19/01/23 ===",
"Unit_hipGetSetDevice_MultiThreaded",
#endif
@@ -1453,27 +1418,6 @@
"Unit_cache_coherency_cpu_gpu",
"Unit_cache_coherency_gpu_gpu",
"=== SWDEV-438556:Below tests failed in stress test on 15/12/23 ===",
"Unit_Coalesced_Group_Getters_Positive_Basic",
"Unit_Coalesced_Group_Getters_Via_Non_Member_Functions_Positive_Basic",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - int",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - unsigned int",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - long",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - unsigned long",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - long long",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - unsigned long long",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - float",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - double",
"Unit_Coalesced_Group_Shfl_Down_Positive_Basic - int",
"Unit_Coalesced_Group_Shfl_Down_Positive_Basic - unsigned int",
"Unit_Coalesced_Group_Shfl_Down_Positive_Basic - long",
"Unit_Coalesced_Group_Shfl_Down_Positive_Basic - unsigned long",
"Unit_Coalesced_Group_Shfl_Down_Positive_Basic - long long",
"Unit_Coalesced_Group_Shfl_Down_Positive_Basic - unsigned long long",
"Unit_Coalesced_Group_Shfl_Down_Positive_Basic - float",
"Unit_Coalesced_Group_Shfl_Down_Positive_Basic - double",
"Unit_Coalesced_Group_Sync_Positive_Basic - uint8_t",
"Unit_Coalesced_Group_Sync_Positive_Basic - uint16_t",
"Unit_Coalesced_Group_Sync_Positive_Basic - uint32_t",
"=== SWDEV-439298: Below test failing in CQE staging ===",
"Unit_hipCGMultiGridGroupType_Barrier",
"=== SWDEV-443630 : Below test failed in stress test on 19/01/24 ===",
@@ -1515,6 +1459,26 @@
"Unit_safeAtomicMin_Positive_SameAddress - float",
"=== SWDEV-454220 : Below test hanged in stress test on 22/03/24 ===",
"Unit_hipExtLaunchKernel_Positive_Basic",
"=== SWDEV-454220 Below test fail in stress test 03/29/24 ===",
"Unit_Coalesced_Group_Shfl_Positive_Basic - int",
"Unit_Coalesced_Group_Shfl_Positive_Basic - unsigned int",
"Unit_Coalesced_Group_Shfl_Positive_Basic - long",
"Unit_Coalesced_Group_Shfl_Positive_Basic - unsigned long",
"Unit_Coalesced_Group_Shfl_Positive_Basic - long long",
"Unit_Coalesced_Group_Shfl_Positive_Basic - unsigned long long",
"Unit_Coalesced_Group_Shfl_Positive_Basic - float",
"Unit_Coalesced_Group_Shfl_Positive_Basic - double",
"Unit_Coalesced_Group_Sync_Positive_Basic - uint8_t",
"Unit_Coalesced_Group_Sync_Positive_Basic - uint16_t",
"Unit_Coalesced_Group_Sync_Positive_Basic - uint32_t",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - int",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - unsigned int",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - long",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - unsigned long",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - long long",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - unsigned long long",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - float",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - double",
"=== Below tests fail in stress test 03/29/24 ===",
"Unit_Warp_Ballot_Positive_Basic",
"Unit_Warp_Vote_Any_Positive_Basic",
projects/hip-tests/catch/unit/cooperativeGrps/coalesced_group.cc
+161 -95
View File
@@ -31,43 +31,61 @@ THE SOFTWARE.
namespace cg = cooperative_groups;
template <unsigned int warp_size, typename BaseType = cg::coalesced_group>
template <typename BaseType = cg::coalesced_group>
static __global__ void coalesced_group_size_getter(unsigned int* sizes, uint64_t active_mask) {
const cg::thread_block_tile<warp_size> tile =
cg::tiled_partition<warp_size>(cg::this_thread_block());
#if (__GFX8__ || __GFX9__)
constexpr unsigned int ksize = 64;
#else
constexpr unsigned int ksize = 32;
#endif
const cg::thread_block_tile<ksize> tile =
cg::tiled_partition<ksize>(cg::this_thread_block());
if (active_mask & (static_cast<uint64_t>(1) << tile.thread_rank())) {
BaseType active = cg::coalesced_threads();
sizes[thread_rank_in_grid()] = active.size();
}
}
template <unsigned int warp_size, typename BaseType = cg::coalesced_group>
template <typename BaseType = cg::coalesced_group>
static __global__ void coalesced_group_thread_rank_getter(unsigned int* thread_ranks,
uint64_t active_mask) {
const cg::thread_block_tile<warp_size> tile =
cg::tiled_partition<warp_size>(cg::this_thread_block());
#if (__GFX8__ || __GFX9__)
constexpr unsigned int ksize = 64;
#else
constexpr unsigned int ksize = 32;
#endif
const cg::thread_block_tile<ksize> tile =
cg::tiled_partition<ksize>(cg::this_thread_block());
if (active_mask & (static_cast<uint64_t>(1) << tile.thread_rank())) {
BaseType active = cg::coalesced_threads();
thread_ranks[thread_rank_in_grid()] = active.thread_rank();
}
}
template <unsigned int warp_size>
static __global__ void coalesced_group_non_member_size_getter(unsigned int* sizes,
uint64_t active_mask) {
const cg::thread_block_tile<warp_size> tile =
cg::tiled_partition<warp_size>(cg::this_thread_block());
#if (__GFX8__ || __GFX9__)
constexpr unsigned int ksize = 64;
#else
constexpr unsigned int ksize = 32;
#endif
const cg::thread_block_tile<ksize> tile =
cg::tiled_partition<ksize>(cg::this_thread_block());
if (active_mask & (static_cast<uint64_t>(1) << tile.thread_rank())) {
cg::coalesced_group active = cg::coalesced_threads();
sizes[thread_rank_in_grid()] = cg::group_size(active);
}
}
template <unsigned int warp_size>
static __global__ void coalesced_group_non_member_thread_rank_getter(unsigned int* thread_ranks,
uint64_t active_mask) {
const cg::thread_block_tile<warp_size> tile =
cg::tiled_partition<warp_size>(cg::this_thread_block());
#if (__GFX8__ || _GFX9__)
constexpr unsigned int ksize = 64;
#else
constexpr unsigned int ksize = 32;
#endif
const cg::thread_block_tile<ksize> tile =
cg::tiled_partition<ksize>(cg::this_thread_block());
if (active_mask & (static_cast<uint64_t>(1) << tile.thread_rank())) {
cg::coalesced_group active = cg::coalesced_threads();
thread_ranks[thread_rank_in_grid()] = cg::thread_rank(active);
@@ -82,14 +100,14 @@ static unsigned int get_active_thread_count(uint64_t active_mask, unsigned int p
return active_thread_count;
}
static uint64_t get_active_mask(unsigned int test_case) {
static uint64_t get_active_mask(unsigned int test_case, size_t warp_size) {
uint64_t active_mask = 0;
switch (test_case) {
case 0: // 1st thread
active_mask = 1;
break;
case 1: // last thread
active_mask = static_cast<uint64_t>(1) << (kWarpSize - 1);
active_mask = static_cast<uint64_t>(1) << (warp_size - 1);
break;
case 2: // all threads
active_mask = 0xFFFFFFFFFFFFFFFF;
@@ -120,20 +138,18 @@ static uint64_t get_active_mask(unsigned int test_case) {
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Coalesced_Group_Getters_Positive_Basic") {
int device;
hipDeviceProp_t device_properties;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&device_properties, device));
if (!device_properties.cooperativeLaunch) {
HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!");
return;
}
size_t warp_size = static_cast<size_t>(device_properties.warpSize);
const auto blocks = GenerateBlockDimensionsForShuffle();
const auto threads = GenerateThreadDimensionsForShuffle();
auto test_case = GENERATE(range(0, 4));
uint64_t active_mask = get_active_mask(test_case);
uint64_t active_mask = get_active_mask(test_case, warp_size);
INFO("Grid dimensions: x " << blocks.x << ", y " << blocks.y << ", z " << blocks.z);
INFO("Block dimensions: x " << threads.x << ", y " << threads.y << ", z " << threads.z);
INFO("Coalesced group mask: " << active_mask);
@@ -146,29 +162,28 @@ TEST_CASE("Unit_Coalesced_Group_Getters_Positive_Basic") {
HIP_CHECK(hipMemset(uint_arr_dev.ptr(), 0, grid.thread_count_ * sizeof(unsigned int)));
// Launch Kernel
coalesced_group_size_getter<kWarpSize><<<blocks, threads>>>(uint_arr_dev.ptr(), active_mask);
coalesced_group_size_getter<<<blocks, threads>>>(uint_arr_dev.ptr(), active_mask);
HIP_CHECK(hipMemcpy(uint_arr.ptr(), uint_arr_dev.ptr(),
grid.thread_count_ * sizeof(*uint_arr.ptr()), hipMemcpyDeviceToHost));
HIP_CHECK(hipMemset(uint_arr_dev.ptr(), 0, grid.thread_count_ * sizeof(unsigned int)));
HIP_CHECK(hipDeviceSynchronize());
coalesced_group_thread_rank_getter<kWarpSize>
<<<blocks, threads>>>(uint_arr_dev.ptr(), active_mask);
coalesced_group_thread_rank_getter<<<blocks, threads>>>(uint_arr_dev.ptr(), active_mask);
// Verify coalesced_group.size() values
unsigned int coalesced_size = 0;
const auto partitions_in_block = (grid.threads_in_block_count_ + kWarpSize - 1) / kWarpSize;
const auto partitions_in_block = (grid.threads_in_block_count_ + warp_size - 1) / warp_size;
for (int i = 0; i < grid.thread_count_; i++) {
const auto rank_in_block = grid.thread_rank_in_block(i).value();
const int rank_in_partition = rank_in_block % kWarpSize;
const int rank_in_partition = rank_in_block % 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 coalesced group size must be recalculated
if (rank_in_block == (partitions_in_block - 1) * kWarpSize) {
if (rank_in_block == (partitions_in_block - 1) * warp_size) {
unsigned int partition_size =
grid.threads_in_block_count_ - (partitions_in_block - 1) * kWarpSize;
grid.threads_in_block_count_ - (partitions_in_block - 1) * warp_size;
coalesced_size = get_active_thread_count(active_mask, partition_size);
} else if (rank_in_block == 0) {
coalesced_size = get_active_thread_count(active_mask, kWarpSize);
coalesced_size = get_active_thread_count(active_mask, warp_size);
}
if (active_mask & (static_cast<uint64_t>(1) << rank_in_partition)) {
if (uint_arr.ptr()[i] != coalesced_size) {
@@ -185,7 +200,7 @@ TEST_CASE("Unit_Coalesced_Group_Getters_Positive_Basic") {
unsigned int coalesced_rank = 0;
for (int i = 0; i < grid.thread_count_; i++) {
const auto rank_in_block = grid.thread_rank_in_block(i).value();
const int rank_in_partition = rank_in_block % kWarpSize;
const int rank_in_partition = rank_in_block % warp_size;
if (rank_in_partition == 0) coalesced_rank = 0;
if (active_mask & (static_cast<uint64_t>(1) << rank_in_partition)) {
@@ -217,15 +232,13 @@ TEST_CASE("Unit_Coalesced_Group_Getters_Via_Base_Type_Positive_Basic") {
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&device_properties, device));
if (!device_properties.cooperativeLaunch) {
HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!");
return;
}
size_t warp_size = static_cast<size_t>(device_properties.warpSize);
const auto blocks = GenerateBlockDimensionsForShuffle();
const auto threads = GenerateThreadDimensionsForShuffle();
auto test_case = GENERATE(range(0, 4));
uint64_t active_mask = get_active_mask(test_case);
uint64_t active_mask = get_active_mask(test_case, warp_size);
INFO("Grid dimensions: x " << blocks.x << ", y " << blocks.y << ", z " << blocks.z);
INFO("Block dimensions: x " << threads.x << ", y " << threads.y << ", z " << threads.z);
INFO("Coalesced group mask: " << active_mask);
@@ -239,30 +252,28 @@ TEST_CASE("Unit_Coalesced_Group_Getters_Via_Base_Type_Positive_Basic") {
HIP_CHECK(hipMemset(uint_arr_dev.ptr(), 0, grid.thread_count_ * sizeof(unsigned int)));
// Launch Kernel
coalesced_group_size_getter<kWarpSize, cg::thread_group>
<<<blocks, threads>>>(uint_arr_dev.ptr(), active_mask);
coalesced_group_size_getter<<<blocks, threads>>>(uint_arr_dev.ptr(), active_mask);
HIP_CHECK(hipMemcpy(uint_arr.ptr(), uint_arr_dev.ptr(),
grid.thread_count_ * sizeof(*uint_arr.ptr()), hipMemcpyDeviceToHost));
HIP_CHECK(hipMemset(uint_arr_dev.ptr(), 0, grid.thread_count_ * sizeof(unsigned int)));
HIP_CHECK(hipDeviceSynchronize());
coalesced_group_thread_rank_getter<kWarpSize, cg::thread_group>
<<<blocks, threads>>>(uint_arr_dev.ptr(), active_mask);
coalesced_group_thread_rank_getter<<<blocks, threads>>>(uint_arr_dev.ptr(), active_mask);
// Verify coalesced_group.size() values
unsigned int coalesced_size = 0;
const auto partitions_in_block = (grid.threads_in_block_count_ + kWarpSize - 1) / kWarpSize;
const auto partitions_in_block = (grid.threads_in_block_count_ + warp_size - 1) / warp_size;
for (int i = 0; i < grid.thread_count_; i++) {
const auto rank_in_block = grid.thread_rank_in_block(i).value();
const int rank_in_partition = rank_in_block % kWarpSize;
const int rank_in_partition = rank_in_block % 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 coalesced group size must be recalculated
if (rank_in_block == (partitions_in_block - 1) * kWarpSize) {
if (rank_in_block == (partitions_in_block - 1) * warp_size ) {
unsigned int partition_size =
grid.threads_in_block_count_ - (partitions_in_block - 1) * kWarpSize;
grid.threads_in_block_count_ - (partitions_in_block - 1) * warp_size;
coalesced_size = get_active_thread_count(active_mask, partition_size);
} else if (rank_in_block == 0) {
coalesced_size = get_active_thread_count(active_mask, kWarpSize);
coalesced_size = get_active_thread_count(active_mask, warp_size);
}
if (active_mask & (static_cast<uint64_t>(1) << rank_in_partition)) {
if (uint_arr.ptr()[i] != coalesced_size) {
@@ -279,7 +290,7 @@ TEST_CASE("Unit_Coalesced_Group_Getters_Via_Base_Type_Positive_Basic") {
unsigned int coalesced_rank = 0;
for (int i = 0; i < grid.thread_count_; i++) {
const auto rank_in_block = grid.thread_rank_in_block(i).value();
const int rank_in_partition = rank_in_block % kWarpSize;
const int rank_in_partition = rank_in_block % warp_size;
if (rank_in_partition == 0) coalesced_rank = 0;
if (active_mask & (static_cast<uint64_t>(1) << rank_in_partition)) {
@@ -311,15 +322,12 @@ TEST_CASE("Unit_Coalesced_Group_Getters_Via_Non_Member_Functions_Positive_Basic"
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&device_properties, device));
if (!device_properties.cooperativeLaunch) {
HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!");
return;
}
size_t warp_size = static_cast<size_t>(device_properties.warpSize);
const auto blocks = GenerateBlockDimensionsForShuffle();
const auto threads = GenerateThreadDimensionsForShuffle();
auto test_case = GENERATE(range(0, 4));
uint64_t active_mask = get_active_mask(test_case);
uint64_t active_mask = get_active_mask(test_case, warp_size);
INFO("Grid dimensions: x " << blocks.x << ", y " << blocks.y << ", z " << blocks.z);
INFO("Block dimensions: x " << threads.x << ", y " << threads.y << ", z " << threads.z);
INFO("Coalesced group mask: " << active_mask);
@@ -333,30 +341,28 @@ TEST_CASE("Unit_Coalesced_Group_Getters_Via_Non_Member_Functions_Positive_Basic"
HIP_CHECK(hipMemset(uint_arr_dev.ptr(), 0, grid.thread_count_ * sizeof(unsigned int)));
// Launch Kernel
coalesced_group_non_member_size_getter<kWarpSize>
<<<blocks, threads>>>(uint_arr_dev.ptr(), active_mask);
coalesced_group_non_member_size_getter<<<blocks, threads>>>(uint_arr_dev.ptr(), active_mask);
HIP_CHECK(hipMemcpy(uint_arr.ptr(), uint_arr_dev.ptr(),
grid.thread_count_ * sizeof(*uint_arr.ptr()), hipMemcpyDeviceToHost));
HIP_CHECK(hipMemset(uint_arr_dev.ptr(), 0, grid.thread_count_ * sizeof(unsigned int)));
HIP_CHECK(hipDeviceSynchronize());
coalesced_group_non_member_thread_rank_getter<kWarpSize>
<<<blocks, threads>>>(uint_arr_dev.ptr(), active_mask);
coalesced_group_non_member_thread_rank_getter<<<blocks, threads>>>(uint_arr_dev.ptr(), active_mask);
// Verify coalesced_group.size() values
unsigned int coalesced_size = 0;
const auto partitions_in_block = (grid.threads_in_block_count_ + kWarpSize - 1) / kWarpSize;
const auto partitions_in_block = (grid.threads_in_block_count_ + warp_size - 1) / warp_size;
for (int i = 0; i < grid.thread_count_; i++) {
const auto rank_in_block = grid.thread_rank_in_block(i).value();
const int rank_in_partition = rank_in_block % kWarpSize;
const int rank_in_partition = rank_in_block % 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 coalesced group size must be recalculated
if (rank_in_block == (partitions_in_block - 1) * kWarpSize) {
if (rank_in_block == (partitions_in_block - 1) * warp_size) {
unsigned int partition_size =
grid.threads_in_block_count_ - (partitions_in_block - 1) * kWarpSize;
grid.threads_in_block_count_ - (partitions_in_block - 1) * warp_size;
coalesced_size = get_active_thread_count(active_mask, partition_size);
} else if (rank_in_block == 0) {
coalesced_size = get_active_thread_count(active_mask, kWarpSize);
coalesced_size = get_active_thread_count(active_mask, warp_size);
}
if (active_mask & (static_cast<uint64_t>(1) << rank_in_partition)) {
if (uint_arr.ptr()[i] != coalesced_size) {
@@ -373,7 +379,7 @@ TEST_CASE("Unit_Coalesced_Group_Getters_Via_Non_Member_Functions_Positive_Basic"
unsigned int coalesced_rank = 0;
for (int i = 0; i < grid.thread_count_; i++) {
const auto rank_in_block = grid.thread_rank_in_block(i).value();
const int rank_in_partition = rank_in_block % kWarpSize;
const int rank_in_partition = rank_in_block % warp_size;
if (rank_in_partition == 0) coalesced_rank = 0;
if (active_mask & (static_cast<uint64_t>(1) << rank_in_partition)) {
@@ -385,11 +391,16 @@ TEST_CASE("Unit_Coalesced_Group_Getters_Via_Non_Member_Functions_Positive_Basic"
}
}
template <typename T, unsigned int warp_size>
template <typename T>
__global__ void coalesced_group_shfl_up(T* const out, const unsigned int delta,
const uint64_t active_mask) {
const cg::thread_block_tile<warp_size> tile =
cg::tiled_partition<warp_size>(cg::this_thread_block());
#if (__GFX8__ || __GFX9__)
constexpr unsigned int ksize = 64;
#else
constexpr unsigned int ksize = 32;
#endif
const cg::thread_block_tile<ksize> tile =
cg::tiled_partition<ksize>(cg::this_thread_block());
if (active_mask & (static_cast<uint64_t>(1) << tile.thread_rank())) {
cg::coalesced_group active = cg::coalesced_threads();
T var = static_cast<T>(active.thread_rank());
@@ -398,14 +409,22 @@ __global__ void coalesced_group_shfl_up(T* const out, const unsigned int delta,
}
template <typename T> void CoalescedGroupShflUpTestImpl() {
int device;
hipDeviceProp_t device_properties;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&device_properties, device));
size_t warp_size = static_cast<size_t>(device_properties.warpSize);
const auto blocks = GenerateBlockDimensionsForShuffle();
const auto threads = GenerateThreadDimensionsForShuffle();
auto test_case = GENERATE(range(0, 4));
uint64_t active_mask = get_active_mask(test_case);
uint64_t active_mask = get_active_mask(test_case, warp_size);
INFO("Grid dimensions: x " << blocks.x << ", y " << blocks.y << ", z " << blocks.z);
INFO("Block dimensions: x " << threads.x << ", y " << threads.y << ", z " << threads.z);
INFO("Coalesced group mask: " << active_mask);
unsigned int active_thread_count = get_active_thread_count(active_mask, kWarpSize);
unsigned int active_thread_count = get_active_thread_count(active_mask, warp_size);
auto delta = GENERATE(range(static_cast<size_t>(0), kWarpSize));
delta = delta % active_thread_count;
@@ -416,14 +435,14 @@ template <typename T> void CoalescedGroupShflUpTestImpl() {
LinearAllocGuard<T> arr_dev(LinearAllocs::hipMalloc, alloc_size);
LinearAllocGuard<T> arr(LinearAllocs::hipHostMalloc, alloc_size);
coalesced_group_shfl_up<T, kWarpSize><<<blocks, threads>>>(arr_dev.ptr(), delta, active_mask);
coalesced_group_shfl_up<T><<<blocks, threads>>>(arr_dev.ptr(), delta, active_mask);
HIP_CHECK(hipMemcpy(arr.ptr(), arr_dev.ptr(), alloc_size, hipMemcpyDeviceToHost));
HIP_CHECK(hipDeviceSynchronize());
unsigned int coalesced_rank = 0;
for (int i = 0; i < grid.thread_count_; i++) {
const auto rank_in_block = grid.thread_rank_in_block(i).value();
const int rank_in_partition = rank_in_block % kWarpSize;
const int rank_in_partition = rank_in_block % warp_size;
if (rank_in_partition == 0) coalesced_rank = 0;
if (active_mask & (static_cast<uint64_t>(1) << rank_in_partition)) {
int target = coalesced_rank - delta;
@@ -454,11 +473,16 @@ TEMPLATE_TEST_CASE("Unit_Coalesced_Group_Shfl_Up_Positive_Basic", "", int, unsig
CoalescedGroupShflUpTestImpl<TestType>();
}
template <typename T, unsigned int warp_size>
template <typename T>
__global__ void coalesced_group_shfl_down(T* const out, const unsigned int delta,
const uint64_t active_mask) {
const cg::thread_block_tile<warp_size> tile =
cg::tiled_partition<warp_size>(cg::this_thread_block());
#if (__GFX8__ || __GFX9__)
constexpr unsigned int ksize = 64;
#else
constexpr unsigned int ksize = 32;
#endif
const cg::thread_block_tile<ksize> tile =
cg::tiled_partition<ksize>(cg::this_thread_block());
if (active_mask & (static_cast<uint64_t>(1) << tile.thread_rank())) {
cg::coalesced_group active = cg::coalesced_threads();
T var = static_cast<T>(active.thread_rank());
@@ -467,14 +491,22 @@ __global__ void coalesced_group_shfl_down(T* const out, const unsigned int delta
}
template <typename T> void CoalescedGroupShflDownTest() {
int device;
hipDeviceProp_t device_properties;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&device_properties, device));
size_t warp_size = static_cast<size_t>(device_properties.warpSize);
const auto blocks = GenerateBlockDimensionsForShuffle();
const auto threads = GenerateThreadDimensionsForShuffle();
auto test_case = GENERATE(range(0, 4));
uint64_t active_mask = get_active_mask(test_case);
uint64_t active_mask = get_active_mask(test_case, warp_size);
INFO("Grid dimensions: x " << blocks.x << ", y " << blocks.y << ", z " << blocks.z);
INFO("Block dimensions: x " << threads.x << ", y " << threads.y << ", z " << threads.z);
INFO("Coalesced group mask: " << active_mask);
unsigned int active_thread_count = get_active_thread_count(active_mask, kWarpSize);
unsigned int active_thread_count = get_active_thread_count(active_mask, warp_size);
auto delta = GENERATE(range(static_cast<size_t>(0), kWarpSize));
delta = delta % active_thread_count;
@@ -485,25 +517,25 @@ template <typename T> void CoalescedGroupShflDownTest() {
LinearAllocGuard<T> arr_dev(LinearAllocs::hipMalloc, alloc_size);
LinearAllocGuard<T> arr(LinearAllocs::hipHostMalloc, alloc_size);
coalesced_group_shfl_down<T, kWarpSize><<<blocks, threads>>>(arr_dev.ptr(), delta, active_mask);
coalesced_group_shfl_down<T><<<blocks, threads>>>(arr_dev.ptr(), delta, active_mask);
HIP_CHECK(hipMemcpy(arr.ptr(), arr_dev.ptr(), alloc_size, hipMemcpyDeviceToHost));
HIP_CHECK(hipDeviceSynchronize());
unsigned int coalesced_rank = 0;
unsigned int coalesced_size = 0;
const auto partitions_in_block = (grid.threads_in_block_count_ + kWarpSize - 1) / kWarpSize;
const auto partitions_in_block = (grid.threads_in_block_count_ + warp_size - 1) / warp_size;
for (int i = 0; i < grid.thread_count_; i++) {
const auto rank_in_block = grid.thread_rank_in_block(i).value();
const int rank_in_partition = rank_in_block % kWarpSize;
const int rank_in_partition = rank_in_block % warp_size;
if (rank_in_partition == 0) coalesced_rank = 0;
// If the number of threads in a block is not a multiple of warp size, the
// last warp will have inactive threads and coalesced group size must be recalculated
if (rank_in_block == (partitions_in_block - 1) * kWarpSize) {
if (rank_in_block == (partitions_in_block - 1) * warp_size) {
unsigned int partition_size =
grid.threads_in_block_count_ - (partitions_in_block - 1) * kWarpSize;
grid.threads_in_block_count_ - (partitions_in_block - 1) * warp_size;
coalesced_size = get_active_thread_count(active_mask, partition_size);
} else if (rank_in_block == 0) {
coalesced_size = get_active_thread_count(active_mask, kWarpSize);
coalesced_size = get_active_thread_count(active_mask, warp_size);
}
if (active_mask & (static_cast<uint64_t>(1) << rank_in_partition)) {
int target = coalesced_rank + delta;
@@ -533,11 +565,16 @@ TEMPLATE_TEST_CASE("Unit_Coalesced_Group_Shfl_Down_Positive_Basic", "", int, uns
CoalescedGroupShflDownTest<TestType>();
}
template <typename T, unsigned int warp_size>
template <typename T>
__global__ void coalesced_group_shfl(T* const out, uint8_t* target_lanes,
const uint64_t active_mask) {
const cg::thread_block_tile<warp_size> tile =
cg::tiled_partition<warp_size>(cg::this_thread_block());
#if (__GFX8__ || __GFX9__)
constexpr unsigned int ksize = 64;
#else
constexpr unsigned int ksize = 32;
#endif
const cg::thread_block_tile<ksize> tile =
cg::tiled_partition<ksize>(cg::this_thread_block());
if (active_mask & (static_cast<uint64_t>(1) << tile.thread_rank())) {
cg::coalesced_group active = cg::coalesced_threads();
T var = static_cast<T>(active.thread_rank());
@@ -547,14 +584,22 @@ __global__ void coalesced_group_shfl(T* const out, uint8_t* target_lanes,
}
template <typename T> void CoalescedGroupShflTest() {
int device;
hipDeviceProp_t device_properties;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&device_properties, device));
size_t warp_size = static_cast<size_t>(device_properties.warpSize);
const auto blocks = GenerateBlockDimensionsForShuffle();
const auto threads = GenerateThreadDimensionsForShuffle();
auto test_case = GENERATE(range(0, 4));
uint64_t active_mask = get_active_mask(test_case);
uint64_t active_mask = get_active_mask(test_case, warp_size);
INFO("Grid dimensions: x " << blocks.x << ", y " << blocks.y << ", z " << blocks.z);
INFO("Block dimensions: x " << threads.x << ", y " << threads.y << ", z " << threads.z);
INFO("Coalesced group mask: " << active_mask);
unsigned int active_thread_count = get_active_thread_count(active_mask, kWarpSize);
unsigned int active_thread_count = get_active_thread_count(active_mask, warp_size);
CPUGrid grid(blocks, threads);
const auto alloc_size = grid.thread_count_ * sizeof(T);
@@ -572,31 +617,36 @@ template <typename T> void CoalescedGroupShflTest() {
HIP_CHECK(hipMemcpy(target_lanes_dev.ptr(), target_lanes.ptr(),
active_thread_count * sizeof(uint8_t), hipMemcpyHostToDevice));
coalesced_group_shfl<T, kWarpSize>
<<<blocks, threads>>>(arr_dev.ptr(), target_lanes_dev.ptr(), active_mask);
coalesced_group_shfl<T><<<blocks, threads>>>(arr_dev.ptr(), target_lanes_dev.ptr(), active_mask);
HIP_CHECK(hipMemcpy(arr.ptr(), arr_dev.ptr(), alloc_size, hipMemcpyDeviceToHost));
HIP_CHECK(hipDeviceSynchronize());
unsigned int coalesced_rank = 0;
unsigned int coalesced_size = 0;
const auto partitions_in_block = (grid.threads_in_block_count_ + kWarpSize - 1) / kWarpSize;
const auto partitions_in_block = (grid.threads_in_block_count_ + warp_size - 1) / warp_size;
for (int i = 0; i < grid.thread_count_; i++) {
const auto rank_in_block = grid.thread_rank_in_block(i).value();
const int rank_in_partition = rank_in_block % kWarpSize;
const int rank_in_partition = rank_in_block % warp_size;
if (rank_in_partition == 0) coalesced_rank = 0;
// If the number of threads in a block is not a multiple of warp size, the
// last warp will have inactive threads and coalesced group size must be recalculated
if (rank_in_block == (partitions_in_block - 1) * kWarpSize) {
if (rank_in_block == (partitions_in_block - 1) * warp_size) {
unsigned int partition_size =
grid.threads_in_block_count_ - (partitions_in_block - 1) * kWarpSize;
grid.threads_in_block_count_ - (partitions_in_block - 1) * warp_size;
coalesced_size = get_active_thread_count(active_mask, partition_size);
} else if (rank_in_block == 0) {
coalesced_size = get_active_thread_count(active_mask, kWarpSize);
coalesced_size = get_active_thread_count(active_mask, warp_size);
}
if (active_mask & (static_cast<uint64_t>(1) << rank_in_partition)) {
auto target = target_lanes.ptr()[coalesced_rank];
if (target >= coalesced_size) target = 0;
if (target >= coalesced_size) {
#if HT_NVIDIA
target = 0;
#else
target %= coalesced_size;
#endif
}
if (arr.ptr()[i] != target) {
REQUIRE(arr.ptr()[i] == target);
}
@@ -632,14 +682,21 @@ template <typename T> static inline T GenerateRandomInteger(const T min, const T
return dist(GetRandomGenerator());
}
template <bool use_global, size_t warp_size, typename T>
template <bool use_global, typename T>
__global__ void coalesced_group_sync_check(T* global_data, unsigned int* wait_modifiers,
const uint64_t active_mask) {
#if (__GFX8__ || __GFX9__)
constexpr unsigned int ksize = 64;
#else
constexpr unsigned int ksize = 32;
#endif
extern __shared__ uint8_t shared_data[];
T* const data = use_global ? global_data : reinterpret_cast<T*>(shared_data);
const auto tid = cg::this_grid().thread_rank();
const auto block = cg::this_thread_block();
const cg::thread_block_tile<warp_size> partition = cg::tiled_partition<warp_size>(block);
const cg::thread_block_tile<ksize> partition = cg::tiled_partition<ksize>(block);
const auto data_idx = [&block](unsigned int i) { return use_global ? i : (i % block.size()); };
@@ -678,11 +735,20 @@ __global__ void coalesced_group_sync_check(T* global_data, unsigned int* wait_mo
}
template <bool global_memory, typename T> void CoalescedGroupSyncTest() {
int device;
hipDeviceProp_t device_properties;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&device_properties, device));
size_t warp_size = static_cast<size_t>(device_properties.warpSize);
const auto randomized_run_count = GENERATE(range(0, cmd_options.cg_iterations));
const auto blocks = GenerateBlockDimensionsForShuffle();
const auto threads = GenerateThreadDimensionsForShuffle();
auto test_case = GENERATE(range(0, 4));
uint64_t active_mask = get_active_mask(test_case);
uint64_t active_mask = get_active_mask(test_case, warp_size);
INFO("Grid dimensions: x " << blocks.x << ", y " << blocks.y << ", z " << blocks.z);
INFO("Block dimensions: x " << threads.x << ", y " << threads.y << ", z " << threads.z);
INFO("Coalesced group mask: " << active_mask);
@@ -715,16 +781,16 @@ template <bool global_memory, typename T> void CoalescedGroupSyncTest() {
HIP_CHECK(hipMemcpy(wait_modifiers_dev.ptr(), wait_modifiers.ptr(),
grid.thread_count_ * sizeof(unsigned int), hipMemcpyHostToDevice));
coalesced_group_sync_check<global_memory, kWarpSize><<<blocks, threads, shared_memory_size>>>(
coalesced_group_sync_check<global_memory><<<blocks, threads, shared_memory_size>>>(
arr_dev.ptr(), wait_modifiers_dev.ptr(), active_mask);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(arr.ptr(), arr_dev.ptr(), alloc_size, hipMemcpyDeviceToHost));
HIP_CHECK(hipDeviceSynchronize());
for (int i = 0; i < grid.thread_count_; i++) {
const auto rank_in_block = grid.thread_rank_in_block(i).value();
const int rank_in_partition = rank_in_block % kWarpSize;
const int rank_in_partition = rank_in_block % warp_size;
if (active_mask & (static_cast<uint64_t>(1) << rank_in_partition)) {
if (arr.ptr()[i] != 1) {
REQUIRE(arr.ptr()[i] == 1);
projects/hip-tests/catch/unit/cooperativeGrps/cooperative_groups_common.hh
+1 -4
View File
@@ -23,11 +23,7 @@ THE SOFTWARE.
#include <hip/hip_cooperative_groups.h>
namespace {
#if (!__GFX8__ && !__GFX9__) || HT_NVIDIA
constexpr size_t kWarpSize = 32;
#else
constexpr size_t kWarpSize = 64;
#endif
constexpr int kMaxGPUs = 8;
} // namespace
@@ -59,6 +55,7 @@ static __device__ void busy_wait(unsigned long long wait_period) {
}
}
template <class T> bool CheckDimensions(unsigned int device, T kernel, dim3 blocks, dim3 threads) {
hipDeviceProp_t props;
int max_blocks_per_sm = 0;
projects/hip-tests/catch/unit/cooperativeGrps/thread_block_tile.cc
-24
View File
@@ -121,9 +121,6 @@ template <bool dynamic, size_t... tile_sizes> void BlockPartitionGettersBasicTes
*/
TEST_CASE("Unit_Thread_Block_Tile_Getters_Positive_Basic") {
BlockPartitionGettersBasicTest<false, 2, 4, 8, 16, 32>();
#if HT_AMD && (__GFX8__ || __GFX9__)
BlockPartitionGettersBasicTest<false, 64>();
#endif
}
/**
@@ -141,9 +138,6 @@ TEST_CASE("Unit_Thread_Block_Tile_Getters_Positive_Basic") {
*/
TEST_CASE("Unit_Thread_Block_Tile_Dynamic_Getters_Positive_Basic") {
BlockPartitionGettersBasicTest<true, 2, 4, 8, 16, 32>();
#if HT_AMD && (__GFX8__ || __GFX9__)
BlockPartitionGettersBasicTest<true, 64>();
#endif
}
@@ -201,9 +195,6 @@ template <typename T, size_t... tile_sizes> void BlockTileShflUpTest() {
TEMPLATE_TEST_CASE("Unit_Thread_Block_Tile_Shfl_Up_Positive_Basic", "", int, unsigned int, long,
unsigned long, long long, unsigned long long, float, double) {
BlockTileShflUpTest<TestType, 2, 4, 8, 16, 32>();
#if HT_AMD && (__GFX8__ || __GFX9__)
BlockTileShflUpTest<TestType, 64>();
#endif
}
@@ -273,9 +264,6 @@ template <typename T, size_t... tile_sizes> void BlockTileShflDownTest() {
TEMPLATE_TEST_CASE("Unit_Thread_Block_Tile_Shfl_Down_Positive_Basic", "", int, unsigned int, long,
unsigned long, long long, unsigned long long, float, double) {
BlockTileShflDownTest<TestType, 2, 4, 8, 16, 32>();
#if HT_AMD && (__GFX8__ || __GFX9__)
BlockTileShflDownTest<TestType, 64>();
#endif
}
@@ -340,9 +328,6 @@ template <typename T, size_t... tile_sizes> void BlockTileShflXORTest() {
TEMPLATE_TEST_CASE("Unit_Thread_Block_Tile_Shfl_XOR_Positive_Basic", "", int, unsigned int, long,
unsigned long, long long, unsigned long long, float, double) {
BlockTileShflXORTest<TestType, 2, 4, 8, 16, 32>();
#if HT_AMD && (__GFX8__ || __GFX9__)
BlockTileShflXORTest<TestType, 64>();
#endif
}
template <typename T, size_t tile_size>
@@ -423,9 +408,6 @@ template <typename T, size_t... tile_sizes> void BlockTileShflTest() {
TEMPLATE_TEST_CASE("Unit_Thread_Block_Tile_Shfl_Positive_Basic", "", int, unsigned int, long,
unsigned long, long long, unsigned long long, float, double) {
BlockTileShflTest<TestType, 2, 4, 8, 16, 32>();
#if HT_AMD && (__GFX8__ || __GFX9__)
BlockTileShflTest<TestType, 64>();
#endif
}
@@ -540,15 +522,9 @@ template <bool global_memory, typename T, size_t... tile_sizes> void BlockTileSy
TEMPLATE_TEST_CASE("Unit_Thread_Block_Tile_Sync_Positive_Basic", "", uint8_t, uint16_t, uint32_t) {
SECTION("Global memory") {
BlockTileSyncTest<true, TestType, 2, 4, 8, 16, 32>();
#if HT_AMD && (__GFX8__ || __GFX9__)
BlockTileSyncTest<true, TestType, 64>();
#endif
}
SECTION("Shared memory") {
BlockTileSyncTest<false, TestType, 2, 4, 8, 16, 32>();
#if HT_AMD && (__GFX8__ || __GFX9__)
BlockTileSyncTest<true, TestType, 64>();
#endif
}
}