From feaa0bbe438f8fc1a243ee05ea3fb030f212d9fc Mon Sep 17 00:00:00 2001 From: kjayapra-amd Date: Wed, 6 Sep 2023 05:14:57 -0400 Subject: [PATCH] SWDEV-420149 - Fixing unsafe atomics change for targets other than gfx90a. Change-Id: Ifa26fd1c51505c5e86e121233c8352e3bde846ea [ROCm/hip-tests commit: a363fa2833bb3bf67575fe3658c744685b725110] --- .../AtomicAdd_Coherent_withunsafeflag.cc | 10 ++++--- .../catch/unit/deviceLib/BuiltIns_fadd.cc | 18 ++++++++++--- ...safeAtomicAdd_Coherent_withnounsafeflag.cc | 9 +++++-- .../unsafeAtomicAdd_Coherent_withoutflag.cc | 9 +++++-- ...unsafeAtomicAdd_Coherent_withunsafeflag.cc | 9 +++++-- .../unit/deviceLib/unsafeAtomicAdd_RTC.cc | 27 ++++++++++++++----- 6 files changed, 63 insertions(+), 19 deletions(-) diff --git a/projects/hip-tests/catch/unit/deviceLib/AtomicAdd_Coherent_withunsafeflag.cc b/projects/hip-tests/catch/unit/deviceLib/AtomicAdd_Coherent_withunsafeflag.cc index b9fffb5de5..ceb9530439 100644 --- a/projects/hip-tests/catch/unit/deviceLib/AtomicAdd_Coherent_withunsafeflag.cc +++ b/projects/hip-tests/catch/unit/deviceLib/AtomicAdd_Coherent_withunsafeflag.cc @@ -88,9 +88,13 @@ TEMPLATE_TEST_CASE("Unit_AtomicAdd_CoherentwithUnsafeflag", "", "global_atomic_add_f64"); REQUIRE(testResult == true); } - - REQUIRE(A_h[0] == INITIAL_VAL); - REQUIRE(result[0] == 0); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + REQUIRE(A_h[0] == INITIAL_VAL); + REQUIRE(result[0] == 0); + } else { + REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL); + REQUIRE(result[0] == INITIAL_VAL); + } HIP_CHECK(hipHostFree(A_h)); HIP_CHECK(hipHostFree(result)); } diff --git a/projects/hip-tests/catch/unit/deviceLib/BuiltIns_fadd.cc b/projects/hip-tests/catch/unit/deviceLib/BuiltIns_fadd.cc index 5303afb519..47a533dc56 100644 --- a/projects/hip-tests/catch/unit/deviceLib/BuiltIns_fadd.cc +++ b/projects/hip-tests/catch/unit/deviceLib/BuiltIns_fadd.cc @@ -80,8 +80,13 @@ TEST_CASE("Unit_BuiltInAtomicAdd_CoherentGlobalMem") { HIP_CHECK(hipGetLastError()); std::cout << "test 1" << std::endl; HIP_CHECK(hipDeviceSynchronize()); - REQUIRE(A_h[0] == INITIAL_VAL); - REQUIRE(*result_h == 0); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + REQUIRE(A_h[0] == INITIAL_VAL); + REQUIRE(*result_h == 0); + } else { + REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL); + REQUIRE(*result_h == INITIAL_VAL); + } HIP_CHECK(hipHostFree(A_h)); HIP_CHECK(hipFree(result)); } @@ -202,8 +207,13 @@ TEST_CASE("Unit_BuiltInAtomicAdd_CoherentGlobalMemWithRtc") { nullptr, nullptr, config_d)); HIP_CHECK(hipDeviceSynchronize()); HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost)); - REQUIRE(A_h[0] == INITIAL_VAL); - REQUIRE(*B_h == 0); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + REQUIRE(A_h[0] == INITIAL_VAL); + REQUIRE(*B_h == 0); + } else { + REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL); + REQUIRE(*B_h == INITIAL_VAL); + } HIP_CHECK(hipHostFree(A_h)); HIP_CHECK(hipFree(result)); free(B_h); diff --git a/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withnounsafeflag.cc b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withnounsafeflag.cc index 91b2fc433f..159bda85f5 100644 --- a/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withnounsafeflag.cc +++ b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withnounsafeflag.cc @@ -89,8 +89,13 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentwithnoUnsafeflag", "", REQUIRE(testResult == true); } - REQUIRE(A_h[0] == INITIAL_VAL); - REQUIRE(result[0] == 0); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + REQUIRE(A_h[0] == INITIAL_VAL); + REQUIRE(result[0] == 0); + } else { + REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL); + REQUIRE(result[0] == INITIAL_VAL); + } HIP_CHECK(hipHostFree(A_h)); HIP_CHECK(hipHostFree(result)); } diff --git a/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withoutflag.cc b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withoutflag.cc index 9999c39344..bd43b1b776 100644 --- a/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withoutflag.cc +++ b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withoutflag.cc @@ -89,8 +89,13 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_Coherentwithoutflag", "", REQUIRE(testResult == true); } - REQUIRE(A_h[0] == INITIAL_VAL); - REQUIRE(result[0] == 0); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + REQUIRE(A_h[0] == INITIAL_VAL); + REQUIRE(result[0] == 0); + } else { + REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL); + REQUIRE(result[0] == INITIAL_VAL); + } HIP_CHECK(hipHostFree(A_h)); HIP_CHECK(hipHostFree(result)); } diff --git a/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withunsafeflag.cc b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withunsafeflag.cc index cbe614995f..d1b117d115 100644 --- a/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withunsafeflag.cc +++ b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withunsafeflag.cc @@ -90,8 +90,13 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentwithUnsafeflag", "", REQUIRE(testResult == true); } - REQUIRE(A_h[0] == INITIAL_VAL); - REQUIRE(result[0] == 0); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + REQUIRE(A_h[0] == INITIAL_VAL); + REQUIRE(result[0] == 0); + } else { + REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL); + REQUIRE(result[0] == INITIAL_VAL); + } HIP_CHECK(hipHostFree(A_h)); HIP_CHECK(hipHostFree(result)); } diff --git a/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_RTC.cc b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_RTC.cc index 6a96dbe32c..0639321c2e 100644 --- a/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_RTC.cc +++ b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_RTC.cc @@ -129,8 +129,13 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentRTCnounsafeatomicflag", "", HIP_CHECK(hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, config_d)); HIP_CHECK(hipDeviceSynchronize()); - REQUIRE(A_h[0] == INITIAL_VAL); - REQUIRE(*result == 0); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + REQUIRE(A_h[0] == INITIAL_VAL); + REQUIRE(*result == 0); + } else { + REQUIRE(A_h[0] == INITIAL_VAL + INCREMENT_VAL); + REQUIRE(*result == INITIAL_VAL); + } HIP_CHECK(hipHostFree(A_h)); HIP_CHECK(hipHostFree(result)); } @@ -221,8 +226,13 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentRTCunsafeatomicflag", "", HIP_CHECK(hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, config_d)); HIP_CHECK(hipDeviceSynchronize()); - REQUIRE(A_h[0] == INITIAL_VAL); - REQUIRE(*result == 0); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + REQUIRE(A_h[0] == INITIAL_VAL); + REQUIRE(*result == 0); + } else { + REQUIRE(A_h[0] == INITIAL_VAL + INCREMENT_VAL); + REQUIRE(*result == INITIAL_VAL); + } HIP_CHECK(hipHostFree(A_h)); HIP_CHECK(hipHostFree(result)); } @@ -309,8 +319,13 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentRTCwithoutflag", "", HIP_CHECK(hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, config_d)); HIP_CHECK(hipDeviceSynchronize()); - REQUIRE(A_h[0] == INITIAL_VAL); - REQUIRE(*result == 0); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + REQUIRE(A_h[0] == INITIAL_VAL); + REQUIRE(*result == 0); + } else { + REQUIRE(A_h[0] == INITIAL_VAL + INCREMENT_VAL); + REQUIRE(*result == INITIAL_VAL); + } HIP_CHECK(hipHostFree(A_h)); HIP_CHECK(hipHostFree(result)); }