Topic domain
GPU Computing
Throughput-oriented execution and its programming models.
- Articles
- 80
- Clusters
- 8
- Foundational
- 24
- Intermediate
- 33
- Advanced
- 23
Warps, divergence, and occupancy.
- Warp Granularity: Why Thirty-Two Threads Share One Program Counter
- Divergence and the Reconvergence Stack
- What Occupancy Actually Counts
- Predication and Branch Selection
- Active-Mask Accounting and Issue Slot Efficiency
- Register Allocation and the Occupancy Ceiling
GPU Memory
10Coalescing, shared memory, and hierarchies.
- The Address Spaces a GPU Kernel Can Touch
- How a Warp Turns Thirty-Two Loads into Memory Transactions
- Stride, Misalignment, and the Cost of a Bad Layout
- Bank Conflicts Inside GPU Shared Memory
- Tiling Through Shared Memory and Overlapping the Load
- Register Pressure and the Spill into Local Memory
Launch geometry, synchronization, and patterns.
- Thread Hierarchy and Global Index Arithmetic in a Kernel Launch
- Tail Handling and Bounds Checks in Partial Tiles
- Grid-Stride Loops and Launch Size Independence
- Block Barriers and Divergent Synchronization
- Tiled Stencils and Shared Memory Halos
- Tree Reductions and Floating-Point Summation Order
Rasterization, shading, and raster state.
- Rasterization Rules: Pixel Centers and the Top-Left Fill Convention
- Homogeneous Clipping and Why Rasterizers Use a Guard Band
- Perspective-Correct Interpolation and the Reciprocal-W Trick
- Quad Shading, Helper Invocations, and Where Derivatives Come From
- Early Depth Testing and the Order of Per-Fragment Operations
- Multisample Coverage and What the Resolve Step Actually Does
Compute APIs
10Vendor and portable acceleration interfaces.
- The Contract a Compute API Encodes
- Launch Geometry and the Arithmetic of Thread Indexing
- OpenCL Objects, Contexts, and Runtime Compilation
- Command Queues, Events, and Cross-Stream Ordering
- Host-Device Transfers and the Cost of Crossing the Bus
- PTX, SASS, and What a Fat Binary Contains
Roofline reasoning and bottleneck diagnosis.
- The Roofline Bound and Arithmetic Intensity
- Achieved Bandwidth Versus the Quoted Peak
- Occupancy as a Latency-Hiding Budget
- Wave Quantization and the Tail Effect
- Warp Stall Reasons and Cycle Attribution
- Memory-Level Parallelism and Latency Tolerance
Partitioning, passthrough, and multi-tenancy.
- The Interception Surface of a Virtual GPU
- PCI Passthrough and the Boundaries an IOMMU Actually Draws
- SR-IOV Virtual Functions and Static Partitioning of GPU Engines
- Paravirtual GPUs and the Mediated Device Lifecycle
- Fixed-Slice Partitioning of Shared Engines Inside One Device
- Time-Slicing and What a GPU Context Switch Costs
Host-device orchestration and unified memory.
- Pinned Host Memory and Why Pageable Transfers Cannot Be Asynchronous
- Unified Virtual Addressing and How a Pointer Becomes Usable on a Device
- Stream Ordering: What a Queue Guarantees and What It Does Not
- GPU Page Faults and Demand Migration in Unified Memory
- Host-Mapped Memory over PCIe: Zero-Copy Access and Its Latency Wall
- System-Scope Atomics and Memory Ordering Across the Bus
All GPU Computing articles
- Oct 15, 2026 Bank Conflicts Inside GPU Shared Memory How a shared-memory address selects one of thirty-two banks, when the hardware serves a warp in a single pass, and why padding a tile by one column removes a thirty-two-way serialization.
- Sep 2, 2026 Streams, Events, and Concurrent Kernel Execution What a stream actually orders, why two kernels can overlap or fail to, and how events turn a stream of launches into a dependency graph.
- Sep 1, 2026 Latency Hiding and the Arithmetic of Warp Counts Deriving how many resident warps are needed to cover a memory stall from instruction count and latency, and where the Little's law model overstates what interleaving can deliver.
- Aug 27, 2026 Stride, Misalignment, and the Cost of a Bad Layout How stride and alignment interact to add sectors to every warp request, why a row-major array accessed by column collapses bandwidth, and what a shared-memory transpose costs in return.
- Aug 26, 2026 Register Pressure and the Spill into Local Memory Why the register file caps how many threads an SM can host, how the compiler trades registers for scattered memory traffic when it runs short, and why the spilled addresses still coalesce.
- Aug 20, 2026 Time-Slicing and What a GPU Context Switch Costs Preemption on a GPU is not one operation but several, each at a different granularity, and the cost of switching decides whether oversubscription is worth it.
- Aug 18, 2026 Fixed-Slice Partitioning of Shared Engines Inside One Device Splitting a GPU into isolated instances without separate PCI functions requires slicing the caches and memory controllers, and the slices that cannot be cut set the floor on granularity.
- Aug 17, 2026 Raster State Validation and Why Draw Order Is the Only Ordering Primitive How per-draw state forms an opaque pipeline object, why changing render targets mid-command-buffer forces a flush of binned work, and why the API exposes almost no way to sequence fragments directly.
- Aug 16, 2026 Stream Ordering: What a Queue Guarantees and What It Does Not The ordering contract of a command stream, how host-side submission stays asynchronous, and why two streams give no ordering at all without an explicit dependency.
- Jul 27, 2026 SR-IOV Virtual Functions and Static Partitioning of GPU Engines A single PCI function becomes many by replicating the parts that carry identity, and the GPU functions that cannot be replicated become the hard limits on how finely a card can be split.
- Jul 24, 2026 The Roofline Bound and Arithmetic Intensity How the ratio of arithmetic to bytes moved places a kernel under either a compute ceiling or a bandwidth ceiling, and why the ridge point differs between devices.
- Jul 23, 2026 Winding Order, Backface Culling, and the Two-Sided Lighting Problem How the signed area of a projected triangle identifies a front face, why culling is a rasterizer-level rejection rather than a vertex operation, and what happens to normal direction when a face is seen from behind.
- Jul 16, 2026 OpenCL Objects, Contexts, and Runtime Compilation The platform-to-queue object graph in OpenCL, what a program build actually produces, and how vendor portability breaks at kernel compile time rather than at run time.
- Jul 15, 2026 Single-Pass Prefix Scan with Decoupled Look-Back How a scan produces every prefix in one launch without a chain of kernels, what each block publishes and waits for, and why the algorithm stays linear in total work.
- May 22, 2026 Warp Granularity: Why Thirty-Two Threads Share One Program Counter How a block of threads is carved into warps, why the hardware tracks one instruction stream per group of thirty-two lanes, and what that granularity costs at the edges.
- May 4, 2026 PCI Passthrough and the Boundaries an IOMMU Actually Draws Why handing a GPU to a guest requires an isolation unit that no sibling device can bypass, and what changes when the device caches its own translations.
- Apr 19, 2026 Pinned Host Memory and Why Pageable Transfers Cannot Be Asynchronous Why a copy engine cannot be handed an ordinary host pointer, what the driver does instead, and how the staging buffer turns an asynchronous copy into a synchronous one.
- Apr 18, 2026 Conservative Rasterization for Thin and Subpixel Geometry How expanding the coverage test to "the pixel touches the triangle" guarantees that no primitive is dropped, what the inner coverage variant adds, and which artifacts the expanded test creates.
- Apr 3, 2026 System-Scope Atomics and Memory Ordering Across the Bus Why an atomic operation on device memory is not automatically visible to the host, what the scope qualifier changes, and how release and acquire order a flag against its payload.
- Apr 1, 2026 Block Barriers and Divergent Synchronization What __syncthreads() actually orders within a block, why a barrier reached by only some threads is undefined, and how warp-level primitives replace it where blocks are too coarse.
- Mar 30, 2026 Perspective-Correct Interpolation and the Reciprocal-W Trick Why screen-space linear interpolation of a clip-space attribute is wrong, how dividing by w before the rasterizer and dividing back after the vector unit restores correctness, and what it costs.
- Mar 25, 2026 Host-Device Transfers and the Cost of Crossing the Bus Pinned versus pageable host memory, why a blocking copy costs more than its bandwidth suggests, and where the staging copy inside the runtime comes from.
- Mar 20, 2026 PTX, SASS, and What a Fat Binary Contains The two-stage compilation path from CUDA C to machine code, why PTX survives an architecture change and SASS does not, and how the loader picks a cubin from an embedded archive.
- Mar 12, 2026 Live Migration of a Guest That Holds a GPU Moving a running guest with an attached accelerator requires capturing device state, channel state, and page state in an order that leaves neither side with a stale view.
- Mar 5, 2026 Occupancy as a Latency-Hiding Budget How resident warps are rationed by registers, shared memory, and block granularity, and how to compute the number of concurrent requests a given occupancy can sustain.
- Feb 9, 2026 Address Translation Across Host and Device A GPU access may traverse three translation stages, and where the stages live determines whether a shared address space is cheap, whether faults are recoverable, and what a migration can move.
- Feb 8, 2026 Capability Negotiation Across Architectures Choosing a backend and a kernel variant at load time: what to query, how to key a selection cache, and where feature detection diverges from the specification.
- Jan 22, 2026 Warp Stall Reasons and Cycle Attribution How a profiler turns idle issue slots into named stall categories, what each category actually means at the scheduler, and how to read the breakdown without over-interpreting it.
- Jan 12, 2026 Driver-Level Runtime Compilation and Module Loading The chain from kernel source to a loadable module: intermediate collection, per-device lowering, cached artefacts, and the failure modes at each stage of the load.
- Jan 2, 2026 Clock Derating and the Roofline That Moves Why the ceiling a kernel is measured against depends on power state and duration, and how sustained clocks change which side of the ridge a kernel occupies.
- Dec 22, 2025 Kernel Fusion and the Cost of a Launch Boundary What a separate kernel launch actually costs in time and traffic, when merging two kernels into one is a net win, and which shapes cannot be fused.
- Oct 21, 2025 Atomic Contention and the L2 as a Serialization Point Why a global atomic executes at the cache slice that owns its address, how a hot counter turns a parallel machine into a queue, and which reductions avoid the queue entirely.
- Oct 14, 2025 Thread Hierarchy and Global Index Arithmetic in a Kernel Launch How a grid of blocks and threads is flattened into buffer indices, where the arithmetic overflows, and what the launch dimensions deliberately do not guarantee.
- Aug 28, 2025 Tree Reductions and Floating-Point Summation Order Why a pairwise tree sum is not the same number as a linear sum, how much error the tree removes, and how shuffle-based and shared-memory reduction stages compose.
- Aug 12, 2025 Warp Specialization and Producer-Consumer Warps How assigning different code to different warps of a block removes divergence at the source, what the handoff between producer and consumer warps costs, and when it loses.
- Aug 11, 2025 Partitioning Work Across Devices and the Cost of Every Dependency How to split a computation across several devices without drowning it in synchronization, and why the cost of a dependency is set by the slowest participant.
- Aug 8, 2025 Divergence and the Reconvergence Stack How a warp executes both sides of a branch when its lanes disagree, what the per-warp mask stack records, and why reconvergence is not always where the source suggests.
- Jul 30, 2025 What Occupancy Actually Counts Why occupancy is a ratio of resident warps to a hardware maximum rather than a utilization measure, and how registers, shared memory, and block limits each cap it.
- Jun 17, 2025 Tail Handling and Bounds Checks in Partial Tiles Why the last block of a launch is the one that returns early, how predicated access differs from an early return, and what goes wrong when barriers and zero-length buffers meet an oversized grid.
- Jun 17, 2025 Unified Virtual Addressing and How a Pointer Becomes Usable on a Device Why a single address space spanning host and device memory lets an API infer a pointer location from its value, what that inference depends on, and where it fails.
- May 19, 2025 Rasterization Rules: Pixel Centers and the Top-Left Fill Convention Why a triangle becomes a set of fragments at sample points rather than pixel areas, and how the top-left rule keeps shared edges from being drawn twice or not at all.
- May 15, 2025 Grid-Wide Barriers and Persistent Work Queues How a barrier that spans every block is built from atomics and counters, what residency it requires, and how a persistent grid replaces the launch queue with a device-side one.
- Apr 29, 2025 The Contract a Compute API Encodes Kernels, launch geometry, memory objects, and queues: the four obligations every GPU acceleration interface must satisfy, and why each one is spelled differently across vendors.
- Apr 14, 2025 The Address Spaces a GPU Kernel Can Touch How global, local, shared, and constant memory differ in scope, lifetime, and backing store, and why the compiler must know which space a pointer refers to.
- Apr 6, 2025 Quad Shading, Helper Invocations, and Where Derivatives Come From Why fragment shaders execute in 2x2 quads, what a derivative instruction computes when neighbouring pixels are not covered, and how helper invocations keep implicit gradients defined.
- Mar 20, 2025 Register Allocation and the Occupancy Ceiling How per-thread register counts are rounded into per-warp reservations, why a few extra registers can cost half the resident warps, and what forcing a lower count actually costs.
- Mar 17, 2025 Device Memory Oversubscription and Page Migration Selling more device memory than exists requires evicting pages to host memory, and the cost of that depends entirely on how the device detects and services a miss.
- Mar 15, 2025 Host-Mapped Memory over PCIe: Zero-Copy Access and Its Latency Wall What happens when a kernel dereferences host memory directly across the bus, why write-combining changes the cost, and the access patterns that make it worth doing.
- Mar 7, 2025 Command Queues, Events, and Cross-Stream Ordering What an in-order queue actually guarantees, how events turn streams into a dependency graph, and why the same code is correct under CUDA streams and wrong under Vulkan barriers.
- Feb 20, 2025 Grid-Stride Loops and Launch Size Independence Why the grid-stride loop survives shrunken launches and changing problem sizes, what it costs against a one-element-per-thread kernel, and when it is the wrong shape.
- Feb 4, 2025 Achieved Bandwidth Versus the Quoted Peak Why a measured stream rate falls short of the number on the datasheet, how to count the bytes a kernel really moved, and which part of the gap is structural.
- Feb 3, 2025 GPU Page Faults and Demand Migration in Unified Memory How a device MMU fault suspends a warp, what the driver does with the faulting address, and why migration granularity decides whether oversubscription works at all.
- Jan 22, 2025 Deferred Shading and the Tile-Binned Alternative Why recording surface attributes instead of shaded colour defers the expensive lighting work, and how binning fragments per screen tile reaches the same locality without the bandwidth.
- Jan 18, 2025 Tiled Stencils and Shared Memory Halos How an overlapping tile plus a barrier turns repeated global reads into repeated shared-memory reads, how halo overhead scales with tile size, and what register blocking adds.
- Jan 15, 2025 Row-Buffer Locality Inside GPU DRAM How DRAM banks, open-row policies, and bank groups shape what a memory controller can schedule, and why many concurrent streams lower achieved bandwidth without moving more bytes.
- Oct 10, 2024 The Interception Surface of a Virtual GPU A GPU presents a register file, a bus-mastering DMA engine, and an interrupt source at once; sharing one means deciding which of the three the host keeps control of.
- Oct 1, 2024 Residency Hints and the Arithmetic of Managed-Memory Thrashing How preferred-location and accessed-by advice changes where a managed page lives, why advice is not a guarantee, and how to tell a capacity problem from a policy problem.
- Sep 27, 2024 Atomics, Fences, and Memory Scopes in Device Code What a device-scope atomic orders, why relaxed operations reorder across a barrier, and how scope qualifiers change what a fence must flush.
- Sep 24, 2024 Launch Geometry and the Arithmetic of Thread Indexing How a grid extent becomes a linear identifier, what the block and grid dimension limits are, and why a non-multiple block size leaves a guard predicate in every kernel.
- Aug 20, 2024 Constant Memory and the Read-Only Load Path How a 64 KB host-written bank gets a fast path to the pipelines, why uniform reads inside a warp are almost free while divergent ones serialize, and what the non-coherent load promises the compiler.
- Aug 18, 2024 Portable Single-Source Layers and Where They Leak How SYCL, OpenMP offload, and Kokkos present one source file for host and device, and which abstractions fail to compile, fail to perform, or silently diverge in results.
- Jul 23, 2024 Divergence Amplification in Nested Control Flow Why nesting multiplies serialized path count rather than adding to it, how the mask stack depth bounds what the hardware can track, and which rearrangements collapse the product.
- Jul 22, 2024 Tiling Through Shared Memory and Overlapping the Load The reuse arithmetic that decides a tile size, why a register tile on top of a shared tile cuts scratchpad traffic again, and how an asynchronous copy lets the next tile load while the current one is computed.
- Jun 28, 2024 Predication and Branch Selection When a compiler converts a conditional into predicated instructions or a select, and the cost comparison that decides whether branching or straight-line code is cheaper for a warp.
- May 13, 2024 Issue-Limited Kernels and the Cost of Divergence When the instruction stream rather than the memory system sets the runtime, how warp divergence turns one instruction into several, and how to tell the two apart with counters.
- Dec 14, 2023 Active-Mask Accounting and Issue Slot Efficiency How to compute the fraction of lane capacity a warp actually delivers from its active masks, and why instruction issue rather than lane count is the resource that runs out.
- Nov 27, 2023 Multisample Coverage and What the Resolve Step Actually Does How one shading result is shared across several coverage samples, why resolve averages coverage rather than colour, and where alpha-to-coverage and centroid sampling fit.
- Nov 23, 2023 Early Depth Testing and the Order of Per-Fragment Operations Which per-fragment operations happen before the fragment shader and which after, why stencil and depth tests can be hoisted only when the shader permits it, and what that costs on overdraw.
- Nov 20, 2023 Memory-Level Parallelism and Latency Tolerance How many independent memory requests a kernel keeps in flight determines whether it can saturate bandwidth, and how unrolling and load reordering raise that count without adding warps.
- Sep 26, 2023 Wave Quantization and the Tail Effect Why a grid that fills the device one-and-a-bit times wastes the surplus, how to compute the number of waves, and which grid shapes avoid a long final wave.
- Sep 5, 2023 Per-Lane Program Counters and Independent Thread Scheduling What changes when each lane carries its own program counter and call stack, how divergence stops being a strict two-phase serialization, and why implicit reconvergence disappears.
- Aug 12, 2023 Page Table Mirroring Between CPU and Device How the kernel lets a device driver snapshot a process page table, why the snapshot needs an invalidation callback, and what breaks when the two disagree.
- Jul 23, 2023 L2 Slices, Address Hashing, and Partition Camping Why the last-level GPU cache is physically distributed across memory controllers, how addresses are spread over it, and how a stride that resonates with the mapping leaves most of the memory system idle.
- Jul 9, 2023 Tenant Isolation and the Channels Performance Interference Travels Through Separating address spaces does not separate throughput; the shared resources that remain determine how much one tenant can slow another and which controls can bound it.
- Jun 14, 2023 Peer-to-Peer Transfers and the Topology That Decides Whether They Work What peer access actually establishes in the device page tables, why the path between two devices is not the path through the host, and how to read the topology.
- Mar 28, 2023 Paravirtual GPUs and the Mediated Device Lifecycle How a host-side scheduler multiplexes one physical GPU across guests without giving any of them direct access, and what the create-configure-assign-teardown sequence has to guarantee.
- Mar 22, 2023 How a Warp Turns Thirty-Two Loads into Memory Transactions The request the memory system actually receives from a warp-wide load, why the unit of transfer is a 32-byte sector rather than a cache line, and how a partial write forces a read first.
- Sep 29, 2022 Homogeneous Clipping and Why Rasterizers Use a Guard Band How triangles are clipped against the view volume in clip space before the perspective divide, why clipping only the near plane is usually enough, and what the guard band buys.
- Apr 22, 2022 Atomics and Contention in Counters and Aggregates Why a single global atomic counter serialises a kernel, how per-warp and per-block aggregation cuts that traffic, and which failure modes survive the optimisation.
- Apr 6, 2022 Sector Counting and Why Cache Hit Rate Misleads How memory transactions are accounted in sectors, why a high hit rate can coexist with poor performance, and how to read the traffic counters that separate the two.