GPU Programming and Kernel Optimization

References

A synthesis of selected conference talks and technical references. Citations link to the source material; they do not imply that every talk on this subject is included.

1. Roofline Performance Model — NERSC

   The Roofline model relates attainable arithmetic throughput to compute capacity, memory bandwidth, and arithmetic intensity: operations performed per byte moved at a specified memory level. Low-intensity work can be limited by moving data rather than executing arithmetic. Reuse can increase intensity by doing more work per transferred byte. Roofline is an upper performance envelope; measured execution can lie below it because of additional inefficiencies. The same workload can fall in different limiting regimes on different hardware.

2. Triton: Matrix Multiplication

   The tutorial multiplies an M-by-K matrix by a K-by-N matrix into M-by-N output tiles. Each program accumulates successive K chunks into FP32 partial outputs, zero-fills masked K-edge loads, and masks final stores. Addresses use explicit dimension strides. Its optional activation runs before conversion to FP16, so its rounding boundary differs from applying activation to an already stored FP16 result. Autotuning varies tile dimensions, warp counts, and pipeline stages, keyed by M, N, and K. An upstream test compares a 512-by-512 FP16 case against torch.matmul with absolute tolerance 0.01 and zero relative tolerance.

3. A Set of Level 3 Basic Linear Algebra Subprograms

   The March 1990 Level 3 BLAS paper organized reusable numerical interfaces around matrix-matrix operations. Its motivation was that vector-at-a-time algorithms often moved too much data on machines with hierarchical memory. Partitioning matrices into blocks allowed repeated calculations while operands remained in cache or local memory. Parallelism could occur both between independent blocks and within each block. BLAS means Basic Linear Algebra Subprograms: standard numerical building blocks whose implementations can be tailored to hardware while callers retain a common interface.

4. CUDA Programming Guide: Programming Model

   The host is the CPU and its memory; the device is the GPU and its memory. A kernel is a function invoked for GPU execution. Its launch organizes threads into blocks and blocks into a grid. Each block executes on one streaming multiprocessor, or SM, whose resources support computation and cooperation. A grid can contain far more blocks than can execute simultaneously, and ordinary blocks cannot assume scheduling order or depend on another block becoming resident. CUDA groups threads into 32-thread warps. Under single-instruction, multiple-thread execution, threads follow the same kernel code but may take different branches; inactive lanes are masked while a branch executes.

5. PyTorch: torch.Storage

   A regular tensor is a multidimensional array described by underlying storage, element type, shape, strides, and storage offset. Shape gives each dimension's size; a stride describes the step needed to advance one element along a dimension. Storage contains bytes, while this metadata determines how those bytes are interpreted. Multiple tensors can share the same storage, so distinct tensor objects do not necessarily mean independent buffers.

6. Matrix Multiplication Background User's Guide

   GEMM computes C=αAB+βC; ordinary multiplication uses α=1 and β=0. For A shaped M×K and B shaped K×N, each of the M×N outputs combines K products, conventionally counted as 2MNK floating-point operations. The guide's ideal FP16 traffic estimate gives intensity MNK/(MK+KN+MN) operations per byte. Larger output tiles improve operand reuse but produce fewer independent blocks. Its fixed 256×128-tile A100 example uses M=2304 and K=4096: N=1536 produces 108 tiles, while N=1544 produces 117, leaving nine tiles in an additional scheduling wave. Tile quantization means incomplete output tiles; wave quantization means an incompletely occupied final group of resident blocks.

7. OpenXLA: GPU Architecture Overview

   XLA separates logical tensor shape from physical layout. Layout assignment can represent a transpose as metadata, while conflicting layout requirements can require a copy that physically transposes data. XLA fusion groups multiple operations into one GPU kernel and passes internal intermediates through registers or shared memory instead of materializing them in high-bandwidth device memory. This establishes why an operator graph need not map one-to-one onto launches or intermediate arrays.

8. Faster Parallel Reductions on Kepler

   The implementation first accumulates several input elements per thread in registers, then combines thread partials through warp shuffle reductions. One lane per warp stores a partial sum in shared memory; after a block barrier, the first warp reduces those partials. For larger arrays, one kernel writes one result per block and a subsequent kernel combines them. This separates ownership of intermediate outputs from the synchronization needed to consume them and reduces inter-thread communication by doing more work locally.

9. CUDA C++ Best Practices Guide

   Kernel launches return before GPU work completes; CPU timing must account for completion. CUDA events record device timestamps when reached in a stream, and the example waits for the stop event before reading elapsed time. Extra synchronization can alter execution. Bandwidth is data transferred per unit time. Occupancy is active warps per multiprocessor divided by its supported maximum; registers, shared memory, and block size constrain residency, and higher occupancy need not improve performance. Amdahl's law bounds fixed-workload improvement by the fraction accelerated: unchanged work remains even if the optimized portion becomes arbitrarily fast.

10. What every AI engineer needs to know about GPUs

    Concurrency hides waiting by letting other work progress, while parallelism increases how much work can progress simultaneously.

11. What every AI engineer needs to know about GPUs

    The speaker recommends organizing GPU work to exploit growing bandwidth, because latency improvements face stronger physical limits.

12. HIP C++ Language Extensions

    HIP calls the number of threads in an execution group warpSize; on AMD hardware this is the wavefront size. Portable code must query or specialize for the target rather than assume 32. The inspected documentation specifies 64 for gfx9 and 32 for gfx10 and later HIP targets. Lane masks and collective operations must accommodate the applicable width.

13. Khronos: cl_khr_subgroups

    OpenCL sub-groups are implementation-controlled groups of work-items within a work-group, with their own collective operations and synchronization primitives. They may map to efficient hardware execution structures. The specification distinguishes subgroup membership from independent forward-progress guarantees, which are an optional device property in OpenCL 2.1.

14. Brook for GPUs: Stream Computing on Graphics Hardware

    Stanford's Brook work presented at SIGGRAPH 2004 addressed the difficulty of expressing ordinary computations through graphics APIs, textures and triangles. Brook supplied streams of records, kernels applied across their elements, and reductions combining records. Its compiler and runtime mapped these abstractions onto programmable graphics hardware. Kernel bodies retained temporary values in registers instead of requiring a separate memory round trip for every elementary operation. This connected a more accessible programming interface with locality, rather than treating GPU acceleration as arithmetic throughput alone.

15. Scalable Parallel Programming with CUDA

    Nickolls, Buck, Garland and Skadron's March/April 2008 account dates CUDA's release to 2007. It describes graphics-API limitations as motivation for a general parallel-computing interface. CUDA's central abstractions were hierarchical thread groups, shared memories and barriers: programmers decomposed problems into independently schedulable blocks, then cooperating threads within each block. Separating the logical decomposition from the number of physical processors let the runtime distribute blocks across GPUs with different available parallelism.

16. Roofline: An Insightful Visual Performance Model for Multicore Architectures

    Williams, Waterman and Patterson's April 2009 article proposed Roofline as an understandable bound-and-bottleneck model amid increasingly diverse multicore architectures. Its operational intensity counted operations per byte reaching DRAM after cache filtering, allowing locality improvements to change the bound. The authors measured sustainable memory bandwidth with optimized microbenchmarks. Their aim was to guide improvements, rather than predict every detail of runtime.

17. Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations

    The June 2019 Triton paper addressed operations that existing vendor libraries did not efficiently support, leaving researchers dependent on expert-written kernels. It made statically shaped multidimensional tiles the programming unit. Programmers expressed tile operations and the work assigned to program instances; compiler passes handled intra-tile parallelization and details such as coalescing and synchronization. This offered a middle ground between calling a fixed numerical primitive and explicitly coordinating individual CUDA threads, while still requiring more implementation work than a high-level matrix expression.

18. FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness

    FlashAttention organizes attention computation into tiles to reduce transfers between GPU high-bandwidth memory and on-chip SRAM. Its central optimization is reducing data movement while computing exact attention, rather than approximating attention to reduce the set of interactions. This distinguishes an implementation optimization from changing the model's attention pattern or replacing the decoding algorithm. The paper also describes a separate block-sparse extension, which is approximate and should not be confused with the exact algorithm.

19. CUDA Programming Guide: Writing SIMT Kernels

    Global allocations persist until freed and can hold inputs and outputs across launches. Shared memory is a block-accessible, programmer-managed on-chip scratchpad; registers hold thread-private state. Local memory is thread-private but physically resides in device memory; register spilling can add device-memory writes and reads. Coalescing concerns how a warp's addresses share memory transactions, not merely whether every thread reads a value. In the documented 32-byte-segment model, aligned contiguous four-byte loads across 32 lanes require four transactions; scattered addresses can require 32. Shared memory has separate bank constraints: the transpose example's 32-by-32 float tile has column-access conflicts, removed by padding its row length to 33.

20. GPU Performance Background User's Guide — NVIDIA

    GPU execution uses parallel processing units and a memory hierarchy. A simple performance model compares time to move the required bytes with time to perform the required arithmetic; with sufficient overlap, the slower term dominates. Arithmetic intensity is operations per byte accessed. Comparing it with the hardware's compute-to-bandwidth ratio helps identify a possible compute or memory bottleneck. But small workloads, insufficient parallelism, extra memory reads, and launch latency can invalidate the estimate. Changing batch size can change reuse and intensity, so model size and peak arithmetic throughput alone cannot predict request latency.

21. Running LLMs locally: Practical LLM Performance on DGX Spark — Mozhgan Kabiri chimeh, NVIDIA

    Memory capacity determines whether a model fits, but data-movement efficiency remains a separate constraint on throughput.

22. Nsight Systems User Guide

    Nsight Systems distinguishes a CPU-side CUDA API trace from GPU workload activity. The latter includes kernel executions and memory operations such as host-to-device copies, organized by GPU context and stream. This provides evidence for separating submission intervals from actual device execution and locating transfers or gaps between GPU operations.

23. PyTorch: Tensor Views

    Creating a tensor view changes interpretation of existing storage without moving its data. A transpose can therefore be cheap to create while producing a non-contiguous tensor whose later access pattern affects performance. Calling contiguous on a non-contiguous tensor creates a contiguous copy; on an already contiguous tensor it returns the tensor itself. Views also share mutations with their base tensors.

24. CUDA Programming Guide: Register Spilling and Shared-Memory Banks

    When thread-private state needs more registers than are available, the compiler can spill values into local memory, which physically resides in device memory; using the values later adds reads. In the documented shared-memory layout, successive 32-bit words map to successive banks. Distinct locations requested from one bank by a warp require serialized service; reading the same location can broadcast. A 32-word row stride maps column accesses to the same bank. Padding to 33 changes that mapping so the transpose example accesses different banks.

25. CUDA C++ Programming Guide 13.0.3: Multi-Stage Asynchronous Data Copies

    An asynchronous copy immediately followed by a wait does not overlap that transfer with computation. CUDA's two-stage example allocates separate shared-memory regions, submits the next copy and computes using the preceding completed region. Producer threads acquire and commit stages; consumers wait before accessing a stage and release it after use. If every stage is occupied, acquisition blocks until consumers release resources. The example primes the first transfer and drains the last computation, making startup, steady-state overlap and completion distinct phases.

26. NVIDIA Compute Sanitizer

    Compute Sanitizer separates memory-access checking, shared-memory race detection, initialization checks, and synchronization checks. Its block-reduction example shows a reader racing with other threads that populate shared memory; a block barrier orders those phases. A second example still races within a warp despite having a later block barrier, requiring warp-level synchronization at the earlier producer-consumer boundary. The synchronization examples include invalid divergent barrier participation.

27. CUDA Programming Guide: C/C++ Language Extensions

    An atomic read-modify-write protects an update to its target location. CUDA's legacy atomic functions provide relaxed ordering and scoped atomicity, without introducing fences for unrelated accesses. Unsuffixed atomicAdd has device scope; block and system variants have corresponding scopes and conditions. CUDA's newer atomic interfaces permit explicit memory-order and thread-scope selection.

28. CUDA C++ Programming Guide 13.0: Streams, Synchronization, and Events

    A CUDA stream orders its commands. Independent streams provide no general relative execution order. An event recorded in a specified stream completes after preceding work in that stream; cudaStreamWaitEvent makes later commands in a waiting stream delay until the event completes. Application to the existing split-K example: launch the combining kernel after its producers in one stream, or establish dependencies on every producer stream before combining. Merely recording one producer event does not establish completion of unrelated streams.

29. NVIDIA: Floating Point and IEEE 754

    Finite floating-point formats cannot represent every real value exactly, so operations round their results. Addition is consequently non-associative: changing parentheses can change the rounded answer even when every operation follows IEEE 754. NVIDIA supplies a worked single-precision example with different results for two addition orders. Fused multiply-add computes a product plus an addend with one rounding step instead of rounding the product and sum separately. Its cancellation example shows that fused evaluation can retain information lost by separate operations. Infinity and NaN have reserved encodings.

30. NVIDIA TensorRT: Accuracy Considerations

    FP32 means single-precision 32-bit floating point; FP16 is 16-bit half precision. BF16 is Brain Floating Point, a 16-bit format with a wider exponent field but fewer fraction bits than FP16. Format choice therefore changes representable range and precision. Overflow concerns results beyond a format's finite range; the guide describes FP16 overflow producing infinity and subsequent errors leading to NaN, meaning not-a-number. Storage format and the precision used for intermediate arithmetic can be selected separately rather than inferred from the input type.

31. PyTorch: Numerical Accuracy

    Matching input dtype does not fully specify a matrix operation's numerical behavior. PyTorch documents FP16 and BF16 GEMM implementations that occasionally truncate intermediate accumulations to reduced precision, potentially producing infinity even when the final mathematical result fits. Backend controls can disable these reductions. Its attention math backend instead upcasts FP16/BF16 inputs to FP32 for intermediate computation, then downcasts outputs, trading additional memory and potentially slower execution for accuracy. Batched and individually evaluated matrix products can also differ numerically despite representing the same mathematics.

32. PyTorch: torch.testing

    For finite real values, assert_close accepts an element when |actual−expected| ≤ atol + rtol·|expected|. Absolute tolerance therefore controls acceptance near a zero reference, while relative tolerance scales with reference magnitude. Infinities must match exactly; NaNs match only when explicitly enabled. Device, dtype, layout, and optionally strides have separate checks. A numerical closeness check therefore has both value and representation semantics.

33. Quantization and Training of Neural Networks for Efficient Integer-Arithmetic-Only Inference

    Affine quantization represents a real value approximately as r=S(q-Z), where q is an integer, S>0 is scale, and Z represents real zero. Encoding rounds to a discrete level and clips outside the representable range. The paper uses eight-bit weights and activations with wider accumulators and integer rescaling, enabling integer matrix products instead of floating-point inference. Training simulates quantization while retaining floating-point optimization, and activation ranges are estimated during training. Accuracy depends on ranges, rounding, and quantization-aware adaptation; deployment speed depends on efficient kernels and hardware support for the resulting arithmetic.

34. Low Level Technicals of LLMs

    Compare intermediate layer outputs against a reference implementation using a log L2 error plot, but establish the reference before automating variant searches.

35. AI Kernel Generation: What's Working, What's Not, What's Next

    Gate optimization on compilation, execution, and correctness, then use actual hardware measurements to guide further synthesis under human supervision.

36. PyTorch: Benchmark Utils

    PyTorch's benchmark Timer performs warmups for lazy initialization, controls CPU threadpool size, and synchronizes asynchronous accelerator work when necessary. It emphasizes repeated measurements to expose run-to-run variation and support median summaries. Statement and setup code are separate, and result labels and environment metadata help organize comparisons.

37. AI Kernel Generation: What's Working, What's Not, What's Next

    A useful kernel benchmark must define floating-point correctness and control input size, execution timing, warm-up, and caching effects.

38. CUDA Graphs — NVIDIA CUDA Programming Guide

    CUDA graphs represent operations as nodes and dependencies as edges. Stream capture records submitted work into a graph instead of executing it during capture. Instantiation validates and prepares an executable graph, which can then be launched repeatedly. Repeated graph launch reduces per-operation host setup and launch overhead; it does not remove the GPU computations. Capture coverage, graph updates and unsupported operations still constrain applicability.

39. Getting Started with CUDA Graphs

    Alan Gray's September 2019 experiment kept a 500,000-element floating-point scaling kernel unchanged on Tesla V100 with CUDA 10.1 and 512 threads per block. Device execution took 2.9 microseconds. Across 1,000 iterations of 20 launches, synchronizing after every kernel produced 9.6 microseconds per kernel including overhead; synchronizing after each group reduced this to 3.8 microseconds. Replaying a graph containing the same kernels yielded 3.4 microseconds, including amortized graph setup. The approximately 400-microsecond setup cost was paid once. The improvement changed submission and synchronization, not arithmetic.

40. Nsight Compute Profiling Guide

    Nsight Compute exposes memory requests, transactions, register-spill activity, and sampled warp states. A short-scoreboard stall can reflect shared-memory dependencies, special math, or branching, so the guide recommends checking memory evidence before diagnosing bank conflicts. A not-selected warp is eligible but another warp issued; a high count can indicate sufficient latency-hiding work. Profiling can serialize launches, replay kernels, adjust clocks, and flush caches between passes. Replayed device-memory state does not restore hardware cache contents, and disabling flushing changes reproducibility assumptions.

41. FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioning

    FlashAttention-2 addressed limitations remaining after reducing attention's off-chip traffic. Its forward pass distributes query-row tiles across independent thread blocks, providing more parallel work when batch size and head count are small. Within a block, assigning different query rows to warps avoids combining their partial outputs through shared memory. It also delays output normalization to reduce non-matrix arithmetic. Larger tiles remain a tradeoff: they reduce shared-memory accesses but increase registers and shared-memory demand; excessive register use can spill, and excessive shared-memory use can prevent launch.

42. CUTLASS: Efficient GEMM in CUDA

    CUTLASS nests threadblock, warp, and instruction tiles to reuse operands through shared memory and registers while accumulating output regions. Larger tiles reduce repeated global fetches but can waste work at boundaries or leave too few blocks. Parallel split-K divides the dot-product reduction dimension among blocks, then combines partial outputs using workspace and a second kernel. The epilogue applies elementwise work to accumulated results and can exchange data through shared memory for efficient output stores. Software pipelining double-buffers shared-memory tiles and register fragments to overlap movement with computation. The Hopper producer-consumer design waits for empty buffers before filling them and for filled buffers before consumption.

43. NVIDIA Launches Revolutionary Volta GPU Platform, Fueling Next Era of AI and High Performance Computing

    NVIDIA announced the Volta architecture and its first Volta processor, Tesla V100, on May 10, 2017. The announcement introduced Tensor Cores alongside CUDA cores to accelerate AI workloads.

44. NVIDIA Tesla V100 GPU Architecture

    NVIDIA motivated Volta's Tensor Cores by growing neural-network matrix workloads. The documented hardware operation computes D=A×B+C on 4×4 matrices, with FP16 multiplicands and FP16 or FP32 accumulation matrices. Larger matrix products are assembled from these operations. CUDA 9 exposed a different programming granularity: warp-level matrix operations spanning 16×16 matrices across 32 threads. cuBLAS and cuDNN also exposed Tensor Core support. Hardware operation size, programmer-visible fragments and application matrix shape are therefore distinct.

45. CUDA C++ Programming Guide 13.0: Warp Matrix and Shuffle Functions

    WMMA fragments distribute matrix elements across a warp using an unspecified internal mapping. mma_sync computes D=A·B+C collectively; every warp thread must participate with matching parameters. Supported shapes, input types, accumulator types, and layouts are constrained. The documented load requires 256-bit pointer alignment and leading dimensions divisible by eight half elements or four float elements. BF16 fragments require float accumulators. Separately, shuffle intrinsics exchange register values among participating lanes without shared memory; they do not supply a memory barrier.

46. What every AI engineer needs to know about GPUs

    The talk describes prompt processing as a better fit for GPU computation than decoding, where repeatedly moving model weights can dominate.

47. Luminal - Search-Based Deep Learning Compilers - Joe Fioti

    Fusion can remove intermediate global-memory writes and reads between dependent operations.

48. cuBLAS: cublasLtEpilogue_t

    cuBLASLt exposes matrix-result postprocessing through epilogues. ReLU applies max(x,0) independently to each result. The bias epilogue broadcasts a packed vector whose length matches output-matrix rows across all columns. Combined modes apply bias followed by ReLU or GELU. Auxiliary variants can additionally materialize activation-related values, showing that intermediate retention depends on required outputs.

49. FusionStitching: Boosting Execution Efficiency of Memory Intensive Computations for DL Workloads

    FusionStitching identifies conflicting memory layouts, parallelization choices and on-chip resource requirements as reasons naive fusion can slow execution. Its implementation transfers intermediates through registers for thread-local composition, warp reductions for suitable collective patterns, or shared memory for block-level composition. Fusion candidates are rejected when resource requirements cannot be satisfied or generated execution is unprofitable. Its evaluation weighs saved launches and off-chip traffic against fused-kernel execution time; complex patterns can require generating and timing a candidate rather than trusting a simple traffic model.

50. Luminal - Search-Based Deep Learning Compilers - Joe Fioti

    Luminal performs buffer reuse after search by assigning the same storage to intermediate buffers with nonoverlapping lifetimes.

51. Triton: Fused Softmax

    Softmax converts a row of scores into normalized exponential weights. The tutorial subtracts the row maximum, exponentiates, sums and divides. Its separate-operation example counts 5MN+2M element reads and 3MN+2M writes; the fused design reads MN elements and writes MN while retaining row intermediates on chip. It pads the computational row to a power of two, masks invalid loads with negative infinity and masks stores. The wrapper inspects register and shared-memory use to limit resident programs. An upstream test uses the irregular shape 1823×781 and compares against torch.softmax.

52. PyTorch: torch.Tensor.scatter_add_

    Scatter-add accumulates every source contribution into its indexed destination, including contributions sharing a destination. This differs from choosing one writer's value. PyTorch separately warns that CUDA execution may be nondeterministic: a specified summation result does not imply a fixed floating-point accumulation order.

53. CUDA Pro Tip: Optimized Filtering with Warp-Aggregated Atomics

    Warp aggregation combines participating threads' increments before one elected thread updates a shared counter, then distributes distinct output positions. In the article's K80/CUDA 8.0.61 filtering experiment on 100×2^20 integers, increasing the passing fraction from 5% to 50% reduced ordinary global-atomic bandwidth from about 55 to 8 GiB/s. The November 2017 update explains that CUDA 9 automatically applies aggregation in many cases, narrowing the manual optimization's advantage. A competing prefix-sum implementation preserves input order, while the atomic filter does not, so their performance comparison also involves different output contracts.

54. Root Mean Square Layer Normalization

    Zhang and Sennrich's NeurIPS 2019 paper proposed RMSNorm to reduce normalization overhead by omitting mean subtraction. It rescales each component by the vector's root mean square, sqrt(sum(x_i²)/n), then applies a learned componentwise gain. Unlike LayerNorm, it does not center the vector. For a kernel, the defining shared statistic is therefore a sum of squares rather than both a mean and centered variance. The paper investigated whether this simpler operation retained useful model behavior across its tested tasks.

55. PyTorch: RMSNorm

    PyTorch specifies RMSNorm as y_i=x_i·γ_i/sqrt(ε+sum(x_j²)/n), reducing over the trailing dimensions named by normalized_shape. Epsilon is added inside the square root for numerical stability. A single normalized dimension produces a rowwise reduction followed by rescaling; the output retains the input shape. Optional learned gains are per element of the normalized shape.

56. Low Level Technicals of LLMs

    The speaker presents normalization as a training-stability aid but warns that implementing its backward kernel is substantially more involved than its forward computation.

57. Online Normalizer Calculation for Softmax

    Milakov and Gimelshein's 2018 algorithm maintains a running maximum m and exponential sum d. On receiving x, it updates m'=max(m,x) and d'=d·exp(m−m')+exp(x−m'). Rescaling preserves earlier contributions when the maximum changes. Combining maximum and denominator calculation removes one input pass; producing all outputs still requires another pass. Their FP32 Tesla V100 experiment with CUDA 9.1 reported approximately 1.3× improvement over three-pass safe softmax at vector length 4000 and batch size 4000. With only ten vectors, the one-block-per-vector implementation underutilized the GPU and showed smaller gains.

58. FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness

    FlashAttention's tiled softmax combines block maxima and exponential sums after rescaling them to a common maximum. Weighted value outputs must receive corresponding rescaling before combination. Algorithm 1 carries row maxima, normalization sums and output accumulators instead of storing the full pairwise score and probability matrices. Its theorem retains O(N²d) arithmetic for sequence length N and head dimension d, while requiring O(N) additional memory beyond inputs and output. Backward computation regenerates score tiles from saved inputs and normalization statistics.

59. FlashAttention README: Causal-Mask and Small-Query Behavior

    FlashAttention's documented version-2.1 behavior aligns causal masks to the bottom-right corner when query and key lengths differ. Its five-query, two-key example fully masks the first three query rows and specifies zero outputs for them. Earlier version-2.0 alignment differs. The version-2.2 notes also describe splitting small-query attention across thread blocks and using a separate combining kernel, providing a concrete example of one attention operation requiring multiple kernels.

60. FlashAttention CUDA Softmax Implementation

    The CUDA softmax helper explicitly handles a maximum of negative infinity, which can occur when every score is masked. It substitutes zero for the subtraction's maximum term to avoid negative infinity minus negative infinity producing NaN. The online update separately retains the running maximum and, with its infinity-check specialization enabled, guards the rescaling calculation. Final normalization avoids dividing by a zero denominator. These branches show why the ordinary finite-score recurrence needs additional implementation logic for empty valid rows.

61. FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-precision

    The 2024 paper separates producer warps loading operand tiles from consumers computing with them, and overlaps matrix operations with softmax across iterations. On H100 80GB SXM5, its non-causal FP16 ablation uses batch 4, sequence length 8448, 16 heads and head dimension 128. Table 2 reports 3.538 ms with both techniques, 4.021 ms without matrix–softmax pipelining, and 4.105 ms without warp specialization. A three-stage variant performed worse than two stages: compiler scheduling failed to create the intended additional overlap, while extra live intermediates required more registers and smaller tiles.

62. Accelerating Embedding Lookups with cuEmbed

    An embedding lookup retrieves table rows selected by integer indices and may combine them by summation or another operation. cuEmbed demonstrates that irregular row selection can coexist with coalesced accesses within each row: different warps need not read neighboring rows. Its implementation increases memory-level parallelism—the number of memory operations in flight—using multiple row loads per thread, vectorized accesses and prefetched indices. Published experiments compare uniform and power-law index distributions. Frequently reused rows can remain in L1/L2 caches, so the rate of processing requested embedding bytes can exceed physical HBM bandwidth without moving that many bytes through HBM.

63. PyTorch: torch.Tensor.scatter_

    Scatter assignment writes source values to destinations selected by an index tensor. PyTorch warns that when several source values target the same destination, one value is selected nondeterministically; ordinary scatter assignment does not promise to combine them. Its gradient can also be incorrect for duplicate destinations. Index validity and duplicate-destination behavior therefore belong in the operation's contract before parallelizing or reordering writes.

64. cuEmbed: Embedding Operations and Index Transformations

    cuEmbed defines hotness as the number of table rows looked up for each batch sample. Forward lookup groups indices by sample. Embedding-gradient accumulation instead combines contributions destined for each table row, so its backward interface requires repeated table indices to be contiguous. The supplied transformation sorts table indices while carrying corresponding sample identifiers and optional weights. Compressed outputs additionally retain an inverse mapping to original table rows. Reordering therefore changes execution grouping while preserving the identities needed to interpret contributions and outputs.

65. Special topics in Kernels, RL, Reward Hacking in Agents

    Try Torch.Compile first and compare it with handwritten kernels on the actual workload and PyTorch version.

66. AI Kernel Generation: What's Working, What's Not, What's Next

    The reported agent struggled with highly optimized primitives and increasing problem complexity; it was not shown to replace expert algorithm design.

67. AI Kernel Generation: What's Working, What's Not, What's Next

    Equivalent high-level expressions can exploit better-optimized backend primitives or reduce operation launches without introducing custom kernels.

68. PyTorch: Dealing with Recompilations

    Compiled implementations can depend on assumptions about input shapes and exact scalar values. PyTorch documents recompilation when these assumptions change, and dynamic-shape options that can reduce some recompilations. Previously compiled variants remain reusable when their guards pass; after configured limits, unmatched calls can execute eagerly instead of triggering further compilation. Splitting a large graph can reduce the amount recompiled while imposing an execution-performance cost. These mechanisms make workload variability and compilation amortization part of the implementation choice.

69. Triton: triton.autotune

    Triton's autotuner evaluates a supplied list of kernel configurations when designated key arguments change. Configuration evaluation runs the kernel repeatedly, so an operation that updates existing values can apply those updates repeatedly during tuning. The API provides reset_to_zero and restore_value controls for this state. It also exposes configuration pruning, a benchmark function and optional disk caching of tuning timings. Selecting a configuration therefore entails real execution and setup costs, separate from subsequent steady-state calls.

70. Luminal - Search-Based Deep Learning Compilers - Joe Fioti

    Luminal constructs equivalent kernel alternatives with simple rewrite rules, profiles candidates, and uses search to select fast implementations.

71. Low Level Technicals of LLMs

    The speaker favors Triton as a compilation intermediary, but portability still depends on working compiler and attention-library backends.

72. Your Coding Agent Should Do AI System Engineering

    Kernel distribution needs explicit hardware and software compatibility metadata, not just kernel source.

73. From model weights to API endpoint with TensorRT-LLM

    Optimization profiles should reflect the intended workload because input shapes guide specialized kernel selection.

74. Vector Search Benchmark[eting]

    Benchmarketing can select favorable scenarios and aggregate results so that one specialized optimization appears to establish general superiority.

75. Validity of the Single Processor Approach to Achieving Large Scale Computing Capabilities

    Gene M. Amdahl's 1967 paper, written at IBM, challenged the assumption that adding parallel arithmetic units necessarily produces proportionate application gains. He emphasized sequential data-management work and irregularities such as boundaries, variable computation and storage rearrangement. His argument was that improving parallel processing alone leaves consequential work unchanged. This supplies the historical motivation for measuring the complete operation sequence after accelerating one kernel.

76. Vector Search Benchmark[eting]

    A flawed benchmark can suggest where a system may be strong, even when its overall conclusion is unreliable.

77. Luminal - Search-Based Deep Learning Compilers - Joe Fioti

    Luminal prepares a queue of kernels and submits them together to reduce repeated CPU dispatch delays.

78. Low Level Technicals of LLMs

    Compare corresponding layer outputs against a reference implementation and repeat the comparison at different precisions.

79. Why Rust is the Ideal Language for Vibe-Coding

    Tests generated after implementation may reproduce implementation details, and passing tests do not generally establish correctness for every input.

80. What every AI engineer needs to know about GPUs

    Evaluate arithmetic work relative to memory traffic, not operation count alone: GPUs favor substantial computation on data that has already been loaded.