Update external interop code description

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 <https://github.com/ROCm/rocm-examples/tree/develop/HIP-Basic/vulkan_interop>`_
 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
-

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,18 +556,19 @@ 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;
     desc.flags                       = 0;
 
     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,24 +1173,28 @@ 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<<<dim3(ceiling_div(grid_width, tile_size),
                                ceiling_div(grid_height, tile_size)),
                           dim3(tile_size, tile_size),
                           0,
                           this->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.