From b515c67a1f6d47101b6a8496ab04fdfef58fd0ed Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Sat, 18 Jul 2026 14:06:06 -0700 Subject: [PATCH 1/2] [UR][HIP] Implement async device allocation Port the CUDA adapter's stream-ordered (async) allocation support to HIP. urEnqueueUSMDeviceAllocExp and urEnqueueUSMFreeExp are implemented using the HIP stream-ordered memory allocator (hipMallocAsync/hipFreeAsync) on the device's default memory pool, mirroring the CUDA implementation but using ScopedDevice instead of ScopedContext. Host and shared async allocations remain unsupported, matching CUDA. Register the four async-alloc entry points in the adapter's urGetEnqueueExpProcAddrTable; they were previously stubbed but never wired into the DDI table, so the loader returned UR_RESULT_ERROR_UNINITIALIZED. UR_DEVICE_INFO_ASYNC_USM_ALLOCATIONS_SUPPORT_EXP now queries hipDeviceAttributeMemoryPoolsSupported so the device advertises the ext_oneapi_async_memory_alloc aspect when supported. The default-pool async allocation e2e tests (async_alloc, ooo_queue_async_alloc) pass on an AMD Instinct MI300A (gfx942). Tests that construct a sycl memory_pool still require the USM memory pool experimental API (urUSMPool*Exp), which remains unimplemented on HIP. Co-authored-by: Cursor --- .../source/adapters/hip/async_alloc.cpp | 91 ++++++++++++++++--- .../source/adapters/hip/device.cpp | 5 + .../adapters/hip/ur_interface_loader.cpp | 4 + 3 files changed, 86 insertions(+), 14 deletions(-) diff --git a/unified-runtime/source/adapters/hip/async_alloc.cpp b/unified-runtime/source/adapters/hip/async_alloc.cpp index 9052477f15b85..d731017f365d8 100644 --- a/unified-runtime/source/adapters/hip/async_alloc.cpp +++ b/unified-runtime/source/adapters/hip/async_alloc.cpp @@ -1,4 +1,4 @@ -//===--------- async_alloc.cpp - CUDA Adapter -----------------------------===// +//===--------- async_alloc.cpp - HIP Adapter ------------------------------===// // // // Part of the LLVM Project, under the Apache License v2.0 with LLVM @@ -9,31 +9,94 @@ #include -UR_APIEXPORT ur_result_t urEnqueueUSMDeviceAllocExp( - ur_queue_handle_t, ur_usm_pool_handle_t, const size_t, - const ur_exp_async_usm_alloc_properties_t *, uint32_t, - const ur_event_handle_t *, void **, ur_event_handle_t *) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +#include "common.hpp" +#include "context.hpp" +#include "enqueue.hpp" +#include "event.hpp" +#include "queue.hpp" +#include "usm.hpp" + +UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMDeviceAllocExp( + ur_queue_handle_t hQueue, [[maybe_unused]] ur_usm_pool_handle_t hPool, + const size_t size, const ur_exp_async_usm_alloc_properties_t *, + uint32_t numEventsInWaitList, const ur_event_handle_t *phEventWaitList, + void **ppMem, ur_event_handle_t *phEvent) try { + std::unique_ptr RetImplEvent{nullptr}; + + ScopedDevice Active(hQueue->getDevice()); + uint32_t StreamToken; + ur_stream_guard Guard; + hipStream_t HIPStream = hQueue->getNextComputeStream( + numEventsInWaitList, phEventWaitList, Guard, &StreamToken); + + UR_CHECK_ERROR(enqueueEventsWait(hQueue, HIPStream, numEventsInWaitList, + phEventWaitList)); + + if (phEvent) { + RetImplEvent = std::make_unique( + UR_COMMAND_ENQUEUE_USM_DEVICE_ALLOC_EXP, hQueue, HIPStream, + StreamToken); + UR_CHECK_ERROR(RetImplEvent->start()); + } + + // Allocate from the device's default stream-ordered memory pool. The HIP + // adapter does not expose a native pool for ur_usm_pool_handle_t, so any + // provided pool falls back to the default pool. + UR_CHECK_ERROR(hipMallocAsync(ppMem, size, HIPStream)); + + if (phEvent) { + UR_CHECK_ERROR(RetImplEvent->record()); + *phEvent = RetImplEvent.release(); + } + + return UR_RESULT_SUCCESS; +} catch (ur_result_t Err) { + return Err; } -UR_APIEXPORT ur_result_t urEnqueueUSMSharedAllocExp( +UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMSharedAllocExp( ur_queue_handle_t, ur_usm_pool_handle_t, const size_t, const ur_exp_async_usm_alloc_properties_t *, uint32_t, const ur_event_handle_t *, void **, ur_event_handle_t *) { return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } -UR_APIEXPORT ur_result_t urEnqueueUSMHostAllocExp( +UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMHostAllocExp( ur_queue_handle_t, ur_usm_pool_handle_t, const size_t, const ur_exp_async_usm_alloc_properties_t *, uint32_t, const ur_event_handle_t *, void **, ur_event_handle_t *) { return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } -UR_APIEXPORT ur_result_t urEnqueueUSMFreeExp(ur_queue_handle_t, - ur_usm_pool_handle_t, void *, - uint32_t, - const ur_event_handle_t *, - ur_event_handle_t *) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMFreeExp( + ur_queue_handle_t hQueue, [[maybe_unused]] ur_usm_pool_handle_t hPool, + void *pMem, uint32_t numEventsInWaitList, + const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) try { + std::unique_ptr RetImplEvent{nullptr}; + + ScopedDevice Active(hQueue->getDevice()); + uint32_t StreamToken; + ur_stream_guard Guard; + hipStream_t HIPStream = hQueue->getNextComputeStream( + numEventsInWaitList, phEventWaitList, Guard, &StreamToken); + + UR_CHECK_ERROR(enqueueEventsWait(hQueue, HIPStream, numEventsInWaitList, + phEventWaitList)); + + if (phEvent) { + RetImplEvent = std::make_unique( + UR_COMMAND_ENQUEUE_USM_FREE_EXP, hQueue, HIPStream, StreamToken); + UR_CHECK_ERROR(RetImplEvent->start()); + } + + UR_CHECK_ERROR(hipFreeAsync(pMem, HIPStream)); + + if (phEvent) { + UR_CHECK_ERROR(RetImplEvent->record()); + *phEvent = RetImplEvent.release(); + } + + return UR_RESULT_SUCCESS; +} catch (ur_result_t Err) { + return Err; } diff --git a/unified-runtime/source/adapters/hip/device.cpp b/unified-runtime/source/adapters/hip/device.cpp index 42ed5cbeaae2c..364916dea9d44 100644 --- a/unified-runtime/source/adapters/hip/device.cpp +++ b/unified-runtime/source/adapters/hip/device.cpp @@ -999,6 +999,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, return ReturnValue(false); case UR_DEVICE_INFO_ASYNC_BARRIER: return ReturnValue(false); + case UR_DEVICE_INFO_ASYNC_USM_ALLOCATIONS_SUPPORT_EXP: { + int MemPoolsSupported = + getAttribute(hDevice, hipDeviceAttributeMemoryPoolsSupported); + return ReturnValue(static_cast(MemPoolsSupported)); + } case UR_DEVICE_INFO_IL_VERSION: return ReturnValue(""); diff --git a/unified-runtime/source/adapters/hip/ur_interface_loader.cpp b/unified-runtime/source/adapters/hip/ur_interface_loader.cpp index 2b15e2f9f8aa2..72a9a0686d093 100644 --- a/unified-runtime/source/adapters/hip/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/hip/ur_interface_loader.cpp @@ -499,6 +499,10 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetEnqueueExpProcAddrTable( pDdiTable->pfnTimestampRecordingExp = urEnqueueTimestampRecordingExp; pDdiTable->pfnNativeCommandExp = urEnqueueNativeCommandExp; + pDdiTable->pfnUSMDeviceAllocExp = urEnqueueUSMDeviceAllocExp; + pDdiTable->pfnUSMSharedAllocExp = urEnqueueUSMSharedAllocExp; + pDdiTable->pfnUSMHostAllocExp = urEnqueueUSMHostAllocExp; + pDdiTable->pfnUSMFreeExp = urEnqueueUSMFreeExp; pDdiTable->pfnCommandBufferExp = urEnqueueCommandBufferExp; pDdiTable->pfnKernelLaunchWithArgsExp = urEnqueueKernelLaunchWithArgsExp; pDdiTable->pfnGraphExp = urEnqueueGraphExp; From f7b51d0570173f4ca3f213d1e2f48f5bf604da36 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Sat, 18 Jul 2026 14:34:53 -0700 Subject: [PATCH 2/2] [UR][HIP] Implement USM memory pool experimental API Port the CUDA adapter's native USM memory pool support to HIP, backed by the HIP stream-ordered memory allocator (hipMemPool_t). A ur_usm_pool_handle_t may now wrap either a UMF pool (existing behaviour) or a native hipMemPool_t. Implement the experimental pool entry points: urUSMPoolCreateExp, urUSMPoolDestroyExp, urUSMPoolGetDefaultDevicePoolExp, urUSMPoolGetInfoExp, urUSMPoolSetInfoExp and urUSMPoolTrimToExp using hipMemPoolCreate/Destroy, hipDeviceGetDefaultMemPool, hipMemPoolGet/SetAttribute and hipMemPoolTrimTo. urUSMPoolGetDevicePoolExp and urUSMPoolSetDevicePoolExp remain unsupported, matching CUDA. urEnqueueUSMDeviceAllocExp now allocates from the supplied native pool via hipMallocFromPoolAsync when a pool is provided, falling back to the device's default pool otherwise. The USM memory pool e2e tests (memory_pool, async_alloc_from_pool, async_alloc_from_default_pool, ooo_queue_async_alloc_from_pool) pass on an AMD Instinct MI300A (gfx942). Co-authored-by: Cursor --- .../source/adapters/hip/async_alloc.cpp | 21 +- unified-runtime/source/adapters/hip/usm.cpp | 208 ++++++++++++++++-- unified-runtime/source/adapters/hip/usm.hpp | 18 ++ 3 files changed, 215 insertions(+), 32 deletions(-) diff --git a/unified-runtime/source/adapters/hip/async_alloc.cpp b/unified-runtime/source/adapters/hip/async_alloc.cpp index d731017f365d8..68caa384bda65 100644 --- a/unified-runtime/source/adapters/hip/async_alloc.cpp +++ b/unified-runtime/source/adapters/hip/async_alloc.cpp @@ -17,10 +17,10 @@ #include "usm.hpp" UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMDeviceAllocExp( - ur_queue_handle_t hQueue, [[maybe_unused]] ur_usm_pool_handle_t hPool, - const size_t size, const ur_exp_async_usm_alloc_properties_t *, - uint32_t numEventsInWaitList, const ur_event_handle_t *phEventWaitList, - void **ppMem, ur_event_handle_t *phEvent) try { + ur_queue_handle_t hQueue, ur_usm_pool_handle_t hPool, const size_t size, + const ur_exp_async_usm_alloc_properties_t *, uint32_t numEventsInWaitList, + const ur_event_handle_t *phEventWaitList, void **ppMem, + ur_event_handle_t *phEvent) try { std::unique_ptr RetImplEvent{nullptr}; ScopedDevice Active(hQueue->getDevice()); @@ -39,10 +39,15 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMDeviceAllocExp( UR_CHECK_ERROR(RetImplEvent->start()); } - // Allocate from the device's default stream-ordered memory pool. The HIP - // adapter does not expose a native pool for ur_usm_pool_handle_t, so any - // provided pool falls back to the default pool. - UR_CHECK_ERROR(hipMallocAsync(ppMem, size, HIPStream)); + // Allocate from the specified device pool if provided. Otherwise, allocate + // from the current device's default pool. + if (hPool) { + assert(hPool->usesHipPool()); + UR_CHECK_ERROR( + hipMallocFromPoolAsync(ppMem, size, hPool->getHipPool(), HIPStream)); + } else { + UR_CHECK_ERROR(hipMallocAsync(ppMem, size, HIPStream)); + } if (phEvent) { UR_CHECK_ERROR(RetImplEvent->record()); diff --git a/unified-runtime/source/adapters/hip/usm.cpp b/unified-runtime/source/adapters/hip/usm.cpp index 78bbe9434816a..e771c48d65995 100644 --- a/unified-runtime/source/adapters/hip/usm.cpp +++ b/unified-runtime/source/adapters/hip/usm.cpp @@ -413,6 +413,52 @@ bool ur_usm_pool_handle_t_::hasUMFPool(umf_memory_pool_t *umf_pool) { HostMemPool.get() == umf_pool; } +ur_usm_pool_handle_t_::ur_usm_pool_handle_t_(ur_context_handle_t Context, + ur_device_handle_t Device, + ur_usm_pool_desc_t *PoolDesc) + : handle_base(), Context{Context}, Device{Device} { + if (!(PoolDesc->flags & UR_USM_POOL_FLAG_USE_NATIVE_MEMORY_POOL_EXP)) + throw UR_RESULT_ERROR_INVALID_ARGUMENT; + + hipMemPoolProps MemPoolProps{}; + size_t threshold = 0; + + const void *pNext = PoolDesc->pNext; + while (pNext != nullptr) { + const ur_base_desc_t *BaseDesc = static_cast(pNext); + switch (BaseDesc->stype) { + case UR_STRUCTURE_TYPE_USM_POOL_LIMITS_DESC: { + const ur_usm_pool_limits_desc_t *Limits = + reinterpret_cast(BaseDesc); + MemPoolProps.maxSize = Limits->maxPoolableSize; + maxSize = Limits->maxPoolableSize; + threshold = Limits->minDriverAllocSize; + break; + } + default: { + throw UR_RESULT_ERROR_INVALID_ARGUMENT; + } + } + pNext = BaseDesc->pNext; + } + + MemPoolProps.allocType = hipMemAllocationTypePinned; + MemPoolProps.location.id = Device->getIndex(); + MemPoolProps.location.type = hipMemLocationTypeDevice; + UR_CHECK_ERROR(hipMemPoolCreate(&HIPMemPool, &MemPoolProps)); + + // Release threshold is not a property when creating a pool. + // It must be set separately. + UR_CHECK_ERROR(urUSMPoolSetInfoExp(this, + UR_USM_POOL_INFO_RELEASE_THRESHOLD_EXP, + &threshold, 8 /*uint64_t*/)); +} + +ur_usm_pool_handle_t_::ur_usm_pool_handle_t_(ur_context_handle_t Context, + ur_device_handle_t Device, + hipMemPool_t HIPMemPool) + : handle_base(), Context{Context}, Device{Device}, HIPMemPool(HIPMemPool) {} + UR_APIEXPORT ur_result_t UR_APICALL urUSMPoolCreate( /// [in] handle of the context object ur_context_handle_t Context, @@ -488,34 +534,135 @@ bool checkUSMAlignment(uint32_t &alignment, const ur_usm_desc_t *pUSMDesc) { (alignment == 0 || ((alignment & (alignment - 1)) == 0))); } -UR_APIEXPORT ur_result_t UR_APICALL urUSMPoolCreateExp(ur_context_handle_t, - ur_device_handle_t, - ur_usm_pool_desc_t *, - ur_usm_pool_handle_t *) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +UR_APIEXPORT ur_result_t UR_APICALL +urUSMPoolCreateExp(ur_context_handle_t Context, ur_device_handle_t Device, + ur_usm_pool_desc_t *pPoolDesc, ur_usm_pool_handle_t *pPool) { + // This entry point only supports native mem pools. + if (!(pPoolDesc->flags & UR_USM_POOL_FLAG_USE_NATIVE_MEMORY_POOL_EXP)) + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + try { + *pPool = reinterpret_cast( + new ur_usm_pool_handle_t_(Context, Device, pPoolDesc)); + } catch (ur_result_t e) { + return e; + } catch (...) { + return UR_RESULT_ERROR_UNKNOWN; + } + return UR_RESULT_SUCCESS; } -UR_APIEXPORT ur_result_t UR_APICALL urUSMPoolDestroyExp(ur_context_handle_t, - ur_device_handle_t, - ur_usm_pool_handle_t) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +UR_APIEXPORT ur_result_t UR_APICALL +urUSMPoolDestroyExp(ur_context_handle_t hContext, ur_device_handle_t hDevice, + ur_usm_pool_handle_t hPool) { + UR_ASSERT(std::find(hContext->getDevices().begin(), + hContext->getDevices().end(), + hDevice) != hContext->getDevices().end(), + UR_RESULT_ERROR_INVALID_CONTEXT); + ScopedDevice Active(hDevice); + + try { + UR_CHECK_ERROR(hipMemPoolDestroy(hPool->getHipPool())); + } catch (ur_result_t Err) { + return Err; + } catch (...) { + return UR_RESULT_ERROR_UNKNOWN; + } + + return UR_RESULT_SUCCESS; } -UR_APIEXPORT ur_result_t UR_APICALL urUSMPoolSetInfoExp(ur_usm_pool_handle_t, - ur_usm_pool_info_t, - void *, size_t) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +UR_APIEXPORT ur_result_t UR_APICALL urUSMPoolSetInfoExp( + ur_usm_pool_handle_t hPool, ur_usm_pool_info_t propName, void *pPropValue, + size_t) { + hipMemPoolAttr attr; + + // All current values are expected to be of size uint64_t. + switch (propName) { + case UR_USM_POOL_INFO_RELEASE_THRESHOLD_EXP: + attr = hipMemPoolAttrReleaseThreshold; + break; + case UR_USM_POOL_INFO_RESERVED_HIGH_EXP: + attr = hipMemPoolAttrReservedMemHigh; + break; + case UR_USM_POOL_INFO_USED_HIGH_EXP: + attr = hipMemPoolAttrUsedMemHigh; + break; + default: + return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION; + } + try { + UR_CHECK_ERROR( + hipMemPoolSetAttribute(hPool->getHipPool(), attr, pPropValue)); + } catch (ur_result_t Err) { + return Err; + } catch (...) { + return UR_RESULT_ERROR_UNKNOWN; + } + return UR_RESULT_SUCCESS; } UR_APIEXPORT ur_result_t UR_APICALL urUSMPoolGetDefaultDevicePoolExp( - ur_context_handle_t, ur_device_handle_t, ur_usm_pool_handle_t *) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + ur_context_handle_t hContext, ur_device_handle_t hDevice, + ur_usm_pool_handle_t *pPool) { + UR_ASSERT(std::find(hContext->getDevices().begin(), + hContext->getDevices().end(), + hDevice) != hContext->getDevices().end(), + UR_RESULT_ERROR_INVALID_CONTEXT); + ScopedDevice Active(hDevice); + + try { + hipMemPool_t HipPool; + UR_CHECK_ERROR( + hipDeviceGetDefaultMemPool(&HipPool, hDevice->getIndex())); + + *pPool = reinterpret_cast( + new ur_usm_pool_handle_t_(hContext, hDevice, HipPool)); + } catch (ur_result_t Err) { + return Err; + } catch (...) { + return UR_RESULT_ERROR_UNKNOWN; + } + + return UR_RESULT_SUCCESS; } -UR_APIEXPORT ur_result_t UR_APICALL urUSMPoolGetInfoExp(ur_usm_pool_handle_t, - ur_usm_pool_info_t, - void *, size_t *) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +UR_APIEXPORT ur_result_t UR_APICALL +urUSMPoolGetInfoExp(ur_usm_pool_handle_t hPool, ur_usm_pool_info_t propName, + void *pPropValue, size_t *pPropSizeRet) { + hipMemPoolAttr attr; + + switch (propName) { + case UR_USM_POOL_INFO_RELEASE_THRESHOLD_EXP: + attr = hipMemPoolAttrReleaseThreshold; + break; + case UR_USM_POOL_INFO_RESERVED_CURRENT_EXP: + attr = hipMemPoolAttrReservedMemCurrent; + break; + case UR_USM_POOL_INFO_RESERVED_HIGH_EXP: + attr = hipMemPoolAttrReservedMemHigh; + break; + case UR_USM_POOL_INFO_USED_CURRENT_EXP: + attr = hipMemPoolAttrUsedMemCurrent; + break; + case UR_USM_POOL_INFO_USED_HIGH_EXP: + attr = hipMemPoolAttrUsedMemHigh; + break; + default: + return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION; + } + + uint64_t value = 0; + UR_CHECK_ERROR( + hipMemPoolGetAttribute(hPool->getHipPool(), attr, (void *)&value)); + + if (pPropValue) { + *(size_t *)pPropValue = value; + } + if (pPropSizeRet) { + *(size_t *)pPropSizeRet = sizeof(size_t); + } + + return UR_RESULT_SUCCESS; } UR_APIEXPORT ur_result_t UR_APICALL urUSMPoolGetDevicePoolExp( @@ -528,11 +675,24 @@ UR_APIEXPORT ur_result_t UR_APICALL urUSMPoolSetDevicePoolExp( return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } -UR_APIEXPORT ur_result_t UR_APICALL urUSMPoolTrimToExp(ur_context_handle_t, - ur_device_handle_t, - ur_usm_pool_handle_t, - size_t) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +UR_APIEXPORT ur_result_t UR_APICALL +urUSMPoolTrimToExp(ur_context_handle_t hContext, ur_device_handle_t hDevice, + ur_usm_pool_handle_t hPool, size_t minBytesToKeep) { + UR_ASSERT(std::find(hContext->getDevices().begin(), + hContext->getDevices().end(), + hDevice) != hContext->getDevices().end(), + UR_RESULT_ERROR_INVALID_CONTEXT); + ScopedDevice Active(hDevice); + + try { + UR_CHECK_ERROR(hipMemPoolTrimTo(hPool->getHipPool(), minBytesToKeep)); + } catch (ur_result_t Err) { + return Err; + } catch (...) { + return UR_RESULT_ERROR_UNKNOWN; + } + + return UR_RESULT_SUCCESS; } UR_APIEXPORT ur_result_t UR_APICALL urUSMContextMemcpyExp(ur_context_handle_t, diff --git a/unified-runtime/source/adapters/hip/usm.hpp b/unified-runtime/source/adapters/hip/usm.hpp index 8f032fd1d553a..9df8197b30fb4 100644 --- a/unified-runtime/source/adapters/hip/usm.hpp +++ b/unified-runtime/source/adapters/hip/usm.hpp @@ -15,10 +15,13 @@ usm::DisjointPoolAllConfigs InitializeDisjointPoolConfig(); +// A ur_usm_pool_handle_t can represent different types of memory pools. It may +// sit on top of a UMF pool or a hipMemPool_t, but not both. struct ur_usm_pool_handle_t_ : ur::hip::handle_base { ur::RefCount RefCount; ur_context_handle_t Context = nullptr; + ur_device_handle_t Device = nullptr; usm::DisjointPoolAllConfigs DisjointPoolConfigs = usm::DisjointPoolAllConfigs(); @@ -27,10 +30,25 @@ struct ur_usm_pool_handle_t_ : ur::hip::handle_base { umf::pool_unique_handle_t SharedMemPool; umf::pool_unique_handle_t HostMemPool; + hipMemPool_t HIPMemPool{0}; + size_t maxSize = 0; + ur_usm_pool_handle_t_(ur_context_handle_t Context, ur_usm_pool_desc_t *PoolDesc); + // Explicit device pool backed by a native hipMemPool_t. + ur_usm_pool_handle_t_(ur_context_handle_t Context, ur_device_handle_t Device, + ur_usm_pool_desc_t *PoolDesc); + + // Explicit device default pool. + ur_usm_pool_handle_t_(ur_context_handle_t Context, ur_device_handle_t Device, + hipMemPool_t HIPMemPool); + bool hasUMFPool(umf_memory_pool_t *umf_pool); + + // To be used if ur_usm_pool_handle_t represents a hipMemPool_t. + bool usesHipPool() const { return HIPMemPool != hipMemPool_t{0}; }; + hipMemPool_t getHipPool() { return HIPMemPool; }; }; // Implements memory allocation via driver API for USM allocator interface