Zero-copy host-device sharing in OpenCL with clSVMAlloc

Problem and useful takeaway

You have a heterogeneous CPU+GPU workload where the host and device repeatedly read and write the same working set. Copying through clCreateBuffer and clEnqueueReadBuffer adds latency and doubles memory traffic. OpenCL Shared Virtual Memory lets you allocate one virtual range with clSVMAlloc and let both host and device kernels access it directly. The practical decision is not whether SVM exists, but which SVM flags your devices actually support and whether you can tolerate explicit coherence.

Desired outcome

Zero-copy data sharing where the host pointer is usable by device kernels without an explicit buffer object, with coherence controlled by allocation flags and explicit synchronization. The host writes or reads the same address range the device uses, and you avoid per-launch map/unmap copies when fine-grain coherence is available.

Prerequisites

• OpenCL 2.0 or later runtime and driver on the host.
• A device that reports SVM support via CL_DEVICE_SVM_CAPABILITIES.
• Unified virtual addressing or coherent SVM on the target platform.

Check the runtime you will run against. Query the device before allocating. Do not assume fine-grain coherence on discrete GPUs; it is common on integrated CPU+GPU and some recent discrete parts, but not guaranteed.

Focused procedure

1. Query SVM capabilities

Run on the host process with read access to the OpenCL context. Create a context and command queue for the target device, then query capabilities.

cl_device_svm_capabilities caps;
clGetDeviceInfo(device, CL_DEVICE_SVM_CAPABILITIES, sizeof(caps), &caps, NULL);

Interpret the bitmask. You need at least CL_DEVICE_SVM_COARSE_GRAIN_BUFFER for basic SVM. CL_DEVICE_SVM_FINE_GRAIN_BUFFER indicates host-device coherence without explicit map. CL_DEVICE_SVM_ATOMICS indicates device atomics on SVM pointers are supported.

2. Choose allocation flags

Pick flags that match the capability and your kernel needs. A common safe choice for read‑write sharing is coarse‑grain with explicit map. If fine‑grain is present and you want zero‑copy, use CL_MEM_SVM_FINE_GRAIN_BUFFER.

Add CL_MEM_SVM_ATOMICS only if caps include atomics and your kernel uses atomic operations on SVM memory. Atomics can incur performance penalties on integrated devices.

3. Allocate SVM memory

Allocate on the host with clSVMAlloc. The returned pointer is host‑accessible and can be passed to kernels as a pointer argument.

size_t bytes = N * sizeof(float);
cl_mem_flags flags = CL_MEM_SVM_FINE_GRAIN_BUFFER; // or COARSE_GRAIN_BUFFER
void *ptr = clSVMAlloc(context, flags, bytes, 0);
if (!ptr) { /* fallback */ }

Risk: allocation can fail if the driver does not support the requested flags. Do not mix SVM pointers with classic cl_mem buffers in the same kernel on implementations with undefined behavior for mixed access.

4. Use the pointer in kernels

Declare the kernel argument as a pointer to SVM memory. The kernel can read and write through it directly.

__kernel void work(__global float *data) { /* ... */ }

Set the argument with clSetKernelArgSVMPointer. No clCreateBuffer is needed.

5. Synchronize for coherence

With coarse‑grain SVM you must map/unmap or use clEnqueueSVMMap/Unmap and memory fences before host access. With fine‑grain coherent SVM, host writes are visible to the device after appropriate OpenCL memory fences and command queue completion. Insert clFinish or event wait before host reads results.

Expected checks

• Device capability check returns a non-zero bitmask containing the flags you require.
• clSVMAlloc returns a non‑null pointer and the runtime reports no error.
• A small test kernel writes a known pattern to SVM memory. After queue completion, host reads the same address range and verifies the pattern matches the reference.
• Observe synchronization behavior under coarse vs fine grain modes. Coarse grain requires explicit map/unmap; fine grain should allow immediate host visibility after fence.

Check return values for clSVMAlloc and kernel launch. If allocation fails, trigger fallback logic rather than proceeding.

Recovery options

• If SVM is unsupported, fall back to traditional clCreateBuffer with explicit clEnqueueWriteBuffer and clEnqueueReadBuffer. This preserves correctness at the cost of copies.
• If fine‑grain is unsupported but coarse‑grain is available, allocate with CL_MEM_SVM_COARSE_GRAIN_BUFFER and use clEnqueueSVMMap/Unmap around host access. This keeps a single allocation but adds map overhead.
• If atomics are required but not supported, remove CL_MEM_SVM_ATOMICS and redesign the kernel to avoid SVM atomics, or fall back to buffer objects with device‑side atomics.

Limitations

SVM capability varies widely across vendors and generations. Fine‑grain coherent support is not guaranteed on discrete GPUs. Mixing SVM pointers with classic OpenCL buffers in the same kernel can cause undefined behavior on some implementations. Atomics on SVM require explicit device support and may incur performance penalties on integrated devices.