Does clFinish guarantee host-pointer coherence for CL_MEM_USE_HOST_PTR buffers on out-of-order queues?

Direct Answer

No. clFinish alone does not establish host-pointer coherence for a CL_MEM_USE_HOST_PTR buffer after a kernel write on an out-of-order queue, regardless of the CL_DEVICE_HOST_UNIFIED_MEMORY value. The OpenCL 1.2 memory model requires an explicit device-to-host synchronization command—clEnqueueMapBuffer, clEnqueueReadBuffer, or clEnqueueReadBufferRect—ordered after the kernel via an event wait list or barrier. clFinish only guarantees that previously enqueued commands have completed on the device; it does not by itself perform the buffer synchronization that makes the host pointer safe to read.

Required Synchronization Sequence

1. Capture the event K returned by the kernel enqueue (clEnqueueNDRangeKernel) on the out-of-order queue.
2. Enqueue a map or read command for the same buffer with K in its event_wait_list:

   cl_event map_evt;
   clEnqueueMapBuffer(queue, buffer, CL_FALSE, CL_MAP_READ, 0, size, 1, &K, &map_evt, &err);

3. Wait for the map/read to complete: clWaitForEvent(map_evt) or clFinish(queue) after the map/read enqueue.
4. Access the host pointer (or the mapped pointer returned by clEnqueueMapBuffer).
5. Before reusing the buffer in device commands, enqueue clEnqueueUnmapMemObject (if mapped) or ensure the read command has completed.

Step 2 is essential: on an out-of-order queue, program order does not order the host-facing command after the kernel. The event wait list (or a clEnqueueBarrierWithWaitList) creates the required happens-before edge.

Effect of CL_DEVICE_HOST_UNIFIED_MEMORY

The required sequence does not change. When CL_DEVICE_HOST_UNIFIED_MEMORY == CL_TRUE, the device and host share physical memory, so some drivers appear to make the host pointer visible immediately after clFinish. This is implementation-specific behavior, not a portable guarantee. A driver update, a different device on the same platform, or a discrete GPU (CL_FALSE) will break code that relies on clFinish alone. Treat CL_DEVICE_HOST_UNIFIED_MEMORY as a performance hint only.

Portability Hazard

Code that reads host_ptr after only clFinish works on some integrated GPUs and CPUs today but will silently corrupt data on discrete GPUs or after driver changes. The OpenCL 1.2 specification (section 5.2.1, 5.3) is explicit: the host must not access the memory region while a command using the buffer is in flight, and coherence is established only by the map/read/write commands themselves.

Verification Checklist

• Query CL_DEVICE_HOST_UNIFIED_MEMORY once at startup; log it but do not branch synchronization logic on it.
• On each target device/driver, compare a clFinish-only host read against the event-ordered map/read sequence using a known kernel write pattern.
• Confirm completion with clGetEventInfo(event, CL_EVENT_COMMAND_EXECUTION_STATUS, ...) returning CL_COMPLETE before dereferencing the host pointer.