diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux index 3d7a078aaa..97a2c2a0a8 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux @@ -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", diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/coalesced_group.cc b/projects/hip-tests/catch/unit/cooperativeGrps/coalesced_group.cc index 9956134dab..4ff7e7d510 100644 --- a/projects/hip-tests/catch/unit/cooperativeGrps/coalesced_group.cc +++ b/projects/hip-tests/catch/unit/cooperativeGrps/coalesced_group.cc @@ -31,43 +31,61 @@ THE SOFTWARE. namespace cg = cooperative_groups; -template +template static __global__ void coalesced_group_size_getter(unsigned int* sizes, uint64_t active_mask) { - const cg::thread_block_tile tile = - cg::tiled_partition(cg::this_thread_block()); + #if (__GFX8__ || __GFX9__) + constexpr unsigned int ksize = 64; + #else + constexpr unsigned int ksize = 32; + #endif + const cg::thread_block_tile tile = + cg::tiled_partition(cg::this_thread_block()); if (active_mask & (static_cast(1) << tile.thread_rank())) { BaseType active = cg::coalesced_threads(); sizes[thread_rank_in_grid()] = active.size(); } } -template +template static __global__ void coalesced_group_thread_rank_getter(unsigned int* thread_ranks, uint64_t active_mask) { - const cg::thread_block_tile tile = - cg::tiled_partition(cg::this_thread_block()); + #if (__GFX8__ || __GFX9__) + constexpr unsigned int ksize = 64; + #else + constexpr unsigned int ksize = 32; + #endif + const cg::thread_block_tile tile = + cg::tiled_partition(cg::this_thread_block()); if (active_mask & (static_cast(1) << tile.thread_rank())) { BaseType active = cg::coalesced_threads(); thread_ranks[thread_rank_in_grid()] = active.thread_rank(); } } -template static __global__ void coalesced_group_non_member_size_getter(unsigned int* sizes, uint64_t active_mask) { - const cg::thread_block_tile tile = - cg::tiled_partition(cg::this_thread_block()); + #if (__GFX8__ || __GFX9__) + constexpr unsigned int ksize = 64; + #else + constexpr unsigned int ksize = 32; + #endif + const cg::thread_block_tile tile = + cg::tiled_partition(cg::this_thread_block()); if (active_mask & (static_cast(1) << tile.thread_rank())) { cg::coalesced_group active = cg::coalesced_threads(); sizes[thread_rank_in_grid()] = cg::group_size(active); } } -template static __global__ void coalesced_group_non_member_thread_rank_getter(unsigned int* thread_ranks, uint64_t active_mask) { - const cg::thread_block_tile tile = - cg::tiled_partition(cg::this_thread_block()); + #if (__GFX8__ || _GFX9__) + constexpr unsigned int ksize = 64; + #else + constexpr unsigned int ksize = 32; + #endif + const cg::thread_block_tile tile = + cg::tiled_partition(cg::this_thread_block()); if (active_mask & (static_cast(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(1) << (kWarpSize - 1); + active_mask = static_cast(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(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<<>>(uint_arr_dev.ptr(), active_mask); + coalesced_group_size_getter<<>>(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 - <<>>(uint_arr_dev.ptr(), active_mask); + coalesced_group_thread_rank_getter<<>>(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(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(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(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 - <<>>(uint_arr_dev.ptr(), active_mask); + coalesced_group_size_getter<<>>(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 - <<>>(uint_arr_dev.ptr(), active_mask); + coalesced_group_thread_rank_getter<<>>(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(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(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(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 - <<>>(uint_arr_dev.ptr(), active_mask); + coalesced_group_non_member_size_getter<<>>(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 - <<>>(uint_arr_dev.ptr(), active_mask); + coalesced_group_non_member_thread_rank_getter<<>>(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(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(1) << rank_in_partition)) { @@ -385,11 +391,16 @@ TEST_CASE("Unit_Coalesced_Group_Getters_Via_Non_Member_Functions_Positive_Basic" } } -template +template __global__ void coalesced_group_shfl_up(T* const out, const unsigned int delta, const uint64_t active_mask) { - const cg::thread_block_tile tile = - cg::tiled_partition(cg::this_thread_block()); + #if (__GFX8__ || __GFX9__) + constexpr unsigned int ksize = 64; + #else + constexpr unsigned int ksize = 32; + #endif + const cg::thread_block_tile tile = + cg::tiled_partition(cg::this_thread_block()); if (active_mask & (static_cast(1) << tile.thread_rank())) { cg::coalesced_group active = cg::coalesced_threads(); T var = static_cast(active.thread_rank()); @@ -398,14 +409,22 @@ __global__ void coalesced_group_shfl_up(T* const out, const unsigned int delta, } template void CoalescedGroupShflUpTestImpl() { + + int device; + hipDeviceProp_t device_properties; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&device_properties, device)); + + size_t warp_size = static_cast(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(0), kWarpSize)); delta = delta % active_thread_count; @@ -416,14 +435,14 @@ template void CoalescedGroupShflUpTestImpl() { LinearAllocGuard arr_dev(LinearAllocs::hipMalloc, alloc_size); LinearAllocGuard arr(LinearAllocs::hipHostMalloc, alloc_size); - coalesced_group_shfl_up<<>>(arr_dev.ptr(), delta, active_mask); + coalesced_group_shfl_up<<>>(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(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(); } -template +template __global__ void coalesced_group_shfl_down(T* const out, const unsigned int delta, const uint64_t active_mask) { - const cg::thread_block_tile tile = - cg::tiled_partition(cg::this_thread_block()); + #if (__GFX8__ || __GFX9__) + constexpr unsigned int ksize = 64; + #else + constexpr unsigned int ksize = 32; + #endif + const cg::thread_block_tile tile = + cg::tiled_partition(cg::this_thread_block()); if (active_mask & (static_cast(1) << tile.thread_rank())) { cg::coalesced_group active = cg::coalesced_threads(); T var = static_cast(active.thread_rank()); @@ -467,14 +491,22 @@ __global__ void coalesced_group_shfl_down(T* const out, const unsigned int delta } template void CoalescedGroupShflDownTest() { + + int device; + hipDeviceProp_t device_properties; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&device_properties, device)); + + size_t warp_size = static_cast(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(0), kWarpSize)); delta = delta % active_thread_count; @@ -485,25 +517,25 @@ template void CoalescedGroupShflDownTest() { LinearAllocGuard arr_dev(LinearAllocs::hipMalloc, alloc_size); LinearAllocGuard arr(LinearAllocs::hipHostMalloc, alloc_size); - coalesced_group_shfl_down<<>>(arr_dev.ptr(), delta, active_mask); + coalesced_group_shfl_down<<>>(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(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(); } -template +template __global__ void coalesced_group_shfl(T* const out, uint8_t* target_lanes, const uint64_t active_mask) { - const cg::thread_block_tile tile = - cg::tiled_partition(cg::this_thread_block()); + #if (__GFX8__ || __GFX9__) + constexpr unsigned int ksize = 64; + #else + constexpr unsigned int ksize = 32; + #endif + const cg::thread_block_tile tile = + cg::tiled_partition(cg::this_thread_block()); if (active_mask & (static_cast(1) << tile.thread_rank())) { cg::coalesced_group active = cg::coalesced_threads(); T var = static_cast(active.thread_rank()); @@ -547,14 +584,22 @@ __global__ void coalesced_group_shfl(T* const out, uint8_t* target_lanes, } template void CoalescedGroupShflTest() { + + int device; + hipDeviceProp_t device_properties; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&device_properties, device)); + + size_t warp_size = static_cast(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 void CoalescedGroupShflTest() { HIP_CHECK(hipMemcpy(target_lanes_dev.ptr(), target_lanes.ptr(), active_thread_count * sizeof(uint8_t), hipMemcpyHostToDevice)); - coalesced_group_shfl - <<>>(arr_dev.ptr(), target_lanes_dev.ptr(), active_mask); + coalesced_group_shfl<<>>(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(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 static inline T GenerateRandomInteger(const T min, const T return dist(GetRandomGenerator()); } -template +template __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(shared_data); const auto tid = cg::this_grid().thread_rank(); const auto block = cg::this_thread_block(); - const cg::thread_block_tile partition = cg::tiled_partition(block); + const cg::thread_block_tile partition = cg::tiled_partition(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 void CoalescedGroupSyncTest() { + + int device; + hipDeviceProp_t device_properties; + + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&device_properties, device)); + + size_t warp_size = static_cast(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 void CoalescedGroupSyncTest() { HIP_CHECK(hipMemcpy(wait_modifiers_dev.ptr(), wait_modifiers.ptr(), grid.thread_count_ * sizeof(unsigned int), hipMemcpyHostToDevice)); - coalesced_group_sync_check<<>>( + coalesced_group_sync_check<<>>( 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(1) << rank_in_partition)) { if (arr.ptr()[i] != 1) { REQUIRE(arr.ptr()[i] == 1); diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/cooperative_groups_common.hh b/projects/hip-tests/catch/unit/cooperativeGrps/cooperative_groups_common.hh index d12128187f..aeec9e5fcc 100644 --- a/projects/hip-tests/catch/unit/cooperativeGrps/cooperative_groups_common.hh +++ b/projects/hip-tests/catch/unit/cooperativeGrps/cooperative_groups_common.hh @@ -23,11 +23,7 @@ THE SOFTWARE. #include 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 bool CheckDimensions(unsigned int device, T kernel, dim3 blocks, dim3 threads) { hipDeviceProp_t props; int max_blocks_per_sm = 0; diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/thread_block_tile.cc b/projects/hip-tests/catch/unit/cooperativeGrps/thread_block_tile.cc index 15d9e4602f..8ae266b0e4 100644 --- a/projects/hip-tests/catch/unit/cooperativeGrps/thread_block_tile.cc +++ b/projects/hip-tests/catch/unit/cooperativeGrps/thread_block_tile.cc @@ -121,9 +121,6 @@ template void BlockPartitionGettersBasicTes */ TEST_CASE("Unit_Thread_Block_Tile_Getters_Positive_Basic") { BlockPartitionGettersBasicTest(); -#if HT_AMD && (__GFX8__ || __GFX9__) - BlockPartitionGettersBasicTest(); -#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(); -#if HT_AMD && (__GFX8__ || __GFX9__) - BlockPartitionGettersBasicTest(); -#endif } @@ -201,9 +195,6 @@ template 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(); -#if HT_AMD && (__GFX8__ || __GFX9__) - BlockTileShflUpTest(); -#endif } @@ -273,9 +264,6 @@ template 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(); -#if HT_AMD && (__GFX8__ || __GFX9__) - BlockTileShflDownTest(); -#endif } @@ -340,9 +328,6 @@ template 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(); -#if HT_AMD && (__GFX8__ || __GFX9__) - BlockTileShflXORTest(); -#endif } template @@ -423,9 +408,6 @@ template 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(); -#if HT_AMD && (__GFX8__ || __GFX9__) - BlockTileShflTest(); -#endif } @@ -540,15 +522,9 @@ template void BlockTileSy TEMPLATE_TEST_CASE("Unit_Thread_Block_Tile_Sync_Positive_Basic", "", uint8_t, uint16_t, uint32_t) { SECTION("Global memory") { BlockTileSyncTest(); -#if HT_AMD && (__GFX8__ || __GFX9__) - BlockTileSyncTest(); -#endif } SECTION("Shared memory") { BlockTileSyncTest(); -#if HT_AMD && (__GFX8__ || __GFX9__) - BlockTileSyncTest(); -#endif } }