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

SWDEV-514141 - Fix zero clock rate issues (#4)

Browse Source 
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
This commit is contained in:
Sang, Tao
2025-06-11 11:41:25 -04:00
committed by GitHub
parent c4284d50c5
commit ec8ff45a1d
15 changed files with 332 additions and 581 deletions
Download Patch File Download Diff File Expand all files Collapse all files
catch/include/hip_test_common.hh
+16
View File
@@ -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 \
catch/multiproc/CMakeLists.txt
+20 -20
View File
@@ -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()
catch/multiproc/hipMemCoherencyTstMProc.cc
+49 -124
View File
@@ -39,37 +39,15 @@
#include <sys/mman.h>
#include <sys/wait.h>
#include <chrono>
#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<void **>(&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<void **>(&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<int*>(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<int*>(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);
}
}
catch/perftests/graph/parallelGraph.cc
+4 -49
View File
@@ -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 <typename T>
__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<hipGraphNode_t> 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<void *>(&NElem),
reinterpret_cast<void *>(&clkRate)};
void *kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kernelNodeParams.func = reinterpret_cast<void *>(vectorADD<int>);
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));
catch/perftests/stream/hipPerfDeviceConcurrency.cc
+4
View File
@@ -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
catch/perftests/stream/hipPerfStreamConcurrency.cc
+4 -1
View File
@@ -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
catch/unit/clock/hipClockCheck.cc
+9 -3
View File
@@ -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;
}
catch/unit/memory/CMakeLists.txt
+1
View File
@@ -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()
catch/unit/memory/hipMemPoolApi.cc
+81 -137
View File
@@ -21,7 +21,7 @@
1) This testcase verifies the basic scenario - supported on
all devices
*/
#include "mempool_common.hh"
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#include <hip_test_checkers.hh>
@@ -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(&notified, 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<void**>(&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<void*>(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<void*>(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(&notified, 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<void**>(&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<void*>(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(&notified, 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<void**>(&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<void*>(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(&notified1, sizeof(unsigned int)));
HIP_CHECK(hipHostMalloc(&notified2, 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<void**>(&C), numElements * sizeof(float), mem_pool, stream));
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&C), numElements * sizeof(float), mem_pool, stream1));
int value = 0;
SECTION("Disallow Opportunistic - No Reuse") {
numElements = 8 * 1024 * 1024;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&A), numElements * sizeof(float), mem_pool, stream));
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&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<void*>(A), stream));
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(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<void*>(B), stream2));
@@ -466,24 +424,20 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") {
SECTION("Allow Opportunistic - Reuse") {
numElements = 8 * 1024 * 1024;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&A), numElements * sizeof(float), mem_pool, stream));
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&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<void*>(A), stream));
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(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<void*>(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<void**>(&A), numElements * sizeof(float), mem_pool, stream));
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&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<void*>(A), stream));
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(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<void*>(B), stream2));
}
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(C), stream));
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(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(&notified, 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<void**>(&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<void*>(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<void*>(C), stream));
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipHostFree(notified));
}
catch/unit/memory/hipMemPoolSetGetAttribute.cc
+61 -97
View File
@@ -19,7 +19,6 @@
#include "mempool_common.hh"
#include <resource_guards.hh>
#include <utils.hh>
/**
* @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(&notified1, sizeof(unsigned int)));
HIP_CHECK(hipHostMalloc(&notified2, 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<void*>(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<void*>(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<void**>(&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<void*>(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<void*>(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<void*>(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<void*>(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(&notified1, sizeof(unsigned int)));
HIP_CHECK(hipHostMalloc(&notified2, 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<void*>(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<void*>(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<void*>(alloc_mem3), stream1.stream()));
HIP_CHECK(hipEventDestroy(event));
HIP_CHECK(hipHostFree(notified1));
HIP_CHECK(hipHostFree(notified2));
}
/**
catch/unit/memory/hipMemPoolTrimTo.cc
+7 -10
View File
@@ -22,7 +22,6 @@
#include <resource_guards.hh>
#include <utils.hh>
/**
* @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(&notified, 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<void*>(alloc_mem2), stream.stream()));
HIP_CHECK(hipHostFree(notified));
}
static bool thread_results[NUMBER_OF_THREADS];
catch/unit/memory/hipMemoryAllocateCoherent.cc
-85
View File
@@ -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 <hip_test_common.hh>
#include <hip_test_kernels.hh>
#include <hip_test_checkers.hh>
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<void *>(hostRes)));
}
catch/unit/memory/hipSVMCommon.h
+3 -3
View File
@@ -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 <vector>
#include <string>
@@ -137,5 +137,5 @@ inline void align_free(void* ptr) {
#endif
}
#endif // #ifndef __COMMON_H__
#endif // #ifndef __HIPSVMCOMMON_H__
catch/unit/memory/mempool_common.hh
+28 -37
View File
@@ -92,6 +92,13 @@ template <typename T> __global__ void kernel_500ms_gfx11(T* host_res, int clk_ra
#endif
}
template <typename T> __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 <typename F> void MallocMemPoolAsync_OneAlloc(F malloc_func, const MemPools mempool_type) {
int device_id = 0;
HIP_CHECK(hipSetDevice(device_id));
@@ -102,7 +109,9 @@ template <typename F> 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(&notified, sizeof(unsigned int)));
*notified = 0;
const auto allocation_size = GENERATE(kPageSize / 2, kPageSize, kPageSize * 2);
LinearAllocGuard<int> host_alloc(LinearAllocs::hipHostMalloc, allocation_size);
MemPoolGuard mempool(mempool_type, device_id);
@@ -114,16 +123,8 @@ template <typename F> 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 <typename F> 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 <typename F> 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 <typename F>
@@ -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(&notified, sizeof(unsigned int)));
*notified = 0;
const auto allocation_size = GENERATE(kPageSize / 2, kPageSize, kPageSize * 2);
LinearAllocGuard<int> 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 <typename F> void MallocMemPoolAsync_Reuse(F malloc_func, const MemPools mempool_type) {
@@ -250,7 +249,9 @@ template <typename F> 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(&notified, sizeof(unsigned int)));
*notified = 0;
MemPoolGuard mempool(mempool_type, device_id);
int *alloc_mem1, *alloc_mem2, *alloc_mem3;
@@ -265,16 +266,8 @@ template <typename F> 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 <typename F> 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 <typename F> void MallocMemPoolAsync_Reuse(F malloc_func, const MemPool
REQUIRE(allocation_size2 == value64);
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(alloc_mem3), stream.stream()));
HIP_CHECK(hipHostFree(notified));
}
// definitions
catch/unit/streamperthread/hipStreamPerThrdCompilerOptn.cc
+45 -15
View File
@@ -26,7 +26,7 @@ THE SOFTWARE.
#include <ctime>
#include <hip_test_common.hh>
#include "hip/hip_cooperative_groups.h"
#include <utils.hh>
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<char *>(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(&notified, 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(&notified, 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(&notified, 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));