Using Semaphore and Memory Sharing Extensions for Vulkan Interop with NVIDIA OpenCL

OpenCL is evolving the way it can interoperate with other APIs, such as Vulkan. This post gives you a tour of the new style of OpenCL interop, which is already supported in the latest NVIDIA drivers. We include downloadable sample code so you can try this new functionality today.

The need for a new way to interop

Developers often use OpenCL for compute together with other APIs, such as OpenGL, to access functionality including graphics rendering. OpenCL has long enabled the sharing of implicit buffer and image objects with OpenGL, OpenGL ES, EGL, Direct3D 10, and Direct3D 11 through extensions:

cl_khr_gl_sharing
cl_khr_gl_event
cl_khr_egl_image
cl_khr_egl_event
cl_khr_d3d10_sharing
cl_khr_d3d11_sharing

New generation GPU APIs such as Vulkan use explicit references to external memory together with semaphores to coordinate access to shared resources. Until now, there have been no OpenCL extensions to enable external memory and semaphore sharing with this new class of API.

Interop between OpenCL and Vulkan has been in strong demand for both mobile and desktop platforms. NVIDIA has closely worked with the Khronos OpenCL Working Group to release a set of provisional cross-vendor KHR extensions. The extensions enable applications to efficiently share data between OpenCL and APIs such as Vulkan, with significantly increased flexibility compared to previous-generation interop APIs using implicit resources.

This set of new external memory and semaphore sharing extensions provide a generic framework that enables OpenCL to import external memory and semaphore handles exported by external APIs, using a methodology that will be familiar to Vulkan developers. OpenCL then uses those semaphores to synchronize the external runtime, coordinating the use of shared memory.

Diagram shows how OpenCL imports memory and semaphore handles from Vulkan, and uses semaphores to synchronize memory ownership and access. — *Figure 1. Interoperability relationship between OpenCL and Vulkan software*

External API-specific interop extensions can then be added to handle the details of interacting with specific APIs. Vulkan interop is available today, and additional APIs, such as DirectX 12, are planned.

The OpenCL new external semaphore and memory sharing functionality includes separate sets of carefully structured extensions.

Semaphore extensions

This set of extensions adds the ability to create OpenCL semaphore objects from OS-specific semaphore handles.

cl_khr_semaphore—Represents semaphores with wait and signal. This is a new class of OpenCL objects.
cl_khr_external_semaphore—Extends cl_khr_semaphore with mechanisms for importing and exporting external semaphores, similar to VK_KHR_external_semaphore.

The following extensions extend cl_khr_external_semaphore with handle-type-specific behavior:

cl_khr_external_semaphore_opaque_fd—Shares external semaphores using Linux fd handles with reference transference, similar to VK_KHR_external_semaphore_fd.
cl_khr_external_semaphore_win32—Shares external semaphores using win32 NT and KMT handles with reference transference, similar to VK_KHR_external_semaphore_win32.

Memory extensions

These extensions add the ability to create OpenCL memory objects from OS-specific memory handles. They have a similar design to the Vulkan external memory extension VK_KHR_external_memory.

cl_khr_external_memory—Imports external memory from other APIs.

The following extensions extend cl_khr_external_memory with handle-type-specific behavior:

cl_khr_external_memory_opaque_fd—Shares external memory using Linux fd handles, similar to VK_KHR_external_memory_fd.
cl_khr_external_memory_win32—Shares external memory using win32 NT and KMT handles, similar to VK_KHR_external_memory_win32.

Using OpenCL

The typical interop use case consists of the following steps.

Check if the required support is available:

Check if the required extensions cl_khr_external_semaphore and cl_khr_external_memory are supported by the underlying OpenCL platform and devices with clGetPlatformInfo and clGetDeviceInfo.
To be able to use Win32 semaphore and memory handles, check if the cl_khr_external_semaphore_win32_khr and cl_khr_external_memory_win32_khr extensions are present.
To be able to use FD semaphore and memory handles, check if the cl_khr_external_semaphore_opaque_fd_khr and cl_khr_external_memory_opaque_fd_khr extensions are present. This can also be checked by querying the supported handle types.

Importing external semaphores requires cl_khr_external_semaphore. If cl_khr_external_semaphore_opaque_fd is supported, you can import external semaphores exported by Vulkan using FD handles in OpenCL with clCreateSemaphoreWithPropertiesKHR.

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

Importing images requires cl_khr_external_memory and support for images. In OpenCL, import external semaphores exported by Vulkan using Win32 handles with clCreateSemaphoreWithPropertiesKHR.

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

In OpenCL, import external memory exported by Vulkan using the FD handle as buffer memory with 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);

In OpenCL, import external memory exported by Vulkan as image memory using clCreateImageWithProperties.

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

Synchronize between OpenCL and Vulkan using semaphore wait and signal.

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