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

需要一種新的互操作方式

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

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 ）使用對外部內(nèi)存的顯式引用以及信號量來協(xié)調(diào)對共享資源的訪問。到目前為止，還沒有 OpenCL 擴(kuò)展來支持外部內(nèi)存和信號量與這類新的 API 共享。

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

這組新的外部內(nèi)存和信號量共享擴(kuò)展提供了一個(gè)通用框架，使 OpenCL 能夠使用 Vulkan 開發(fā)人員熟悉的方法導(dǎo)入外部 API 導(dǎo)出的外部內(nèi)存和信號量句柄。然后， OpenCL 使用這些信號量來同步外部運(yùn)行時(shí)，協(xié)調(diào)共享內(nèi)存的使用。

圖 1 。 OpenCL 與 Vulkan 軟件的互操作關(guān)系

然后可以添加特定于 API 的外部互操作擴(kuò)展，以處理與特定 API 交互的細(xì)節(jié)。 Vulkan 互操作現(xiàn)在可用，并計(jì)劃使用其他 API ，如 DirectX 12 。

OpenCL 新的外部信號量和內(nèi)存共享功能包括單獨(dú)的一組精心構(gòu)造的擴(kuò)展。

信號量擴(kuò)展

這組擴(kuò)展增加了從操作系統(tǒng)特定的信號量句柄創(chuàng)建 OpenCL 信號量對象的能力。

cl_khr_semaphore – 表示帶有等待和信號的信號量。這是一個(gè)新的 OpenCL 對象類。

cl_khr_external_semaphore – 使用導(dǎo)入和導(dǎo)出外部信號量的機(jī)制擴(kuò)展cl_khr_semaphore，類似于 VK_KHR_external_semaphore 。

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

cl_khr_external_semaphore_opaque_fd – 使用帶有引用傳輸?shù)?Linux fd 句柄共享外部信號量，類似于 VK_KHR_external_semaphore_fd 。

cl_khr_external_semaphore_win32 – 與 VK_KHR_external_semaphore_win32 類似，使用 win32 NT 和 KMT 句柄與引用轉(zhuǎn)移共享外部信號量。

內(nèi)存擴(kuò)展

這些擴(kuò)展增加了從操作系統(tǒng)特定的內(nèi)存句柄創(chuàng)建 OpenCL 內(nèi)存對象的能力。它們的設(shè)計(jì)與 Vulkan 外部存儲器擴(kuò)展 VK_KHR_external_memory 。 類似

cl_khr_external_memory – 從其他 API 導(dǎo)入外部內(nèi)存。

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

cl_khr_external_memory_opaque_fd – 使用 Linux fd 句柄共享外部內(nèi)存，類似于 VK_KHR_external_memory_fd 。

cl_khr_external_memory_win32 – 使用 win32 NT 和 KMT 句柄共享外部內(nèi)存，類似于 VK_KHR_external_memory_win32 。

使用 OpenCL

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

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

檢查底層 OpenCL 平臺和帶有clGetPlatformInfo和clGetDeviceInfo的設(shè)備是否支持所需的擴(kuò)展cl_khr_external_semaphore和cl_khr_external_memory。

為了能夠使用 Win32 信號量和內(nèi)存句柄，請檢查cl_khr_external_semaphore_win32_khr和cl_khr_external_memory_win32_khr擴(kuò)展是否存在。

為了能夠使用 FD 信號量和內(nèi)存句柄，請檢查cl_khr_external_semaphore_opaque_fd_khr和cl_khr_external_memory_opaque_fd_khr擴(kuò)展是否存在。這也可以通過查詢支持的句柄類型來檢查。

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

// 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);

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

// 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 導(dǎo)出的外部內(nèi)存作為緩沖內(nèi)存與clCreateBufferWithProperties一起導(dǎo)入。

// 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 導(dǎo)出的外部內(nèi)存作為圖像內(nèi)存導(dǎo)入。

// 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'

關(guān)于作者

Nikhil Joshi 目前在NVIDIA 管理 OpenCL 驅(qū)動(dòng)程序團(tuán)隊(duì)。他還代表 NVIDIA 參加 Khronos OpenCL 工作組。他在 NVIDIA 的計(jì)算團(tuán)隊(duì)工作了 10 多年，致力于不同的計(jì)算 API ，包括 CUDA 、 Renderscript 和 OpenCL

Rekha Mukund 是 NVIDIA 計(jì)算組的產(chǎn)品經(jīng)理，負(fù)責(zé)為汽車、 Jetson 和 Android 平臺開發(fā) CUDA Tegra 產(chǎn)品。她還負(fù)責(zé)管理 NVIDIA SimNet 產(chǎn)品和 OpenCL 計(jì)劃。在加入 NVIDIA 之前， Rekha 在付費(fèi)電視技術(shù)領(lǐng)域與思科合作了八年多。她是英國大學(xué)計(jì)算機(jī)科學(xué)學(xué)院的金牌獲得者，他是印度國家級乒乓球運(yùn)動(dòng)員和狂熱的旅行者。

審核編輯：郭婷