86 docs tagged with "cuda"

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.

> 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.

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.

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.

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.

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.

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.

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.

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.

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.

This page has not been written yet.

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.

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.

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.

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.

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.

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.

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.

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.

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.

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.

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.

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.

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 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.