SWDEV-245532 - HIP - Vulkan interop semaphores

Change-Id: I89be5ee84d4728d76e1987d5016c944c0dbc9a96

hipamd/include/hip/amd_detail/hip_runtime_api.h

@@ -365,6 +365,64 @@ typedef struct hipExternalMemoryBufferDesc_st {
 
 typedef void* hipExternalMemory_t;
 
+typedef enum hipExternalSemaphoreHandleType_enum {
+  hipExternalSemaphoreHandleTypeOpaqueFd = 1,
+  hipExternalSemaphoreHandleTypeOpaqueWin32 = 2,
+  hipExternalSemaphoreHandleTypeOpaqueWin32Kmt = 3,
+  hipExternalSemaphoreHandleTypeD3D12Fence = 4
+} hipExternalSemaphoreHandleType;
+
+
+typedef struct hipExternalSemaphoreHandleDesc_st {
+  hipExternalSemaphoreHandleType type;
+  union {
+    int fd;
+    struct {
+      void* handle;
+      const void* name;
+    } win32;
+  } handle;
+  unsigned int flags;
+} hipExternalSemaphoreHandleDesc;
+
+typedef void* hipExternalSemaphore_t;
+
+typedef struct hipExternalSemaphoreSignalParams_st {
+  struct {
+    struct {
+      unsigned long long value;
+    } fence;
+
+    struct {
+      unsigned long long key;
+    } keyedMutex;
+    unsigned int reserved[12];
+  } params;
+
+  unsigned int flags;
+  unsigned int reserved[16];
+} hipExternalSemaphoreSignalParams;
+
+/**
+ * External semaphore wait parameters, compatible with driver type
+ */
+typedef struct hipExternalSemaphoreWaitParams_st {
+  struct {
+    struct {
+      unsigned long long value;
+    } fence;
+
+    struct {
+      unsigned long long key;
+      unsigned int timeoutMs;
+    } keyedMutex;
+    unsigned int reserved[10];
+  } params;
+
+  unsigned int flags;
+  unsigned int reserved[16];
+} hipExternalSemaphoreWaitParams;
+
 #if __HIP_HAS_GET_PCH
 /**
  * Internal use only. This API may change in the future
@@ -1587,6 +1645,74 @@ hipError_t hipEventQuery(hipEvent_t event);
  */
 hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void* ptr);
 
+
+
+/**
+ *  @brief Imports an external semaphore.
+ *
+ *  @param[out] extSem_out  External semaphores to be waited on
+ *  @param[in] semHandleDesc Semaphore import handle descriptor
+ *
+ *  @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue
+ *
+ *  @see
+ */
+hipError_t hipImportExternalSemaphore(hipExternalSemaphore_t* extSem_out,
+                                      const hipExternalSemaphoreHandleDesc* semHandleDesc);
+
+
+
+/**
+ *  @brief Signals a set of external semaphore objects.
+ *
+ *  @param[in] extSem_out  External semaphores to be waited on
+ *  @param[in] paramsArray Array of semaphore parameters
+ *  @param[in] numExtSems Number of semaphores to wait on
+ *  @param[in] stream Stream to enqueue the wait operations in
+ *
+ *  @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue
+ *
+ *  @see
+ */
+
+hipError_t hipSignalExternalSemaphoresAsync(const hipExternalSemaphore_t* extSemArray,
+                                            const hipExternalSemaphoreSignalParams* paramsArray,
+                                            unsigned int numExtSems, hipStream_t stream);
+
+
+
+/**
+ *  @brief Waits on a set of external semaphore objects
+ *
+ *  @param[in] extSem_out  External semaphores to be waited on
+ *  @param[in] paramsArray Array of semaphore parameters
+ *  @param[in] numExtSems Number of semaphores to wait on
+ *  @param[in] stream Stream to enqueue the wait operations in
+ *
+ *  @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue
+ *
+ *  @see
+ */
+hipError_t hipWaitExternalSemaphoresAsync(const hipExternalSemaphore_t* extSemArray,
+                                              const hipExternalSemaphoreWaitParams* paramsArray,
+                                              unsigned int numExtSems, hipStream_t stream);
+
+
+
+/**
+ *  @brief Destroys an external semaphore object and releases any references to the underlying resource. Any outstanding signals or waits must have completed before the semaphore is destroyed.
+
+ *
+ *  @param[in] extSem handle to an external memory object
+
+ *
+ *  @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue
+ *
+ *  @see
+ */
+hipError_t hipDestroyExternalSemaphore(hipExternalSemaphore_t extSem);
+
+ 
 /**
 *  @brief Imports an external memory object.
 *

hipamd/rocclr/hip_hcc.def.in

@@ -287,3 +287,7 @@ hipStreamIsCapturing
 hipStreamBeginCapture
 hipStreamEndCapture
 hipGraphExecDestroy
+hipImportExternalSemaphore
+hipSignalExternalSemaphoresAsync
+hipWaitExternalSemaphoresAsync
+hipDestroyExternalSemaphore

hipamd/rocclr/hip_hcc.map.in

@@ -281,6 +281,10 @@ global:
     hipStreamBeginCapture;
     hipStreamEndCapture;
     hipGraphExecDestroy;
+    hipImportExternalSemaphore;
+    hipSignalExternalSemaphoresAsync;
+    hipWaitExternalSemaphoresAsync;
+    hipDestroyExternalSemaphore;
     extern "C++" {
     hip_impl::hipLaunchKernelGGLImpl*;
     hip_impl::demangle*;

hipamd/rocclr/hip_memory.cpp

@@ -158,6 +158,98 @@ hipError_t hipDestroyExternalMemory(hipExternalMemory_t extMem) {
   HIP_RETURN(hipSuccess);
 }
 
+
+hipError_t hipImportExternalSemaphore(hipExternalSemaphore_t* extSem_out,
+                                      const hipExternalSemaphoreHandleDesc* semHandleDesc)
+{
+  HIP_INIT_API(hipImportExternalSemaphore, extSem_out, semHandleDesc);
+  if (extSem_out == nullptr || semHandleDesc == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  amd::Device* device = hip::getCurrentDevice()->devices()[0];
+
+#ifdef _WIN32
+  if (device->importExtSemaphore(extSem_out, semHandleDesc->handle.win32.handle)) {
+#else
+  if (device->importExtSemaphore(
+          extSem_out, semHandleDesc->handle.fd)) {
+#endif
+    HIP_RETURN(hipSuccess);
+  }
+  HIP_RETURN(hipErrorInvalidValue);
+}
+
+
+hipError_t hipSignalExternalSemaphoresAsync(
+    const hipExternalSemaphore_t* extSemArray, const hipExternalSemaphoreSignalParams* paramsArray,
+    unsigned int numExtSems, hipStream_t stream )
+{
+  HIP_INIT_API(hipSignalExternalSemaphoresAsync, extSemArray, paramsArray, numExtSems, stream);
+  if (extSemArray == nullptr || paramsArray == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  amd::HostQueue* queue = hip::getQueue(stream);
+  const amd::Device& device = queue->vdev()->device();
+
+  for (unsigned int i = 0; i < numExtSems; i++) {
+    if (extSemArray[i] != nullptr) {
+      amd::ExternalSemaphoreCmd* command =
+          new amd::ExternalSemaphoreCmd(*queue, extSemArray[i], paramsArray[i].params.fence.value,
+                                        amd::ExternalSemaphoreCmd::COMMAND_SIGNAL_EXTSEMAPHORE);
+      if (command == nullptr) {
+        return hipErrorOutOfMemory;
+      }
+      command->enqueue();
+      command->release();
+    } else {
+      HIP_RETURN(hipErrorInvalidValue);
+    }
+  }
+
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipWaitExternalSemaphoresAsync(const hipExternalSemaphore_t* extSemArray,
+        const hipExternalSemaphoreWaitParams* paramsArray,
+        unsigned int numExtSems, hipStream_t stream)
+{
+  HIP_INIT_API(hipWaitExternalSemaphoresAsync, extSemArray, paramsArray, numExtSems,
+               stream);
+  if (extSemArray == nullptr || paramsArray == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  amd::HostQueue* queue = hip::getQueue(stream);
+  const amd::Device& device = queue->vdev()->device();
+
+  for (unsigned int i = 0; i < numExtSems; i++) {
+    if (extSemArray[i] != nullptr) {
+      amd::ExternalSemaphoreCmd* command =
+          new amd::ExternalSemaphoreCmd(*queue, extSemArray[i], paramsArray[i].params.fence.value,
+                                        amd::ExternalSemaphoreCmd::COMMAND_WAIT_EXTSEMAPHORE);
+      if (command == nullptr) {
+        return hipErrorOutOfMemory;
+      }
+      command->enqueue();
+      command->release();
+    } else {
+      HIP_RETURN(hipErrorInvalidValue);
+    }
+  }
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipDestroyExternalSemaphore(hipExternalSemaphore_t extSem)
+{
+  HIP_INIT_API(hipDestroyExternalSemaphore, extSem);
+  if (extSem == nullptr ) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  amd::Device* device = hip::getCurrentDevice()->devices()[0];
+  device->DestroyExtSemaphore(extSem);
+  HIP_RETURN(hipSuccess);
+}
+
+
 // ================================================================================================
 hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
 {