SWDEV-245532 - HIP - Vulkan interop

Change-Id: Iba1ef8112e318b4f099da5a4a4602e0dae7de9e3


[ROCm/hip commit: a1b321bba4]

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
  *

projects/hip/rocclr/hip_hcc.def.in

@@ -273,3 +273,6 @@ hipMemcpyParam2DAsync
 __gnu_h2f_ieee
 __gnu_f2h_ieee
 hipExtStreamGetCUMask
+hipImportExternalMemory
+hipExternalMemoryGetMappedBuffer
+hipDestroyExternalMemory

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*;

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<amd::BufferVk*>(extMem);
+  const device::Memory* devMem = buf->getDeviceMemory(*hip::getCurrentDevice()->devices()[0]);
+  if (devMem != nullptr) {
+    *devPtr = reinterpret_cast<void*>(devMem->virtualAddress());
+  }
+
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipDestroyExternalMemory(hipExternalMemory_t extMem) {
+  HIP_INIT_API(hipDestroyExternalMemory, extMem);
+
+  if (extMem == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  reinterpret_cast<amd::BufferVk*>(extMem)->release();
+
+  HIP_RETURN(hipSuccess);
+}
+
 // ================================================================================================
 hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
 {