From ccff5216f0fc796071c448d5b1ad00edc502b0a9 Mon Sep 17 00:00:00 2001 From: pghafari Date: Thu, 22 Apr 2021 20:21:01 -0400 Subject: [PATCH] SWDEV-245532 - HIP - Vulkan interop semaphores Change-Id: I89be5ee84d4728d76e1987d5016c944c0dbc9a96 --- .../include/hip/amd_detail/hip_runtime_api.h | 126 ++++++++++++++++++ hipamd/rocclr/hip_hcc.def.in | 4 + hipamd/rocclr/hip_hcc.map.in | 4 + hipamd/rocclr/hip_memory.cpp | 92 +++++++++++++ 4 files changed, 226 insertions(+) diff --git a/hipamd/include/hip/amd_detail/hip_runtime_api.h b/hipamd/include/hip/amd_detail/hip_runtime_api.h index 79bc84a84c..b64f6e9ec1 100644 --- a/hipamd/include/hip/amd_detail/hip_runtime_api.h +++ b/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. * diff --git a/hipamd/rocclr/hip_hcc.def.in b/hipamd/rocclr/hip_hcc.def.in index 1e5007e0ba..559a08289d 100755 --- a/hipamd/rocclr/hip_hcc.def.in +++ b/hipamd/rocclr/hip_hcc.def.in @@ -287,3 +287,7 @@ hipStreamIsCapturing hipStreamBeginCapture hipStreamEndCapture hipGraphExecDestroy +hipImportExternalSemaphore +hipSignalExternalSemaphoresAsync +hipWaitExternalSemaphoresAsync +hipDestroyExternalSemaphore diff --git a/hipamd/rocclr/hip_hcc.map.in b/hipamd/rocclr/hip_hcc.map.in index 41332efcac..e7d761d820 100755 --- a/hipamd/rocclr/hip_hcc.map.in +++ b/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*; diff --git a/hipamd/rocclr/hip_memory.cpp b/hipamd/rocclr/hip_memory.cpp index df5f83e727..20555a0edc 100755 --- a/hipamd/rocclr/hip_memory.cpp +++ b/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) {