Developer Blog

How to share Vulkan and OpenCL workloads on Android with Qualcomm Adreno GPUs

Written by

Zimu Guan

Oct 22, 2026

Android GPU pipelines may use Vulkan for rendering while relying on existing OpenCL kernels for compute, machine learning, denoising or super-resolution. Moving intermediate results through the CPU adds unnecessary copying and synchronization overhead.

This tutorial shows how to share GPU memory and coordinate workloads directly between Vulkan and OpenCL on Android with Adreno GPUs, including the external-memory handles, ownership transfers and sync-FD semaphores required for a complete Vulkan-to-OpenCL-to-Vulkan workflow.

What you will build: A Vulkan → OpenCL → Vulkan pipeline that shares GPU memory directly, runs an OpenCL kernel on the shared resource and returns the result to Vulkan without CPU staging.

Note for the team: The configuration and feature support described here generally apply to Snapdragon 8 Elite Gen 5 processor and later products using the target Android driver stack. For accuracy, this article assumes support for the required capabilities. Actual support remains device- and driver-dependent, so applications should treat the runtime capability queries as authoritative.

Android Interop Extensions for Adreno GPUs

Vulkan/OpenCL interop isn't one extension or one handle type. Memory sharing and synchronization use different external-object mechanisms. For Adreno GPUs on Android, those choices look like this:

Purpose	Vulkan side	OpenCL side
Querying external support	VK_KHR_external_memory_capabilities, VK_KHR_external_semaphore_capabilities	OpenCL platform/device extension queries
Matching devices and drivers	VkPhysicalDeviceIDProperties::{deviceUUID, driverUUID}	cl_khr_device_uuid, CL_DEVICE_UUID_KHR, CL_DRIVER_UUID_KHR
Cross-API memory sharing	VK_KHR_external_memory with handle-specific extensions (see below)	cl_khr_external_memory with handle-specific extensions (see below)
Cross-API synchronization	VK_KHR_external_semaphore, VK_KHR_external_semaphore_fd, VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD_BIT	cl_khr_external_semaphore, cl_khr_external_semaphore_sync_fd, CL_SEMAPHORE_HANDLE_SYNC_FD_KHR

Cross-API memory sharing here is a bit complex since there are multiple options available, as shown in the table below.

Cross-API memory sharing approach	Vulkan side	OpenCL side	Typical use and restriction
Generic opaque fd	VK_KHR_external_memory_fd, VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT_KHR	cl_khr_external_memory_opaque_fd, CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR	Buffers and linear-tiling images. Does not provide reliable tiling inference for optimal-tiling images.
Qualcomm Vulkan opaque fd	Same Vulkan extensions and handle type as generic opaque fd	cl_qcom_external_memory_vulkan_opaque_fd, CL_EXTERNAL_MEMORY_HANDLE_VULKAN_OPAQUE_FD_QCOM	Vulkan images with format and internal memory layout inference. Qualcomm-specific.
Android Hardware Buffer	VK_ANDROID_external_memory_android_hardware_buffer, VK_EXTERNAL_MEMORY_HANDLE_TYPE_ANDROID_HARDWARE_BUFFER_BIT_ANDROID	cl_khr_external_memory_android_hardware_buffer	Android-native sharing with format, dimensions, and data layout inference.
DMA-BUF	VK_KHR_external_memory_fd, VK_EXT_external_memory_dma_buf, VK_EXTERNAL_MEMORY_HANDLE_TYPE_DMA_BUF_BIT_EXT	cl_khr_external_memory_dma_buf	Existing low-level Linux allocator, media, and display pipelines. Image format and layout require an additional agreement.

For sharing buffers and linear images, a generic opaque fd is usually sufficient. Optimal-tiling images, however, are trickier because the tiling layouts are implementation-defined, but an fd alone does not describe it.

Android Hardware Buffer handles this in an Android native way and is supported by multiple vendors. Qualcomm Technologies also provides a Vulkan opaque-fd extension that supplies the additional information needed for Vulkan images. This extension can be integrated into an existing opaque-fd implementation easily with a few lines of code changes.

DMA-BUF sits at a lower level. It is more common in system, media, camera, and display pipelines than in application-level Vulkan/OpenCL interop.

The rest of this article will focus on the generic and Qualcomm Technologies’ Vulkan opaque-fd approaches. We'll look at AHB-based Vulkan/OpenCL interop in a follow-up post - please stay tuned.

One more setup detail is required before using opaque-fd shared memory objects: Vulkan and OpenCL must refer to the same physical GPU and compatible driver. To achieve this, compare both:

VkPhysicalDeviceIDProperties::deviceUUID and VkPhysicalDeviceIDProperties::driverUUID with CL_DEVICE_UUID_KHR and CL_DRIVER_UUID_KHR respectively, as exposed by cl_khr_device_uuid.

Matching only the device UUID identifies the same GPU but is not sufficient for opaque-fd compatibility. We won't cover that matching code here.

Memory Sharing between APIs

Figure 1: Opaque-fd external-memory sharing between Vulkan and OpenCL.

Memory sharing gives Vulkan and OpenCL separate API objects that refer to the same GPU allocation. In our opaque-fd memory flow, Vulkan exports the allocation as VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT_KHR. OpenCL can then import that fd with either the generic opaque-fd handle or the Vulkan opaque-fd handle from Qualcomm Technologies, depending on the resource:

Vulkan resource imported via opaque fd	CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR	CL_EXTERNAL_MEMORY_HANDLE_VULKAN_OPAQUE_FD_QCOM
Buffer	Valid	Valid, equivalent behavior for buffers
Linear-tiling Vulkan image	Valid	Valid, with image format/layout inferred
Optimal-tiling Vulkan image	Usually not suitable for Adreno GPUs on Android	Required

The handle CL_EXTERNAL_MEMORY_HANDLE_VULKAN_OPAQUE_FD_QCOM isn't only for optimal-tiling images. For buffers, it behaves like the generic opaque-fd handle. For images, it lets the implementation infer the Vulkan image format and internal layout.

The distinction matters when importing an optimal-tiling Vulkan image. On Adreno, CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR generally appears in CL_DEVICE_EXTERNAL_MEMORY_IMPORT_ASSUME_LINEAR_IMAGES_HANDLE_TYPES_KHR, so OpenCL treats images imported through that handle as linear. Using it for an optimal-tiling Vulkan image may make Vulkan and OpenCL interpret the same allocation differently, resulting in undefined behavior.

CL_EXTERNAL_MEMORY_HANDLE_VULKAN_OPAQUE_FD_QCOM avoids this mismatch by allowing OpenCL to infer the Vulkan image format and internal layout.

Check out the Qualcomm OpenCL SDK for the declarations and full specification of cl_qcom_external_memory_vulkan_opaque_fd.

Now let's look at how memory sharing is set up and used in practice.

Setup: Generic Opaque-FD Import

We'll start with CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR, the generic opaque-fd handle. On Adreno, this is a valid choice for buffers and linear-tiling images.

The setup has four steps:

Create the Vulkan resource with external memory enabled.
Allocate exportable Vulkan memory and bind it to the resource.
Export the memory fd from Vulkan.
Import the fd into OpenCL and create a cl_mem.

Creating an Exportable Vulkan Resource

For a buffer, add VkExternalMemoryBufferCreateInfo to VkBufferCreateInfo::pNext:

VkExternalMemoryBufferCreateInfo external_buffer_info = {};
external_buffer_info.sType = VK_STRUCTURE_TYPE_EXTERNAL_MEMORY_BUFFER_CREATE_INFO;
external_buffer_info.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT_KHR;

VkBufferCreateInfo buffer_info = {};
buffer_info.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO;
buffer_info.pNext = &external_buffer_info;
buffer_info.size = size;
buffer_info.usage = usage;

vkCreateBuffer(device, &buffer_info, nullptr, &buffer);

For an image, use VkExternalMemoryImageCreateInfo:

VkExternalMemoryImageCreateInfo external_image_info = {};
external_image_info.sType = VK_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMAGE_CREATE_INFO;
external_image_info.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT_KHR;

VkImageCreateInfo image_info = {};
image_info.sType = VK_STRUCTURE_TYPE_IMAGE_CREATE_INFO;
image_info.pNext = &external_image_info;
image_info.imageType = VK_IMAGE_TYPE_2D;
image_info.format = VK_FORMAT_R8G8B8A8_UNORM;
image_info.extent = {width, height, 1};
image_info.mipLevels = 1;
image_info.arrayLayers = 1;
image_info.samples = VK_SAMPLE_COUNT_1_BIT;
image_info.tiling = VK_IMAGE_TILING_LINEAR;
image_info.usage = usage;

vkCreateImage(device, &image_info, nullptr, &image);

This generic import is for linear-tiling images. Optimal-tiling images are covered later.

Allocating Exportable Memory

Creating the resource isn't enough. The memory bound to the resource must also be exportable:

VkExportMemoryAllocateInfoKHR export_info = {};
export_info.sType = VK_STRUCTURE_TYPE_EXPORT_MEMORY_ALLOCATE_INFO_KHR;
export_info.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT_KHR;

VkMemoryAllocateInfo alloc_info = {};
alloc_info.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO;
alloc_info.pNext = &export_info;
alloc_info.allocationSize = memory_requirements.size;
alloc_info.memoryTypeIndex = memory_type_index;

Some external buffers and images may also require dedicated allocation. To check, query the external-memory properties for the opaque-fd handle type and inspect externalMemoryFeatures.

For buffers, use vkGetPhysicalDeviceExternalBufferProperties.

For images, use vkGetPhysicalDeviceImageFormatProperties2 with VkPhysicalDeviceExternalImageFormatInfo in the input chain and VkExternalImageFormatProperties in the output chain.

If externalMemoryFeatures includes VK_EXTERNAL_MEMORY_FEATURE_DEDICATED_ONLY_BIT, add VkMemoryDedicatedAllocateInfo to the allocation chain:

VkMemoryDedicatedAllocateInfo dedicated_info = {};

// external_memory_features is returned by the corresponding buffer or image capability query.
if ((external_memory_features &
     VK_EXTERNAL_MEMORY_FEATURE_DEDICATED_ONLY_BIT) != 0) {
    dedicated_info.sType = VK_STRUCTURE_TYPE_MEMORY_DEDICATED_ALLOCATE_INFO;
    dedicated_info.image = image; // for image allocation
    dedicated_info.buffer = VK_NULL_HANDLE;

    // for buffer allocation:
    // dedicated_info.image = VK_NULL_HANDLE;
    // dedicated_info.buffer = buffer;

    export_info.pNext = &dedicated_info;
}

Then allocate the memory and bind it:

vkAllocateMemory(device, &alloc_info, nullptr, &memory);
// for image
vkBindImageMemory(device, image, memory, 0);
// for buffer
// vkBindBufferMemory(device, buffer, memory, 0);

Exporting the FD from Vulkan

After the memory is allocated, export a POSIX fd from the VkDeviceMemory:

int fd = -1;

VkMemoryGetFdInfoKHR fd_info = {};
fd_info.sType = VK_STRUCTURE_TYPE_MEMORY_GET_FD_INFO_KHR;
fd_info.handleType = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT_KHR;
fd_info.memory = memory;

vkGetMemoryFdKHR(device, &fd_info, &fd);

The fd refers to the exported Vulkan memory allocation and can now be imported into OpenCL.

Importing the FD into OpenCL

OpenCL imports the exported fd when creating the memory object, using the following property list:

cl_mem_properties props[] = {
    CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR,
    (cl_mem_properties)fd,
    CL_MEM_DEVICE_HANDLE_LIST_KHR,
    (cl_mem_properties)cl_device,
    CL_MEM_DEVICE_HANDLE_LIST_END_KHR,
    0
};

For a buffer, create a cl_mem buffer:

cl_int err = CL_SUCCESS;
cl_mem cl_buffer = clCreateBufferWithProperties(
    /*context*/      context,
    /*properties*/   props,
    /*flags*/        flags,
    /*size*/         size,
    /*host_ptr*/     nullptr,
    /*errcode_ret*/  &err);

For a linear-tiling image, query the Vulkan subresource layout and pass its rowPitch to OpenCL before creating a cl_mem image. Leaving image_row_pitch at zero would make it implementation-defined for an external-memory image, and might not match Vulkan's padding.

// query the Vulkan subresource layout
VkImageSubresource subresource = {};
subresource.aspectMask = VK_IMAGE_ASPECT_COLOR_BIT;
subresource.mipLevel = 0;
subresource.arrayLayer = 0;

VkSubresourceLayout vulkan_layout = {};
vkGetImageSubresourceLayout(device, image, &subresource, &vulkan_layout);

// create OpenCL image
cl_image_format format = {};
format.image_channel_order = CL_RGBA;
format.image_channel_data_type = CL_UNORM_INT8;

cl_image_desc desc = {};
desc.image_type        = CL_MEM_OBJECT_IMAGE2D;
desc.image_width       = width;
desc.image_height      = height;
desc.image_depth       = 0; // only used for 3D image.
desc.image_array_size  = 0; // only used for 1D/2D image array 

// If image_row_pitch is zero and the image is created from an external
// memory handle, then the image row pitch is implementation-defined.
// Vulkan may pad each row. Using zero here would let OpenCL choose an
// implementation-defined pitch that may not match Vulkan's rowPitch.
desc.image_row_pitch   = (size_t)vulkan_layout.rowPitch;

// Not used for a non-array 2D image. For an image array, use the queried
// arrayPitch; for a 3D image, use depthPitch.
desc.image_slice_pitch = 0; 

desc.num_mip_levels    = 0; // must be 0 unless cl_khr_mipmap_image is supported
desc.num_samples       = 0; // must be 0
desc.buffer            = nullptr;

cl_mem cl_image = clCreateImageWithProperties(
    /*context*/      context,
    /*properties*/   props,
    /*flags*/        flags,
    /*image_format*/ &format,
    /*image_desc*/   &desc,
    /*host_ptr*/     nullptr,
    /*errcode_ret*/  &err);

The generic opaque-fd handle is suitable here only for linear-tiling Vulkan images. For an optimal-tiling image, keep the same Vulkan memory export workflow, but use the Qualcomm Technologies’ Vulkan opaque-fd import handle on the OpenCL side.

Setup: Qualcomm Technologies’ Vulkan Opaque-FD Import

The cl_qcom_external_memory_vulkan_opaque_fd extension adds CL_EXTERNAL_MEMORY_HANDLE_VULKAN_OPAQUE_FD_QCOM, an OpenCL import handle for Vulkan allocations exported by the Qualcomm Vulkan driver. The extension can be used more broadly, but the reason we need it in this post is optimal-tiling image import.

The flow is almost the same as the generic opaque-fd import: create an external Vulkan resource, allocate exportable memory, export the fd, and create a cl_mem from that fd. We'll start on the Vulkan side, where the external-memory setup is unchanged and the image itself is created as optimal tiling.

Creating the Vulkan Image

Create the Vulkan image as an external optimal-tiling image:

VkExternalMemoryImageCreateInfo external_image_info = {};
external_image_info.sType = VK_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMAGE_CREATE_INFO;
external_image_info.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT_KHR;

VkImageCreateInfo image_info = {};
image_info.sType = VK_STRUCTURE_TYPE_IMAGE_CREATE_INFO;
image_info.pNext = &external_image_info;
image_info.imageType = VK_IMAGE_TYPE_2D;
image_info.format = VK_FORMAT_R8G8B8A8_UNORM;
image_info.extent = {width, height, 1};
image_info.mipLevels = 1;
image_info.arrayLayers = 1;
image_info.samples = VK_SAMPLE_COUNT_1_BIT;
image_info.tiling = VK_IMAGE_TILING_OPTIMAL; // Different: optimal tiling.
image_info.usage = usage;

vkCreateImage(device, &image_info, nullptr, &image);

Then allocate exportable memory, bind it, and export the fd as before. Once we have that fd, the part that is specific to Qualcomm Technologies happens during OpenCL import.

Importing with the Qualcomm extension

Compared with the generic linear-tiling image import, the OpenCL object creation looks almost the same. The two important differences are the handle type and the null image_format:

cl_mem_properties props[] = {
    // Different: use the Qualcomm Vulkan opaque-fd handle on the OpenCL side.
    CL_EXTERNAL_MEMORY_HANDLE_VULKAN_OPAQUE_FD_QCOM,
    (cl_mem_properties)fd,
    CL_MEM_DEVICE_HANDLE_LIST_KHR,
    (cl_mem_properties)cl_device,
    CL_MEM_DEVICE_HANDLE_LIST_END_KHR,
    0
};

cl_image_desc desc = {};
desc.image_type        = CL_MEM_OBJECT_IMAGE2D;
desc.image_width       = width;
desc.image_height      = height;
desc.image_depth       = 0;
desc.image_array_size  = 0;
desc.image_row_pitch   = 0; // must be zero
desc.image_slice_pitch = 0; // must be zero
desc.num_mip_levels    = 0;
desc.num_samples       = 0;
desc.buffer            = nullptr;

cl_mem cl_image = clCreateImageWithProperties(
    /*context*/      context,
    /*properties*/   props,
    /*flags*/        flags,
    /*image_format*/ nullptr, // Different: must be nullptr, image format is inferred from the Vulkan image.
    /*image_desc*/   &desc,
    /*host_ptr*/     nullptr,
    /*errcode_ret*/  &err);

There are a few more image-specific rules for cl_qcom_external_memory_vulkan_opaque_fd to keep in mind:

cl_image_desc must match the Vulkan image.
Row pitch and slice pitch must be zero.
The image type must be CL_MEM_OBJECT_IMAGE2D, CL_MEM_OBJECT_IMAGE2D_ARRAY, or CL_MEM_OBJECT_IMAGE3D.
The Vulkan/OpenCL format pair must be listed in the extension specification. Importing an unlisted format may result in undefined behavior.
Partially resident Vulkan images are not supported.

See the cl_qcom_external_memory_vulkan_opaque_fd specification included with the Qualcomm OpenCL SDK for full details.

Render Loop: Acquiring and Releasing External Memory

During each frame, access to the shared memory must be explicitly transferred between Vulkan and OpenCL. Vulkan uses queue-family ownership barriers to release the shared resources to VK_QUEUE_FAMILY_EXTERNAL and acquire them back, while OpenCL uses clEnqueueAcquireExternalMemObjectsKHR and clEnqueueReleaseExternalMemObjectsKHR.

After the Vulkan producer work, release the image to VK_QUEUE_FAMILY_EXTERNAL:

// record Vulkan work that produces the shared resource
// record a barrier that releases ownership of the shared resource to the external queue family
VkImageMemoryBarrier release_to_opencl{};
release_to_opencl.sType               = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER;
release_to_opencl.srcAccessMask       = access_mask_before_release;
release_to_opencl.dstAccessMask       = 0;
release_to_opencl.oldLayout           = image_layout_before_release;
release_to_opencl.newLayout           = VK_IMAGE_LAYOUT_GENERAL;
release_to_opencl.srcQueueFamilyIndex = vulkan_queue_family_index;
release_to_opencl.dstQueueFamilyIndex = VK_QUEUE_FAMILY_EXTERNAL;
release_to_opencl.image               = shared_image;
release_to_opencl.subresourceRange    = {VK_IMAGE_ASPECT_COLOR_BIT, 0, 1, 0, 1};

vkCmdPipelineBarrier(cmd, stage_mask_before_release, VK_PIPELINE_STAGE_BOTTOM_OF_PIPE_BIT,
                     0, 0, nullptr, 0, nullptr, 1, &release_to_opencl);

// submit the command buffer with proper cross-API synchronization

With proper cross-API synchronization after this barrier, OpenCL can acquire the imported objects, use them, and release them:

// assume an in-order OpenCL command queue
cl_mem objects[] = {cl_input, cl_output};
cl_uint num_objects = 2;

clEnqueueAcquireExternalMemObjectsKHR(
    /*command_queue*/           queue,
    /*num_mem_objects*/         num_objects,
    /*mem_objects*/             objects,
    /*num_events_in_wait_list*/ 0,
    /*event_wait_list*/         nullptr,
    /*event*/                   nullptr);

clEnqueueNDRangeKernel(...);

clEnqueueReleaseExternalMemObjectsKHR(
    /*command_queue*/           queue,
    /*num_mem_objects*/         num_objects,
    /*mem_objects*/             objects,
    /*num_events_in_wait_list*/ 0,
    /*event_wait_list*/         nullptr,
    /*event*/                   nullptr);

Once the OpenCL work and memory release have been properly synchronized with Vulkan, Vulkan can acquire ownership back:

// record a barrier that acquires ownership of the shared image from the external queue family
VkImageMemoryBarrier acquire_from_opencl{};
acquire_from_opencl.sType               = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER;
acquire_from_opencl.srcAccessMask       = 0;
acquire_from_opencl.dstAccessMask       = access_mask_after_acquire;
acquire_from_opencl.oldLayout           = VK_IMAGE_LAYOUT_GENERAL;
acquire_from_opencl.newLayout           = image_layout_after_acquire;
acquire_from_opencl.srcQueueFamilyIndex = VK_QUEUE_FAMILY_EXTERNAL;
acquire_from_opencl.dstQueueFamilyIndex = vulkan_queue_family_index;
acquire_from_opencl.image               = shared_image;
acquire_from_opencl.subresourceRange    = {VK_IMAGE_ASPECT_COLOR_BIT, 0, 1, 0, 1};

vkCmdPipelineBarrier(cmd, VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT, stage_mask_after_acquire,
                     0, 0, nullptr, 0, nullptr, 1, &acquire_from_opencl);

// record Vulkan work that consumes the shared resource
// submit the command buffer with proper cross-API synchronization

Shared buffers follow the same ownership sequence using VkBufferMemoryBarrier. Use the Vulkan queue family and VK_QUEUE_FAMILY_EXTERNAL in the same directions, and set offset and size to the shared buffer range.

The Vulkan ownership barriers and the OpenCL acquire/release commands make the shared resources available to the receiving API. They don't, by themselves, order Vulkan work relative to OpenCL work. External semaphores provide that cross-API execution ordering.

Synchronization between APIs

Synchronization is the other half of the interop. Once both APIs can refer to the same memory, we still need to order the GPU work so neither API reads stale data or writes while the other API is still using the resource.

Figure 2: Sync-fd and opaque-fd external-semaphore lifetime models.

Different GPU implementations may expose different semaphore handle types for interop. A workflow based on a semaphore opaque-fd handle, such as the Unix implementation in the Khronos OpenCL interoperability sample, isn't equivalent to the sync-fd approach used on Android. The handle type changes the lifetime and ownership model of the semaphore payload.

Semaphore handle types	What the fd represents	Lifetime model	Android support
VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD_BIT, CL_SEMAPHORE_HANDLE_OPAQUE_FD_KHR	A handle to a semaphore payload/state shared by both API objects	Persistent linkage after setup	No
VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD_BIT, CL_SEMAPHORE_HANDLE_SYNC_FD_KHR	A Linux sync-file / Android fence payload for one synchronization point	One-shot; export/import again for the next synchronization point	Yes

With a semaphore opaque-fd handle, the fd is a handle to an underlying semaphore payload. After one API exports it and the other API imports it, the Vulkan semaphore and the OpenCL semaphore refer to the same signal state. That relationship is persistent: set up the shared semaphore objects once, then use them in the frame loop.

A sync fd is different. A sync fd is a Linux sync-file / Android fence for one synchronization point. Importing it doesn't permanently link the Vulkan and OpenCL semaphore objects. It transfers the current wait/signaled state as a temporary payload. Once that payload is exported, it should be treated as consumed on the exporting side.

This is why sync-fd payloads need export/import at each handoff. For every Vulkan-to-OpenCL or OpenCL-to-Vulkan synchronization point, export a fresh sync fd, import it into the other API's semaphore, then wait or signal that semaphore in the normal queue work.

With that one-shot model in mind, we'll consider a Vulkan -> OpenCL -> Vulkan frame throughout the rest of this post. This flow requires two synchronization directions: Vulkan to OpenCL and OpenCL back to Vulkan.

Setup: Create Cross-API Semaphores

We'll use one semaphore pair for each direction. Each pair has a Vulkan-side semaphore object and an OpenCL-side semaphore object, so this example creates two Vulkan semaphores and two OpenCL semaphores.

Creating Vulkan Semaphores

Create the Vulkan semaphore that will be exported to OpenCL with sync-fd export capability:

VkExportSemaphoreCreateInfo export_info = {};
export_info.sType = VK_STRUCTURE_TYPE_EXPORT_SEMAPHORE_CREATE_INFO;
export_info.handleTypes = VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD_BIT;

VkSemaphoreCreateInfo semaphore_info = {};
semaphore_info.sType = VK_STRUCTURE_TYPE_SEMAPHORE_CREATE_INFO;
semaphore_info.pNext = &export_info;

vkCreateSemaphore(device, &semaphore_info, nullptr, &vk_sema_vk_to_cl);

Create another Vulkan semaphore that will receive the OpenCL signal:

VkSemaphoreCreateInfo semaphore_info = {};
semaphore_info.sType = VK_STRUCTURE_TYPE_SEMAPHORE_CREATE_INFO;

vkCreateSemaphore(device, &semaphore_info, nullptr, &vk_sema_cl_to_vk);

Creating OpenCL Semaphores

Create the Vulkan-to-OpenCL semaphore with -1, a special placeholder for an already-signaled sync-fd payload. The import succeeds, and the semaphore receives a temporary payload as if a valid fd had been provided:

int init_fd = -1; // -1 fd is treated like a valid sync fd referring to an object that has already signaled

cl_semaphore_properties_khr import_props[] = {
    CL_SEMAPHORE_TYPE_KHR,
    CL_SEMAPHORE_TYPE_BINARY_KHR,
    CL_SEMAPHORE_HANDLE_SYNC_FD_KHR,
    (cl_semaphore_properties_khr)init_fd,
    CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR,
    (cl_semaphore_properties_khr)cl_device,
    CL_SEMAPHORE_DEVICE_HANDLE_LIST_END_KHR,
    0
};

cl_semaphore_khr cl_sema_vk_to_cl = clCreateSemaphoreWithPropertiesKHR(context, import_props, &err);

Create the OpenCL-to-Vulkan semaphore with sync-fd export enabled:

cl_semaphore_properties_khr export_props[] = {
    CL_SEMAPHORE_TYPE_KHR,
    CL_SEMAPHORE_TYPE_BINARY_KHR,
    CL_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR,
    CL_SEMAPHORE_HANDLE_SYNC_FD_KHR,
    CL_SEMAPHORE_EXPORT_HANDLE_TYPES_LIST_END_KHR,
    CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR,
    (cl_semaphore_properties_khr)cl_device,
    CL_SEMAPHORE_DEVICE_HANDLE_LIST_END_KHR,
    0
};

cl_semaphore_khr cl_sema_cl_to_vk = clCreateSemaphoreWithPropertiesKHR(context, export_props, &err);

At this point, we've created the reusable semaphore objects, but we haven't exchanged a valid sync fd. Because sync-fd payloads are one-shot, a fresh fd will be exported and imported for each cross-API synchronization point in the render loop, not here.

Render Loop: Export, Import, Wait, and Signal

Now let's move from setup to the render loop.

The semaphore objects created during setup are reused across frames, but the sync-fd payloads are one-shot. Each time work crosses from one API to the other, the signaling side exports a new sync fd and the waiting side imports it before waiting.

Our Vulkan -> OpenCL -> Vulkan frame therefore requires two handoffs:

Vulkan -> OpenCL:
    Vulkan work signals a semaphore
    Vulkan exports a sync fd
    OpenCL re-imports the fd
    OpenCL work waits on the semaphore

OpenCL -> Vulkan:
    OpenCL work signals a semaphore
    OpenCL exports a sync fd
    Vulkan imports the fd
    Vulkan work waits on the semaphore

Vulkan to OpenCL

After submitting the Vulkan work with the Vulkan-to-OpenCL semaphore set to be signaled, export the semaphore payload as a sync fd. At this point, the semaphore may already be signaled or may have a signal operation pending execution; both states are valid for export.

// Submit Vulkan work that produces the shared input with the 
// ownership-release barriers, and queues a signal operation 
// for vk_sema_vk_to_cl before exporting its payload.

int fd = -1;

VkSemaphoreGetFdInfoKHR fd_info = {};
fd_info.sType = VK_STRUCTURE_TYPE_SEMAPHORE_GET_FD_INFO_KHR;
fd_info.handleType = VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD_BIT;
fd_info.semaphore = vk_sema_vk_to_cl;

vkGetSemaphoreFdKHR(device, &fd_info, &fd);

Then re-import the fd into the existing OpenCL semaphore:

clReImportSemaphoreSyncFdKHR(
    /*sema_object*/ cl_sema_vk_to_cl,
    /*properties*/  nullptr,
    /*fd*/          fd);

OpenCL waits on the re-imported semaphore before acquiring shared memory:

clEnqueueWaitSemaphoresKHR(
    /*command_queue*/           queue,
    /*num_sema_objects*/        1,
    /*sema_objects*/            &cl_sema_vk_to_cl,
    /*sema_payload_list*/       nullptr,
    /*num_events_in_wait_list*/ 0,
    /*event_wait_list*/         nullptr,
    /*event*/                   nullptr);

Next, enqueue the acquire, kernel dispatch, and release commands as shown in the memory-sharing section. They will execute after the semaphore wait completes.

OpenCL to Vulkan

Enqueue the OpenCL-to-Vulkan semaphore signal so that it executes after the kernel dispatch and memory release:

clEnqueueSignalSemaphoresKHR(
    /*command_queue*/           queue,
    /*num_sema_objects*/        1,
    /*sema_objects*/            &cl_sema_cl_to_vk,
    /*sema_payload_list*/       nullptr,
    /*num_events_in_wait_list*/ 0,
    /*event_wait_list*/         nullptr,
    /*event*/                   nullptr);

Then flush the OpenCL queue:

clFlush(queue);

This is important, because enqueuing an OpenCL command doesn't necessarily submit it to the device immediately; the command may remain queued until an explicit or implicit flush. Before exporting the sync fd, make sure the wait, acquire, kernel dispatch, release, and signal commands have been issued to the device. clFlush() does that without blocking until completion.

Now export the OpenCL semaphore as a sync fd:

int fd = -1;

clGetSemaphoreHandleForTypeKHR(
    /*sema_object*/ cl_sema_cl_to_vk,
    /*device*/      cl_device,
    /*handle_type*/ CL_SEMAPHORE_HANDLE_SYNC_FD_KHR,
    /*handle_size*/ sizeof(fd),
    /*handle_ptr*/  &fd,
    /*size_ret*/    nullptr);

Import that fd into Vulkan:

VkImportSemaphoreFdInfoKHR import_info = {};
import_info.sType = VK_STRUCTURE_TYPE_IMPORT_SEMAPHORE_FD_INFO_KHR;
import_info.semaphore = vk_sema_cl_to_vk;
import_info.flags = VK_SEMAPHORE_IMPORT_TEMPORARY_BIT;
import_info.handleType = VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD_BIT;
import_info.fd = fd;

vkImportSemaphoreFdKHR(device, &import_info);

The next Vulkan queue submission waits on vk_sema_cl_to_vk, so Vulkan doesn't read the shared output until OpenCL has released it and signaled completion.

Putting It Together

We've looked at memory sharing and synchronization separately. Now let's organize the full workflow around what we create once, and what we refresh every frame.

During setup:

1. Match the Vulkan physical device and driver with OpenCL by comparing both the device UUID and driver UUID.

2. Query both APIs for the external-memory and external-semaphore capabilities required by the selected handle types and resource configuration.

3. Create the Vulkan buffer or image with external memory enabled.

4. Allocate exportable Vulkan memory, bind it to the resource, and export it as VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT_KHR.

5. Import that fd into OpenCL:

For buffers and linear-tiling images, either CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR or CL_EXTERNAL_MEMORY_HANDLE_VULKAN_OPAQUE_FD_QCOM can be used.
For optimal-tiling Vulkan images, use CL_EXTERNAL_MEMORY_HANDLE_VULKAN_OPAQUE_FD_QCOM.

6. Create sync-fd semaphore objects for the two synchronization directions: use VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD_BIT on Vulkan and CL_SEMAPHORE_HANDLE_SYNC_FD_KHR on OpenCL.

During each frame:

Vulkan submits the producer work and ownership-release barriers, queuing a signal operation for the VK-to-CL VkSemaphore.
Vulkan exports the VK-to-CL VkSemaphore payload as a fresh sync fd.
OpenCL re-imports the sync fd into the VK-to-CL cl_semaphore_khr.
OpenCL enqueues a wait on the imported cl_semaphore_khr.
OpenCL enqueues acquisition of the shared objects.
OpenCL enqueues the kernel dispatches.
OpenCL enqueues release of the shared objects.
OpenCL enqueues a signal on the CL-to-VK cl_semaphore_khr.
OpenCL calls clFlush() to issue the queued commands to the device.
OpenCL exports the CL-to-VK cl_semaphore_khr payload as a fresh sync fd.
Vulkan imports the sync fd into the CL-to-VK VkSemaphore.
Vulkan submits the ownership-acquire barriers and consumer work, waiting on the imported VkSemaphore.

The per-frame pseudocode looks like this:

while (rendering) {
    // Submits the producer work and ownership-release barriers,
    // with a signal operation on the Vulkan-to-OpenCL semaphore.
    submit_vulkan_producer_release_and_signal(vk_sema_vk_to_cl);

    int vk_to_cl_fd = export_vk_semaphore_sync_fd(vk_sema_vk_to_cl);
    clReImportSemaphoreSyncFdKHR(cl_sema_vk_to_cl, nullptr, vk_to_cl_fd);
    
    clEnqueueWaitSemaphoresKHR(queue, 1, &cl_sema_vk_to_cl, nullptr, 0, nullptr, nullptr);

    cl_mem objects[] = {cl_input, cl_output};
    clEnqueueAcquireExternalMemObjectsKHR(queue, 2, objects, 0, nullptr, nullptr);
    clEnqueueNDRangeKernel(...);
    clEnqueueReleaseExternalMemObjectsKHR(queue, 2, objects, 0, nullptr, nullptr);

    clEnqueueSignalSemaphoresKHR(queue, 1, &cl_sema_cl_to_vk, nullptr, 0, nullptr, nullptr);
    
    clFlush(queue);

    int cl_to_vk_fd = export_cl_semaphore_sync_fd(cl_sema_cl_to_vk);
    import_vk_semaphore_sync_fd(vk_sema_cl_to_vk, cl_to_vk_fd);

    // Submits the ownership-acquire barriers and consumer work,
    // with a wait operation on the OpenCL-to-Vulkan semaphore.
    submit_vulkan_wait_acquire_and_consume(vk_sema_cl_to_vk);
    
    // Before the next iteration re-imports vk_sema_cl_to_vk, ensure this
    // Vulkan submission has completed, e.g. by waiting on a per-frame fence.
}

OpenCL Interop sample in Adreno GPU Vulkan code sample framework

Figure 3: Vulkan renders the scene, OpenCL converts it to grayscale, and Vulkan presents the result in the OpenCL interop sample.

The Adreno GPU Vulkan Code Sample Framework includes an opencl_interop sample, which is a small end-to-end version of the flow we've been building up. Vulkan renders the scene color, OpenCL turns it into grayscale, and Vulkan uses the result in the final blit.

To make the memory interop easy to compare, the sample exposes three modes: buffer interop, linear-tiling image interop, and optimal-tiling image interop. The first two use the generic opaque-fd handle. The optimal-tiling image mode uses cl_qcom_external_memory_vulkan_opaque_fd. Around those modes, the sample uses sync-fd semaphores to move execution from Vulkan to OpenCL and back again.

Check out the OpenCL interop sample for the full implementation.

Conclusion

For Adreno GPUs on Android, the Vulkan/OpenCL interop approach covered in this post can be summarized as follows:

For memory sharing, one practical solution is to use opaque fd. Vulkan exports the shared allocation as opaque fd. OpenCL can then import buffers and linear-tiling images with the generic opaque-fd handle when that's enough, or use cl_qcom_external_memory_vulkan_opaque_fd when importing an optimal-tiling Vulkan image.

Use sync fd for cross-API synchronization. Don't treat it like a semaphore opaque-fd handle: sync fd is a one-shot payload, so every cross-API synchronization point needs a fresh export/import operation.

Finally, remember that memory ownership and execution ordering are separate. Vulkan queue-family barriers and OpenCL acquire/release commands establish access to the shared resources on their respective sides; external semaphores order those operations and the Vulkan/OpenCL workloads.