AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-470698 - fix formatting, add format check workflow (#657)

Browse Source
This commit is contained in:
Danylo Lytovchenko
2025-08-20 16:28:06 +02:00
committed by GitHub
parent 5840940caa
commit f7338717ae
1574 changed files with 162972 additions and 199346 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/unit/synchronization/cache_coherency_gpu_gpu.cc
+40 -61
View File
@@ -23,36 +23,31 @@ THE SOFTWARE.
// 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
gpu_spin_loop_or_abort_on_negative_one(unsigned int* address,
unsigned int value) {
__device__ int gpu_spin_loop_or_abort_on_negative_one(unsigned int* address, unsigned int value) {
unsigned int compare;
bool check = false;
do {
compare = value;
check = __hip_atomic_compare_exchange_strong(
address, /*expected=*/ &compare,
/*desired=*/ value, __ATOMIC_ACQUIRE, __ATOMIC_ACQUIRE,
/*scope=*/ __HIP_MEMORY_SCOPE_SYSTEM);
if (compare == -1)
return -1;
check =
__hip_atomic_compare_exchange_strong(address, /*expected=*/&compare,
/*desired=*/value, __ATOMIC_ACQUIRE, __ATOMIC_ACQUIRE,
/*scope=*/__HIP_MEMORY_SCOPE_SYSTEM);
if (compare == -1) return -1;
} while (!check);
return 0;
}
// This kernel requires a single block, single thread dispatch.
__global__ void
gpu_cache0(int *A, int *B, int *X, int *Y, size_t N,
unsigned int *AA1, unsigned int *AA2,
unsigned int *BA1, unsigned int *BA2, unsigned int *cache0_result) {
__global__ void gpu_cache0(int* A, int* B, int* X, int* Y, size_t N, unsigned int* AA1,
unsigned int* AA2, unsigned int* BA1, unsigned int* BA2,
unsigned int* cache0_result) {
for (size_t i = 0; i < N; i++) {
// Store data into A, system fence, and atomically mark flag.
// This guarantees this global write is visible by device 1.
A[i] = X[i];
__hip_atomic_fetch_add(AA1, 1,
__ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM);
__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) {
if (gpu_spin_loop_or_abort_on_negative_one(BA1, i + 1) == -1) {
*cache0_result = -1;
break;
}
@@ -61,17 +56,14 @@ gpu_cache0(int *A, int *B, int *X, int *Y, size_t N,
bool stored_data_matches = (B[i] == Y[i]);
if (!stored_data_matches) {
// 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]);
__hip_atomic_exchange(AA2, -1,
__ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM);
printf("FAIL: at i=%zu, B[i]=%d, which does not match Y[i]=%d.\n", i, B[i], Y[i]);
__hip_atomic_exchange(AA2, -1, __ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM);
*cache0_result = -1;
}
// Otherwise tell the other thread to continue.
__hip_atomic_fetch_add(AA2, 1,
__ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM);
__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) {
if (gpu_spin_loop_or_abort_on_negative_one(BA2, i + 1) == -1) {
*cache0_result = -1;
break;
}
@@ -80,30 +72,25 @@ gpu_cache0(int *A, int *B, int *X, int *Y, size_t N,
}
// This kernel requires a single block, single thread dispatch.
__global__ void
gpu_cache1(int *A, int *B, int *X, int *Y, size_t N,
unsigned int *AA1, unsigned int *AA2,
unsigned int *BA1, unsigned int *BA2, unsigned int *cache1_result) {
__global__ void gpu_cache1(int* A, int* B, int* X, int* Y, size_t N, unsigned int* AA1,
unsigned int* AA2, unsigned int* BA1, unsigned int* BA2,
unsigned int* cache1_result) {
for (size_t i = 0; i < N; i++) {
B[i] = Y[i];
__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) {
__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;
}
bool stored_data_matches = (A[i] == X[i]);
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]);
__hip_atomic_exchange(BA2, -1,
__ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM);
printf("FAIL: at i=%zu, A[i]=%d, which does not match X[i]=%d.\n", i, A[i], X[i]);
__hip_atomic_exchange(BA2, -1, __ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM);
*cache1_result = -1;
}
__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) {
__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;
}
@@ -114,8 +101,8 @@ 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 = nullptr;
unsigned int *cache1_result = nullptr;
unsigned int* cache0_result = nullptr;
unsigned int* cache1_result = nullptr;
size_t N = 1024;
size_t Nbytes = N * sizeof(int);
int numDevices = 0;
@@ -141,13 +128,13 @@ static bool gpu_to_gpu_coherency() {
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(
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));
HIP_CHECK(
hipExtMallocWithFlags(reinterpret_cast<void**>(&B_d), Nbytes, hipDeviceMallocFinegrained));
}
SECTION("With host(SVM) fine grained buffer") {
HIP_CHECK(hipSetDevice(0));
@@ -181,20 +168,16 @@ static bool gpu_to_gpu_coherency() {
unsigned int *AA1_h, *AA2_h, *BA1_h, *BA2_h;
unsigned int *AA1_d, *AA2_d, *BA1_d, *BA2_d;
HIP_CHECK(hipHostMalloc(&AA1_h, sizeof(unsigned int), hipHostMallocCoherent));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&AA1_d),
AA1_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&AA1_d), AA1_h, 0));
*AA1_h = 0;
HIP_CHECK(hipHostMalloc(&AA2_h, sizeof(unsigned int), hipHostMallocCoherent));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&AA2_d),
AA2_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&AA2_d), AA2_h, 0));
*AA2_h = 0;
HIP_CHECK(hipHostMalloc(&BA1_h, sizeof(unsigned int), hipHostMallocCoherent));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&BA1_d),
BA1_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&BA1_d), BA1_h, 0));
*BA1_h = 0;
HIP_CHECK(hipHostMalloc(&BA2_h, sizeof(unsigned int), hipHostMallocCoherent));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&BA2_d),
BA2_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&BA2_d), BA2_h, 0));
*BA2_h = 0;
// Skip the first stream.
@@ -225,17 +208,13 @@ static bool gpu_to_gpu_coherency() {
const unsigned blocks = 1;
const unsigned threadsPerBlock = 1;
HIP_CHECK(hipSetDevice(0));
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);
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);
// Check if launch failed.
HIP_CHECK(hipGetLastError());
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);
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);
HIP_CHECK(hipGetLastError());
// Wait for kernels on both devices.
@@ -248,13 +227,13 @@ static bool gpu_to_gpu_coherency() {
HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(B_h, B_d, Nbytes, hipMemcpyDeviceToHost));
for (size_t i = 0; i < N; i++) {
for (size_t i = 0; i < N; i++) {
REQUIRE(A_h[i] == (100000000 + i));
REQUIRE(B_h[i] == (300000000 + i));
}
// Free all the device and host memory allocated.
if(deviceFineGrain) {
if (deviceFineGrain) {
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(B_d));
} else {