使用信號量和內存共享擴展與 NVIDIA OpenCL 進行 Vulkan 互操作

OpenCL 正在改進與其他 API （如 Vulkan ）的互操作方式。本文向您介紹了最新的 OpenCL 互操作風格，最新的 NVIDIA 驅動程序已經支持這種風格。我們提供了可下載的示例代碼，所以您今天可以嘗試這個新功能。

需要一種新的互操作方式

開發人員通常將 OpenCL for compute 與其他 API （如 OpenGL ）一起使用，以訪問包括圖形渲染在內的功能。 OpenCL 長期以來一直支持通過擴展與 OpenGL 、 OpenGL ES 、 EGL 、 Direct3D 10 和 Direct3D 11 共享隱式緩沖區和圖像對象：

cl_khr_gl_sharing
cl_khr_gl_event
cl_khr_egl_image
cl_khr_egl_event
cl_khr_d3d10_sharing
cl_khr_d3d11_sharing

新一代 GPU API （如 Vulkan ）使用對外部內存的顯式引用以及信號量來協調對共享資源的訪問。到目前為止，還沒有 OpenCL 擴展來支持外部內存和信號量與這類新的 API 共享。

OpenCL 和 Vulkan 之間的互操作在移動和桌面平臺上都有很強的需求。 NVIDIA 與 Khronos OpenCL 工作組密切合作，發布了一套臨時跨供應商的 KHR 擴展。這些擴展使應用程序能夠在 OpenCL 和 Vulkan 等 API 之間高效地共享數據，與使用隱式資源的前一代互操作 API 相比，靈活性顯著提高。

這組新的外部內存和信號量共享擴展提供了一個通用框架，使 OpenCL 能夠使用 Vulkan 開發人員熟悉的方法導入外部 API 導出的外部內存和信號量句柄。然后， OpenCL 使用這些信號量來同步外部運行時，協調共享內存的使用。

Diagram shows how OpenCL imports memory and semaphore handles from Vulkan, and uses semaphores to synchronize memory ownership and access. — *圖 1 。 OpenCL 與 Vulkan 軟件的互操作關系*

然后可以添加特定于 API 的外部互操作擴展，以處理與特定 API 交互的細節。 Vulkan 互操作現在可用，并計劃使用其他 API ，如 DirectX 12 。

OpenCL 新的外部信號量和內存共享功能包括單獨的一組精心構造的擴展。

信號量擴展

這組擴展增加了從操作系統特定的信號量句柄創建 OpenCL 信號量對象的能力。

cl_khr_semaphore – 表示帶有等待和信號的信號量。這是一個新的 OpenCL 對象類。
cl_khr_external_semaphore – 使用導入和導出外部信號量的機制擴展cl_khr_semaphore，類似于 VK_KHR_external_semaphore 。

以下擴展使用特定于句柄類型的行為擴展cl_khr_external_semaphore：

cl_khr_external_semaphore_opaque_fd – 使用帶有引用傳輸的 Linux fd 句柄共享外部信號量，類似于 VK_KHR_external_semaphore_fd 。
cl_khr_external_semaphore_win32 – 與 VK_KHR_external_semaphore_win32 類似，使用 win32 NT 和 KMT 句柄與引用轉移共享外部信號量。

內存擴展

這些擴展增加了從操作系統特定的內存句柄創建 OpenCL 內存對象的能力。它們的設計與 Vulkan 外部存儲器擴展 VK_KHR_external_memory . 類似

cl_khr_external_memory – 從其他 API 導入外部內存。

以下擴展使用特定于句柄類型的行為擴展cl_khr_external_memory：

cl_khr_external_memory_opaque_fd – 使用 Linux fd 句柄共享外部內存，類似于 VK_KHR_external_memory_fd 。
cl_khr_external_memory_win32 – 使用 win32 NT 和 KMT 句柄共享外部內存，類似于 VK_KHR_external_memory_win32 。

使用 OpenCL

典型的互操作用例包括以下步驟。

檢查所需的支持是否可用：

檢查底層 OpenCL 平臺和帶有clGetPlatformInfo和clGetDeviceInfo的設備是否支持所需的擴展cl_khr_external_semaphore和cl_khr_external_memory。
為了能夠使用 Win32 信號量和內存句柄，請檢查cl_khr_external_semaphore_win32_khr和cl_khr_external_memory_win32_khr擴展是否存在。
為了能夠使用 FD 信號量和內存句柄，請檢查cl_khr_external_semaphore_opaque_fd_khr和cl_khr_external_memory_opaque_fd_khr擴展是否存在。這也可以通過查詢支持的句柄類型來檢查。

導入外部信號量需要cl_khr_external_semaphore。如果支持cl_khr_external_semaphore_opaque_fd，則可以使用clCreateSemaphoreWithPropertiesKHR和 OpenCL 中的 FD 句柄導入 Vulkan 導出的外部信號量。

// Get cl_devices of the platform. clGetDeviceIDs(..., &devices, &deviceCount); // Create cl_context with just first device clCreateContext(..., 1, devices, ...); // Obtain fd/win32 or similar handle for external semaphore to be imported from the other API. int fd = getFdForExternalSemaphore();// Create clSema of type cl_semaphore_khr usable on the only available device assuming the semaphore was imported from the same device. cl_semaphore_properties_khr sema_props[] = {(cl_semaphore_properties_khr)CL_SEMAPHORE_TYPE_KHR, (cl_semaphore_properties_khr)CL_SEMAPHORE_TYPE_BINARY_KHR, (cl_semaphore_properties_khr)CL_SEMAPHORE_HANDLE_OPAQUE_FD_KHR, (cl_semaphore_properties_khr)fd, 0}; int errcode_ret = 0; cl_semaphore_khr clSema = clCreateSemaphoreWithPropertiesKHR(context, sema_props, &errcode_ret);

導入圖像需要cl_khr_external_memory和對圖像的支持。在 OpenCL 中，通過clCreateSemaphoreWithPropertiesKHR使用 Win32 句柄導入 Vulkan 導出的外部信號量。

// Get cl_devices of the platform. clGetDeviceIDs(..., &devices, &deviceCount); // Create cl_context with just first device clCreateContext(..., 1, devices, ...); // Obtain fd/win32 or similar handle for external semaphore to be imported from the other API. void *handle = getWin32HandleForExternalSemaphore();  // Create clSema of type cl_semaphore_khr usable on the only available device assuming the semaphore was imported from the same device. cl_semaphore_properties_khr sema_props[] = {(cl_semaphore_properties_khr)CL_SEMAPHORE_TYPE_KHR, (cl_semaphore_properties_khr)CL_SEMAPHORE_TYPE_BINARY_KHR, (cl_semaphore_properties_khr)CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KHR, (cl_semaphore_properties_khr)handle, 0}; int errcode_ret = 0; cl_semaphore_khr clSema = clCreateSemaphoreWithPropertiesKHR(context, sema_props, &errcode_ret);

在 OpenCL 中，使用 FD 句柄將 Vulkan 導出的外部內存作為緩沖內存與clCreateBufferWithProperties一起導入。

// Get cl_devices of the platform. 
clGetDeviceIDs(..., &devices, &deviceCount); 
 
// Create cl_context with just first device 
clCreateContext(..., 1, devices, ...); 
 
// Obtain fd/win32 or similar handle for external memory to be imported from other API. 
int fd = getFdForExternalMemory(); 
 
// Create extMemBuffer of type cl_mem from fd. 
cl_mem_properties_khr extMemProperties[] = 
{ (cl_mem_properties_khr)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR, 
 (cl_mem_properties_khr)fd,
0

}; 
cl_mem extMemBuffer = clCreateBufferWithProperties(/*context*/ clContext, 
 /*properties*/ extMemProperties, 
 /*flags*/ 0, 
 /*size*/ size, 
 /*host_ptr*/ NULL, 
 /*errcode_ret*/ &errcode_ret);

在 OpenCL 中，使用clCreateImageWithProperties將 Vulkan 導出的外部內存作為圖像內存導入。

// Create img of type cl_mem. Obtain fd/win32 or similar handle for external memory to be imported from other API. int fd = getFdForExternalMemory(); // Set cl_image_format based on external image info cl_image_format clImgFormat = { }; clImageFormat.image_channel_order = CL_RGBA; clImageFormat.image_channel_data_type = CL_UNORM_INT8; // Set cl_image_desc based on external image info size_t clImageFormatSize; cl_image_desc image_desc = { }; image_desc.image_type = CL_MEM_OBJECT_IMAGE2D_ARRAY; image_desc.image_width = width; image_desc.image_height = height; image_desc.image_depth = depth; cl_mem_properties_khr extMemProperties[] = { (cl_mem_properties_khr)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR, (cl_mem_properties_khr)fd, 0 }; cl_mem img = clCreateImageWithProperties(/*context*/ clContext, /*properties*/ extMemProperties, /*flags*/ 0, /*image_format*/ &clImgFormat, /*image_desc*/ &image_desc, /*errcode_ret*/ &errcode_ret)

使用信號量 wait 和 signal 在 OpenCL 和 Vulkan 之間同步。

// Create clSema using one of the above examples of external semaphore creation. int errcode_ret = 0; cl_semaphore_khr clSema = clCreateSemaphoreWithPropertiesKHR(context, sema_props, &errcode_ret); while (true) { // (not shown) Signal the semaphore from the other API, // Wait for the semaphore in OpenCL clEnqueueWaitSemaphoresKHR( /*command_queue*/ command_queue, /*num_sema_objects*/ 1, /*sema_objects*/ &clSema, /*num_events_in_wait_list*/ 0, /*event_wait_list*/ NULL, /*event*/ NULL); clEnqueueNDRangeKernel(command_queue, ...); clEnqueueSignalSemaphoresKHR(/*command_queue*/ command_queue, /*num_sema_objects*/ 1, /*sema_objects*/ &clSema, /*num_events_in_wait_list*/ 0, /*event_wait_list*/ NULL, /*event*/ NULL); // (not shown) Launch work in the other API that waits on 'clSema'