From 5242decbc07bcd2b55c0f8d06cc9c698574d3eab Mon Sep 17 00:00:00 2001 From: pghafari Date: Fri, 19 Mar 2021 18:36:13 -0400 Subject: [PATCH] SWDEV-245532 - HIP - Vulkan interop Change-Id: Iba1ef8112e318b4f099da5a4a4602e0dae7de9e3 [ROCm/hip commit: a1b321bba4993ef16385f557250a88ec5d53a88b] --- .../include/hip/amd_detail/hip_runtime_api.h | 72 ++++++++++++++++++- projects/hip/rocclr/hip_hcc.def.in | 3 + projects/hip/rocclr/hip_hcc.map.in | 3 + projects/hip/rocclr/hip_memory.cpp | 53 ++++++++++++++ 4 files changed, 128 insertions(+), 3 deletions(-) diff --git a/projects/hip/include/hip/amd_detail/hip_runtime_api.h b/projects/hip/include/hip/amd_detail/hip_runtime_api.h index 777f9377da..7739c3b1d0 100644 --- a/projects/hip/include/hip/amd_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/amd_detail/hip_runtime_api.h @@ -333,6 +333,38 @@ typedef struct hipLaunchParams_t { hipStream_t stream; ///< Stream identifier } hipLaunchParams; + +typedef enum hipExternalMemoryHandleType_enum { + hipExternalMemoryHandleTypeOpaqueFd = 1, + hipExternalMemoryHandleTypeOpaqueWin32 = 2, + hipExternalMemoryHandleTypeOpaqueWin32Kmt = 3, + hipExternalMemoryHandleTypeD3D12Heap = 4, + hipExternalMemoryHandleTypeD3D12Resource = 5, + hipExternalMemoryHandleTypeD3D11Resource = 6, + hipExternalMemoryHandleTypeD3D11ResourceKmt = 7, +} hipExternalMemoryHandleType; + +typedef struct hipExternalMemoryHandleDesc_st { + hipExternalMemoryHandleType type; + union { + int fd; + struct { + void *handle; + const void *name; + } win32; + } handle; + unsigned long long size; + unsigned int flags; +} hipExternalMemoryHandleDesc; + +typedef struct hipExternalMemoryBufferDesc_st { + unsigned long long offset; + unsigned long long size; + unsigned int flags; +} hipExternalMemoryBufferDesc; + +typedef void* hipExternalMemory_t; + #if __HIP_HAS_GET_PCH /** * Internal use only. This API may change in the future @@ -1254,11 +1286,8 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback * not execute until the defined wait condition is true. * * hipStreamWaitValueGte: waits until *ptr&mask >= value - * * hipStreamWaitValueEq : waits until *ptr&mask == value - * * hipStreamWaitValueAnd: waits until ((*ptr&mask) & value) != 0 - * * hipStreamWaitValueNor: waits until ~((*ptr&mask) | (value&mask)) != 0 * * @note when using 'hipStreamWaitValueNor', mask is applied on both 'value' and '*ptr'. @@ -1551,6 +1580,43 @@ hipError_t hipEventQuery(hipEvent_t event); */ hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void* ptr); +/** +* @brief Imports an external memory object. +* +* @param[out] extMem_out Returned handle to an external memory object +* @param[in] memHandleDesc Memory import handle descriptor +* +* @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue +* +* @see +*/ +hipError_t hipImportExternalMemory(hipExternalMemory_t* extMem_out, const hipExternalMemoryHandleDesc* memHandleDesc); + +/** +* @brief Maps a buffer onto an imported memory object. +* +* @param[out] devPtr Returned device pointer to buffer +* @param[in] extMem Handle to external memory object +* @param[in] bufferDesc Buffer descriptor +* +* @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue +* +* @see +*/ +hipError_t hipExternalMemoryGetMappedBuffer(void **devPtr, hipExternalMemory_t extMem, const hipExternalMemoryBufferDesc *bufferDesc); + + +/** +* @brief Destroys an external memory object. +* +* @param[in] extMem External memory object to be destroyed +* +* @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue +* +* @see +*/ +hipError_t hipDestroyExternalMemory(hipExternalMemory_t extMem); + /** * @brief Allocate memory on the default accelerator * diff --git a/projects/hip/rocclr/hip_hcc.def.in b/projects/hip/rocclr/hip_hcc.def.in index 1a4e1a8eb9..eb18971dda 100755 --- a/projects/hip/rocclr/hip_hcc.def.in +++ b/projects/hip/rocclr/hip_hcc.def.in @@ -273,3 +273,6 @@ hipMemcpyParam2DAsync __gnu_h2f_ieee __gnu_f2h_ieee hipExtStreamGetCUMask +hipImportExternalMemory +hipExternalMemoryGetMappedBuffer +hipDestroyExternalMemory \ No newline at end of file diff --git a/projects/hip/rocclr/hip_hcc.map.in b/projects/hip/rocclr/hip_hcc.map.in index b9a2d64301..cf433b18f5 100755 --- a/projects/hip/rocclr/hip_hcc.map.in +++ b/projects/hip/rocclr/hip_hcc.map.in @@ -267,6 +267,9 @@ global: hipTexObjectGetResourceDesc; hipTexObjectGetResourceViewDesc; hipTexObjectGetTextureDesc; + hipImportExternalMemory; + hipExternalMemoryGetMappedBuffer; + hipDestroyExternalMemory; extern "C++" { hip_impl::hipLaunchKernelGGLImpl*; hip_impl::demangle*; diff --git a/projects/hip/rocclr/hip_memory.cpp b/projects/hip/rocclr/hip_memory.cpp index 0faf531dba..de29366213 100755 --- a/projects/hip/rocclr/hip_memory.cpp +++ b/projects/hip/rocclr/hip_memory.cpp @@ -25,6 +25,7 @@ #include "platform/context.hpp" #include "platform/command.hpp" #include "platform/memory.hpp" +#include "amdocl/cl_vk_amd.hpp" // ================================================================================================ amd::Memory* getMemoryObject(const void* ptr, size_t& offset) { @@ -105,6 +106,58 @@ hipError_t ihipFree(void *ptr) return hipErrorInvalidValue; } +hipError_t hipImportExternalMemory(hipExternalMemory_t* extMem_out, const hipExternalMemoryHandleDesc* memHandleDesc) { + HIP_INIT_API(hipImportExternalMemory, extMem_out, memHandleDesc); + + size_t sizeBytes = memHandleDesc->size; + amd::Context& amdContext = *hip::getCurrentDevice()->asContext(); + + amd::BufferVk* pBufferVk = nullptr; +#ifdef _WIN32 + pBufferVk = new (amdContext) amd::BufferVk(amdContext, sizeBytes, memHandleDesc->handle.win32.handle); +#else + pBufferVk = new (amdContext) amd::BufferVk(amdContext, sizeBytes, memHandleDesc->handle.fd); +#endif + + if (!pBufferVk) { + HIP_RETURN(hipErrorOutOfMemory); + } + + if (!pBufferVk->create()) { + pBufferVk->release(); + HIP_RETURN(hipErrorOutOfMemory); + } + *extMem_out = pBufferVk; + + HIP_RETURN(hipSuccess); +} + +hipError_t hipExternalMemoryGetMappedBuffer(void **devPtr, hipExternalMemory_t extMem, const hipExternalMemoryBufferDesc *bufferDesc) { + HIP_INIT_API(hipExternalMemoryGetMappedBuffer, devPtr, extMem, bufferDesc); + + if (extMem == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + amd::BufferVk *buf = reinterpret_cast(extMem); + const device::Memory* devMem = buf->getDeviceMemory(*hip::getCurrentDevice()->devices()[0]); + if (devMem != nullptr) { + *devPtr = reinterpret_cast(devMem->virtualAddress()); + } + + HIP_RETURN(hipSuccess); +} + +hipError_t hipDestroyExternalMemory(hipExternalMemory_t extMem) { + HIP_INIT_API(hipDestroyExternalMemory, extMem); + + if (extMem == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + reinterpret_cast(extMem)->release(); + + HIP_RETURN(hipSuccess); +} + // ================================================================================================ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {