OpenCL fine-grained SVM buffer mapping coherence boundary

Coherence guarantees you can rely on

OpenCL 2.0 defines fine‑grained SVM as memory that can be accessed by both host and device without explicit copies. The specification states:

• clEnqueueSVMMap establishes a host‑side synchronization point.
• clEnqueueSVMUnmap establishes a device‑side synchronization point.
• There is no global memory‑consistency model; each device’s cache hierarchy is independent.

In practice this means that:

• A host write followed by a kernel read is *not* automatically visible – you must enforce ordering.
• Similarly, a kernel write followed by a host read needs explicit device‑to‑host synchronization.

Is map/unmap alone enough?

Map/unmap are *permission boundaries*, but they do not guarantee that data has propagated through device caches. The spec does not require the implementation to flush or invalidate caches automatically. Therefore:

• After a kernel writes to an SVM buffer you must issue a device‑side barrier or a marker that waits for the kernel to finish before the host reads.
• After a host writes and before the device reads you must wait for the map operation to complete (usually implicit in an in‑order queue) or explicitly use an event dependency.

Deterministic visibility when SVM is optional (OpenCL 3.0)

OpenCL 3.0 makes fine‑grained SVM optional. Portable code should:

1. Query CL_DEVICE_SVM_FINE_GRAIN_BUFFER with clGetDeviceInfo.
2. If the flag is set, use the fine‑grained pattern described above.
3. Otherwise fall back to a classic buffer + clEnqueueCopyBuffer path.

Regardless of the path, you need to enforce ordering with the same primitives: clEnqueueBarrierWithWaitList, clEnqueueMarkerWithWaitList, or event dependencies. Using clFinish after each kernel is the simplest, though it can be expensive.

Portable pattern for a single host‑to‑device write, kernel compute, device‑to‑host read

cl_event map_evt, kernel_evt, unmap_evt;

// 1. Host writes to SVM
clEnqueueSVMMap(queue, svm_ptr, 0, size, CL_MAP_WRITE, 0, NULL, &map_evt);
// ... host writes data here ...
clEnqueueSVMUnmap(queue, svm_ptr, 0, NULL, &unmap_evt);

// 2. Kernel that reads and writes the same SVM buffer
clSetKernelArgSVMPointer(kernel, 0, svm_ptr);
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 1, &unmap_evt, &kernel_evt);

// 3. Device‑to‑host visibility
clEnqueueBarrierWithWaitList(queue, 1, &kernel_evt, NULL);

// 4. Host reads back
clEnqueueSVMMap(queue, svm_ptr, 0, size, CL_MAP_READ, 0, NULL, &map_evt);
// ... host reads data ...
clEnqueueSVMUnmap(queue, svm_ptr, 0, NULL, NULL);
clFinish(queue);

Key points:

• Each clEnqueueSVMMap or clEnqueueSVMUnmap is paired with an event that you can use to build a dependency chain.
• Explicit barriers or markers are required after kernel execution before the host can safely read the buffer.
• In an out‑of‑order queue you must attach the events; otherwise the implicit ordering of an in‑order queue already guarantees the correct sequence.

What to ask for diagnosis

If you are still seeing stale data, confirm whether you are using an out‑of‑order command queue. In that case you must explicitly chain the events shown above; otherwise the implicit ordering of an in‑order queue is sufficient.