ThunderKittens: Embedded DSL for Fast AI Kernels

The Hardware-Software Chasm in Accelerated AI Infrastructure

In modern enterprise foundation model serving, high-throughput model pre-training, and real-time automated workflows powered by AI Fetchly, computational performance is directly dictated by how efficiently mathematical models map onto underlying accelerator hardware. Whether orchestrating distributed web extraction pipelines via AiFetchly Web Scraping, serving massive multimodal document copilots through real-time data fetching, or streaming long-context reasoning tokens, raw compute density governs infrastructure economics.

Over the past decade, semiconductor architecture has undergone a radical transformation. Traditional graphics processing units (GPUs) operated predominantly on Single Instruction, Multiple Threads (SIMT) vector pipelines, where programmers reasoned about individual threads executing floating-point operations across a 1D grid. However, the exponential explosion of deep learning workloads spurred hardware vendors to introduce dedicated matrix multiply units: Tensor Cores.

On modern datacenter architectures such as the NVIDIA Hopper H100 SXM5 and Blackwell B200, Tensor Cores do not merely augment general compute—they are the compute. On an H100 GPU, 16-bit Tensor Cores represent approximately 94% of the total available floating-point operations per second (FLOPs), delivering up to 1,979 TFLOPS of FP16/BF16 tensor compute with structured sparsity (or 989.5 dense TFLOPS), compared to a meager 60 TFLOPS across general-purpose FP32 vector ALUs. Consequently, if an AI kernel fails to route its computational throughput through Tensor Cores, its hardware utilization is capped at less than 6% of the GPU’s potential.

Yet, despite this reality, developing high-performance AI kernels that saturate modern hardware has become an intimidating engineering challenge. Hand-written CUDA kernels routinely underperform theoretical ceilings by 30% to 50%, even for standard operations like softmax attention, linear attention, and state-space models. FlashAttention-2 suffered a 47% performance regression when first ported from Ampere A100 to Hopper H100, and it took more than two years of intensive specialized systems research to release FlashAttention-3.

The Polarization of Modern GPU Programming

To understand why this software bottleneck persists, we must analyze the two extremes of the current kernel engineering landscape:

1. The Compiler Abstraction Wall (Triton, TVM, XLA): High-level Python-based domain-specific compilers like OpenAI Triton provide intuitive tensor-level block abstractions. While Triton works remarkably well for standard matrix multiplications, its compiler backend relies on automated heuristics for register allocation, instruction scheduling, and memory swizzling. When implementing non-standard architectures—such as recurrent state space models (Mamba-2), long convolutions (Hyena, FlashFFTConv), or fine-grained KV cache pruning—Triton frequently spills registers to slow local memory, fails to overlap asynchronous copy engines with compute, and cannot expose emerging microarchitectural hardware features like Hopper’s Tensor Memory Accelerator (TMA) or Warp-Group Matrix Multiply and Accumulate (WGMMA) instructions.
2. The C++ Template Metaprogramming Labyrinth (CUTLASS / CuTe): At the opposite extreme lies NVIDIA’s CUTLASS and its layout algebra sub-library, CuTe. CUTLASS provides bare-metal control over every byte of shared memory, every register allocation, and every asynchronous pipeline barrier. However, CUTLASS achieves this through tens of thousands of lines of deeply nested C++ template metaprogramming. Developing or modifying a custom kernel in CUTLASS requires deep expertise in CUDA layout algebra, nested iterators, and hardware barrier tokens. Even well-funded AI research labs treat CUTLASS codebases as brittle, impenetrable artifacts that only a handful of specialized systems engineers can navigate.

The Hopper Architectural Shift: WGMMA, TMA, and Asynchronous Execution

The release of NVIDIA’s Hopper (SM90) architecture widened this programming chasm by fundamentally changing the hardware execution model:

• Warp-Group Matrix Multiply (WGMMA): Prior architectures (Volta, Ampere) executed Tensor Core operations at the warp level (32 threads) via mma.sync. Hopper introduced WGMMA, which executes across an entire warp group—four contiguous warps (128 threads) working cooperatively on a single instruction. WGMMA operates asynchronously: the warp group issues a 2D matrix multiply instruction (wgmma.mma_async), and the hardware executes the calculation in the background while the issuing threads remain free to perform other work. Crucially, the accumulator registers are distributed across all 128 threads in non-trivial layout slices.
• Tensor Memory Accelerator (TMA): In previous generations, moving data from global High-Bandwidth Memory (HBM) into shared memory (SRAM) required threads to load data into registers (ld.global) and write it into shared memory (st.shared), or use asynchronous copy instructions (cp.async) that still required thread-level address calculation and loop unrolling. Hopper introduces TMA, a dedicated hardware DMA engine that transfers multi-dimensional tensor tiles directly between global memory and shared memory with zero register consumption and zero ALU address arithmetic.
• Hardware Asynchronous Barriers (mbarrier): Because memory transfers and tensor calculations execute out-of-band across decoupled hardware units, software thread synchronization (__syncthreads()) causes severe pipeline stalls. Hopper replaces software barriers with transaction-tracking hardware barriers (mbarrier), allowing threads to arrive, check transaction byte counts, and wait asynchronously.

When kernel developers attempt to stitch these asynchronous hardware primitives together manually, they face combinatorial complexity: managing multi-stage circular shared memory buffers, calculating layout swizzles to eliminate 32-bank memory serialization, and coordinating producer-consumer worker warps.

To resolve this crisis, researchers at Stanford University’s Hazy Research lab—including Benjamin F. Spector, Simran Arora, Aaryan Singhal, Daniel Y. Fu, and Christopher Ré—introduced ThunderKittens (TK) (arXiv:2410.20399, open-sourced on GitHub).

ThunderKittens demonstrates that developers do not need hundreds of thousands of lines of template abstractions to achieve peak GPU efficiency. By choosing a small, opinionated set of abstractions centered around 2D tiles, hardware-managed layouts, and an asynchronous Load-Compute-Store-Finish (LCSF) pipeline template, ThunderKittens allows developers to author kernels that match or exceed CuBLAS and FlashAttention-3 performance in fewer than 100 lines of readable, PyTorch-like C++.

The ThunderKittens Design Philosophy: Tiles, Layouts, and Asynchrony

ThunderKittens is an embedded C++ domain-specific language (EDSL) that lives directly inside standard CUDA source files. Unlike compiler-based frameworks, ThunderKittens requires no custom toolchains or Python codegen stages; it compiles directly with nvcc and grants the developer unrestrained access to native C++ features, inline PTX assembly, and CUDA runtime APIs.

At its core, ThunderKittens reorganizes the GPU programming model around three foundational principles that mirror the physical GPU hierarchy:

The Analytical Cost Model of GPU Execution

To understand how ThunderKittens structures its abstractions, the authors formulate a hardware execution cost model. In a modern GPU kernel, the total execution cost $C_\text{Overall}$ is governed by the overlap between memory transactions, arithmetic compute pipelines, and synchronization overheads:

$$C_\text{Overall} = \max\Big(\underbrace{C_\text{HBM}, C_\text{L2}, C_\text{L1}, C_\text{Shared}}{\text{Memory Latency & Bandwidth}}, \underbrace{C\text{Tensor}, C_\text{ALU}, C_\text{FMA}, C_\text{XU}}{\text{Compute Throughput}}\Big) + \underbrace{C\text{Setup} + C_\text{Sync}}_{\text{Overhead Latency}}$$

In an ideal kernel execution, memory movement occurs concurrently with arithmetic calculation, such that the actual execution time equals the maximum of the memory and compute components:

$$C_\text{Ideal} = \max(C_\text{Memory}, C_\text{Compute})$$

However, naive kernel implementations suffer from three critical degradations:

1. Serialization of Memory and Compute: Compute warps stall idly waiting for global memory loads to complete ($C_\text{Memory} + C_\text{Compute}$).
2. Shared Memory Bank Conflicts: Simultaneous thread accesses to identical shared memory banks serialize access, multiplying $C_\text{Shared}$ by up to 8x or 16x.
3. Synchronization Bubbles: Blocking barriers (__syncthreads()) force all execution units to drain their instruction pipelines, inflating $C_\text{Sync}$.

ThunderKittens systematically targets each term in this cost equation: it optimizes $C_\text{Shared}$ through automatic compile-time layout swizzling, drives $C_\text{Setup}$ and $C_\text{Sync}$ toward zero via persistent grids and hardware mbarriers, and enforces $\max(C_\text{Memory}, C_\text{Compute})$ overlap through its warp-specialized LCSF template.

Warp-Level Parallelism: The Tile Abstraction

In standard CUDA, programmers write scalar operations executed by single threads:

In ThunderKittens, the basic primitive is never a scalar thread. The fundamental unit of data is a Tile—a 2D matrix structure owned collaboratively by a warp (32 threads) or a warp group (128 threads):

• Register Tiles (kittens::rt): Data structures residing directly within physical GPU registers. Register tiles are parameterized by data type (float, bf16, half), spatial dimensions (height $M$ and width $N$, which must be multiples of 16), and memory layout (row_layout or col_layout). For example:
  • rt_bf<16, 64> represents a 16x64 tile of Brain Floating Point (BF16) values distributed across the register file of a single warp.
  • rt_fl<64, 64> represents a 64x64 float accumulator tile owned by a 4-warp warp group.
• Shared Memory Tiles (kittens::st): Data structures allocated within on-chip shared memory (SRAM). Shared tiles are parameterized by data type and shape:
  • st_bf<64, 64> represents a 64x64 BF16 tile staged in shared memory, ready to be consumed by WGMMA instructions or populated by TMA asynchronous copies.
• Global Memory Descriptors (kittens::gl): Multi-dimensional tensor descriptors in HBM (representing 4D tensors such as ${B, H, N, D}$), which configure hardware TMA address-generation tables automatically at compile time.

Eliminating Shared Memory Bank Conflicts via Automatic Swizzling

Shared memory on NVIDIA GPUs is divided into 32 distinct 4-byte (32-bit) banks. In any given clock cycle, each bank can service a 4-byte read or write. If two or more threads in a warp attempt to access different memory addresses that map to the same physical bank, a bank conflict occurs, and the hardware serializes the requests sequentially.

Consider a 16x64 tile of 16-bit values (BF16) in shared memory:

• Naive Row-Major Layout: Thread 0 accesses row 0 column 0 (bank 0). Thread 16 accesses row 1 column 0 (which, at 64 elements $\times$ 2 bytes = 128 bytes offset, maps exactly $128 \div 4 = 32$ banks later, landing back on bank 0). Loading a 16x64 tile into Tensor Core registers under a naive row-major layout triggers catastrophic 8-way bank conflicts, cutting effective shared memory bandwidth by 87.5%!
• Padded Layout: Developers frequently attempt to solve bank conflicts by adding dummy padding columns (e.g., allocating a 16x68 buffer). While this alters bank alignment, padding wastes precious shared memory capacity and breaks hardware TMA support, which requires strictly power-of-two contiguous memory alignments.

ThunderKittens solves this systematically by introducing mathematical swizzling. Instead of storing elements linearly, shared memory address offsets are permuted using bitwise XOR operations on row and column coordinate bits:

$$\text{Bank Index} = (\text{Logical Column}) \oplus (\text{Logical Row} \gg \text{Shift})$$

ThunderKittens provides three compile-time swizzled layouts: 32-byte, 64-byte, and 128-byte swizzling. The compiler automatically selects the widest swizzle supported by the tile geometry. By mapping data along diagonal XOR trajectories across physical banks, ThunderKittens completely eliminates bank conflicts during Tensor Core loads and WGMMA instructions without requiring a single byte of memory padding.

PyTorch-like Operands on Tiles

In ThunderKittens, tensor operations are expressed at the tile level using intuitive mathematical primitives that mirror PyTorch:

Notice how mma_ABt specifies an in-register Tensor Core matrix multiplication with transpose on matrix $B$. ThunderKittens enforces strict compile-time type and layout checking: if $A$ is not in row-major layout or $B$ is not in column-major layout, the CUDA compiler produces a descriptive static assertion error rather than compiling invalid assembly that fails silently at runtime.

The Block-Level LCSF Asynchronous Pipeline Template

At the thread-block level, the defining architectural innovation of ThunderKittens is the Load-Compute-Store-Finish (LCSF) template.

Traditional CUDA kernels interleave memory loads and arithmetic computations sequentially within a single thread loop:

This lock-step alternation destroys GPU throughput. Modern datacenter GPUs feature independent hardware subsystems: the TMA copy engine can transfer gigabytes of data over the high-bandwidth memory bus while the Tensor Cores execute floating-point matrix multiplications at full speed.

Producer-Consumer Warp Specialization

The ThunderKittens LCSF template coordinates execution by dividing the warps within a thread block into specialized roles:

1. Producer Warps (Load / Store Workers): Warp 0 does not perform floating-point matrix math. Its sole responsibility is to orchestrate memory movement. It interacts with the Tensor Memory Accelerator (TMA) to issue asynchronous bulk copy instructions from HBM into circular shared memory buffers.
2. Consumer Warps (Compute Workers): Warps 1 through 4 (forming a 128-thread warp group) focus exclusively on compute. They wait for shared memory buffers to become populated, stream data into registers or issue WGMMA matrix multiply instructions directly against shared memory, perform necessary elementwise reductions (such as softmax or activations), and write results to output buffers.
3. Multi-Stage Circular Buffering: The LCSF template maintains an $N$-stage circular buffer in shared memory. While the consumer warps are crunching data in stage $k$, the producer warp is concurrently loading input data for stage $k+1$ and writing completed outputs from stage $k-1$ back to global memory.

Empirical Scaling of Multi-Stage Buffers

The impact of buffer depth on arithmetic throughput is profound. In their empirical benchmarks on an NVIDIA H100 GPU running a standard matrix multiplication ($M=N=K=4096$), the authors evaluated efficiency as the number of pipeline buffer stages $N$ varied from 1 to 4:

Pipeline Buffer Stages ($N$)	H100 GEMM Throughput (TFLOPS)	Memory Latency Overlap	Efficiency vs. Single-Stage
1 Stage (Synchronous Baseline)	260 TFLOPS	0% (Strict serial alternation)	1.00x (Baseline)
2 Stages (Double Buffering)	484 TFLOPS	~50% (Next tile loads during compute)	1.86x
3 Stages (Triple Buffering)	683 TFLOPS	~85% (HBM transfer bubbles absorbed)	2.63x
4 Stages (Quad Buffering)	760 TFLOPS	~98% (Near-perfect compute saturation)	2.92x

With a single buffer stage, the kernel achieves only 260 TFLOPS because Tensor Cores stall whenever data is read from global memory. By scaling to a 4-stage circular pipeline managed by hardware transaction barriers, throughput nearly triples to 760 TFLOPS—saturating the physical limits of the H100 SM.

Writing an Attention Kernel in the LCSF Template

To see how cleanly this abstraction manifests in code, examine how a complete FlashAttention-style forward step is expressed in ThunderKittens. The developer simply implements four concise methods:

In fewer than 50 lines of code, this implementation achieves:

• Asynchronous hardware TMA transfers with zero register overhead.
• Hardware mbarrier transaction tracking.
• WGMMA asynchronous Tensor Core execution across 128 threads.
• Shared memory bank-conflict-free 128-byte swizzled memory layouts.
• Perfect overlap between global memory transfers and matrix multiply compute.

Grid-Level Optimization: Persistent Grids and L2 Cache Reuse

Beyond the warp and block levels, high-performance GPU execution requires optimizing the Grid level—how thread blocks are scheduled and dispatched across the GPU’s streaming multiprocessors (SMs).

An NVIDIA H100 SXM5 GPU contains 132 physical SMs. Thread blocks dispatched to these SMs share a massive 50 MB on-chip L2 cache (12 TB/s bandwidth) and an 80 GB High-Bandwidth Memory (HBM3) pool (3.35 TB/s bandwidth).

ThunderKittens introduces two vital grid-level optimizations: Persistent Grids and L2-Aware Block Launch Scheduling.

Persistent Grids: Amortizing Block Launch Latency

When a standard CUDA kernel is launched with thousands of thread blocks, the GPU hardware scheduler must repeatedly allocate register files, configure shared memory partitions, and map thread blocks to available SMs. This setup and tear-down cycle introduces significant latency overhead ($C_\text{Setup}$). Furthermore, if a grid size does not divide evenly into 132 SMs (e.g., launching 133 blocks), the GPU executes the first 132 blocks at 100% occupancy and then runs a trailing “wave” with only 1 active block—wasting 99.2% of the GPU’s compute during the second wave!

ThunderKittens addresses this through a Persistent Grid architecture. Instead of launching thousands of transient thread blocks, ThunderKittens launches exactly 132 persistent thread blocks—one per physical SM. Each block enters a persistent while loop that pulls work chunks atomically from a global task queue. While one chunk executes its finish() stage, the producer warp pre-emptively loads the input tiles for the next work item into shared memory.

The performance gains across varying matrix dimensions ($M=N=4096$) demonstrate how persistent grids prevent latency collapse on small and medium GEMM workloads:

Matrix Dimension ($K$)	Standard Launch (TFLOPS)	Persistent Launch (TFLOPS)	CuBLAS (TFLOPS)	Persistent Speedup
$K = 64$	93 TFLOPS	108 TFLOPS	69 TFLOPS	+16.1% over standard (+56% vs CuBLAS)
$K = 128$	161 TFLOPS	184 TFLOPS	133 TFLOPS	+14.3% over standard (+38% vs CuBLAS)
$K = 256$	271 TFLOPS	309 TFLOPS	242 TFLOPS	+14.0% over standard (+28% vs CuBLAS)
$K = 512$	414 TFLOPS	450 TFLOPS	407 TFLOPS	+8.7% over standard (+10% vs CuBLAS)
$K = 1024$	565 TFLOPS	600 TFLOPS	633 TFLOPS	+6.2% over standard (Matches CuBLAS)

L2 Cache Reuse: Spatial Block Traversal Scheduling

In large-scale matrix operations and attention mechanisms, multiple thread blocks require access to overlapping slices of global memory. For example, during attention forward computation, multiple query blocks across sequence length $N$ access identical Key and Value tensors.

If thread blocks are launched in naive row-major order, by the time the GPU reaches the next batch of query blocks, the previously loaded Key/Value tiles have already been evicted from the 50 MB L2 cache, forcing the memory subsystem to re-fetch identical data across the slow 3.35 TB/s HBM bus.

ThunderKittens incorporates spatial scheduling swizzles (such as Hilbert space-filling curves and 8-row blocked swizzles) to order block launches such that blocks sharing identical weights execute concurrently across adjacent SMs:

Workload	Block Scheduling Order	Consumed HBM Bandwidth	H100 Efficiency (TFLOPS)	Performance Impact
GEMM ($M=N=K=16384$)	Swizzled ${8, N, M/8}$	982 GB/s	805 TFLOPS	2.05x Speedup (L2 hit rate maximized)
GEMM ($M=N=K=16384$)	Naive ${N, M}$	3,070 GB/s	392 TFLOPS	Memory bandwidth thrashing
Attention Forward ($D=128$)	Swizzled ${N, H, B}$	213 GB/s	600 TFLOPS	1.21x Speedup (Zero cache thrashing)
Attention Forward ($D=128$)	Naive ${B, H, N}$	2,390 GB/s	494 TFLOPS	Extreme HBM bus saturation

By optimizing block launch order alone, ThunderKittens reduces consumed HBM bandwidth from 3,070 GB/s to 982 GB/s on large matrix multiplies, instantly boosting throughput from 392 TFLOPS to 805 TFLOPS!

Empirical Benchmarks: ThunderKittens vs. Frontier Frameworks

To evaluate the real-world performance of ThunderKittens, the Stanford team benchmarked custom TK kernels against the most heavily optimized production baselines available, including CuBLAS, FlashAttention-2, FlashAttention-3, CUTLASS 3.x, and Triton, on an NVIDIA H100 SXM5 80GB GPU.

Deep Dive: Attention Backward Superiority

While FlashAttention-3 achieves exceptional forward performance on H100 GPUs, its backward pass exhibits an unexpected performance ceiling. Profiling reveals that FlashAttention-3’s backward kernel, authored in CUTLASS, experiences shared memory bank conflicts when computing query, key, and value gradients simultaneously.

In contrast, ThunderKittens automatically enforces 128-byte swizzled memory layouts on all backward gradient shared memory tiles. The profiling comparison on NVIDIA H100 SXM5 reveals the architectural difference:

Profiling Metric	FlashAttention-3 Backward	ThunderKittens Backward	Difference / Advantage
Tensor Core Utilization	61.2%	58.2%	Comparable arithmetic density
Issue Slot Utilization	25.1%	34.8%	+38.6% higher instruction throughput
Throughput (TPS)	328 GB/s	490 GB/s	+49.4% higher memory throughput
HBM Stalls (Cycles)	1.83 cycles	1.63 cycles	10.9% lower memory stall time
Shared Memory Stalls (Cycles)	0.92 cycles	0.14 cycles	6.57x reduction in shared memory stalls!
End-to-End Backward Speed	Baseline	+28.9% Faster	Resolves CUTLASS bank conflicts

Because ThunderKittens reduces shared memory stalls from 0.92 cycles per instruction down to 0.14 cycles, the kernel delivers a massive 28.9% end-to-end speedup over FlashAttention-3 on the backward pass.

Emerging Architectures: Mamba-2 and Linear Attention

Where ThunderKittens truly shines is on emerging, non-standard architectures where established vendor libraries (like cuDNN) offer no off-the-shelf kernels:

• Mamba-2 Chunk State-Space Models (SSMs): Mamba-2 replaces traditional softmax attention with state-space duality (SSD). In Triton, writing the 1-semiseparable matrix transformation requires scalar loops that fail to utilize Hopper’s WGMMA hardware. ThunderKittens maps the semi-separable matrix operations directly onto 16x16 register tiles, running the core chunk recurrence on Tensor Cores to achieve 360 TFLOPS—an 8.0x speedup over the official Triton implementation.
• Linear Attention (Based & Hedgehog): Linear attention replaces quadratic $O(N^2)$ softmax attention with Taylor expansions or feature maps $\phi(Q)\phi(K)^T V$. In standard CUDA and Triton, maintaining the cumulative state matrix causes severe register pressure. ThunderKittens schedules the state accumulation across register tiles, achieving 350 TFLOPS—a 14x acceleration over baseline implementations.

Framework Footprint and Maintainability

Beyond raw speed, ThunderKittens radically reduces code bloat and cognitive overhead for engineering teams:

Metric / Attribute	NVIDIA CUTLASS 3.x	OpenAI Triton	Stanford ThunderKittens
Core Library Size	~22.0 MB	~12.6 MB	\< 1.0 MB (Header-only C++)
Lines of Code for Attention	2,500+ lines	~150 lines	~85 lines
Compile Time (Single Kernel)	45 - 120 seconds	JIT runtime overhead	5 - 12 seconds (nvcc)
Hardware Primitives Exposed	TMA, WGMMA, mbarrier	Indirect via IR	Direct C++ Tile API
Non-GEMM Architecture Support	Extremely difficult	Moderate (often slow)	Intuitive & Extensible

Production Implementation: Building a Custom ThunderKittens Extension

To deploy ThunderKittens kernels in production AI infrastructure—such as vLLM, SGLang, HuggingFace transformers, or internal inference runtimes—engineers wrap the C++ kernel in a PyTorch C++ extension using pybind11.

Below is a complete, production-ready walkthrough showing how to author, bind, and invoke a ThunderKittens-accelerated linear attention kernel.

Step 1: The ThunderKittens Kernel (linear_attn_tk.cu)

Step 2: The PyTorch C++ Binding (binding.cpp)

Step 3: PyTorch Build Configuration (setup.py)

Step 4: Python Benchmark & Verification (benchmark.py)

Production Readiness, Best Practices, and Future Trajectory

Deploying custom ThunderKittens kernels within enterprise serving clusters requires observing crucial systems engineering practices:

1. Compile-Time Architecture Specialization (sm_90a)

Hopper-specific instructions—particularly WGMMA (warpgroup::mma_AB) and TMA (tma::load_async)—require compiling with architecture flag -arch=sm_90a (not generic -arch=sm_90). The a suffix enables asynchronous architecture extensions. If deploying on Ada Lovelace (RTX 4090 / L40) or Ampere (A100), configure ThunderKittens to use warp-level mma_sync and cp.async pipelines instead of Hopper-exclusive warp groups.

2. Register Budgeting and Occupancy Tuning

An NVIDIA H100 SM contains 64,512 32-bit registers, capped at 255 registers per thread. When launching a 128-thread warp group, a single consumer block consumes $128 \times 255 = 32,640$ registers—exactly half of the SM’s physical register file.

• To maintain high occupancy, ensure your consumer warps do not exceed 128 registers per thread.
• ThunderKittens enforces strict register typing: compile-time dimensions in global descriptors (gl) are encoded into the instruction cache (I-cache) rather than consuming registers, preventing accidental register spills to L1/L2 cache.

3. ThunderKittens 2.0 & Cross-Hardware Evolution

As foundation models expand beyond standard FP16/BF16 into microscopic precisions (FP8, FP4, and 1-bit ternary networks), the ThunderKittens project has expanded aggressively:

• ThunderKittens 2.0 & FP8/NVFP4: Introduced native Microscaling FP8 (MXFP8) and NVFP4 tensor core primitives for NVIDIA Blackwell (B200 / GB200) architectures, exposing Blackwell’s fifth-generation Tensor Core generation instructions (TCGEN05).
• ParallelKittens (Multi-GPU): Extends tile primitives across NVLink network fabrics, enabling tile-granularity tensor communications that overlap cross-GPU All-Gather collectives with local Tensor Core compute.
• HipKittens & Apple MLX (ThunderMittens): Ported tile abstractions to AMD CDNA3 (MI300X) via ROCm/HIP and Apple Silicon unified memory via Metal/MLX, proving that the 2D tile abstraction transcends vendor-specific hardware boundaries.

By bridging the divide between compiler-level accessibility and bare-metal hardware saturation, ThunderKittens demystifies AI kernel engineering—transforming GPU acceleration from alien technology into an accessible, maintainable engineering craft.