diff --git a/projects/hip-tests/catch/unit/atomics/atomicCAS.cc b/projects/hip-tests/catch/unit/atomics/atomicCAS.cc index f89779b13b..bdeb555dde 100644 --- a/projects/hip-tests/catch/unit/atomics/atomicCAS.cc +++ b/projects/hip-tests/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; diff --git a/projects/hip-tests/catch/unit/atomics/atomicCAS_system.cc b/projects/hip-tests/catch/unit/atomics/atomicCAS_system.cc index 2b39d04415..0014cea46a 100644 --- a/projects/hip-tests/catch/unit/atomics/atomicCAS_system.cc +++ b/projects/hip-tests/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; diff --git a/projects/hip-tests/catch/unit/atomics/unsafeAtomicAdd.cc b/projects/hip-tests/catch/unit/atomics/unsafeAtomicAdd.cc index 9c0989e7e3..1fde5861c6 100644 --- a/projects/hip-tests/catch/unit/atomics/unsafeAtomicAdd.cc +++ b/projects/hip-tests/catch/unit/atomics/unsafeAtomicAdd.cc @@ -128,19 +128,25 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_Positive_Multi_Kernel", "", float, doub template ::value || - std::is_same::value, + std::is_same::value || + std::is_same::value || + std::is_same::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 val; if constexpr (std::is_same::value) { val = __float22half2_rn(float2{1.0f, 2.0f}); - } else { + } else if constexpr (std::is_same::value) { val = __float22bfloat162_rn(float2{1.0f, 2.0f}); + } else if constexpr (std::is_same::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::value) { hout = __half22float2(dout); - } else { + } else if constexpr (std::is_same::value) { hout = __bfloat1622float2(dout); + } else if constexpr (std::is_same::value) { + hout.x = 32.0f; + hout.y = __half2float(dout); + } else { + hout.x = 32.0f; + hout.y = __bfloat162float(dout); } REQUIRE(hout.x == 32.0f);