Skip to content

test(hip-tests): detect XNACK via the pageable memory access attribute #13430

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
litvaOo merged 2 commits into from
Oct 21, 2026
Merged

test(hip-tests): detect XNACK via the pageable memory access attribute #13430

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
885c9dc
test(hip-tests): detect XNACK via the pageable memory access attribute
iassiour Oct 14, 2026
00da8b4
test(hip-tests): drop redundant pageable query and document isXnackOn
iassiour Oct 14, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
13 changes: 9 additions & 4 deletions projects/hip-tests/catch/include/hip_test_common.hh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -524,13 +524,18 @@ inline bool checkConcurrentKernels(int num_devices) {
return true;
}

// The runtime reports pageable memory access iff XNACK is on; stall-only XNACK GPUs never
// report xnack+ in gcnArchName.
inline bool isXnackOn() {
hipDeviceProp_t prop;
#if HT_AMD
int device = 0;
int pageable = 0;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
return gfxName.find("xnack+") != std::string::npos;
HIP_CHECK(hipDeviceGetAttribute(&pageable, hipDeviceAttributePageableMemoryAccess, device));
return pageable != 0;
#else
return false;
#endif
}

inline bool areWarpMatchFunctionsSupported() {
Expand Down
24 changes: 4 additions & 20 deletions projects/hip-tests/catch/multiproc/hipMemCoherencyTstMProc.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -91,11 +91,7 @@ bool static TstCoherency(int* Ptr, bool HmmMem) {
HIP_TEST_CASE(Unit_malloc_CoherentTst) {
CHECK_PCIE_ATOMIC_SUPPORT
CHECK_MANAGED_MEMORY_SUPPORT
hipDeviceProp_t prop;
HIPCHECK(hipGetDeviceProperties(&prop, 0));
char* p = NULL;
p = strstr(prop.gcnArchName, "xnack+");
if (p) {
if (HipTest::isXnackOn()) {
int *Ptr = nullptr, SIZE = sizeof(int);
bool HmmMem = true;

Expand All @@ -116,11 +112,7 @@ HIP_TEST_CASE(Unit_malloc_CoherentTst) {
#if HT_AMD
HIP_TEST_CASE(Unit_malloc_CoherentTstWthAdvise) {
CHECK_MANAGED_MEMORY_SUPPORT
hipDeviceProp_t prop;
HIPCHECK(hipGetDeviceProperties(&prop, 0));
char* p = NULL;
p = strstr(prop.gcnArchName, "xnack+");
if (p) {
if (HipTest::isXnackOn()) {
int *Ptr = nullptr, SIZE = sizeof(int);

// Allocating hipMallocManaged() memory
Expand All @@ -146,11 +138,7 @@ HIP_TEST_CASE(Unit_malloc_CoherentTstWthAdvise) {
HIP_TEST_CASE(Unit_mmap_CoherentTst) {
CHECK_PCIE_ATOMIC_SUPPORT
CHECK_MANAGED_MEMORY_SUPPORT
hipDeviceProp_t prop;
HIPCHECK(hipGetDeviceProperties(&prop, 0));
char *p = NULL;
p = strstr(prop.gcnArchName, "xnack+");
if (p) {
if (HipTest::isXnackOn()) {
bool HmmMem = true;
int *Ptr =
reinterpret_cast<int *>(mmap(NULL, sizeof(int), PROT_READ | PROT_WRITE,
Expand All @@ -170,11 +158,7 @@ HIP_TEST_CASE(Unit_mmap_CoherentTst) {
#if HT_AMD
HIP_TEST_CASE(Unit_mmap_CoherentTstWthAdvise) {
CHECK_MANAGED_MEMORY_SUPPORT
hipDeviceProp_t prop;
HIPCHECK(hipGetDeviceProperties(&prop, 0));
char *p = NULL;
p = strstr(prop.gcnArchName, "xnack+");
if (p) {
if (HipTest::isXnackOn()) {
int SIZE = sizeof(int);
int *Ptr = reinterpret_cast<int *>(mmap(NULL, SIZE, PROT_READ | PROT_WRITE,
MAP_PRIVATE | MAP_ANONYMOUS, 0, 0));
Expand Down
9 changes: 2 additions & 7 deletions projects/hip-tests/catch/stress/memory/hipHmmOvrSubscriptionTst.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,13 +22,8 @@ HIP_TEST_CASE(Stress_HMM_OverSubscriptionTst) {
const bool hmm = HipTest::isManagedMemorySupportedOnDevice(0);

bool shouldRun = []() -> bool {
#if HT_AMD // For AMD this gcn arch needs to have xnack+
int device = 0;
hipDeviceProp_t props{};
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&props, device));
std::string arch(props.gcnArchName);
return arch.find("xnack+") != std::string::npos;
#if HT_AMD // For AMD this needs XNACK on
return HipTest::isXnackOn();
#else // For CUDA this depends on SM and attribute check should be fine
return true;
#endif
Expand Down
3 changes: 1 addition & 2 deletions projects/hip-tests/catch/stress/memory/hipHostRegisterStress.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,9 +42,8 @@ static __global__ void Inc(uint8_t* Ad) {
HIP_TEST_CASE(Stress_hipHostRegister_Oversubscription) {
hipDeviceProp_t prop;
HIP_CHECK(hipGetDeviceProperties(&prop, 0));
std::string arch = prop.gcnArchName;
#if HT_AMD
if (std::string::npos == arch.find("xnack+")) {
if (!HipTest::isXnackOn()) {
HIP_SKIP_TEST(HipTest::SkipReason::kGpuXnackNotEnabled);
}
#endif
Expand Down
24 changes: 4 additions & 20 deletions projects/hip-tests/catch/unit/memory/hipMemAdvise_old.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -242,19 +242,9 @@ HIP_TEST_CASE(Unit_hipMemAdvise_TstAccessedByFlg4) {
successfully without hang or segfault*/
#if __linux__ && HT_AMD
HIP_TEST_CASE(Unit_hipMemAdvise_TstAlignedAllocMem) {
// The following code block checks for xnack+
// so as to skip if the device is not xnack+
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);

if (gfxName.find("xnack+") != std::string::npos) {
int pageMemAccess = 0;
HIP_CHECK(hipDeviceGetAttribute(&pageMemAccess, hipDeviceAttributePageableMemoryAccess, 0));
WARN("hipDeviceAttributePageableMemoryAccess:" << pageMemAccess);
if (HipTest::isManagedMemorySupportedOnDevice(0) && (pageMemAccess == 1)) {
// Skip if XNACK is off
if (HipTest::isXnackOn()) {
if (HipTest::isManagedMemorySupportedOnDevice(0)) {
int *Mllc = nullptr, MemSz = 4096 * 4, NumElms = 4096, InitVal = 123;
// Mllc = reinterpret_cast<(int *)>(aligned_alloc(4096, MemSz));
Mllc = reinterpret_cast<int*>(aligned_alloc(4096, 4096 * 4));
Expand Down Expand Up @@ -290,13 +280,7 @@ HIP_TEST_CASE(Unit_hipMemAdvise_TstAlignedAllocMem_XNACK) {
HIP_SKIP_TEST("cannot set XNACK via environment.");
}

hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);

if (gfxName.find("xnack+") != std::string::npos) {
if (HipTest::isXnackOn()) {
hip::SpawnProc proc("hipMemAdviseTstAlignedAllocMem", true);
REQUIRE(proc.run() == 0);
} else {
Expand Down
11 changes: 1 addition & 10 deletions projects/hip-tests/catch/unit/virtualMemoryManagement/hipGetProcAddressVmmApis.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,16 +31,7 @@ HIP_TEST_CASE(Unit_hipGetProcAddress_VMM) {
HIP_SKIP_TEST(HipTest::SkipReason::kVmmUnsupported);
}

hipDeviceProp_t devProp;
int device;
bool xnackEnabled = false;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&devProp, device));
std::string gfxName(devProp.gcnArchName);

if (gfxName.find("xnack+") != std::string::npos) {
xnackEnabled = true;
}
const bool xnackEnabled = HipTest::isXnackOn();

void* hipMemGetAllocationGranularity_ptr = nullptr;
void* hipMemCreate_ptr = nullptr;
Expand Down
Loading