API Reference#
code/simulator.h is a single header of 3791 lines that exposes 590 documented symbols. Every page below maps to one // ---- Section ---- banner in the header, so the documentation and the source keep the same shape.
Kind |
Count |
|---|---|
Functions |
490 |
Macros |
38 |
Structs |
22 |
Enums |
18 |
Classes |
8 |
Constructors |
8 |
Constants |
6 |
265 of them have at least one worked example in the repository’s test corpus, reachable through the View links.
Portable spelling and kernel launch#
Section |
Symbols |
Summary |
|---|---|---|
1 |
Real GPU path — hand everything to the actual CUDA runtime. |
|
12 |
blocks run one at a time in random order; a block’s threads run as fibers synchronised at __syncthreads(), visited in a fresh random order … |
|
7 |
A cooperative launch guarantees every block is resident at once, which is what makes this_grid().sync() a legal grid-wide barrier. |
|
5 |
Streams are NOT synchronous. |
Execution model and debugging#
Section |
Symbols |
Summary |
|---|---|---|
4 |
One instance each (single-threaded/cooperative, so the block scheduler simply rewrites these right before switching into a given thread’s … |
|
2 |
||
Manual block / thread execution order (for reproducible debugging) |
20 |
By default blocks and threads run in a RANDOM order (see the header overview) — great for finding order-dependent bugs, but you can’t … |
12 |
Install a callback the scheduler runs at the start of every block and after each __syncthreads() release. |
|
Sanitizer-lite: memcheck / initcheck / racecheck / synccheck |
17 |
A CPU-single-thread simulator sees the whole interleaving, so it can flag bugs a real GPU hides — approximating NVIDIA compute-sanitizer … |
Device-side language and intrinsics#
Section |
Symbols |
Summary |
|---|---|---|
37 |
||
29 |
Cache-hinted loads/stores are plain pointer access in the one-address-space sim; the hint (streaming / last-use / global) has no effect on … |
|
48 |
||
11 |
Single OS thread + cooperative fibers: fibers only ever switch at an explicit __syncthreads()/warp sync, never mid-expression, so a plain … |
|
15 |
Lanes 0..31 of a warp are consecutive linear thread ids (tid % 32 == laneId). |
|
30 |
A shim for the widely-used cooperative_groups API so kernels written against it run on the simulator: group sync maps to the block/warp … |
|
35 |
this_grid().sync() is a GRID-WIDE barrier. |
CUDA runtime API#
Section |
Symbols |
Summary |
|---|---|---|
2 |
||
6 |
A plain CPU heap, but tracked so the host/device split is real: device pointers come from cudaMalloc and are meant to travel through … |
|
10 |
The sim reports a single device that mimics the sm_75 target (compute 7.5, 2 GiB), and — since every operation is already synchronous and … |
|
5 |
Set |
|
Streams / events (deferred + replayed; see “Asynchronous stream queues”) |
10 |
A stream owns a queue of deferred ops; a synchronisation point replays them. |
Unified-memory hints (advisory; no-ops in the single-address-space sim) |
7 |
The sim keeps host and device in one address space, so prefetching and advice change nothing — they exist only so unified-memory code … |
14 |
||
8 |
||
2 |
||
7 |
||
Zero-copy / pinned mapping (one address space: device ptr == host ptr) |
3 |
|
2 |
||
2 |
||
12 |
||
2 |
Portable spelling: a kernel launching a child writes LAUNCH_STREAM(child, grid, block, cudaStreamFireAndForget, args…) which is the real … |
|
3 |
||
Unified-memory stream attachment (advisory in one address space) |
4 |
|
Portable UM prefetch/advise helper (CUDA-13 signature under nvcc) |
2 |
A repo portable-spelling helper (like LAUNCH_STREAM): call uvmPrefetch/uvmAdvise in a test and it works on both toolchains. |
Textures and surfaces#
Section |
Symbols |
Summary |
|---|---|---|
10 |
Textures and surfaces (sim only) A CPU model of the bindless texture API: cudaArray-backed storage, a channel descriptor, and texture … |
|
9 |
||
15 |
||
7 |
Integer fetch into a linear resource: no addressing, no filtering (real tex1Dfetch is the same). |
|
Surface reads / writes (byte offsets on x, like real surfaces) |
5 |
Library subsets#
Section |
Symbols |
Summary |
|---|---|---|
20 |
||
47 |
One address space means a device_vector IS a host vector; the algorithms are the <algorithm>/<numeric> ones. |
|
41 |
fp16 / bf16 surrogate (sim only) — stands in for <cuda_fp16.h> / <cuda_bf16.h> Values are STORED as float and rounded through the real … |
|
12 |
CUDA graphs (sim only) — stream capture replayed as a recorded op list Capture records the deferred ops a stream would have queued; … |
|
11 |
cuRAND subset (sim only) — a deterministic CPU RNG curand_init/curand/curand_uniform/curand_normal on a small counter-based generator … |
|
7 |
cuda::barrier / async copies (sim only) The sim copies synchronously and has no memory pipeline, so an async copy is a plain memcpy and a … |
|
19 |
The multi-stage async-copy API. The sim copies synchronously and has no memory pipeline, so acquire/commit/wait/release are bookkeeping … |
Other#
Section |
Symbols |
Summary |
|---|---|---|
1 |