From 478cee0f68504b804d6c4804ab2a8c9d6da5b73e Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi <53337087+satyanveshd@users.noreply.github.com> Date: Thu, 6 Nov 2025 12:07:32 +0530 Subject: [PATCH] SWDEV-559525 - Add the HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE attribute support (#1647) * SWDEV-559525 - Add the HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE attribute implementation * Update indentation in hip_memory.cpp --- projects/clr/hipamd/src/hip_memory.cpp | 24 +++- .../unit/memory/hipPointerGetAttribute.cc | 109 +++++++++++++++++- 2 files changed, 125 insertions(+), 8 deletions(-) diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 3a20376e29..dddd54e448 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -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(data) = 0; + } else if ((memObj->getMemFlags() & kManagedAlloc) == kManagedAlloc) { + // managed allocation + *reinterpret_cast(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(data) = 0; + } + } else { + // device pointer, allocated using cudaMalloc + *reinterpret_cast(data) = 1; + } + } else { + // must be a normal host pointer or virtual memory not backed to a physical memory + *reinterpret_cast(data) = 0; + status = hipErrorInvalidValue; + } break; } case HIP_POINTER_ATTRIBUTE_RANGE_START_ADDR: { diff --git a/projects/hip-tests/catch/unit/memory/hipPointerGetAttribute.cc b/projects/hip-tests/catch/unit/memory/hipPointerGetAttribute.cc index 1e8336a25f..97fc15237e 100644 --- a/projects/hip-tests/catch/unit/memory/hipPointerGetAttribute.cc +++ b/projects/hip-tests/catch/unit/memory/hipPointerGetAttribute.cc @@ -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 #include 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(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(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(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(&A_d), &pitch_A, width, NUM_H)); + HIP_CHECK(hipPointerGetAttribute(&datatype, HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE, + reinterpret_cast(A_d))); + + REQUIRE(datatype == 1); + } +#if HT_AMD + SECTION("Malloc Array Allocation") { + CHECK_IMAGE_SUPPORT + hipArray_t B_d; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + 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(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(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(ptrA))); + + REQUIRE(datatype == 0); + } + +}