AI/rocm-systems
2
0
Bifurcation 0
Code Tickets Demandes d'ajout Actions Paquets Projets Publications Wiki Activité

SWDEV-329590 - Addresses coherency test failures on Navi21 (#2652)

Parcourir la source 
- Tests using HSA_XNACK=1 restricted to MI200
- Added checks to verify HMM and Pageable memory access attributes to avoid failures

Change-Id: Ic0b107264378ce0c4f0aab770c941ae2b57342c2
Cette révision appartient à :
ROCm CI Service Account
2022-05-17 11:31:19 +05:30
révisé par GitHub
Parent ee541eaef1
révision 80e010bd0e
2 fichiers modifiés avec 242 ajouts et 218 suppressions
Télécharger le Fichier Patch Télécharger le Fichier des Différences Développer tous les fichiers Réduire tous les fichiers
tests/catch/multiproc/hipMemCoherencyTstMProc.cc
+165 -149
Voir le fichier
@@ -113,9 +113,8 @@ TEST_CASE("Unit_malloc_CoherentTst") {
REQUIRE(false); REQUIRE(false);
} }
/* GpuId[0] for gfx906 exists--> 1 for yes and 0 for no /* GpuId[0] for gfx90a exists--> 1 for yes and 0 for no*/
GpuId[0] for gfx908 exists--> 1 for yes and 0 for no*/ int GpuId[1] = {0};
int GpuId[2] = {0, 0};
p = fork(); p = fork();
if (p < 0) { if (p < 0) {
@@ -128,8 +127,8 @@ TEST_CASE("Unit_malloc_CoherentTst") {
// Read string from child and close reading end. // Read string from child and close reading end.
read(fd1[0], GpuId, 2 * sizeof(int)); read(fd1[0], GpuId, 2 * sizeof(int));
close(fd1[0]); close(fd1[0]);
if ((GpuId[0] == 1) || (GpuId[1] == 1)) { if (GpuId[0] == 0) {
WARN("This test is not applicable on MI60 & MI100." WARN("This test is applicable for MI200."
"Skipping the test!!"); "Skipping the test!!");
exit(0); exit(0);
} }
@@ -138,16 +137,11 @@ TEST_CASE("Unit_malloc_CoherentTst") {
hipDeviceProp_t prop; hipDeviceProp_t prop;
HIPCHECK(hipGetDeviceProperties(&prop, 0)); HIPCHECK(hipGetDeviceProperties(&prop, 0));
char *p = NULL; char *p = NULL;
p = strstr(prop.gcnArchName, "gfx906"); p = strstr(prop.gcnArchName, "gfx90a");
if (p) { if (p) {
WARN("gfx906 gpu found on this system!!"); WARN("gfx90a gpu found on this system!!");
GpuId[0] = 1; GpuId[0] = 1;
} }
p = strstr(prop.gcnArchName, "gfx908");
if (p) {
WARN("gfx908 gpu found on this system!!");
GpuId[1] = 1;
}
// Write concatenated string and close writing end // Write concatenated string and close writing end
write(fd1[1], GpuId, 2 * sizeof(int)); write(fd1[1], GpuId, 2 * sizeof(int));
close(fd1[1]); close(fd1[1]);
@@ -208,9 +202,8 @@ TEST_CASE("Unit_malloc_CoherentTstWthAdvise") {
REQUIRE(false); REQUIRE(false);
} }
/* GpuId[0] for gfx906 exists--> 1 for yes and 0 for no /* GpuId[0] for gfx90a exists--> 1 for yes and 0 for no */
GpuId[0] for gfx908 exists--> 1 for yes and 0 for no*/ int GpuId[1] = {0};
int GpuId[2] = {0, 0};
p = fork(); p = fork();
if (p < 0) { if (p < 0) {
@@ -223,8 +216,8 @@ TEST_CASE("Unit_malloc_CoherentTstWthAdvise") {
// Read string from child and close reading end. // Read string from child and close reading end.
read(fd1[0], GpuId, 2 * sizeof(int)); read(fd1[0], GpuId, 2 * sizeof(int));
close(fd1[0]); close(fd1[0]);
if ((GpuId[0] == 1) || (GpuId[1] == 1)) { if (GpuId[0] == 0) {
WARN("This test is not applicable on MI60 & MI100." WARN("This test is applicable for MI200."
"Skipping the test!!"); "Skipping the test!!");
exit(0); exit(0);
} }
@@ -233,16 +226,11 @@ TEST_CASE("Unit_malloc_CoherentTstWthAdvise") {
hipDeviceProp_t prop; hipDeviceProp_t prop;
HIPCHECK(hipGetDeviceProperties(&prop, 0)); HIPCHECK(hipGetDeviceProperties(&prop, 0));
char *p = NULL; char *p = NULL;
p = strstr(prop.gcnArchName, "gfx906"); p = strstr(prop.gcnArchName, "gfx90a");
if (p) { if (p) {
WARN("gfx906 gpu found on this system!!"); WARN("gfx90a gpu found on this system!!");
GpuId[0] = 1; GpuId[0] = 1;
} }
p = strstr(prop.gcnArchName, "gfx908");
if (p) {
WARN("gfx908 gpu found on this system!!");
GpuId[1] = 1;
}
// Write concatenated string and close writing end // Write concatenated string and close writing end
write(fd1[1], GpuId, 2 * sizeof(int)); write(fd1[1], GpuId, 2 * sizeof(int));
close(fd1[1]); close(fd1[1]);
@@ -305,9 +293,8 @@ TEST_CASE("Unit_mmap_CoherentTst") {
REQUIRE(false); REQUIRE(false);
} }
/* GpuId[0] for gfx906 exists--> 1 for yes and 0 for no /* GpuId[0] for gfx90a exists--> 1 for yes and 0 for no */
GpuId[0] for gfx908 exists--> 1 for yes and 0 for no*/ int GpuId[1] = {0};
int GpuId[2] = {0, 0};
p = fork(); p = fork();
if (p < 0) { if (p < 0) {
@@ -320,8 +307,8 @@ TEST_CASE("Unit_mmap_CoherentTst") {
// Read string from child and close reading end. // Read string from child and close reading end.
read(fd1[0], GpuId, 2 * sizeof(int)); read(fd1[0], GpuId, 2 * sizeof(int));
close(fd1[0]); close(fd1[0]);
if ((GpuId[0] == 1) || (GpuId[1] == 1)) { if (GpuId[0] == 0) {
WARN("This test is not applicable on MI60 & MI100." WARN("This test is not applicable for MI200."
"Skipping the test!!"); "Skipping the test!!");
exit(0); exit(0);
} }
@@ -330,16 +317,11 @@ TEST_CASE("Unit_mmap_CoherentTst") {
hipDeviceProp_t prop; hipDeviceProp_t prop;
HIPCHECK(hipGetDeviceProperties(&prop, 0)); HIPCHECK(hipGetDeviceProperties(&prop, 0));
char *p = NULL; char *p = NULL;
p = strstr(prop.gcnArchName, "gfx906"); p = strstr(prop.gcnArchName, "gfx90a");
if (p) { if (p) {
WARN("gfx906 gpu found on this system!!"); WARN("gfx90a gpu found on this system!!");
GpuId[0] = 1; GpuId[0] = 1;
} }
p = strstr(prop.gcnArchName, "gfx908");
if (p) {
WARN("gfx908 gpu found on this system!!");
GpuId[1] = 1;
}
// Write concatenated string and close writing end // Write concatenated string and close writing end
write(fd1[1], GpuId, 2 * sizeof(int)); write(fd1[1], GpuId, 2 * sizeof(int));
close(fd1[1]); close(fd1[1]);
@@ -403,9 +385,8 @@ TEST_CASE("Unit_mmap_CoherentTstWthAdvise") {
REQUIRE(false); REQUIRE(false);
} }
/* GpuId[0] for gfx906 exists--> 1 for yes and 0 for no /* GpuId[0] for gfx90a exists--> 1 for yes and 0 for no */
GpuId[0] for gfx908 exists--> 1 for yes and 0 for no*/ int GpuId[1] = {0};
int GpuId[2] = {0, 0};
p = fork(); p = fork();
if (p < 0) { if (p < 0) {
@@ -418,8 +399,8 @@ TEST_CASE("Unit_mmap_CoherentTstWthAdvise") {
// Read string from child and close reading end. // Read string from child and close reading end.
read(fd1[0], GpuId, 2 * sizeof(int)); read(fd1[0], GpuId, 2 * sizeof(int));
close(fd1[0]); close(fd1[0]);
if ((GpuId[0] == 1) || (GpuId[1] == 1)) { if (GpuId[0] == 0) {
WARN("This test is not applicable on MI60 & MI100." WARN("This test is applicable for MI200."
"Skipping the test!!"); "Skipping the test!!");
exit(0); exit(0);
} }
@@ -428,16 +409,11 @@ TEST_CASE("Unit_mmap_CoherentTstWthAdvise") {
hipDeviceProp_t prop; hipDeviceProp_t prop;
HIPCHECK(hipGetDeviceProperties(&prop, 0)); HIPCHECK(hipGetDeviceProperties(&prop, 0));
char *p = NULL; char *p = NULL;
p = strstr(prop.gcnArchName, "gfx906"); p = strstr(prop.gcnArchName, "gfx90a");
if (p) { if (p) {
WARN("gfx906 gpu found on this system!!"); WARN("gfx90a gpu found on this system!!");
GpuId[0] = 1; GpuId[0] = 1;
} }
p = strstr(prop.gcnArchName, "gfx908");
if (p) {
WARN("gfx908 gpu found on this system!!");
GpuId[1] = 1;
}
// Write concatenated string and close writing end // Write concatenated string and close writing end
write(fd1[1], GpuId, 2 * sizeof(int)); write(fd1[1], GpuId, 2 * sizeof(int));
close(fd1[1]); close(fd1[1]);
@@ -497,7 +473,7 @@ TEST_CASE("Unit_mmap_CoherentTstWthAdvise") {
#if HT_AMD #if HT_AMD
TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg1") { TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg1") {
if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) { if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) {
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!"); WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
REQUIRE(false); REQUIRE(false);
} }
int stat = 0; int stat = 0;
@@ -538,7 +514,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg1") {
#if HT_AMD #if HT_AMD
TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg2") { TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg2") {
if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) { if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) {
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!"); WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
REQUIRE(false); REQUIRE(false);
} }
int stat = 0; int stat = 0;
@@ -579,7 +555,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg2") {
#if HT_AMD #if HT_AMD
TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg3") { TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg3") {
if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) { if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) {
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!"); WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
REQUIRE(false); REQUIRE(false);
} }
int stat = 0; int stat = 0;
@@ -587,7 +563,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg3") {
int *Ptr = nullptr, *PtrD = nullptr, SIZE = sizeof(int); int *Ptr = nullptr, *PtrD = nullptr, SIZE = sizeof(int);
YES_COHERENT = false; YES_COHERENT = false;
// Allocating hipHostMalloc() memory // Allocating hipHostMalloc() memory
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNumaUser)); HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNumaUser));
*Ptr = 4; *Ptr = 4;
hipStream_t strm; hipStream_t strm;
HIP_CHECK(hipStreamCreate(&strm)); HIP_CHECK(hipStreamCreate(&strm));
@@ -620,7 +596,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg3") {
#if HT_AMD #if HT_AMD
TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg4") { TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg4") {
if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) { if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) {
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!"); WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
REQUIRE(false); REQUIRE(false);
} }
int stat = 0; int stat = 0;
@@ -662,33 +638,43 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg4") {
#if HT_AMD #if HT_AMD
TEST_CASE("Unit_hipHostMalloc_WthEnv1") { TEST_CASE("Unit_hipHostMalloc_WthEnv1") {
if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) { if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) {
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!"); WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
REQUIRE(false); REQUIRE(false);
} }
int stat = 0; int stat = 0, Pageable = 0;
if (fork() == 0) { // child process
int *Ptr = nullptr, SIZE = sizeof(int); HIP_CHECK(hipDeviceGetAttribute(&Pageable,
bool HmmMem = false; hipDeviceAttributePageableMemoryAccess, 0));
YES_COHERENT = false; INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
// Allocating hipHostMalloc() memory
HIP_CHECK(hipHostMalloc(&Ptr, SIZE)); if (Pageable) {
*Ptr = 4; if (fork() == 0) { // child process
TstCoherency(Ptr, HmmMem); int *Ptr = nullptr, SIZE = sizeof(int);
if (YES_COHERENT) { bool HmmMem = false;
// exit() with code 10 which indicates pass YES_COHERENT = false;
HIP_CHECK(hipHostFree(Ptr)); // Allocating hipHostMalloc() memory
exit(10); HIP_CHECK(hipHostMalloc(&Ptr, SIZE));
} else { *Ptr = 4;
// exit() with code 9 which indicates fail TstCoherency(Ptr, HmmMem);
HIP_CHECK(hipHostFree(Ptr)); if (YES_COHERENT) {
exit(9); // exit() with code 10 which indicates pass
} HIP_CHECK(hipHostFree(Ptr));
} else { // parent process exit(10);
wait(&stat); } else {
int Result = WEXITSTATUS(stat); // exit() with code 9 which indicates fail
if (Result != 10) { HIP_CHECK(hipHostFree(Ptr));
REQUIRE(false); 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 #endif
@@ -700,33 +686,43 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1") {
#if HT_AMD #if HT_AMD
TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg1") { TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg1") {
if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) { if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) {
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!"); WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
REQUIRE(false); REQUIRE(false);
} }
int stat = 0; int stat = 0, Pageable = 0;
if (fork() == 0) { // child process
int *Ptr = nullptr, SIZE = sizeof(int); HIP_CHECK(hipDeviceGetAttribute(&Pageable,
bool HmmMem = false; hipDeviceAttributePageableMemoryAccess, 0));
YES_COHERENT = false; INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
// Allocating hipHostMalloc() memory
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocPortable)); if (Pageable) {
*Ptr = 1; if (fork() == 0) { // child process
TstCoherency(Ptr, HmmMem); int *Ptr = nullptr, SIZE = sizeof(int);
if (YES_COHERENT) { bool HmmMem = false;
// exit() with code 10 which indicates pass YES_COHERENT = false;
HIP_CHECK(hipHostFree(Ptr)); // Allocating hipHostMalloc() memory
exit(10); HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocPortable));
} else { *Ptr = 1;
// exit() with code 9 which indicates fail TstCoherency(Ptr, HmmMem);
HIP_CHECK(hipHostFree(Ptr)); if (YES_COHERENT) {
exit(9); // exit() with code 10 which indicates pass
} HIP_CHECK(hipHostFree(Ptr));
} else { // parent process exit(10);
wait(&stat); } else {
int Result = WEXITSTATUS(stat); // exit() with code 9 which indicates fail
if (Result != 10) { HIP_CHECK(hipHostFree(Ptr));
REQUIRE(false); 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 #endif
@@ -737,33 +733,43 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg1") {
#if HT_AMD #if HT_AMD
TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg2") { TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg2") {
if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) { if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) {
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!"); WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
REQUIRE(false); REQUIRE(false);
} }
int stat = 0; int stat = 0, Pageable = 0;
if (fork() == 0) { // child process
int *Ptr = nullptr, SIZE = sizeof(int); HIP_CHECK(hipDeviceGetAttribute(&Pageable,
bool HmmMem = false; hipDeviceAttributePageableMemoryAccess, 0));
YES_COHERENT = false; INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
// Allocating hipHostMalloc() memory
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocWriteCombined)); if (Pageable) {
*Ptr = 4; if (fork() == 0) { // child process
TstCoherency(Ptr, HmmMem); int *Ptr = nullptr, SIZE = sizeof(int);
if (YES_COHERENT) { bool HmmMem = false;
// exit() with code 10 which indicates pass YES_COHERENT = false;
HIP_CHECK(hipHostFree(Ptr)); // Allocating hipHostMalloc() memory
exit(10); HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocWriteCombined));
} else { *Ptr = 4;
// exit() with code 9 which indicates fail TstCoherency(Ptr, HmmMem);
HIP_CHECK(hipHostFree(Ptr)); if (YES_COHERENT) {
exit(9); // exit() with code 10 which indicates pass
} HIP_CHECK(hipHostFree(Ptr));
} else { // parent process exit(10);
wait(&stat); } else {
int Result = WEXITSTATUS(stat); // exit() with code 9 which indicates fail
if (Result != 10) { HIP_CHECK(hipHostFree(Ptr));
REQUIRE(false); 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 #endif
@@ -775,33 +781,43 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg2") {
#if HT_AMD #if HT_AMD
TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg3") { TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg3") {
if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) { if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) {
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!"); WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
REQUIRE(false); REQUIRE(false);
} }
int stat = 0; int stat = 0, Pageable = 0;
if (fork() == 0) { // child process
int *Ptr = nullptr, SIZE = sizeof(int); HIP_CHECK(hipDeviceGetAttribute(&Pageable,
bool HmmMem = false; hipDeviceAttributePageableMemoryAccess, 0));
YES_COHERENT = false; INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
// Allocating hipHostMalloc() memory
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNumaUser)); if (Pageable) {
*Ptr = 1; if (fork() == 0) { // child process
TstCoherency(Ptr, HmmMem); int *Ptr = nullptr, SIZE = sizeof(int);
if (YES_COHERENT) { bool HmmMem = false;
// exit() with code 10 which indicates pass YES_COHERENT = false;
HIP_CHECK(hipHostFree(Ptr)); // Allocating hipHostMalloc() memory
exit(10); HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNumaUser));
} else { *Ptr = 1;
// exit() with code 9 which indicates fail TstCoherency(Ptr, HmmMem);
HIP_CHECK(hipHostFree(Ptr)); if (YES_COHERENT) {
exit(9); // exit() with code 10 which indicates pass
} HIP_CHECK(hipHostFree(Ptr));
} else { // parent process exit(10);
wait(&stat); } else {
int Result = WEXITSTATUS(stat); // exit() with code 9 which indicates fail
if (Result != 10) { HIP_CHECK(hipHostFree(Ptr));
REQUIRE(false); 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 #endif
tests/catch/unit/memory/hipMemCoherencyTst.cc
+77 -69
Voir le fichier
@@ -91,54 +91,39 @@ static void TstCoherency(int *Ptr, bool HmmMem) {
} }
} }
static int HmmAttrPrint() {
int managed = 0;
INFO("The following are the attribute values related to HMM for"
" device 0:\n");
HIP_CHECK(hipDeviceGetAttribute(&managed,
hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
INFO("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
HIP_CHECK(hipDeviceGetAttribute(&managed,
hipDeviceAttributeConcurrentManagedAccess, 0));
INFO("hipDeviceAttributeConcurrentManagedAccess: " << managed);
HIP_CHECK(hipDeviceGetAttribute(&managed,
hipDeviceAttributePageableMemoryAccess, 0));
INFO("hipDeviceAttributePageableMemoryAccess: " << managed);
HIP_CHECK(hipDeviceGetAttribute(&managed,
hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
INFO("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:"
<< managed);
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
0));
INFO("hipDeviceAttributeManagedMemory: " << managed);
return managed;
}
/* Test case description: The following test validates if fine grain /* Test case description: The following test validates if fine grain
behavior is observed or not with memory allocated using hipHostMalloc()*/ behavior is observed or not with memory allocated using hipHostMalloc()*/
// The following tests are disabled for Nvidia as they are not consistently // The following tests are disabled for Nvidia as they are not consistently
// passing // passing
#if HT_AMD #if HT_AMD
TEST_CASE("Unit_hipHostMalloc_CoherentTst") { TEST_CASE("Unit_hipHostMalloc_CoherentTst") {
int *Ptr = nullptr, SIZE = sizeof(int); int *Ptr = nullptr, SIZE = sizeof(int), Pageable = 0;
bool HmmMem = false; bool HmmMem = false;
YES_COHERENT = false; YES_COHERENT = false;
// 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(hipDeviceGetAttribute(&Pageable,
HIP_CHECK(hipHostFree(Ptr)); hipDeviceAttributePageableMemoryAccess, 0));
REQUIRE(YES_COHERENT); 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");
}
} }
#endif #endif
@@ -149,12 +134,19 @@ TEST_CASE("Unit_hipHostMalloc_CoherentTst") {
// passing // passing
#if HT_AMD #if HT_AMD
TEST_CASE("Unit_hipMallocManaged_CoherentTst") { TEST_CASE("Unit_hipMallocManaged_CoherentTst") {
int *Ptr = nullptr, SIZE = sizeof(int); int *Ptr = nullptr, SIZE = sizeof(int), Pageable = 0, managed = 0;
bool HmmMem = true; bool HmmMem = true;
YES_COHERENT = false; YES_COHERENT = false;
int managed = HmmAttrPrint(); HIP_CHECK(hipDeviceGetAttribute(&Pageable,
if (managed == 1) { hipDeviceAttributePageableMemoryAccess, 0));
INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
0));
INFO("hipDeviceAttributeManagedMemory: " << managed);
if (managed == 1 && Pageable == 1) {
// Allocating hipMallocManaged() memory // Allocating hipMallocManaged() memory
SECTION("hipMallocManaged with hipMemAttachGlobal flag") { SECTION("hipMallocManaged with hipMemAttachGlobal flag") {
HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachGlobal)); HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachGlobal));
@@ -166,8 +158,8 @@ TEST_CASE("Unit_hipMallocManaged_CoherentTst") {
HIP_CHECK(hipFree(Ptr)); HIP_CHECK(hipFree(Ptr));
REQUIRE(YES_COHERENT); REQUIRE(YES_COHERENT);
} else { } else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " SUCCEED("GPU 0 doesn't support ManagedMemory or PageableMemoryAccess"
"attribute. Hence skipping the testing with Pass result.\n"); "device attribute. Hence skipping the test with Pass result.\n");
} }
} }
#endif #endif
@@ -175,30 +167,40 @@ TEST_CASE("Unit_hipMallocManaged_CoherentTst") {
/* Test case description: The following test validates if memory access is fine /* Test case description: The following test validates if memory access is fine
with memory allocated using hipMallocManaged() and CoarseGrain Advise*/ with memory allocated using hipMallocManaged() and CoarseGrain Advise*/
TEST_CASE("Unit_hipMallocManaged_CoherentTstWthAdvise") { TEST_CASE("Unit_hipMallocManaged_CoherentTstWthAdvise") {
int *Ptr = nullptr, SIZE = sizeof(int); int *Ptr = nullptr, SIZE = sizeof(int), managed = 0;
YES_COHERENT = false; YES_COHERENT = false;
// Allocating hipMallocManaged() memory
SECTION("hipMallocManaged with hipMemAttachGlobal flag") { HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachGlobal)); 0));
} INFO("hipDeviceAttributeManagedMemory: " << managed);
SECTION("hipMallocManaged with hipMemAttachHost flag") {
HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachHost)); if (managed == 1) {
} // Allocating hipMallocManaged() memory
SECTION("hipMallocManaged with hipMemAttachGlobal flag") {
HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachGlobal));
}
SECTION("hipMallocManaged with hipMemAttachHost flag") {
HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachHost));
}
#if HT_AMD #if HT_AMD
HIP_CHECK(hipMemAdvise(Ptr, SIZE, hipMemAdviseSetCoarseGrain, 0)); HIP_CHECK(hipMemAdvise(Ptr, SIZE, hipMemAdviseSetCoarseGrain, 0));
#endif #endif
// Initializing Ptr memory with 9 // Initializing Ptr memory with 9
*Ptr = 9; *Ptr = 9;
hipStream_t strm; hipStream_t strm;
HIP_CHECK(hipStreamCreate(&strm)); HIP_CHECK(hipStreamCreate(&strm));
SquareKrnl<<<1, 1, 0, strm>>>(Ptr); SquareKrnl<<<1, 1, 0, strm>>>(Ptr);
HIP_CHECK(hipStreamSynchronize(strm)); HIP_CHECK(hipStreamSynchronize(strm));
if (*Ptr == 81) { if (*Ptr == 81) {
YES_COHERENT = true; YES_COHERENT = true;
}
HIP_CHECK(hipFree(Ptr));
HIP_CHECK(hipStreamDestroy(strm));
REQUIRE(YES_COHERENT);
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the test with Pass result.\n");
} }
HIP_CHECK(hipFree(Ptr));
HIP_CHECK(hipStreamDestroy(strm));
REQUIRE(YES_COHERENT);
} }
@@ -226,12 +228,18 @@ TEST_CASE("Unit_hipMalloc_CoherentTst") {
hipExtMallocWithFlags()*/ hipExtMallocWithFlags()*/
#if HT_AMD #if HT_AMD
TEST_CASE("Unit_hipExtMallocWithFlags_CoherentTst") { TEST_CASE("Unit_hipExtMallocWithFlags_CoherentTst") {
int *Ptr = nullptr, SIZE = sizeof(int), InitVal = 9; int *Ptr = nullptr, SIZE = sizeof(int), InitVal = 9, Pageable = 0, managed = 0;
bool FineGrain = true; bool FineGrain = true;
YES_COHERENT = false; YES_COHERENT = false;
int managed = HmmAttrPrint(); HIP_CHECK(hipDeviceGetAttribute(&Pageable,
if (managed == 1) { hipDeviceAttributePageableMemoryAccess, 0));
INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
0));
INFO("hipDeviceAttributeManagedMemory: " << managed);
if (managed == 1 && Pageable == 1) {
// Allocating hipExtMallocWithFlags() memory with flags // Allocating hipExtMallocWithFlags() memory with flags
SECTION("hipExtMallocWithFlags with hipDeviceMallocFinegrained flag") { SECTION("hipExtMallocWithFlags with hipDeviceMallocFinegrained flag") {
HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast<void**>(&Ptr), SIZE*2, HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast<void**>(&Ptr), SIZE*2,
@@ -264,8 +272,8 @@ TEST_CASE("Unit_hipExtMallocWithFlags_CoherentTst") {
HIP_CHECK(hipFree(Ptr)); HIP_CHECK(hipFree(Ptr));
REQUIRE(YES_COHERENT); REQUIRE(YES_COHERENT);
} else { } else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " SUCCEED("GPU 0 doesn't support ManagedMemory or PageableMemoryAccess"
"attribute. Hence skipping the testing with Pass result.\n"); "device attribute. Hence skipping the test with Pass result.\n");
} }
} }
#endif #endif