diff --git a/docs/how-to/hip_runtime_api/external_interop.rst b/docs/how-to/hip_runtime_api/external_interop.rst index d72a2aa4f6..d007390814 100644 --- a/docs/how-to/hip_runtime_api/external_interop.rst +++ b/docs/how-to/hip_runtime_api/external_interop.rst @@ -46,6 +46,7 @@ Proper memory management ensures stability and efficient resource utilization. Example =============================================================================== + ROCm examples include a `HIP--Vulkan interoperation example `_ demonstrates how to perform interoperation between HIP and Vulkan. @@ -58,24 +59,56 @@ these initial steps, the kernel executes the sine wave computation, and Vulkan continuously updates the window framebuffer to display the computed data until the window is closed. -The Vulkan external descriptor is defined in the following way: +The following code converts a Vulkan memory handle to its equivalent HIP handle. +The input VkDeviceMemory and the created HIP memory represents the same physical +area of GPU memory, through the handles of each respective API. Writing to the +buffer in one API will allow us to read the results through the other. Note that +access to the buffer should be synchronized between the APIs, for example using +queue syncs or semaphores. + +.. literalinclude:: ../../tools/example_codes/external_interop.hip + :start-after: // [Sphinx vulkan memory to hip start] + :end-before: // [Sphinx vulkan memory to hip end] + :language: cpp + +The Vulkan semaphore is converted to HIP semaphore shown in the following +example. Signaling on the semaphore in one API will allow the other API to wait +on it, which is how we can guarantee synchronized access to resources in a +cross-API manner. .. literalinclude:: ../../tools/example_codes/external_interop.hip - :start-after: // [Sphinx external descriptor start] - :end-before: // [Sphinx external descriptor end] + :start-after: // [Sphinx semaphore convert start] + :end-before: // [Sphinx semaphore convert end] :language: cpp - -The Vulkan semaphore is converted to HIP semaphore shown in the following example: + +When the HIP external memory is exported from Vulkan and imported to HIP, it +is not yet ready for use. To actually use the memory, we need to map it to a +pointer so that we may pass it to the kernel so that it can be read from and +written to. The external memory map to HIP in the following example: .. literalinclude:: ../../tools/example_codes/external_interop.hip - :start-after: // [Sphinx semaphore use in kernel start] - :end-before: // [Sphinx semaphore use in kernel end] + :start-after: // [Sphinx map external memory start] + :end-before: // [Sphinx map external memory end] :language: cpp - -The external memory is used in the following example: + +Wait for buffer is ready and not under modification at Vulkan side: + +.. literalinclude:: ../../tools/example_codes/external_interop.hip + :start-after: // [Sphinx wait semaphore start] + :end-before: // [Sphinx wait semaphore end] + :language: cpp + +The sinewave kernel implementation: + +.. literalinclude:: ../../tools/example_codes/external_interop.hip + :start-after: [Sphinx sinewave kernel start] + :end-before: // [Sphinx sinewave kernel end] + :language: cpp + +Signal to Vulkan that we are done with the buffer and that it can proceed with +rendering: .. literalinclude:: ../../tools/example_codes/external_interop.hip - :start-after: // [Sphinx external memory use in kernel start] - :end-before: // [Sphinx external memory use in kernel end] + :start-after: // [Sphinx signal semaphore start] + :end-before: // [Sphinx signal semaphore end] :language: cpp - diff --git a/docs/tools/example_codes/external_interop.hip b/docs/tools/example_codes/external_interop.hip index b8ee73159c..a037170809 100644 --- a/docs/tools/example_codes/external_interop.hip +++ b/docs/tools/example_codes/external_interop.hip @@ -430,7 +430,7 @@ VkBuffer create_buffer(const graphics_context& ctx, hipExternalMemory_t memory_to_hip(const graphics_context& ctx, const VkDeviceMemory memory, const VkDeviceSize size) { - // [Sphinx external descriptor start] + // [Sphinx vulkan memory to hip start] // Prepare the HIP external semaphore descriptor with the platform-specific // handle type that we wish to import. This value should correspond to the // handleTypes field set in VkExportMemoryAllocateInfoKHR while creating the @@ -468,7 +468,7 @@ hipExternalMemory_t hipExternalMemory_t hip_memory; HIP_CHECK(hipImportExternalMemory(&hip_memory, &desc)); return hip_memory; - // [Sphinx external descriptor end] + // [Sphinx vulkan memory to hip end] } /// \brief Utility function to create a Vulkan semaphore. @@ -511,7 +511,7 @@ VkSemaphore create_semaphore(const graphics_context& ctx, const bool external = /// \see create_semaphore for creating such a semaphore. hipExternalSemaphore_t semaphore_to_hip(const graphics_context& ctx, const VkSemaphore sema) { - // [Sphinx semaphore use in kernel start] + // [Sphinx semaphore import start] // Prepare the HIP external semaphore descriptor with the platform-specific handle type // that we wish to import. This value should correspond to the handleTypes field set in // the VkExportSemaphoreCreateInfoKHR structure that was passed to Vulkan when creating @@ -547,7 +547,7 @@ hipExternalSemaphore_t semaphore_to_hip(const graphics_context& ctx, const VkSem // Import the native semaphore to HIP to create a HIP external semaphore. hipExternalSemaphore_t hip_sema; HIP_CHECK(hipImportExternalSemaphore(&hip_sema, &desc)); - // [Sphinx semaphore use in kernel end] + // [Sphinx semaphore import end] return hip_sema; } @@ -556,7 +556,7 @@ hipExternalSemaphore_t semaphore_to_hip(const graphics_context& ctx, const VkSem /// so that we may pass it to the kernel so that it can be read from and written to. void* map_hip_external_memory(const hipExternalMemory_t mem, const VkDeviceSize size) { - // [Sphinx external memory use in kernel start] + // [Sphinx map external memory start] hipExternalMemoryBufferDesc desc = {}; desc.offset = 0; desc.size = size; @@ -564,10 +564,11 @@ void* map_hip_external_memory(const hipExternalMemory_t mem, const VkDeviceSize void* ptr; HIP_CHECK(hipExternalMemoryGetMappedBuffer(&ptr, mem, &desc)); - // [Sphinx external memory use in kernel end] + // [Sphinx map external memory end] return ptr; } +// [Sphinx sinewave kernel start] /// \brief The main HIP kernel for this example - computes a simple sine wave over a /// 2-dimensional grid of points. /// \param height_map - the grid of points to compute a sine wave for. It is expected to be @@ -586,6 +587,7 @@ __global__ void sinewave_kernel(float* height_map, const float time) height_map[x * grid_width + y] = sinf(u * freq + time) * cosf(v * freq + time); } } +// [Sphinx sinewave kernel end] /// \brief In order to increase efficiency, we pipeline the rendering process. This allows us to render /// the next frame already while another frame is being presented by Vulkan. The \p frame structure @@ -1147,11 +1149,13 @@ struct renderer // on it then. if(this->frame_index != 0) { + // [Sphinx wait semaphore start] hipExternalSemaphoreWaitParams wait_params = {}; HIP_CHECK(hipWaitExternalSemaphoresAsync(&this->hip_buffer_ready, &wait_params, 1, this->hip_stream)); + // [Sphinx wait semaphore end] } #else // If semaphores are not supported or not used, then we need to perform a full queue @@ -1169,6 +1173,7 @@ struct renderer // computation over a 2D-grid. constexpr size_t tile_size = 8; + // [Sphinx kernel call start] // Launch the HIP kernel to advance the simulation. sinewave_kernel<<hip_stream>>>(this->hip_height_buffer, time); HIP_CHECK(hipGetLastError()); + // [Sphinx kernel call end] // Signal to Vulkan that we are done with the buffer and that it can proceed // with rendering. #if USE_EXTERNAL_SEMAPHORES == 1 && USE_SIGNAL_SEMAPHORE == 1 // If semaphores are supported and used, signal the semaphore that indicates // that the simulation has finished. + // [Sphinx signal semaphore start] hipExternalSemaphoreSignalParams signal_params = {}; HIP_CHECK(hipSignalExternalSemaphoresAsync(&this->hip_simulation_finished, &signal_params, 1, this->hip_stream)); + // [Sphinx signal semaphore end] #else // If semaphores are not used or not supported, we need to again perform a full // queue sync from the HIP side this time.