From bcff12dedc2cb0430f31ce84290b306731ee6613 Mon Sep 17 00:00:00 2001 From: Rahul Manocha Date: Sat, 9 Mar 2024 00:20:12 +0000 Subject: [PATCH] [SWDEV-438556][SWDEV-449685] - Fix for CG MultiGrid Group Tests Change-Id: Iaf3d422e57465fefe9adb1fc7c01a480ada37db3 --- catch/hipTestMain/config/config_amd_linux | 3 --- catch/include/utils.hh | 17 +++++++++++++++++ catch/unit/cooperativeGrps/multi_grid_group.cc | 12 ++++++------ 3 files changed, 23 insertions(+), 9 deletions(-) diff --git a/catch/hipTestMain/config/config_amd_linux b/catch/hipTestMain/config/config_amd_linux index 7a6b4acdad..987881ca85 100644 --- a/catch/hipTestMain/config/config_amd_linux +++ b/catch/hipTestMain/config/config_amd_linux @@ -1410,9 +1410,6 @@ "Unit_hipFreeMipmappedArrayMultiTArray - char", "Unit_hipFreeMipmappedArrayMultiTArray - int", "Unit_hipIpcGetMemHandle_Positive_Unique_Handles_Reused_Memory", - "Unit_Multi_Grid_Group_Getters_Positive_Basic", - "Unit_Multi_Grid_Group_Getters_Positive_Base_Type", - "Unit_Multi_Grid_Group_Getters_Positive_Non_Member_Functions", "Unit_Coalesced_Group_Getters_Via_Base_Type_Positive_Basic", "Unit_Coalesced_Group_Shfl_Positive_Basic - int", "Unit_Coalesced_Group_Shfl_Positive_Basic - unsigned int", diff --git a/catch/include/utils.hh b/catch/include/utils.hh index 3855308a42..c195dbd292 100644 --- a/catch/include/utils.hh +++ b/catch/include/utils.hh @@ -69,6 +69,23 @@ static inline void ArrayAllOf(const T* arr, uint32_t count, F value_gen) { } } +template +static inline void ArrayInRange(const T* arr, uint32_t count,const T minval,const T maxval) { + for (auto i = 0u; i < count; ++i) { + if(arr[i] < minval) + { + INFO("Mismatch at index: " << i); + REQUIRE(arr[i] > minval); + } + else if(arr[i] > maxval) + { + INFO("Mismatch at index: " << i); + REQUIRE(arr[i] < maxval); + } + } +} + + template void PitchedMemoryVerify(T* const ptr, const size_t pitch, const size_t width, const size_t height, const size_t depth, F expected_value_generator) { diff --git a/catch/unit/cooperativeGrps/multi_grid_group.cc b/catch/unit/cooperativeGrps/multi_grid_group.cc index 0beedc53e6..a7490bb360 100644 --- a/catch/unit/cooperativeGrps/multi_grid_group.cc +++ b/catch/unit/cooperativeGrps/multi_grid_group.cc @@ -247,8 +247,8 @@ TEST_CASE("Unit_Multi_Grid_Group_Getters_Positive_Basic") { HIP_CHECK(hipSetDevice(i)); // Verify multi_grid_group.thread_rank() values const auto multi_grid_thread0_rank = multi_grid.thread0_rank_in_multi_grid(i); - ArrayAllOf(uint_arr[i].ptr(), multi_grid.grids_[i].thread_count_, - [rank_0 = multi_grid_thread0_rank](uint32_t j) { return rank_0 + j; }); + ArrayInRange(uint_arr[i].ptr(), multi_grid.grids_[i].thread_count_, multi_grid_thread0_rank, + multi_grid_thread0_rank + multi_grid.grids_[i].thread_count_); HIP_CHECK(hipMemcpy(uint_arr[i].ptr(), uint_arr_dev[i].ptr(), multi_grid.grids_[i].thread_count_ * sizeof(*uint_arr[i].ptr()), hipMemcpyDeviceToHost)); @@ -400,8 +400,8 @@ TEST_CASE("Unit_Multi_Grid_Group_Getters_Positive_Base_Type") { HIP_CHECK(hipSetDevice(i)); // Verify multi_grid_group.thread_rank() values const auto multi_grid_thread0_rank = multi_grid.thread0_rank_in_multi_grid(i); - ArrayAllOf(uint_arr[i].ptr(), multi_grid.grids_[i].thread_count_, - [rank_0 = multi_grid_thread0_rank](uint32_t j) { return rank_0 + j; }); + ArrayInRange(uint_arr[i].ptr(), multi_grid.grids_[i].thread_count_, multi_grid_thread0_rank, + multi_grid_thread0_rank + multi_grid.grids_[i].thread_count_); HIP_CHECK(hipMemcpy(uint_arr[i].ptr(), uint_arr_dev[i].ptr(), multi_grid.grids_[i].thread_count_ * sizeof(*uint_arr[i].ptr()), hipMemcpyDeviceToHost)); @@ -508,8 +508,8 @@ TEST_CASE("Unit_Multi_Grid_Group_Getters_Positive_Non_Member_Functions") { HIP_CHECK(hipDeviceSynchronize()); // Verify multi_grid_group.thread_rank() values const auto multi_grid_thread0_rank = multi_grid.thread0_rank_in_multi_grid(i); - ArrayAllOf(uint_arr[i].ptr(), multi_grid.grids_[i].thread_count_, - [rank_0 = multi_grid_thread0_rank](uint32_t j) { return rank_0 + j; }); + ArrayInRange(uint_arr[i].ptr(), multi_grid.grids_[i].thread_count_, multi_grid_thread0_rank, + multi_grid_thread0_rank + multi_grid.grids_[i].thread_count_); } }