CL_OUT_OF_RESOURCES in Command Queue during kernel enqueue

Direct Answer

When clEnqueueNDRangeKernel returns CL_OUT_OF_RESOURCES, the runtime is telling you that it cannot allocate the per‑work‑group resources required to launch the kernel. This is a launch‑time failure, not a signal that the command queue itself is full or that a global‑memory buffer failed to allocate. The error is generic: it could be due to register pressure, a fixed __local buffer that exceeds the device’s local memory, or insufficient scratch/constant cache for the chosen work‑group size. If the failure persists after the queue is quiesced, the device or driver has a hard limit that cannot be satisfied by any smaller work‑group size that still meets the kernel’s required work‑group size constraints.

Key Takeaways

• Global‑memory allocation failures are reported by clCreateBuffer/clCreateImage as CL_MEM_OBJECT_ALLOCATION_FAILURE, not CL_OUT_OF_RESOURCES.
• Local‑memory over‑allocation may produce CL_OUT_OF_RESOURCES, CL_INVALID_WORK_GROUP_SIZE, or even a build‑time CL_BUILD_PROGRAM_FAILURE depending on the ICD.
• An in‑order queue that blocks on a pending command will not return CL_OUT_OF_RESOURCES for queue pressure; the queue simply stalls until the pending command completes.

Practical Steps for Your Application

1. Quiesce the Queue: Call clFinish(queue) or wait on an event that depends on all prior enqueues. This guarantees that no pending commands are consuming resources.
2. Query Device Limits: Use clGetDeviceInfo for CL_DEVICE_LOCAL_MEM_SIZE and CL_DEVICE_MAX_WORK_GROUP_SIZE. After clSetKernelArg, query the kernel’s CL_KERNEL_LOCAL_MEM_SIZE and CL_KERNEL_WORK_GROUP_SIZE via clGetKernelWorkGroupInfo.
3. Check Work‑Group Feasibility: If the fixed __local buffer size exceeds CL_DEVICE_LOCAL_MEM_SIZE or the required work‑group size is larger than CL_DEVICE_MAX_WORK_GROUP_SIZE, shrinking the work‑group will not help. In that case, consider moving the buffer to global memory or restructuring the kernel.
4. Retry with a Smaller Work‑Group: If the current size is within device limits but the launch still fails, try a smaller work‑group that still satisfies the kernel’s minimum size. Re‑enqueue and observe whether the error clears.
5. Verify Buffer Visibility: After a successful launch, depend on the kernel event or call clFinish before performing any host‑side buffer read. Use blocking map/read or an event‑dependent copy to guarantee memory consistency.

When to Treat the Failure as Persistent

If the error persists after the queue is fully drained and all device limits are respected, the kernel’s resource demand exceeds what the device can provide. At that point, a hard failure is appropriate: either redesign the kernel, reduce its register usage, or target a device with larger local memory or register file.

Missing Diagnostic Detail

To refine the recommendation, please provide the size of the __local buffer you are allocating per work‑group (or the total local memory usage). Knowing whether the local memory requirement is fixed or work‑group‑size dependent will help determine if a simple work‑group size reduction will succeed.