OpenCL 2.0 Fine‑Grained SVM Guide

Learn how to allocate, use, and verify fine‑grained SVM buffers in OpenCL 2.0 to share pointers between host and kernel without explicit map/unmap calls.

Desired Outcome

Allocate a fine‑grained SVM buffer, write a test pattern from the host, launch a kernel that increments each element, and read the updated values directly from the host pointer without using clEnqueueMapBuffer or clEnqueueSVMMap.

Prerequisites

OpenCL 2.0‑capable GPU or CPU device and driver.
Development environment with OpenCL headers and linkable library (e.g., libOpenCL).
Basic knowledge of C/C++ and OpenCL API concepts (context, command queue, program, kernel).

Procedure

Query SVM capabilities

Call clGetDeviceInfo for CL_DEVICE_SVM_CAPABILITIES and check that the CL_DEVICE_SVM_FINE_GRAIN_BUFFER bit is set. Also retrieve CL_DEVICE_SVM_ALIGNMENT to know the required alignment.

cl_uint svmCaps;
clGetDeviceInfo(device, CL_DEVICE_SVM_CAPABILITIES, sizeof(svmCaps), &svmCaps, NULL);
if (!(svmCaps & CL_DEVICE_SVM_FINE_GRAIN_BUFFER)) {
    /* fine‑grained SVM not supported */
}
size_t alignment;
clGetDeviceInfo(device, CL_DEVICE_SVM_ALIGNMENT, sizeof(alignment), &alignment, NULL);

Create an SVM‑capable context

When creating the context, pass the CL_CONTEXT_SVM_CAPABLE flag (value CL_TRUE) in the properties list.

cl_context_properties props[] = {
    CL_CONTEXT_PLATFORM, (cl_context_properties)platform,
    CL_CONTEXT_SVM_CAPABLE, CL_TRUE,
    0
};
cl_context context = clCreateContextFromType(props, CL_DEVICE_TYPE_GPU, NULL, NULL, &err);

Allocate fine‑grained SVM memory

Use clSVMAlloc with the fine‑grained flag and the alignment obtained earlier.

void *svmPtr = clSVMAlloc(context,
                          CL_MEM_SVM_FINE_GRAIN_BUFFER,
                          numElements * sizeof(cl_int),
                          alignment);
if (!svmPtr) {
    /* allocation failed */
}

Initialize data on the host

Because the allocation is fine‑grained, the host can write directly to the returned pointer.

cl_int *data = (cl_int *)svmPtr;
for (size_t i = 0; i < numElements; ++i) {
    data[i] = (cl_int)i;   /* pattern 0,1,2,... */
}

Build and create the kernel

Compile the program with -cl-std=CL2.0 to enable SVM‑related language features. A simple kernel that increments each element:
```
__kernel void inc(__global int *ptr, unsigned int count) {
    unsigned int gid = get_global_id(0);
    if (gid < count) {
        ptr[gid] = ptr[gid] + 1;
    }
}
```
Create the program from source, build it, and create the kernel.

Set kernel arguments and enqueue

Pass the SVM pointer as the first argument; no extra mapping is needed.

clSetKernelArg(kernel, 0, sizeof(void *), &svmPtr);
clSetKernelArg(kernel, 1, sizeof(unsigned int), &numElements);
size_t global = numElements;
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, NULL, NULL);

Synchronize and verify

For fine‑grained SVM the host sees changes automatically after the kernel finishes; nevertheless, insert a blocking wait or a flush/finish to ensure completion.
```
clFinish(queue);   /* or clWaitForEvents with an event from the enqueue */
```
Now read back the values directly from svmPtr and compare to the expected pattern (each element incremented by one).
```
bool ok = true;
for (size_t i = 0; i < numElements; ++i) {
    if (data[i] != (cl_int)(i + 1)) {
        ok = false;
        break;
    }
}
```

Cleanup

Free the SVM allocation and release OpenCL objects.

clSVMFree(context, svmPtr);
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseCommandQueue(queue);
clReleaseContext(context);

Expected Checks

Every OpenCL call returns CL_SUCCESS (check the error code variable).
CL_DEVICE_SVM_CAPABILITIES contains CL_DEVICE_SVM_FINE_GRAIN_BUFFER.
The returned pointer from clSVMAlloc is non‑NULL and properly aligned (you can verify with reinterpret_cast<uintptr_t>(svmPtr) % alignment == 0).
After clFinish, the host‑side values match the expected incremented pattern.

Recovery Options and Limitations

Missing fine‑grained support: If the capability bit is not present, fall back to coarse‑grained SVM (CL_MEM_SVM_COARSE_GRAIN_BUFFER) which requires explicit clEnqueueSVMMap/clEnqueueSVMUnmap for host access, or to a regular buffer with clCreateBuffer and clEnqueueMapBuffer/clEnqueueUnmapMemObject.
Alignment faults: Mis‑aligned pointers can cause undefined behavior. Always query CL_DEVICE_SVM_ALIGNMENT and use it as the alignment argument to clSVMAlloc. If allocation fails with CL_OUT_OF_RESOURCES, try reducing the size or increasing alignment.
Mixing SVM and regular buffers: When a kernel accesses both SVM pointers and regular buffer objects, ensure proper synchronization because the consistency models differ. Use clEnqueueSVMMap/clEnqueueSVMUnmap or clEnqueueMapBuffer as needed, or finish the queue before switching access.
Kernel build errors: Inspect the build log with clGetProgramBuildInfo and CL_PROGRAM_BUILD_LOG. Adjust the source or compile options (e.g., ensure -cl-std=CL2.0 is present).
Rolling back state: The only state‑changing step in this guide is the SVM allocation. If allocation succeeds but later steps fail, call clSVMFree to release the memory before exiting or retrying.

Practical Verification

To confirm that the SVM buffer is truly fine‑grained, you can run a second test where the host modifies the data while the kernel is executing (using separate threads or overlapped command queues) and observe that the kernel sees the updates without explicit mapping. This is optional but demonstrates the implicit consistency model.

Using OpenCL 2.0 Fine‑Grained Shared Virtual Memory for Host‑Kernel Data Sharing