OpenCL enqueue barrier vs enqueue marker - what's the difference?

The OpenCL specification has the API functions:

_{(I'm looking at v2.2 but they exist in v3.0 as well.)}

Not only do they have an identical signature, they also have very similar descriptions. What is the difference between these two?

Solution

My answer assumes the common case where OpenCL command queues are in-order queues. Out of order queues are less common and require a more fine grained synchronization.

A simple way to see the difference between a marker and a barrier is by looking at the most typical use case scenario of each one of them. From a point of view of an asynchronous program:

A marker is typically used by a producer.
A barrier is typically used by a consumer.

So you would typically us them in scenarios where commands on one queue can't begin execution before a command on a different queue has finished.

Example: Assume you have 2 command queues: queue1 and queue2. And two associated kernels: kernel1 and kernel2 respectively. kernel1 writes to a buffer (cl_mem) that kernel2 reads.

A sequence of command recording (on a single thread) could look like that:

queue1.enqueue(kernel1) 
event1 = queue1.enqueue_marker() // event 1 will signal only after kernel 1 is done
queue2.enqueue_barrier(wait_list = event1) // kernel 2 will not start before event1 is signaled
queue2.enqueue(kernel2)

The reason you rarely see markers and barriers used in real code is that each clEnqueue*** API call implicitly act as a marker and a barrier. You typically use barriers and markers if you don't have access to the raw clEnqueue*** calls.

For example:

cl_int clEnqueueNDRangeKernel ( cl_command_queue command_queue,
    cl_kernel kernel,
    cl_uint work_dim,
    const size_t *global_work_offset,
    const size_t *global_work_size,
    const size_t *local_work_size,
    cl_uint num_events_in_wait_list,
    const cl_event *event_wait_list, // barrier
    cl_event *event) // marker

The returned event marks the end of the enqueued command (and also all previous commands, in the case of in-order queue). You may provide the returned event to a consumer command (clEnqueue***) on a different queue (as wait list) to be used a a barrier.