SWDEV-420149 - Fixing unsafe atomics change for targets other than gfx90a.

Change-Id: Ifa26fd1c51505c5e86e121233c8352e3bde846ea


[ROCm/hip-tests commit: a363fa2833]

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

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

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

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

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

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