libcuda — CUDA Runtime API Interception

CUDA Runtime API Interception Layer

Implements libcudart.so as a drop-in replacement for the NVIDIA CUDA runtime. Every CUDA API call from the application is intercepted and redirected into the GPGPU-Sim cycle-level simulator — memory allocation, data transfer, kernel launch, and synchronisation all run on the simulated GPU.

Source Files

CUDA

API Layer

80+

Functions Intercepted

C++

Language

LD_PRELOAD API Interception Kernel Launch Binary Parsing

💡 What is libcuda?

Analogy: Think of libcuda as a customs checkpoint at a border crossing. Every CUDA call from your application tries to "cross the border" to the GPU. Instead of reaching real hardware, libcuda intercepts the call and routes it to GPGPU-Sim's simulated GPU — inspecting every packet (argument), stamping it (creating kernel_info_t), and forwarding it to the simulation pipeline.

Architecture

flowchart LR
  subgraph App["CUDA Application"]
    A1["cudaMalloc()"]
    A2["cudaMemcpy()"]
    A3["cudaLaunch()"]
  end
  subgraph Libcuda["libcuda / cuda_runtime_api.cc"]
    B1["Memory Manager\ng_mallocPtr_Size"]
    B2["Stream Manager\nCUstream_st"]
    B3["Kernel Resolver\nCUctx_st::m_kernel_lookup"]
    B4["Binary Loader\ncuobjdump parser"]
  end
  subgraph Core["GPGPU-Sim Core"]
    C1["gpgpu_sim::cycle()"]
    C2["stream_manager"]
  end
  A1 --> B1
  A2 --> B2
  A3 --> B3
  B4 --> B3
  B1 --> C2
  B2 --> C2
  B3 --> C2
  C2 --> C1

Data Flow — Kernel Launch Path

flowchart TD
  R["__cudaRegisterFatBinary()"] --> E["Extract PTX via cuobjdump"]
  E --> S["Build symbol_table"]
  S --> RF["__cudaRegisterFunction()\nMap hostFun → function_info*"]
  RF --> CC["cudaConfigureCall()\nPush kernel_config onto stack"]
  CC --> SA["cudaSetupArgument()\nAppend gpgpu_ptx_sim_arg"]
  SA --> CL["cudaLaunch()\nResolve function_info*"]
  CL --> KI["Create kernel_info_t\n(grid, block, entry, stream)"]
  KI --> SM["stream_manager::push()\nstream_operation(kernel)"]
  SM --> SIM["gpgpu_sim_thread_concurrent\ngpu→launch(kernel)"]
  SIM --> CY["gpgpu_sim::cycle()\nShader pipeline → Memory"]

🚶 Step-by-Step Narrated Walkthrough

Follow a single cudaLaunch(myKernel) call as it flows through the libcuda pipeline:

Fat Binary Registration — At program start, __cudaRegisterFatBinary() extracts PTX from the CUDA fat binary using cuobjdump and builds a symbol_table.
Function Registration — __cudaRegisterFunction() maps the host-side function pointer to a function_info* in CUctx_st::m_kernel_lookup.
Configure Call — cudaConfigureCall(gridDim, blockDim, sharedMem, stream) pushes a kernel_config onto g_cuda_launch_stack.
Setup Arguments — Each cudaSetupArgument() appends a gpgpu_ptx_sim_arg (pointer, size, offset) to the top config's argument list.
Launch Resolution — cudaLaunch(hostFun) looks up function_info* via the host function pointer and pops the kernel config from the stack.
Kernel Info Creation — A kernel_info_t is built with grid/block dims, the function entry point, arguments, and the target stream ID.
Texture Snapshot — Texture binding state is copied into kernel_info_t so the kernel sees a consistent view.
Stream Enqueue — stream_manager::push() places a stream_operation of type stream_kernel_launch into the command queue.
Simulation Thread — gpgpu_sim_thread_concurrent() dequeues the operation and calls gpu→launch(kernel).
Cycle Simulation — gpgpu_sim::cycle() runs the shader pipeline and memory hierarchy until all CTAs complete.

Key Classes & Structures

#	Class / Struct	File	Description
1	`gpgpu_context`	gpgpu_context.h	Master singleton context. Owns pointers to `cuda_runtime_api`, `cuda_sim`, `GPGPUsim_ctx*`. Entry points: `synchronize()`, `start_sim_thread()`.
2	`cuda_runtime_api`	cuda_api_object.h	Runtime state: `fatbinmap`, `name_symtab`, `pinned_memory`, `g_mallocPtr_Size`. Key methods: `cuobjdumpInit()`, `extract_code_using_cuobjdump()`.
3	`CUctx_st`	cuda_api_object.h	CUDA context. Links host functions to device code via `m_kernel_lookup: map<const void, function_info>`. Stores per-binary `m_code` symbol tables.
4	`_cuda_device_id`	cuda_api_object.h	Device abstraction wrapping `gpgpu_sim*`. Methods: `num_devices()`, `get_device(n)`, `get_prop()`.
5	`kernel_config`	cuda_api_object.h	Pending launch state: `m_GridDim`, `m_BlockDim`, `m_args` (arg list), `m_stream`.
6	`CUevent_st`	stream_manager.h	CUDA event object for timing and synchronisation.
7	`CUstream_st`	stream_manager.h	CUDA stream — ordered queue of GPU operations.
8	`gpgpu_ptx_sim_arg`	abstract_hardware_model.h	Kernel argument descriptor: `m_start`, `m_nbytes`, `m_offset`.
9	`function_info`	ptx_ir.h	Decoded PTX kernel representation; looked up by name or host function pointer.
10	`symbol_table`	ptx_ir.h	PTX symbol table — one per loaded binary; maps names to definitions.
11	`glbmap_entry`	cuda_api_object.h	OpenGL interop buffer mapping: `m_bufferObj`, `m_devPtr`, `m_size`.

Key API Functions

Category	Function	Purpose
Device	`cudaGetDeviceCount()`	Returns 1 (simulated GPU)
Device	`cudaGetDeviceProperties()`	Fills `cudaDeviceProp` from sim config
Memory	`cudaMalloc()`	`gpu_malloc()` — increments from `GLOBAL_HEAP_START`
Memory	`cudaMemcpy()`	Dispatches H2D/D2H/D2D via `stream_operation`
Memory	`cudaMemcpyAsync()`	Queues memcpy on specified stream
Memory	`cudaFree()`	No-op (no garbage collection in simulator)
Memory	`cudaMallocHost()`	Pinned memory — host `malloc` + tracking map
Launch	`cudaConfigureCall()`	Pushes `kernel_config` onto `g_cuda_launch_stack`
Launch	`cudaSetupArgument()`	Appends `gpgpu_ptx_sim_arg` to config args
Launch	`cudaLaunch()`	Resolves `function_info`, creates `kernel_info_t`, enqueues
Sync	`cudaDeviceSynchronize()`	Global barrier — waits for all streams
Stream	`cudaStreamCreate()`	Creates `CUstream_st` object
Event	`cudaEventRecord()`	Records event in stream timeline
Register	`__cudaRegisterFatBinary()`	Extracts PTX from fat binary using cuobjdump
Register	`__cudaRegisterFunction()`	Maps host function pointer → device kernel name

Module Interactions

Interacts With	Direction	Mechanism	Purpose
CUDA Application	← inbound	`LD_PRELOAD` / link-time	Intercepts all `cuda*()` API calls
cuobjdump_to_ptxplus	→ outbound	Process invocation	Converts fat binary to PTXPlus when enabled
cuda-sim (PTX parser)	→ outbound	`gpgpu_ptx_sim_load_ptx_from_*`	Loads extracted PTX into functional simulator
stream_manager	→ outbound	`push(stream_operation)`	Enqueues kernel launches and memcpy operations
gpgpu_sim	→ outbound	`gpu→launch(kernel)`	Hands off kernels for cycle-level simulation
gpgpusim_entrypoint	↔ both	Thread synchronisation	Manages simulation thread lifecycle

Configuration & Environment Variables

Variable	Type	Description
`CUDA_INSTALL_PATH`	Env	Path to CUDA toolkit (required for cuobjdump)
`GPGPUSIM_ROOT`	Env	Simulator root directory
`PTX_SIM_MODE_FUNC`	Env	`1` = functional-only; `0` = performance simulation
`CUOBJDUMP_SIM_FILE`	Env	Override binary path for cuobjdump extraction
`PTX_SIM_USE_PTX_FILE`	Env	Use pre-extracted PTX file instead of fat binary
`-gpgpu_ptx_use_cuobjdump`	Config	Enable binary parsing via cuobjdump
`-gpgpu_ptx_convert_to_ptxplus`	Config	Enable PTXPlus translation mode
`-gpgpu_ptx_force_max_capability`	Config	Force a specific SM compute capability

Code Snippets

Kernel Launch — cudaLaunch() core logic

// cuda_runtime_api.cc — simplified cudaLaunch flow
cudaError_t cudaLaunch(const char *hostFun) {
  // 1. Resolve device function from host pointer
  CUctx_st *ctx = GPGPUSim_Context();
  function_info *entry = ctx->get_kernel(hostFun);

  // 2. Pop kernel_config from launch stack
  kernel_config config = g_cuda_launch_stack.back();
  g_cuda_launch_stack.pop_back();

  // 3. Create kernel_info_t with grid/block dims
  kernel_info_t *grid = new kernel_info_t(
      config.grid_dim(), config.block_dim(),
      entry, config.get_args());

  // 4. Enqueue into stream_manager
  stream_operation op(grid, sim_mode, config.get_stream());
  g_stream_manager->push(op);
  return cudaSuccess;
}

Memory Allocation — GPU address space

// Device memory allocation tracking
std::map<unsigned long long, size_t> g_mallocPtr_Size;

void *gpu_malloc(size_t size) {
  unsigned long long addr = m_dev_malloc;  // starts at GLOBAL_HEAP_START
  m_dev_malloc += size;
  // Align to 256 bytes
  if (size % 256) m_dev_malloc += (256 - size % 256);
  g_mallocPtr_Size[addr] = size;
  return (void *)addr;
}

Directory Map

Source tree for libcuda

libcuda/
├── cuda_runtime_api.cc    ← 7,260 lines — implements 80+ cuda*() API functions
├── cuda_api_object.h      ← Declares CUctx_st, _cuda_device_id, kernel_config
├── cuda_api.h             ← Public header for cuda runtime API declarations
├── gpgpu_context.h        ← Central gpgpu_context singleton class
├── cuobjdump.h            ← cuobjdump parsing declarations
├── cuobjdump.l            ← Flex lexer for cuobjdump output
├── cuobjdump.y            ← Bison parser for cuobjdump output
├── Makefile               ← Builds libcudart.so
└── CMakeLists.txt         ← CMake build configuration

Glossary

Term	Definition
Fat Binary	NVIDIA compiled binary containing PTX and/or cubin for multiple GPU architectures
cubin	Compiled CUDA binary (SASS code) for a specific GPU compute capability
Host Function Pointer	CPU-side address of a `__global__` kernel stub; used as lookup key to find device code
Symbol Table	Maps PTX symbols (functions, variables) to their definitions and memory locations
Pinned Memory	Non-pageable host memory enabling faster DMA transfers to device
Stream Operation	Queued GPU command (kernel launch, memcpy, event) in a CUDA stream