diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux index b209f608a7..f310749271 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux @@ -1174,8 +1174,6 @@ "Unit_dynamic_loading_device_kernels_from_library", "Unit_hipApiDynamicLoad", "Unit_hipMalloc_gpptest", - "Unit_cache_coherency_cpu_gpu", - "Unit_cache_coherency_gpu_gpu", "=== SWDEV-438556:Below tests failed in stress test on 15/12/23 ===", "=== SWDEV-439298: Below test failing in CQE staging ===", "Unit_hipCGMultiGridGroupType_Barrier", diff --git a/projects/hip-tests/catch/unit/synchronization/CMakeLists.txt b/projects/hip-tests/catch/unit/synchronization/CMakeLists.txt index 50018e5c43..6d6b55e0ef 100644 --- a/projects/hip-tests/catch/unit/synchronization/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/synchronization/CMakeLists.txt @@ -10,12 +10,15 @@ add_custom_target(memcpyInt.hsaco COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOA ${HIP_PATH}/${CMAKE_INSTALL_LIBDIR}/../../include --rocm-path=${ROCM_PATH}) # only for AMD if(HIP_PLATFORM MATCHES "amd") +# There are problems in Windows: __hip_atomicsXXX() won't work as expected +if(NOT WIN32) set(AMD_SRC cache_coherency_cpu_gpu.cc cache_coherency_gpu_gpu.cc ) set(TEST_SRC ${TEST_SRC} ${AMD_SRC}) endif() +endif() hip_add_exe_to_target(NAME synchronizationTests TEST_SRC ${TEST_SRC} diff --git a/projects/hip-tests/catch/unit/synchronization/cache_coherency_cpu_gpu.cc b/projects/hip-tests/catch/unit/synchronization/cache_coherency_cpu_gpu.cc index 0d4c9c0136..d458864dbd 100644 --- a/projects/hip-tests/catch/unit/synchronization/cache_coherency_cpu_gpu.cc +++ b/projects/hip-tests/catch/unit/synchronization/cache_coherency_cpu_gpu.cc @@ -21,8 +21,6 @@ THE SOFTWARE. #include #include -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(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_kernel(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(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) { *dresult = -1; @@ -65,13 +63,13 @@ gpu_kernel(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(AA2), -1, - __ATOMIC_RELEASE, __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES); + __hip_atomic_exchange(AA2, -1, + __ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM); *dresult = -1; } // Otherwise tell the other thread to continue. - __opencl_atomic_fetch_add(reinterpret_cast(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) { *dresult = -1; @@ -130,29 +128,38 @@ cpu_thread(int *A, int *B, int *X, int *Y, size_t N, static bool cpu_to_gpu_coherency() { int *A_d, *B_d, *X_d, *Y_d; int *A_res, *A_h, *B_h, *X_h, *Y_h; - unsigned int hresult, dresult; + unsigned int hresult = 0; + unsigned int *dresult = nullptr; size_t N = 1024; size_t Nbytes = N * sizeof(int); int numDevices = 0; + int deviceFineGrain = 0; HIP_CHECK(hipGetDeviceCount(&numDevices)); if (numDevices < 1) { HipTest::HIP_SKIP_TEST("Skipping because devices < 1"); - return 0; - } - - // Skip this test if feature is not supported. - static int device0 = 0; - hipDeviceProp_t props; - HIP_CHECK(hipGetDeviceProperties(&props, device0)); - if (strncmp(props.gcnArchName, "gfx90a", 6) != 0 && - strncmp(props.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"); + return true; + } + fprintf(stderr, "info: allocate device mem (%zu bytes) on device 0\n", Nbytes); + HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast(&A_d), + Nbytes, hipDeviceMallocFinegrained)); + } + SECTION("With host(SVM) fine grained buffer") { + HIP_CHECK(hipHostMalloc(&A_d, Nbytes)); + } + A_h = A_d; + + HIP_CHECK(hipHostMalloc(&dresult, sizeof(unsigned int))); + *dresult = 0; // Allocate Host Side Memory. Coherent Fine-grained Memory for array B. - 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); HIP_CHECK(hipHostMalloc(&B_h, Nbytes, (hipHostMallocCoherent | hipHostMallocMapped))); HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&B_d), B_h, 0)); @@ -188,18 +195,10 @@ static bool cpu_to_gpu_coherency() { *BA2_h = 0; // Skip the first stream, ensure stream is non-blocking. - hipStream_t stream[2]; - HIP_CHECK(hipStreamCreate(&stream[0])); + hipStream_t stream; HIP_CHECK(hipSetDevice(0)); - HIP_CHECK(hipStreamCreateWithFlags(&stream[1], hipStreamNonBlocking)); + HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); - // Allocate Device Side Memory. Coherent Fine-grained Memory for array A. - printf("info: allocate device 0 mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0); - hipError_t status = hipExtMallocWithFlags(reinterpret_cast(&A_d), - Nbytes, hipDeviceMallocFinegrained); - REQUIRE(status == hipSuccess); - // SVM memory - host pointer is the same as device pointer to array A. - A_h = A_d; HIP_CHECK(hipMalloc(&X_d, Nbytes)); HIP_CHECK(hipMalloc(&Y_d, Nbytes)); @@ -210,21 +209,21 @@ static bool cpu_to_gpu_coherency() { const unsigned blocks = 1; const unsigned threadsPerBlock = 1; hipLaunchKernelGGL(gpu_kernel, dim3(blocks), dim3(threadsPerBlock), - 0, stream[1], + 0, stream, A_d, B_d, X_d, Y_d, N, - AA1_d, AA2_d, BA1_d, BA2_d, &dresult); + AA1_d, AA2_d, BA1_d, BA2_d, dresult); // Check if launch failed. HIP_CHECK(hipGetLastError()); - REQUIRE(dresult == 0); // Do not sync the launched stream, instead run the cpu_thread. std::thread host_thread(cpu_thread, A_h, B_h, X_h, Y_h, N, AA1_h, AA2_h, BA1_h, BA2_h, &hresult); - host_thread.detach(); - REQUIRE(hresult == 0); // Wait for Device side to finish. - HIP_CHECK(hipStreamSynchronize(stream[1])); + HIP_CHECK(hipStreamSynchronize(stream)); + host_thread.join(); + REQUIRE(*dresult == 0); + REQUIRE(hresult == 0); // Evaluate the resultant arrays A and B. A_res = reinterpret_cast(malloc(Nbytes)); @@ -237,7 +236,11 @@ static bool cpu_to_gpu_coherency() { } // Free all the device and host memory allocated. - HIP_CHECK(hipFree(A_d)); + if (deviceFineGrain) { + HIP_CHECK(hipFree(A_d)); + } else { + HIP_CHECK(hipHostFree(A_d)); + } HIP_CHECK(hipFree(X_d)); HIP_CHECK(hipFree(Y_d)); HIP_CHECK(hipHostFree(AA1_h)); @@ -245,10 +248,11 @@ static bool cpu_to_gpu_coherency() { HIP_CHECK(hipHostFree(BA1_h)); HIP_CHECK(hipHostFree(BA2_h)); HIP_CHECK(hipHostFree(B_h)); + HIP_CHECK(hipHostFree(dresult)); free(X_h); free(Y_h); free(A_res); - + HIP_CHECK(hipStreamDestroy(stream)); return true; } diff --git a/projects/hip-tests/catch/unit/synchronization/cache_coherency_gpu_gpu.cc b/projects/hip-tests/catch/unit/synchronization/cache_coherency_gpu_gpu.cc index 42df8266ad..4eac137232 100644 --- a/projects/hip-tests/catch/unit/synchronization/cache_coherency_gpu_gpu.cc +++ b/projects/hip-tests/catch/unit/synchronization/cache_coherency_gpu_gpu.cc @@ -21,8 +21,6 @@ THE SOFTWARE. #include #include -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(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(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(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(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(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(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(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(&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(&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(malloc(Nbytes)); HIP_CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess); B_h = reinterpret_cast(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(&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(&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; }