原书索引（Index）

Symbols

#pragma unroll, 326, 333

unrolling too many iterations, 333

100-trillion parameter models, 11

NVIDIA supercomputers in a rack, 13-15

reason for discussing, 11

scaling toward, 939-943

3FS (see Fire-Flyer File System (3FS; Deep‐Seek))

<<< >>> (chevron syntax)

CUDA-specific launch parameters, 205

how many threads, 203

__global__, 203

__launch_bounds__, 229-231, 310

occupancy optimization, 311

__ldg(), 273

__shared__ regions, 405

__shfl_sync, 271-273

__syncthreads(), 396

~680-billion parameter models, 9-11

Aachieved occupancy, 222, 297

inspecting, 297-305, 482

kernel compute throughput versus GPU

FLOPS, 299

kernel memory throughput versus

HBM, 299

parallelism to increase, 222

activation checkpointing, 572

FSDP automatic checkpointing, 574-577

adaptive batching, 893-899

adaptive routing, 695, 911

admission control (see early rejection)

AI agents that self-improve, 931-934

AI factories via NVL72 building blocks, 43

AI Futures Project, 931-934

AI optimizing AI systems, 936-939

AI supercomputers in a rack (NVIDIA), 13-15

31-33

AI systems performance engineers, 3

GPUs near 100% utilized, 47

NVLink usage, 47

scaling distributed training and inference, 5

transparency and reproducibility, 6

air cooling versus liquid cooling, 45-46

all-to-all communication, 692

about, 23

cache coherency, 24

CPU + GPU superchip, 23-33

GPU Tensor Cores and Transformer

Engine, 29-31

many GPUs as one, 33-43

multi-GPU programming, 38-39

multirack and storage communication, 41

NVIDIA Blackwell dual-die GPU, 26-28

NVIDIA Grace CPU, 26

NVL72, 33

NVLink and NVSwitch, 34-37

SHARP for in-network aggregations, 40-41

streaming multiprocessors, threads, warps,

benchmarking and profiling, 4

cross-team collaboration, 6

future capabilities of AI, 940

generating returns well above cost, 18

managing resources efficiently, 5

overview, 3

performance monitoring, 47

mixture-of-experts models, 693

AlphaTensor (DeepMind), 924

Amazon FSx for Lustre, 174

performance improvement verification, 175

Amazon S3, 175

AMP (see automatic mixed precision (AMP;PyTorch))

AOT Autograd, 612

FX Graph, 609, 612

PyTorch compiler pipeline, 606

AOTInductor (PyTorch), 615

application performance management (APM),

application-level inference optimizations,

740-757debouncing and request coalescing, 755

model cascading and tiered deployment,

750-752prefix caching, 743-749

prompt cleansing, 741

prompt compression, 740

streaming responses, 752-755

token output limits and timeouts, 756

Apptainer, 91

architecture of GPUs, 191-203

maximum threads per SM and block, 229

memory hierarchy, 212-218

parallelism, 191

cuTile, 59

Data Parallel versus Distributed Data

Parallel, 132-137

fully sharded data parallel, 133

PTX, 365

Special Function Unit, 194

streaming multiprocessors, 192

warps, 32, 192-194

thread block clusters, 196

thread blocks, 195-199

thread block size, 200

threads, 195-199

threads-per-block and blocks-per-grid,

200-202arithmetic intensity, 233, 337

batching increasing, 348

CUTLASS for optimal, 360-364

data reuse increasing, 339

distributed shared memory increasing, 406

increasing, 339, 344-347

kernel fusion increasing, 344-347

|

multilevel tiling, 340

software prefetching, 341

asynchronous execution with streams, 106

asynchronous memory transfers with CUT‐LASS, 360

AsyncTP (asynchronous tensor parallel), 590

aten operations

FX Graph, 609

PrimTorch IR decomposing, 613

PyTorch compiler pipeline, 606

atomic work queues, 482-489

dynamic work allocation, 483

CUDA Graphs device initiated, 501

attention

AI-produced CUDA kernel for attention,

automated performance tests, 4

AI optimizing AI systems, 936-939

automatic mixed precision (AMP; PyTorch),

355-356

AutoParallel (PyTorch), 590

autotuning

AI-assisted autotuning, 936-939

static kernels transformed into adaptive, 858

TorchInductor, 544, 617-620

transformer self-attention and MLP paths,

B

B200 (NVIDIA Blackwell), 1, 27, 201, 217

low versus high arithmetic intensity, 338

prefill and decode phase differences, 893

pruning increasing, 349

PyTorch, 350

recomputation of cheap expressionsincreasing, 350

Roofline performance model, 337-338, 521

lower precision with Tensor Cores, 351

software prefetching, 341

tiling increasing, 260, 339

atomic counters, 483

Nsight Compute, 483-484

atomic queues, 485-489

926-929bottleneck fixed by FlashAttention, 15

context parallelism communicating acrosspartitions, 679

PyTorch optimized attention mechanisms,

853-858

Triton kernels, 643

B300 (NVIDIA Blackwell Ultra), 12, 49

bandwidthTest of CPU-GPU memory datatransfer, 70

batchingadaptive batching, 893-899

arithmetic intensity increased, 348

comparing static, dynamic, continuous, 721

continuous batching, 720

dynamic batching, 718-720

static batching, 718

benchmarks

AI systems performance engineer, 4

bandwidthTest, 70

CUDA Graphs impact, 494

data pipelines, 183

I/O continuous profiling and tuning,

184-188dynamic parallelism, 511

GPU clock and, 84

kernels generated by TorchInductor, 553

MLPerf, 7

performance benchmarking, 592-598

PyTorch performance heads-up display,

594-595best-fit with coalescing (BFC), 880

BF16, 356

Blackwell GPU (NVIDIA) 100-trillion parameter models, 12

architecture

B200, 201, 217

B300, 49

FP32 and INT32 pipelines merged, 328

Blackwell B200 GPU, 27, 201, 217

appearing as one GPU, 27

chiplet approach, 27

Data Center GPU Manager, 47

I/O monitoring, 177

inference monitoring, 706

dual-die GPU, 26-28

appearing as one GPU, 27

chiplet approach, 27

remote memory access latency, 62

FP32 80 TFLOP theoretical, 261

global memory access, 247

ILP configuration, 332

many GPUs as one, 33-43

multi-GPU programming, 38-39

(see also multiple GPUs)

31-33

TFLOPS, 255

Blackwell HBM3e, 27

Blackwell NVL72 (see NVL72)

Blackwell Ultra B300 100-trillion parameter models, 12

NVL72 drop-in upgrade, 49

blocksPerGrid, 203, 205

launch parameter, 207-209

BlueField-3 DPU (NVIDIA), 41

book supplemental material online, xxiiibook web page, xxivbottlenecksattention mechanisms optimized by

PyTorch, 553

communication versus compute, 188

experts overloaded in MoE, 676

GPU bottlenecks, 288-293

multirack and storage communication,

SHARP for in-network aggregations,

40-41memory, 25

NVL72, 33

NVLink and NVSwitch, 34-37

roofline model, 233-237

shared memory of 32 banks, 269

streaming multiprocessors, threads, warps,

about, 288

data pipeline profiling and tuning, 289

iteratively profiling, 301-303

kernel optimization, 303-305

Nsight Compute and Roofline analysis,

Nsight Systems timeline view, 288

PyTorch profiler via Kineto, 291-293

warp stall reasons, 293-297

identifying bottlenecks, 4

I/O continuous profiling and tuning,

184-188

NFS servers, 174

tuning parameters, 175

NVIDIA mechanical sympathy, 16

profilingexample MoE model, 531-543

Linux perf on CPU and GPU, 539-543

Nsight Systems and NVTX Timelines,

535-537

|

NVTX markers in profiling tools,

528-531tools for profiling, 528-531

shared filesystems, 174

SHARP and adaptive routing minimizing,

bucketing (PyTorch), 107

butterfly schedules, 692

C

C++

CUDA C++ Pipeline API, 279-282

hardware-software codesign, 282

libraries in CUDA Toolkit, 59

C-states and CPU frequency, 75

cache coherency, 24

NVLink-C2C, 24, 38

caching mechanismscaching CUDA kernels, 891

metrics for cache hits and misses, 703

prefix caching, 743-749

prewarming caches, 889-893

read-only data caches, 273-278

const __restrict__, 273, 276, 277

pitfall of not stating read-only, 273

vLLM LMCache, 703, 815

callback function between GPU and CPU, 447

CC (compute capability) and CUDA Toolkitversion, 58

CDU (Coolant Distribution Unit), 45

obsolete

CFS (Completely Fair Scheduler; Linux), 73

ChatGPT (OpenAI) stateful conversation, 744

checklist for AI systems performanceadvanced tuning strategies, 978-979

algorithmic tricks, 978-979

arithmetic optimizations, 975-977

cost optimization, 947-950

CUDA tuning optimizations, 968-973

data processing pipelines, 964-966

GPU programming, 968-973

GPU resource management and scheduling,

959-961

I/O optimization, 961-963

inference and serving, 982-984

multinode, 984-985

optimizations

|

953-954checkpoints written to disk, 75

chevron syntax (<<< >>>)

CUDA-specific launch parameters, 205

how many threads, 203

child-kernel launch limit (CUDA), 507

chiplet approach of dual-die GPUs, 27

appearing as one GPU, 27

Chrome tracing with Perfetto viewer, 528, 531

chunked prefill, 723, 893-899

CI (see continuous integration (CI))cloud storage caches, 174

performance improvement verification, 175

cluster pools, 765-780

cluster scheduler, 47

cluster size maximum, portable, 406

co-packaged optics (CPO), 43

cold start training strategy, 10

CollNet communication algorithm, 130

communication (see distributed networkingcommunication)compilerscommonly used compilers, 56

compiling versus writing custom kernels,

optimizationsnetwork optimization, 980-982

operating system and driver optimizations,

957-959performance profiling, debugging, andmonitoring, 966-968

power management, 986-988

reduced/mixed precision, 975-977

reproducibility and documentation, 950-952

system architecture and hardware planning,

952-953thermal management, 986-988

unified CPU-GPU “Superchip” architecture,

float8 not provided by CUDA, 248

__global__, 203

instruction-level parallelism exposed,

325-334compiler hints, 326

__launch_bounds__, 229-231, 310

occupancy optimization, 311

nvcc CUDA compiler, 57

(see also nvcc CUDA compiler)predication, 317

PTX, 203, 616

generations

optimizationinline PTX code, 365-369

JIT compiling, 59

streaming assembler, 365-369

generations

TorchInductor, 616

Triton, 640

PyTorch compiler, 60, 543

compilesmart compilers, 934-936

Triton, 60

(see also Triton (OpenAI))vectorized memory access, 247-254

vectorized memory operations, 248

obsolete

Completely Fair Scheduler (CFS; Linux), 73

compressing data, 176

compute capability (CC) and CUDA Toolkitversion, 58

Compute Sanitizer (NVIDIA), 231

compute-sanitizer CLI, 231

NVIDIA Tools Extension support, 231

concurrency and synchronization, 561

four primary tools, 232

inference correctness, 715

compute-bound GPU stalls, 301-303

kernel optimization, 303-305

compute-sanitizer CLI, 231

NVIDIA Tools Extension support, 231

concurrency with CUDA streams, 555-561

compute overlapping with data transfers,

433-435concurrency, 556-558

launching on intended stream, 561

mixture-of-experts models, 560

NVIDIA Compute Sanitizer, 561

profiling when introducing concurrency,

stream synchronization with events,

558-560const __restrict__, 273, 276, 277

“Contact” (Phish), xxv

containerd with NVIDIA Container Toolkit, 89

containers, 88

NVIDIA container runtime injectinglibraries, 89

NVIDIA Container Toolkit and CUDAcompatibility, 88

reducing image size, 91

context parallelism (CP), 670, 678

Context Parallel (PyTorch), 554

hybrid parallelism, 681

continuous batching, 720

continuous integration (CI)

Computer Sanitizer integrated, 232

performance benchmarking, 592-598

594-595continuous prewarming, 889-893

continuous scheduling, 721

Coolant Distribution Unit (CDU), 45

cooling via air versus liquid, 45-46

cooperative groups (CGs), 395-401

when to combine, 401

thread block clusters contrasted with, 402

Cooperative Groups API (CUDA), 272, 396

coordinating thread block clusters, 408-410

cooperative thread array (CTA), 195

(see also thread blocks)runtime operation binding, 810

correctness issues in inference, 715-717

costs

generatingcost optimization overview, 947-950

NVIDIA Container Toolkit, 88

CUDA compatibility, 88

GPU performance, 88

runtime optimizations for GPUs, 88-91

container overlay filesystem overhead,

captured

PyTorch performance heads-up display,

Cooperative Groups API, 272, 396

coordinating thread block clusters,

408-410described, 402

launching a kernel in cooperative mode, 396

grid size for launch, 397

reserving all SMs in the grid, 401

persistent kernels via grid sync, 398-401

|

model parameter-count explosion, 1

100-trillion parameter models, 11

trainingcold start strategy, 10

DeepSeek-R1, 2

DeepSeek-V3, 9

Google Gemini Ultra, 2

OpenAI GPT-4, 2

CP (see context parallelism (CP))

CPO (co-packaged optics), 43

CPU + GPU superchip, 23-33

NVIDIA Blackwell dual-die GPU, 26-28

appearing as one GPU, 27

Transformer

NVIDIA Grace CPU, 26

NVLink-C2C, 24

streaming multiprocessors, threads, warps,

31-33

CPU and OS configuration, 61-77

about, 61

CPU frequency and C-states, 75

filesystem caching and write-back, 75

NUMA awareness and CPU pinning, 62-69

CPU + GPU superchip CPU-to-GPUdata transfers, 69

CPU pinning, 62-69

NUMA nodes, 62

NUMA-friendly memory allocation andpinning, 70-72

efficiency of pinned memory, 70

max locked memory setting, 72

OS limit on pinned memory, 71

pinned memory for data loaders, 71

scheduler and interrupt affinity, 73

irqbalance daemon, 74

transparent hugepages, 72

tuning host CPU memory allocator, 76

virtual memory and swapping, 74

CPU DRAM

CPU pinning, 62-69

bandwidth

NIXL GPU-to-GPU direct transfer, 152

CUDA programming flow, 191

CUDA Unified Memory, 86

programming CUDA, 218-221

Unified CPU-GPU Memory, 24

GDS, 70, 167-172, 573

|

superchips’ high-speed extension of

GPU memory, 128

NIXL offloading KV cache to, 149, 157-159

parameters offloaded to, 573

peer-to-peer DMA avoiding, 581

swapping, 87

CPU host and GPU device communication via

CUDA events, 447

callback function, 447

CPython Frame Evaluation API, 607

CTA (cooperative thread array), 195

(see also thread blocks)

CTRL approach to prompt trimming, 743

CUBIN, 58

cuBLAS

CUDA Toolkit, 58, 61

PyTorch delegating tasks to, 61

CUDAchild-kernel launch limit, 507

compatibility across GPU generations, 58,

host CPU memory allocator tuned, 76

large slower extension of GPU memory, 25

CUDA Graph launched, 492-496

float8 not provided, 248

kernels created in Python, 56

kernels executing asynchronously, 206

global fault flag for illegal operations,

programmingasynchronous memory allocation,

210-212

CUDA streams, 211

flow between CPU and GPU, 191

__global__, 203

GPU memory hierarchy, 212-218

high occupancy and GPU utilization,

221-229highest-level and most-recent APIsavailable, 283

kernel errors, 206

kernel inputs in 1D, 210

kernel inputs in 2D and 3D, 209

kernels for parallel work, 203-206

launch parameters, 207-209

maximum threadsPerBlock compile timeparameter, 229-231

memory pools, 210-212

minimum thread blocks resident on each

SM, 229-231

NVIDIA Compute Sanitizer, 231

occupancy tuning with

__launch_bounds__, 229-231

roofline model, 233-237

Unified Memory, 218-221

Python-based frameworks, 60

Python-centric libraries, 638

Unified Memory, 86

programming CUDA, 218-221

Unified CPU-GPU Memory, 24

versions of container libraries and hostdriver, 88

NVIDIA Container Toolkit, 88

CUDA C++ Pipeline API, 279-282

hardware-software codesign, 282

CUDA compiler (see nvcc CUDA compiler)

CUDA eventscommunication from GPU device to CPUhost, 447

CUDA stream synchronization, 446,

558-560cross-stream synchronization, 448

multi-GPU overlap of compute and data,

463-469profiling with, 449

CUDA Graphs, 489-506

about CUDA Graphs, 311, 489

benchmarking, 494

best practices, 567-568

capturing a CUDA Graph, 492-496, 562-565

capturing only a portion of pipeline, 492

device-initiated launch, 497-501

dynamic graph update, 496

warm-up pass necessary, 494

compiler modes that trigger, 618

static shapes required, 618, 622

conditional graph nodes, 501-506

nested, 505

CUDA streams, 467

capturing a CUDA Graph, 492-496

PyTorch, 468

device-initiated launch, 497-501

dependency management, 500

example of use, 500

persistent scheduling, 501

dynamic graph update, 496

inference engines, 490

CUDA Managed Memory (see CUDA Unified

Memory)

CUDA memory allocator, 570

CUDA Memory Inspector, 570

CUDA Ninjas, 926

CUDA Occupancy API, 309, 310, 312

cooperative grids, 400

CUDA Pipeline APIdouble-buffering, 375-380

TMA plus, 282

warp specialization, 385-390

PyTorch, 390

CUDA Profiling Tools Interface (CUPTI), 291

CUDA Program Counter Sampling (Nsight

Compute), 711

CUDA Python library (NVIDIA), 56

CUDA Runtime API callback function, 447

CUDA Runtime library, 57

CUDA streams, 429

kernel launch overhead reduced, 562-569

kernels that use NVSHMEM, 516

launching a CUDA Graph, 492-496

device-initiated launch, 497-501

memory management, 491

memory pools, 566

multiple-GPU collectives captured, 517-520

pitfalls, 495

prewarming, 889-893

PyTorch, 490, 493

CUDA Graph Trees, 568

static memory pools, 491

replaying the graph, 565-566

any stream for replay, 564

TorchInductor, 618

when to use dynamic parallelism instead,

asynchronous execution, 106

compute overlapping with data transfers,

433-435concurrency, 556-558

multiple GPUs, 463-469

concurrency, 555-561

launching on intended stream, 561

profiling when introducing concurrency,

creating in PyTorch, 555

profiling when adding streams, 558

CUDA Graph captured, 492-496

default streams

|

default versus explicit streams, 443

legacy default stream, 441

per-thread default streams, 442

using default streams, 444-446

kernel execution overlapping with, 429-433

launching five kernels on two streams,

429-431warp specialization replaced with CUDAstreams, 431-433

mixture-of-experts models, 560

multi-GPU overlap of compute and data,

463-469multiple streams per GPU, 464

Programmatic Dependent Launch, 469-472

stream-ordered memory allocator, 435-439

multiple GPUs, 463-469

using, 439-441

synchronization with events and callbacks,

446, 558-560

cross-stream synchronization, 448

warp specialization, 449-456

thread block clusters, 456

CUDA Toolkitnvcc CUDA compiler, 57

NVIDIA Compute Sanitizer, 231

NVIDIA software stack, 57

optimized libraries included, 58

C++ and Python libraries, 59

CUDA Unified Memory, 86

programming CUDA, 218-221

Unified CPU-GPU Memory, 24

cudaMemPool, 705

cudart (CUDA Runtime library), 57

CUDA_DEVICE_ORDER, 143

CUDA_VISIBLE_DEVICES, 143

cuDNN, 58

CUPTI (CUDA Profiling Tools Interface), 291

CuPy, 59, 897

cuPyNumeric, 59

current used at typical data center, 44

liquid cooling, 46

custom kernels with OpenAI Triton, 637-658

CUDA C++ custom kernels, 643

registering with PyTorch, 641-643

tuning kernel launch parameters, 643

cuTile, 59

CUDA kernels in Python, 56

CUDA Toolkit, 59

CUTLASS (NVIDIA)

|

D

DALI (Data Loading Library; NVIDIA), 182,

DASH framework, 918

Data Center GPU Manager (DCGM; NVIDIA),

data centerscurrent used at typical data center, 44

liquid cooling, 46

data compression, 176

data input pipeline optimization, 583-585

data parallelism (DP), 670, 677

versus TP, PP, and hybrid, 843-847

data pipelines (see pipelines)

Data Processing Unit (DPU) in NVL72, 41

data reuse via shared memory, 255-264

arithmetic intensity increased by, 339

DataLoader (PyTorch)

CPU affinity for each worker process, 64-69

data input pipeline optimization, 583-585

data loading and preprocessing, 178-181

multiple workers, 167, 178, 181, 583

persistent workers, 181

pinned memory flag, 71, 434, 583

prefetch factor, 584

datasets for training via NeMo Curator, 183

DCGM (see Data Center GPU Manager

(DCGM; NVIDIA))

DDP (see Distributed Data Parallel (DDP;PyTorch))

performance

CUDA kernels in Python, 56

CUDA Toolkit, 59

thread block pairs, 413

TMEM, 352

warp specialization, 360

I/O monitoring, 177

inference monitoring, 706

monitoring straggler nodes or processes,

definition of data parallelism, 132, 677

Paralleldefinition of Distributed Data Parallel,

model weights and data split over GPUs,

debouncing, 755

debuggingcompiler performance issues, 552

DWARF, 540

graph breaks, 628, 636

monitoring graph recompilations, 609,

inference correctness issues, 715-717

NCCL, 145

debugging enabled, 145, 715

NCCL logs, 147, 715

NCCL_DEBUG, 129, 141, 145, 715

test suite, 715

numerical correctness and accuracy,

624-627

NVIDIA Compute Sanitizer, 231

PyTorch compiler, 623

Debugging With Attributed Record Formats

(DWARF), 540

decodeabout decode phase, 759

constrained decoding, 689

disaggregated from prefillabout disaggregation, 663, 665, 760

architecture, 664

batch size optimizations, 893

cluster pools, 765-780

deploying with Kubernetes, 668-670, 714

latency and throughput, 666

load balancing, 834-841

prefill to decode data transfers, 667,

812-817scaling, 665, 795

separate prefill and decode stages,

150-152, 664

disaggregated routing policies, 780-795

capacity-aware routing, 786-788

configurationexample dynamic routing policy in code,

latency-aware routing, 788-791

multipath inference (racing), 791

QoS and early rejection policies, 792-795

routing factors, 781-784

speculative decoding, 791

hardware and parallelism, 820-826

KV cache transfer from prefill, 812-817

connector and data path design, 817-820

advantages of disaggregation, 762-765

cluster pools, 765-780

routing and scheduling policies, 780-795

scalability, 665, 795

DeepEP (Deep Experts Parallelism; DeepSeek),

inline PTX code, 369

DeepGEMM (DeepSeek), 6

DeepMind AlphaTensor, 924

DeepSeek Fire-Flyer File System (3FS), 6,

172-174

DeepSeek Multi-Head Latent Attention (MLA),

DeepSeek-R1 (NVIDIA)automated kernel optimizations, 926-929

chain-of-thought reasoning, 10

custom kernels, advanced optimization,less-capable GPUs, 1

efficiency gains, 1, 2

MoE efficiency, 12

training costs, 2

DeepSeek-V3, 9-11

MoE efficiency, 12

DeepSeek.AIexport restrictions, 1, 9-11

open source GitHub repositories, 6

Open-Source Week, 6, 7

deploymentdisaggregated inference system, 668-670

inference optimizations, 714

KV cache size, 812

zero-copy GPU-to-GPU, 813-817

LLM decode phase as memory bound, 228

optimized decode kernels, 799-805

FlashMLA, 800

FlexDecoding, 802-805

ThunderMLA, 801

scale-out of workers, 891

speculative decoding, 665, 681-689

combining techniques, 688

EAGLE algorithm, 684

interleaving decode steps, 688

Medusa multiple heads, 682

Medusa multiple heads multitoken,

686-688single-model self-speculative decoding,

two-model draft-based, 682-685

why disaggregation, 761-796

|

NVL72 as preintegrated rack appliance, 42

Kubernetes, 42

NVIDIA Base Command Manager, 42

Managementtiered deployment, 752

direct memory access (DMA)

NVIDIA GDS, 168

O_DIRECT, 168

GDS, 170

peer-to-peer DMA, 581

disaggregated prefill and decodeabout, 663, 665, 759

architecture, 664

separate prefill and decode, 150-152

separate prefill and decode stages, 664

speculative decoding, 665

batch size optimizations, 893

deploying with Kubernetes, 668-670, 714

hardware and parallelism, 820-826

KV cache transfer from prefill to decode,

812-817latency and throughput, 666

load balancing, 834

NIXL, 148-161

about NIXL, 148-150

asynchronous API with callbacks,

153-157

Dynamo throughput improved, 159

intelligent interconnect routing, 152, 667

KV cache offloading with NIXL, 157-159

prefill to decode data transfers, 667

separate prefill and decode inferencestages, 150-152

prefill to decode data transfers, 667

connector and data path design, 817-820

KV cache, 812-817

KV cache size, 812

zero-copy GPU-to-GPU, 813-817

prefill-decode interference, 665

routing and scheduling policies, 780-795

capacity-aware routing, 786-788

configurationexample dynamic routing policy in code,

latency-aware routing, 788-791

multipath inference (racing), 791

QoS and early rejection policies, 792-795

|

cluster pools, 765-780

routing and scheduling policies, 780-795

scalability, 795

Distributed Data Parallel (DDP; PyTorch)bucketing, 107

Data Parallel versus, 132-137

scaling via DDP with torch.compile, 586

distributed networking communicationmulti-GPU communication, 125-148

routing factors, 781-784

speculative decoding, 791

scaling, 665

scale-out of workers, 891

why disaggregation, 761-796

advantages of disaggregation, 762-765,

data parallelism as parallelism strategy,

670, 677

definition of Data Parallel, 132, 677

definition of Distributed Data Parallel,

schedulingdefinition of Data Parallel, 132

definition of Distributed Data Parallel,

Distributed Data Parallel strategies,

132-137

NCCL communication algorithms,

129-131

NCCL topology awareness, 126-129

profiling and debugging NCCL, 145

multinode communication pitfalls, 119-125

GLOO backend instead of NCCL,

119-121

GPU memory fragmentation under

UCX/RDMA, 124-125

insufficient bandwidth or misconfigured

NICs, 122

mismatched NCCL versions, 122

straggler nodes or processes, 123

summary, 125

TCP port exhaustion during NCCLbootstrap, 122

multinode connectivity tuned, 117-119

NCCL pitfalls and gotchas, 137-145

communicator hanging, giving errors,shutting down, 143-145

CPU-GPU NUMA-node for NCCLthreads, 142

creating and destroying communicators,

creating communicators too often,

137-140ignoring warnings and errors, 143

overturning or disabling features viaenvironment variables, 140-142

NVIDIA Magnum IO, 113

optimizing for Kubernetes, 95

computationcomparison of no overlap with overlap,

107-112

Distributed Data Parallel strategies,

132-137reducing frequency and volume, 107

streams for asynchronous execution, 106

RDMA, 114-125

distributed shared memory (DSMEM), 341

sharing state between thread blocks, 405

thread block clusters, 402-413

algorithms for parallelizing workloads,

416-418coordinating thread block clusters,

408-410described, 373

global memory traffic reduced, 413-416

launching a thread block cluster, 407,

scratch memory, 406

thread block pair, 410-413

thread block swizzling, 404

threads accessing each other’s sharedmemory, 196, 373, 404

warp specialization with thread blockclusters, 418-425

DistServe, 760

DMA (see direct memory access (DMA))

Dockerno swap and enforcing NUMA affinity, 74

NVIDIA Container Toolkit, 88

RDMA, 114

runtime optimizationsabout containers, 88

container overlay filesystem overhead,

NVIDIA container runtime injectinglibraries, 89

NVIDIA Container Toolkit and CUDAcompatibility, 88

reducing image size, 91

runtime optimizations for GPUs, 88-91

double buffering with CUTLASS, 360

parallelism

DPU (Data Processing Unit) in NVL72, 41

DRAM as shared memory, 255

(see also CPU DRAM)

DSMEM (see distributed shared memory

(DSMEM))

DTensor (PyTorch), 590

dual-die GPUs, 26

appearing as one GPU, 27

GPC “partitions”, 403

programming, 217

remote memory access latency, 62

DualPipe (DeepSeek), 6

mixture-of-experts model DeepSeek-V3, 9

DWARF (Debugging With Attributed Record

Formats), 540

dynamic activation quantization, 919

dynamic batching, 718-720

adaptive batching delay, 719

static batching versus, 718

dynamic CUDA Graph update, 496

dynamic early exit networks, 918

dynamic parallelism (DP), 506-512

profiling before and after a change, 512

child-kernel launch limit, 507

device memory usage monitored, 507

“stack overflow” errors avoided, 507

when to capture CUDA Graph instead, 507

dynamic routinglatency-aware scheduling, 724-725

MoE inference strategies, 691-697

691-693expert replication, 694

load balancing and capacity factors, 693

dynamic schedulingatomic work queues, 482-489

benchmarks, 511

expert communication optimization,

atomic counters, 483

atomic queues, 485-489

dynamic work allocation, 483

|

Nsight Compute, 483-484

dynamic congestion-aware scheduling,

913-915load balancing, 834-841

hotspot

Arrow’s adaptive instance scaling,

837-838dynamic resource scaling, 840

Mooncake adaptive strategies, 839

TetriInfer’s two-level scheduler, 835

dynamic shapes (PyTorch)

TorchDynamo, 553, 609

variable-length sequences via, 620-623

padding

XLA, 658

dynamic shared-memory allocation, 858

Dynamo (see NVIDIA Dynamo; TorchDynamo

(PyTorch))

Eearly rejection, 829-831

QoS, 792-795

ECC mode for GPU memory, 85

EGM (Extended GPU Memory), 24

environment variablesbeing explicit instead of default values, 144

CUDA_DEVICE_ORDER, 143

CUDA_VISIBLE_DEVICES, 143

debugging compiler phases, 636

MALLOC_CONF, 76

NCCL features overturned or disabled,

140-142

NCCL information online, 137

NCCL_ALGO, 127, 131, 907

NCCL_ASYNC_ERROR_HANDLING,

144, 145

NCCL_BUFFSIZE, 140

NCCL_CROSS_NIC, 909

NCCL_DEBUG, 129, 141, 145, 715

NCCL_IGNORE_CPU_AFFINITY, 142

NCCL_MAX_NCHANNELS, 142

NCCL_MIN_NCHANNELS, 142

NCCL_NSOCKS_PERTHREAD, 118, 122,

NCCL_P2P_DISABLE, 140

NCCL_PORT_RANGE, 95

NCCL_PROFILER_PLUGIN, 145

|

EP (see expert parallelism (EP))

EPLB (expert parallelism load balancer; Deep‐Seek), 6

error handling

CUDA kernels, 206

inference system, 731

NCCL, 145

NCCL_ASYNC_ERROR_HANDLING, 144

ExecuTorch, 528, 531

expert collocation, 692

expert parallelism (EP), 670, 674-677

versus TP, PP, and hybrid, 843-847

expert parallelism load balancer (EPLB; Deep‐Seek), 6

expert replication, 694

exponential computation unit, 16

Extended GPU Memory (EGM), 24

Ffailure

NCCL_SHARP_DISABLE, 142

NCCL_SHM_DISABLE, 141

NCCL_SOCKET_IFNAME, 95, 143

NCCL_SOCKET_NTHREADS, 118, 122,

NCCL_TOPO_DUMP_FILE, 129

NCCL_TOPO_FILE, 142

PYTORCH_CUDA_ALLOC_CONF, 86,

TCMAL‐LOC_MAX_TOTAL_THREAD_CACHE_BYTES,TCMALLOC_RELEASE_RATE, 77

TORCHDYNAMO_REPRO_*, 636

TORCHINDUCTOR_BENCH‐MARK_KERNEL, 553

TORCHINDUCTOR_UNIQUE_KER‐NEL_NAMES, 553

TORCH_COMPILE_DEBUG, 636

TORCH_LOGS, 552

compiler debugging, 636

logging options summary, 636

performance hints, 623

updating debugging values to production,

hybrid parallelism, 681

load balancing, 676

model weights and data split over GPUs,

filesystem caching and write-back, 75

GPU failures at scale per Meta, 143

NCCL communicator hanging, givingerrors, shutting down, 143-145

fatbinary model, 58

fault tolerance, 832

Feedforward Neural Network (FFN) activationcheckpointing, 572

Feynman GPU, 51

filesystemscaching and write-back, 75

DeepSeek Fire-Flyer File System, 172-174

FUSE unable to GDS, 173

must be optimized for large concurrent I/O,

NFS, 174

tuning parameters, 175

parallel filesystems, 174

shared filesystems, 174

tuning, 176

Fire-Flyer File System (3FS; DeepSeek), 6,

172-174

FireworksAI, 566

first in, first out (FIFO) scheduler, 724-725

FlashAttention, 15

FlashMLA (Flash Multi-Head Latent Attention;DeepSeek), 6, 800

hardware-software codesign, 15

FlexAttention (PyTorch), 553, 619, 802

nested jagged-layout tensors, 802

PagedAttention, 803

FlexDecoding (PyTorch), 554, 802-805

nested jagged-layout tensors, 802

float8 not provided by CUDA, 248

FP4, FP8, FP16, 29, 356

dynamic precision changes, 848-853

inference quantization, 734

weight-only quantization, 735

FSDP (Fully Sharded Data Parallel; PyTorch),

133, 574-577

automated checkpointing, 572, 574-577

scaling via FSDP with torch.compile,

586-589

SimpleFSDP, 590

TorchTitan, 590

FUSE unable to GDS, 173

fused kernels (see kernel fusion)future hardware from NVIDIA, 49-52

Rubin Ultra and Vera Rubin Ultra, 51

Vera Rubin Superchip, 50

FX Graph

AOT Autograd, 609, 612

AutoParallel, 590

PrimTorch IR, 614

TorchDynamo, 609

G

GB200 (see NVIDIA GB200 NVL72)

GB300 (see NVIDIA GB300 NVL72)

GDS (GPUDirect Storage; NVIDIA), 70,

167-172, 573

Gemini Ultra (Google) training costs, 2

GEMM (general matrix multiply)

AI-discovered algorithms, 924-926

GEMM kernels, 348

CUTLASS, 362

GEMM pipeline, 472-478

PyTorch compiler TorchInductor, 545

roofline analysis, 537-539

thread block clusters, 417

GitHub

DeepSeek repositories, 6

reporting an issue, 625

TorchBench-based performance regression,

GLaM (Google), 676

global memorycaching loads, 273

traffic reduced, 413-416

global memoryarithmetic intensity, 233

coalesced versus uncoalesced access,

Blackwell Ultra and Grace Blackwell Ultra,

doubling every year, 51

Feynman GPU, 51

GPU programming across generations, 58,

dumping graph after each stage, 637

TorchInductor, 612

239-246

GPU memory hierarchy, 212-218

Initcheck of NVIDIA Compute Sanitizer,

matching loads to native transaction size,

pitfall of repeatedly reading same data, 255

|

__global__, 203

Gloo backend should be NCCL, 119-121

goodput as throughput metric, 16

Google Cloud TPU and XLA, 615, 658

Google Gemini Ultra training costs, 2

Google GLaM, 676

Google Switch Transformer MoE efficiency, 12

GPCs (GPU processing clusters), 403

GPT-4 (OpenAI) training costs, 2

GPU Boost (NVIDIA), 84

setting power limit below TDP, 84

underclocking to reduce heat, 84

GPU driverkeeping GPU driver loaded, 87

nvidia-smi, 57

persistence mode for performance, 77

software stack, 57

GPU processing clusters (GPCs), 403

GPUDirect RDMA, 38, 115

multinode, multirack communication,

908-911

NCCL’s native allocator in PyTorch Mem‐Pool, 579

verify true GPUDirectRDMA, 115

Storage

GPUsarchitecture, 191-203

maximum threads per SM and block,

memory hierarchy, 212-218

Special Function Unit, 194

streaming multiprocessors, 192

thread block clusters, 196

thread block size, 200

thread blocks, 195-199

threads, 195-199

threads-per-block and blocks-per-grid,

200-202

Unified CPU-GPU Memory, 24

warps, 32, 192-194

bottlenecks profiled and diagnosed, 288-293

about, 288

utilizationdata pipeline profiling and tuning, 289

iteratively profiling, 301-303

kernel optimization, 303-305

monitoring in production, 705

|

Nsight Compute and Roofline analysis,

Nsight Systems timeline view, 288, 291

PyTorch profiler via Kineto, 291-293

warp stall reasons, 293-297

compressed data, 176

maximumcontainer runtime optimizations, 88-91

about containers, 88

container overlay filesystem overhead,

NVIDIA container runtime injectinglibraries, 89

NVIDIA Container Toolkit, 88

NVIDIA Container Toolkit and CUDAcompatibility, 88

reducing image size, 91

bandwidth

CUDA events for communication with CPUhost, 447

callback function, 447

custom kernels with OpenAI Triton,

637-658registering with PyTorch, 641-643

tuning kernel launch parameters, 643

distributed multi-GPU communication,

125-148

FLOPs outpacing memory bandwidth, 229

GPU failures at scale per Meta, 143

GPU-NIC affinity forced, 910

GPU-to-GPU memory sharing with

NVSHMEM, 513-517

idle time measurement, 177

Kubernetes scheduler via NVIDIA deviceplugin, 91

memory access patternsavoiding shared memory, 271-273

avoiding shared-memory bank conflicts,

264-271coalesced versus uncoalesced globalmemory access, 239-246

read-only data caches, 273-278

read-only data caches pitfall, 273

symmetric memory, 582

Tensor Memory Accelerator for tilefetch, 278-283

tiling and data reuse via shared memory,

255-264vectorized memory access, 247-254

warp shuffle intrinsics, 271

memory fragmentation under UCX/RDMA,

124-125multiple GPU orchestration, 512-521

(see also multiple GPUs)

GPU-to-GPU memory sharing, 513-517

NVIDIA SHMEM library, 513-517

parallelismadding two vectors sequentially and inparallel, 222-228

architecture of GPU, 191

child-kernel launch limit, 507

cuTile, 59

Data Parallel versus Distributed Data

Parallel, 132-137

dynamic parallelism, 506-512

fully sharded data parallel, 133

PTX, 365

thread block cluster algorithms, 416-418

power limits, 47, 730

process-GPU mapping, 902-904

programming CUDA, 221-229

memory hierarchy, 212-218

reading data directly from storage devices,

167-172runtime settings for performance, 77-87

ECC mode for memory, 85

GPU clock speeds, 84

oversubscription

GPU persistence mode, 77

keeping GPU driver loaded, 87

MIG, 81-84

MPS, 78-80

utilization versus latency, 729

virtualization via Multi-Instance GPU, 5, 94

Grace Blackwell focus but not limit, 14

CPU and GPU linked by NVLink-C2C, 62,

Grace Blackwell GB200 NVL72, 811, 900

Grace Blackwell GB200 Superchip, 23

Grace Blackwell Superchip, 13

memory, 25, 128

memory

Grace Blackwell Ultra GB300 NVL72, 49, 900

Grace CPU, 26

Grace Blackwell Superchip, 13

memory, 25

NVL72, 33

Grace Hopper GH200, 23

Grace Hopper Superchip NIXL offloading KVcache, 157-159

gradient checkpointing, 572

(see also activation checkpointing)

Grafana dashboardsinference monitoring, 705, 706

NVLink/NVSwitch monitoring, 902

graph breaks, 627-635

explain() output, 552, 589

forcing error to be raised at compile, 608

grids, 195-199

tricks using CPU-GPU unified memory,

avoiding whenever possible, 608, 627

conditionals as tensor operations, 630

data-dependent branches, 630

marking functions and code blocks assafe, 632

minimizing graph recompilations, 632

causes, 633-635

sequence not captured into a singlegraph, 627

shape mismatches, 553

tips for handling common causes,

633-635unsupported operations in model, 552

complex graphs requiring graph breaks, 627

debugging, 628, 636

monitoring graph recompilations, 609,

captured

PyTorch compiler trace, 533

TorchDynamo, 611

blocks-per-grid sizes, 200-202

blocksPerGrid parameter, 203, 207-209

cooperative groupscooperative kernel launch grid size, 397,

persistent kernels via grid sync, 398-401

maximum dimensions, 202

|

H

H100 versus NVL72, 39

half-precision, 356

hardware

AI supercomputer-in-a-box overview, 23

cache coherency, 24

CPU + GPU superchip, 23-33

GPU Tensor Cores and Transformer

Engine, 29-31

many GPUs as one, 33-43

multi-GPU programming, 38-39

multirack and storage communication,

NVIDIA Blackwell dual-die GPU, 26-28

NVIDIA Grace CPU, 26

NVL72, 33

(see also NVL72)streaming multiprocessors, threads,warps, 31-33

memory-optimized

DeepSeek.AI export restrictions, 1, 9-11

disaggregated prefill and decode, 820-826

future evolution, 939-943

future hardware from NVIDIA, 49-52

Blackwell Ultra and Grace Blackwell

Ultra, 49

doubling every year, 51

Feynman GPU, 51

Rubin Ultra and Vera Rubin Ultra, 51

Vera Rubin Superchip, 50

GPU programming across generations, 58

GPU to Kubernetes scheduler, 91

hardware-software codesign, 15-16

CUDA Pipeline API plus TMA, 282

DeepSeek MLA, 15

homogeneous to avoid straggler nodes orprocesses, 123

occupancy limited by, 200, 201

optimizations open source, 6

Page Migration Engine, 86

ROI of upgrading, 48

roofline model, 233-237

HBM (high-bandwidth memory), 940

hiding latency (see latency hiding)high occupancy, 199

maintaining high occupancy, 221-229

impact of occupancy, 222-228

programming CUDA, 221-229

|

high-bandwidth memory (HBM), 940

Holistic Trace Analysis (HTA; PyTorch), 528,

multi-GPU profiling, 591

Hopper H100, 1

HTA (see Holistic Trace Analysis (HTA;PyTorch))hugepages, 72

transparent hugepages, 72

Hugging Face Transformers library

CTRL, 743

KV cache offloading, 864

KV cache quantized, 867

implementation

Hugging Face Ultra-Scale Playbook, 935

hybrid parallelism, 679-681

versus TP versus PP, 843-847

hyperthreading, 69

I

networking

IBTA (InfiniBand Trade Association), 38

ILP (see instruction-level parallelism (ILP))in-flight batching, 720

Inductor (see TorchInductor (PyTorch))inferencebatchingarithmetic intensity increased, 348

comparing static, dynamic, continuous,

continuous batching, 720

dynamic batching, 718-720

latency-aware scheduling, 724-725

stall-free scheduling, 723

static batching, 718

continuous scheduling, 721

CUDA Graphs and inference engines, 490

disaggregated prefill and decode, 663

(see also disaggregated prefill anddecode)dynamic routing strategies for MoE models,

691-697expert communication optimization,

691-693expert replication, 694

load balancing and capacity factors, 693

Dynamo throughput improved by NIXL,

massive inference clusters, 678

megakernels, 394

multinode inference, 663

disaggregated

Kubernetes for deploying, 668-670, 714

NIXL disaggregated inference, 148-161

about NIXL, 126, 148-150

asynchronous API with callbacks,

153-157intelligent interconnect routing, 152

KV cache offloading with NIXL, 157-159

separate prefill and decode inferencestages, 150-152

inferenceadaptive batching, 893-899

adaptive parallelism strategies, 843-847

chunked prefill scheduling, 893-899

dynamic activation quantization, 919

dynamic early exit networks, 918

dynamic memory-allocation switching,

880-886dynamic precision changes, 848-853

dynamic shared memory allocation,

858-862dynamic token pruning, 919

edge-oriented MoE memory budgeting,

input-aware layer skipping, 918

kernel autotuning, 853-858

KV cache compression and policyswitching, 867-875

MoE speculative expert routing, 918

topology-awareoccupancy-aware kernel selection,

858-862prewarming graphs and caches, 889-893

runtime kernel patching, 886-889

speculative KV prefetching, 862-867

tuning AI at runtime via RL agents, 875

optimizations at the application level,

740-757debouncing and request coalescing, 755

705-709observe-hypothesize-tune loop, 702-705

profiling with Nsight, 709-712

runtime tuning via RL agents, 875

systems-level optimizations, 725

troubleshooting, 712

InfiniBand (IB) networkingcontainer access to InfiniBand devices, 114

counters on nodes to spot stragglers, 123

enabling GPUDirect RDMA, 96

host channel adapter, 121

RDMA support, 115

SHARP, 146

InfiniBand Trade Association (IBTA), 38

inputvariable-length sequences, 620-623

model cascading and tiered deployment,

750-752prefix caching, 743-749

prompt cleansing, 741

prompt compression, 740

streaming responses, 752-755

token output limits and timeouts, 756

optimizations at the systems level, 725

GPU utilization versus latency, 729

KV cache, 732

memory, 732

computationpower and thermal constraints, 730

quantization for real-time inference,

733-739about quantization, 733

activation quantization, 737

quantizationpost-training quantization workflow,

quant-dequant steps into executiongraph, 739

reducing precision, 734

weight-only quantization, 735

structured sparsity for inference workloads,

tuning performancecorrectness issue debugging, 715-717

full-stack optimizations, 713-715

monitoring system metrics and counters,

dynamic shapes, 620-623

|

padding, 621

paddingvery long inputscontext parallelism, 678

prompt compression, 740

stall-free scheduling, 723

instruction scheduling with PTX, 366

instruction-level parallelism (ILP)exploitingcompiler hints, 326, 333

interleaving independent operations, 332

unrolling loops, 332

exposing, 325-334

compiler hints, 326, 333

occupancy and ILP, 326, 330, 333

instructionsincreasing for latency-bound GPUs, 303

increasing ILP with PTX, 367

performed

PyTorch compiler not guaranteeing, 321

register pressure, 333

warp stall reasons, 296, 297

INT8, 357

inter-kernel pipeliningabout inter-kernel pipelining, 429

compute overlapping with data transfers viastreams, 433-435

concurrency, 556-558

multiple GPUs, 463-469

CUDA stream ordered memory allocator,

435-439using, 439-441

using with multiple GPUs, 463-469

default streamsdefault versus explicit streams, 443

legacy default streams, 441

per-thread default streams, 442

using default streams, 444-446

kernel execution overlapped with CUDAstreams, 429-433

launching five kernels on two streams,

429-431warp specialization replaced with CUDAstreams, 431-433

multi-GPU overlap of compute and data,

463-469

|

cross-stream synchronization, 448

warp specialization

CUDA streams, 449-456

thread block clusters and CUDAstreams, 456

interrupt requests (IRQs)irqbalance daemon, 74

scheduler and interrupt affinity, 73

intra-kernel pipelining, 374

CUDA Pipeline API, 385-390

CUDA Pipeline API in PyTorch, 390

CUDA streams, 449-456

thread block clusters, 418-425

when to use warp specialization versusdouble-buffering, 385

bottlenecks

Programmatic Dependent Launch, 469-472

specializationsynchronization with events and callbacks,

about intra-kernel pipelining, 374

comparing techniques, 374, 389, 421-425

cooperative groups, 395-401

Cooperative Groups API, 272, 396,

408-410described, 402

launching a kernel in cooperative mode,

396, 401

persistent kernels via grid sync, 398-401

thread block clusters versus, 402

when to combine with persistent kernels,

double-buffering with CUDA Pipeline API,

375-380persistent kernels, 391-395

common workloads for, 393

cooperative groups via grid sync,

398-401megakernels, 394

occupancy, 393

warp specialization, 395

when to combine with cooperativegroups, 401

thread block clusters, 373

(see also thread block clusters)warp specialization, 381-385

continuous profiling and tuning workflow,

184-188data pipeline tuning, 178

data loading and preprocessing, 178-181

NeMo Curator for training datasets, 183

NVIDIA Data Loading Library, 182

scaling out workers as GPUs scaled, 181

I/O isolation in Kubernetes, 98

monitoring, 177-178

NVIDIA Magnum IO, 113

optimizing GPU-based storagecloud storage caches, 174

compressing data, 176

DeepSeek Fire-Flyer File System,

172-174fast storage and data locality, 163

filesystem optimized, 166

multiqueue block I/O scheduler, 166

NFS servers, 174

NFS servers tuning parameters, 175

NVIDIA GDS, 167-172

NVMe and filesystem tuned forthroughput, 166

parallel filesystems, 174

read ahead, 166

read size tuned, 165

sequential versus random reads, 165-166

shared filesystems, 174

striping files, 175

tuning filesystems, 176

irqbalance daemon, 74

iteration-level scheduling, 720

Jjemalloc, 76

job scheduling with Kubernetes and SLURM,

K

K8s (see Kubernetes (K8s))kernel autotuning (see autotuning)kernel caching, 891

kernel fusionarithmetic intensity increased, 344

custom fusion less necessary with PyTorch,

FlexAttention, 619, 802

memory-bound workloads, 522

TorchInductor, 313, 344, 610

fused attention kernels, 391

GEMM kernel prologue and epiloguefusion, 545

kernel launch overhead reduced with CUDA

Graphs, 562-569

best practices for CUDA Graphs, 567-568

capturing a CUDA Graph, 562-565

replaying the graph, 565-566

any stream for replay, 564

kernel optimization, 303-305

automated, 926-929

reinforcement learning approach, 929-931

runtime kernel patching, 886-889

Kimi K2 (Moonshot AI), 943

Kineto (PyTorch profiler), 291, 528, 529

NCCL profiler plugin, 145

Kubernetes (K8s)about, 91

cluster scheduler, 47

deploying a disaggregated inference system,

668-670, 714

host machine should be tuned, 98

I/O isolation, 98

improving resource guarantees, 96

job scheduling, 93

memory isolation, 97

memory limit, 97

MIG devices in device plugin, 84

MIG mode, 94

no swap and enforcing NUMA affinity, 74

NVIDIA Container Toolkit, 88

NVIDIA device plugin for Kubernetes, 91

NVIDIA Kubernetes GPU Operator, 91

NVL72 deployment, 42

optimizing network communication, 95

orchestrating containers with Topology

Manager, 92

orchestration jitter reduced, 96

RDMA, 114

runtime optimizationsabout containers, 88

container overlay filesystem overhead,

NVIDIA container runtime injectinglibraries, 89

NVIDIA Container Toolkit and CUDAcompatibility, 88

reducing image size, 91

runtime optimizations for GPUs, 88-91

|

topology awareness, 91

Topology Manager, 92

Kubernetes Topology Manager (NVIDIA), 92

KV cache, 149

compression and policy switching, 867-875

inference tuning, 732

NIXL streamlining transfers, 149

intelligent interconnect for transfers, 152

offloading to CPU memory, 149,

157-159prefill and decode inference stages,

150-152prefill to decode data transfers, 667,

812-817prefill phase, 759

prefix merge events, 703

speculative KV prefetching, 862-867

KV cache offloading, 864

superchip memory bandwidth, 811

transfer from prefill to decode, 667, 812-817

connector and data path design, 817-820

KV cache size, 812

zero-copy GPU-to-GPU, 813-817

tuning, 805-812

disaggregated KV cache pool, 805-808

memory layout, 810

reuse and prefix sharing, 808-810

vLLM PagedAttention, 722, 732

KV Cache Manager (NVIDIA Dynamo),

157-159

Llatencydisaggregation of prefill and decode, 666

GPU utilization versus, 729

latency-aware scheduling and dynamicrouting, 724-725

latency-bound GPU stalls, 301-303

kernel optimization, 303-305

latency-sensitive training workflows, 75

maximum latency for dynamic batching,

profiling when streaming enabled, 755

latency hiding, 32, 221

inter-kernelhigher occupancy to better hide latency, 306

high occupancy plus high ILP, 333

low achieved occupancy, 297

|

within a kernel (see intra-kernel pipelining)

__launch_bounds__, 229-231, 310

occupancy optimization, 311

LazyLLM, 919

__ldg(), 273

Linux

CPU DRAM as CPU NUMA memory, 62

GPU driver, 57

GPU servers typically Linux, 56

memory swapping avoided, 74

NUMA balancing, 63

numactl

CPU pinning, 62, 63

forcing memory allocation from specific

NUMA node, 62, 70

process CPU and memory policiesapplied on launch, 70

NVMe

DeepSeek Fire-Flyer File System,

schedulerscompletely fair queueing obsolete, 166

Completely Fair Scheduler, 73

multiqueue block I/O scheduler, 166

none scheduler, 166

setting for, 166

Managementcluster scheduler, 47

liquid cooling versus air cooling, 45-46

Little Red Whiting Hood, xxv

LLM request nonuniformity, 705

LMCache (vLLM), 703

disaggregated prefill and decode, 760, 815

load balancingdynamic scheduling, 834-841

172-174fast storage and data locality, 163

GPUDirect Storage, 70, 167-172

parameters offloaded to, 573

sequential versus random reads, 165-166

tuning for throughput, 166

XFS optimized, 166

perf, 528, 530

example MoE model, 539-543

NVIDIA Performance Monitoring Unit,

hotspot

Arrow’s adaptive instance scaling,

837-838dynamic resource scaling, 840

Mooncake adaptive strategies, 839

TetriInfer’s two-level scheduler, 835

expert parallelism, 676

MoE expert rebalancing, 911

temporal load balancing, 913

logginginference monitoring, 708

MLPerf for performance benchmarks,

595-598

NCCL logs, 147, 715

PyTorch compiler, 636

TORCH_LOGS, 552, 636

logging options summary, 636

performance hints, 623

loop unrolling, 332, 344

M

Magnum IO (NVIDIA), 113, 542

MALLOC_CONF, 76

massive modelsinference clusters, 678

parallelism strategies for serving, 670-681

hybrid parallelism, 679-681

matmul (matrix multiplication)

AsyncTP, 590

DeepGEMM library FP8-optimized, 6

persistent Triton kernel, 646-653

mechanical sympathy, 15-16

Medusa multiple head speculative decoding,

multitoken decoding, 686-688

megakernels, 394

Mellanox, 41

memoryasynchronous memory transfers with CUT‐LASS, 360

Blackwell GPU, 25

close to CPU, close to GPU, 70

CPU-GPU interconnects on superchips, 128

performance, 25

CUDA stream ordered memory allocator,

435-439multiple GPUs, 463-469

using with CUDA streams, 439-441

CUDA Unified Memory, 86

programming CUDA, 218-221

Unified CPU-GPU Memory, 24

distributed shared memory (see distributedshared memory (DSMEM))dynamic memory-allocation switching,

880-886

ECC mode for GPU memory, 85

edge-oriented MoE memory budgeting, 919

fragmentation monitoring, 705

dynamic memory-allocation switching,

880-886

GDS, 70, 167-172, 573

global memory traffic reduced, 413-416

(see also global memory)

GPU FLOPs outpacing memory bandwidth,

GPU memory access patternsavoiding shared memory, 271-273

avoiding shared-memory bank conflicts,

264-271coalesced versus uncoalesced globalmemory access, 239-246

read-only data caches, 273-278

read-only data caches pitfall, 273

symmetric memory, 582

Tensor Memory Accelerator for tilefetch, 278-283

tiling and data reuse via shared memory,

255-264vectorized memory access, 247-254

oversubscription

GPU memory fragmentation under UCX/

RDMA, 124-125

GPU-to-GPU memory sharing with

NVSHMEM, 513-517

Grace Blackwell Superchip, 25, 128

Grace CPU, 25

high-bandwidth memory, 940

hugepages, 72

transparent hugepages, 72

inference tuning, 732

kernel memory throughput versus HBMmemory bandwidth, 299

Kubernetesmemory isolation, 97

memory limit, 97

KV cachememory layouts optimized, 810

offloaded to CPU memory, 157-159

|

memory-bound workload example, 228

NUMA-friendly allocation and pinning,

70-72efficiency of pinned memory, 70

max locked memory setting, 72, 74

OS limit on pinned memory, 71

pinned memory for data loaders, 71

out of memory (OOM) errors

CPU + GPU architecture, 26

max locked memory setting, 72

memory fragmentation or excessivecaching, 86

monitoring GPU memory usage, 87

profiling and tuning via PyTorch profiler,

569-582activation checkpointing, 572

CUDA memory allocator tuned, 570

FSDP automatic checkpointing, 574-577

FSDP with tensor parallel and pipelineparallel, 577

topologiesoffloading parameters to CPU and

NVMe, 573

peer-to-peer DMA, 581

pluggable memory allocators, 578-581

programming CUDAasynchronous memory allocation,

210-212

CUDA streams, 211

GPU memory hierarchy, 212-218

memory pools, 210-212

PyTorch memory allocator plugin, 578-581

recomputation versus memory, 350

scratch memory, 406

DSMEM, 407

(see also distributed shared memory

(DSMEM))shared-memory exchanges disabled, 141

(see also shared memory (SMEM))swapping avoided, 74

max locked memory setting, 74

symmetric memory, 582

tuning host CPU memory allocator, 76

Unified CPU-GPU Memory, 24

memory-bound GPU stalls, 301-303

kernel optimization, 303-305

MIG (Multi-Instance GPU; NVIDIA)

GPU driver runtime settings, 81-84

|

profile naming, 81

single GPU split into smaller GPUs, 47

when to use, 83

mixed precisionautomatic mixed precision, 355-356

reduced precision, 356

Tensor Cores, 351-360

mixture-of-experts (MoE) modelsabout MoE models, 531

all experts active simultaneously, 675

all-to-all communication, 693

CUDA streams, 560

DeepEP communication library, 6

DeepSeek-V3, 9-11

requirementsspeculative expert routing, 918

Attention

GPU virtualization, 5

Kubernetes MIG mode, 94

persistence mode recommended for GPU,

DualPipe parallelism algorithm, 9

dynamic routing strategies for inference,

691-697expert communication optimization,

691-693expert replication, 694

load balancing and capacity factors, 693

edge-oriented memory budgeting, 919

expert parallelism, 670, 674-677

load balancing, 676

model weights and data split over GPUs,

expert rebalancing and regrouping, 911

hybrid parallelism, 679-681

MoE efficiency, 12

open MoE rivaling best closed models,

multinode inference needed, 663

parallelism strategies, 670-681

profiling, 531-543

Linux perf on CPU and GPU, 539-543

Nsight Systems and NVTX Timelines,

535-537

PyTorch profiler, 532-535

roofline analysis for GEMM kernels,

MLPerf (MLCommons), 7

caution on per-GPU results, 8

Logging, 595-598

MLPerf Inference, 7

MLPerf Training, 7

MMA API (CUDA), 358-360

model cascading, 750-752

MoE (see mixture-of-experts (MoE) models)monitoring

CPU utilization, 182

filesystem I/O during training, 176

GPU memory usage, 87

graph recompilations, 609, 632

inference system metrics and counters,

705-709correctness issue debugging, 715-717

production environment, 712

memory fragmentation, 705

model rambling, 757

multiple-GPU NVLink utilization, 901

network throughput, 122

performance monitoring, 47

bottlenecks

GPUs near 100% utilized, 47

I/O continuous profiling and tuning,

184-188

NVLink usage, 47

power monitoring, 47

storage I/O, 177-178

straggler nodes or processes, 123

Mooncakeadaptive strategies, 839

early rejection, 829

Moonshot AI Kimi K2, 943

MPS (NVIDIA), 78-80

multi-GPU (see multiple GPUs)

Multi-Instance GPU (see MIG (Multi-Instance

GPU; NVIDIA))multidie GPU GPC “partition”, 403

(see also dual-die GPUs)multilevel tiling, 340

multinode connectivity tuned, 117-119

multinode communication pitfalls, 119-125

GLOO backend instead of NCCL,

119-121

GPU memory fragmentation under

UCX/RDMA, 124-125

insufficient bandwidth or misconfigured

NICs, 122

mismatched NCCL versions, 122

straggler nodes or processes, 123

summary, 125

TCP port exhaustion during NCCLbootstrap, 122

NCCL with UCX, 581

multinode inference, 663

disaggregated inference (see disaggregatedprefill and decode)

Kubernetes for deploying, 668-670, 714

multiple GPUscapturing collectives via NCCL and CUDA

Graphs, 517-520

congestion- and topology-aware scheduling,

899-917adaptive process-GPU mapping, 902-904

dynamic congestion-aware scheduling,

913-915

GPUDirect RDMA, 908-911

MoE expert rebalancing and regrouping,

communication

NCCL collective communication,

905-908

constraints

NVSwitch and fine-tuned scheduling,

915-917telemetry and monitoring, 901

many GPUs as one, 33-43

multimillion-GPU clusters, 939-943

n-GPU scaling pattern, 520

NCCL for distributed communication,

125-148orchestrating across, 512-521

GPU-to-GPU memory sharing, 513-517

NVIDIA SHMEM library, 513-517

symmetric memory, 582

overlap of compute and data, 463-469

multiple streams per GPU, 464

PyTorch, 468

peer-to-peer DMA, 581

profiling with HTA, 591

programming multiple GPUs, 38-39

PyTorch pluggable memory allocators,

578-581

|

stream-ordered memory allocator, 463-469

multirack and storage communication, 41

Nn-GPU scaling pattern, 520

Nagle’s algorithm, 754

NCCL (NVIDIA Collective Communications

Library)about NCCL, 5, 125, 905

backend instead of GLOO, 119-121

bootstrap TCP port exhaustion, 122

CUDA Toolkit, 58

debugging, 145

inference correctness issue, 715

NCCL debugging enabled, 715

NCCL logs, 147, 715

test suite, 715

dedicated high-priority streams for networktransfers, 464

ring or tree algorithm, 465-466, 905-907

distributed multi-GPU communication,

125-148about, 125

all-reduces in parallelism, 132

CollNet communication algorithm, 130

communication algorithms, 129-131

many-to-many collectives, 149, 150, 160

NCCL_ALGO environment variable,

127, 131

optimizing, 905-908

parallel aggregated tree communicationalgorithm, 130

ring communication algorithm, 129

topology awareness, 126-129

tree communication algorithm, 129

wave scheduling of collectives, 907

distributed multi-GPU communicationsring or tree algorithm, 905-907

rotating ring endpoints, 907

GPU-NIC affinity forced, 910

maximum TCP socket connections percommunicator, 141

memory allocator in PyTorch, 578

multiple NICs in parallel, 909

multiple-GPU collectives captured, 517-520

NIXL versus, 160

many-to-many collectives by NCCL,

149, 150, 160

persistent user buffers, 148

|

NCCL logs, 147

RDMA container access to devices, 114

SHARP and, 146

UCX for multinode topologies, 581

user buffer registration, 148

versions mismatched, 122

NCCL_ALGO, 127, 131, 907

NCCL_ASYNC_ERROR_HANDLING, 144,

NCCL_BUFFSIZE, 140

NCCL_CROSS_NIC, 909

NCCL_DEBUG, 129, 141, 145, 715

NCCL logs, 147, 715

NCCL_IGNORE_CPU_AFFINITY, 142

NCCL_MAX_NCHANNELS, 142

NCCL_MIN_NCHANNELS, 142

NCCL_NSOCKS_PERTHREAD, 118, 122, 141

NCCL_P2P_DISABLE, 140

NCCL_PORT_RANGE, 95

NCCL_PROFILER_PLUGIN, 145

NCCL_SHARP_DISABLE, 142

NCCL_SHM_DISABLE, 141

NCCL_SOCKET_IFNAME, 95, 143

NCCL_SOCKET_NTHREADS, 118, 122, 141

NCCL_TOPO_DUMP_FILE, 129

NCCL_TOPO_FILE, 142

ncu for Nsight Computeprofiler, 288, 528, 530

Roofline analysis, 290, 537-539

nested jagged-layout tensors (NJT), 802

nested tensors (PyTorch), 351

networking communication (see distributednetworking communication)

NFS servers, 174

pitfalls and gotchas, 137-145

communicator hanging, giving errors,shutting down, 143-145

CPU-GPU NUMA-node for NCCLthreads, 142

creating and destroying communicators,

creating communicators too often,

137-140ignoring warnings and errors, 143

overturning or disabling features viaenvironment variables, 140-142

versions mismatched, 122

profiling, 145

tuning parameters, 175

NIC-GPU affinity forced, 910

NIXL (NVIDIA Inference Xfer Library)about NIXL, 5, 126, 148-150

dedicated high-priority streams for networktransfers, 466

disaggregated inference, 148-161

about NIXL, 148-150

asynchronous API with callbacks,

153-157

Dynamo throughput improved with

NIXL, 159

intelligent interconnect routing, 152, 667

KV cache offloading with NIXL, 157-159

prefill to decode data transfers, 667

separate prefill and decode inferencestages, 150-152

NCCL versus, 160

one-to-one transfers by NIXL, 132, 149,

150, 160

NVIDIA Dynamo core component, 149

Dynamo throughput improved, 159

Nsight Computeachieved occupancy, 297-305

kernel compute throughput versus GPU

FLOPS, 299

kernel memory throughput versus HBMmemory bandwidth, 299

underutilization, 482

ALU pipe busy, 295

CLI ncuprofiler, 288, 528, 530

Roofline analysis, 290, 537-539

CUDA Program Counter Sampling, 711

ILP configuration, 332

impact of parallel versus serial operation,

inference profiling, 711

CUDA Program Counter Sampling, 711

kernel shift from memory to computebound, 353

launch statistics, 418

Memory Workload Analysisatomic transactions counter, 483-484

coalesced versus uncoalesced memoryaccesses, 241

memory workload analysis, 856

Occupancy

Limited by Registers, 333

low achieved FLOPS, 300

example MoE model, 537-539

Source Counterslow achieved FLOPS, 300

warp divergence profiled, 318

Speed of Light analysis, 351

Stall

Compute Unit Busy, 295

Exec Dependency, 295, 334

Idle, 295

Long Scoreboard, 293

Math Pipe Throttle, 295

Memory Dependency, 295

Memory Throttle, 294

No Eligible, 295

Not Selected, 294

other reasons, 295-297

Short Scoreboard, 294

Texture Throttle, 295

Triton kernel support, 657

underutilization metrics, 482

updated to latest GPUs, 293

Warp State, 318

Nsight Systemsabout, 288

CLI nsysprofiler, 288, 528, 530, 535-537

timeline view, 288, 535-537

communication versus compute bottleneck,

NVTX markers, 288, 535-537

Triton kernel support, 657

updated to latest GPUs, 293

nsys as CLI for Nsight Systems

Occupancy Calculator, 309

occupancy limiters, 297

parallelism strategies profiled, 680

per-kernel metrics, 277

metrics

Roofline analysis, 290

I/O monitoring, 177

impact of parallel versus serial operation,

inference profiling, 709-710

NVLink profiling, 129, 902

parallelism strategies profiled, 680

Python backtrace sampling, 292

PyTorch focused mode, 292

timeline view, 277, 288, 291, 535-537, 856

|

profiler, 288, 528, 530, 535-537

timeline view, 288, 535-537

NUMA nodes, 62

NUMA awareness and CPU pinning, 62-69

CPU pinning, 62-69

NUMA-friendly memory allocation andpinning, 70-72

numactl

CPU pinning, 62, 63

forcing memory allocation from specific

NUMA node, 62, 70

process CPU and memory policies appliedon launch, 70

numerical correctness and accuracy debugging,

624-627numpy replacement cuPyNumeric, 59

NV-HBI (High-Bandwidth Interface), 27

nvcc CUDA compilercompatibility across GPU generations, 58,

-cubin option, 59

CUDA Toolkit in software stack, 57

fatbinary, 58

float8 not provided, 248

__global__, 203

__launch_bounds__, 229-231

optimizing for occupancy, 229-231

per-thread default stream enabled, 442

predication, 317

NVIDIA

AI supercomputers in a rack, 13-15

CPU + GPU superchip, 23-33

GPU Tensor Cores and Transformer

Engine, 29-31

many GPUs as one, 33-43

multi-GPU programming, 38-39

multirack and storage communication,

NVIDIA Blackwell dual-die GPU, 26-28

NVIDIA Grace CPU, 26

NVL72, 33

(see also NVL72)

NVLink and NVSwitch, 34-37

SHARP for in-network aggregations,

40-41streaming multiprocessors, threads,warps, 31-33

future hardware, 49-52

|

Blackwell Ultra and Grace Blackwell

Ultra, 49

doubling every year, 51

Feynman GPU, 51

Rubin Ultra and Vera Rubin Ultra, 51

Vera Rubin Superchip, 50

stock drop on DeepSeek-R1 costs, 2

NVIDIA Base Command Manager, 42

NVIDIA Blackwell B200, 1, 27, 201, 217

NVIDIA Blackwell B300 Ultra, 12, 49

NVIDIA Blackwell dual-die GPU, 26-28

remote memory access latency, 62

NVIDIA Collective Communications Library

Communications

NVIDIA Compute Sanitizer, 231

NVIDIA Tools Extension support, 231

concurrency and synchronization, 561

four primary tools, 232

inference correctness, 715

NVIDIA Container Toolkit, 88

NVIDIA CPU + GPU superchip, 23-33

31-33

NVIDIA Data Center GPU Manager (see Data

Center GPU Manager (DCGM; NVIDIA))

NVIDIA Data Loading Library (DALI), 182,

Manager

NVIDIA Dynamo with core component NIXL,

improving throughput, 159

NVIDIA GB200 NVL72, 13, 33, 811

P2P-capable NVLink, 581

training throughput per GPU, 7

NVIDIA GB300 NVL72, 13, 33, 900

Grace Blackwell Ultra Superchip, 49

NVIDIA GDS (see GDS (GPUDirect Storage;NVIDIA))

compute-sanitizer CLI, 231

GPU Tensor Cores and Transformer

Engine, 29-31

NVIDIA Blackwell dual-die GPU, 26-28

appearing as one GPU, 27

remote memory access latency, 62

NVIDIA Grace CPU, 26

NVLink-C2C, 24

streaming multiprocessors, threads, warps,

NVIDIA Grace Blackwell GB200 and GB300,

13-15

NVIDIA Grace CPU, 26

NVIDIA H800 GPU, 1, 9-11

NVIDIA Inference Xfer Library (see NIXL

(NVIDIA Inference Xfer Library))

NVIDIA Kubernetes GPU Operator, 91

NVIDIA Magnum IO, 113, 542

NVIDIA Management Library (NVML)

GPU-NIC affinity forced, 910

NVLink telemetry and monitoring, 901

power limit monitoring, 300

NVIDIA NeMo, 183

NeMo Curator, 183

NVIDIA NIXL (see NIXL (NVIDIA Inference

Xfer Library))

NVIDIA Performance Monitoring Unit (PMU),

NVIDIA SHMEM (NVSHMEM), 38, 513-517

CUDA Graphs, 516

when to use, 516

NVIDIA software stack, 56

about, 61

CUDA kernels created in Python, 56

CUDA Runtime and Toolkit, 57

Toolkit C++ and Python libraries, 59

GPU driver, 57

nvidia-smi, 57

PyTorch introduction, 60

NVIDIA Tools Extension (see NVTX (NVIDIA

Tools Extension))

NVIDIA Topology-Aware GPU Selection

(NVTAGS), 903

NVIDIA Transformer Engine (see Transformer

Engine (TE; NVIDIA))

NVIDIA Warp

CUDA kernels in Python, 56

CUDA Toolkit, 59

nvidia-peermem driver, 38

nvidia-persistenced daemon, 77

nvidia-smi, 57

nvJPEG, 176

NVL72, 13, 33

AI factory building block, 43

Blackwell Ultra and Grace Blackwell Ultraupgrades, 49

delivered as preintegrated rack appliance, 42

Kubernetes, 42

NVIDIA Base Command Manager, 42

900-901telemetry and monitoring, 901

NVLink

NCCL over dedicated streams, 465-466

NVL in names, 13

NVL72, 33, 43

NVSwitch, 34-37

900-901telemetry and monitoring, 901

Transformer Engine in modern GPUs, 16

NVLink-C2C (chip-to-chip), 24

Linux treatment of memory, 62

NVMe (Linux)

DeepSeek Fire-Flyer File System, 172-174

fast storage and data locality, 163

GPUDirect Storage, 70, 167-172

parameters offloaded to, 573

sequential versus random reads, 165-166

tuning for throughput, 166

XFS optimized, 166

NVMe over Fabrics (NVMe-oF), 41

Management

GPU-to-GPU bandwidth, 39, 93

H100 comparison, 39

inference throughput per GPU, 7

multi-GPU programming, 38-39

multirack and storage communication, 41

Data Processing Units, 41

NVLink and NVSwitch, 34-37, 43

performance monitoring, 47

power requirements, 43

liquid cooling versus air cooling, 45-46

sharing and scheduling workloads, 47

topology and bandwidth constraints,

fine-tuned scheduling and, 915-917

Magnum IO support, 114

NVIDIA SHARP, 40-41

NVIDIA supercomputer on a rack, 13

NVL72, 43

telemetry and monitoring, 901

topology and bandwidth constraints,

900-901

P2P-capable, 581

single domain for data tensor pipeline, 14

topology and bandwidth constraints,

cache coherency, 24, 38

Grace Blackwell CPU and GPU linked, 62,

|

NVML (see NVIDIA Management Library

(NVML))

NVSHMEM (see NVIDIA SHMEM

(NVSHMEM))

NVSwitch, 34-37

fine-tuned scheduling and, 915-917

Magnum IO support, 114

NVIDIA SHARP, 40-41

NVIDIA supercomputer on a rack, 13

NVL72, 43

telemetry and monitoring, 901

topology and bandwidth constraints,

900-901

NVTAGS (NVIDIA Topology-Aware GPU

Selection), 903

NVTX (NVIDIA Tools Extension)compute-sanitizer CLI supporting, 231

markers, 528

Holistic Trace Analysis, 591

profiling tools that use markers, 528-531

Nsight Systemsexample MoE model, 535-537

timeline view, 288, 535-537

O

Object Storage Targets (OSTs), 175

occupancyachieved occupancy and GPU utilization,

297-305definition of achieved occupancy, 297

kernel compute throughput versus GPU

FLOPS, 299

kernel memory throughput versus HBMmemory bandwidth, 299

configuring launch parameters, 207

CUDA Occupancy API, 309, 310

occupancy optimization, 312

definition of occupancy, 221, 306

hardware limiting, 200, 201

high occupancy, 199

high ILP with, 333

maintaining high occupancy, 221-229

ILP and, 326, 330, 333

impact of occupancy, 222-228

increasing occupancy and per-threadresources, 230, 306, 312

occupancy-aware kernel selection, 858-862

persistent kernels, 393

register trade-off, 231

|

active warps per scheduler, 307

CUDA Occupancy API, 310, 312

eligible warps per cycle, 307

__launch_bounds__, 229-231, 310, 311

right occupancy for workload, 307

techniques, 308

online resources (see resources online)

OOM (out of memory) errors

CPU + GPU architecture, 26

max locked memory setting, 72

cachingmonitoring GPU memory usage, 87

OOM killer (Linux), 97

Open Collective Foundation, 942

Open Infra Index (DeepSeek), 7

Open-Source Week (DeepSeek), 6, 7

OpenAI

ChatGPT stateful conversation, 744

GPT-4 training costs, 2

Triton, 638

617, 641-643

smart compilers, 934-936

TorchInductor basis, 60, 544

tuning kernel launch parameters, 643

warp specialization, 645

Triton Proton Profiler, 657

operating system, 56

operational intensity (see arithmetic intensity)optical networking, 43

optimization of AI systems by AI, 936-939

OS and CPU configuration, 61-77

tuning, 306

accessing shared memory, 641

advanced kernel implementations,

645-656

CUDA kernels in Python, 56

custom kernels, 637-658

double-buffering pipelining, 653-656

language and compiler, 60

Nsight Compute and Systems support,

open source, 60

persistent matmul kernel, 646-653

profiling with Triton Proton Profiler,

registering custom kernels with PyTorch,

about, 61

CPU frequency and C-states, 75

filesystem caching and write-back, 75

NUMA awareness and CPU pinning, 62-69

CPU + GPU superchip CPU-to-GPUdata transfers, 69

CPU pinning, 62-69

NUMA nodes, 62

NUMA-friendly memory allocation andpinning, 70-72

efficiency of pinned memory, 70

max locked memory setting, 72

OS limit on pinned memory, 71

pinned memory for data loaders, 71

scheduler and interrupt affinity, 73

irqbalance daemon, 74

transparent hugepages, 72

tuning host CPU memory allocator, 76

virtual memory and swapping, 74

OSTs (Object Storage Targets), 175

out of memory (OOM) errors

CPU + GPU architecture, 26

max locked memory setting, 72

cachingmonitoring GPU memory usage, 87

outputstreaming responses, 752-755

structured outputs, 689

token output limits and timeouts, 756

overlapping communication and computationautomating, 935

CUDA streams, 433-435

concurrency, 556-558

multiple GPUs, 463-469

distributed networking communication,

104-112

Distributed Data Parallel strategies,

132-137inference tuning, 725-729

computationsoffloaded KV cache, 864

PyTorch pattern, 557

O_DIRECTbypassing page cache, 75, 168

GDS, 170

P

P2P (peer-to-peer) DMA, 581

P2P (peer-to-peer) GPU copies, 140

padding shared arrays, 267

swizzling as alternative, 270

page cache

DeepSeek Fire-Flyer File System bypassing,

latency-sensitive training workflows, 75

O_DIRECT bypassing, 75, 168

Page Migration Engine (PME; NVIDIA), 86

page-cache size, 75

page-locked memory, 70

PagedAttention (vLLM), 722, 732, 779

FlexDecoding, 803

parallel aggregated tree (PAT), 130

parallel filesystems, 174

parallel thread execution (see PTX (parallelthread execution))parallelismadaptive strategies for ultrascale inference,

parallelarchitecture of GPU, 191

cuTile, 59

Paralleldynamic parallelism, 506-512

benchmarks, 511

child-kernel launch limit, 507

device memory usage monitored, 507

profiling before and after a change, 512

“stack overflow” errors avoided, 507

when to capture CUDA Graph instead,

fully sharded data parallel, 133

PTX, 365

executionspeculative decoding, 665, 681-689

strategies, 670-681

context parallelism, 670, 678

data parallelism, 670, 677

expert parallelism, 670, 674-677

hybrid parallelism, 679-681

model weights and data split over GPUs,

pipeline parallelism, 670, 673

profiling, 680

summary table, 670

tensor parallelism, 670, 672

|

thread block cluster algorithms, 416-418

parameters offloaded to CPU and NVMe, 573

partitioned global address space (PGAS), 38

PAT (parallel aggregated tree), 130

PCIe, 465-466

PDUs (power distribution units), 44

peer-to-peer (see P2P (peer-to-peer) DMA)per-thread default streams (PTDS), 442

perf (Linux), 528, 530

example MoE model, 539-543

NVIDIA Performance Monitoring Unit, 542

Perfetto Trace Viewer, 528, 530, 591

performance boostscompressing data, 176

CUDA Graphs, 494

conditional graph nodes, 506

disaggregating prefill and decode, 665

dynamic batching, 718-720

high occupancy with high ILP, 333

kernel optimization, 303-305

max-autotune compile mode, 606

memory pinning, 71

network reads eliminated, 176

pitfall of repeatedly reading same globalmemory data, 255

running many warps concurrently, 221

stream 0 legacy default stream avoided, 442

TORCH_LOGS performance hints, 623

performanceperformance metricsabout ultimate goal of performance tuning,

automated performance tests, 4

benchmarking and profiling, 4

(see also benchmarks; profiling)cache hits and misses, 703

bottlenecksdisaggregation of prefill and decode, 666

goodput as throughput metric, 16

GPU idle time measurement, 177

GPU utilization percentage, 729

GPUs with NVIDIA Container Toolkit, 88

inference tuningerror handling, 731

full-stack optimizations, 713-715

|

Nsight Compute Roofline analysis, 290

RPS (requests per second), 719

thread block size, 200

TPOT (time per output token)disaggregation of prefill and decode, 666

prefill and decode phases, 759

TTFT (time to first token)constrained outputs, 690

context parallelism, 679

disaggregation of prefill and decode, 666

prefill and decode phases, 759

performance regression continuous integration,

592-598persistence mode for GPU, 77

nvidia-persistenced daemon, 77

recommended when using MIG, 95

persistent kernels, 391-395

when to combine, 401

megakernels, 394

occupancy, 393

warp specialization, 395

persistent threads (see persistent kernels)

PGAS (partitioned global address space), 38

pipeline parallel (PP), 670, 673

monitoring system metrics and counters,

705-709observe-hypothesize-tune loop, 702-705

profiling with Nsight, 709-712

troubleshooting, 712

I/O continuous profiling and tuning,

184-188latency hiding as important tool, 32

MLPerf Logging, 595-598

performance benchmarking, 592-598

capturedperformance monitoring, 47

GPUs near 100% utilized, 47

NVLink usage, 47

prefix-merge metrics, 703

PyTorch performance heads-up display,

594-595roofline model, 233-237

common workloads for, 393

cooperative groups via grid sync, 398-401

FSDP with, 577

hybrid parallelism, 681

PyTorch compiler with, 589

tensor parallel versus, 674

versus TP versus hybrid, 843-847

pipelinescomparison of no pipeline with pipeline,

107-112

CUDA C++ Pipeline API, 279-282

hardware-software codesign, 282

CUDA Graphs, 492

(see also CUDA Graphs)data compression, 176

data pipeline tuning, 178

data loading and preprocessing, 178-181

NeMo Curator for training datasets, 183

NVIDIA Data Loading Library, 182

scaling out workers as GPUs scaled, 181

distributed data parallel strategies, 132-137

inter-kernel pipeliningabout inter-kernel pipelining, 429

compute overlapping with data transfersvia streams, 433-435, 463-469,

556-558cross-stream synchronization, 448

allocator

allocator

allocatordefault stream usage, 444-446

default versus explicit streams, 443

kernel execution overlapped with CUDAstreams, 429-433

launching five kernels on two CUDAstreams, 429-431

legacy default streams, 441

multi-GPU overlap of compute and data,

463-469

PDL with thread block clusters and warpspecialization, 472-478

per-thread default streams, 442

Programmatic Dependent Launch,

callbackswarp specialization and CUDA streams,

449-456warp specialization replaced with CUDAstreams, 431-433

clusters

kernelswhen to combine persistent kernels andcooperative groups, 401

when to use warp specialization versusdouble-buffering, 385

monitoring I/O, 177-178

computationreducing frequency and volume, 107

streams for asynchronous execution, 106

Triton double-buffering, 653-656

PMU (NVIDIA Performance Monitoring Unit),

PoC (point of coherency), 217

POD-Attention, 810

Podman with NVIDIA Container Toolkit, 89

point of coherency (PoC), 217

portable cluster size maximum, 406

power and thermal constraints on inference,

power distribution units (PDUs), 44

power monitoring, 47

setting slightly below TDP, 84

power requirements

GPU persistence mode, 78

NVL72, 43

liquid cooling versus air cooling, 45-46

power monitoring, 47

power-saving mode for CPUs, 75

PP (see pipeline parallel (PP))

intra-kernel pipeliningabout intra-kernel pipelining, 374

comparing techniques, 374, 389, 421-425

cooperative groups, 402

(see also cooperative groups (CGs))double-buffering with Pipeline API,

375-380persistent kernels, 391-395

PyTorch, 390

thread block clusters, 373

(see also thread block clusters)warp specialization, 381-385

warp specialization and CUDA streams,

NVIDIA Management Library, 300

power limits, 47, 730

|

#pragma unroll, 326, 333

unrolling too many iterations, 333

precisiondynamic precision changes, 848-853

FP4, FP8, FP16, 29, 356

inference quantization, 734

weight-only quantization, 735

mixed precisionautomatic mixed precision, 355-356

reduced precision, 356

Tensor Cores, 351-360

quantization (see quantization)

Predibase, 929-931

predication, 317

warp divergence avoided, 317

warp divergence minimized, 319-323

prefillabout prefill phase, 759

chunked prefill, 723, 893-899

context parallelism, 670, 678

Context Parallel (PyTorch), 554

disaggregated from decodeabout disaggregation, 663, 665, 760

architecture, 664

batch size optimizations, 893

cluster pools, 765-780

deploying with Kubernetes, 668-670, 714

latency and throughput, 666

load balancing, 834-841

prefill to decode data transfers, 667,

812-817scaling, 665

separate prefill and decode stages,

150-152, 664

disaggregated routing policies, 780-795

capacity-aware routing, 786-788

configurationexample dynamic routing policy in code,

latency-aware routing, 788-791

multipath inference (racing), 791

QoS and early rejection policies, 792-795

routing factors, 781-784

speculative decoding, 791

FlashAttention, 15

hardware and parallelism, 820-826

hybrid prefill with GPU-CPU collaboration,

826-828

|

advantages of disaggregation, 762-765

prefill workers, 768-771

routing and scheduling policies, 780-795

scalability, 795

prefill-decode interference, 665

prefix caching, 743-749

SGLang RadixAttention, 746-747

prewarming graphs and caches, 889-893

PrimTorch Intermediate Representation (IR),

FX graph output, 614

PyTorch compiler pipeline, 606

process-GPU mapping, 902-904

profiling, 4

KV cache transfer to decode, 812-817

connector and data path design, 817-820

KV cache size, 812

zero-copy GPU-to-GPU, 813-817

optimized kernels via FlexDecoding,

802-805scale-out of workers, 891

why disaggregation, 761-796

achieved occupancy, 297-305

kernel compute throughput versus GPU

FLOPS, 299

kernel memory throughput versus HBMmemory bandwidth, 299

before and after a change, 512

compiler performance issues, 552

CUDA events for, 449

DataLoader, 181

example mixture-of-experts model, 531-543

about MoE models, 531

Linux perf on CPU and GPU, 539-543

Nsight Systems and NVTX Timelines,

535-537

PyTorch profiler, 532-535

roofline analysis for GEMM kernels,

537-539

GPU bottlenecks, 288-293

utilization

CUDA Profiling Tools Interface, 291

iteratively profiling, 301-303

kernel optimization, 303-305

Nsight Compute, 288

Nsight Compute and Roofline analysis,

Nsight Systems, 288

Nsight Systems timeline view, 288

PyTorch profiler via Kineto, 291-293

warp stall reasons, 293-297

graph break recompilations monitored, 609,

Holistic Trace Analysis, 591

inference, 709-712

production environment, 712

I/O continuous profiling and tuning,

bottleneckskernels generated by TorchInductor, 553

latency with streaming enabled, 755

memory via PyTorch profiler, 569-582

activation checkpointing, 572

CUDA memory allocator tuned, 570

FSDP automatic checkpointing, 574-577

FSDP with tensor parallel and pipelineparallel, 577

topologiesoffloading parameters, 573

peer-to-peer DMA, 581

pluggable memory allocators, 578-581

multi-GPU profiling

HTA, 591

Nsight Systems, 902

NCCL, 145

NCCL logs, 147

Nsight Compute occupancy limiters, 297

Limited by Registers, 333

Nsight Compute warp stall reasons, 293-297

Stall: Compute Unit Busy, 295

Stall: Exec Dependency, 295, 334

Stall: Idle, 295

Stall: Long Scoreboard, 293

Stall: Math Pipe Throttle, 295

Stall: Memory Dependency, 295

Stall: Memory Throttle, 294

Stall: No Eligible, 295

Stall: Not Selected, 294

Stall: Short Scoreboard, 294

Stall: Texture Throttle, 295

stalling for other reasons, 295-297

NVTX markers, 528

profiling tools that use markers, 528-531

occupancy tuning, 307

parallelism strategies, 680

tools for profiling, 528-531

Triton Proton Profiler, 657

warp divergence, 318

workload under insufficient bandwidth, 122

Programmatic Dependent Launch (PDL; NVI‐DIA), 469-472

specializationprogrammingautomated code optimization, 934-936

CUDAasynchronous memory allocation,

210-212

CUDA streams, 211

flow between CPU and GPU, 191

__global__, 203

GPU memory hierarchy, 212-218

high occupancy and GPU utilization,

221-229highest-level and most-recent APIsavailable, 283

kernel errors, 206

kernel inputs in 1D, 210

kernel inputs in 2D and 3D, 209

kernels for parallel work, 203-206

launch parameters, 207-209

maximum threadsPerBlock compile timeparameter, 229-231

memory pools, 210-212

occupancy tuning with

__launch_bounds__, 229-231

roofline model, 233-237

Unified Memory, 218-221

instruction-level parallelismcompiler hints, 333

interleaving independent operations, 332

unrolling loops, 332, 344

kernel fusion, 344-347

loop unrolling, 332, 344

MMA API, 358-360

multiple GPUs, 38-39

NVIDIA Compute Sanitizer, 231

NVLink-C2C CPU to GPU memory access,

parallelizing workloads with threadblockclusters, 416-418

PTX, 352, 365

code from compiler, 58, 203

|

compatibility across GPU generations,

58, 203

inline PTX code, 365-369

inline PTX code in DeepSeek DeepEP,

streaming assembler, 365-369

streaming assembler changing acrossgenerations, 367

Triton, 640

Python-centric CUDA libraries, 638

smart compilers, 934-936

Triton programming model, 638

Prometheusinference monitoring, 705, 706

alert manager, 716

Data Center GPU Manager metrics, 706

NVLink/NVSwitch monitoring, 902

prompt cleansing, 741

CTRL approach, 743

prompt compression, 740

prompt servers, 768

prompts extremely longcontext parallelism, 678

prompt compression, 740

stall-free scheduling, 723

pruning, 347

dynamic token pruning, 919

TorchAO, 555

PTDS (per-thread default streams), 442

PTX (parallel thread execution), 365, 616

code from compiler, 58, 203, 616

generations

JIT compiling, 59

inline PTX code, 365-369

optimizationstreaming assembler, 365-369

generations

TMEM and CUTLASS, 352

TorchInductor, 616

Triton, 640

ptxas warning, 229

Pythonconditionals as tensor operations, 611

CUDA kernels created, 56

CUDA libraries, 638

CUDA Toolkit libraries, 59

|

(see also PyTorch)

Nsight Systems backtrace sampling, 292

Triton language and compiler, 60

PyTorcharithmetic intensity, 350

attention mechanisms optimized, 553

automatic mixed precision, 355-356

backend support, 119

numpy replacement cuPyNumeric, 59

frameworks built on CUDA, 60

verifying which backend, 121

compiler (see PyTorch compiler)

CUDA Graphs, 490, 493

static memory pools, 491

custom kernel registration, 641-643

DataLoader

CPU affinity for each worker process,

64-69data input pipeline optimization,

583-585data loading and preprocessing, 178-181

multiple workers, 167, 178, 181, 583

persistent workers, 181

pinned memory flag, 71, 434, 583

prefetch factor, 584

Distributed Data Parallel communicationbucketing, 107

data parallel, 132

Data Parallel versus Distributed Data

Parallel, 132-137

definition of Data Parallel, 132

definition of Distributed Data Parallel,

model parallel, 132

DistributedSampler, 164

filesystem caching, 75

GPU idle time measurement, 177

intra-kernel pipelining, 390

introduction, 60

compiler stack, 60, 543, 605

CUDA complexity abstracted away, 61

tensor operations using GPUs, 60

TorchInductor backend, 60

memory allocator best-fit with coalescing,

memory allocator plugin, 578-581

memory caching allocator, 211

memory profiler, 528, 530

NCCL collectives running on dedicatesstreams, 464

nested tensors, 351

NVLink domain, 14

persistent kernels and megakernels, 394

profiler (see PyTorch profiler)sparse format from trained dense models,

streamscreating, 555

default versus explicit streams, 443

multiple GPUs with multiple streams,

transferprofiling when adding streams, 558

symmetric memory, 582

PyTorch Architecture Optimization (torchao),

pruning, 555

quantization, 555

INT8 quantization, 358

sparsity, 555

PyTorch compiler (torch.compile), 543-553

about the compiler, 60, 543

arithmetic intensity, 350

AsyncTP enabled, 590

compilation pipeline, 606-627

about the compilation pipeline, 606

AOT Autograd, 612

PrimTorch IR, 613

TorchDynamo, 607-612

TorchInductor, 615

compiler modes, 549-551

CUDA Graphs triggered, 618

eager mode, 623

max-autotune mode, 606

static shapes required for CUDA Graphs,

618, 622

compiler stack, 60, 543, 605

compiling versus writing custom kernels,

CUDA Graph Trees, 568

debugging, 623, 627, 636

numerical correctness and accuracy,

624-627

TORCH_LOGS, 552, 623

disabling, 623

eager mode, 623

profiling and benchmarking kernelsgenerated, 553

Transformer Engine, 619

using, 544-548

vectorized memory operations, 248

PyTorch DataLoader (see DataLoader

(PyTorch))

Distributed

PyTorch Fully Sharded Data Parallel (see FSDP

(Fully Sharded Data Parallel; PyTorch))

PyTorch Holistic Trace Analysis (see Holistic

Trace Analysis (HTA; PyTorch))

PyTorch Kineto, 291, 528, 529

NCCL profiler plugin, 145

PyTorch Mega-Cache, 544

PyTorch profiler

CUDA Performance Tools Interface, 291

example mixture-of-experts model, 532-535

GPU bottlenecks, 291-293

Kineto library, 291, 528, 529

dynamic shapes, 620-623

export, 615

kernel fusion, 313, 344, 610

launch overhead reduced via CUDA

Graphs, 491

Mega-Cache, 544

per-thread default stream enabled, 442

performance hints, 623

pipeline parallel with, 589

performancequantization with torchao, 555

scaling with DDP, 586

scaling with FSDP, 586-589

SimpleFSDP, 590

smart compilers, 934-936

tensor parallel with, 589

TorchInductor, 245, 544, 545

NCCL profiler plugin, 145

memory profiler, 528, 530

memory profiling and tuning, 569-582

activation checkpointing, 572

CUDA memory allocator tuned, 570

FSDP automatic checkpointing, 574-577

FSDP with tensor parallel and pipelineparallel, 577

topologies

|

offloading parameters, 573

peer-to-peer DMA, 581

pluggable memory allocators, 578-581

NCCL profiler plugin, 145

Nsight Systems PyTorch focused mode, 292

PyTorch TorchDynamo (see TorchDynamo

(PyTorch))

PyTorch TorchEval (see TorchEval (PyTorch))

PyTorch TorchInductor (see TorchInductor

(PyTorch))

PYTORCH_CUDA_ALLOC_CONF, 86, 435

Q

QoS (quality of service), 831

dynamic routing policies, 785

early rejection policies, 792-795

quad-die Rubin GPU, 51

GPC “partitions”, 403

quality of service (see QoS (quality of service))quantizationdynamic activation quantization, 919

inference optimization, 733-739

about quantization, 733

activation quantization, 737

quantizationpost-training quantization workflow,

quant-dequant steps into executiongraph, 739

reducing precision, 734

weight-only quantization, 735

KV compression, 867

TorchAO, 555

Quantum InfiniBand (NVIDIA), 43

R

R300 Rubin Ultra GPU, 51

RadixAttention (SGLang), 723

prefix caching, 746

random versus sequential reads, 165-166

RDMA (remote direct memory access), 38,

114-125

DeepSeek Fire-Flyer File System, 173

GPUDirect RDMA, 38, 115

multinode, multirack communication,

908-911

NCCL’s native allocator in PyTorch

MemPool, 579

|

verify true GPUDirect RDMA, 115

InfiniBand support, 115

zero-copy GPU-to-GPU transfer, 813-817

RDMA over Converged Ethernet (RoCE), 38

read-only data caches, 273-278

const __restrict__, 273, 276, 277

pitfall of not stating read-only, 273

reading data

DeepSeek Fire-Flyer File System, 172-174

network reads eliminated, 176

read ahead, 166

sequential versus random pattern, 165-166

filesystem optimized, 166

read size tuned, 165

recomputation versus memory trade-off, 350

reduced precision, 356

registers

GPUDirect RDMA, 38

memory hierarchy, 32

occupancy trade-off, 231

register pressure from ILP, 333

register spilling, 230

reinforcement learning (RL)agents tuning AI at runtime, 875

GPU kernel optimization via, 929-931

remote direct memory access (see RDMA

(remote direct memory access))reproducibility, 6, 950-952

request coalescing, 755

requests per second (RPS), 719

resources onlinebook supplemental material, xxiiibook web page, xxiv

FireworksAI CUDA Graphs post, 566

GPU failures at scale per Meta, 143

NCCL environment variables, 137

PyTorch overlapping computation and datatransfer, 557

ring communication algorithm, 129

RL (see reinforcement learning (RL))

RoCE (RDMA over Converged Ethernet), 38

enabling GPUDirect RDMA, 96

Roofline analysis (Nsight Compute), 290

example MoE model, 537-539

roofline model, 233-237

roofline-guided scheduling, 521

Roofline performance model, 337

arithmetic intensity, 337-338, 521

lower precision with Tensor Cores, 351

RPS (requests per second), 719

Rubin Ultra GPU (R300), 51

runtime optimization, 714

S

SASS (see streaming assembler (SASS))

Reductionarchical Aggregation and Reduction

Protocol; NVIDIA))scaled dot product attention (SDPA; PyTorch),

scaling, 5

disaggregated prefill and decode, 795

prefill and decode workers, 891

n-GPU scaling pattern, 520

PyTorch distributed libraries, 585-592

DDP with torch.compile, 586

FSDP with torch.compile, 586-589

scheduler and interrupt affinity, 73

schedulingbutterfly schedules, 692

cluster scheduler, 47

multipleadaptive process-GPU mapping, 902-904

dynamic congestion-aware scheduling,

913-915

GPUDirect RDMA, 908-911

MoE expert rebalancing and regrouping,

communication

NCCL collective communication,

905-908

NVSwitch and fine-tuned scheduling,

915-917telemetry and monitoring, 901

continuous scheduling, 721

dynamic batching merged with, 722

dynamic scheduling and load balancing,

congestion-awaredynamic scheduling with atomic workqueues, 482-489

atomic counters, 483

atomic queues, 485-489

atomic queues allowing dynamic workallocation, 483

Nsight Compute, 483-484

first in, first out (FIFO) scheduler, 724-725

in-kernel persistent scheduling, 501

iteration-level scheduling, 720

Kubernetes job scheduling, 93

latency-aware scheduling and dynamicrouting, 724-725

Linux schedulerscompletely fair queueing obsolete, 166

Completely Fair Scheduler, 73

multiqueue block I/O scheduler, 166

none scheduler, 166

setting for, 166

roofline-guided scheduling, 521

sharing workloads, 47

stall-free scheduling, 723

thread block scheduling, 403

wave scheduling of NCCL collectives, 907

scratch memory, 406

SDPA (scaled dot product attention; PyTorch),

sequence parallelism (see context parallelism

(CP))sequential versus random reads, 165-166

SFU (see Special Function Unit (SFU))

SGLangdynamic batching merged with continuousscheduling, 722

RadixAttention, 723

megakernel decode throughput, 394

continuousprefix caching, 746-747

speculative decoding, 683

shardscombining small files into large shards, 165

DeepSeek Fire-Flyer File System shardingmetadata, 173

sharding data across nodes, 164

shared filesystems, 174

shared memory (SMEM)avoiding shared memory, 271-273

avoiding shared-memory bank conflicts,

264-271padding shared arrays, 267

Blackwell GPU shared memory, 269

|

data reuse, 255-264

distributed shared memory, 341

DRAM as, 255

dynamic shared-memory allocation,

858-862occupancy balanced with resource usage,

PyTorch, 350

tiling, 255-264, 339

CUTLASS, 360

Triton programming, 641

__shared__ regions, 405

sharing and scheduling workloads, 47

SHARP (Scalable Hierarchical Aggregation and

Reduction Protocol; NVIDIA), 40-41

NCCL and, 146

offloading aggregation operations to switch,

__shfl_sync, 271

shifted all-to-all schedule, 692

silicon photonics, 43

Simple Linux Utility for Resource Management

(see SLURM (Simple Linux Utility for

Resource Management))

SimpleFSDP, 590

SIMT (single instruction, multiple threads)model, 32, 638

simultaneous multithreading (SMT), 69

single instruction, multiple threads (SIMT)model, 32, 638

single-kernel tiling, 646-653

single-program multiple-data (SPMD), 638

Singularity, 91

sleep state for CPU, 76

slicing a GPU (see MIG (Multi-Instance GPU;NVIDIA))

Slinky, 93

SLO-aware request managementearly rejection, 829-831

fault tolerance, 832

quality of service, 831

SLURM (Simple Linux Utility for Resource

Management), 42

cluster scheduler, 47

job scheduling, 93

SM-resident limits, 202

SmoothQuant, 737

SMs (streaming multiprocessors), 31-33

|

latency hiding, 32

warp schedulers, 327

SMT (simultaneous multithreading), 69

softwarehardware-software codesign, 15-16

automatingsoftware stack, 56

nvidia-smi, 57

PyTorch introduction, 60

sparse modelsmixture-of-experts as, 9, 12

reducing effective compute requirements, 9,

sparsityoptimized sparse GEMM kernels, 348

pruning, 347

inference workloads, 349

TorchAO, 555

SpeCache, 862

Special Function Unit (SFU), 194

Spectrum-X800 Ethernet, 43

speculative decoding, 665, 681-689

maximum

GPU architecture, 192-194

threads and registers per SM, 200

warps, 192-194

high occupancy, 199

programming CUDA, 221-229

minimum thread blocks resident, 229-231

SM-resident limits, 202

Tensor Cores, 29-31

threads executed in warps, 32

CUDA Pipeline API plus TMA, 282

DeepSeek Multi-Head Latent Attention,

about, 61

CUDA kernels created in Python, 56

CUDA Runtime and Toolkit, 57

Toolkit C++ and Python libraries, 59

GPU driver, 57

arithmetic intensity increased, 349

structured sparsity, 347-349

combining techniques, 688

interleaving decode steps, 688

Medusa multiple heads, 682

multitoken decoding, 686-688

single-model self-speculative decoding, 686

two-model draft-based, 682-685

EAGLE algorithm, 684

speculative KV prefetching, 862

speculative MoE expert routing, 918

SPMD (single-program multiple-data), 638

stall-free scheduling, 723

static batching, 718

static indexing replaced with atomic-drivenwork queue, 489

storage communication, 41

storage I/O optimizationscloud storage caches, 174

bottleneckscompressing data, 176

continuous profiling and tuning workflow,

184-188data pipeline tuning, 178

data loading and preprocessing, 178-181

NeMo Curator for training datasets, 183

NVIDIA Data Loading Library, 182

scaling out workers as GPUs scaled, 181

DeepSeek Fire-Flyer File System, 172-174

fast storage and data locality, 163

monitoring I/O, 177-178

NFS servers, 174

tuning parameters, 175

NVIDIA GDS, 167-172

NVMe and filesystem tuned for throughput,

multiqueue block I/O scheduler, 166

object storage, 175

sequential versus random reads, 165-166

filesystem optimized, 166

read size tuned, 165

shared filesystems, 174

striping files, 175

tuning filesystems, 176

straggler effect MoE models, 676

stream 0, 441

streaming assembler (SASS), 365-369

changing per GPU architecture generations,

streaming multiprocessors (see SMs (streamingmultiprocessors))streaming responses, 752-755

streams for asynchronous execution, 106

(see also CUDA streams)strided indexing, 240

striping files, 175

structured outputs, 689

structured sparsity, 347-349

surface object APIs, 277

swizzling, 270

thread block swizzling, 404

symmetric memory, 582

SymPy, 621

about dynamic shapes, 620

__syncthreads(), 396

systemctl, 78

inferenceerror handling, 731

GPU utilization versus latency, 729

KV cache, 732

memory, 732

computationpower and thermal constraints, 730

Ttcmalloc, 76

TCMAL‐LOC_MAX_TOTAL_THREAD_CACHE_BYTES,TCMALLOC_RELEASE_RATE, 77

TDP (thermal design power), 84

temporal load balancing, 913

Tensor Cores, 29-31

enabling, 355

theoretical throughput of NVL72 rack, 29

TMEM, 352-354

Transformer Engine, 358

very high active utilization, 295

Tensor Memory Accelerator (TMA), 353

capabilities, 279

memoryglobal memory traffic reduced, 413-416

hardware-software codesign, 282

starting a TMA transfer, 278

tensor parallel (TP), 670, 672

CUTLASS, 360-364

mixed precision, 351-360

MMA API, 358-360

TF32, 355

|

AsyncTP, 590

FSDP with, 577

hybrid parallelism, 681

pipeline parallel versus, 674

PyTorch compiler with, 589

versus PP versus hybrid, 843-847

TensorBoard trace viewer, 591

TensorRT-LLM (NVIDIA), 490

TetriInfer’s two-level scheduler, 835

texture object APIs, 277

TF32, 355

enabling, 355

TFLOPS (teraFLOPS), 255

thermal constraints on inference, 730

thermal design power (TDP), 84

Thompson, Martin, 15

THP (transparent hugepages), 72

thread block clusters, 196, 373

cooperative groups contrasted with, 402

described, 402

distributed shared memory, 402-413

algorithms for parallelizing workloads,

416-418coordinating thread block clusters,

408-410global memory traffic reduced, 413-416

launching a thread block cluster, 407,

scratch memory, 406

sharing state between thread blocks, 405

thread block clusters described, 373

thread block pair, 410-413

thread block swizzling, 404

threads accessing each other’s sharedmemory, 196, 373, 404

threads using thread block clusters and

DSMEM, 404

warp specialization with thread blockclusters, 418-425

intra-kernel pipelining with inter-kerneloverlap, 472-478

launching a thread block cluster, 407, 418

sharing state between thread blocks, 405

thread block pairs, 407, 410-413

tiling, 341-343

up to 16 thread blocks, 343

thread blocks, 195-199

cooperative mode, 396

coscheduled on same GPC, 402

|

(see also thread block clusters)thread block size, 200

tiling, 255

memorycooperative groups, 395-401

distributed shared memory, 404

GPU architecture, 195-199

intrawarp communication, 271-273

programming CUDAchevron syntax (<<< >>>), 203

maximum threadsPerBlock compile timeparameter, 229-231

threadsPerBlock, 203, 205

threadsPerBlock as launch parameter,

warp divergence, 199

threadsPerBlock, 203, 205

launch parameter, 207-209

throughput metricsdisaggregation of prefill and decode, 666

FP4, FP8, FP16, 29

“goodput”, 16

NVL72 rack Tensor Core throughput, 29

distributed shared memory, 402-413

sharing state between thread blocks, 405

thread block clusters described, 373

threads accessing each other’s sharedmemory, 196, 404

threads using thread block clusters and

DSMEM, 404

distributed shared memory andthreads accessing each other’s sharedmemory, 373

programming CUDAblocksPerGrid, 203, 205

blocksPerGrid as launch parameter,

207-209minimum thread blocks resident on each

SM, 229-231

threadsPerBlock, 203, 205

threadsPerBlock as launch parameter,

207-209sharing state between thread blocks, 405

thread block pairs, 407, 410-413

207-209tiling, 255

warps, 32, 192-194

NVMe and filesystem tuned for throughput,

ThunderMLA (Stanford), 801

TileIR, 263

tilingarithmetic intensity increased, 260, 339

multilevel tiling, 340

double-buffered cooperative tiling, 375-380

multilevel tiling, 340

PyTorch, 350

shared memory for tiling, 255-264, 339

CUTLASS, 360

single-kernel tiling, 646-653

software prefetching, 341

thread block clusters, 341-343

tiles, 255

TMA fetching from global into sharedmemory, 278-283

time per output token (see TPOT (time peroutput token))time to first token (see TTFT (time to firsttoken))

TLB (see translation lookaside buffer (TLB))

TMA (see Tensor Memory Accelerator (TMA))

TMEM, 352-354

double buffering with CUTLASS, 360

Transformer Engine, 358

toolkit (see CUDA Toolkit)torch.compile (see PyTorch compiler

(torch.compile))torch.profiler (see PyTorch profiler)

TorchAO, 555

pruning, 555

quantization, 555

INT8 quantization, 358

sparsity, 555

TorchBench, 592

TorchDynamo (PyTorch), 544, 607-612

debugging numerical accuracy, 624

dynamic shapes, 553, 609

explain() output, 553, 628

all-reduce-related warnings, 589

graph breaks, 552, 589

graph breaks debugged, 628

FX Graph output, 609

dumping graph after each stage, 637

graph breaks, 611

operations

PyTorch compiler pipeline, 606

PyTorch compiler stack, 60, 543, 605

stances for errors or fallbacks, 610

TORCH_LOGS, 552

TORCHDYNAMO_REPRO_*, 636

TorchEval (PyTorch), 528, 531, 585

TorchInductor (PyTorch), 615

minifier tool, 624

kernels

PrimTorch IR simplifying operations, 614

generated

PyTorch compiler pipeline, 606

PyTorch compiler stack, 60, 543, 605

SimpleFSDP, 590

symbolic shapes, 615

Triton based, 60, 544, 616, 637

TORCHINDUCTOR_BENCHMARK_KER‐NEL, 553

TORCHINDUCTOR_UNIQUE_KER‐NEL_NAMES, 553

TorchTitan (PyTorch), 590

AsyncTP, 590

SimpleFSDP reducing memory usage, 590

TORCH_COMPILE_DEBUG, 636

TORCH_LOGS, 552

forcing error to be raised at compile, 608

minifier tool, 624

profiling and debugging compiler issues,

arithmetic intensity increased, 339

automatic mixed precision, 355-356

autotuning, 544, 617-620

max-autotune compile mode, 545, 606,

CUDA Graphs, 618

debug mode, 624, 637

inspecting TorchInductor code, 623

FX Graph, 612

introduction, 60

IR, 615

kernel fusion, 313, 344, 610

fused attention kernels, 391

GEMM kernel prologue and epilogue,

compiler debugging, 636

logging options summary, 636

performance hints, 623

|

TP (see tensor parallel (TP))

TPOT (time per output token)disaggregation of prefill and decode, 666

prefill and decode phases, 759

training 100-trillion parameter models, 11

loading training data, 13

cold start strategy, 10

costs

DeepSeek-R1, 2

DeepSeek-V3, 9

Google Gemini Ultra, 2

OpenAI GPT-4, 2

datasets via NVIDIA NeMo Curator, 183

Transformer Engine (TE; NVIDIA), 29-31

balancing accuracy and precision, 29

installing only if calling TE modulesdirectly, 619

mechanical sympathy, 16

modern GPUs using, 16

PyTorch compiler, 619

TMEM, 358

translation lookaside buffer (TLB), 72

transparency

AI systems performance engineers, 6

DeepSeek Open-Source Week, 7

transparent hugepages (THP), 72

tree communication algorithm, 129

tree parallelism communication algorithm, 130

Triton (OpenAI)advanced kernel implementations, 645-656

double-buffering pipelining, 653-656

persistent matmul kernel, 646-653

warp specialization, 645

CUDA kernels in Python, 56

custom kernels, 637-658

language and compiler, 60

open source, 60

programmingaccessing shared memory, 641

autotuning Triton kernels, 643

custom kernels, 637-658

Nsight Compute and Systems support,

profiling with Triton Proton Profiler,

programming model, 638

registering custom kernels with PyTorch,

|

617, 641-643

tuning kernel launch parameters, 643

Proton Profiler, 657

smart compilers, 934-936

TorchInductor basis, 60, 544, 616, 637

troubleshooting inference, 712

TTFT (time to first token)constrained outputs, 690

context parallelism, 679

disaggregation of prefill and decode, 666

prefill and decode phases, 759

speculative KV prefetching, 862-867

U

UCX (Unified Communication X; NVIDIA),

underutilization GPU stalls, 301-303

kernel optimization, 303-305

Nsight Compute metrics, 301

Unified Communication X (UCX; NVIDIA),

Unified CPU-GPU Memory, 24

available

NCCL release notes, 144

Nsight Systems and Nsight Compute withlatest GPUs, 293

NVIDIA continually updating libraries, 364

ROI of upgrading hardware, 48

update but verify, 187

utilization-maximization policy, 897

Vvariable-length sequences via dynamic shapes,

620-623vectorized memory access, 247-254

Vera Rubin Superchip (VR200), 50

vLLMconstrained decoding, 690

CUDA Graphs, 490

disaggregated prefill, 667, 669

dynamic batching merged with continuousscheduling, 722

PagedAttention, 722

LMCache, 703

disaggregated prefill and decode, 760,

megakernel decode throughput, 394

NIXL-vLLM integration, 667

NVLink domain, 14

PagedAttention, 722, 732, 779

parallel decoding, 688

prefix caching, 743, 809

speculative decoding, 683

troubleshooting inference, 713

VR200 (Vera Rubin Superchip), 50

W

Warp (see NVIDIA Warp)warp stall reasons, 293-297

warps, 32

efficiency via PyTorch, 323-324

efficient intrawarp communication, 323

execution efficiency improved, 314-324

GPU architecture, 192-194

high occupancy, 199

ILP increased for latency-bound GPUs, 303

latency hiding, 32

maintaining high occupancy, 221

shuffle intrinsics avoiding shared memory,

271-273stalls or frequent pauses, 293-297

execution-dependency stalls, 295

memory-related stalls, 293

warp divergence, 199, 314

avoiding, 315

causes of, 314

efficient intrawarp communication, 323

improving warp execution efficiency,

314-324predication to minimize, 319-323

profiling and detecting, 318

PyTorch considerations for efficiency,

323-324warp specialization

CUDA streams and warp specialization,

449-456

CUTLASS, 360

inter-kernel pipelining with CUDAstreams, 431-433

inter-kernel pipelining with CUDAstreams and thread block clusters,

intra-kernel pipelining, 381-385

intra-kernel pipelining with CUDA

Pipeline API, 385-390

intra-kernel pipelining with inter-kerneloverlap, 472-478

intra-kernel pipelining with persistentkernels, 395

intra-kernel pipelining with thread blockclusters, 418-425

Triton, 645

web page for book, xxivweight-splitting parallelism strategies, 671

workloadsdynamic work queues, 482-489

atomic counters, 483

atomic queues, 483, 485-489

Nsight Compute, 483-484

memory-bound workload example, 228

persistent kernels, 393

profiling under unsufficient bandwidth, 122

right occupancy for workload, 307

sharing and scheduling, 47

workloadswrite-back, 75

X

XFSoptimizing, 166

tuning for throughput, 166

XLA, 658

Z

ZeRO Stage-3 in FSDP, 574

zero-copy GPU-to-GPU transfer, 813-817

zero-copy registration, 148

activating, 658

Google Cloud TPU, 615, 658

PyTorch compiler pipeline, 606, 609, 612

|