AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

[SWDEV-438556][SWDEV-449685] - Fix for CG MultiGrid Group Tests

Przeglądaj źródła 
Change-Id: Iaf3d422e57465fefe9adb1fc7c01a480ada37db3
This commit is contained in:
Rahul Manocha
2024-03-09 00:20:12 +00:00
zatwierdzone przez Rakesh Roy
rodzic 4a80eee9f1
commit bcff12dedc
3 zmienionych plików z 23 dodań i 9 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
catch/hipTestMain/config/config_amd_linux
-3
Wyświetl plik
@@ -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",
catch/include/utils.hh
+17
Wyświetl plik
@@ -69,6 +69,23 @@ static inline void ArrayAllOf(const T* arr, uint32_t count, F value_gen) {
}
}
template <typename T>
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 <typename T, typename F>
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) {
catch/unit/cooperativeGrps/multi_grid_group.cc
+6 -6
Wyświetl plik
@@ -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_);
}
}