Core Infrastructure — Hardware Abstraction & Simulation Entry Point

Hardware Abstraction & Simulation Entry Point

The foundational layer of GPGPU-Sim: hardware abstraction types (warps, instructions, memory accesses), the simulation entry point and threading model, CUDA stream management, configuration parsing, and debug/trace infrastructure. Every other module in the simulator builds on these primitives.

Files

5.3k

Lines of Code

30+

Classes / Structs

C++

Language

Hardware Abstraction CUDA Streams Config Parser Debug / Trace

💡 What Is This Module?

Think of Core Infrastructure as the foundation and wiring of a building. The foundation (abstract_hardware_model) defines the structural shapes—walls, floors, load-bearing beams—that every room in the building must respect. The main electrical panel (gpgpusim_entrypoint) powers the entire structure and decides when the lights come on. The plumbing manifold (stream_manager) routes work orders to the right room. The control panel (option_parser) lets you reconfigure the building’s behaviour without tearing down walls. And the diagnostic wiring (debug + trace) lets you inspect any circuit at runtime.

No concrete timing logic lives here—this module defines the vocabulary (kernels, warps, instructions, memory accesses) that the timing model, cache hierarchy, and interconnect all speak.

Architecture

flowchart TD
  subgraph Runtime["CUDA / OpenCL Runtime"]
    LIBCUDA["libcuda / libopencl"]
  end

  subgraph Core["Core Infrastructure (src/)"]
    AHM["abstract_hardware_model\n(kernel_info_t, warp_inst_t,\nmem_access_t, simt_stack,\ncore_t, gpgpu_t)"]
    ENTRY["gpgpusim_entrypoint\n(GPGPUsim_ctx, sim threads,\ninit & sync)"]
    SM["stream_manager\n(CUstream_st, stream_operation,\nCUevent_st, round-robin)"]
    OPT["option_parser\n(OptionParser, OptionRegistry,\nParseFile / ParseCmdLine)"]
    DBG["debug\n(brk_pt, PTX debugger,\nwatchpoints)"]
    TRC["trace\n(DPRINTF, DTRACE,\ntrace_streams)"]
    STAT["statwrapper\n(Stats histograms,\nBookSim wrapper)"]
  end

  LIBCUDA -->|"creates kernel_info_t\npushes stream_operation"| SM
  SM -->|"do_operation → gpu->launch()"| ENTRY
  ENTRY -->|"creates gpgpu_sim\ndrives cycle()"| AHM
  OPT -->|"feeds core_config,\ngpgpu_sim_config"| ENTRY
  OPT -->|"feeds gpgpu_functional_sim_config"| AHM
  AHM -->|"warp_inst_t flows to\nshader pipeline"| TIMING["Timing Model\n(gpgpu-sim/)"]
  AHM -->|"mem_access_t flows to\ncache hierarchy"| CACHE["Cache / Memory\n(gpgpu-sim/)"]
  DBG -->|"inspects core_t,\nglobal memory"| AHM
  TRC -->|"DPRINTF used by"| TIMING
  STAT -->|"histogram data"| TIMING

Data Flow

sequenceDiagram
  participant App as CUDA Application
  participant Lib as libcuda
  participant SM as stream_manager
  participant Entry as gpgpusim_entrypoint
  participant GPU as gpgpu_sim
  participant Core as shader_core_ctx
  participant Mem as Memory Subsystem

  App->>Lib: cudaLaunchKernel()
  Lib->>Lib: Create kernel_info_t (grid/block dims, function_info)
  Lib->>SM: push(stream_operation{kernel_launch, stream})
  SM->>SM: Enqueue into CUstream_st FIFO
  Note over Entry: Sim thread spins waiting for work
  Entry->>SM: operation(&sim_cycles) — round-robin front()
  SM->>SM: do_operation() — check launch latency
  SM->>GPU: gpu->launch(kernel_info_t)
  GPU->>Core: Distribute CTAs to shader cores
  Core->>Core: Fetch → Decode → Issue warp_inst_t
  Core->>Core: simt_stack::update() — handle divergence
  Core->>Mem: generate_mem_accesses() → coalesce → mem_access_t
  Mem-->>Core: Data returns / cache hits
  Core-->>GPU: Kernel CTAs complete
  GPU-->>SM: register_finished_kernel()
  SM-->>Lib: record_next_done() — unblock host

🚶 Step-by-Step: Tracing a Kernel Launch from API Call to Execution

Follow a single cudaLaunchKernel() call as it flows through Core Infrastructure:

Application calls cudaLaunchKernel() — The CUDA runtime intercept in libcuda captures the call, resolving the kernel function pointer to a function_info* object.
Create kernel_info_t — libcuda allocates a kernel_info_t with grid/block dimensions, a unique UID, stream ID, argument list (gpgpu_ptx_sim_arg), and a snapshot of bound textures.
Push to stream_manager — A stream_operation of type stream_kernel_launch is constructed and pushed into the target CUstream_st’s FIFO. A mutex protects the queue.
Simulation thread wakes — In gpgpu_sim_thread_concurrent(), the spin-loop detects !empty_protected() and calls g_the_gpu->init() to reset cycle counters.
Stream scheduling (stream_manager::front()) — The round-robin scheduler tries stream_zero first, then iterates named streams from m_last_stream. It selects the first non-busy, non-empty stream.
Operation dispatch (do_operation()) — For a kernel launch, the code checks gpu->can_start_kernel() and waits for m_launch_latency to count down to zero (simulating CPU→GPU transfer delay).
gpu->launch(kernel) — The kernel is handed to gpgpu_sim, which maps it to a grid-to-stream entry and begins distributing CTAs across shader cores based on resource availability (registers, shared memory, max CTAs per SM).
Warp creation & SIMT stack init — Each CTA spawns warps. Per-warp simt_stack objects are initialized with the kernel entry PC and a full active mask (all 32 threads active).
Fetch / Decode / Issue cycle — Each gpgpu_sim::cycle() tick drives shader cores. Instructions are fetched, decoded into warp_inst_t objects carrying opcode, operands, latency, and active mask.
Branch divergence (simt_stack::update()) — When a branch diverges, threads are partitioned by next_pc. The PDOM reconvergence algorithm pushes taken/not-taken paths onto the stack with the post-dominator as the reconvergence PC.
Memory coalescing (memory_coalescing_arch()) — For global/local loads and stores, per-thread addresses are mapped to segment-aligned transactions (128B → 64B → 32B reduction). Each reduced transaction becomes a mem_access_t pushed into the access queue.
Kernel completion & cleanup — When all CTAs finish, gpgpu_sim calls stream_manager::register_finished_kernel(). The stream pops the operation via record_next_done(), unblocking any host-side synchronization.

Key Components

Class / Struct	File	Description
`kernel_info_t`	abstract_hardware_model.h	Launched kernel representation: grid/block dims, UID, stream ID, CTA iteration state, parent/child for CDP, launch latency, texture snapshots.
`warp_inst_t`	abstract_hardware_model.h	Instruction bound to a warp: active mask, per-thread addresses, memory access queue (`m_accessq`), coalescing logic, atomic callbacks.
`inst_t`	abstract_hardware_model.h	Base decoded instruction: PC, opcode (`op_type`), barrier info, operand types, latency, initiation interval, register mapping.
`simt_stack`	abstract_hardware_model.h	PDOM-based SIMT reconvergence stack per warp. Implements `launch()`, `update()`, `get_active_mask()`.
`mem_access_t`	abstract_hardware_model.h	Single coalesced memory transaction: address, type (GLOBAL_ACC_R, etc.), size, byte/sector masks, warp mask.
`core_config`	abstract_hardware_model.h	Per-core config: `warp_size`, shared-memory banks, coalescing arch, cache line sizes, max issue per warp.
`gpgpu_t`	abstract_hardware_model.h	Base GPU class: global/tex/surf memory spaces, `gpu_malloc`, `memcpy_to/from_gpu`, texture binding maps.
`core_t`	abstract_hardware_model.h	Abstract SM core: holds SIMT stacks, thread info pointers, `execute_warp_inst_t()`, `updateSIMTStack()`.
`register_set`	abstract_hardware_model.h	Multi-slot pipeline register: `has_free()`, `has_ready()`, `move_in()`, `move_out_to()`. Supports sub-core model.
`memory_space_t`	abstract_hardware_model.h	Type-safe memory-space tag: `global_space`, `shared_space`, `local_space`, `const_space`, `tex_space`, with bank index.
`mem_fetch_interface`	abstract_hardware_model.h	Abstract push interface for memory requests: `full()`, `push()`.
`transaction_info`	abstract_hardware_model.h	Per-transaction coalescing state: 32B chunk bitmask, byte mask, active thread mask.
`GPGPUsim_ctx`	gpgpusim_entrypoint.h	Top-level sim context: semaphores, GPU config pointer, GPU object, stream manager, simulation flags.
`stream_manager`	stream_manager.h	Global stream manager: stream list, grid-to-stream map, `stream_zero`, round-robin `front()`, `operation()` dispatch.
`CUstream_st`	stream_manager.h	Per-stream FIFO of `stream_operation`s, mutex-protected, with pending flag.
`stream_operation`	stream_manager.h	Queued operation (memcpy, kernel launch, event, wait-event) with addresses, sizes, and kernel pointer.
`CUevent_st`	stream_manager.h	CUDA event: UID, blocking flag, update count, simulated cycle, wallclock timestamp.
`OptionParser`	option_parser.cc	Template `Register<T>()`, `ParseCommandLine()`, `ParseFile()`, `ParseString()`, `Print()`.
`OptionRegistry<T>`	option_parser.cc	Templated concrete registry with `fromString()` specializations for bool, char*, string.
`brk_pt`	debug.h	Breakpoint/watchpoint: fileline + thread_uid or address + value. Used by the interactive PTX debugger.
`Trace` namespace	trace.h	Per-subsystem tracing: `enabled`, `sampling_core`, `trace_streams_enabled[]`, `init()`.

Deep Dives

🚀 Memory Coalescing

How warp_inst_t::memory_coalescing_arch() transforms 32 per-thread addresses into minimal memory transactions.

Segment selection: 32B sectors vs 128B lines
Transaction gathering by subwarp
128B → 64B → 32B reduction
Atomic conflict splitting

🔀 SIMT Stack & Reconvergence

The PDOM-based reconvergence algorithm in simt_stack::update() — how warps handle branch divergence.

Thread partitioning by next_pc
Post-dominator reconvergence point
Stack push/pop for taken/not-taken paths
CALL/RET handling for function calls

⚙ Kernel Dispatch Pipeline

From cudaLaunchKernel() to gpu->launch() — the full journey through stream scheduling.

kernel_info_t creation and argument binding
Stream FIFO enqueue and round-robin scheduling
Launch latency simulation
CTA-to-SM distribution

🔧 Configuration System

The OptionParser template registry — how 200+ knobs are parsed from files and command lines.

Recursive -config file parsing
Type-safe OptionRegistry<T> templates
Comment stripping and tokenization
core_config / gpgpu_sim_config derivation

Directory Map

Source tree for Core Infrastructure (src/ root)

src/
├── abstract_hardware_model.h    ← 48 KB — Hardware abstraction types, enums, base classes
├── abstract_hardware_model.cc   ← 45 KB — Memory coalescing, SIMT stack, kernel mgmt
├── gpgpusim_entrypoint.h        ← 2.7 KB — GPGPUsim_ctx simulation context
├── gpgpusim_entrypoint.cc       ← 15 KB — Sim threads, init, synchronization loop
├── stream_manager.h             ← 8.7 KB — CUDA stream / event type definitions
├── stream_manager.cc            ← 17 KB — Stream scheduling, operation dispatch
├── option_parser.h              ← 2.8 KB — C-style option parser API
├── option_parser.cc             ← 17 KB — Template-based option registry, file parsing
├── debug.h                      ← 3.0 KB — brk_pt class (breakpoints / watchpoints)
├── debug.cc                     ← 7.9 KB — Interactive PTX debugger loop
├── trace.h                      ← 3.6 KB — Tracing macros (DPRINTF, DTRACE)
├── trace.cc                     ← 2.3 KB — Trace stream initialization
├── trace_streams.tup            ←         X-macro enum for trace subsystems
├── statwrapper.h                ← 0.3 KB — Statistics wrapper API
├── statwrapper.cc               ← 1.1 KB — BookSim Stats class wrapper
└── tr1_hash_map.h               ← 2.2 KB — Portability shim for unordered_map

Configuration Knobs

Knob	Type	Default	Description
`-gpgpu_ptx_use_cuobjdump`	BOOL	1	Use cuobjdump for PTX/SASS extraction from binaries
`-gpgpu_ptx_convert_to_ptxplus`	BOOL	0	Convert SASS to PTXPlus and run PTXPlus
`-gpgpu_ptx_force_max_capability`	UINT32	0	Force maximum compute capability
`-gpgpu_ptx_inst_debug_to_file`	BOOL	0	Dump instruction debug output to file
`-gpgpu_ptx_inst_debug_file`	CSTR	inst_debug.txt	Debug output filename
`-gpgpu_experimental_lib_support`	BOOL	0	Extract code from CUDA libraries
`-checkpoint_option`	INT32	0	Enable checkpointing (0 = off)
`-checkpoint_kernel`	INT32	1	Checkpoint during which kernel invocation
`-resume_option`	INT32	0	Enable resume from checkpoint (0 = off)
`warp_size`	unsigned	32	Threads per warp
`gpgpu_coalesce_arch`	int	13	Coalescing architecture (13=GT200, 20=Fermi, 40=Maxwell+)
`num_shmem_bank`	unsigned	16	Number of shared memory banks per SM
`gpgpu_shmem_size`	unsigned	—	Shared memory size per SM (bytes)
`gpgpu_max_insn_issue_per_warp`	unsigned	—	Max instruction issues per warp per cycle

Code Examples

Concurrent Simulation Main Loop

The heart of the simulator — the thread that drives cycle-by-cycle execution (gpgpusim_entrypoint.cc):

void *gpgpu_sim_thread_concurrent(void *ctx_ptr) {
  gpgpu_context *ctx = (gpgpu_context *)ctx_ptr;
  do {
    // Spin until work appears in any stream
    while (ctx->the_gpgpusim->g_stream_manager->empty_protected() &&
           !ctx->the_gpgpusim->g_sim_done)
      ;
    ctx->the_gpgpusim->g_sim_active = true;
    ctx->the_gpgpusim->g_the_gpu->init();
    bool active = false, sim_cycles = false;
    do {
      // Dispatch next stream operation; break if GPU idle
      if (ctx->the_gpgpusim->g_stream_manager->operation(&sim_cycles) &&
          !ctx->the_gpgpusim->g_the_gpu->active())
        break;
      // Performance simulation: one GPU cycle
      if (ctx->the_gpgpusim->g_the_gpu->active()) {
        ctx->the_gpgpusim->g_the_gpu->cycle();
        sim_cycles = true;
        ctx->the_gpgpusim->g_the_gpu->deadlock_check();
      }
      active = ctx->the_gpgpusim->g_the_gpu->active() ||
               !(ctx->the_gpgpusim->g_stream_manager->empty_protected());
    } while (active && !ctx->the_gpgpusim->g_sim_done);
    ctx->the_gpgpusim->g_sim_active = false;
  } while (!ctx->the_gpgpusim->g_sim_done);
}

Memory Coalescing — Transaction Reduction

How per-thread addresses are coalesced into minimal memory transactions (abstract_hardware_model.cc):

void warp_inst_t::memory_coalescing_arch(bool is_write,
                                         mem_access_type access_type) {
  unsigned segment_size = /* 32, 64, or 128 based on arch & data_size */;
  for (unsigned subwarp = 0; subwarp < warp_parts; subwarp++) {
    std::map<new_addr_type, transaction_info> subwarp_transactions;
    // Step 1: gather per-thread addresses into segment-aligned transactions
    for (unsigned thread ...) {
      new_addr_type block_address =
          line_size_based_tag_func(addr, segment_size);
      transaction_info &info = subwarp_transactions[block_address];
      info.chunks.set(chunk);
      info.active.set(thread);
      info.bytes.set(idx + i);
    }
    // Step 2: reduce transaction sizes (128→64→32 if half unused)
    for (auto &t : subwarp_transactions)
      memory_coalescing_arch_reduce_and_send(
          is_write, access_type, t.second, t.first, segment_size);
  }
}

Stream Operation Dispatch

How the stream manager dispatches queued CUDA operations (stream_manager.cc):

bool stream_operation::do_operation(gpgpu_sim *gpu) {
  switch (m_type) {
    case stream_memcpy_host_to_device:
      gpu->memcpy_to_gpu(m_device_address_dst, m_host_address_src, m_cnt);
      m_stream->record_next_done();
      break;
    case stream_kernel_launch:
      if (gpu->can_start_kernel() && m_kernel->m_launch_latency == 0) {
        gpu->set_cache_config(m_kernel->name());
        gpu->launch(m_kernel);
      } else {
        if (m_kernel->m_launch_latency) m_kernel->m_launch_latency--;
        return false;  // not ready yet
      }
      break;
    case stream_event:
      m_event->update(gpu->gpu_tot_sim_cycle, wallclock);
      m_stream->record_next_done();
      break;
    case stream_wait_event:
      if (m_event->num_updates() >= m_cnt) m_stream->record_next_done();
      else return false;
      break;
  }
  m_done = true;
  return true;
}

Glossary

Term	Definition
CTA	Cooperative Thread Array — CUDA thread block; group of threads sharing shared memory and barriers.
Warp	Group of 32 threads executing in lockstep (SIMD). The fundamental scheduling unit.
SIMT	Single Instruction, Multiple Threads — NVIDIA's execution model allowing threads to diverge.
PDOM	Post-Dominator — the reconvergence point where divergent branches must rejoin.
PTX	Parallel Thread Execution — NVIDIA's virtual ISA, input to the simulator's functional model.
Coalescing	Combining per-thread memory requests in a warp into fewer, wider transactions.
Sector	A 32-byte portion of a cache line; used in Maxwell+ architectures.
CDP	CUDA Dynamic Parallelism — launching kernels from within GPU kernels.