Async Command Queue Deep-Dive

Design Overview

The async queue bridges the OpenCL runtime layer and the Core::Queue execution engine. It uses four static map tracking tables to keep OpenCL objects alive while their associated commands are in flight.

graph TD
  CL[cl* API call] -->|creates Command| AE[asyncEnqueue]
  AE -->|enqueue| CQ[Core::Queue]
  AE -->|creates| EV[_cl_event]
  AE -->|stores in| EM[eventMap]
  CQ -->|on complete| AR[asyncQueueRelease]
  AR -->|fires| CB[Event callbacks]
  AR -->|releases| MO[Retained mem objects]
  AR -->|releases| KR[Retained kernels]

Tracking Maps

Four module-scoped maps index by Command*:

• memObjectMap — retained cl_mem objects per command
• kernelMap — the retained cl_kernel for kernel commands
• eventMap — the cl_event created for the command
• waitListMap — retained wait-list events the command depends on

asyncEnqueue — Creating Events

When the runtime enqueues any command (buffer copy, kernel launch, etc.), asyncEnqueue() performs three tasks:

void asyncEnqueue(cl_command_queue queue, cl_command_type type,
                  Command* cmd, cl_uint numEvents,
                  const cl_event* waitList, cl_event* eventOut) {
  // 1. Retain & record each wait-list event
  for (unsigned i = 0; i < numEvents; i++) {
    cmd->waitList.push_back(waitList[i]->event);
    waitListMap[cmd].push_back(waitList[i]);
    clRetainEvent(waitList[i]);
  }

  // 2. Forward to the Core queue — returns an Event*
  Event* event = queue->queue->enqueue(cmd);

  // 3. Wrap in a cl_event and store in the map
  cl_event _event     = new _cl_event;
  _event->dispatch    = m_dispatchTable;
  _event->context     = queue->context;
  _event->queue       = queue;
  _event->type        = type;
  _event->event       = event;
  _event->refCount    = 1;
  eventMap[cmd]       = _event;

  if (eventOut) { clRetainEvent(_event); *eventOut = _event; }
}

asyncQueueRetain — Object Lifetime

Two overloads exist: one for cl_mem and one for cl_kernel. The kernel overload also retains every cl_mem referenced via kernel->memArgs.

void asyncQueueRetain(Command* cmd, cl_kernel kernel) {
  clRetainKernel(kernel);
  kernelMap[cmd] = kernel;

  // Also retain all memory objects bound as kernel arguments
  for (auto itr = kernel->memArgs.begin();
       itr != kernel->memArgs.end(); itr++) {
    asyncQueueRetain(cmd, itr->second);
  }
}

Release & Callback Flow

sequenceDiagram
  participant CQ as Core::Queue
  participant AQR as asyncQueueRelease
  participant CB as User Callbacks
  CQ->>AQR: command completed
  AQR->>AQR: release memObjectMap entries
  AQR->>AQR: release kernelMap + delete clone
  AQR->>CB: fire event callbacks
  AQR->>AQR: release waitListMap events
  AQR->>AQR: clReleaseEvent(self)

asyncQueueRelease is called from the releaseCommand() helper in runtime.cpp. It walks all four maps, fires the user-registered CL_CALLBACK list, then releases every retained object.

Command Cascading

Each Command struct has an execBefore list. When releaseCommand() runs, it recurses through that list, releasing dependent commands first. This enables implicit barriers and ordered-queue semantics.

Event Lifecycle Detail

An event transitions through several states during its lifetime. The Core::Event tracks execution state internally, while the _cl_event wrapper holds the user-visible reference count and callbacks.

// releaseCommand helper from runtime.cpp
void releaseCommand(oclgrind::Command* command) {
  if (command) {
    asyncQueueRelease(command);

    // Release dependent commands recursively
    while (!command->execBefore.empty()) {
      oclgrind::Command* cmd = command->execBefore.front();
      command->execBefore.pop_front();
      releaseCommand(cmd);
    }
    delete command;
  }
}

Memory Retention Example

When a buffer read/write is enqueued, the runtime retains the cl_mem argument so the buffer cannot be freed while the transfer is in flight:

// Inside clEnqueueReadBuffer (simplified)
Command* cmd    = new Command();
cmd->type       = Command::READ_BUFFER;
cmd->readBuffer = {mem->address + offset, ptr, size};
asyncQueueRetain(cmd, mem);     // keeps cl_mem alive
asyncEnqueue(queue, CL_COMMAND_READ_BUFFER,
             cmd, numEvents, waitList, eventOut);

Callback Ordering

User callbacks registered via clSetEventCallback are stored in a std::list on the _cl_event. They are fired in registration order during asyncQueueRelease, after all retained objects have been released. This matches the OpenCL specification requirement that callbacks observe the completion state.

Thread Safety Considerations

The four tracking maps are module-scoped static variables with no mutex protection. This is safe because Oclgrind processes all commands on a single thread — the Core::Queue executes commands sequentially. Applications using multiple host threads must externally synchronise their enqueue calls.

Kernel Command Special Handling

For kernel enqueues, the runtime clones the Kernel object to capture the argument bindings at enqueue time. The clone is stored in a KernelCommand sub-struct. On release, the async queue:

1. Releases the cloned kernel (delete ((KernelCommand*)cmd)->kernel)
2. Releases the user-facing cl_kernel via clReleaseKernel
3. Releases all cl_mem arguments previously retained

This ensures the user can safely modify kernel arguments or release the kernel object immediately after clEnqueueNDRangeKernel.

Wait-List Event Dependencies

Each retained wait-list event is pushed into cmd->waitList (for the Core queue) and into waitListMap[cmd] (for later release). This dual tracking ensures the Core queue can check completion status, while the async queue layer can balance the clRetainEvent calls with matching releases once the command finishes.

Summary of Object Flow

The complete lifecycle for every enqueued command follows this pattern: retain → enqueue → execute → release → callbacks. Understanding this flow is key to debugging reference-count leaks in applications tested under Oclgrind.

Command Types Handled

Edge Cases & Common Pitfalls

Troubleshooting

Reference count leak detected

If Oclgrind reports objects not freed at context destruction, check that every clRetainEvent is balanced by a clReleaseEvent. The async queue adds one implicit retain per event; user code must release its own references.

Callbacks fire in unexpected order

Callbacks are invoked in registration order per event, but events themselves complete in enqueue order (the queue is serial). If you register callbacks across multiple events, their interleaving depends on command completion order.

Related Topics

• OpenCL API Layer — the cl* functions that call into asyncEnqueue
• ICD Dispatch Table — how API calls are routed to the runtime
• Simulation Lifecycle — the direct-execution alternative that bypasses the async queue
• Runtime Module Overview — parent page with links to all runtime sub-topics