diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux index 8caafa5e70..4705846baa 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux @@ -1447,9 +1447,6 @@ "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_Warp_Ballot_Positive_Basic", - "Unit_Warp_Vote_Any_Positive_Basic", - "Unit_Warp_Vote_All_Positive_Basic", "=== SWDEV-443630 - Below tests failed in stress test on 19/01/23 ===", "Unit_hipGetSetDevice_MultiThreaded", #endif @@ -1536,9 +1533,6 @@ "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_Warp_Ballot_Positive_Basic", - "Unit_Warp_Vote_Any_Positive_Basic", - "Unit_Warp_Vote_All_Positive_Basic", "=== SWDEV-439298: Below test failing in CQE staging ===", "Unit_hipCGMultiGridGroupType_Barrier", "=== SWDEV-443630 : Below test failed in stress test on 19/01/24 ===", @@ -1580,6 +1574,10 @@ "Unit_safeAtomicMin_Positive_SameAddress - float", "=== SWDEV-454220 : Below test hanged in stress test on 22/03/24 ===", "Unit_hipExtLaunchKernel_Positive_Basic", + "=== Below tests fail in stress test 03/29/24 ===", + "Unit_Warp_Ballot_Positive_Basic", + "Unit_Warp_Vote_Any_Positive_Basic", + "Unit_Warp_Vote_All_Positive_Basic", #endif #if defined NAVI21 "=== SWDEV-445961: These tests hang in PSDB stress test on 09/02/2024 ===", diff --git a/projects/hip-tests/catch/unit/warp/warp_all.cc b/projects/hip-tests/catch/unit/warp/warp_all.cc index 73d7bf782d..1e038b05c0 100644 --- a/projects/hip-tests/catch/unit/warp/warp_all.cc +++ b/projects/hip-tests/catch/unit/warp/warp_all.cc @@ -41,7 +41,7 @@ static bool check_if_all(uint64_t predicate_mask, uint64_t active_mask, size_t p return true; } -__global__ void kernel_all(int* const out, const uint64_t* const active_masks, uint64_t predicate) { +__global__ void kernel_all(uint64_t* const out, const uint64_t* const active_masks, uint64_t predicate) { if (deactivate_thread(active_masks)) { return; } @@ -49,12 +49,13 @@ __global__ void kernel_all(int* const out, const uint64_t* const active_masks, u const auto grid = cg::this_grid(); const auto warp = cg::tiled_partition(cg::this_thread_block(), warpSize); - out[grid.thread_rank()] = __all((predicate & (static_cast(1) << warp.thread_rank()))); + int pred = MASK_SHIFT(predicate, warp.thread_rank()); + out[grid.thread_rank()] = __all(pred); } -class WarpAll : public WarpVoteTest { +class WarpAll : public WarpVoteTest { public: - void launch_kernel(int* const arr_dev, const uint64_t* const active_masks) { + void launch_kernel(uint64_t* const arr_dev, const uint64_t* const active_masks) { auto test_case = GENERATE(range(0, 5)); predicate_mask_ = get_predicate_mask(test_case, this->warp_size_); INFO("Predicate mask: " << predicate_mask_); @@ -62,7 +63,7 @@ class WarpAll : public WarpVoteTest { predicate_mask_); } - void validate(const int* const arr) { + void validate(const uint64_t* const arr) { ArrayAllOf(arr, this->grid_.thread_count_, [this](unsigned int i) -> std::optional { const auto rank_in_block = this->grid_.thread_rank_in_block(i).value(); const auto rank_in_warp = rank_in_block % this->warp_size_; diff --git a/projects/hip-tests/catch/unit/warp/warp_any.cc b/projects/hip-tests/catch/unit/warp/warp_any.cc index 19c21a24f8..fe5c7c96f6 100644 --- a/projects/hip-tests/catch/unit/warp/warp_any.cc +++ b/projects/hip-tests/catch/unit/warp/warp_any.cc @@ -32,7 +32,7 @@ THE SOFTWARE. namespace cg = cooperative_groups; -__global__ void kernel_any(int* const out, const uint64_t* const active_masks, uint64_t predicate) { +__global__ void kernel_any(uint64_t* const out, const uint64_t* const active_masks, uint64_t predicate) { if (deactivate_thread(active_masks)) { return; } @@ -40,12 +40,13 @@ __global__ void kernel_any(int* const out, const uint64_t* const active_masks, u const auto grid = cg::this_grid(); const auto warp = cg::tiled_partition(cg::this_thread_block(), warpSize); - out[grid.thread_rank()] = __any((predicate & (static_cast(1) << warp.thread_rank()))); + int pred = MASK_SHIFT(predicate, warp.thread_rank()); + out[grid.thread_rank()] = __any(pred); } -class WarpAny : public WarpVoteTest { +class WarpAny : public WarpVoteTest { public: - void launch_kernel(int* const arr_dev, const uint64_t* const active_masks) { + void launch_kernel(uint64_t* const arr_dev, const uint64_t* const active_masks) { auto test_case = GENERATE(range(0, 5)); predicate_mask_ = get_predicate_mask(test_case, this->warp_size_); INFO("Predicate mask: " << predicate_mask_); @@ -53,7 +54,7 @@ class WarpAny : public WarpVoteTest { predicate_mask_); } - void validate(const int* const arr) { + void validate(const uint64_t* const arr) { ArrayAllOf(arr, this->grid_.thread_count_, [this](unsigned int i) -> std::optional { const auto rank_in_block = this->grid_.thread_rank_in_block(i).value(); const auto rank_in_warp = rank_in_block % this->warp_size_; diff --git a/projects/hip-tests/catch/unit/warp/warp_ballot.cc b/projects/hip-tests/catch/unit/warp/warp_ballot.cc index 168c794dbb..7fa2b0d562 100644 --- a/projects/hip-tests/catch/unit/warp/warp_ballot.cc +++ b/projects/hip-tests/catch/unit/warp/warp_ballot.cc @@ -40,8 +40,8 @@ __global__ void kernel_ballot(uint64_t* const out, const uint64_t* const active_ const auto grid = cg::this_grid(); const auto warp = cg::tiled_partition(cg::this_thread_block(), warpSize); - out[grid.thread_rank()] = - __ballot((predicate & (static_cast(1) << warp.thread_rank()))); + int pred = MASK_SHIFT(predicate, warp.thread_rank()); + out[grid.thread_rank()] = __ballot(pred); } class WarpBallot : public WarpVoteTest { diff --git a/projects/hip-tests/catch/unit/warp/warp_common.hh b/projects/hip-tests/catch/unit/warp/warp_common.hh index e0870492f7..9fe207c030 100644 --- a/projects/hip-tests/catch/unit/warp/warp_common.hh +++ b/projects/hip-tests/catch/unit/warp/warp_common.hh @@ -25,6 +25,9 @@ THE SOFTWARE. #include #include +#define MASK_SHIFT(x, n) \ + (x & (static_cast(1) << n)) >> n + const unsigned long long Every5thBit = 0x1084210842108421; const unsigned long long Every9thBit = 0x8040201008040201; const unsigned long long Every5thBut9th = Every5thBit & ~Every9thBit; @@ -37,7 +40,6 @@ inline __device__ bool deactivate_thread(const uint64_t* const active_masks) { const auto warps_per_block = (block.size() + warpSize - 1) / warpSize; const auto block_rank = (blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x; const auto idx = block_rank * warps_per_block + block.thread_rank() / warpSize; - return !(active_masks[idx] & (static_cast(1) << warp.thread_rank())); }