Oclgrind runtime — OpenCL ICD Runtime

OpenCL 1.2 API ICD Dispatch Async Queue LD_PRELOAD Injection

Table of Contents

1 — Architecture & API Flow

Every OpenCL API call enters through the ICD dispatch table, which maps the function pointer to the Oclgrind implementation in runtime.cpp. The runtime function validates arguments, creates or updates the appropriate wrapper struct, delegates to the core simulation engine, and returns the result to the caller.

flowchart TD
    APP["OpenCL Application"]
    ICD["ICD Loader
libOpenCL.so"]
    DISP["m_dispatchTable[256]
icd.h / runtime.cpp"]
    RT["runtime.cpp
cl* implementations"]
    AQ["async_queue
asyncEnqueue / asyncQueueRelease"]
    CORE_CTX["core::Context"]
    CORE_Q["core::Queue"]
    CORE_K["core::Kernel"]
    CORE_M["core::Memory"]
    CORE_P["core::Program"]

    APP -->|"clEnqueue*()"| ICD
    ICD -->|"dispatch[N]"| DISP
    DISP --> RT
    RT -->|"create wrappers
validate args"| RT
    RT -->|"kernel / mem ops"| AQ
    AQ -->|"queue→enqueue(cmd)"| CORE_Q
    RT -->|"context mgmt"| CORE_CTX
    RT -->|"buffer alloc / free"| CORE_M
    RT -->|"build / compile"| CORE_P
    CORE_Q -->|"dispatch kernel"| CORE_K
    CORE_K -->|"read / write"| CORE_M

    style APP fill:#1a2332,stroke:#58a6ff,color:#c9d1d9
    style ICD fill:#1a2332,stroke:#d29922,color:#c9d1d9
    style DISP fill:#1c2128,stroke:#58a6ff,color:#c9d1d9
    style RT fill:#1c2128,stroke:#3fb950,color:#c9d1d9
    style AQ fill:#1c2128,stroke:#d29922,color:#c9d1d9
    style CORE_CTX fill:#0d1117,stroke:#58a6ff,color:#c9d1d9
    style CORE_Q fill:#0d1117,stroke:#58a6ff,color:#c9d1d9
    style CORE_K fill:#0d1117,stroke:#58a6ff,color:#c9d1d9
    style CORE_M fill:#0d1117,stroke:#58a6ff,color:#c9d1d9
    style CORE_P fill:#0d1117,stroke:#58a6ff,color:#c9d1d9

Figure 1 — OpenCL API call flow from application through ICD dispatch into Core

2 — Async Command Queue

The AsyncQueue subsystem handles all enqueue operations. When a clEnqueue* call is made, the runtime creates a core::Command object, retains any referenced cl_mem or cl_kernel objects, attaches the event wait list, and hands it off to core::Queue::enqueue(). On completion, asyncQueueRelease fires callbacks, releases retained objects, and cleans up event references. Four static maps track the lifecycle: memObjectMap, kernelMap, eventMap, and waitListMap.

sequenceDiagram
    participant App as Application
    participant RT as runtime.cpp
    participant AQ as async_queue
    participant CQ as core::Queue
    participant EV as cl_event

    App->>RT: clEnqueueNDRangeKernel()
    RT->>RT: validate args, clone kernel
    RT->>AQ: asyncQueueRetain(cmd, kernel)
    AQ->>AQ: clRetainKernel + retain memArgs
    RT->>AQ: asyncEnqueue(queue, type, cmd, waitList)
    AQ->>AQ: attach wait-list events
    AQ->>CQ: queue→enqueue(cmd)
    CQ-->>AQ: return core::Event*
    AQ->>EV: create _cl_event wrapper
    AQ-->>RT: event ready
    RT-->>App: CL_SUCCESS + event

    Note over CQ: Command executes asynchronously

    CQ->>AQ: command complete callback
    AQ->>AQ: asyncQueueRelease(cmd)
    AQ->>AQ: release cl_mem, cl_kernel
    AQ->>EV: fire user callbacks
    AQ->>AQ: release wait-list events

Figure 2 — Sequence diagram of an async kernel enqueue through the command queue

3 — Key Structs & Types

4 — Directory Map

5 — Module Interactions

6 — Configuration & Environment Variables

7 — Code Examples

7.1 — Platform & Device Initialization (runtime.cpp)

The ICD entry point lazily creates the singleton platform and device, reading memory sizes from environment variables:

CL_API_ENTRY cl_int CL_API_CALL clIcdGetPlatformIDsKHR(
  cl_uint num_entries, cl_platform_id* platforms, cl_uint* num_platforms)
{
  REGISTER_API;
  if (platforms && num_entries < 1)
    ReturnError(NULL, CL_INVALID_VALUE);

  if (!m_platform) {
    m_platform = new _cl_platform_id;
    m_platform->dispatch = m_dispatchTable;

    m_device = new _cl_device_id;
    m_device->dispatch = m_dispatchTable;
    m_device->globalMemSize  = oclgrind::getEnvInt(
      "OCLGRIND_GLOBAL_MEM_SIZE", DEFAULT_GLOBAL_MEM_SIZE, false);
    m_device->localMemSize   = oclgrind::getEnvInt(
      "OCLGRIND_LOCAL_MEM_SIZE", DEFAULT_LOCAL_MEM_SIZE, false);
    m_device->maxWGSize      = oclgrind::getEnvInt(
      "OCLGRIND_MAX_WGSIZE", DEFAULT_MAX_WGSIZE, false);
  }
  if (platforms)     platforms[0] = m_platform;
  if (num_platforms) *num_platforms = 1;
  return CL_SUCCESS;
}

7.2 — Async Enqueue Lifecycle (async_queue.cpp)

When a command completes, asyncQueueRelease tears down retained references and fires user callbacks:

void asyncQueueRelease(Command* cmd)
{
  // Release retained memory objects
  if (memObjectMap.find(cmd) != memObjectMap.end()) {
    list<cl_mem> memObjects = memObjectMap[cmd];
    while (!memObjects.empty()) {
      clReleaseMemObject(memObjects.front());
      memObjects.pop_front();
    }
    memObjectMap.erase(cmd);
  }

  // Release kernel (if this was a kernel command)
  if (cmd->type == Command::KERNEL) {
    clReleaseKernel(kernelMap[cmd]);
    kernelMap.erase(cmd);
    delete ((KernelCommand*)cmd)->kernel;
  }

  // Fire user event callbacks
  cl_event event = eventMap[cmd];
  for (auto &cb : event->callbacks)
    cb.first(event, event->event->state, cb.second);

  // Release wait-list events
  for (auto &e : waitListMap[cmd])
    clReleaseEvent(e);
  waitListMap.erase(cmd);
  clReleaseEvent(event);
}

7.3 — Launcher LD_PRELOAD Injection (oclgrind.cpp)

The oclgrind binary constructs preload paths and uses execvp to replace itself with the target application:

// Construct LD_PRELOAD with the Oclgrind runtime library
string ldPreload = libdir + "/liboclgrind-rt." + LIB_EXTENSION;
const char* oldLdPreload = getenv(LD_PRELOAD_ENV);
if (oldLdPreload) {
  ldPreload += ":";
  ldPreload += oldLdPreload;
}

setEnvironment(LD_LIBRARY_PATH_ENV, ldLibraryPath.c_str());
setEnvironment(LD_PRELOAD_ENV, ldPreload.c_str());

// Replace process with target application
if (execvp(appArgs[0], appArgs) == -1) {
  cerr << "[Oclgrind] Failed to launch target application" << endl;
  exit(1);
}

8 — Deep-Dive Pages