Oclgrind — Virtual OpenCL Device Simulator

flowchart TD
  subgraph Runtime["Runtime Layer (ICD)"]
    RT["runtime.cpp
OpenCL API impl"]
    AQ["AsyncQueue
Command scheduling"]
    ICD["icd.h
ICD dispatch table"]
  end
  subgraph Core["Core Simulation Engine"]
    CTX["Context
Plugin manager"]
    PRG["Program
LLVM compilation"]
    KRN["Kernel
Argument binding"]
    KI["KernelInvocation
NDRange orchestrator"]
    WG["WorkGroup
Barrier sync"]
    WI["WorkItem
IR interpreter"]
    MEM["Memory
Address space mgr"]
  end
  subgraph Plugins["Analysis Plugins"]
    MC["MemCheck"]
    RD["RaceDetector"]
    UI["Uninitialized"]
    IC["InstructionCounter"]
    ID["InteractiveDebugger"]
    LG["Logger"]
  end
  subgraph KernelCLI["Kernel CLI"]
    SIM["Simulation
Standalone runner"]
  end

  RT -->|"clEnqueueNDRange"| AQ
  AQ -->|"dispatch"| CTX
  CTX -->|"compile"| PRG
  CTX -->|"invoke"| KI
  KI -->|"spawn"| WG
  WG -->|"execute"| WI
  WI -->|"load/store"| MEM
  CTX -->|"notify hooks"| MC
  CTX -->|"notify hooks"| RD
  CTX -->|"notify hooks"| UI
  CTX -->|"notify hooks"| IC
  CTX -->|"notify hooks"| ID
  CTX -->|"notify hooks"| LG
  SIM -->|"standalone exec"| CTX
  ICD -->|"dispatch"| RT

sequenceDiagram
  participant App as Host Application
  participant RT as Runtime (ICD)
  participant CTX as Context
  participant PRG as Program
  participant KI as KernelInvocation
  participant WG as WorkGroup
  participant WI as WorkItem
  participant MEM as Memory
  participant PLG as Plugins

  App->>RT: clCreateProgramWithSource()
  RT->>CTX: Create Program
  CTX->>PRG: Compile OpenCL to LLVM IR
  PRG-->>CTX: Module ready
  App->>RT: clEnqueueNDRangeKernel()
  RT->>CTX: Queue kernel dispatch
  CTX->>PLG: notifyKernelBegin()
  CTX->>KI: Create invocation (global/local sizes)
  KI->>WG: Spawn work-groups (parallel or sequential)
  WG->>WI: Create work-items per group
  loop For each instruction
    WI->>WI: Interpret LLVM IR instruction
    WI->>MEM: Load/Store memory
    MEM-->>WI: Data returned
    WI->>CTX: notifyInstructionExecuted()
    CTX->>PLG: Forward to all plugins
  end
  WG->>CTX: notifyWorkGroupComplete()
  KI-->>CTX: All groups done
  CTX->>PLG: notifyKernelEnd()
  CTX-->>RT: Execution complete
  RT-->>App: Event signaled

sequenceDiagram
  participant RT as Runtime (ICD)
  participant CTX as Core::Context
  participant MEM as Core::Memory
  participant PLG as Plugins
  participant KRN as Kernel CLI

  RT->>CTX: clEnqueueNDRangeKernel
  Note over RT,CTX: Runtime translates OpenCL API to Core calls
  CTX->>MEM: allocateBuffer / load / store
  Note over CTX,MEM: Memory validates all accesses
  CTX->>PLG: notifyMemoryLoad / notifyMemoryStore
  Note over CTX,PLG: Every memory op forwarded to active plugins
  PLG-->>CTX: logError (if violation detected)
  KRN->>CTX: Simulation::run (standalone path)
  Note over KRN,CTX: CLI bypasses Runtime, talks directly to Core