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

SWDEV-476125 - Fix synchronization tests

Quellcode durchsuchen 
Fix issues of
    Unit_cache_coherency_cpu_gpu
    Unit_cache_coherency_gpu_gpu
Enable them for all devices.

Change-Id: I19ba2084fdddd7b173edddb4e9c1b16cf7a97314


[ROCm/hip-tests commit: 352932738a]
Dieser Commit ist enthalten in:
taosang2
2024-07-31 18:54:21 -04:00
committet von Rakesh Roy
Ursprung 51620a1bb6
Commit 2bfe34ad91
4 geänderte Dateien mit 119 neuen und 96 gelöschten Zeilen
Patch-Datei herunterladen Vergleichs-Datei herunterladen Alle Dateien ausklappen Alle Dateien einklappen
projects/hip-tests/catch/unit/synchronization/cache_coherency_gpu_gpu.cc
+71 -53
Datei anzeigen
@@ -21,8 +21,6 @@ THE SOFTWARE.
#include <hip_test_kernels.hh>
#include <hip_test_common.hh>
typedef _Atomic(unsigned int) atomic_uint;
// Helper function to spin on address until address equals value.
// If the address holds the value of -1, abort because the other thread failed.
__device__ int
@@ -32,10 +30,10 @@ gpu_spin_loop_or_abort_on_negative_one(unsigned int* address,
bool check = false;
do {
compare = value;
check = __opencl_atomic_compare_exchange_strong(
reinterpret_cast<atomic_uint*>(address), /*expected=*/ &compare,
check = __hip_atomic_compare_exchange_strong(
address, /*expected=*/ &compare,
/*desired=*/ value, __ATOMIC_ACQUIRE, __ATOMIC_ACQUIRE,
/*scope=*/ __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES);
/*scope=*/ __HIP_MEMORY_SCOPE_SYSTEM);
if (compare == -1)
return -1;
} while (!check);
@@ -51,8 +49,8 @@ gpu_cache0(int *A, int *B, int *X, int *Y, size_t N,
// Store data into A, system fence, and atomically mark flag.
// This guarantees this global write is visible by device 1.
A[i] = X[i];
__opencl_atomic_fetch_add(reinterpret_cast<atomic_uint*>(AA1), 1,
__ATOMIC_RELEASE, __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES);
__hip_atomic_fetch_add(AA1, 1,
__ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM);
// Wait on device 1's global write to B.
if (gpu_spin_loop_or_abort_on_negative_one(BA1, i+1) == -1) {
*cache0_result = -1;
@@ -65,13 +63,13 @@ gpu_cache0(int *A, int *B, int *X, int *Y, size_t N,
// If the data does not match, alert other thread and abort.
printf("FAIL: at i=%zu, B[i]=%d, which does not match Y[i]=%d.\n",
i, B[i], Y[i]);
__opencl_atomic_exchange(reinterpret_cast<atomic_uint*>(AA2), -1,
__ATOMIC_RELEASE, __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES);
__hip_atomic_exchange(AA2, -1,
__ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM);
*cache0_result = -1;
}
// Otherwise tell the other thread to continue.
__opencl_atomic_fetch_add(reinterpret_cast<atomic_uint*>(AA2), 1,
__ATOMIC_RELEASE, __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES);
__hip_atomic_fetch_add(AA2, 1,
__ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM);
// Wait on kernel gpu_cache1 to finish checking X is stored in A.
if (gpu_spin_loop_or_abort_on_negative_one(BA2, i+1) == -1) {
*cache0_result = -1;
@@ -88,8 +86,8 @@ gpu_cache1(int *A, int *B, int *X, int *Y, size_t N,
unsigned int *BA1, unsigned int *BA2, unsigned int *cache1_result) {
for (size_t i = 0; i < N; i++) {
B[i] = Y[i];
__opencl_atomic_fetch_add(reinterpret_cast<atomic_uint*>(BA1), 1,
__ATOMIC_RELEASE, __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES);
__hip_atomic_fetch_add(BA1, 1,
__ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM);
if (gpu_spin_loop_or_abort_on_negative_one(AA1, i+1) == -1) {
*cache1_result = -1;
break;
@@ -99,12 +97,12 @@ gpu_cache1(int *A, int *B, int *X, int *Y, size_t N,
if (!stored_data_matches) {
printf("FAIL: at i=%zu, A[i]=%d, which does not match X[i]=%d.\n",
i, A[i], X[i]);
__opencl_atomic_exchange(reinterpret_cast<atomic_uint*>(BA2), -1,
__ATOMIC_RELEASE, __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES);
__hip_atomic_exchange(BA2, -1,
__ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM);
*cache1_result = -1;
}
__opencl_atomic_fetch_add(reinterpret_cast<atomic_uint*>(BA2), 1,
__ATOMIC_RELEASE, __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES);
__hip_atomic_fetch_add(BA2, 1,
__ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM);
if (gpu_spin_loop_or_abort_on_negative_one(AA2, i+1) == -1) {
*cache1_result = -1;
break;
@@ -116,32 +114,54 @@ gpu_cache1(int *A, int *B, int *X, int *Y, size_t N,
static bool gpu_to_gpu_coherency() {
int *A_d, *B_d, *X_d0, *X_d1, *Y_d0, *Y_d1;
int *A_h, *B_h, *X_h, *Y_h;
unsigned int cache0_result, cache1_result;
unsigned int *cache0_result = nullptr;
unsigned int *cache1_result = nullptr;
size_t N = 1024;
size_t Nbytes = N * sizeof(int);
int numDevices = 0;
int numTestDevices = 2;
int deviceFineGrain = 0;
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices < numTestDevices) {
HipTest::HIP_SKIP_TEST("Skipping because devices < 2");
return 0;
}
// Skip this test if either device does not support this feature.
hipDeviceProp_t props0, props1;
HIP_CHECK(hipGetDeviceProperties(&props0, 0));
HIP_CHECK(hipGetDeviceProperties(&props1, 1));
if ((strncmp(props0.gcnArchName, "gfx90a", 6) != 0 ||
strncmp(props1.gcnArchName, "gfx90a", 6) != 0) &&
(strncmp(props0.gcnArchName, "gfx940", 6) != 0 ||
strncmp(props1.gcnArchName, "gfx940", 6) != 0)) {
printf("info: skipping test on devices other than gfx90a and gfx940.\n");
return true;
}
SECTION("With device fine grained buffer") {
HIP_CHECK(hipDeviceGetAttribute(&deviceFineGrain, hipDeviceAttributeFineGrainSupport, 0));
if (deviceFineGrain == 0) {
HipTest::HIP_SKIP_TEST("The test skipped due to deviceFineGrain = 0 on device 0");
return true;
}
HIP_CHECK(hipDeviceGetAttribute(&deviceFineGrain, hipDeviceAttributeFineGrainSupport, 1));
if (deviceFineGrain == 0) {
HipTest::HIP_SKIP_TEST("The test skipped due to deviceFineGrain = 0 on device 1");
return true;
}
HIP_CHECK(hipSetDevice(0));
HIP_CHECK(hipDeviceEnablePeerAccess(1, 0));
fprintf(stderr, "info: allocate device mem (%zu bytes) on device 0\n", Nbytes);
HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast<void**>(&A_d),
Nbytes, hipDeviceMallocFinegrained));
HIP_CHECK(hipSetDevice(1));
HIP_CHECK(hipDeviceEnablePeerAccess(0, 0));
fprintf(stderr, "info: allocate device mem (%zu bytes) on device 1\n", Nbytes);
HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast<void**>(&B_d),
Nbytes, hipDeviceMallocFinegrained));
}
SECTION("With host(SVM) fine grained buffer") {
HIP_CHECK(hipSetDevice(0));
HIP_CHECK(hipHostMalloc(&A_d, Nbytes));
HIP_CHECK(hipSetDevice(1));
HIP_CHECK(hipHostMalloc(&B_d, Nbytes));
}
HIP_CHECK(hipSetDevice(0));
HIP_CHECK(hipHostMalloc(&cache0_result, sizeof(unsigned int)));
HIP_CHECK(hipHostMalloc(&cache1_result, sizeof(unsigned int)));
*cache0_result = 0;
*cache1_result = 0;
// Allocate Host Side Memory.
printf("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
fprintf(stderr, "info: allocate host mem (%zu bytes)\n", Nbytes);
A_h = reinterpret_cast<int*>(malloc(Nbytes));
HIP_CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess);
B_h = reinterpret_cast<int*>(malloc(Nbytes));
@@ -183,27 +203,15 @@ static bool gpu_to_gpu_coherency() {
// Set-up Device 0.
HIP_CHECK(hipSetDevice(0));
// Enable P2P access to Device 1.
HIP_CHECK(hipDeviceEnablePeerAccess(1, 0));
HIP_CHECK(hipStreamCreateWithFlags(&stream[1], hipStreamNonBlocking));
// Allocating Coherent Memory for Array A_d on Device 0.
printf("info: allocate device 0 mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
hipError_t status = hipExtMallocWithFlags(reinterpret_cast<void**>(&A_d),
Nbytes, hipDeviceMallocFinegrained);
REQUIRE(status == hipSuccess);
HIP_CHECK(hipMalloc(&X_d0, Nbytes));
HIP_CHECK(hipMalloc(&Y_d0, Nbytes));
// Set-up Device 1.
HIP_CHECK(hipSetDevice(1));
// Enable P2P access to Device 0.
HIP_CHECK(hipDeviceEnablePeerAccess(0, 0));
HIP_CHECK(hipStreamCreateWithFlags(&stream[2], hipStreamNonBlocking));
// Allocating Coherent Memory for Array B_d on Device 1.
printf("info: allocate device 1 mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
status = hipExtMallocWithFlags(reinterpret_cast<void**>(&B_d),
Nbytes, hipDeviceMallocFinegrained);
REQUIRE(status == hipSuccess);
HIP_CHECK(hipMalloc(&X_d1, Nbytes));
HIP_CHECK(hipMalloc(&Y_d1, Nbytes));
@@ -220,21 +228,21 @@ static bool gpu_to_gpu_coherency() {
hipLaunchKernelGGL(gpu_cache0, dim3(blocks), dim3(threadsPerBlock),
0, stream[1],
A_d, B_d, X_d0, Y_d0, N,
AA1_d, AA2_d, BA1_d, BA2_d, &cache0_result);
AA1_d, AA2_d, BA1_d, BA2_d, cache0_result);
// Check if launch failed.
HIP_CHECK(hipGetLastError());
REQUIRE(cache0_result == 0);
HIP_CHECK(hipSetDevice(1));
hipLaunchKernelGGL(gpu_cache1, dim3(blocks), dim3(threadsPerBlock),
0, stream[2],
A_d, B_d, X_d1, Y_d1, N,
AA1_d, AA2_d, BA1_d, BA2_d, &cache1_result);
AA1_d, AA2_d, BA1_d, BA2_d, cache1_result);
HIP_CHECK(hipGetLastError());
REQUIRE(cache1_result == 0);
// Wait for kernels on both devices.
HIP_CHECK(hipStreamSynchronize(stream[1]));
HIP_CHECK(hipStreamSynchronize(stream[2]));
REQUIRE(*cache0_result == 0);
REQUIRE(*cache1_result == 0);
// Evaluate the resultant arrays A and B.
HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
@@ -246,8 +254,13 @@ static bool gpu_to_gpu_coherency() {
}
// Free all the device and host memory allocated.
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(B_d));
if(deviceFineGrain) {
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(B_d));
} else {
HIP_CHECK(hipHostFree(A_d));
HIP_CHECK(hipHostFree(B_d));
}
HIP_CHECK(hipFree(X_d0));
HIP_CHECK(hipFree(Y_d0));
HIP_CHECK(hipFree(X_d1));
@@ -256,11 +269,16 @@ static bool gpu_to_gpu_coherency() {
HIP_CHECK(hipHostFree(AA2_h));
HIP_CHECK(hipHostFree(BA1_h));
HIP_CHECK(hipHostFree(BA2_h));
HIP_CHECK(hipHostFree(cache0_result));
HIP_CHECK(hipHostFree(cache1_result));
free(A_h);
free(B_h);
free(X_h);
free(Y_h);
for (int i = 0; i < 3; i++) {
HIP_CHECK(hipStreamDestroy(stream[i]));
}
return true;
}