Implementing and Verifying OpenCL 2.2 Shared Virtual Memory for Zero‑Copy Host‑Device Data Sharing

Desired outcome

Enable zero‑copy data sharing between the host and an OpenCL 2.2 device using Shared Virtual Memory (SVM) so that the same pointer can be accessed by both sides without explicit copy commands.

Prerequisites

• An OpenCL platform that implements version 2.2 or newer.
• A device that reports non‑zero CL_DEVICE_SVM_CAPABILITIES.
• A C compiler capable of building OpenCL kernels with -cl-std=CL2.2.
• Basic familiarity with the OpenCL API (context, command queue, program, kernel).

Procedure

1. Detect SVM support

cl_uint svmCaps = 0;
cl_int err = clGetDeviceInfo(device,
                             CL_DEVICE_SVM_CAPABILITIES,
                             sizeof(svmCaps),
                             &svmCaps,
                             NULL);
if (err != CL_SUCCESS || svmCaps == 0) {
    fprintf(stderr, \"SVM not supported on this device\\n\");
    /* fallback to regular buffers */
}

The capability flags indicate which SVM variants are supported: coarse‑grained, fine‑grained, fine‑grained system, and atomic support.

2. Allocate an SVM buffer

size_t size = 1024 * 1024; /* 1 MiB */
void *svmPtr = clSVMAlloc(context,
                          CL_MEM_READ_WRITE |
                          CL_MEM_SVM_FINE_GRAIN_BUFFER,
                          size,
                          64); /* 64‑byte alignment required for fine‑grained SVM */
if (svmPtr == NULL) {
    fprintf(stderr, \"SVM allocation failed\\n\");
    /* fallback */
}

The alignment argument must be at least 64 bytes when fine‑grained SVM is used; otherwise the behaviour is undefined.

3. Initialize the buffer on the host

Because fine‑grained SVM allows the host to access the memory directly, you can write a pattern after ensuring visibility with a map or a fence.

/* Make sure the host sees the latest contents */
clEnqueueSVMMap(queue,
                CL_TRUE,
                CL_MAP_WRITE,
                svmPtr,
                size,
                0,
                NULL,
                NULL);
/* Fill with a known pattern, e.g., increasing 32‑bit integers */
cl_uint *data = (cl_uint *)svmPtr;
for (size_t i = 0; i < size / sizeof(cl_uint); ++i) {
    data[i] = (cl_uint)i;
}
clEnqueueSVMUnmap(queue,
                  svmPtr,
                  0,
                  NULL,
                  NULL);

The map/unmap pair guarantees coherency; alternatively a fine‑grained SVM system with atomics could rely on explicit memory fences.

4. Write a simple kernel that increments each element

__kernel void inc(__global uint *ptr, uint count)
{
    uint gid = get_global_id(0);
    if (gid < count) {
        ptr[gid] = ptr[gid] + 1;
    }
}

Save this source as inc.cl and compile with -cl-std=CL2.2.

5. Set kernel arguments and enqueue

cl_kernel kernel = clCreateKernel(program, \"inc\", &err);
clSetKernelArg(kernel, 0, sizeof(void *), &svmPtr);
cl_uint count = (cl_uint)(size / sizeof(cl_uint));
clSetKernelArg(kernel, 1, sizeof(count), &count);

size_t global = size / sizeof(cl_uint);
cl_event kernelEvent = NULL;
err = clEnqueueNDRangeKernel(queue,
                             kernel,
                             1,
                             NULL,
                             &global,
                             NULL,
                             0,
                             NULL,
                             &kernelEvent);
if (err != CL_SUCCESS) {
    fprintf(stderr, \"Kernel enqueue failed: %d\\n\", err);
}

Because the kernel receives an SVM pointer, no cl_mem object is needed.

6. Ensure host sees the results

After kernel completion, map the SVM for reading or rely on fine‑grained coherency if a fence was issued.

clWaitForEvents(1, &kernelEvent);
clEnqueueSVMMap(queue,
                CL_TRUE,
                CL_MAP_READ,
                svmPtr,
                size,
                0,
                NULL,
                NULL);
/* Verify */
cl_uint *result = (cl_uint *)svmPtr;
bool ok = true;
for (size_t i = 0; i < size / sizeof(cl_uint); ++i) {
    if (result[i] != ((cl_uint)i + 1)) {
        ok = false;
        break;
    }
}
clEnqueueSVMUnmap(queue,
                  svmPtr,
                  0,
                  NULL,
                  NULL);
if (!ok) {
    fprintf(stderr, \"Verification failed\\n\");
}

7. Performance check – confirm zero‑copy

Query profiling information on the kernel event to ensure no copy commands were issued.

cl_ulong start, end;
clGetEventProfilingInfo(kernelEvent,
                        CL_PROFILING_COMMAND_START,
                        sizeof(start),
                        &start,
                        NULL);
clGetEventProfilingInfo(kernelEvent,
                        CL_PROFILING_COMMAND_END,
                        sizeof(end),
                        &end,
                        NULL);
double kernelTime = (end - start) * 1e-6; /* milliseconds */
printf(\"Kernel execution time: %.3f ms\\n\", kernelTime);
/* Additionally, you can enable tracing or look at the command queue
   to verify that no CL_COMMAND_COPY_BUFFER events appear. */

Expected checks

• Return code of clGetDeviceInfo for SVM capabilities is CL_SUCCESS and the value is non‑zero.
• clSVMAlloc returns a non‑NULL pointer.
• Kernel enqueue returns CL_SUCCESS and the event completes without error.
• Verification loop shows every element incremented by one.
• Profiling shows a reasonable kernel time and no copy‑related commands in the queue trace.

Recovery options (fallback)

• If SVM capabilities are zero or clSVMAlloc fails, allocate a regular buffer with clCreateBuffer.
• Copy data to the device with clEnqueueCopyBuffer before kernel execution and copy back after.
• Use the same kernel but pass a cl_mem argument instead of an SVM pointer.

Limitations

• Fine‑grained SVM requires the kernel to be compiled with the CL_KERNEL_REQUIRE_SVM_FINE_GRAIN_BUFFER attribute; otherwise the kernel will not see host‑side updates.
• The maximum size of a single SVM allocation is limited by CL_DEVICE_MAX_MEM_ALLOC_SIZE; query this value to avoid allocation failures.
• SVM memory is not automatically coherent; host writes must be followed by a map/unmap or a memory fence before the kernel starts, and kernel writes must be flushed before the host reads.
• Not all OpenCL 2.2 implementations expose fine‑grained SVM; some only support coarse‑grained, which still requires explicit mapping for host access.