ICD Dispatch Table Deep-Dive

What is ICD?

The Installable Client Driver (ICD) model lets multiple OpenCL implementations coexist. An ICD loader (e.g. libOpenCL.so) dispatches each cl* call through a per-platform function-pointer table. Every OpenCL struct's first field is void* dispatch, pointing to this table.

graph LR
  App[Host Application] -->|clCreateBuffer| Loader[ICD Loader libOpenCL.so]
  Loader -->|dispatch table[42]| OG[oclgrind-rt.so]
  Loader -->|dispatch table[42]| NV[nvidia-opencl.so]
  Loader -->|dispatch table[42]| AMD[amdocl64.so]

The Rename Macro Block

When built with -DOCLGRIND_ICD, icd.h renames every public cl* symbol to _cl*. This avoids linker clashes with the ICD loader — Oclgrind's functions are only reachable through the dispatch table.

// icd.h — symbol renaming (excerpt)
#ifdef OCLGRIND_ICD
#define clCreateContext      _clCreateContext
#define clCreateBuffer       _clCreateBuffer
#define clEnqueueNDRangeKernel _clEnqueueNDRangeKernel
// ... ~90 more definitions
#endif

Struct Layout — dispatch First

Every struct in icd.h starts with void* dispatch. The ICD loader casts the handle, reads the first pointer, and indexes into it. Oclgrind sets this pointer to m_dispatchTable in every new _cl_* allocation.

// Creating a new event — dispatch is always first
cl_event _event     = new _cl_event;
_event->dispatch    = m_dispatchTable;  // ICD slot pointer
_event->context     = queue->context;
_event->queue       = queue;
_event->type        = type;
_event->event       = event;            // Core::Event*
_event->refCount    = 1;

oclgrind.cpp — The Injection Launcher

The oclgrind executable does not implement OpenCL. It launches the user's application with LD_PRELOAD (Linux), DYLD_INSERT_LIBRARIES (macOS), or DLL injection (Windows) pointing to liboclgrind-rt.

sequenceDiagram
  participant User as $ oclgrind ./myapp
  participant OG as oclgrind.cpp
  participant ENV as Environment
  participant Target as myapp
  User->>OG: parseArguments()
  OG->>OG: map --data-races → OCLGRIND_DATA_RACES=1
  OG->>ENV: setenv(LD_PRELOAD, liboclgrind-rt.so)
  OG->>ENV: setenv(LD_LIBRARY_PATH, ...)
  OG->>Target: execvp(appArgs)
  Note over Target: Loads with oclgrind-rt intercepting OpenCL

Windows DLL Injection

On Windows, oclgrind.cpp takes a different approach: it creates the target process in a suspended state, allocates memory in the child process for the DLL path, writes it, and launches a remote thread calling LoadLibraryA. After the DLL loads, it calls initOclgrind via another remote thread, then resumes the main thread.

Platform Enumeration

The ICD loader discovers Oclgrind by calling clIcdGetPlatformIDsKHR (exported from the runtime library). This function lazily creates the singleton m_platform and m_device, reading memory sizes from environment variables like OCLGRIND_GLOBAL_MEM_SIZE.

Platform Info Queries

clGetPlatformInfo returns metadata about the Oclgrind platform. The implementation is a switch statement over param_name that returns static strings and version numbers:

// clGetPlatformInfo excerpt
switch (param_name) {
  case CL_PLATFORM_NAME:    data = "Oclgrind";       break;
  case CL_PLATFORM_VENDOR:  data = "Oclgrind";       break;
  case CL_PLATFORM_VERSION:
    data = "OpenCL 3.0 (Oclgrind " PACKAGE_VERSION ")";
    break;
  case CL_PLATFORM_EXTENSIONS:
    data = "cl_khr_icd";
    break;
  case CL_PLATFORM_ICD_SUFFIX_KHR:
    data = "oclg";
    break;
  // ...
}

Device Initialisation Detail

The singleton m_device is created alongside the platform. Its memory sizes are populated from environment variables with defaults:

• OCLGRIND_GLOBAL_MEM_SIZE → 128 MiB
• OCLGRIND_CONSTANT_MEM_SIZE → 64 KiB
• OCLGRIND_LOCAL_MEM_SIZE → 32 KiB
• OCLGRIND_MAX_WGSIZE → 1024

These values are read once via oclgrind::getEnvInt() and cached in the device struct for the lifetime of the process.

Linux vs macOS vs Windows

The injection mechanism differs by platform:

Library Path Resolution

getLibDirPath() in oclgrind.cpp resolves the path to the library directory by reading /proc/self/exe (Linux), _NSGetExecutablePath (macOS), or GetModuleFileName (Windows), stripping the binary name, then appending /lib (or staying in-place during testing).

ICD Dispatch Table Layout

The table is a 256-slot void* array. Each slot corresponds to the ordinal defined by the Khronos ICD specification. For example slot 0 is clGetPlatformIDs, slot 4 is clCreateContext, etc. Oclgrind fills every slot it supports; unsupported slots are NULL (the ICD loader returns CL_INVALID_OPERATION if called).

How Dispatch Works at Runtime

When the ICD loader receives a call like clCreateBuffer(ctx, ...), it dereferences ctx->dispatch (which is m_dispatchTable), indexes to the clCreateBuffer slot, and calls the function pointer. This indirection costs one memory load per API call — negligible for a simulator.

Supported OpenCL Version

Oclgrind advertises OpenCL 3.0 at both the platform and device level. The ICD suffix is "oclg" — this is used by the loader to distinguish Oclgrind from other drivers in /etc/OpenCL/vendors/.

The cl_khr_icd Extension

The only extension Oclgrind reports is cl_khr_icd. This extension defines two entry points:

• clIcdGetPlatformIDsKHR — discovers the platform
• clGetExtensionFunctionAddressForPlatform — looks up extension functions by name

Oclgrind's implementation of the latter returns NULL for everything except "clIcdGetPlatformIDsKHR" and "clGetPlatformInfo".

Security Implications

The LD_PRELOAD injection technique requires that the target binary is not setuid/setgid — the dynamic linker ignores LD_PRELOAD for privileged executables. On Windows, the remote-thread DLL injection requires the launching process to have PROCESS_ALL_ACCESS rights on the child, which is granted automatically since oclgrind.exe creates the child itself.

Testing Mode

When the environment variable OCLGRIND_TESTING is set, getLibDirPath() skips the ../lib suffix and looks for the runtime library in the same directory as the executable. This allows the test suite to run against an uninstalled build.

Edge Cases & Common Pitfalls

Troubleshooting

Application does not see Oclgrind platform

Run clinfo under the same environment. If Oclgrind is missing, either the .icd file is absent or LD_PRELOAD was not set. On Linux, verify with ls /etc/OpenCL/vendors/ and check that the path inside the .icd file resolves to an existing liboclgrind-rt.so.

Windows injection fails with access denied

If oclgrind.exe cannot open the child process with PROCESS_ALL_ACCESS, anti-virus software may be blocking remote thread creation. Add an exception for oclgrind.exe or run from an elevated command prompt.

Symbol clashes when not using ICD mode

If Oclgrind is built without -DOCLGRIND_ICD, the cl* symbols are exported directly and will clash with the system ICD loader. Always build with ICD mode enabled when installing alongside other OpenCL drivers.

Platform Discovery Flow

When an application calls clGetPlatformIDs, the ICD loader scans /etc/OpenCL/vendors/*.icd for vendor files. Each file contains the path to a shared library. The loader opens the library and calls clIcdGetPlatformIDsKHR to retrieve the platform handle. Oclgrind's implementation lazily creates the singleton m_platform and m_device at this point, reading memory configuration from environment variables.

Vendor file registration

To register without using LD_PRELOAD, create /etc/OpenCL/vendors/oclgrind.icd containing the absolute path to liboclgrind-rt.so. This approach is preferred for system-wide installation and avoids the security restrictions of LD_PRELOAD on setuid binaries.

Related Topics

• OpenCL API Layer — the cl* implementations behind the dispatch table
• Async Command Queue — how enqueued commands flow after dispatch
• CLI Interface — the oclgrind-kernel tool that bypasses ICD entirely
• Runtime Module Overview — parent page with links to all runtime sub-topics