AI/rocm-systems
2
0
Fork 0
Code Issues Pull-Requests Actions Pakete Projekte Releases Wiki Aktivität

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

Quellcode durchsuchen 
Change-Id: Icacb6e94c64c5d4ce3f15c6d4a50fe8506cedfa9


[ROCm/hip-tests commit: c7236e3e2e]
Dieser Commit ist enthalten in:
Rahul Manocha
2024-12-20 12:32:04 -08:00
committet von Rakesh Roy
Ursprung b7f9ee2330
Commit 6cde56faa8
3 geänderte Dateien mit 23 neuen und 9 gelöschten Zeilen
Patch-Datei herunterladen Vergleichs-Datei herunterladen Alle Dateien ausklappen Alle Dateien einklappen
projects/hip-tests/catch/unit/atomics/atomicCAS.cc
+3 -2
Datei anzeigen
@@ -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;
projects/hip-tests/catch/unit/atomics/atomicCAS_system.cc
+3 -3
Datei anzeigen
@@ -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;
projects/hip-tests/catch/unit/atomics/unsafeAtomicAdd.cc
+17 -4
Datei anzeigen
@@ -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);