@@ -24,6 +24,10 @@ Supported CUDA APIs
2424* **CudaExternalMemoryGetMappedBuffer **
2525* **cudaExternalMemoryGetMappedMipmappedArray **
2626* **cudaImportExternalMemory **
27+ * **cudaDestroyExternalSemaphore **
28+ * **cudaImportExternalSemaphore **
29+ * **cudaSignalExternalSemaphoresAsync **
30+ * **cudaWaitExternalSemaphoresAsync **
2731
2832Examples
2933********
@@ -98,55 +102,105 @@ CUDA
98102
99103.. code-block :: none
100104
101- // Create a resource using Vulkan APIs
105+ VkDevice device = {};
106+
107+ // Create a texture resource using Vulkan APIs
102108 VkImage vkTexture;
109+ VkImageCreateInfo imageCreateInfo = {}; // init and fill texture descriptor
110+ vkCreateImage(device, &imageCreateInfo, nullptr, &vkTexture);
111+
112+ // Create a semaphore resource using Vulkan APIs
113+ VkSemaphore vkSemaphore;
114+ VkSemaphoreCreateInfo semaphoreInfo = {}; // init and fill semaphore descriptor
115+ vkCreateSemaphore(device, &semaphoreInfo, nullptr, &vkSemaphore);
103116
104- cudaExternalMemoryHandleDesc memHandleDesc; cudaExternalMemoryMipmappedArrayDesc mipmapDesc;
117+ cudaStream_t stream;
118+ cudaStreamCreate(&stream);
119+
120+ cudaExternalMemoryHandleDesc memHandleDesc;
121+ cudaExternalMemoryMipmappedArrayDesc mipmapDesc;
105122 cudaExternalMemory_t externalMemory;
123+ cudaExternalSemaphoreHandleDesc semHandleDesc;
124+ cudaExternalSemaphore_t externalSemaphore;
125+ cudaExternalSemaphoreWaitParams waitParams;
126+ cudaExternalSemaphoreSignalParams signalParams;
127+
106128
107- // Import the memory from external resource (Vulkan )
129+ // Import the memory from external resource (vkTexture )
108130 cudaImportExternalMemory(&externalMemory, &memHandleDesc);
131+ // Import the semaphore from external resource (vkSemaphore)
132+ cudaImportExternalSemaphore(&externalSemaphore, &semHandleDesc);
133+
134+ // Wait on the semaphore using external resource
135+ cudaWaitExternalSemaphoresAsync(&externalSemaphore, &waitParams, 1, stream);
109136
110137 // Get the mapped array from the CUDA resource
111138 cudaMipmappedArray_t cudaMipmappedArray = nullptr;
112- cudaExternalMemoryGetMappedMipmappedArray(&cudaMipmappedArray,
113- externalMemory,
114- &mipmapDesc);
139+ cudaExternalMemoryGetMappedMipmappedArray(&cudaMipmappedArray, externalMemory, &mipmapDesc);
140+
141+ // Signal the semaphore using external resource
142+ cudaSignalExternalSemaphoresAsync(&externalSemaphore, &signalParams, 1, stream);
115143
116144 // Retrieve the tex data as a cudaArray from cudaMipmappedArray
117145 cudaArray_t cudaArr;
118146 cudaGetMipmappedArrayLevel(&cudaArr, cudaMipmappedArray, 0);
119147
120- // Destroy the CUDA resource
148+ // Destroy the CUDA resources
121149 cudaDestroyExternalMemory(externalMemory);
150+ cudaDestroyExternalSemaphore(externalSemaphore);
122151
123152 Migrated Code
124153
125154.. code-block :: none
126155
127- // Create a resource using Vulkan APIs
156+ VkDevice device = {};
157+
158+ // Create a texture resource using Vulkan APIs
128159 VkImage vkTexture;
160+ VkImageCreateInfo imageCreateInfo = {}; // init and fill texture descriptor
161+ vkCreateImage(device, &imageCreateInfo, nullptr, &vkTexture);
162+
163+ // Create a semaphore resource using Vulkan APIs
164+ VkSemaphore vkSemaphore;
165+ VkSemaphoreCreateInfo semaphoreInfo = {}; // init and fill semaphore descriptor
166+ vkCreateSemaphore(device, &semaphoreInfo, nullptr, &vkSemaphore);
167+
168+ dpct::queue_ptr stream;
169+ stream = dpct::get_current_device().create_queue();
129170
130171 dpct::experimental::external_mem_handle_desc memHandleDesc;
131172 dpct::experimental::external_mem_img_desc mipmapDesc;
132173 sycl::ext::oneapi::experimental::external_mem externalMemory;
174+ dpct::experimental::external_sem_handle_desc semHandleDesc;
175+ dpct::experimental::external_sem_wrapper_ptr externalSemaphore;
176+ dpct::experimental::external_sem_params waitParams;
177+ dpct::experimental::external_sem_params signalParams;
178+
133179
134- // Import the memory from external resource (Vulkan )
180+ // Import the memory from external resource (vkTexture )
135181 dpct::experimental::import_external_memory(&externalMemory, &memHandleDesc));
182+ // Import the semaphore from external resource (vkSemaphore)
183+ dpct::experimental::import_external_semaphore(&externalSemaphore, &semHandleDesc);
184+
185+ // Wait on the semaphore using external resource
186+ dpct::experimental::wait_external_semaphore(&externalSemaphore, &waitParams, 1, stream);
136187
137188 // Get the mapped array from the CUDA resource
138189 dpct::experimental::image_mem_wrapper_ptr cudaMipmappedArray = nullptr;
139- cudaMipmappedArray = new dpct::experimental::image_mem_wrapper(
140- externalMemory,
141- &mipmapDesc);
190+ cudaMipmappedArray = new dpct::experimental::image_mem_wrapper(externalMemory, &mipmapDesc);
191+
192+ // Signal the semaphore using external resource
193+ dpct::experimental::signal_external_semaphore(&externalSemaphore, &signalParams, 1, stream);
142194
143195 // Retrieve the tex data as a cudaArray from cudaMipmappedArray
144196 dpct::experimental::image_mem_wrapper_ptr cudaArr;
145197 cudaArr = cudaMipmappedArray->get_mip_level(0);
146198
147- // Destroy the CUDA resource
199+ // Destroy the CUDA resources
148200 sycl::ext::oneapi::experimental::release_external_memory(
149201 externalMemory, dpct::get_in_order_queue());
202+ delete externalSemaphore;
203+
150204
151205 Additional Migration Examples
152206-----------------------------
0 commit comments