CUDA Tools & Further Reading

nvcc — CUDA compiler driver current

compile · orchestrates the build; produces the fatbinary

Splits host/device code, runs the device toolchain (cicc → ptxas → fatbinary), hands host code to your system C++ compiler, and embeds device code in the host object.

nvcc -arch=sm_120 app.cu -o app

NVCC guide · Understanding PTX (blog)

ptxas — PTX optimizing assembler current

compile · PTX → architecture-specific SASS; where register allocation happens

Invoked automatically by nvcc. Its most-used "tool" role is reporting resource pressure: pass -v to see per-kernel register / shared / local / constant usage — the number that governs occupancy (Lesson 2). Cap registers with -maxrregcount=N.

nvcc -Xptxas=-v -arch=sm_120 app.cu   # or the modern: nvcc -res-usage ...

nvcc compiler switches

nvrtc — runtime compilation library

compile · compiles CUDA C++ source strings to PTX/cubin at runtime

A library (libnvrtc), not a CLI — you link it and call its C API to generate or specialize kernels dynamically without shipping nvcc. The JIT engine under PyTorch, CuPy, OptiX, and similar codegen-heavy stacks.

nvrtcCompileProgram(prog, nOpts, opts);  nvrtcGetPTX(prog, ptx);

NVRTC docs · NVIDIA/jitify helper

nvlink · fatbinary — device linker & fatbin packer current

compile · separate device-code linking; bundling per-arch cubins

nvlink links relocatable device objects when you use -rdc=true (calling __device__ functions across translation units). fatbinary packs multiple per-arch cubins/PTX into the fatbin nvcc embeds. Both are normally internal to nvcc. Runtime counterpart: the nvFatbin library (CUDA 12.4+) builds fatbins programmatically.

nvcc -rdc=true a.cu b.cu -o app        # nvcc calls nvlink internally

NVCC guide (separate compilation)

cuobjdump — dump SASS/PTX/ELF current

inspect · extracts device code from host binaries and cubins

Finds the embedded fatbins for you, so it works on your final executable — no need to isolate a cubin first.

cuobjdump -sass ./app        # disassembled SASS
cuobjdump --list-elf ./app   # which architectures are bundled

CUDA Binary Utilities

nvdisasm — cubin disassembler + CFG current

inspect · richer disassembly of a standalone cubin

More detail than cuobjdump — control-flow graphs, register liveness, source-line correlation — but reads only cubin files, not host binaries. Pair with -lineinfo for source lines.

nvdisasm -cfg kernel.cubin   # emits a DOT control-flow graph

CUDA Binary Utilities

cu++filt · nvprune — demangler & fatbin pruner current

inspect · readable kernel names; shrink device binaries

cu++filt is the CUDA analog of GNU c++filt — turns a mangled symbol like _Z1fIiEbl back into a readable signature. nvprune strips a library down to only the architectures you need (common on the big static math libs).

cu++filt _Z1fIiEbl
nvprune -arch sm_120 libcublas_static.a -o lib120.a

CUDA Binary Utilities

GNU binutils on CUDA binaries — readelf · nm · objdump current

inspect · a CUDA host executable is just an ELF file

Device code lives in the .nv_fatbin section; a hidden pre-main constructor calls __cudaRegisterFatBinary/__cudaRegisterFunction to map host stubs to device kernels. Standard tools reveal all of it.

readelf -S ./app | grep nv_fatbin      # locate the embedded device code
nm ./app | grep __cudaRegister         # the registration stubs

"What happens when you run a GPU kernel" (Fergus Finn)

CuAssembler — SASS assembler third-party

inspect/edit · NVIDIA ships no public SASS assembler — this fills the gap

The de-facto community tool for round-tripping nvdisasm output back into a cubin, so you can hand-tune machine code or microbenchmark specific SASS. Powerful but fragile across driver/arch versions — verify support for your GPU (Blackwell/sm_120) before relying on it.

cloudcores/CuAssembler

nsys — Nsight Systems current

tracing (CUPTI) · whole-app CPU+GPU timeline, low overhead

Shows API calls, kernel launches, memcpies, overlap, gaps, and NVTX ranges on one timeline, with the host launch linked to its GPU kernel. The first profiler you reach for — it finds what is slow and where the gaps are.

nsys profile --stats=true --trace=cuda,nvtx,osrt -o report ./app

Nsight Systems User Guide

ncu — Nsight Compute current

hardware counters + PC sampling · per-kernel deep dive, high overhead (kernel replay)

Occupancy, memory-vs-compute roofline, warp stall reasons, cache/sector efficiency, per-source-line counters. Replays each kernel to read the limited HW counters, so point it at one kernel. Build with -lineinfo for source correlation.

ncu --set full -k mykernel -c 1 -o profile ./app

Kernel Profiling Guide · CLI reference

Nsight Graphics — current, graphics only

graphics trace · D3D / Vulkan / OpenGL frame & shader profiling

Listed to rule it out: it's for graphics pipelines, not CUDA-C compute. For compute, use Nsight Systems/Compute. Relevant only if your workload is DirectCompute / graphics-queue shaders.

Nsight Graphics

PyTorch Profiler — torch.profiler framework

tracing (Kineto → CUPTI) · in-framework kernel/op timeline for PyTorch

Not part of your general-CUDA path, but ubiquitous in ML: the first tool for model profiling before dropping to nsys/ncu. Exports Chrome/TensorBoard traces.

torch.profiler docs

compute-sanitizer — correctness suite current

binary instrumentation · patches device code and checks every access

Four checkers: memcheck (out-of-bounds / misaligned), racecheck (shared races — Lesson 6), initcheck (uninitialized global reads), synccheck (illegal __syncthreads). Finds the bug a lucky run hides, and pins a sticky fault (Lesson 8) to its real kernel + address. Replaced the removed cuda-memcheck.

compute-sanitizer --tool memcheck ./app     # or racecheck / initcheck / synccheck

Compute Sanitizer manual · Hunt Bugs with Compute Sanitizer (blog)

cuda-gdb — device debugger current

debugger back-end · halts warps; reads state (needs -G)

GDB extended for device code: breakpoints inside kernels, single-step warps, inspect per-thread registers and memory, switch focus across grids/blocks/threads. Reads state; no counters.

cuda-gdb --args ./app   # break mykernel · run · cuda block 0 thread 0 · info cuda warps

CUDA-GDB manual

Nsight Visual Studio Code Edition current

editor integration · front-end to cuda-gdb + CUDA IntelliSense

Free VS Code extension for editor-integrated GPU-thread debugging, including remote / WSL / container dev — the modern replacement for the retired Nsight Eclipse Edition.

Nsight VS Code Edition docs

CUPTI — CUDA Profiling Tools Interface library

tracing + counters · the layer every CUDA profiler sits on

Provides tracing (Activity + Callback APIs) and counter profiling (PC Sampling, PM Sampling, SASS metrics). nsys, Nsight Compute, PyTorch/Kineto, HPCToolkit all build on it. You use it directly only if you're writing a profiler.

CUPTI docs

NVTX — NVIDIA Tools Extension library

annotation · produces no data — labels tool timelines with your semantics

Wrap code in named ranges so Nsight Systems labels its timeline "dataloader", "forward", "solve" instead of raw API calls. Header-only C, with C++ and Python bindings. Tiny effort, huge readability win when profiling.

nvtx3::scoped_range r{"forward_pass"};   // Python: with nvtx.annotate("forward"):

NVIDIA/NVTX · Custom timelines with NVTX (blog)

Occupancy API — launch-config helper library

runtime query · picks block sizes for max occupancy

The programmatic replacement for the deprecated CUDA_Occupancy_Calculator.xls: ask the runtime for the block size that maximizes resident warps for your kernel. Nsight Compute also has an interactive occupancy section.

cudaOccupancyMaxPotentialBlockSize(&minGrid, &blockSize, myKernel, 0, 0);

Occupancy API (blog)

nvidia-smi + NVML current

driver telemetry · util %, memory, power, temp, clocks, processes

NVML is the C library; nvidia-smi the CLI on top. Sampled management data, not perf counters — and "GPU-Util" is time-busy, not saturation (one active SM reads 100%). Nearly every third-party monitor reads NVML.

nvidia-smi --query-gpu=utilization.gpu,memory.used,power.draw --format=csv -l 1

nvidia-smi docs · NVML API

DCGM · dcgmi + DCGM-Exporter current, datacenter

telemetry + sampled counters · cluster-scale health & metrics

Data Center GPU Manager: continuous low-overhead profiling metrics (SM activity, tensor/FP pipe use, DRAM/PCIe/NVLink throughput) without kernel replay. dcgmi is the CLI; DCGM-Exporter feeds Prometheus → Grafana, standard for Kubernetes GPU fleets. Overkill for a single dev box, but the answer to "how do datacenters monitor GPUs."

dcgmi dmon -e 203,252        # GPU util, mem util

DCGM-Exporter docs · Monitoring GPUs in Kubernetes (blog)

nvtop · nvitop · gpustat — live dashboards third-party

driver telemetry (NVML) · friendlier than nvidia-smi for day-to-day watching

nvtop: htop-style curses UI with per-process view and interactive kill. nvitop: the most feature-rich, plus a Python API — the common "best daily pick." gpustat: lightweight one-liner for scripts and shared boxes. All read NVML, so telemetry only — not for tuning.

pip install nvitop && nvitop      # or: nvtop  /  gpustat -i 1

nvitop · nvtop · gpustat