SWDEV-492165 - Add new test types for atomicCAS and unsafeAtomicAdd

Change-Id: Icacb6e94c64c5d4ce3f15c6d4a50fe8506cedfa9

catch/unit/atomics/atomicCAS.cc

@@ -62,7 +62,8 @@ THE SOFTWARE.
  * ------------------------
  *    - HIP_VERSION >= 5.2
  */
-TEMPLATE_TEST_CASE("Unit_atomicCAS_Positive", "", int, unsigned int, unsigned long long TYPES) {
+TEMPLATE_TEST_CASE("Unit_atomicCAS_Positive", "", int, unsigned int, unsigned long long,
+                                                  unsigned short int TYPES) {
   int warp_size = 0;
   HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
   const auto cache_line_size = 128u;
@@ -107,7 +108,7 @@ TEMPLATE_TEST_CASE("Unit_atomicCAS_Positive", "", int, unsigned int, unsigned lo
  *    - HIP_VERSION >= 5.2
  */
 TEMPLATE_TEST_CASE("Unit_atomicCAS_Positive_Multi_Kernel", "", int, unsigned int,
-                   unsigned long long TYPES) {
+                   unsigned long long, unsigned short int TYPES) {
   int warp_size = 0;
   HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
   const auto cache_line_size = 128u;

catch/unit/atomics/atomicCAS_system.cc

@@ -61,7 +61,7 @@ THE SOFTWARE.
  *    - HIP_VERSION >= 5.2
  */
 TEMPLATE_TEST_CASE("Unit_atomicCAS_system_Positive_Peer_GPUs", "", int, unsigned int,
-                   unsigned long long TYPES) {
+                   unsigned long long, unsigned short int TYPES) {
   int warp_size = 0;
   HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
   const auto cache_line_size = 128u;
@@ -111,7 +111,7 @@ TEMPLATE_TEST_CASE("Unit_atomicCAS_system_Positive_Peer_GPUs", "", int, unsigned
  *    - HIP_VERSION >= 5.2
  */
 TEMPLATE_TEST_CASE("Unit_atomicCAS_system_Positive_Host_And_GPU", "", int, unsigned int,
-                   unsigned long long TYPES) {
+                   unsigned long long, unsigned short int TYPES) {
   int warp_size = 0;
   HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
   const auto cache_line_size = 128u;
@@ -161,7 +161,7 @@ TEMPLATE_TEST_CASE("Unit_atomicCAS_system_Positive_Host_And_GPU", "", int, unsig
  *    - HIP_VERSION >= 5.2
  */
 TEMPLATE_TEST_CASE("Unit_atomicCAS_system_Positive_Host_And_Peer_GPUs", "", int, unsigned int,
-                   unsigned long long TYPES) {
+                   unsigned long long, unsigned short int TYPES) {
   int warp_size = 0;
   HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
   const auto cache_line_size = 128u;

catch/unit/atomics/unsafeAtomicAdd.cc

@@ -128,19 +128,25 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_Positive_Multi_Kernel", "", float, doub
 
 template <typename Type,
           std::enable_if_t<std::is_same<Type, __half2>::value ||
-                               std::is_same<Type, __hip_bfloat162>::value,
+                               std::is_same<Type, __hip_bfloat162>::value ||
+                               std::is_same<Type, __hip_bfloat16>::value ||
+                               std::is_same<Type, __half>::value,
                            bool> = true>
 __global__ void unsafe_add_kernel(Type* ptr, Type val) {
   (void)unsafeAtomicAdd(ptr, val);
 }
 
-TEMPLATE_TEST_CASE("Unit_unsafe_atomic_add_half_and_bfloat", "", __half2, __hip_bfloat162) {
+TEMPLATE_TEST_CASE("Unit_unsafe_atomic_add_half_and_bfloat", "", __half2, __hip_bfloat162, __half, __hip_bfloat16) {
   auto kernel = unsafe_add_kernel<TestType>;
   TestType val;
   if constexpr (std::is_same<TestType, __half2>::value) {
     val = __float22half2_rn(float2{1.0f, 2.0f});
-  } else {
+  } else if constexpr (std::is_same<TestType, __hip_bfloat162>::value) {
     val = __float22bfloat162_rn(float2{1.0f, 2.0f});
+  } else if constexpr (std::is_same<TestType, __half>::value) {
+    val = __float2half(float{2.0f});
+  } else {
+    val = __float2bfloat16(float{2.0f});
   }
 
   TestType* out;
@@ -150,11 +156,18 @@ TEMPLATE_TEST_CASE("Unit_unsafe_atomic_add_half_and_bfloat", "", __half2, __hip_
 
   TestType dout;
   HIP_CHECK(hipMemcpy(&dout, out, sizeof(TestType), hipMemcpyDeviceToHost));
+
   float2 hout;
   if constexpr (std::is_same<TestType, __half2>::value) {
     hout = __half22float2(dout);
-  } else {
+  } else if constexpr (std::is_same<TestType, __hip_bfloat162>::value) {
     hout = __bfloat1622float2(dout);
+  } else if constexpr (std::is_same<TestType, __half>::value) {
+    hout.x = 32.0f;
+    hout.y  = __half2float(dout);
+  } else {
+    hout.x = 32.0f;
+    hout.y = __bfloat162float(dout);
   }
 
   REQUIRE(hout.x == 32.0f);