From ec8ff45a1d9e3ccbab74a58feae84c91668ed982 Mon Sep 17 00:00:00 2001 From: "Sang, Tao" Date: Wed, 11 Jun 2025 11:41:25 -0400 Subject: [PATCH] SWDEV-514141 - Fix zero clock rate issues (#4) 1.Remove clock functions from some tests that don't need them. 2.In some memory pool tests and coherency tests, timer-based kernel delay isn't reliable, use pinned host based notification instead. 3.Add CHECK_PCIE_ATOMICS_SUPPORT before some tests. 4.catch/unit/memory/hipMemoryAllocateCoherent.cc is removed as it is useless and originally excluded in building. 5.Some tests can still pass even if clock rate =0, thus they will be kept as is. 6.Some logic and format improvement in some tests. Change-Id: I6b3c6bf54c61cffd45cd6f17c75998f751b75725 --- catch/include/hip_test_common.hh | 16 ++ catch/multiproc/CMakeLists.txt | 40 ++-- catch/multiproc/hipMemCoherencyTstMProc.cc | 173 ++++---------- catch/perftests/graph/parallelGraph.cc | 53 +---- .../stream/hipPerfDeviceConcurrency.cc | 4 + .../stream/hipPerfStreamConcurrency.cc | 5 +- catch/unit/clock/hipClockCheck.cc | 12 +- catch/unit/memory/CMakeLists.txt | 1 + catch/unit/memory/hipMemPoolApi.cc | 218 +++++++----------- .../unit/memory/hipMemPoolSetGetAttribute.cc | 158 +++++-------- catch/unit/memory/hipMemPoolTrimTo.cc | 17 +- .../unit/memory/hipMemoryAllocateCoherent.cc | 85 ------- catch/unit/memory/hipSVMCommon.h | 6 +- catch/unit/memory/mempool_common.hh | 65 +++--- .../hipStreamPerThrdCompilerOptn.cc | 60 +++-- 15 files changed, 332 insertions(+), 581 deletions(-) delete mode 100644 catch/unit/memory/hipMemoryAllocateCoherent.cc diff --git a/catch/include/hip_test_common.hh b/catch/include/hip_test_common.hh index 86fc4d8836..d2cd7ec4e1 100644 --- a/catch/include/hip_test_common.hh +++ b/catch/include/hip_test_common.hh @@ -314,6 +314,15 @@ inline bool isImageSupported() { return imageSupport != 0; } +inline bool isPcieAtomicsSupported() { + int pcieAtomics = 1; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIPCHECK(hipDeviceGetAttribute(&pcieAtomics, hipDeviceAttributeHostNativeAtomicSupported, + device)); + return pcieAtomics != 0; +} + inline bool areWarpMatchFunctionsSupported() { int matchFunctionsSupported = 1; #if HT_NVIDIA @@ -500,6 +509,13 @@ class BlockingContext { return; \ } +// This must be called in host-device memory conherency tests +#define CHECK_PCIE_ATOMICS_SUPPORT \ + if (!HipTest::isPcieAtomicsSupported()) { \ + INFO("Pcie atomics is not support on the device. Skipped."); \ + return; \ + } + // This must be called in the beginning of warp test app's main() to indicate warp match functions // are supported. #define CHECK_WARP_MATCH_FUNCTIONS_SUPPORT \ diff --git a/catch/multiproc/CMakeLists.txt b/catch/multiproc/CMakeLists.txt index 629f37372c..688e74a6ba 100644 --- a/catch/multiproc/CMakeLists.txt +++ b/catch/multiproc/CMakeLists.txt @@ -1,24 +1,24 @@ # Common Tests set(TEST_SRC - childMalloc.cc - hipDeviceComputeCapabilityMproc.cc - hipDeviceGetPCIBusIdMproc.cc - hipDeviceTotalMemMproc.cc - hipGetDeviceAttributeMproc.cc - hipGetDeviceCountMproc.cc - hipGetDevicePropertiesMproc.cc - hipSetGetDeviceMproc.cc - hipIpcMemAccessTest.cc - hipMallocConcurrencyMproc.cc - hipMemCoherencyTstMProc.cc - hipIpcEventHandle.cc - deviceAllocationMproc.cc - hipNoGpuTsts.cc - hipMemGetInfoMProc.cc + childMalloc.cc + hipDeviceComputeCapabilityMproc.cc + hipDeviceGetPCIBusIdMproc.cc + hipDeviceTotalMemMproc.cc + hipGetDeviceAttributeMproc.cc + hipGetDeviceCountMproc.cc + hipGetDevicePropertiesMproc.cc + hipSetGetDeviceMproc.cc + hipIpcMemAccessTest.cc + hipMallocConcurrencyMproc.cc + hipMemCoherencyTstMProc.cc + hipIpcEventHandle.cc + deviceAllocationMproc.cc + hipNoGpuTsts.cc + hipMemGetInfoMProc.cc ) if(UNIX) - add_custom_target(dummy_kernel.code + add_custom_target(dummy_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/dummy_kernel.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../multiproc/dummy_kernel.code @@ -30,18 +30,18 @@ endif() # the last argument linker libraries is required for this test but optional to the function if(HIP_PLATFORM MATCHES "nvidia") -hip_add_exe_to_target(NAME MultiProc + hip_add_exe_to_target(NAME MultiProc TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests LINKER_LIBS nvrtc) + set_target_properties(MultiProc PROPERTIES COMPILE_FLAGS -arch=sm_70) elseif(HIP_PLATFORM MATCHES "amd") -hip_add_exe_to_target(NAME MultiProc + hip_add_exe_to_target(NAME MultiProc TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests LINKER_LIBS hiprtc) endif() if(UNIX) - add_dependencies(build_tests dummy_kernel.code) + add_dependencies(build_tests dummy_kernel.code) endif() - diff --git a/catch/multiproc/hipMemCoherencyTstMProc.cc b/catch/multiproc/hipMemCoherencyTstMProc.cc index 0de5b75cef..f91b8b15f6 100644 --- a/catch/multiproc/hipMemCoherencyTstMProc.cc +++ b/catch/multiproc/hipMemCoherencyTstMProc.cc @@ -39,37 +39,15 @@ #include #include #include +#include "../unit/memory/hipSVMCommon.h" -__global__ void CoherentTst(int *ptr, int PeakClk) { +__global__ void CoherentTst(int *ptr, volatile unsigned int *expired) { // Incrementing the value by 1 - int64_t GpuFrq = int64_t(PeakClk) * 1000; - int64_t StrtTck = clock64(); - #if HT_AMD atomicAdd_system(ptr, 1); - #else - atomicAdd(ptr, 1); - #endif - // The following while loop checks the value in ptr for around 3-4 seconds - while ((clock64() - StrtTck) <= (3 * GpuFrq)) { - #if HT_AMD - if (atomicCAS_system(ptr, 3, 4) == 3) break; - #else - if (atomicCAS(ptr, 3, 4) == 3) break; - #endif - } -} - -__global__ void CoherentTst_gfx11(int *ptr, int PeakClk) { -#if HT_AMD - // Incrementing the value by 1 - int64_t GpuFrq = int64_t(PeakClk) * 1000; - int64_t StrtTck = clock_function(); - atomicAdd_system(ptr, 1); - // The following while loop checks the value in ptr for around 3-4 seconds - while ((clock_function() - StrtTck) <= (3 * GpuFrq)) { + // The following while loop checks the value until expiration. + while (*expired == 0) { if (atomicCAS_system(ptr, 3, 4) == 3) break; } -#endif } __global__ void SquareKrnl(int *ptr) { @@ -77,40 +55,25 @@ __global__ void SquareKrnl(int *ptr) { *ptr = (*ptr) * (*ptr); } -// The variable below will work as signal to decide pass/fail -static bool YES_COHERENT = false; - // The function tests the coherency of allocated memory -static void TstCoherency(int *Ptr, bool HmmMem) { - int *Dptr = nullptr, peak_clk; +// Return false on failure, true on success. +bool static TstCoherency(int *Ptr, bool HmmMem) { + using namespace std::chrono_literals; + int *Dptr = nullptr; hipStream_t strm; HIP_CHECK(hipStreamCreate(&strm)); // storing value 1 in the memory created above *Ptr = 1; - // Getting gpu frequency - if (IsGfx11()) { - HIPCHECK(hipDeviceGetAttribute(&peak_clk, - hipDeviceAttributeWallClockRate, 0)); - } else { - HIPCHECK(hipDeviceGetAttribute(&peak_clk, - hipDeviceAttributeClockRate, 0)); - } + unsigned int *expired = nullptr; + HIP_CHECK(hipHostMalloc(&expired, sizeof(unsigned int))); // hipHostMallocCoherent by defaut + *expired = 0; if (!HmmMem) { - HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&Dptr), - Ptr, 0)); - if (IsGfx11()) { - CoherentTst_gfx11<<<1, 1, 0, strm>>>(Dptr, peak_clk); - } else { - CoherentTst<<<1, 1, 0, strm>>>(Dptr, peak_clk); - } + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&Dptr), Ptr, 0)); + CoherentTst<<<1, 1, 0, strm>>>(Dptr, expired); } else { - if (IsGfx11()) { - CoherentTst_gfx11<<<1, 1, 0, strm>>>(Ptr, peak_clk); - } else { - CoherentTst<<<1, 1, 0, strm>>>(Ptr, peak_clk); - } + CoherentTst<<<1, 1, 0, strm>>>(Ptr, expired); } // looping until the value is 2 for 3 seconds std::chrono::steady_clock::time_point start = @@ -119,14 +82,20 @@ static void TstCoherency(int *Ptr, bool HmmMem) { std::chrono::steady_clock::now() - start).count() < 3) { if (*Ptr == 2) { *Ptr += 1; + std::this_thread::sleep_for(200ms); // Make sure kernel gets updated Dptr break; } } + *expired = 1; // Notify kernel loop to exit HIP_CHECK(hipStreamSynchronize(strm)); HIP_CHECK(hipStreamDestroy(strm)); + HIP_CHECK(hipHostFree(expired)); + if (*Ptr == 4) { - YES_COHERENT = true; + return true; } + fprintf(stderr, "TstCoherency: *Ptr=%u\b", *Ptr); + return false; } /* Test case description: The following test validates if fine grain @@ -134,6 +103,7 @@ static void TstCoherency(int *Ptr, bool HmmMem) { // The following test is failing on Nvidia platform hence disabled it for now #if HT_AMD TEST_CASE("Unit_malloc_CoherentTst") { + CHECK_PCIE_ATOMICS_SUPPORT hipDeviceProp_t prop; HIPCHECK(hipGetDeviceProperties(&prop, 0)); char *p = NULL; @@ -146,12 +116,12 @@ TEST_CASE("Unit_malloc_CoherentTst") { if (managed == 1) { int *Ptr = nullptr, SIZE = sizeof(int); bool HmmMem = true; - YES_COHERENT = false; + // Allocating hipMallocManaged() memory Ptr = reinterpret_cast(malloc(SIZE)); - TstCoherency(Ptr, HmmMem); + auto ret = TstCoherency(Ptr, HmmMem); free(Ptr); - REQUIRE(YES_COHERENT); + REQUIRE(ret); } } else { HipTest::HIP_SKIP_TEST("GPU is not xnack enabled hence skipping the test...\n"); @@ -175,7 +145,7 @@ TEST_CASE("Unit_malloc_CoherentTstWthAdvise") { 0)); if (managed == 1) { int *Ptr = nullptr, SIZE = sizeof(int); - YES_COHERENT = false; + // Allocating hipMallocManaged() memory Ptr = reinterpret_cast(malloc(SIZE)); *Ptr = 4; @@ -197,6 +167,7 @@ TEST_CASE("Unit_malloc_CoherentTstWthAdvise") { // The following test is failing on Nvidia platform hence disabling it for now #if HT_AMD TEST_CASE("Unit_mmap_CoherentTst") { + CHECK_PCIE_ATOMICS_SUPPORT hipDeviceProp_t prop; HIPCHECK(hipGetDeviceProperties(&prop, 0)); char *p = NULL; @@ -214,14 +185,12 @@ TEST_CASE("Unit_mmap_CoherentTst") { WARN("Mapping Failed\n"); REQUIRE(false); } - // Initializing the value with 1 - *Ptr = 1; - TstCoherency(Ptr, HmmMem); + auto ret = TstCoherency(Ptr, HmmMem); int err = munmap(Ptr, sizeof(int)); if (err != 0) { WARN("munmap failed\n"); } - REQUIRE(YES_COHERENT); + REQUIRE(ret); } } else { HipTest::HIP_SKIP_TEST("GPU is not xnack enabled hence skipping the test...\n"); @@ -286,7 +255,6 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg1") { int stat = 0; if (fork() == 0) { int *Ptr = nullptr, *PtrD = nullptr, SIZE = sizeof(int); - YES_COHERENT = false; // Allocating hipHostMalloc() memory HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocPortable)); *Ptr = 4; @@ -327,7 +295,6 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg2") { int stat = 0; if (fork() == 0) { int *Ptr = nullptr, *PtrD = nullptr, SIZE = sizeof(int); - YES_COHERENT = false; // Allocating hipHostMalloc() memory HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocWriteCombined)); *Ptr = 4; @@ -368,7 +335,6 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg3") { int stat = 0; if (fork() == 0) { int *Ptr = nullptr, *PtrD = nullptr, SIZE = sizeof(int); - YES_COHERENT = false; // Allocating hipHostMalloc() memory HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNumaUser)); *Ptr = 4; @@ -409,7 +375,6 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg4") { int stat = 0; if (fork() == 0) { int *Ptr = nullptr, *PtrD = nullptr, SIZE = sizeof(int); - YES_COHERENT = false; // Allocating hipHostMalloc() memory HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNonCoherent)); *Ptr = 4; @@ -449,28 +414,18 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1") { REQUIRE(false); } int stat = 0; - if (fork() == 0) { // child process + CHECK_PCIE_ATOMICS_SUPPORT int *Ptr = nullptr, SIZE = sizeof(int); bool HmmMem = false; - YES_COHERENT = false; // Allocating hipHostMalloc() memory HIP_CHECK(hipHostMalloc(&Ptr, SIZE)); - *Ptr = 4; - TstCoherency(Ptr, HmmMem); - if (YES_COHERENT) { - // exit() with code 10 which indicates pass - HIP_CHECK(hipHostFree(Ptr)); - exit(10); - } else { - // exit() with code 9 which indicates fail - HIP_CHECK(hipHostFree(Ptr)); - exit(9); - } + auto ret = TstCoherency(Ptr, HmmMem); + HIP_CHECK(hipHostFree(Ptr)); + exit(ret ? EXIT_SUCCESS : EXIT_FAILURE); } else { // parent process wait(&stat); - int Result = WEXITSTATUS(stat); - if (Result != 10) { + if (WEXITSTATUS(stat) != EXIT_SUCCESS) { REQUIRE(false); } } @@ -488,28 +443,18 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg1") { REQUIRE(false); } int stat = 0; - if (fork() == 0) { // child process + CHECK_PCIE_ATOMICS_SUPPORT int *Ptr = nullptr, SIZE = sizeof(int); bool HmmMem = false; - YES_COHERENT = false; // Allocating hipHostMalloc() memory HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocPortable)); - *Ptr = 1; - TstCoherency(Ptr, HmmMem); - if (YES_COHERENT) { - // exit() with code 10 which indicates pass - HIP_CHECK(hipHostFree(Ptr)); - exit(10); - } else { - // exit() with code 9 which indicates fail - HIP_CHECK(hipHostFree(Ptr)); - exit(9); - } + auto ret = TstCoherency(Ptr, HmmMem); + HIP_CHECK(hipHostFree(Ptr)); + exit(ret ? EXIT_SUCCESS : EXIT_FAILURE); } else { // parent process wait(&stat); - int Result = WEXITSTATUS(stat); - if (Result != 10) { + if (WEXITSTATUS(stat) != EXIT_SUCCESS) { REQUIRE(false); } } @@ -526,28 +471,18 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg2") { REQUIRE(false); } int stat = 0; - if (fork() == 0) { // child process + CHECK_PCIE_ATOMICS_SUPPORT int *Ptr = nullptr, SIZE = sizeof(int); bool HmmMem = false; - YES_COHERENT = false; // Allocating hipHostMalloc() memory HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocWriteCombined)); - *Ptr = 4; - TstCoherency(Ptr, HmmMem); - if (YES_COHERENT) { - // exit() with code 10 which indicates pass - HIP_CHECK(hipHostFree(Ptr)); - exit(10); - } else { - // exit() with code 9 which indicates fail - HIP_CHECK(hipHostFree(Ptr)); - exit(9); - } + auto ret = TstCoherency(Ptr, HmmMem); + HIP_CHECK(hipHostFree(Ptr)); + exit(ret ? EXIT_SUCCESS : EXIT_FAILURE); } else { // parent process wait(&stat); - int Result = WEXITSTATUS(stat); - if (Result != 10) { + if (WEXITSTATUS(stat) != EXIT_SUCCESS) { REQUIRE(false); } } @@ -564,28 +499,18 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg3") { REQUIRE(false); } int stat = 0; - if (fork() == 0) { // child process + CHECK_PCIE_ATOMICS_SUPPORT int *Ptr = nullptr, SIZE = sizeof(int); bool HmmMem = false; - YES_COHERENT = false; // Allocating hipHostMalloc() memory HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNumaUser)); - *Ptr = 1; - TstCoherency(Ptr, HmmMem); - if (YES_COHERENT) { - // exit() with code 10 which indicates pass - HIP_CHECK(hipHostFree(Ptr)); - exit(10); - } else { - // exit() with code 9 which indicates fail - HIP_CHECK(hipHostFree(Ptr)); - exit(9); - } + auto ret = TstCoherency(Ptr, HmmMem); + HIP_CHECK(hipHostFree(Ptr)); + exit(ret ? EXIT_SUCCESS : EXIT_FAILURE); } else { // parent process wait(&stat); - int Result = WEXITSTATUS(stat); - if (Result != 10) { + if (WEXITSTATUS(stat) != EXIT_SUCCESS) { REQUIRE(false); } } diff --git a/catch/perftests/graph/parallelGraph.cc b/catch/perftests/graph/parallelGraph.cc index d8eac866e3..d03d1b8441 100644 --- a/catch/perftests/graph/parallelGraph.cc +++ b/catch/perftests/graph/parallelGraph.cc @@ -39,37 +39,19 @@ const unsigned int kNumNode = 5; * - Launches an executable graph in the specified stream. */ unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); -__device__ void Delay(uint32_t time, const uint32_t ticks_per_ms) { - while (time--) { -#if HT_AMD - uint64_t start = wall_clock64(); - while (wall_clock64() - start < ticks_per_ms) { - __builtin_amdgcn_s_sleep(10); - } -#endif -#if HT_NVIDIA - uint64_t start = clock64(); - while (clock64() - start < ticks_per_ms) { - } -#endif - } -} template -__global__ void vectorADD(const T *A_d, const T *B_d, T *C_d, size_t NELEM, - int clockrate) { +__global__ void vectorADD(const T *A_d, const T *B_d, T *C_d, size_t NELEM) { size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); size_t stride = blockDim.x * gridDim.x; for (size_t i = offset; i < NELEM; i += stride) { C_d[i] = A_d[i] + B_d[i]; } - Delay(1, clockrate); } /** * Test Description * ------------------------ * - Create the graph with multiple parallel branches. - * - Introduce some delay in the kernel. * - Calculate the time taken to graph execution. * Test source * ------------------------ @@ -79,14 +61,6 @@ __global__ void vectorADD(const T *A_d, const T *B_d, T *C_d, size_t NELEM, * - HIP_VERSION >= 6.4 */ TEST_CASE("Unit_hipGraph_Performance_Improvement_ParallelGraph") { - int clkRate; -#if HT_AMD - HIP_CHECK( - hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0)); -#endif -#if HT_NVIDIA - HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0)); -#endif hipGraphNode_t memCpy1, memCpy2, memCpy3; std::vector kNode(kNumNode); hipGraph_t graph; @@ -107,8 +81,7 @@ TEST_CASE("Unit_hipGraph_Performance_Improvement_ParallelGraph") { for (int i = 0; i < kNumNode; i++) { hipKernelNodeParams kernelNodeParams{}; - void *kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem), - reinterpret_cast(&clkRate)}; + void *kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; kernelNodeParams.func = reinterpret_cast(vectorADD); kernelNodeParams.gridDim = dim3(blocks); kernelNodeParams.blockDim = dim3(threadsPerBlock); @@ -166,15 +139,6 @@ TEST_CASE("Unit_hipGraph_Performance_Improvement_ParallelGraph") { * - HIP_VERSION >= 6.4 */ TEST_CASE("Unit_hipGraph_Performance_With_Stream_Operations") { - int clkRate; -#if HT_AMD - HIP_CHECK( - hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0)); -#endif -#if HT_NVIDIA - HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0)); -#endif - unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); hipStream_t stream; HIP_CHECK(hipStreamCreate(&stream)); @@ -187,7 +151,7 @@ TEST_CASE("Unit_hipGraph_Performance_With_Stream_Operations") { HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyDefault, stream)); for (int i = 0; i < kNumNode; i++) { hipLaunchKernelGGL(vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, - stream, A_d, B_d, C_d, NElem, clkRate); + stream, A_d, B_d, C_d, NElem); } HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDefault, stream)); HIP_CHECK(hipStreamSynchronize(stream)); @@ -218,15 +182,6 @@ TEST_CASE("Unit_hipGraph_Performance_With_Stream_Operations") { */ TEST_CASE("Unit_hipGraph_Performance_With_Stream_Capture") { - int clkRate; -#if HT_AMD - HIP_CHECK( - hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0)); -#endif -#if HT_NVIDIA - HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0)); -#endif - unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); hipGraph_t graph; hipStream_t stream, streamForGraph; @@ -239,7 +194,7 @@ TEST_CASE("Unit_hipGraph_Performance_With_Stream_Capture") { HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyDefault, stream)); for (int i = 0; i < kNumNode; i++) { hipLaunchKernelGGL(vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, - stream, A_d, B_d, C_d, NElem, clkRate); + stream, A_d, B_d, C_d, NElem); } HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDefault, stream)); HIP_CHECK(hipStreamEndCapture(stream, &graph)); diff --git a/catch/perftests/stream/hipPerfDeviceConcurrency.cc b/catch/perftests/stream/hipPerfDeviceConcurrency.cc index c0ea3e4638..b07c9f49ed 100644 --- a/catch/perftests/stream/hipPerfDeviceConcurrency.cc +++ b/catch/perftests/stream/hipPerfDeviceConcurrency.cc @@ -133,6 +133,10 @@ bool hipPerfDeviceConcurrency::run(unsigned int testCase, int numGpus) { int clkFrequency = 0; HIP_CHECK(hipDeviceGetAttribute(&clkFrequency, hipDeviceAttributeClockRate, i)); + if (clkFrequency == 0) { + std::cout << "clkFrequency = 0, set it to 1000000\n"; + clkFrequency = 1000000; + } clkFrequency =(unsigned int)clkFrequency/1000; // Maximum iteration count diff --git a/catch/perftests/stream/hipPerfStreamConcurrency.cc b/catch/perftests/stream/hipPerfStreamConcurrency.cc index 2c86c27503..ba4a04aa9d 100644 --- a/catch/perftests/stream/hipPerfStreamConcurrency.cc +++ b/catch/perftests/stream/hipPerfStreamConcurrency.cc @@ -245,7 +245,10 @@ bool hipPerfStreamConcurrency::run(unsigned int testCase, HIP_CHECK(hipDeviceGetAttribute(&clkFrequency, hipDeviceAttributeClockRate, deviceId)); - + if (clkFrequency == 0) { + std::cout << "clkFrequency = 0, set it to 1000000\n"; + clkFrequency = 1000000; + } clkFrequency =(unsigned int)clkFrequency/1000; // Maximum iteration count diff --git a/catch/unit/clock/hipClockCheck.cc b/catch/unit/clock/hipClockCheck.cc index 0c6cdee14c..9fa876e97c 100644 --- a/catch/unit/clock/hipClockCheck.cc +++ b/catch/unit/clock/hipClockCheck.cc @@ -120,7 +120,10 @@ TEST_CASE("Unit_hipClock64_Positive_Basic") { HIP_CHECK(hipSetDevice(0)); int clock_rate = 0; // in kHz HIP_CHECK(hipDeviceGetAttribute(&clock_rate, hipDeviceAttributeClockRate, 0)); - + if (clock_rate == 0) { + HipTest::HIP_SKIP_TEST("hipDeviceAttributeClockRate returns 0"); + return; + } if (IsGfx11()) { HipTest::HIP_SKIP_TEST("Issue with clock64() function on gfx11 devices!"); return; @@ -149,7 +152,10 @@ TEST_CASE("Unit_hipClock_Positive_Basic") { HIP_CHECK(hipSetDevice(0)); int clock_rate = 0; // in kHz HIP_CHECK(hipDeviceGetAttribute(&clock_rate, hipDeviceAttributeClockRate, 0)); - + if (clock_rate == 0) { + HipTest::HIP_SKIP_TEST("hipDeviceAttributeClockRate returns 0"); + return; + } if (IsGfx11()) { HipTest::HIP_SKIP_TEST("Issue with clock() function on gfx11 devices!"); return; @@ -180,7 +186,7 @@ TEST_CASE("Unit_hipWallClock64_Positive_Basic") { HIP_CHECK(hipDeviceGetAttribute(&clock_rate, hipDeviceAttributeWallClockRate, 0)); if (!clock_rate) { - HipTest::HIP_SKIP_TEST("hipDeviceAttributeWallClockRate is not supported"); + HipTest::HIP_SKIP_TEST("hipDeviceAttributeWallClockRate returns 0"); return; } diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index ade91ddd86..6beae6009a 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -276,4 +276,5 @@ hip_add_exe_to_target(NAME SVMAtomicTest if(HIP_PLATFORM MATCHES "nvidia") set_target_properties(SVMAtomicTest PROPERTIES COMPILE_FLAGS -arch=sm_70) set_target_properties(MemoryTest1 PROPERTIES COMPILE_FLAGS -arch=sm_70) + set_target_properties(MemoryTest2 PROPERTIES COMPILE_FLAGS -arch=sm_70) endif() diff --git a/catch/unit/memory/hipMemPoolApi.cc b/catch/unit/memory/hipMemPoolApi.cc index ee85e4fc88..b2f47a7673 100644 --- a/catch/unit/memory/hipMemPoolApi.cc +++ b/catch/unit/memory/hipMemPoolApi.cc @@ -21,7 +21,7 @@ 1) This testcase verifies the basic scenario - supported on all devices */ - +#include "mempool_common.hh" #include #include #include @@ -105,36 +105,6 @@ TEST_CASE("Unit_hipMemPoolApi_Basic") { HIP_CHECK(hipStreamDestroy(stream)); } -constexpr auto wait_ms = 500; - -__global__ void kernel500ms(float* hostRes, int clkRate) { - int tid = threadIdx.x + blockIdx.x * blockDim.x; - hostRes[tid] = tid + 1; - __threadfence_system(); - // expecting that the data is getting flushed to host here! - uint64_t start = clock64()/clkRate, cur; - if (clkRate > 1) { - do { cur = clock64()/clkRate-start;}while (cur < wait_ms); - } else { - do { cur = clock64()/start;}while (cur < wait_ms); - } -} - -__global__ void kernel500ms_gfx11(float* hostRes, int clkRate) { -#if HT_AMD - int tid = threadIdx.x + blockIdx.x * blockDim.x; - hostRes[tid] = tid + 1; - __threadfence_system(); - // expecting that the data is getting flushed to host here! - uint64_t start = clock_function()/clkRate, cur; - if (clkRate > 1) { - do { cur = clock_function()/clkRate-start;}while (cur < wait_ms); - } else { - do { cur = clock_function()/start;}while (cur < wait_ms); - } -#endif -} - TEST_CASE("Unit_hipMemPoolApi_BasicAlloc") { int mem_pool_support = 0; HIP_CHECK(hipSetDevice(0)); @@ -144,6 +114,9 @@ TEST_CASE("Unit_hipMemPoolApi_BasicAlloc") { SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); return; } + unsigned int *notified = nullptr; + HIP_CHECK(hipHostMalloc(¬ified, sizeof(unsigned int))); + *notified = 0; initMemPoolProps(); hipMemPool_t mem_pool; HIP_CHECK(hipMemPoolCreate(&mem_pool, &kPoolProps)); @@ -159,16 +132,8 @@ TEST_CASE("Unit_hipMemPoolApi_BasicAlloc") { HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&C), numElements * sizeof(float), mem_pool, stream)); int blocks = 1024; - int clkRate; hipMemPoolAttr attr; - if (IsGfx11()) { - HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0)); - kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate); - } else { - HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0)); - - kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate); - } + notifiedKernel<<<32, blocks, 0, stream>>>(B, notified); HIP_CHECK(hipFreeAsync(reinterpret_cast(B), stream)); @@ -176,6 +141,9 @@ TEST_CASE("Unit_hipMemPoolApi_BasicAlloc") { std::uint64_t res_before_sync = 0; HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &res_before_sync)); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified = 1; // Notify kernel loop to exit + HIP_CHECK(hipStreamSynchronize(stream)); std::uint64_t res_after_sync = 0; @@ -223,6 +191,7 @@ TEST_CASE("Unit_hipMemPoolApi_BasicAlloc") { HIP_CHECK(hipMemPoolDestroy(mem_pool)); HIP_CHECK(hipFreeAsync(reinterpret_cast(C), stream)); HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipHostFree(notified)); } TEST_CASE("Unit_hipMemPoolApi_BasicTrim") { @@ -232,6 +201,9 @@ TEST_CASE("Unit_hipMemPoolApi_BasicTrim") { SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); return; } + unsigned int *notified = nullptr; + HIP_CHECK(hipHostMalloc(¬ified, sizeof(unsigned int))); + *notified = 0; initMemPoolProps(); hipMemPool_t mem_pool; HIP_CHECK(hipMemPoolCreate(&mem_pool, &kPoolProps)); @@ -247,15 +219,7 @@ TEST_CASE("Unit_hipMemPoolApi_BasicTrim") { HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&C), numElements * sizeof(float), mem_pool, stream)); int blocks = 2; - int clkRate; - if (IsGfx11()) { - HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0)); - kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate); - } else { - HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0)); - - kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate); - } + notifiedKernel<<<32, blocks, 0, stream>>>(B, notified); hipMemPoolAttr attr; attr = hipMemPoolAttrReleaseThreshold; @@ -279,6 +243,8 @@ TEST_CASE("Unit_hipMemPoolApi_BasicTrim") { // Trim must be a nop because execution isn't done REQUIRE(res_before_trim == res_after_trim); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified = 1; // Notify kernel loop to exit HIP_CHECK(hipStreamSynchronize(stream)); std::uint64_t res_after_sync = 0; @@ -311,6 +277,7 @@ TEST_CASE("Unit_hipMemPoolApi_BasicTrim") { HIP_CHECK(hipMemPoolDestroy(mem_pool)); HIP_CHECK(hipFreeAsync(reinterpret_cast(C), stream)); HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipHostFree(notified)); } TEST_CASE("Unit_hipMemPoolApi_BasicReuse") { @@ -320,6 +287,9 @@ TEST_CASE("Unit_hipMemPoolApi_BasicReuse") { SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); return; } + unsigned int *notified = nullptr; + HIP_CHECK(hipHostMalloc(¬ified, sizeof(unsigned int))); + *notified = 0; initMemPoolProps(); hipMemPool_t mem_pool; HIP_CHECK(hipMemPoolCreate(&mem_pool, &kPoolProps)); @@ -335,16 +305,7 @@ TEST_CASE("Unit_hipMemPoolApi_BasicReuse") { HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&C), numElements * sizeof(float), mem_pool, stream)); int blocks = 2; - int clkRate; - - if (IsGfx11()) { - HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0)); - kernel500ms_gfx11<<<32, blocks, 0, stream>>>(A, clkRate); - } else { - HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0)); - - kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate); - } + notifiedKernel<<<32, blocks, 0, stream>>>(A, notified); hipMemPoolAttr attr; // Not a real free, since kernel isn't done @@ -355,16 +316,17 @@ TEST_CASE("Unit_hipMemPoolApi_BasicReuse") { // Runtime must reuse the pointer REQUIRE(A == B); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified = 1; // Notify kernel loop to exit // Make a sync before the second kernel launch to make sure memory B isn't gone HIP_CHECK(hipStreamSynchronize(stream)); // Second kernel launch with new memory - if (IsGfx11()) { - kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate); - } else { - kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate); - } + *notified = 0; + notifiedKernel<<<32, blocks, 0, stream>>>(B, notified); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified = 1; // Notify kernel loop to exit HIP_CHECK(hipStreamSynchronize(stream)); attr = hipMemPoolAttrUsedMemCurrent; @@ -387,6 +349,7 @@ TEST_CASE("Unit_hipMemPoolApi_BasicReuse") { HIP_CHECK(hipMemPoolDestroy(mem_pool)); HIP_CHECK(hipFreeAsync(reinterpret_cast(C), stream)); HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipHostFree(notified)); } TEST_CASE("Unit_hipMemPoolApi_Opportunistic") { @@ -396,33 +359,32 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") { SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); return; } + unsigned int *notified1 = nullptr, *notified2 = nullptr; + HIP_CHECK(hipHostMalloc(¬ified1, sizeof(unsigned int))); + HIP_CHECK(hipHostMalloc(¬ified2, sizeof(unsigned int))); + *notified1 = 0; + *notified2 = 0; initMemPoolProps(); hipMemPool_t mem_pool; HIP_CHECK(hipMemPoolCreate(&mem_pool, &kPoolProps)); hipMemPoolAttr attr; int blocks = 2; - int clkRate; - if (IsGfx11()) { - HIPCHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0)); - } else { - HIPCHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0)); - } float *A, *B, *C; - hipStream_t stream, stream2; + hipStream_t stream1, stream2; // Create 2 async non-blocking streams - HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); + HIP_CHECK(hipStreamCreateWithFlags(&stream1, hipStreamNonBlocking)); HIP_CHECK(hipStreamCreateWithFlags(&stream2, hipStreamNonBlocking)); size_t numElements = 1024; - HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&C), numElements * sizeof(float), mem_pool, stream)); + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&C), numElements * sizeof(float), mem_pool, stream1)); int value = 0; SECTION("Disallow Opportunistic - No Reuse") { numElements = 8 * 1024 * 1024; - HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&A), numElements * sizeof(float), mem_pool, stream)); + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&A), numElements * sizeof(float), mem_pool, stream1)); // Disable all default pool states attr = hipMemPoolReuseFollowEventDependencies; @@ -432,15 +394,13 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") { attr = hipMemPoolReuseAllowInternalDependencies; HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value)); - // Run kernel for 500 ms in the first stream - if (IsGfx11()) { - kernel500ms_gfx11<<<32, blocks, 0, stream>>>(A, clkRate); - } else { - kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate); - } + // Run kernel in the first stream + notifiedKernel<<<32, blocks, 0, stream1>>>(A, notified1); // Not a real free, since kernel isn't done - HIP_CHECK(hipFreeAsync(reinterpret_cast(A), stream)); + HIP_CHECK(hipFreeAsync(reinterpret_cast(A), stream1)); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified1 = 1; // Notify kernel loop to exit // Sleep for 1 second GPU should be idle by now std::this_thread::sleep_for(std::chrono::milliseconds(1000)); @@ -451,14 +411,12 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") { // Without Opportunistic state runtime must allocate another buffer REQUIRE(A != B); - // Run kernel with the new memory in the second stream - if (IsGfx11()) { - kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate); - } else { - kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate); - } + // Run kernel with the new memory in the second streamn + notifiedKernel<<<32, blocks, 0, stream2>>>(B, notified2); - HIP_CHECK(hipStreamSynchronize(stream)); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified2 = 1; // Notify kernel loop to exit + HIP_CHECK(hipStreamSynchronize(stream1)); HIP_CHECK(hipStreamSynchronize(stream2)); HIP_CHECK(hipFreeAsync(reinterpret_cast(B), stream2)); @@ -466,24 +424,20 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") { SECTION("Allow Opportunistic - Reuse") { numElements = 8 * 1024 * 1024; - HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&A), numElements * sizeof(float), mem_pool, stream)); + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&A), numElements * sizeof(float), mem_pool, stream1)); value = 1; attr = hipMemPoolReuseAllowOpportunistic; // Enable Opportunistic HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value)); - // Run kernel for 500 ms in the first stream - if (IsGfx11()) { - HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0)); - kernel500ms_gfx11<<<32, blocks, 0, stream>>>(A, clkRate); - } else { - HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0)); - kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate); - } + // Run kernel in the first stream + notifiedKernel<<<32, blocks, 0, stream1>>>(A, notified1); // Not a real free, since kernel isn't done - HIP_CHECK(hipFreeAsync(reinterpret_cast(A), stream)); + HIP_CHECK(hipFreeAsync(reinterpret_cast(A), stream1)); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified1 = 1; // Notify kernel loop to exit // Sleep for 1 second GPU should be idle by now std::this_thread::sleep_for(std::chrono::milliseconds(1000)); @@ -495,13 +449,12 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") { REQUIRE(A == B); // Run kernel with the new memory in the second stream - if (IsGfx11()) { - kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate); - } else { - kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate); - } + notifiedKernel<<<32, blocks, 0, stream2>>>(B, notified2); - HIP_CHECK(hipStreamSynchronize(stream)); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified2 = 1; // Notify kernel loop to exit + + HIP_CHECK(hipStreamSynchronize(stream1)); HIP_CHECK(hipStreamSynchronize(stream2)); HIP_CHECK(hipFreeAsync(reinterpret_cast(B), stream2)); @@ -509,23 +462,18 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") { SECTION("Allow Opportunistic - No Reuse") { numElements = 8 * 1024 * 1024; - HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&A), numElements * sizeof(float), mem_pool, stream)); + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&A), numElements * sizeof(float), mem_pool, stream1)); value = 1; attr = hipMemPoolReuseAllowOpportunistic; // Enable Opportunistic HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value)); - // Run kernel for 500 ms in the first stream - - if (IsGfx11()) { - kernel500ms_gfx11<<<32, blocks, 0, stream>>>(A, clkRate); - } else { - kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate); - } + // Run kernel in the first stream + notifiedKernel<<<32, blocks, 0, stream1>>>(A, notified1); // Not a real free, since kernel isn't done - HIP_CHECK(hipFreeAsync(reinterpret_cast(A), stream)); + HIP_CHECK(hipFreeAsync(reinterpret_cast(A), stream1)); numElements = 8 * 1024 * 1024; // Allocate memory for the second stream @@ -534,22 +482,23 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") { REQUIRE(A != B); // Run kernel with the new memory in the second stream - if (IsGfx11()) { - kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate); - } else { - kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate); - } + notifiedKernel<<<32, blocks, 0, stream2>>>(B, notified2); - HIP_CHECK(hipStreamSynchronize(stream)); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified1 = 1; // Notify kernel loop to exit + *notified2 = 1; // Notify kernel loop to exit + HIP_CHECK(hipStreamSynchronize(stream1)); HIP_CHECK(hipStreamSynchronize(stream2)); HIP_CHECK(hipFreeAsync(reinterpret_cast(B), stream2)); } - HIP_CHECK(hipFreeAsync(reinterpret_cast(C), stream)); + HIP_CHECK(hipFreeAsync(reinterpret_cast(C), stream1)); HIP_CHECK(hipMemPoolDestroy(mem_pool)); - HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipStreamDestroy(stream1)); HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipHostFree(notified1)); + HIP_CHECK(hipHostFree(notified2)); } TEST_CASE("Unit_hipMemPoolApi_Default") { @@ -559,7 +508,9 @@ TEST_CASE("Unit_hipMemPoolApi_Default") { SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); return; } - + unsigned int *notified = nullptr; + HIP_CHECK(hipHostMalloc(¬ified, sizeof(unsigned int))); + *notified = 0; hipMemPool_t mem_pool; HIP_CHECK(hipDeviceGetDefaultMemPool(&mem_pool, 0)); @@ -574,16 +525,7 @@ TEST_CASE("Unit_hipMemPoolApi_Default") { HIP_CHECK(hipMallocAsync(reinterpret_cast(&C), numElements * sizeof(float), stream)); int blocks = 2; - int clkRate; - - if (IsGfx11()) { - HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0)); - kernel500ms_gfx11<<<32, blocks, 0, stream>>>(A, clkRate); - } else { - HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0)); - - kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate); - } + notifiedKernel<<<32, blocks, 0, stream>>>(A, notified); hipMemPoolAttr attr; // Not a real free, since kernel isn't done @@ -595,17 +537,18 @@ TEST_CASE("Unit_hipMemPoolApi_Default") { REQUIRE(A == B); // Make a sync before the second kernel launch to make sure memory B isn't gone + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified = 1; // Notify kernel loop to exit HIP_CHECK(hipStreamSynchronize(stream)); // Second kernel launch with new memory - if (IsGfx11()) { - kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate); - } else { - kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate); - } + *notified = 0; + notifiedKernel<<<32, blocks, 0, stream>>>(B, notified); HIP_CHECK(hipFreeAsync(reinterpret_cast(B), stream)); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified = 1; // Notify kernel loop to exit HIP_CHECK(hipStreamSynchronize(stream)); std::uint64_t value64 = 0; @@ -626,4 +569,5 @@ TEST_CASE("Unit_hipMemPoolApi_Default") { HIP_CHECK(hipFreeAsync(reinterpret_cast(C), stream)); HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipHostFree(notified)); } diff --git a/catch/unit/memory/hipMemPoolSetGetAttribute.cc b/catch/unit/memory/hipMemPoolSetGetAttribute.cc index 32b596e19d..d6bfa47ee2 100644 --- a/catch/unit/memory/hipMemPoolSetGetAttribute.cc +++ b/catch/unit/memory/hipMemPoolSetGetAttribute.cc @@ -19,7 +19,6 @@ #include "mempool_common.hh" #include #include - /** * @addtogroup hipMemPoolSetAttribute hipMemPoolSetAttribute * @{ @@ -120,20 +119,18 @@ TEST_CASE("Unit_hipMemPoolSetGetAttribute_Positive_MemBasic") { TEST_CASE("Unit_hipMemPoolSetAttribute_Opportunistic") { int device_id = 0; HIP_CHECK(hipSetDevice(device_id)); - checkMempoolSupported(device_id) + unsigned int* notified1 = nullptr, *notified2 = nullptr; + HIP_CHECK(hipHostMalloc(¬ified1, sizeof(unsigned int))); + HIP_CHECK(hipHostMalloc(¬ified2, sizeof(unsigned int))); + *notified1 = 0; + *notified2 = 0; + MemPoolGuard mempool(MemPools::created, device_id); hipMemPoolAttr attr; int blocks = 2; - int clk_rate; - if (IsGfx11()) { - HIPCHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeWallClockRate, 0)); - } else { - HIPCHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeClockRate, 0)); - } - int *alloc_mem1, *alloc_mem2, *alloc_mem3; // Create 2 async non-blocking streams @@ -158,16 +155,12 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_Opportunistic") { attr = hipMemPoolReuseAllowInternalDependencies; HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value)); - // Run kernel for 500 ms in the first stream - if (IsGfx11()) { - kernel_500ms_gfx11<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate); - } else { - kernel_500ms<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate); - } - + // Run kernel in the first stream + notifiedKernel<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, notified1); // Not a real free, since kernel isn't done HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem1), stream1.stream())); - + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified1 = 1; // Sleep for 1 second GPU should be idle by now std::this_thread::sleep_for(std::chrono::milliseconds(1000)); @@ -178,12 +171,10 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_Opportunistic") { REQUIRE(alloc_mem1 != alloc_mem2); // Run kernel with the new memory in the second stream - if (IsGfx11()) { - kernel_500ms_gfx11<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate); - } else { - kernel_500ms<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate); - } + notifiedKernel<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, notified2); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified2 = 1; HIP_CHECK(hipStreamSynchronize(stream1.stream())); HIP_CHECK(hipStreamSynchronize(stream2.stream())); @@ -203,32 +194,27 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_Opportunistic") { attr = hipMemPoolReuseAllowInternalDependencies; HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value)); - // Run kernel for 500 ms in the first stream - if (IsGfx11()) { - kernel_500ms_gfx11<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate); - } else { - kernel_500ms<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate); - } + // Run kernel in the first stream + notifiedKernel<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, notified1); // Not a real free, since kernel isn't done HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem1), stream1.stream())); - + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified1 = 1; // Sleep for 1 second GPU should be idle by now std::this_thread::sleep_for(std::chrono::milliseconds(1000)); - // Allocate memory for the second stream + // Allocate memory for the first stream HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&alloc_mem2), allocation_size, mempool.mempool(), stream1.stream())); - // Without Opportunistic state runtime must allocate another buffer + // Without Opportunistic state runtime must reuse freed buffer REQUIRE(alloc_mem1 == alloc_mem2); - // Run kernel with the new memory in the second stream - if (IsGfx11()) { - kernel_500ms_gfx11<<<32, blocks, 0, stream1.stream()>>>(alloc_mem2, clk_rate); - } else { - kernel_500ms<<<32, blocks, 0, stream1.stream()>>>(alloc_mem2, clk_rate); - } + // Run kernel with the new memory in the first stream + notifiedKernel<<<32, blocks, 0, stream1.stream()>>>(alloc_mem2, notified2); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified2 = 1; HIP_CHECK(hipStreamSynchronize(stream1.stream())); HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem2), stream1.stream())); @@ -243,19 +229,15 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_Opportunistic") { attr = hipMemPoolReuseAllowOpportunistic; // Enable Opportunistic HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value)); - - // Run kernel for 500 ms in the first stream - if (IsGfx11()) { - HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeWallClockRate, 0)); - kernel_500ms_gfx11<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate); - } else { - HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeClockRate, 0)); - kernel_500ms<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate); - } + // Run kernel in the first stream + notifiedKernel<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, notified1); // Not a real free, since kernel isn't done HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem1), stream1.stream())); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified1 = 1; // Notifiy kernel to exit after 500 ms + // Sleep for 1 second GPU should be idle by now std::this_thread::sleep_for(std::chrono::milliseconds(1000)); @@ -266,11 +248,10 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_Opportunistic") { REQUIRE(alloc_mem1 == alloc_mem2); // Run kernel with the new memory in the second stream - if (IsGfx11()) { - kernel_500ms_gfx11<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate); - } else { - kernel_500ms<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate); - } + notifiedKernel<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, notified2); + + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified2 = 1; // Notifiy kernel to exit after 500 ms HIP_CHECK(hipStreamSynchronize(stream1.stream())); HIP_CHECK(hipStreamSynchronize(stream2.stream())); @@ -288,13 +269,8 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_Opportunistic") { // Enable Opportunistic HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value)); - // Run kernel for 500 ms in the first stream - - if (IsGfx11()) { - kernel_500ms_gfx11<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate); - } else { - kernel_500ms<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate); - } + // Run kernel in the first stream + notifiedKernel<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, notified1); // Not a real free, since kernel isn't done HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem1), stream1.stream())); @@ -307,12 +283,11 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_Opportunistic") { REQUIRE(alloc_mem1 != alloc_mem2); // Run kernel with the new memory in the second stream - if (IsGfx11()) { - kernel_500ms_gfx11<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate); - } else { - kernel_500ms<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate); - } + notifiedKernel<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, notified2); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified1 = 1; + *notified2 = 1; HIP_CHECK(hipStreamSynchronize(stream1.stream())); HIP_CHECK(hipStreamSynchronize(stream2.stream())); @@ -320,6 +295,8 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_Opportunistic") { } HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem3), stream1.stream())); + HIP_CHECK(hipHostFree(notified1)); + HIP_CHECK(hipHostFree(notified2)); } /** @@ -343,12 +320,12 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_EventDependencies") { hipMemPoolAttr attr; int blocks = 2; - int clk_rate; - if (IsGfx11()) { - HIPCHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeWallClockRate, 0)); - } else { - HIPCHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeClockRate, 0)); - } + + unsigned int* notified1 = nullptr, *notified2 = nullptr; + HIP_CHECK(hipHostMalloc(¬ified1, sizeof(unsigned int))); + HIP_CHECK(hipHostMalloc(¬ified2, sizeof(unsigned int))); + *notified1 = 0; + *notified2 = 0; int *alloc_mem1, *alloc_mem2, *alloc_mem3; @@ -371,17 +348,11 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_EventDependencies") { value = 1; attr = hipMemPoolReuseFollowEventDependencies; - // Enable Opportunistic + // Enable Opportunistic- HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value)); - // Run kernel for 500 ms in the first stream - if (IsGfx11()) { - HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeWallClockRate, 0)); - kernel_500ms_gfx11<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate); - } else { - HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeClockRate, 0)); - kernel_500ms<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate); - } + // Run kernel in the first stream + notifiedKernel<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, notified1); // Not a real free, since kernel isn't done HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem1), stream1.stream())); @@ -396,12 +367,11 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_EventDependencies") { REQUIRE(alloc_mem1 == alloc_mem2); // Run kernel with the new memory in the second stream - if (IsGfx11()) { - kernel_500ms_gfx11<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate); - } else { - kernel_500ms<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate); - } + notifiedKernel<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, notified2); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified1 = 1; + *notified2 = 1; HIP_CHECK(hipStreamSynchronize(stream1.stream())); HIP_CHECK(hipStreamSynchronize(stream2.stream())); @@ -418,17 +388,10 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_EventDependencies") { // Enable Opportunistic HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value)); - // Run kernel for 500 ms in the first stream - - if (IsGfx11()) { - kernel_500ms_gfx11<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate); - } else { - kernel_500ms<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate); - } - + // Run kernel in the first stream + notifiedKernel<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, notified1); // Not a real free, since kernel isn't done HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem1), stream1.stream())); - HIP_CHECK(hipEventRecord(event, stream1.stream())); HIP_CHECK(hipStreamWaitEvent(stream2.stream(), event, 0)); @@ -440,12 +403,11 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_EventDependencies") { REQUIRE(alloc_mem1 != alloc_mem2); // Run kernel with the new memory in the second stream - if (IsGfx11()) { - kernel_500ms_gfx11<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate); - } else { - kernel_500ms<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate); - } + notifiedKernel<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, notified2); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified1 = 1; + *notified2 = 1; HIP_CHECK(hipStreamSynchronize(stream1.stream())); HIP_CHECK(hipStreamSynchronize(stream2.stream())); @@ -454,6 +416,8 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_EventDependencies") { HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem3), stream1.stream())); HIP_CHECK(hipEventDestroy(event)); + HIP_CHECK(hipHostFree(notified1)); + HIP_CHECK(hipHostFree(notified2)); } /** diff --git a/catch/unit/memory/hipMemPoolTrimTo.cc b/catch/unit/memory/hipMemPoolTrimTo.cc index b1007de94a..9e58320c71 100644 --- a/catch/unit/memory/hipMemPoolTrimTo.cc +++ b/catch/unit/memory/hipMemPoolTrimTo.cc @@ -22,7 +22,6 @@ #include #include - /** * @addtogroup hipMemPoolTrimTo hipMemPoolTrimTo * @{ @@ -72,6 +71,9 @@ TEST_CASE("Unit_hipMemPoolTrimTo_Positive_Basic") { int device_id = 0; HIP_CHECK(hipSetDevice(device_id)); checkMempoolSupported(device_id) + unsigned int *notified = nullptr; + HIP_CHECK(hipHostMalloc(¬ified, sizeof(unsigned int))); + *notified = 0; const size_t allocation_size1 = kPageSize * kPageSize * 2; const size_t allocation_size2 = kPageSize / 2; @@ -87,15 +89,7 @@ TEST_CASE("Unit_hipMemPoolTrimTo_Positive_Basic") { mempool.mempool(), stream.stream())); int blocks = 2; - int clk_rate; - if (IsGfx11()) { - HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeWallClockRate, 0)); - kernel_500ms_gfx11<<<32, blocks, 0, stream.stream()>>>(alloc_mem1, clk_rate); - } else { - HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeClockRate, 0)); - - kernel_500ms<<<32, blocks, 0, stream.stream()>>>(alloc_mem1, clk_rate); - } + notifiedKernel<<<32, blocks, 0, stream.stream()>>>(alloc_mem1, notified); hipMemPoolAttr attr; attr = hipMemPoolAttrReleaseThreshold; @@ -119,6 +113,8 @@ TEST_CASE("Unit_hipMemPoolTrimTo_Positive_Basic") { // Trim must be a nop because execution isn't done REQUIRE(res_before_trim == res_after_trim); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified = 1; HIP_CHECK(hipStreamSynchronize(stream.stream())); std::uint64_t res_after_sync = 0; @@ -149,6 +145,7 @@ TEST_CASE("Unit_hipMemPoolTrimTo_Positive_Basic") { REQUIRE((allocation_size1 + allocation_size2) == value64); HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem2), stream.stream())); + HIP_CHECK(hipHostFree(notified)); } static bool thread_results[NUMBER_OF_THREADS]; diff --git a/catch/unit/memory/hipMemoryAllocateCoherent.cc b/catch/unit/memory/hipMemoryAllocateCoherent.cc deleted file mode 100644 index 8ba85043ff..0000000000 --- a/catch/unit/memory/hipMemoryAllocateCoherent.cc +++ /dev/null @@ -1,85 +0,0 @@ -/* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR -IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -/* -This testcase verifies the following scenario -1. Allocating the memory and modifying it coherently -*/ - -#include -#include -#include - -constexpr auto wait_sec = 5000; - -__global__ void Kernel(float* hostRes, int clkRate) { - int tid = threadIdx.x + blockIdx.x * blockDim.x; - hostRes[tid] = tid + 1; - __threadfence_system(); - // expecting that the data is getting flushed to host here! - uint64_t start = clock64()/clkRate, cur; - if (clkRate > 1) { - do { cur = clock64()/clkRate-start;}while (cur < wait_sec); - } else { - do { cur = clock64()/start;}while (cur < wait_sec); - } -} - -__global__ void Kernel_gfx11(float* hostRes, int clkRate) { -#if HT_AMD - int tid = threadIdx.x + blockIdx.x * blockDim.x; - hostRes[tid] = tid + 1; - __threadfence_system(); - // expecting that the data is getting flushed to host here! - uint64_t start = clock_function()/clkRate, cur; - if (clkRate > 1) { - do { cur = clock_function()/clkRate-start;}while (cur < wait_sec); - } else { - do { cur = clock_function()/start;}while (cur < wait_sec); - } -#endif -} - -TEST_CASE("Unit_hipHostMalloc_CoherentAccess") { - int blocks = 2; - float* hostRes; - HIP_CHECK(hipHostMalloc(&hostRes, blocks * sizeof(float), - hipHostMallocMapped)); - hostRes[0] = 0; - hostRes[1] = 0; - int clkRate; - if (IsGfx11()) { - HIPCHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0)); - } else { - HIPCHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0)); - } - std::cout << clkRate << std::endl; - auto Kernel_used = IsGfx11() ? Kernel_gfx11 : Kernel; - hipLaunchKernelGGL(HIP_KERNEL_NAME(Kernel_used), dim3(1), dim3(blocks), - 0, 0, hostRes, clkRate); - HIP_CHECK(hipGetLastError()); - int eleCounter = 0; - while (eleCounter < blocks) { - // blocks until the value changes - while (hostRes[eleCounter] == 0) {printf("waiting for counter inc\n");} - eleCounter++; - } - HIP_CHECK(hipHostFree(reinterpret_cast(hostRes))); -} - diff --git a/catch/unit/memory/hipSVMCommon.h b/catch/unit/memory/hipSVMCommon.h index 5b7433fc12..f1bfdc057c 100644 --- a/catch/unit/memory/hipSVMCommon.h +++ b/catch/unit/memory/hipSVMCommon.h @@ -17,8 +17,8 @@ * Modifications Copyright (C)2023 Advanced * Micro Devices, Inc. All rights reserved. */ -#ifndef __COMMON_H__ -#define __COMMON_H__ +#ifndef __HIPSVMCOMMON_H__ +#define __HIPSVMCOMMON_H__ #include #include @@ -137,5 +137,5 @@ inline void align_free(void* ptr) { #endif } -#endif // #ifndef __COMMON_H__ +#endif // #ifndef __HIPSVMCOMMON_H__ diff --git a/catch/unit/memory/mempool_common.hh b/catch/unit/memory/mempool_common.hh index 05876fe721..50d0a2f569 100644 --- a/catch/unit/memory/mempool_common.hh +++ b/catch/unit/memory/mempool_common.hh @@ -92,6 +92,13 @@ template __global__ void kernel_500ms_gfx11(T* host_res, int clk_ra #endif } +template __global__ void notifiedKernel(T* host_res, volatile unsigned int* notified) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + host_res[tid] = tid + 1; + __threadfence_system(); + while (*notified == 0) { } +} + template void MallocMemPoolAsync_OneAlloc(F malloc_func, const MemPools mempool_type) { int device_id = 0; HIP_CHECK(hipSetDevice(device_id)); @@ -102,7 +109,9 @@ template void MallocMemPoolAsync_OneAlloc(F malloc_func, const MemP SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); return; } - + unsigned int *notified = nullptr; + HIP_CHECK(hipHostMalloc(¬ified, sizeof(unsigned int))); + *notified = 0; const auto allocation_size = GENERATE(kPageSize / 2, kPageSize, kPageSize * 2); LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); MemPoolGuard mempool(mempool_type, device_id); @@ -114,16 +123,8 @@ template void MallocMemPoolAsync_OneAlloc(F malloc_func, const MemP stream.stream())); int blocks = 16; - int clk_rate; hipMemPoolAttr attr; - if (IsGfx11()) { - HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeWallClockRate, 0)); - kernel_500ms_gfx11<<<32, blocks, 0, stream.stream()>>>(alloc_mem, clk_rate); - } else { - HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeClockRate, 0)); - - kernel_500ms<<<32, blocks, 0, stream.stream()>>>(alloc_mem, clk_rate); - } + notifiedKernel<<<32, blocks, 0, stream.stream()>>>(alloc_mem, notified); const auto element_count = allocation_size / sizeof(int); constexpr auto thread_count = 1024; @@ -140,6 +141,7 @@ template void MallocMemPoolAsync_OneAlloc(F malloc_func, const MemP attr = hipMemPoolAttrReservedMemCurrent; std::uint64_t res_before_sync = 0; HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &res_before_sync)); + *notified = 1; HIP_CHECK(hipStreamSynchronize(stream.stream())); std::uint64_t res_after_sync = 0; @@ -153,6 +155,7 @@ template void MallocMemPoolAsync_OneAlloc(F malloc_func, const MemP REQUIRE(0 == used_mem); ArrayFindIfNot(host_alloc.host_ptr(), expected_value, element_count); + HIP_CHECK(hipHostFree(notified)); } template @@ -166,7 +169,9 @@ void MallocMemPoolAsync_TwoAllocs(F malloc_func, const MemPools mempool_type) { SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); return; } - + unsigned int *notified = nullptr; + HIP_CHECK(hipHostMalloc(¬ified, sizeof(unsigned int))); + *notified = 0; const auto allocation_size = GENERATE(kPageSize / 2, kPageSize, kPageSize * 2); LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); MemPoolGuard mempool(mempool_type, device_id); @@ -181,16 +186,8 @@ void MallocMemPoolAsync_TwoAllocs(F malloc_func, const MemPools mempool_type) { stream.stream())); int blocks = 16; - int clk_rate; hipMemPoolAttr attr; - if (IsGfx11()) { - HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeWallClockRate, 0)); - kernel_500ms_gfx11<<<32, blocks, 0, stream.stream()>>>(alloc_mem1, clk_rate); - } else { - HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeClockRate, 0)); - - kernel_500ms<<<32, blocks, 0, stream.stream()>>>(alloc_mem1, clk_rate); - } + notifiedKernel<<<32, blocks, 0, stream.stream()>>>(alloc_mem1, notified); const auto element_count = allocation_size / sizeof(int); constexpr auto thread_count = 1024; @@ -211,6 +208,7 @@ void MallocMemPoolAsync_TwoAllocs(F malloc_func, const MemPools mempool_type) { attr = hipMemPoolAttrReservedMemCurrent; std::uint64_t res_before_sync = 0; HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &res_before_sync)); + *notified = 1; HIP_CHECK(hipStreamSynchronize(stream.stream())); std::uint64_t res_after_sync = 0; @@ -238,6 +236,7 @@ void MallocMemPoolAsync_TwoAllocs(F malloc_func, const MemPools mempool_type) { REQUIRE(0 == used_mem); ArrayFindIfNot(host_alloc.host_ptr(), expected_value, element_count); + HIP_CHECK(hipHostFree(notified)); } template void MallocMemPoolAsync_Reuse(F malloc_func, const MemPools mempool_type) { @@ -250,7 +249,9 @@ template void MallocMemPoolAsync_Reuse(F malloc_func, const MemPool SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); return; } - + unsigned int *notified = nullptr; + HIP_CHECK(hipHostMalloc(¬ified, sizeof(unsigned int))); + *notified = 0; MemPoolGuard mempool(mempool_type, device_id); int *alloc_mem1, *alloc_mem2, *alloc_mem3; @@ -265,16 +266,8 @@ template void MallocMemPoolAsync_Reuse(F malloc_func, const MemPool stream.stream())); int blocks = 2; - int clk_rate; - if (IsGfx11()) { - HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeWallClockRate, 0)); - kernel_500ms_gfx11<<<32, blocks, 0, stream.stream()>>>(alloc_mem1, clk_rate); - } else { - HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeClockRate, 0)); - - kernel_500ms<<<32, blocks, 0, stream.stream()>>>(alloc_mem1, clk_rate); - } + notifiedKernel<<<32, blocks, 0, stream.stream()>>>(alloc_mem1, notified); hipMemPoolAttr attr; // Not a real free, since kernel isn't done @@ -286,15 +279,12 @@ template void MallocMemPoolAsync_Reuse(F malloc_func, const MemPool REQUIRE(alloc_mem1 == alloc_mem2); // Make a sync before the second kernel launch to make sure memory B isn't gone + *notified = 1; HIP_CHECK(hipStreamSynchronize(stream.stream())); - + *notified = 0; // Second kernel launch with new memory - if (IsGfx11()) { - kernel_500ms_gfx11<<<32, blocks, 0, stream.stream()>>>(alloc_mem2, clk_rate); - } else { - kernel_500ms<<<32, blocks, 0, stream.stream()>>>(alloc_mem2, clk_rate); - } - + notifiedKernel<<<32, blocks, 0, stream.stream()>>>(alloc_mem2, notified); + *notified = 1; HIP_CHECK(hipStreamSynchronize(stream.stream())); attr = hipMemPoolAttrUsedMemCurrent; @@ -315,6 +305,7 @@ template void MallocMemPoolAsync_Reuse(F malloc_func, const MemPool REQUIRE(allocation_size2 == value64); HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem3), stream.stream())); + HIP_CHECK(hipHostFree(notified)); } // definitions diff --git a/catch/unit/streamperthread/hipStreamPerThrdCompilerOptn.cc b/catch/unit/streamperthread/hipStreamPerThrdCompilerOptn.cc index 9590d6943f..6da1cc3c16 100644 --- a/catch/unit/streamperthread/hipStreamPerThrdCompilerOptn.cc +++ b/catch/unit/streamperthread/hipStreamPerThrdCompilerOptn.cc @@ -26,7 +26,7 @@ THE SOFTWARE. #include #include #include "hip/hip_cooperative_groups.h" - +#include namespace cg = cooperative_groups; namespace DefltStrmPT { @@ -44,7 +44,7 @@ namespace DefltStrmPT { } // namespace DefltStrmPT __device__ int64_t globalInDStrmPT[1024 * 1024]; -__device__ int SigComplte = 0; +__managed__ int SigComplte = 0; // Kernel codes __global__ void DefltStrmPT_Square(int64_t *C_d, int64_t N) { @@ -68,6 +68,12 @@ __global__ void Wait_Kernel3(int clockrate, uint64_t WaitSecs, } } +static __global__ void notifiedKernel(volatile unsigned int *notified, int PassSignal = 0) { + while (*notified == 0) {} // wait until notified to exit. + if (PassSignal) { + SigComplte = 1; + } +} __global__ void DefltStrmPT_Test_gws(uint* buf, uint bufSize, int64_t* tmpBuf, int64_t* result) { extern __shared__ int64_t tmp[]; @@ -216,7 +222,7 @@ void PerThrdDefltStrm_Memset3D(int Async) { HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent)); A_h = reinterpret_cast(malloc(sizeElements)); - REQUIRE(A_h != nullptr); + if (A_h == nullptr) REQUIRE(false); for (size_t i = 0; i < elements; i++) { A_h[i] = 1; @@ -257,16 +263,18 @@ void PerThrdDefltStrm_Memset3D(int Async) { void DefaultPT2_StrmQuery() { - HIP_CHECK(hipDeviceGetAttribute(&(DefltStrmPT::clockrate), - hipDeviceAttributeMemoryClockRate, 0)); + unsigned int *notified = nullptr; + HIP_CHECK(hipHostMalloc(¬ified, sizeof(unsigned int))); + *notified = 0; HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm))); - // StreamQuery with null stream - Wait_Kernel3<<<1, 1>>>(DefltStrmPT::clockrate, 3); - REQUIRE((hipErrorNotReady == hipStreamQuery(0))); // StreamQuery with user created stream - Wait_Kernel3<<<1, 1, 0, DefltStrmPT::Strm>>>(DefltStrmPT::clockrate, 3); + notifiedKernel<<<1, 1, 0, DefltStrmPT::Strm>>>(notified); REQUIRE((hipErrorNotReady == hipStreamQuery(DefltStrmPT::Strm))); + + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified = 1; HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm)); + HIP_CHECK(hipHostFree(notified)); } @@ -285,33 +293,52 @@ void DefaultPT2_StreamSync() { void DefaultPT2_StrmWaitEvent() { + int device; + HIP_CHECK(hipGetDevice(&device)); + if (!DeviceAttributesSupport(device, hipDeviceAttributeManagedMemory)) { + HipTest::HIP_SKIP_TEST("Managed memory is not supported"); + return; + } + hipEvent_t evt; hipStream_t Strm1; + unsigned int *notified = nullptr; + HIP_CHECK(hipHostMalloc(¬ified, sizeof(unsigned int))); + *notified = 0; HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm))); HIP_CHECK(hipStreamCreate(&Strm1)); HIP_CHECK(hipEventCreate(&evt)); - Wait_Kernel3<<<1, 1, 0, DefltStrmPT::Strm>>>(DefltStrmPT::clockrate, 3, 1); + notifiedKernel<<<1, 1, 0, DefltStrmPT::Strm>>>(notified, 1); HIP_CHECK(hipEventRecord(evt, DefltStrmPT::Strm)); HIP_CHECK(hipStreamWaitEvent(Strm1, evt, 0)); - Wait_Kernel3<<<1, 1, 0, Strm1>>>(DefltStrmPT::clockrate, 1); + notifiedKernel<<<1, 1, 0, Strm1>>>(notified); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); // By the time control reaches the below point SigComplte is expected // to be still zero if (SigComplte) { REQUIRE(false); } + *notified = 1; HIP_CHECK(hipStreamSynchronize(Strm1)); HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm)); + if (SigComplte == 0) { + REQUIRE(false); + } HIP_CHECK(hipStreamDestroy(Strm1)); HIP_CHECK(hipEventDestroy(evt)); + HIP_CHECK(hipHostFree(notified)); } void DefaultPT2_EvtQuery() { hipEvent_t evt, evt1; hipError_t err; + unsigned int *notified = nullptr; + HIP_CHECK(hipHostMalloc(¬ified, sizeof(unsigned int))); + *notified = 0; HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm))); HIP_CHECK(hipEventCreate(&evt)); HIP_CHECK(hipEventCreate(&evt1)); - Wait_Kernel3<<<1, 1, 0, DefltStrmPT::Strm>>>(DefltStrmPT::clockrate, 3); + notifiedKernel<<<1, 1, 0, DefltStrmPT::Strm>>>(notified); HIP_CHECK(hipEventRecord(evt, DefltStrmPT::Strm)); err = hipEventQuery(evt); if (err != hipErrorNotReady) { @@ -319,8 +346,10 @@ void DefaultPT2_EvtQuery() { } // Testing for Null or default stream HIP_CHECK(hipEventRecord(evt1, 0)); - std::chrono::time_point start = std::chrono::steady_clock::now(); int Got_hipSuccess = 0; // 0 for no, 1 for yes + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + *notified = 1; // notify to exit + std::chrono::time_point start = std::chrono::steady_clock::now(); while (true) { err = hipEventQuery(evt1); if (err == hipSuccess) { @@ -337,6 +366,7 @@ void DefaultPT2_EvtQuery() { HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm)); HIP_CHECK(hipEventDestroy(evt)); HIP_CHECK(hipEventDestroy(evt1)); + HIP_CHECK(hipHostFree(notified)); } @@ -632,9 +662,9 @@ float DefaultPT2_hipMemcpy2DFromArray() { HIP_CHECK(hipMemcpy2DToArray(Dptr, 0, 0, Hptr_A, DefltStrmPT::width, DefltStrmPT::width, DefltStrmPT::numH, hipMemcpyHostToDevice)); - Wait_Kernel3 <<< 1, 1, 0, DefltStrmPT::Strm >>> (DefltStrmPT::clockrate, + Wait_Kernel3 <<< 1, 1, 0, DefltStrmPT::Strm >>> (DefltStrmPT::clockrate, 1); - HIP_CHECK(hipMemcpy2DFromArray(Hptr_B, DefltStrmPT::width, Dptr, 0, 0, + HIP_CHECK(hipMemcpy2DFromArray(Hptr_B, DefltStrmPT::width, Dptr, 0, 0, DefltStrmPT::width, DefltStrmPT::numH, hipMemcpyDeviceToHost)); HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm)); HIP_CHECK(hipFreeArray(Dptr));