AI/rocm-systems
2
0
Forka 0
Codice Problemi Pull Request Actions Pacchetti Progetti Rilasci Wiki Attività

SWDEV-346579 - atomics test cases for non MI2xx/MI3xx (#3060)

Sfoglia il codice sorgente 
Replace the internal __builtin__ functions with external APIs

Change-Id: Ia3cf37af004b5b9a4833b18a8771a8864772beb4


[ROCm/hip-tests commit: 6a40a263a2]
This commit is contained in:
ROCm CI Service Account
2022-11-17 03:31:05 +05:30
committato da GitHub
parent 29f3e73eb4
commit ab58ccf21a
3 ha cambiato i file con 11 aggiunte e 43 eliminazioni
Scarica il file Patch Scarica il file Diff Expand all files Collapse all files
projects/hip-tests/catch/unit/deviceLib/BuiltIns_fadd.cc
+2 -2
Vedi File
@@ -34,14 +34,14 @@ This testfile verifies __builtin_amdgcn_global_atomic_fadd_f64 API scenarios
#define INITIAL_VAL 5
__global__ void AtomicAdd_GlobalMem(double* addr, double* result) {
double inc_val = 10;
*result = __builtin_amdgcn_global_atomic_fadd_f64(addr, inc_val);
*result = unsafeAtomicAdd(addr, inc_val);
}
static constexpr auto AtomicAddGlobalMem{
R"(
extern "C"
__global__ void AtomicAdd_GlobalMem(double* addr, double* result) {
double inc_val = 10;
*result = __builtin_amdgcn_global_atomic_fadd_f64(addr, inc_val);
*result = unsafeAtomicAdd(addr, inc_val);
}
)"};
/*
projects/hip-tests/catch/unit/deviceLib/BuiltIns_fmax.cc
+4 -20
Vedi File
@@ -39,20 +39,12 @@ __global__ void unsafeAtomicMax_FlatMem(double* addr, double* result) {
__shared__ double int_val;
int_val = 5;
double comp = 10;
if (__builtin_amdgcn_is_shared(
(const __attribute__((address_space(0))) void*)(&int_val)))
*result = __builtin_amdgcn_flat_atomic_fmax_f64(&int_val, comp);
else
*result = __builtin_amdgcn_global_atomic_fmax_f64(&int_val, comp);
*result = unsafeAtomicMax(&int_val, comp);
*addr = int_val;
}
__global__ void unsafeAtomicMax_GlobalMem(double* addr, double* result) {
double comp = 10;
if (__builtin_amdgcn_is_shared(
(const __attribute__((address_space(0))) void*)(addr)))
*result = __builtin_amdgcn_flat_atomic_fmax_f64(addr, comp);
else
*result = __builtin_amdgcn_global_atomic_fmax_f64(addr, comp);
*result = unsafeAtomicMax(addr, comp);
}
static constexpr auto fmaxFlatMem {
R"(
@@ -61,11 +53,7 @@ __global__ void unsafeAtomicMax_FlatMem(double* addr, double* result) {
__shared__ double int_val;
int_val = 5;
double comp = 10;
if (__builtin_amdgcn_is_shared(
(const __attribute__((address_space(0))) void*)(&int_val)))
*result = __builtin_amdgcn_flat_atomic_fmax_f64(&int_val, comp);
else
*result = __builtin_amdgcn_global_atomic_fmax_f64(&int_val, comp);
*result = unsafeAtomicMax(&int_val, comp);
*addr = int_val;
}
)"};
@@ -75,11 +63,7 @@ R"(
extern "C"
__global__ void unsafeAtomicMax_GlobalMem(double* addr, double* result) {
double comp = 10;
if (__builtin_amdgcn_is_shared(
(const __attribute__((address_space(0))) void*)(addr)))
*result = __builtin_amdgcn_flat_atomic_fmax_f64(addr, comp);
else
*result = __builtin_amdgcn_global_atomic_fmax_f64(addr, comp);
*result = unsafeAtomicMax(addr, comp);
}
)"};
projects/hip-tests/catch/unit/deviceLib/BuiltIns_fmin.cc
+5 -21
Vedi File
@@ -41,12 +41,8 @@ __global__ void unsafeAtomicMin_FlatMem(double* addr, double* result) {
__shared__ double int_val;
int_val = 5;
double comp = 10;
if (__builtin_amdgcn_is_shared(
(const __attribute__((address_space(0))) void*)(&int_val)))
*result = __builtin_amdgcn_flat_atomic_fmin_f64(&int_val, comp);
else
*result = __builtin_amdgcn_global_atomic_fmin_f64(&int_val, comp);
*addr = int_val;
*result = unsafeAtomicMin(&int_val, comp);
*addr = int_val;
}
)"};
@@ -55,11 +51,7 @@ R"(
extern "C"
__global__ void unsafeAtomicMin_GlobalMem(double* addr, double* result) {
double comp = 10;
if (__builtin_amdgcn_is_shared(
(const __attribute__((address_space(0))) void*)(addr)))
*result = __builtin_amdgcn_flat_atomic_fmin_f64(addr, comp);
else
*result = __builtin_amdgcn_global_atomic_fmin_f64(addr, comp);
*result = unsafeAtomicMin(addr, comp);
}
)"};
@@ -67,20 +59,12 @@ __global__ void unsafeAtomicMin_FlatMem(double* addr, double* result) {
__shared__ double int_val;
int_val = 5;
double comp = 10;
if (__builtin_amdgcn_is_shared(
(const __attribute__((address_space(0))) void*)(&int_val)))
*result = __builtin_amdgcn_flat_atomic_fmin_f64(&int_val, comp);
else
*result = __builtin_amdgcn_global_atomic_fmin_f64(&int_val, comp);
*result = unsafeAtomicMin(&int_val, comp);
*addr = int_val;
}
__global__ void unsafeAtomicMin_GlobalMem(double* addr, double* result) {
double comp = 10;
if (__builtin_amdgcn_is_shared(
(const __attribute__((address_space(0))) void*)(addr)))
*result = __builtin_amdgcn_flat_atomic_fmin_f64(addr, comp);
else
*result = __builtin_amdgcn_global_atomic_fmin_f64(addr, comp);
*result = unsafeAtomicMin(addr, comp);
}
/*