Using OpenCL 2.0 Shared Virtual Memory for Unified Host‑Device Access

When to choose SVM

If you want to eliminate explicit clEnqueueWriteBuffer and clEnqueueReadBuffer calls, OpenCL 2.0 Shared Virtual Memory (SVM) lets the host and device share a single pointer. Allocate the buffer once with clSVMAlloc, pass that pointer to kernels, and read or write it directly from the host.

Mechanism: coarse‑grained vs fine‑grained SVM

Two SVM modes exist:

• Coarse‑grained – the host and device see the same memory without any explicit synchronization. No clEnqueueSVMMap/clUnmap is needed, but the allocation must be made per device and the memory must not be accessed concurrently without user‑level synchronization.
• Fine‑grained – requires explicit map/unmap (or clEnqueueSVMMemcpy) before the kernel can safely use the buffer. It imposes 64‑byte alignment and gives the driver a chance to handle page faults, which can be useful for large or streaming data.

Both modes require the device to report SVM capability via CL_DEVICE_SVM_CAPABILITIES. If the flag is missing, clSVMAlloc returns CL_INVALID_OPERATION.

Worked example: coarse‑grained SVM

The following minimal program checks for coarse‑grained SVM, allocates a buffer, launches a kernel that increments each element, and reads the result back without explicit copy commands.

#include 
#include 
#include 

#define NUM_ELEMENTS 1024

const char *kernel_src =
"__kernel void inc(__global int *ptr) {\n"
"    size_t i = get_global_id(0);\n"
"    ptr[i] = ptr[i] + 1;\n"
"}\n";

int main(void) {
    cl_int err;
    cl_platform_id platform;
    cl_device_id device;
    cl_context context;
    cl_command_queue queue;
    cl_program program;
    cl_kernel kernel;
    
    /* 1. Find a device that supports coarse‑grained SVM */
    err = clGetPlatformIDs(1, &platform, NULL);
    err |= clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
    cl_device_svm_capabilities caps;
    err = clGetDeviceInfo(device, CL_DEVICE_SVM_CAPABILITIES, sizeof(caps), &caps, NULL);
    if (!(caps & CL_DEVICE_SVM_COARSE_GRAIN_BUFFER)) {
        fprintf(stderr, "Device lacks coarse‑grained SVM support\n");
        return EXIT_FAILURE;
    }
    
    /* 2. Create context and command queue */
    context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
    queue = clCreateCommandQueueWithProperties(context, device, 0, &err);
    
    /* 3. Build kernel from source */
    program = clCreateProgramWithSource(context, 1, &kernel_src, NULL, &err);
    err = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
    kernel = clCreateKernel(program, "inc", &err);
    
    /* 4. Allocate SVM buffer (coarse‑grained, read/write) */
    int *svm_ptr = (int *)clSVMAlloc(context,
                                     CL_MEM_READ_WRITE,
                                     NUM_ELEMENTS * sizeof(int),
                                     0);  /* 0 = default alignment for coarse‑grained */
    if (!svm_ptr) {
        fprintf(stderr, "clSVMAlloc failed\n");
        return EXIT_FAILURE;
    }
    
    /* 5. Initialize data on host (visible to device immediately) */
    for (size_t i = 0; i < NUM_ELEMENTS; ++i) svm_ptr[i] = (int)i;
    
    /* 6. Set kernel argument and enqueue */
    err = clSetKernelArg(kernel, 0, sizeof(void *), &svm_ptr);
    size_t global = NUM_ELEMENTS;
    err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, NULL, NULL);
    
    /* 7. Wait for kernel to finish */
    clFinish(queue);
    
    /* 8. Verify results on host – no explicit read needed */
    int ok = 1;
    for (size_t i = 0; i < NUM_ELEMENTS; ++i) {
        if (svm_ptr[i] != (int)i + 1) { ok = 0; break; }
    }
    printf("Verification %s\n", ok ? "PASSED" : "FAILED");
    
    /* 9. Cleanup */
    clSVMFree(context, svm_ptr);
    clReleaseKernel(kernel);
    clReleaseProgram(program);
    clReleaseCommandQueue(queue);
    clReleaseContext(context);
    return ok ? EXIT_SUCCESS : EXIT_FAILURE;
}

Limits and practical considerations

• Capability check – Always query CL_DEVICE_SVM_CAPABILITIES before calling clSVMAlloc. Skipping this leads to CL_INVALID_OPERATION or, on some drivers, a silent fallback to regular buffers.
• Alignment – Fine‑grained SVM requires 64‑byte alignment; coarse‑grained has looser requirements but still benefits from page‑size alignment for performance.
• Size limits – The total SVM allocation cannot exceed the device’s global memory size. Very large allocations may fail with CL_MEM_OBJECT_ALLOCATION_FAILURE.
• Synchronization – Coarse‑grained SVM does not provide automatic coherence with concurrent command queues or multiple kernels. Use events, barriers, or explicit clEnqueueSVMMemcpy with proper ordering to avoid race conditions.
• Performance trade‑off – For small buffers, coarse‑grained SVM often beats explicit clEnqueueWriteBuffer because it avoids copy overhead. For large, streaming workloads, fine‑grained SVM (with map/unmap) can overlap computation with data movement, but the extra synchronization steps may offset gains.

How to verify SVM works on your hardware

1. Run clGetDeviceInfo for CL_DEVICE_SVM_CAPABILITIES and confirm the presence of either CL_DEVICE_SVM_COARSE_GRAIN_BUFFER or CL_DEVICE_SVM_FINE_GRAIN_BUFFER.
2. Compile and execute the example program above. It should print "Verification PASSED" and return exit code 0.
3. Optionally, replace the kernel with a simple copy and time clEnqueueSVMMemcpy versus clEnqueueWriteBuffer using clGetEventProfilingInfo to see the performance difference on your device.

Common mistakes to avoid

• Assuming all OpenCL 2.0 devices support SVM – many older Intel integrated GPUs or early AMD GCN cards return no SVM capability.
• Calling a kernel that writes to SVM memory without ensuring prior completion of a host write (or vice‑versa) when using fine‑grained SVM without map/unmap.
• Ignoring the return value of clSVMAlloc and proceeding with a null pointer, leading to undefined behavior.
• Allocating SVM from a context that includes multiple devices with differing SVM capabilities; the allocation succeeds only on the least‑capable device, causing failures on others.