Latency bounds for concurrent host-thread submissions to a shared OpenCL out-of-order queue
0 reputation · 28 Jun 2025, 06:55 UTC
0 reputation · 28 Jun 2025, 06:55 UTC
A host application uses a single OpenCL context with an out-of-order command queue. Multiple host threads call clEnqueueNDRangeKernel concurrently, each with its own event, and later wait via clWaitForEvents. Under single-threaded submission the enqueue-to-completion interval is stable, but the design goal is to keep per-command latency predictable as the number of submitting threads grows.
The OpenCL specification permits concurrent submissions from multiple host threads to a shared context, but it does not appear to define an upper bound on the latency a runtime may add when serializing internal queue access or inserting implicit synchronization points. clEnqueueMarker and clEnqueueBarrier create ordering points, yet their latency relative to concurrently submitted commands is also unspecified. Vendor behavior is documented as implementation-defined, so any bound observed on one driver (e.g., Intel Compute Runtime vs. an NVIDIA OpenCL implementation) may not transfer.
No OpenCL version (1.x through 3.0) or widely supported extension provides a guaranteed upper bound on the latency a runtime may add when multiple host threads submit concurrently to a shared queue. The specification guarantees that API calls are thread-safe; it says nothing about their timing, fairness, or how internal serialization scales with thread count. So: the portable strategy is to measure with event profiling, not to rely on a documented bound — and per-thread queues remove your contention on the shared queue object but carry no stronger guarantee.
There is no bounding mechanism, but there is a portable measuring mechanism: event profiling. Create the queue with CL_QUEUE_PROFILING_ENABLE and read timestamps via clGetEventProfilingInfo:
CL_PROFILING_COMMAND_QUEUED — when the host enqueued the commandCL_PROFILING_COMMAND_SUBMIT — when the runtime handed it to the deviceCL_PROFILING_COMMAND_START / END — device executionThe QUEUED → SUBMIT delta is exactly where host-side serialization and runtime-internal locking show up. If that delta grows as you add submitting threads while START → END stays flat, the contention is in submission, not execution. Caveats: timestamp resolution and profiling overhead are implementation-defined, and enabling profiling can itself perturb timing on some drivers.
No. Internal locking, implicit synchronization, and queue management are explicitly implementation-defined. There is no queue creation flag or extension that promises submission cost independent of thread count. clEnqueueMarkerWithWaitList and clEnqueueBarrierWithWaitList define ordering only — their completion latency relative to concurrent submissions is unspecified, and barriers can force the runtime to drain in-flight work, which worsens tail latency under load.
The defensible portable position is: accept that behavior is implementation-defined, set an empirical latency budget per target driver, and re-validate after any runtime update. A bound observed on Intel Compute Runtime tells you nothing reliable about NVIDIA's or AMD's implementation.
The specification treats all queues identically, so per-thread queues come with no additional latency guarantee. What they do change is the practical picture: if the runtime serializes on the queue object, giving each thread its own queue removes that specific contention point. If the serialization is deeper (per-context or per-device locking inside the driver), per-thread queues won't help. This is directly testable: run your N-thread workload against one shared queue, then against N queues, and compare the QUEUED → SUBMIT deltas.
CL_QUEUE_PROFILING_ENABLE.QUEUED, SUBMIT, START, END per event; plot SUBMIT − QUEUED against N.clFlush once per batch instead of relying on implicit flushing.CL_QUEUE_DEVICE capability), moving submission off host threads sidesteps the problem entirely — support is uneven, so query the device first.One note on scope: an out-of-order queue relaxes execution ordering only; it does not relax host-side submission serialization, which is where your observed contention most likely lives. All scaling claims here should be verified on your specific driver versions, since runtime behavior can change between releases.
Use comments to ask for clarification. Post a solution as an answer.
No question comments on this page.