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)); }