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

SWDEV-559525 - Add the HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE attribute support (#1647)

Browse Source 
* SWDEV-559525 - Add the HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE attribute implementation

* Update indentation in hip_memory.cpp
This commit is contained in:
Satyanvesh Dittakavi
2025-11-06 12:07:32 +05:30
committed by GitHub
parent 27f85500f8
commit 478cee0f68
2 changed files with 125 additions and 8 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/src/hip_memory.cpp
+22 -2
View File
@@ -3534,6 +3534,7 @@ hipError_t ihipPointerGetAttributes(void* data, hipPointer_attribute attribute,
hipDeviceptr_t ptr) {
size_t offset = 0;
amd::Memory* memObj = getMemoryObject(ptr, offset);
amd::Memory* vaddr_mem_obj = amd::MemObjMap::FindVirtualMemObj(ptr);
constexpr uint32_t kManagedAlloc = (CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_ALLOC_HOST_PTR);
hipError_t status = hipSuccess;
@@ -3659,8 +3660,27 @@ hipError_t ihipPointerGetAttributes(void* data, hipPointer_attribute attribute,
break;
}
case HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE: {
// TODO: Unclear what to be done for this attribute
status = hipErrorNotSupported;
if (memObj) {
if (getMemoryType(memObj) == hipMemoryTypeHost) {
// host pointer, pinned or registered memory
*reinterpret_cast<int*>(data) = 0;
} else if ((memObj->getMemFlags() & kManagedAlloc) == kManagedAlloc) {
// managed allocation
*reinterpret_cast<int*>(data) = 0;
} else if (vaddr_mem_obj) {
// virtual memory allocation, mapped to a physical memory
if (vaddr_mem_obj->getMemFlags() & CL_MEM_VA_RANGE_AMD) {
*reinterpret_cast<int*>(data) = 0;
}
} else {
// device pointer, allocated using cudaMalloc
*reinterpret_cast<int*>(data) = 1;
}
} else {
// must be a normal host pointer or virtual memory not backed to a physical memory
*reinterpret_cast<int*>(data) = 0;
status = hipErrorInvalidValue;
}
break;
}
case HIP_POINTER_ATTRIBUTE_RANGE_START_ADDR: {
projects/hip-tests/catch/unit/memory/hipPointerGetAttribute.cc
+103 -6
View File
@@ -44,6 +44,16 @@ Functional Scenarios:
behaviour
*/
#define checkVMMSupported(device) { \
int value = 0; \
hipDeviceAttribute_t attr = hipDeviceAttributeVirtualMemoryManagementSupported; \
HIP_CHECK(hipDeviceGetAttribute(&value, attr, device)); \
if (value == 0) { \
printf("Machine does not support VMM. Skipping this test.."); \
return; \
} \
}
#include <hip_test_common.hh>
#include <string>
static constexpr auto NUM_W{16};
@@ -308,12 +318,6 @@ TEST_CASE("Unit_hipPointerGetAttribute_Negative") {
REQUIRE(hipPointerGetAttribute(&data, HIP_POINTER_ATTRIBUTE_P2P_TOKENS,
reinterpret_cast<hipDeviceptr_t>(A_d)) == hipErrorNotSupported);
}
SECTION(
"Pass HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE"
"not supported by HIP") {
REQUIRE(hipPointerGetAttribute(&data, HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE,
reinterpret_cast<hipDeviceptr_t>(A_d)) == hipErrorNotSupported);
}
SECTION(
"Pass HIP_POINTER_ATTRIBUTE_ALLOWED_HANDLE_TYPES"
"not supported by HIP") {
@@ -324,3 +328,96 @@ TEST_CASE("Unit_hipPointerGetAttribute_Negative") {
HIP_CHECK(hipFree(A_d));
free(A_h);
}
/* Allocate memory using different Allocation APIs and check whether
IPC CAPABLE attribute returns correctly */
TEST_CASE("Unit_hipPointerGetAttribute_ipc_capable") {
HIP_CHECK(hipSetDevice(0));
size_t Nbytes = N * sizeof(int);
unsigned int datatype;
SECTION("Malloc Allocation") {
int *A_d;
HIP_CHECK(hipMalloc(&A_d, Nbytes));
HIP_CHECK(hipPointerGetAttribute(&datatype, HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE,
reinterpret_cast<hipDeviceptr_t>(A_d)));
REQUIRE(datatype == 1);
}
size_t pitch_A;
size_t width{NUM_W * sizeof(char)};
SECTION("Malloc Pitch Allocation") {
CHECK_IMAGE_SUPPORT
char* A_d;
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A, width, NUM_H));
HIP_CHECK(hipPointerGetAttribute(&datatype, HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE,
reinterpret_cast<hipDeviceptr_t>(A_d)));
REQUIRE(datatype == 1);
}
#if HT_AMD
SECTION("Malloc Array Allocation") {
CHECK_IMAGE_SUPPORT
hipArray_t B_d;
hipChannelFormatDesc desc = hipCreateChannelDesc<char>();
HIP_CHECK(hipMallocArray(&B_d, &desc, NUM_W, NUM_H, hipArrayDefault));
HIP_CHECK_ERROR(hipPointerGetAttribute(&datatype, HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE,
reinterpret_cast<hipDeviceptr_t>(B_d)),
hipErrorInvalidValue);
HIP_CHECK(hipFreeArray(B_d));
}
SECTION("Malloc 3D Array Allocation") {
CHECK_IMAGE_SUPPORT
int width = 10, height = 10, depth = 10;
hipArray_t arr;
hipChannelFormatDesc channelDesc =
hipCreateChannelDesc(sizeof(float) * 8, 0, 0, 0, hipChannelFormatKindFloat);
HIP_CHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth),
hipArrayDefault));
HIP_CHECK_ERROR(hipPointerGetAttribute(&datatype, HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE,
reinterpret_cast<hipDeviceptr_t>(arr)),
hipErrorInvalidValue);
HIP_CHECK(hipFreeArray(arr));
}
#endif
SECTION("VMM Memory Allocation") {
size_t granularity = 0;
int deviceId = 0;
size_t buffer_size = N * sizeof(int);
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
void* ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
// Set access
hipMemAccessDesc accessDesc = {};
accessDesc.location.type = hipMemLocationTypeDevice;
accessDesc.location.id = device;
accessDesc.flags = hipMemAccessFlagsProtReadWrite;
// Make the address accessible to GPU 0
HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
HIP_CHECK(hipPointerGetAttribute(&datatype, HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE,
reinterpret_cast<hipDeviceptr_t>(ptrA)));
REQUIRE(datatype == 0);
}
}