libopencl — OpenCL Runtime API

OpenCL 1.x Runtime API Translation Layer

Implements libOpenCL.so as a drop-in replacement for the NVIDIA OpenCL driver. OpenCL calls are translated into the same GPGPU-Sim simulator infrastructure used by CUDA — OpenCL source is compiled to PTX via nvopencl_wrapper, then loaded into the simulator's PTX parser for cycle-accurate execution.

Source Files

OpenCL

API Layer

30+

Functions

C++

Language

OpenCL 1.x PTX Compilation NDRange → Grid

💡 What is libopencl?

Analogy: If libcuda is the "customs checkpoint" for CUDA traffic, libopencl is the translator booth next door — it accepts OpenCL-speaking applications, translates their requests into the same internal language (PTX + kernel_info_t), and sends them through the same simulation pipeline. The translation includes compiling OpenCL C source to PTX and mapping OpenCL's NDRange work dimensions to CUDA-style grid/block dimensions.

Architecture

flowchart LR
  subgraph App["OpenCL Application"]
    A1["clCreateBuffer()"]
    A2["clBuildProgram()"]
    A3["clEnqueueNDRange()"]
  end
  subgraph Libopencl["libopencl / opencl_runtime_api.cc"]
    B1["_cl_context\nMemory management"]
    B2["_cl_program\nnvopencl_wrapper"]
    B3["_cl_kernel\nArgument binding"]
    B4["NDRange Mapper\nglobal→grid/block"]
  end
  subgraph Core["GPGPU-Sim Core"]
    C1["gpgpu_opencl_ptx_sim_main_perf()"]
    C2["gpgpu_sim::cycle()"]
  end
  A1 --> B1
  A2 --> B2
  A3 --> B3
  B3 --> B4
  B2 --> B3
  B1 --> C1
  B4 --> C1
  C1 --> C2

Data Flow — OpenCL Kernel Execution

flowchart TD
  S1["clCreateProgramWithSource()\nStore OpenCL C source"] --> S2["clBuildProgram()\nInvoke nvopencl_wrapper"]
  S2 --> S3["nvopencl_wrapper\nOpenCL C → PTX compilation"]
  S3 --> S4["gpgpu_ptx_sim_load_ptx_from_string()\nBuild symbol_table"]
  S4 --> S5["clCreateKernel()\nLookup function_info* by name"]
  S5 --> S6["clSetKernelArg() × N\nBind arguments by index"]
  S6 --> S7["clEnqueueNDRangeKernel()\nMap work dims → grid/block"]
  S7 --> S8["Create kernel_info_t\ngridDim × blockDim"]
  S8 --> S9["gpgpu_opencl_ptx_sim_main_perf()\nCycle-level simulation"]

OpenCL → CUDA Concept Mapping

flowchart LR
  subgraph OpenCL["OpenCL Concepts"]
    OA["cl_platform_id"]
    OB["cl_device_id"]
    OC["cl_context"]
    OD["cl_command_queue"]
    OE["cl_mem"]
    OF["cl_program"]
    OG["cl_kernel"]
    OH["NDRange\nglobal × local"]
  end
  subgraph CUDA["Simulator Internal"]
    CA["Singleton"]
    CB["_cuda_device_id"]
    CC["CUctx_st"]
    CD["CUstream_st"]
    CE["Device pointer"]
    CF["symbol_table"]
    CG["function_info*"]
    CH["gridDim × blockDim"]
  end
  OA --> CA
  OB --> CB
  OC --> CC
  OD --> CD
  OE --> CE
  OF --> CF
  OG --> CG
  OH --> CH

🚶 Step-by-Step Narrated Walkthrough

Follow a single clEnqueueNDRangeKernel() call through the libopencl pipeline:

Platform Discovery — clGetPlatformIDs() returns the singleton g_gpgpu_sim_platform_id. The application sees one GPGPU-Sim platform.
Device Selection — clGetDeviceIDs() returns a _cl_device_id wrapping the simulated gpgpu_sim instance.
Context Creation — clCreateContext() creates a _cl_context binding the device, initialising the memory map m_hostptr_to_cl_mem.
Program Compilation — clCreateProgramWithSource() stores the OpenCL C source string. clBuildProgram() writes it to a temp file, invokes nvopencl_wrapper, reads back PTX, and calls gpgpu_ptx_sim_load_ptx_from_string().
Kernel Creation — clCreateKernel(name) looks up the function_info* in pgm_info::m_kernels by kernel name.
Argument Binding — Each clSetKernelArg(index, size, value) stores an arg_info in the kernel's argument map.
Work Dimension Mapping — clEnqueueNDRangeKernel() computes GridDim = global_work_size / local_work_size for each of 1–3 dimensions.
Kernel Launch — A kernel_info_t is created with the computed grid/block dims and enqueued via gpgpu_opencl_ptx_sim_main_perf().
Simulation — The core gpgpu_sim::cycle() loop executes the kernel to completion.
Synchronisation — clFinish() blocks until all enqueued operations complete on the command queue.

Key Classes & Structures

#	Class / Struct	File	Description
1	`_cl_platform_id`	opencl_runtime_api.cc	Singleton platform with static `m_uid=0`. Global instance `g_gpgpu_sim_platform_id`.
2	`_cl_device_id`	opencl_runtime_api.cc	Device wrapping `gpgpu_sim *m_gpgpu`. Methods: `next()`, `the_device()`.
3	`_cl_context`	opencl_runtime_api.cc	Context holding device ref and memory map `m_hostptr_to_cl_mem`. Methods: `CreateBuffer()`, `lookup_mem()`.
4	`_cl_mem`	opencl_runtime_api.cc	Memory object: `m_device_ptr`, `m_host_ptr`, `m_is_on_host`. Supports USE/ALLOC/COPY host pointer modes.
5	`_cl_command_queue`	opencl_runtime_api.cc	Command queue: `m_context`, `m_device`. Warns if out-of-order property used.
6	`_cl_program`	opencl_runtime_api.cc	Program from source. `m_pgm` map of `pgm_info` per device. Methods: `Build()`, `CreateKernel()`, `get_ptx()`.
7	`pgm_info`	opencl_runtime_api.cc	Per-device program metadata: `m_source`, `m_asm` (compiled PTX), `m_symtab`, `m_kernels`.
8	`_cl_kernel`	opencl_runtime_api.cc	Kernel object: `m_kernel_name`, `m_kernel_impl`, `m_args`. Methods: `SetKernelArg()`, `bind_args()`.
9	`arg_info`	opencl_runtime_api.cc	Per-argument descriptor: index, size, and value pointer for kernel arguments.
10	`function_info`	ptx_ir.h (external)	Decoded PTX kernel representation shared with libcuda.

Key API Functions

Category	Function	Purpose
Platform	`clGetPlatformIDs()`	Returns 1 platform (GPGPU-Sim)
Device	`clGetDeviceIDs()`	Returns simulated GPU device
Device	`clGetDeviceInfo()`	80+ attributes from sim config
Context	`clCreateContext()`	Creates `_cl_context` wrapping device
Queue	`clCreateCommandQueue()`	Creates `_cl_command_queue`
Memory	`clCreateBuffer()`	Allocates on simulated GPU via `gpu_malloc()`
Program	`clCreateProgramWithSource()`	Stores OpenCL C source string
Program	`clBuildProgram()`	Compiles via nvopencl_wrapper → PTX
Kernel	`clCreateKernel()`	Extracts `function_info*` by kernel name
Kernel	`clSetKernelArg()`	Sets argument by index
Execute	`clEnqueueNDRangeKernel()`	Maps NDRange → grid/block, creates `kernel_info_t`, launches
Transfer	`clEnqueueReadBuffer()`	Device → host via `memcpy_from_gpu()`
Transfer	`clEnqueueWriteBuffer()`	Host → device via `memcpy_to_gpu()`
Sync	`clFinish()`	Synchronises command queue

Module Interactions

Interacts With	Direction	Mechanism	Purpose
OpenCL Application	← inbound	`LD_PRELOAD` / link-time	Intercepts all `cl*()` API calls
nvopencl_wrapper	→ outbound	Process invocation	Compiles OpenCL C source to PTX via NVIDIA compiler
cuda-sim (PTX parser)	→ outbound	`gpgpu_ptx_sim_load_ptx_from_string()`	Loads compiled PTX into the functional simulator
gpgpu_sim core	→ outbound	`gpgpu_opencl_ptx_sim_main_perf()`	Launches kernel for cycle-level execution
libcuda	↔ shared infra	Same simulator internals	Both create `kernel_info_t` and share `gpgpu_sim` core

Configuration & Environment Variables

Variable	Type	Description
`OPENCL_REMOTE_GPU_HOST`	Env	Hostname for remote OpenCL compilation (if local GPU unavailable)
`NVOPENCL_LIBDIR`	Env	Path to NVIDIA's OpenCL compiler library
`LD_LIBRARY_PATH`	Env	Must include NVIDIA OpenCL driver libs
`CUDA_INSTALL_PATH`	Env	Required for nvopencl_wrapper compilation

Code Snippets

Work Dimension Mapping — NDRange to Grid/Block

// opencl_runtime_api.cc — clEnqueueNDRangeKernel core logic
// Map OpenCL work dimensions to CUDA grid/block dimensions
dim3 GridDim, BlockDim;

GridDim.x = global_work_size[0] / local_work_size[0];
GridDim.y = (work_dim >= 2) ? global_work_size[1] / local_work_size[1] : 1;
GridDim.z = (work_dim >= 3) ? global_work_size[2] / local_work_size[2] : 1;

BlockDim.x = local_work_size[0];
BlockDim.y = (work_dim >= 2) ? local_work_size[1] : 1;
BlockDim.z = (work_dim >= 3) ? local_work_size[2] : 1;

// Create kernel_info_t and launch
kernel_info_t *grid = new kernel_info_t(GridDim, BlockDim, entry);
gpgpu_opencl_ptx_sim_main_perf(grid);

Program Build — OpenCL Source to PTX

// _cl_program::Build() — simplified
cl_int _cl_program::Build(const char *options) {
  // 1. Write OpenCL source to temp file
  FILE *fp = fopen(src_filename, "w");
  fprintf(fp, "%s", m_pgm[device].m_source.c_str());
  fclose(fp);

  // 2. Invoke nvopencl_wrapper (NVIDIA's compiler)
  std::string cmd = "nvopencl_wrapper " + src_filename;
  if (getenv("OPENCL_REMOTE_GPU_HOST"))
    cmd = "ssh " + remote_host + " " + cmd;
  system(cmd.c_str());

  // 3. Read compiled PTX and load into simulator
  std::string ptx = read_file(ptx_filename);
  gpgpu_ptx_sim_load_ptx_from_string(ptx, source_num);
  return CL_SUCCESS;
}

Directory Map

Source tree for libopencl

libopencl/
├── opencl_runtime_api.cc  ← 1,540 lines — implements 30+ cl*() API functions
│                             Defines all _cl_* classes inline
├── nvopencl_wrapper.cc    ← Wrapper for NVIDIA's OpenCL-to-PTX compiler
├── Makefile               ← Builds libOpenCL.so
└── CMakeLists.txt         ← CMake build configuration

Glossary

Term	Definition
NDRange	N-Dimensional Range — OpenCL's equivalent of CUDA grid + block launch specification
Work Group	OpenCL equivalent of CUDA thread block (CTA) — threads that can synchronise via barriers
Work Item	OpenCL equivalent of a single CUDA thread
nvopencl_wrapper	GPGPU-Sim script that invokes NVIDIA's OpenCL compiler to produce PTX from OpenCL C source
cl_mem	OpenCL memory object representing a buffer on the device (mapped to a simulated GPU address)
Command Queue	Ordered sequence of OpenCL operations — analogous to a CUDA stream