131 docs tagged with "gpu"

Buying a bigger GPU, or more of them, does not buy a proportionally bigger speedup, and the reason has nothing to do with the hardware being slow. It has to do with the fraction of the program that was never made parallel in the first place. Two laws describe the two ways to think about that fraction — one holds the problem size fixed and asks how fast you can finish it, the other holds the time budget fixed and asks how much bigger a problem you can solve — and knowing which one describes your situation changes what "more parallelism" is even supposed to buy you.

A spec sheet lists a GPU as a pile of numbers — core count, clock speed, memory bandwidth — but those numbers only make sense once you know what physical structure they're describing. A GPU is not a bag of independent processors; it's a small number of large, warp-scheduling processors (streaming multiprocessors), each built from smaller replicated pieces, all sharing a common path out to memory. Understanding that structure top to bottom is what turns a spec sheet from marketing copy into something you can reason about.

Every kernel is limited by one of two things before it is limited by anything else: how fast the device can do arithmetic, or how fast the device can move bytes from DRAM. Which one applies is a property of the kernel's own math, computable on paper before you write a line of CUDA, and it determines almost everything about how you should spend optimization effort afterward. The roofline model is the tool that turns "which one applies" into a single number you can compute and a single plot you can place it on.

The classic tiled kernel loop looks like load → syncthreads() → compute → syncthreads() while stage i computes, stage i+1's load can already be in flight.

Some updates can't wait for a barrier — a histogram bin, a running total, a lock-free counter — because the threads touching the same location aren't at a point where syncthreads() or a group sync() even applies; they need the read-modify-write itself to be indivisible. Atomics provide that: a hardware-guaranteed sequence of read, modify, and write on a single memory location that no other thread's atomic on the same location can interleave with. What atomics don't provide is speed for free — how many threads target the same address, not how many threads issue atomics in total, is what determines whether that guarantee is nearly free or a serialization bottleneck.

A benchmark number is easy to produce and easy to produce wrong — every one of the six mistakes below yields a plausible-looking result that is actually measuring something other than the kernel's real performance. This page collects the mechanism behind each mistake and the specific fix, then closes with a checklist meant to be followed literally, not read once and approximated.

The warp-level tools covered so far — divergence handling, independent thread scheduling, the shuffle and vote intrinsics — all operate within a single warp of 32 threads. Most kernels that use shared memory need something coarser: a guarantee that every thread in the whole block, potentially many warps, has reached a point and that everything they wrote before that point is visible to everything they read after it. That guarantee is syncthreads(), and getting its rules exactly right is what stands between a working tiled kernel and one that hangs or reads garbage on some inputs and not others.

A CUDA project stops being a single nvcc invocation the moment it has more than one translation unit, a library dependency, or a need to target more than one GPU architecture, and hand-rolled build scripts get brittle fast at that point. CMake treats CUDA as a first-class language rather than a special case bolted onto a C++ build, which is what makes multi-file projects, per-architecture code generation, and linking against CUDA libraries manageable without duplicating flags across a Makefile.

A GPU's cache hierarchy looks superficially like a CPU's — an L1 per core-analog, a shared L2 behind it — but the access pattern it's optimized for is completely different. A CPU cache is tuned for one thread's temporal and spatial locality; a GPU's L1 and L2 exist to serve tens of thousands of threads issuing memory requests in 32-wide warps, and the granularity at which those requests are actually served is the fact that explains coalescing, wasted bandwidth, and most of what looks like "mysterious" memory performance on a GPU.

> looks like two arbitrary integers, but each one is a real decision with hardware consequences: block size determines how a block's resources are packed into an SM's fixed budgets, and grid size determines how evenly the total work spreads across the GPU's SMs. Picking both well is mostly a matter of a few rules of thumb plus one API that does the hardware-limit arithmetic for you.

NVIDIA's math libraries are tuned per architecture by engineers with access to the SASS scheduler, the microarchitecture team, and hardware that hasn't shipped yet. As a rule of thumb rather than a measured benchmark — the actual gap varies by shape, precision, and architecture generation — a hand-written GEMM that reaches something like 60% of cuBLAS's throughput on the same shapes, on the same architecture, is realistically a good hand-written GEMM. The remaining gap is instruction scheduling, tile-size search, and register allocation tuned per compute capability by people who do nothing else. Programming Tensor Cores makes this same point about wmma kernels specifically; this page generalizes it to the whole library landscape and gives a rule for when to reach for one, and when not to.

Every page in this folder makes its own case, and none of them tells you which one to actually pick — that decision depends on facts about your project that no single page can know. This page collects those facts into four questions, a decision table built from them, and one default answer for the common case of not having a second target yet.

The NCCL page covers the API — communicators, the collective calls, stream integration, grouped calls. This page covers what happens underneath an ncclAllReduce call: the ring and tree algorithms NCCL chooses between, the cost model that explains why ring all-reduce scales the way it does, and how a training loop overlaps communication with compute instead of paying for it serially. It doesn't repeat the API surface — link there for ncclCommInitRank, the collectives table, or ncclGroupStart/ncclGroupEnd.

The rest of this section explains how to do things correctly; this page is the shorter list of specific ways not to. Each entry names a symptom, the mechanism that causes it, and the fix — a checklist to scan before believing a benchmark or a kernel is done, not a tutorial.

Every deployment path covered so far in this folder eventually hands a graph to something that turns it into device code, and up to now that "something" has mostly been a fixed toolkit: TensorRT's builder, OpenVINO's plugins, a vendor NPU SDK. This page steps back and looks at the compilers underneath those toolkits and behind the frameworks themselves — what a graph compiler actually does that a kernel library doesn't, and the handful of stacks (XLA, TVM, MLIR, and torch.compile) that show up across nearly every deployment target in this section.

Compute capability is the single number that determines what a piece of CUDA code can assume about the GPU it runs on — which instructions exist, which tensor-core precisions are available, how big a thread block cluster can be. Getting the build flags around it wrong is one of the most common ways a CUDA binary that worked on the machine it was built on fails, silently or loudly, on someone else's GPU.

Global memory's performance rules assume threads in a warp want different addresses and reward spreading them out into distinct sectors. Constant memory and the read-only data path invert that assumption: they're fast precisely when a warp's threads all want the same data, and the further a kernel's access pattern gets from that, the less either buys over an ordinary global read.

Warp-Level Primitives and Block Synchronization both work, but both lean on implicit context there is no implicit lockstep assumption left to break, because every operation states which threads it applies to.

A modern laptop, phone, or server node contains all three of these, and they are not three points on a speed scale. They are three different answers to the question "how much of this chip should be general-purpose?" The CPU keeps every option open and pays for it in area and energy. The GPU gives up per-thread cleverness to buy arithmetic width, but stays fully programmable — you can still write an arbitrary kernel. The NPU gives up general programmability as well, hard-wiring a small set of tensor operations at fixed precisions, and gets back an energy-per-operation figure neither of the others can approach.

CUB is the layer Thrust is built on and Choosing a Library already named as the tuned building block for reductions, scans, and sorts: a template library of GPU primitives available at three different scopes, chosen depending on whether the call site is host code, a whole kernel, or a handful of cooperating threads inside one. Where Thrust replaces a kernel you'd otherwise write, CUB is what you reach for when you're still writing the kernel yourself and want a tuned, per-architecture primitive as one piece of it.

cuBLAS is NVIDIA's implementation of the BLAS (Basic Linear Algebra Subprograms) interface on the GPU: vector-vector, matrix-vector, and matrix-matrix operations, including the GEMM that Choosing a Library and Programming Tensor Cores both treat as the target hand-written kernels are measured against. The API is small and stable — a handle, a stream, and a handful of call shapes — but it inherits one convention from Fortran BLAS that trips up nearly everyone writing C or C++ against it for the first time.

A single kernel launch costs the CPU roughly 3–10 µs of driver-side work, independent of how much the kernel actually does. A pipeline that issues 50 small kernels per iteration can spend more time launching work than the GPU spends computing it, and Streams and Concurrency doesn't fix that — streams reorder and overlap launches, they don't reduce their count. A CUDA graph captures a whole sequence of operations once and replays it as a single launch, collapsing 50 dispatches into one.

Two different Python packages both answer to "CUDA in Python," and confusing them is the first mistake most people make. NVIDIA's cuda-python is a thin, official binding to the driver and runtime APIs — the same calls this section has been making in C++, now callable from Python. CuPy is a third-party, NumPy-compatible array library built on top of those bindings. Almost nobody wants the first one directly; almost everybody wants the second.

A kernel that reads garbage, writes out of bounds, or produces different output run to run is a different kind of problem from a slow one, and Nsight Compute is the wrong tool for it — a profiler reports how fast something ran, not whether it was correct. cuda-gdb steps through device code the way gdb steps through host code; Compute Sanitizer is a family of runtime checkers that catch specific classes of memory and synchronization bugs without stepping through anything at all.

cuDNN is NVIDIA's library of tuned primitives for deep learning why the first iteration on a new input shape is slow, why results can differ slightly between runs, and why changing a batch size can suddenly change performance by more than the batch size alone would predict.

Beyond dense linear algebra and deep learning, four more CUDA libraries cover the numerical building blocks that show up constantly but rarely justify a hand-written kernel: Fourier transforms, random number generation, sparse linear algebra, and dense/sparse factorizations. Each has its own handle type and its own lifecycle, but — as the closing section here makes explicit — they share more structure with each other, and with cuBLAS, than the four separate APIs first suggest.

CUTLASS is a C++ template library for building GEMM (and convolution) kernels with cuBLAS-class performance out of composable, reusable pieces — tile shapes, memory-movement stages, and epilogues — rather than a single pre-built call. cuBLAS is fast but fixed: cublasSgemm computes C = alpha op(A) op(B) + beta * C and nothing else, in the layouts and precisions it was built for. CUTLASS exists for the shapes and fusions that fall outside that fixed surface — an unusual data type, a custom epilogue, a problem size cuBLAS's kernel selection handles badly — while still generating code tuned close to cuBLAS's own throughput on the same hardware.

Splitting a training job across GPUs means choosing what gets partitioned — the data, the model's layers, or the operations inside a single layer — and each choice trades communication volume against memory savings differently. This page covers the communication mechanics of each strategy: what crosses the interconnect, when, and how much. Distributed Training covers the training-side recipe built on top of these mechanics — ZeRO/FSDP sharding, optimizer state placement, gradient accumulation — and is the page to read for how a framework actually configures and combines them.

Every other page in this folder covers one piece of the deployment problem in depth — a device family, a runtime, a compiler, a quantization technique. This page is the one that ties them together into an order of operations: which decisions to make first, what to check before committing to a target, and what to verify before calling a deployment done. It is deliberately a procedure rather than a survey — the pages it links to already carry the depth, and repeating that depth here would only get it out of sync with the pages that own it.

A host process can see more than one GPU, and CUDA never guesses which one a call should target — it operates against whichever device is current on the calling thread, a piece of state the program has to set itself. Getting this wrong doesn't usually crash; it silently allocates memory or launches kernels on the wrong GPU, or leaves a multi-threaded program with each thread quietly disagreeing about which device it's using. See Error Handling and Checking for what CUDA_CHECK does with the status these calls return.

Every figure in Arithmetic Intensity and the Roofline Model ultimately rests on one number: how many bytes per second a kernel can move between the SMs and DRAM. That number is set by the physical memory technology soldered onto (or stacked next to) the GPU die, and it is very different from the number a datasheet advertises. This page covers what GDDR and HBM actually are, why achieved bandwidth always falls short of the peak figure, and how to measure the one that actually matters for a given kernel.

Shared Memory is scoped to a single block the cluster's combined on-chip shared memory, addressable across block boundaries, without routing through global memory at all.

Most kernels launch from the host with a grid size chosen before any device-side work has happened, which is a poor fit for problems whose parallelism is only known once the GPU has started computing — a mesh that needs refining only in some regions, a tree whose branching factor varies by node, a search whose frontier grows unpredictably. Dynamic parallelism lets a kernel launch further kernels directly from the device, so the grid for the next phase can be sized from data the first phase just produced, without a round-trip through the host.

"NPU" on a phone, laptop, or camera SoC does not name one architecture — it names a family of fixed-function inference engines from different vendors, each with its own SDK, its own supported precisions, and its own idea of how much of the operator set it covers. What Is an NPU covered why this class of hardware exists at all; this page is the vendor-by-vendor reference for the ones you're actually likely to target.

A CUDA API call that fails almost never fails where the mistake happened. Because most of the runtime is asynchronous, a kernel launch or a copy can return immediately with cudaSuccess on the host side while the actual work — and the actual error, if there is one — hasn't executed on the GPU yet. Unchecked, that error surfaces as a failure on some unrelated call several lines or several function calls later, which is why every other page in this section wraps runtime calls in a checking macro rather than trusting a bare return value.

A cudaEventt is a marker that can be dropped into a stream and later queried, waited on, or used to measure elapsed time between two points — it's the mechanism behind both accurate kernel timing and dependencies between streams that don't require the host to get involved. Both uses matter here: naive host-clock timing of GPU work produces numbers that look plausible and are wrong, and coordinating streams without a host round-trip is what makes the concurrency from Streams and Concurrency composable into a real pipeline. See Error Handling and Checking for what CUDACHECK does with the status these calls return.

This page has not been written yet.

Every function and variable in a CUDA source file needs to answer two questions: which processor does this run on or live on, and who is allowed to call or touch it. C++ alone has no way to express that — global, device, host, shared, constant, and their relatives are CUDA's answer, and getting them right is what makes the rest of the language (templates, references, most of the standard library subset) usable across the host/device boundary at all.

Thread Indexing establishes that the fastest-varying array index should track threadIdx.x, so that adjacent threads in a warp touch adjacent addresses. This page is the mechanism that makes that rule matter: how the hardware actually turns a warp's 32 addresses into memory transactions, and why the difference between "adjacent" and "scattered" can be an 8x difference in delivered bandwidth for the exact same amount of useful data.

This page collects the vocabulary the rest of the section assumes, in one alphabetical list rather than grouped by topic, so it works as a lookup target rather than something you read start to finish. Each entry is written to stand alone — you should be able to land here from a search result with no other context and still understand the term — and each ends with a link to the page that develops it properly, with worked examples and the surrounding detail this page deliberately omits.

The Tensor Processing Unit is what happens when the weight-stationary systolic array from Systolic Arrays and Dataflow is scaled up to a datacenter training and inference accelerator, with a compiler stack and an interconnect built around it from the start. It is worth its own page separately from the general dataflow discussion because using a TPU well means accepting a programming model that looks nothing like CUDA: you do not write kernels for it at all.

A GPU is a throughput machine bolted onto a latency machine: thousands of simple cores, oversubscribed with far more threads than can run at once, trading single-thread speed for the ability to hide memory latency behind other work. Almost every performance question in this section reduces to the same one — did you keep the memory system busy, or is the chip sitting idle waiting on a load.

A GPU on a shared cluster isn't just requested and used the way a local one is — a scheduler decides which physical devices a job gets, renumbers them from the job's point of view, and (on Kubernetes) treats them as an indivisible resource unless something extra is configured. Getting any of this wrong tends to look like a correctness bug — a job silently touching the wrong device, or "no GPUs available" on a node that clearly has some — rather than an obvious scheduling error.

The fastest way to make training slower is to leave the GPU waiting on the CPU — and the second fastest way to make it faster, after fixing that, is to stop computing every number with more precision than the task actually needs. Most real speedups come from these two unglamorous facts, not from a cleverer algorithm.

Every transfer that has to bounce through a staging buffer in host memory pays for it twice — once copying into the buffer, once out — and that cost shows up whenever a GPU needs to talk to something other than another GPU on the same NVLink fabric: a network card, an NVMe drive, storage over the network. GPUDirect is NVIDIA's umbrella name for the mechanisms that let those other devices address GPU memory directly instead.

syncthreads() barriers a block; cluster.sync() barriers a cluster; neither reaches every block in a grid. Some algorithms genuinely need that — a multi-pass iterative solver that must finish writing generation *N* everywhere before any block reads generation *N* for generation *N+1*, for instance — and the usual answer, launching a second kernel between the passes, has real cost when the intermediate state is large and expensive to leave and re-establish. Grid-wide synchronization exists for that case, but it is not simply "a bigger syncthreads()": it comes with a hardware constraint that shapes the whole launch around it.

HIP (Heterogeneous-compute Interface for Portability) is AMD's answer to a specific problem driver, compiler, runtime, and the library ecosystem that gives HIP something to link against.

This page has not been written yet.

Fourteen folders is a lot of surface area, and reading them front to back is not the intended use. The section index lists what each folder covers; this page answers the question that list doesn't — what each folder assumes you already know, and what it hands to the folder after it. That is the information you need to enter in the middle, which is what most people do.

Warp Execution and Divergence described divergence as a cost model: a warp pays for every path its lanes take. Before Volta, divergence was also a scheduling model with sharp edges — the hardware tracked one program counter per warp using an explicit reconvergence stack, and code that assumed lockstep execution within a warp could rely on undocumented but consistent scheduling behavior. Volta replaced that mechanism, and the change is why so many older warp-synchronous idioms are now silently broken rather than merely non-portable.

Getting nvcc to compile a .cu file and getting a program to actually run on the GPU are two different problems, and most first-time setup failures come from conflating them. Three separate pieces of software have to agree with each other — a kernel-level driver, a toolkit for building code, and a runtime linked into the binary — and version mismatches between them are the single most common reason "it compiled but won't run" happens.

Occupancy Tuning showed that giving a thread several independent accumulators can substitute for occupancy — the scheduler fills a dependent chain's latency with another chain's independent work instead of waiting on another resident warp. This page generalizes that idea beyond hiding latency for its own sake: instruction-level parallelism, loop unrolling, cheaper instruction choices, and avoiding instruction sequences the hardware doesn't handle well are all ways to make the instructions a thread already issues cost less or overlap better, independent of how many warps are resident.

Everything so far in this section covers bandwidth inside a single GPU — SM to L2, L2 to HBM. The moment a workload needs data on another device, whether that's the host CPU or a second GPU, a completely different and usually much slower link is in the critical path. Which link is available, and at what bandwidth, is not a software choice — it's a property of the physical topology of the machine, and designing a multi-GPU strategy without first knowing that topology is a common source of disappointing scaling.

Every kernel launch has a fixed cost — driver overhead to enqueue it, and a round trip through global memory for its inputs and outputs — that has nothing to do with how much useful arithmetic the kernel does. A chain of small, bandwidth-bound elementwise kernels can spend more of its wall-clock time paying that fixed cost, repeatedly, than doing the actual computation, which is what makes fusing them into a single launch — or avoiding the separate launches altogether — a real optimization rather than a stylistic preference.

Latency and throughput sound like the same idea measured two ways, but a GPU treats them as almost unrelated design targets. Latency is how long one memory request takes to come back; throughput is how many bytes per second the memory system can sustain in steady state. A single DRAM access on a modern GPU takes several hundred nanoseconds — not meaningfully faster than it was a decade ago — yet the same hardware sustains terabytes per second in aggregate. The only way to reconcile a slow individual request with a fast aggregate rate is to have an enormous number of requests outstanding at once, and that single fact is why the CUDA programming model insists you expose thousands of threads instead of a handful.

This page has not been written yet.

This page has not been written yet.

This page has not been written yet.

Global Memory and Coalescing establishes the 32-byte-sector model and works out the arithmetic for coalesced, strided, and AoS access patterns. This page turns that model into an ordered checklist for a memory-bound kernel: which lever to pull first, which to skip, and why the order matters more than any individual technique.

cudaMalloc is the allocator every earlier example reached for, and for a program that allocates once at startup and frees once at exit, it's the right tool. It stops being the right tool the moment allocation moves inside a loop, because cudaMalloc and cudaFree are synchronizing, device-wide operations — they can take tens of microseconds each, which is invisible in a single call and devastating when it happens every iteration of a hot loop. The APIs on this page exist to give allocation a shape that matches how a program actually uses memory: padded for coalescing, ordered in a stream, pooled, or — rarely — managed as raw virtual address space.

Every earlier page in this section has quietly leaned on one synchronization primitive or another — syncthreads(), a cluster.sync(), a pipeline's consumer_wait() — to make writes from one thread visible to reads from another. This page states the rule underneath all of them explicitly: CUDA's memory model is weakly ordered, and without an explicit fence or atomic, there is no guarantee about when, or even whether, one thread's writes become visible to another thread at all. Getting this wrong doesn't usually crash a kernel; it produces a result that's correct most of the time and silently wrong occasionally, which is far worse.

A CUDA kernel does not have one undifferentiated pool of memory to work with — it has six, each with its own scope, lifetime, and performance profile, and picking the wrong one for a given piece of data is one of the most common ways a kernel ends up an order of magnitude slower than it should be. This page is the map: what each space is, who can see it, how long it lives, and the rough latency and bandwidth numbers that make the choice matter. The pages that follow work through each space in depth.

Knowing a kernel's arithmetic intensity puts it on the roofline plot in theory; knowing whether it is actually memory-bound or compute-bound in practice requires measuring the running kernel, because achieved bandwidth and achieved compute throughput are never the datasheet peaks the paper estimate assumed. This page turns the roofline classification from Arithmetic Intensity and the Roofline Model into a concrete diagnostic you run against a profiler, and adds the case the roofline model doesn't represent at all.

Metal is Apple's graphics-and-compute API, and on Apple silicon it sits on top of a hardware fact that has no real CUDA equivalent: the CPU and GPU are not two devices connected by a bus, they are two sets of cores reading the same physical memory. That single fact changes what "offload to the GPU" even means on this hardware, and it's the reason this page exists separately from the rest of the portability folder rather than being folded into a general graphics-API page alongside Vulkan and DirectX Compute.

Nsight Compute organizes hundreds of hardware counters into sections; Speed of Light tells you whether a kernel is memory-bound, compute-bound, or latency-bound, but not which specific resource inside that category is the bottleneck. This page is the metric-by-metric reference for answering that second question — the counters worth reading once Speed of Light has pointed at a direction, what a good value looks like, and what to change when it isn't.

A single GPU is often shared by more processes than it has obvious ways to divide itself among. The default sharing mechanism, time-slicing, works but wastes capacity on small kernels; two other mechanisms exist to do better, and they solve different problems. Multi-Process Service (MPS) lets independent processes' kernels run concurrently instead of merely taking turns; Multi-Instance GPU (MIG) physically partitions the hardware so processes don't share anything at all. Choosing between them means understanding what each one isolates and what it doesn't.

A workload that outgrows one GPU's memory or compute budget needs a second one, and every choice from there — how many processes, how work gets split, where the time actually goes — follows from a small set of rules about how CUDA treats "current device" as thread-local state. Get the discipline wrong and the symptom is rarely a crash; it's silent misallocation onto the wrong device or serialized work that looks like it should overlap.

NCCL (NVIDIA Collective Communications Library) is the library that moves data between GPUs — within a node over NVLink or PCIe, and across nodes over the network — through a small set of collective operations borrowed from the MPI world setting up a communicator, the collectives themselves, and how a call integrates with a stream. Collectives with NCCL covers the harder half — ring versus tree algorithm selection, and overlapping communication with gradient computation in a real training loop — and builds directly on the vocabulary defined here.

Nsight Systems narrows a slow run down to a slow kernel; Nsight Compute is what explains why that one kernel is slow. It replays the kernel with hardware performance counters attached and organizes the results into sections that go from a two-number summary down to per-source-line detail, which is the tool The Optimization Workflow means by "measure first."

Nsight Systems answers "where does the wall-clock time go across CPU, GPU, memory, and the network"; Nsight Compute answers "why is this one kernel slow" — start with Systems, because a kernel that looks slow in isolation is sometimes just waiting behind something else, and no amount of kernel-level tuning fixes a scheduling gap.

CuPy covers array operations, and cp.RawKernel covers the case where you need a real kernel — but that kernel is a CUDA C++ string embedded in a Python file, with no syntax highlighting, no type checking, and no debugger. Numba takes the other route: you write the kernel in Python, decorated with @cuda.jit, and Numba compiles that Python function to PTX at first call.

Marketing names like "Ampere" or "Hopper" map to a compute capability number, and that number — not the marketing name — is what actually gates whether a piece of code compiles or runs. This page walks the generations that matter for code written today, listing only what each one added that changes what you can write or how you must write it, and closes with the canonical table this section's other pages point back at whenever they gate a feature behind a specific compute capability.

Jetson is not a scaled-down GPU with different rules — it is a full CUDA-capable GPU on the same die as an Arm CPU, so everything in folders 03 through 07 of this section applies to it directly memory is physically shared between CPU and GPU, the power envelope is fixed and small, and a second, fixed-function inference engine — the DLA — sits alongside the GPU on the same package.

The Register File and Occupancy derives occupancy as a fixed calculation against three hardware limits — registers, shared memory, and block/thread slots — and works a full example through all three. This page is the tuning counterpart: given that calculation, which levers actually move it, what raising occupancy buys in practice, and — just as important — where it stops buying anything at all.

Getting a model off a training framework and onto an arbitrary piece of inference hardware needs a common description both sides agree on. ONNX is that description: a standardized, framework-neutral way to write down a model graph so that PyTorch, and separately TensorRT, and separately a phone's NPU compiler, can all read the same file and agree on what it means. What actually executes that file is a different question, and conflating the two is the single most common confusion this page exists to clear up.

OpenCL predates CUDA's dominance and was, for years, the only serious vendor-neutral answer to "how do I write one program that runs on GPUs from multiple vendors, plus CPUs, plus other accelerators." It still runs today, still ships in every major GPU driver, and still matters in a specific set of niches — but it lost the mindshare battle for general-purpose GPU compute, and understanding why is as useful as understanding the API itself.

Every portability layer covered so far in this folder still asks for a rewrite: HIP wants CUDA calls translated, SYCL wants kernels re-expressed as lambdas passed to a queue. Directive-based offload takes a different bet — annotate the loop nest you already have, keep one source file that still compiles and runs correctly on the CPU with the annotations ignored, and get working GPU code without restructuring the algorithm. That pitch is genuinely attractive for porting a large, already-correct codebase; it is a worse fit for the few kernels where every last percent of throughput matters, and this page is honest about where that line falls.

Most of the deployment stacks in this folder are strongest on hardware from one vendor, and OpenVINO's vendor is Intel. If the deployment target is an Intel CPU, an Intel integrated GPU, or the NPU built into a recent Intel Core Ultra laptop chip, OpenVINO is usually the toolkit that gets the most performance out of that hardware with the least fighting — it is Intel's own inference stack, tuned against Intel's own silicon, and it is the natural first thing to reach for once the deployment machine is known to be an Intel client device.

Almost every GPU kernel, however specialized, is built from a small set of recurring data-access shapes. Recognizing which pattern a problem is — before writing any code — tells you how parallelizable it is, what its likely performance limiter will be, and often points directly at a library implementation that already exists and is already tuned. This page names those shapes once, and every later applied-kernel page in this knowledge base assumes you already know these names — "this is a reduction" or "this needs a scan" is meant to carry full meaning by the time you reach folder 13.

This page has not been written yet.

Two GPUs in the same box can talk to each other directly, or every byte between them can detour through host memory — the difference is entirely a matter of whether peer access has been enabled, and it has real bandwidth consequences either way. This page covers the API for checking, enabling, and using that direct path; the bandwidths themselves live on Interconnects: PCIe and NVLink.

The SAXPY program in Your First Kernel paid three costs without the source code drawing attention to any of them ordinary malloc'd host memory is pageable, and pageable memory cannot be the source or destination of an asynchronous transfer. Pinned memory removes that restriction, and once transfers are asynchronous, they can be scheduled to overlap with compute instead of paying for it serially.

This page has not been written yet.

Tensor Cores laid out what tensor cores are and what makes a kernel eligible to reach them; this page is about actually writing code that does — the wmma warp-level API, the two facts about it that trip almost everyone up the first time, and why the code you write here is closer to a teaching exercise than to what a shipped GEMM looks like.

Almost every optimization in this section works through the C++ source and trusts nvcc to generate good machine code from it. Occasionally that trust runs out — an instruction has no intrinsic, or the compiler's chosen code sequence needs to be inspected or overridden directly — and the only way forward is reading or writing below the C++ level. This page covers both: reading the SASS a kernel actually compiles to, and, rarely, writing PTX by hand to reach an instruction the compiler won't emit on its own.

PyTorch composes: nearly anything can be built from existing operators. What composition cannot always give you is one kernel. A sequence of PyTorch ops writes every intermediate tensor to global memory and reads it back for the next op, so a chain of cheap elementwise operations spends almost all its time moving data — the problem Kernel Fusion and Launch Overhead covers in general.

An NPU's MAC array (see What Is an NPU?) is built around integer arithmetic first and floating point second, if at all. Getting a model onto that hardware efficiently means converting its weights and activations from floating point into integers in a way that a fixed set of scale and offset numbers can undo well enough that the model still works. That conversion — quantization — is a distinct engineering discipline from anything in ordinary model training, with its own vocabulary, its own failure modes, and its own tooling, and this page is that vocabulary.

Warp Execution and Divergence established the cost model: a warp pays for every path its 32 lanes collectively take, sequentially, so a branch only costs extra when lanes within the same warp disagree. This page is the applied counterpart — given that rule, what actually removes the cost in real kernels.

Summing an array, finding its maximum, counting matches — these collapse many values into one, and doing it efficiently on a GPU means combining values in parallel at every level of the hierarchy rather than serializing down to one thread. Warp-Level Primitives already built the innermost piece, warpReduceSum; this page builds outward from it — warp to block to grid — and then covers the related but distinct problem of a scan, where every intermediate result is wanted, not just the final one.

The Register File and Occupancy covers how a kernel's register usage feeds directly into the occupancy calculation. This page is the other side of that same fact: what registers actually hold, what forces a value out of the register file and into memory instead, and how to tell when that's happened to your kernel.

Arithmetic Intensity and the Roofline Model builds the model from datasheet peaks and a paper estimate of FLOPs and bytes — a first-order filter you can apply before a kernel even runs. This page replaces every number in that estimate with one measured from a real execution: the FLOPs a kernel actually issued, the bytes it actually moved, and the roofs the hardware actually achieves rather than what its spec sheet claims.

Every CUDA C++ example so far — >> launches, cudaMalloc, cudaMemcpy — has gone through the runtime API, the high-level interface linked in as cudart and initialized implicitly the first time a program touches the GPU. Underneath it sits the driver API (cuda.h, linked as cuda), a lower-level, explicit interface that the runtime itself is built on. Almost nothing in application code needs the driver API directly, but understanding what the runtime is hiding explains a class of errors ("invalid device context") that only make sense once you know a context exists at all.

nvcc defaults to compiling each .cu file's device code as a self-contained whole, with every device call resolved and inlined within that one translation unit. That default is invisible right up until device code needs to span files, at which point it fails in a way whole-program C++ intuition doesn't predict.

Global memory is fast in aggregate but every access still pays a round trip through the memory system; when several threads in a block need the same piece of data, or a thread needs to hand a value to another thread in the same block, routing through global memory to do it wastes bandwidth on traffic that never needed to leave the chip. Shared memory exists for exactly that: a small, explicitly-managed, on-chip scratchpad that every thread in a block can read and write, fast enough to use as a staging area rather than just a cache.

Shared memory earns its speed by serving a whole warp in one cycle, but that promise depends on the warp's 32 addresses landing in 32 different pieces of hardware. When they don't, shared memory — normally close behind register speed — degrades to a fraction of it, and the kernel doesn't fail, it just quietly runs slower with no compiler warning to explain why.

Memory Access Optimization stops once accesses are coalesced and vectorized, because coalescing only fixes how efficiently a kernel fetches the bytes it asks for — it does nothing about a kernel that asks DRAM for the same bytes over and over. Shared-memory tiling attacks that second problem directly: load each piece of data into on-chip shared memory once, then let every thread that needs it read it from there instead of going back to DRAM. This page derives why that reuse matters arithmetically and builds the tiled kernel that makes it concrete.

Every processor design answers two questions: how many instruction streams does it execute, and how many data streams does each instruction touch. Those two answers are the whole of Flynn's taxonomy, and they matter here because "GPU" is not a single point on that map — a GPU's arithmetic units are driven by an execution model, SIMT, that is easy to mistake for ordinary SIMD vectorization and behaves differently in exactly the cases that matter for correctness and performance.

This page has not been written yet.

Shared Memory Tiling ends on the bubble its own sgemmTiled kernel still has: every iteration of the tile loop is load → syncthreads() → compute → syncthreads(), and that structure forces the memory system and the ALUs to take turns rather than work at the same time. Software pipelining removes the turn-taking by starting the next tile's load before the current tile's compute has finished, so the two phases run concurrently across iterations instead of serially within one.

This page has not been written yet.

This page has not been written yet.

This page has not been written yet.

Pinned Memory and Host Transfers split a transfer into four chunks, each on its own stream, to overlap H2D copies with kernel execution — but it left unanswered exactly what a stream guarantees, why that overlap can silently fail to happen even with the code written correctly, and how to reason about ordering across streams on purpose. That's this page.

SYCL takes a different route to portability than HIP's near-identical-API strategy: instead of translating CUDA calls one-for-one, it's a Khronos standard for expressing host and device code in single-source, standard C++, with the compiler splitting host and device parts of the same file at compile time. The same .cpp file that launches a kernel also defines it, using ordinary lambdas and templates instead of a separate kernel language. oneAPI is Intel's product built around SYCL — a toolchain, a compiler, and a set of libraries — but SYCL itself is vendor-neutral and has multiple independent implementations.

What Is an NPU established that an NPU's efficiency comes from removing per-instruction overhead, not from having more MAC units than a GPU. This page is about the mechanism that actually achieves that: the systolic array, and the small set of dataflow patterns — weight-stationary, output-stationary, row-stationary — that decide what stays resident in each cell versus what streams past it. Which pattern a piece of hardware picks is not a minor implementation detail; it is the single biggest factor in how much energy that hardware spends moving data around.

A CUDA core executes one scalar fused-multiply-add per thread per cycle. A tensor core executes an entire small matrix-multiply-accumulate in hardware, cooperatively across a warp, in roughly the same number of cycles — which is why a kernel that reaches them can be an order of magnitude faster than the same arithmetic done on CUDA cores, and why so much of applied deep-learning performance work is really about getting a kernel eligible for tensor cores rather than about tuning ordinary FP32 code.

The central fact to hold onto about TensorRT is that it is a compiler, not a runtime library you call into layer by layer. Given a model graph and the exact shapes, precisions, and target GPU you tell it about, it benchmarks candidate kernel implementations for every layer, picks the fastest ones for that specific hardware, fuses what it can, and emits a serialized engine — a compiled artifact, not a portable model file. That one fact explains everything else on this page: why building an engine is slow (it is a search over kernel candidates, not a translation), why the resulting engine is fast (every layer runs the kernel TensorRT found to be fastest on that GPU, for those shapes), and why the engine does not travel to a different GPU, a different TensorRT version, or often even a different driver.

Once you accept that some of your work belongs on a throughput engine, you have to pick one, and the market offers far more options than "NVIDIA or not". There are discrete GPUs on a PCIe slot, GPUs integrated into the same die as the CPU, phone-class GPUs paired with NPUs, datacenter training and inference ASICs reachable only through a compiler, and FPGAs where you describe the datapath yourself. They differ enormously in peak throughput — and that difference is almost never what decides the outcome.

A single .cu file contains two programs wearing one extension: host C++ that runs on the CPU, and device code that has to end up as instructions a specific GPU can execute. nvcc is the tool that splits those apart, compiles each with the right compiler, and glues the results back into one binary — understanding that split is what makes -arch, -code, and the difference between a build that runs everywhere and one that only runs on the GPU it was built for make sense.

Before any of the CUDA syntax in the next section makes sense, one structural fact has to be settled: a discrete GPU is a separate computer. It has its own memory, its own processors, and no automatic view of what the CPU is doing — every byte the GPU touches had to arrive there deliberately, and every result has to leave the same way. This page names that structure once, independent of any specific API, so that the CUDA-specific mechanics in Your First Kernel land on a model you already have rather than a pile of new syntax to memorize.

Tuning a kernel without a loop around the work turns into guessing: try something that sounds plausible, rerun, eyeball whether it got faster, repeat. The workflow that actually converges is narrower than that — profile to find the one resource the kernel is actually waiting on, apply only the fix that targets that resource, re-measure to confirm the fix worked and see what limiter is binding now, and stop once further gains are no longer worth the effort. Every other page in this folder is a toolbox entry for one step of this loop, not a replacement for it.

Every CUDA program in this section so far has assumed an NVIDIA GPU underneath it. That assumption is usually safe in a single research group's cluster and usually false the moment code has to run on a customer's laptop, an AMD-powered supercomputer, or a mobile SoC. "Portability" sounds like a single problem with a single fix — pick a vendor-neutral API and move on — but it is really three separate, increasingly hard problems wearing one name, and confusing them is how teams end up with code that compiles everywhere and runs well nowhere.

Occupancy is defined in the glossary as the ratio of resident warps to the maximum an SM supports, and Latency, Throughput, and Latency Hiding explains why that ratio matters — resident warps are what supply the concurrent memory requests Little's Law demands. This page is about the other half: occupancy is not a tunable dial, it's the output of a fixed calculation against fixed hardware limits, and the register file is usually the tightest of those limits.

The SM is the unit everything about GPU performance is ultimately accounted against: occupancy, register pressure, shared-memory capacity, and warp scheduling are all per-SM quantities. Zooming into one SM explains why a block, once scheduled, stays resident on a single SM for its entire lifetime, and why the resources that limit how many blocks can run concurrently are the ones this page enumerates.

Blocks are independent by design: no portable synchronization between them, no shared on-chip memory, and no guarantee two blocks even run at the same time. That independence is what lets a kernel scale from a laptop GPU to a data-center one, but it also means algorithms that need a little cross-block cooperation — a bit more shared memory than one block's SM can hold, or a barrier across a handful of blocks — have nowhere to turn. A thread block cluster relaxes exactly that restriction, for a small group of blocks the hardware guarantees will be co-resident on the same GPU Processing Cluster (GPC) at the same time.

Every thread in a kernel runs the same code, so the only thing that makes it operate on its piece of the data rather than every thread's piece is the index it computes from its own position in the grid. Getting that formula right — and guarding it correctly — is the one piece of CUDA arithmetic that shows up in essentially every kernel, from SAXPY to the applied kernels later in this section.

A kernel launch like saxpy>>(...) doesn't just start "some threads" — it starts a precisely structured hierarchy, and the shape of that hierarchy is what lets the same compiled kernel run correctly on a small laptop GPU and a data-center accelerator with an order of magnitude more SMs. Understanding the levels of that hierarchy, and which ones can and can't communicate, is the difference between a kernel that scales and one that only happens to work on the GPU it was tested on.

Thrust is a C++ template library, shipped with the CUDA Toolkit, that reproduces the shape of the C++ Standard Template Library on the GPU the host-callable layer above cub's in-kernel primitives.

The training bill arrives, and most of it bought idle GPU time — a genuinely common outcome, and an entirely preventable one. Utilisation, not hourly rate, decides training cost: a half-idle expensive GPU beats neither a well-fed cheap one nor, often, a smaller model trained more efficiently.

Writing a fused kernel in CUDA C++ means writing the fusion and everything around it: the thread-to-element mapping, the shared-memory staging, the vectorized loads, the bank-conflict-free layout. Most of that work is mechanical, most of it is where the bugs live, and none of it is the algorithm you actually wanted to express.

Every allocation covered so far draws a hard line between host and device memory: a pointer is valid on one side or the other, and moving data across the line is an explicit cudaMemcpy the programmer writes and pays for. Unified Memory erases that line for the source code — one pointer, usable from both host and device — while the underlying hardware and driver still have to physically move bytes between two separate memory systems whenever the data is touched from the "wrong" side. Understanding when that migration happens, and how to steer it, is the difference between Unified Memory being a convenience and Unified Memory being a performance trap.

This page has not been written yet.

Vulkan and Direct3D 12 are graphics APIs first, and each exposes a compute pipeline as a sibling of its graphics pipeline rather than a separate product. The reason to reach for one of them is almost never raw compute throughput — it's what happens on either side of the kernel. If a compute pass writes into a buffer or texture that a render pass reads next, doing both in the same API keeps the data on the device, in the API's own memory model, with no cross-API copy and no synchronization handoff between two separate runtimes. That's the entire case for this page: not "Vulkan compute is fast," but "Vulkan compute is already where your renderer lives."

Warps and Warp Schedulers established that a warp scheduler issues one instruction to all 32 lanes of a warp at once. That raises an obvious question: what happens when those 32 threads disagree about which instruction to execute next, because a branch condition evaluated differently across lanes? The answer — the warp executes both outcomes and masks off the lanes that don't apply to each one — is the single most important cost model in CUDA kernel design, and getting it precise is the point of this page.

Independent Thread Scheduling explained why lockstep-dependent, volatile-based tricks for exchanging data within a warp no longer work, and why the fix is a family of intrinsics that carry explicit synchronization and an explicit participant mask. This page is that family: the shuffle intrinsics for moving values directly between lanes' registers, the vote intrinsics for asking a yes/no question of the whole warp, and the canonical reduction pattern built from them.

Every SIMT behavior that looks unusual coming from CPU threading — coalescing, divergence, the fact that occupancy is measured in resident warps rather than resident threads — traces back to one hardware fact what it decides among, how fast it can issue, and how to read its behavior back out of a profiler.

Every compute API so far in this folder assumes a native process with a driver it trusts. A browser tab cannot make that assumption — it runs code from an untrusted origin, on a machine it doesn't control, and has to expose GPU compute without letting a web page read another tab's memory or hang the system. WebGPU is the answer: a browser API, backed by Vulkan, Metal, or D3D12 underneath, that gives web content a compute (and graphics) pipeline shaped like a stripped-down, sandboxed version of those native APIs.

Every accelerator in this section so far has been a variation on "more programmable cores, running in parallel." An NPU (neural processing unit) is a different move entirely: instead of adding parallel general-purpose lanes, it removes almost all of the general-purpose machinery and replaces it with a dataflow of multiply-accumulate (MAC) cells wired specifically for tensor arithmetic. CPU vs GPU vs NPU already introduced this as the third design point; this page works through what that design point actually buys and what it costs.

Most failed GPU ports do not fail because the kernel was slow. They fail because the workload was never shaped like something a GPU accelerates, and the port made that visible only after weeks of work. The kernel itself often does run twenty times faster than the CPU loop it replaced — and the program gets slower anyway, because the time now goes into transfers, synchronization, and the 60% of the runtime that was never offloaded at all.

The useful question is not "why is a GPU faster than a CPU" — it usually isn't. A single CPU core will finish one dependent chain of instructions sooner than any GPU will, and it will do it on branchy, pointer-chasing, irregular code that a GPU handles badly. The real question is how a fixed transistor budget gets spent. A CPU spends most of its area on machinery that makes one instruction stream go fast: out-of-order scheduling, register renaming, branch prediction, and a deep cache hierarchy that hides DRAM latency from a handful of threads. A GPU deletes almost all of that and spends the reclaimed area on arithmetic units, then keeps them busy by oversubscribing the machine with far more threads than can execute in any one cycle.

Every CUDA program, no matter how large, is built from the same five moves: allocate device memory, copy input in, launch a kernel, copy output back, free what was allocated. SAXPY — y = a*x + y, scalar-times-vector-plus-vector — is small enough to show all five in one file without anything else getting in the way. The rest of this page walks the same file section by section.