OpenCL 2.0 Coarse-Grained SVM: Share Buffers Without Explicit Copies

The problem and the practical takeaway

OpenCL 1.x programs often move data with clEnqueueMemcpyBuffer. That pattern works, but it duplicates memory and adds CPU-driven copy overhead on every transfer. If your device supports OpenCL 2.0 coarse-grained shared virtual memory (SVM), you can allocate one buffer whose address is valid on both host and device, pass that pointer directly to a kernel, and avoid explicit copy calls for that buffer.

The takeaway is conditional: zero-copy SVM is only safe after you confirm the device reports CL_DEVICE_SVM_COARSE_GRAIN_BUFFER, and only if the host respects strict map and unmap boundaries around kernel execution.

What coarse-grained SVM means

Shared virtual memory (SVM) lets host and device reference the same virtual addresses. Coarse-grained SVM means the host must explicitly map a region before the device uses it and unmap it before the host touches it again. Fine-grained SVM, by contrast, allows device-visible host writes without explicit mapping, but it requires atomics and host-side locking and is a separate capability.

Requirements before you allocate

• Host and OpenCL device must implement OpenCL 2.0 or later.
• The device must report CL_DEVICE_SVM_COARSE_GRAIN_BUFFER in its CL_DEVICE_SVM_CAPABILITIES query. OpenCL 2.0 alone does not guarantee SVM.
• The OpenCL ICD loader and driver must support SVM on the target device; driver advertisements can be present without a working implementation.

Smallest suitable design

Use a single coarse-grained SVM allocation for the shared region. Do not wrap it in a cl_mem buffer object, and do not call clEnqueueMemcpyBuffer on it.

1. Create an OpenCL context and command queue for a device that passed the SVM capability query.
2. Call clSVMAlloc to allocate the shared region.
3. Call clEnqueueSVMMap to make the region accessible to the host.
4. Pass the SVM pointer directly as a kernel argument.
5. Enqueue the kernel.
6. Call clEnqueueSVMUnmap before the host reads or writes the region.
7. Wait on the unmap event before host access.
8. Free the allocation with clSVMFree when done.

Example host flow

/* Run on the host build machine. Link against the OpenCL ICD loader.
   No special OS permissions are required beyond normal user access to the GPU driver. */

cl_int err;
cl_svm_mem_flags flags = CL_MEM_READ_WRITE;
size_t size = 4096;
void *svm_ptr = clSVMAlloc(context, flags, size, 0);
if (svm_ptr == NULL) {
    /* Rollback: fall back to clCreateBuffer plus clEnqueueMemcpyBuffer. */
}

cl_event map_event;
err = clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_WRITE_INVALIDATE_REGION,
                      svm_ptr, size, 0, NULL, &map_event);
if (err != CL_SUCCESS) { /* mapping failure; do not use svm_ptr */ }

/* Initialize host-side data here if needed while mapped. */

err = clSetKernelArg(kernel, 0, sizeof(void *), &svm_ptr);
if (err != CL_SUCCESS) { /* argument error */ }

cl_event kernel_event;
err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size, &local_size,
                             0, NULL, &kernel_event);

cl_event unmap_event;
err = clEnqueueSVMUnmap(queue, svm_ptr, 0, NULL, &unmap_event);
if (err != CL_SUCCESS) { /* unmapping failure */ }

/* Wait for unmap_event before host reads svm_ptr. Accessing svm_ptr before
   unmap completes is undefined behavior and can corrupt data. */

clSVMFree(context, svm_ptr);

Trust and data boundaries

Because host and device share the same virtual address space for the SVM region, a buggy or malicious kernel can read or write host memory inside the mapped range. Treat the SVM mapping as a trust boundary:

• Isolate untrusted kernels in separate contexts or processes where possible.
• Validate kernel arguments before enqueueing.
• Never store an SVM pointer in a regular cl_mem object or pass it to an OpenCL 1.x device.
• SVM pointers are invalid outside the context that created them.

Operational checks

• Query CL_DEVICE_SVM_CAPABILITIES at context creation and before any clSVMAlloc call.
• Check clSVMAlloc for a NULL return.
• Treat negative return codes from clEnqueueSVMMap and clEnqueueSVMUnmap as mapping failures.
• Wait on the map and unmap events; do not assume completion because the call returned.
• Record any target device where the capability bit is absent or allocation returns NULL.

Failure modes

• The driver advertises OpenCL 2.0 but SVM does not work; capability queries are mandatory.
• Large SVM regions can exhaust contiguous virtual address space on the device.
• Concurrent host writes during kernel execution cause data races.
• Host access to an SVM buffer without unmapping after kernel enqueue is undefined behavior.
• SVM pointers cannot be shared across contexts without per-context allocations or vendor extensions.

Conditions that change the design

Condition	Design implication
Fine-grained SVM with atomics is required	Add host-side locking and confirm device fine-grained SVM support.
Frequent small transfers	Prefer CL_MEM_USE_HOST_PTR buffers with explicit copies instead of SVM.
Multiple devices must share the same data	Use per-context SVM allocations or vendor extensions; a single SVM pointer is context-bound.
Driver lacks working coarse-grained SVM	Roll back to clCreateBuffer plus clEnqueueMemcpyBuffer.

Minimal verification you should run

Build a minimal OpenCL host program that queries clGetDeviceInfo for CL_DEVICE_SVM_CAPABILITIES, allocates with clSVMAlloc, maps, runs a kernel that writes a sentinel value, unmaps, and asserts that the host observes the sentinel. Repeat on every target device and driver version. Inspect clinfo or vendor OpenCL version output on the exact production hardware before committing to the design.

Limitations and practical check

Coarse-grained SVM is optional even on OpenCL 2.0 devices. Do not assume zero-copy behavior from the OpenCL version string alone. The practical check is to run the capability query and a sentinel-write test on each production device; if clSVMAlloc returns NULL or the sentinel is not observed after unmap, fall back to explicit buffer copies.