LARK formally verifies LLM-written GPU kernels — Segment-wise Lean equivalence proofs — What does it mean?

The news. On November 1, 2026, researchers at Imperial College London, the University of Cambridge, the University of Edinburgh and EPFL posted LARK. It formalizes a GPU machine model and the meaning of each GPU operation in Lean 4 once per target language, then has agents prove that generated kernels match PyTorch reference code. Read the paper →

Picture a building inspector who refuses to read every blueprint line by line. Instead, experts certify a parts catalog once: every beam, every bolt, and how each one bears load. The inspector will stamp only a building assembled from that catalog. A fast contractor then builds whatever it likes, and the stamp says the building matches the blueprint, so the contractor is trusted for nothing. LARK works the same way for GPU kernels. The usual check runs a kernel on sample tensors and compares outputs, which in the authors' words "provides evidence but not a guarantee of correctness", and that is too weak when agents, working past the CUDA expertise gap, write kernels faster than people can review them.

What people check, once

The catalog has three parts, all fixed and checked by human experts before any kernel arrives, and none written by agents. The machine model describes a whole launch as one state: program counters, registers for every thread, global, shared and local memory, and the launch geometry. The per-operation semantics say exactly what each instruction does, and the authors measure them against the hardware. A translator composes these into a flat Lean model of an incoming kernel, line by line. Humans verify the operator-level definitions, not each generated line, which is the new trust boundary the paper claims. Memory in the model is unbounded, so an access outside the host's real allocations is outside the guarantee, and the task specification is formalized once per task and tested. The trust boundary also lists the premise and extractor definitions, so it is short but not empty.

Three Lean-checked stages

Stage 1 proves the kernel is free of data races. Kernels work in phases separated by barriers, and the model runs the threads of a phase one after another while real hardware runs them at once. The two agree only if no two threads in a phase write the same cell, or one writes what another reads. Lean checks this footprint condition for every phase and every pair of threads, and a kernel with an unproved lemma is rejected. Stage 2 segments the flat model. An agent rewrites the line-by-line model into a smaller term built from sequences, indexed loops and branches, and Lean checks that this segment-level model is equivalent to the flat one. Stage 3 proves each segment matches its PyTorch counterpart, and chaining the two equivalences gives whole-kernel functional and memory equivalence. Lean's error messages drive the agent loop at each stage.

A worked example: why segment first

Take the paper's running example. Fix three numbers: a matrix multiply of a 64×256 matrix by a 256×512 matrix, 32 blocks of 32 threads, and a loop of 16 slabs of k = 16 (256 / 16). Each block therefore produces 1,024 output cells, and 32 blocks cover all 32,768 cells of the 64×512 result. Inside the loop the flat model has 12 program-counter cases (4 ldmatrix loads and 8 mma instructions). Proving the flat model directly means reasoning about those 12 cases across 16 slabs, or 192 operation instances (illustrative count). The segment-level model replaces them with one loop body under an index j < 16, so the proof handles one body. The paper measures the effect on FlashAttention in PTX with Opus 5.5: the segmented route used 49M tokens for segmentation plus 120M for the equivalence proof, 169M in all, against 220M for the direct proof, which the paper reports as 29% more (the token totals shown here are rounded) and which crossed a token ceiling the default run finished under. Without segments, the agent re-proves the same loop body at every program counter.

What a certificate does not cover

The proof is about exact arithmetic, so floating-point rounding stays with ordinary tests. Each theorem also covers the kernel as adapted: nine of the ten production kernels carry a recorded departure, such as a 64-bit index read as 32 bits. The race check is sufficient but not necessary, so it can reject benign communication between threads as well as a harmful race, and the paper shows this in practice. In llm.c's softmax_forward_kernel7, the third pass reads a clamped address out[min(C-1, k)] that another thread of the same block writes in the same pass, with no barrier between. The loaded value is discarded where it matters, so no test sees the problem. The translator cannot express that address as a simple footprint, so the agent fallback searched for a proof, found a counterexample instead (thread 0 reads out[50256], which thread 80 writes), and the kernel was rejected. The paper's GEMM running example shows the stakes: its kernel listing marks a missing barrier before the shared-memory slab is re-staged, and the fixed version adds a bar.sync.

Goes deeper in: GPU & CUDA → Triton & torch.compile → The CUDA Expertise Gap

Related explainers

• Racer finds 82 GPU kernel bugs — a different route to equivalence: symbolic checking that returns a concrete input separating two kernels.