cl_khr_external_memory(3) Manual Page

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

// Get cl_devices of the platform.
clGetDeviceIDs(..., &devices, &deviceCount);

// Create cl_context with first two devices
clCreateContext(..., 2, devices, ...);

// Create img of type cl_mem usable only on devices[0]

// 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;
image_desc.image_array_size = num_slices;
image_desc.image_row_pitch = width * 8 * 4; // May need alignment
image_desc.image_slice_pitch = image_desc.image_row_pitch * height;
image_desc.num_mip_levels = 1;
image_desc.num_samples = 0;
image_desc.buffer = NULL;

cl_mem_properties_khr extMemProperties[] = {
    (cl_mem_properties_khr)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR,
    (cl_mem_properties_khr)fd,
    (cl_mem_properties_khr)CL_MEM_DEVICE_HANDLE_LIST_KHR,
    (cl_mem_properties_khr)devices[0],
    CL_MEM_DEVICE_HANDLE_LIST_END_KHR,
    0
};

cl_mem img = clCreateImageWithProperties(/*context*/        clContext,
                                         /*properties*/     extMemProperties,
                                         /*flags*/          0,
                                         /*image_format*/   &clImgFormat,
                                         /*image_desc*/     &image_desc,
                                         /*errcode_ret*/    &errcode_ret);

// Use clGetImageInfo to get cl_image_format details.
size_t clImageFormatSize;
clGetImageInfo(img,
               CL_IMAGE_FORMAT,
               sizeof(cl_image_format),
               &clImageFormat,
               &clImageFormatSize);

// Start the main rendering loop

// Create extSem of type cl_semaphore_khr using clCreateSemaphoreWithPropertiesKHR

// Create extMem of type cl_mem using clCreateBufferWithProperties or clCreateImageWithProperties

while (true) {
    // (not shown) Signal the semaphore from the other API

    // Wait for the semaphore in OpenCL, by calling clEnqueueWaitSemaphoresKHR on 'extSem'
    clEnqueueWaitSemaphoresKHR(/*command_queue*/            command_queue,
                               /*num_sema_objects*/         1,
                               /*sema_objects*/             &extSem,
                               /*sema_payload_list*/        NULL,
                               /*num_events_in_wait_list*/  0,
                               /*event_wait_list*/          NULL,
                               /*event*/                    NULL);

    // Launch kernel that accesses extMem
    clEnqueueNDRangeKernel(command_queue, ...);

    // Signal the semaphore in OpenCL
    clEnqueueSignalSemaphoresKHR(/*command_queue*/           command_queue,
                                 /*num_sema_objects*/        1,
                                 /*sema_objects*/            &extSem,
                                 /*sema_payload_list*/       NULL,
                                 /*num_events_in_wait_list*/ 0,
                                 /*event_wait_list*/         NULL,
                                 /*event*/                   NULL);

    // (not shown) Launch work in other API that waits on 'extSem'
}

// Create extSem of type cl_semaphore_khr using
// clCreateSemaphoreWithPropertiesKHR with CL_SEMAPHORE_HANDLE_*_KHR.

// Create extMem1 and extMem2 of type cl_mem using clCreateBufferWithProperties
// or clCreateImageWithProperties

while (true) {
    // (not shown) Signal the semaphore from the other API. Wait for the
    // semaphore in OpenCL, by calling clEnqueueWaitForSemaphore on extSem
    clEnqueueWaitSemaphoresKHR(/*command_queue*/            cq1,
                               /*num_sema_objects*/         1,
                               /*sema_objects*/             &extSem,
                               /*sema_payload_list*/        NULL,
                               /*num_events_in_wait_list*/  0,
                               /*event_wait_list*/          NULL,
                               /*event*/                    NULL);

    // Get explicit ownership of extMem1
    clEnqueueAcquireExternalMemObjectsKHR(/*command_queue*/             cq1,
                                          /*num_mem_objects*/           1,
                                          /*mem_objects*/               &extMem1,
                                          /*num_events_in_wait_list*/   0,
                                          /*event_wait_list*/           NULL,
                                          /*event*/                     NULL);

    // Launch kernel that accesses extMem1 on cq1 on cl_device1
    clEnqueueNDRangeKernel(cq1,  ..., &event1);

    // Launch kernel that accesses both extMem1 and extMem2 on cq2 on cl_device2
    // Migration of extMem1 and extMem2 handles through regular CL memory
    // migration.
    clEnqueueNDRangeKernel(cq2, ..., &event1, &event2);

    // Give up ownership of extMem1 before you signal the semaphore. Handle
    // memory migration here.
    clEnqueueReleaseExternalMemObjectsKHR(/*command_queue*/           cq2
                                          /*num_mem_objects*/         1,
                                          /*mem_objects*/             &extMem1,
                                          /*num_events_in_wait_list*/ 0,
                                          /*event_wait_list*/         NULL,
                                          /*event*/                   NULL);

    // Signal the semaphore from OpenCL
    clEnqueueSignalSemaphoresKHR(/*command_queue*/           cq2,
                                 /*num_sema_objects*/        1,
                                 /*sema_objects*/            &extSem,
                                 /*sema_payload_list*/       NULL,
                                 /*num_events_in_wait_list*/ 0,
                                 /*event_wait_list*/         NULL,
                                 /*event*/                   NULL);

    // (not shown) Launch work in other API that waits on 'extSem'
    // Other API accesses ext1, but not ext2 on device-1
}