OpenCL latency spikes during concurrent out-of-order queue submissions

Question

When utilizing multiple OpenCL command queues configured with CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, a measurable latency gap often emerges between clEnqueueNDRangeKernel and the CL_PROFILING_COMMAND_START timestamp. This delay appears to manifest specifically when multiple queues submit kernels simultaneously, whereas single-queue operations maintain consistent start times.

The OpenCL specification does not define a maximum bound for kernel launch scheduling, leaving the behavior to the driver implementation. It is unclear if this latency is a necessary result of the runtime serializing requests across a shared device context or an overhead from managing event-based synchronization under high contention.

How do different OpenCL runtimes prioritize kernel scheduling when multiple out-of-order queues target the same device? Is there a documented mechanism to minimize the enqueue-to-start delay in concurrent-request environments without switching to blocking in-order queues?

Tasadduq Burney · Answer

Direct answer
Confirmed fact: The OpenCL specification does not define a maximum bound for the time between clEnqueueNDRangeKernel and CL_PROFILING_COMMAND_START, and it does not mandate a scheduling policy for multiple CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE queues targeting the same device. Profiling timestamps CL_PROFILING_COMMAND_QUEUED and CL_PROFILING_COMMAND_START are defined and reliable for measurement.
Likely explanation: Most runtimes serialize kernel launch handling for a shared device context. With several out-of-order queues submitting at once, driver bookkeeping, host thread contention and device resource arbitration can create a measurable enqueue-to-start gap that is not present with a single queue.
How runtimes typically prioritize

Global FIFO or internal batching across all queues to the same device.
Per-queue priority heuristic where supported, e.g., CL_QUEUE_PRIORITY or vendor extensions.
Strict respect for event dependencies even in out-of-order mode; wait lists force ordering.
Throttling when concurrent kernel limits or compute units are exceeded.

Steps to reduce enqueue-to-start delay without in-order queues

Enable profiling on each queue and record QUEUED vs START for every kernel. Compare 1 queue vs N queues to confirm contention.
Check for priority support. Query clGetDeviceInfo for extensions and test CL_QUEUE_PRIORITY creation. If present, place latency-sensitive kernels on a higher priority queue.
Use explicit event dependencies across queues instead of relying on implicit ordering. Chain with clEnqueueMarkerWithWaitList or pass events in the wait list of clEnqueueNDRangeKernel.
Limit active concurrent queues to the device’s practical concurrency. Batch work into fewer queues if the spike correlates with queue count.

Missing diagnostic detail that changes the recommendation: Does your runtime and device expose a queue priority extension, such as CL_QUEUE_PRIORITY or a vendor-specific equivalent, and is it enabled in your build? Priority hints are the only documented way to influence scheduling without adding dependencies.

OpenCL latency spikes during concurrent out-of-order queue submissions

1 answer

Direct answer

How runtimes typically prioritize

Steps to reduce enqueue-to-start delay without in-order queues

Your perspective belongs here.

0 question comments