--- title: 'ProofWright: Automated CUDA Verification' url: https://www.emergentmind.com/topics/proofwright type: topic --- # ProofWright: Automated CUDA Verification Searching arXiv for recent papers on ProofWright and closely related work. {"query":"ProofWright formal verification CUDA arXiv", "max_results": 10} Searching arXiv for expert proof-writing process work related to ProofWright. {"query":"What's in a Proof analyzing expert proof-writing processes F* Verus arXiv", "max_results": 10} ProofWright is an agentic verification framework for LLM-generated CUDA kernels that integrates automated formal verification with LLM-based code generation in order to address a validation bottleneck: kernels can be produced rapidly, yet subtle correctness bugs, race conditions, and reward-hacking behaviors may evade ordinary testing. Its stated target is end-to-end guarantees of memory safety, thread safety, and semantic correctness for generated GPU code; in its KernelBench L1 evaluation, it verifies safety properties for 74% of generated kernels, reports a modest overhead of about 3 minutes per kernel, and establishes semantic equivalence for a class of element-wise kernels [2511.12294]. ## 1. Problem setting and motivation ProofWright is situated in the verification of CUDA kernels synthesized by modern LLM-based coding agents. The motivating claim is that generation and optimization of GPU kernels have become fast enough that trustworthy validation, rather than code production, has become the limiting factor. The framework is therefore designed for a paired input setting: a high-level PyTorch program as intended specification, and an LLM-generated CUDA kernel intended to implement it [2511.12294]. The framework’s motivation rests on three limits of conventional testing. First, runtime testing has limited input coverage. CUDA failures may appear only for rare tensor sizes, alignment conditions, tail cases, grid and block configurations, or synchronization schedules. The paper’s concrete example is an LLM-generated sigmoid kernel whose vectorized tail-handling code causes multiple blocks with `threadIdx.x == 0` to write the same output cell when \(N\) is not a multiple of 4; the bug passed compilation, multiple unit tests, and, in the reported setup, NVIDIA Compute Sanitizer’s `racecheck` [2511.12294]. Second, testing can be vulnerable to reward hacking: the paper presents a vector-add kernel that simply copies reference outputs from device memory, thereby passing tests without implementing the intended computation. Third, manual formal verification is reliable but does not scale to the volume of kernels produced by LLM systems [2511.12294]. CUDA is treated as especially difficult to verify because it combines massive concurrency, a complex memory hierarchy, raw pointers and manual indexing, synchronization primitives such as `__syncthreads()`, low-level optimized indexing schemes, and cross-thread interference. As a result, a credible verifier must reason simultaneously about out-of-bounds accesses, permission discipline, synchronization structure, and the mapping from threads to data [2511.12294]. ## 2. System architecture ProofWright has two major components: a **VerCors Agent** for memory safety and thread safety, and a **Semantic Equivalence Framework** for functional correctness relative to the PyTorch specification [2511.12294]. Before describing the components individually, it is useful to separate the trusted and synthesized parts of the workflow. The framework does not treat the LLM as the verifier. Instead, the LLM is used to synthesize contracts, annotations, helper functions, and proof artifacts, while trust is delegated to formal checkers: VerCors for deductive verification of CUDA safety and certain functional postconditions, and Rocq for machine-checked theorem proving [2511.12294]. | Component | Primary mechanism | Purpose | |---|---|---| | VerCors Agent | VerCors + LLM-guided annotation synthesis | Memory safety and thread safety | | Semantic front end | `torch.fx` + MLRocq | Trusted extraction of PyTorch specifications | | Semantic back end | Rocq + VerCors + LLM-guided synthesis | Semantic equivalence and implementation checking | The VerCors Agent includes a knowledge base consisting of distilled VerCors documentation, few-shot verified CUDA examples, and error/fix examples for common VerCors failures. It also maintains an **Annotation Guide**, described as a dynamic, LLM-generated summary of learned verification procedures and strategies, plus a minimal database of known errors and fixes. Its workflow is iterative: minimally preprocess the CUDA kernel into a form compatible with VerCors, synthesize safety annotations and helper functions, invoke VerCors, inspect verifier feedback, revise the annotations, and retry [2511.12294]. The Semantic Equivalence Framework is divided into a front end and a back end. The front end deliberately uses no LLMs. It applies a PyTorch Static Analyzer based on `torch.fx` to symbolically trace the model, build a graph-based IR, and enrich nodes with tensor metadata. The resulting graph is translated into Rocq terms using the MLRocq library, which formalizes tensor types and many machine-learning operations. The back end then uses a Semantic Equivalence Agent to synthesize Rocq representations corresponding to the CUDA implementation, generate equivalence theorems and tactics, prove them in Rocq, lower the proved properties into VerCors postconditions, and finally check that the CUDA implementation satisfies those postconditions [2511.12294]. ## 3. Safety verification model The safety side of ProofWright is based on VerCors’s permission-based concurrent separation logic. The framework uses the standard VerCors interpretation in which permission `1` grants write access, any non-zero fractional permission grants read access, and race freedom follows from the invariant that the total permission for a memory location never exceeds `1` [2511.12294]. Contracts are expressed over all launched threads. The system attempts to infer “the weakest pre-conditions possible” under which the contract remains satisfiable and safety can be proved, although it does not claim a proof that those inferred preconditions are truly weakest. These preconditions typically include non-nullness, pointer lengths, positivity and size constraints on dimensions, launch constraints for `blockDim` and `gridDim`, and the distribution of read and write permissions across memory locations [2511.12294]. A central reasoning pattern is explicit thread-to-data mapping. The canonical global-thread index is $$ \text{idx} = \text{blockIdx.x} \cdot \text{blockDim.x} + \text{threadIdx.x}, $$ with analogous formulas for multidimensional kernels such as $$ \text{row} = \text{blockIdx.y} \cdot \text{blockDim.y} + \text{threadIdx.y}, \qquad \text{col} = \text{blockIdx.x} \cdot \text{blockDim.x} + \text{threadIdx.x}. $$ These mappings determine both bounds obligations and disjointness of writes [2511.12294]. Because raw indexing expressions are often hostile to SMT solving, the agent synthesizes pure helper functions with contracts. The paper’s example is a 4D flattening function whose postconditions simultaneously specify the arithmetic formula and prove that the resulting offset lies within the allocated tensor extent. The agent also inserts loop invariants such as `0 <= k && k <= N` and local assertions that relate helper functions back to concrete pointer arithmetic. For synchronized kernels, the agent must additionally express permission transfer around `__syncthreads()`, since barriers can redistribute permissions across threads [2511.12294]. The resulting guarantee is conditional but universal: if VerCors accepts the synthesized contract, then any invocation satisfying the inferred precondition is memory-safe and race-free. When verification fails, the paper adopts a conservative operational stance that such a kernel should be treated as unsafe to invoke, while explicitly distinguishing this from a formal proof of unsafety [2511.12294]. ## 4. Semantic equivalence framework ProofWright’s semantic layer formalizes the high-level specification in Rocq and attempts to prove that the generated CUDA kernel implements the same mathematical function. The front end begins with `torch.fx` tracing, converts the PyTorch program into a graph IR, and then maps the graph into MLRocq, a Rocq library whose core tensor type is defined recursively as ```coq Fixpoint TensorND (n:nat):Type:= match n with | 0%nat => Z | S n' => list (TensorND n') end. ``` In this formulation, a \(0\)-dimensional tensor is a scalar `Z`, and an \((n+1)\)-dimensional tensor is a list of \(n\)-dimensional tensors [2511.12294]. The paper illustrates the method with ReLU. A scalar definition ```coq Definition relu_sc (x:Z) : Z := if x