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

SWDEV-308368 - Fix coherent test bugs (#2985)

Browse Source 
1. Remove hipDeviceAttributePageableMemoryAccess checking as it is used to check xnack+ rather than fine-grained synchronization.
Note: hipHostMalloc() will alloc memory with fine-grained access and atomic by default when env HIP_HOST_COHERENT=1. For more information, see hip_programming_guide.
2. Fix some other bugs.

Change-Id: Ieaaed1b15b4d258193d834104b17d9f03a9e4781

[ROCm/hip commit: 15517aaf18]
This commit is contained in:
ROCm CI Service Account
2022-10-19 16:48:44 +05:30
committed by GitHub
parent 3fd7a5dc21
commit 6a87199474
4 changed files with 117 additions and 175 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/tests/catch/multiproc/hipMemCoherencyTstMProc.cc
+94 -138
View File
@@ -40,15 +40,12 @@
__global__ void CoherentTst(int *ptr, int PeakClk) {
// Incrementing the value by 1
int64_t GpuFrq = (PeakClk * 1000);
int64_t GpuFrq = int64_t(PeakClk) * 1000;
int64_t StrtTck = clock64();
atomicAdd(ptr, 1);
// The following while loop checks the value in ptr for around 3-4 seconds
while ((clock64() - StrtTck) <= (3 * GpuFrq)) {
if (*ptr == 3) {
atomicAdd(ptr, 1);
return;
}
if (atomicCAS(ptr, 3, 4) == 3) break;
}
}
@@ -57,8 +54,6 @@ __global__ void SquareKrnl(int *ptr) {
*ptr = (*ptr) * (*ptr);
}
// The variable below will work as signal to decide pass/fail
static bool YES_COHERENT = false;
@@ -641,40 +636,31 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1") {
WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
REQUIRE(false);
}
int stat = 0, Pageable = 0;
int stat = 0;
HIP_CHECK(hipDeviceGetAttribute(&Pageable,
hipDeviceAttributePageableMemoryAccess, 0));
INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
if (Pageable) {
if (fork() == 0) { // child process
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);
}
} else { // parent process
wait(&stat);
int Result = WEXITSTATUS(stat);
if (Result != 10) {
REQUIRE(false);
}
if (fork() == 0) { // child process
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);
}
} else { // parent process
wait(&stat);
int Result = WEXITSTATUS(stat);
if (Result != 10) {
REQUIRE(false);
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
"attribute. Hence skipping the test with Pass result.\n");
}
}
#endif
@@ -689,40 +675,31 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg1") {
WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
REQUIRE(false);
}
int stat = 0, Pageable = 0;
int stat = 0;
HIP_CHECK(hipDeviceGetAttribute(&Pageable,
hipDeviceAttributePageableMemoryAccess, 0));
INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
if (Pageable) {
if (fork() == 0) { // child process
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);
}
} else { // parent process
wait(&stat);
int Result = WEXITSTATUS(stat);
if (Result != 10) {
REQUIRE(false);
}
if (fork() == 0) { // child process
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);
}
} else { // parent process
wait(&stat);
int Result = WEXITSTATUS(stat);
if (Result != 10) {
REQUIRE(false);
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
"attribute. Hence skipping the test with Pass result.\n");
}
}
#endif
@@ -736,45 +713,35 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg2") {
WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
REQUIRE(false);
}
int stat = 0, Pageable = 0;
int stat = 0;
HIP_CHECK(hipDeviceGetAttribute(&Pageable,
hipDeviceAttributePageableMemoryAccess, 0));
INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
if (Pageable) {
if (fork() == 0) { // child process
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);
}
} else { // parent process
wait(&stat);
int Result = WEXITSTATUS(stat);
if (Result != 10) {
REQUIRE(false);
}
if (fork() == 0) { // child process
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);
}
} else { // parent process
wait(&stat);
int Result = WEXITSTATUS(stat);
if (Result != 10) {
REQUIRE(false);
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
"attribute. Hence skipping the test with Pass result.\n");
}
}
#endif
/* Test Case Description: The following test checks if the memory exhibits
fine grain behavior when HIP_HOST_COHERENT is set to 1*/
// The following test is AMD specific test hence skipping for Nvidia
@@ -784,42 +751,31 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg3") {
WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
REQUIRE(false);
}
int stat = 0, Pageable = 0;
int stat = 0;
HIP_CHECK(hipDeviceGetAttribute(&Pageable,
hipDeviceAttributePageableMemoryAccess, 0));
INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
if (Pageable) {
if (fork() == 0) { // child process
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);
}
} else { // parent process
wait(&stat);
int Result = WEXITSTATUS(stat);
if (Result != 10) {
REQUIRE(false);
}
if (fork() == 0) { // child process
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);
}
} else { // parent process
wait(&stat);
int Result = WEXITSTATUS(stat);
if (Result != 10) {
REQUIRE(false);
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
"attribute. Hence skipping the test with Pass result.\n");
}
}
#endif
projects/hip/tests/catch/unit/memory/hipMemAdvise.cc
+5 -5
View File
@@ -777,7 +777,7 @@ TEST_CASE("Unit_hipMemAdvise_TstMemAdvisePrefrdLoc") {
REQUIRE(false);
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
@@ -809,7 +809,7 @@ TEST_CASE("Unit_hipMemAdvise_TstMemAdviseLstPreftchLoc") {
REQUIRE(false);
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
} else {
@@ -847,7 +847,7 @@ TEST_CASE("Unit_hipMemAdvise_TstMemAdviseMultiFlag") {
}
HIP_CHECK(hipFree(Hmm));
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
@@ -929,7 +929,7 @@ TEST_CASE("Unit_hipMemAdvise_ReadMosltyMgpuTst") {
HIP_CHECK(hipFree(Hmm));
HIP_CHECK(hipStreamDestroy(strm));
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
@@ -955,7 +955,7 @@ TEST_CASE("Unit_hipMemAdvise_TstSetUnsetPrfrdLoc") {
REQUIRE(false);
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
projects/hip/tests/catch/unit/memory/hipMemAdviseMmap.cc
+1 -1
View File
@@ -50,7 +50,7 @@ TEST_CASE("Unit_hipMemAdvise_MmapMem") {
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
0));
INFO("hipDeviceAttributeManagedMemory: " << managed);
if ((managed == 1) && (PageableMem == 1)) {
if (PageableMem == 1) {
#ifdef __linux__
// For now this test is enabled only for linux platforms
FILE *fptr;
projects/hip/tests/catch/unit/memory/hipMemCoherencyTst.cc
+17 -31
View File
@@ -89,33 +89,24 @@ static void TstCoherency(int* ptr, bool hmmMem) {
// passing
#if HT_AMD
TEST_CASE("Unit_hipHostMalloc_CoherentTst") {
int *Ptr = nullptr, SIZE = sizeof(int), Pageable = 0;
int *Ptr = nullptr, SIZE = sizeof(int);
bool HmmMem = false;
YES_COHERENT = false;
HIP_CHECK(hipDeviceGetAttribute(&Pageable,
hipDeviceAttributePageableMemoryAccess, 0));
INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
if (Pageable == 1) {
// Allocating hipHostMalloc() memory with hipHostMallocCoherent flag
SECTION("hipHostMalloc with hipHostMallocCoherent flag") {
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocCoherent));
}
SECTION("hipHostMalloc with Default flag") {
HIP_CHECK(hipHostMalloc(&Ptr, SIZE));
}
SECTION("hipHostMalloc with hipHostMallocMapped flag") {
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocMapped));
}
TstCoherency(Ptr, HmmMem);
HIP_CHECK(hipHostFree(Ptr));
REQUIRE(YES_COHERENT);
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
"attribute. Hence skipping the test with Pass result.\n");
// Allocating hipHostMalloc() memory with hipHostMallocCoherent flag
SECTION("hipHostMalloc with hipHostMallocCoherent flag") {
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocCoherent));
}
SECTION("hipHostMalloc with Default flag") {
HIP_CHECK(hipHostMalloc(&Ptr, SIZE));
}
SECTION("hipHostMalloc with hipHostMallocMapped flag") {
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocMapped));
}
TstCoherency(Ptr, HmmMem);
HIP_CHECK(hipHostFree(Ptr));
REQUIRE(YES_COHERENT);
}
#endif
@@ -126,19 +117,14 @@ TEST_CASE("Unit_hipHostMalloc_CoherentTst") {
// passing
#if HT_AMD
TEST_CASE("Unit_hipMallocManaged_CoherentTst") {
int *Ptr = nullptr, SIZE = sizeof(int), Pageable = 0, managed = 0;
int *Ptr = nullptr, SIZE = sizeof(int), managed = 0;
bool HmmMem = true;
YES_COHERENT = false;
HIP_CHECK(hipDeviceGetAttribute(&Pageable,
hipDeviceAttributePageableMemoryAccess, 0));
INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
0));
INFO("hipDeviceAttributeManagedMemory: " << managed);
if (managed == 1 && Pageable == 1) {
if (managed == 1) {
// Allocating hipMallocManaged() memory
SECTION("hipMallocManaged with hipMemAttachGlobal flag") {
HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachGlobal));
@@ -150,7 +136,7 @@ TEST_CASE("Unit_hipMallocManaged_CoherentTst") {
HIP_CHECK(hipFree(Ptr));
REQUIRE(YES_COHERENT);
} else {
SUCCEED("GPU 0 doesn't support ManagedMemory or PageableMemoryAccess"
SUCCEED("GPU 0 doesn't support ManagedMemory "
"device attribute. Hence skipping the test with Pass result.\n");
}
}