Skip to content
VeriTile
Menu

TRITON KERNELS · LEAN 4

Formal Verification for Triton Kernels.

Lean-checked contracts for a typed Triton-style model, with explicit semantics, mathematical specifications, and proof generation using agents.

Theorems cover inputs and parameters that satisfy their stated preconditions.

Contracts cover the translated model and its stated arithmetic assumptions. What our proofs cover →

Triton embedded in Lean

VeriTile embeds Triton in Lean, so users can write kernels in familiar Triton-style syntax. The DSL supplies the formal representation and execution semantics. Lean wrappers and explicit translation choices connect a kernel to its specification; agents can help construct the proof.

TRITON KERNEL

Written by the kernel author

pid     = tl.program_id(axis=0)
offsets = pid * B + tl.arange(0, B)
x       = tl.load(x_ptr + offsets)
y       = tl.load(y_ptr + offsets)
output  = x + y
tl.store(out_ptr + offsets, output)

INTERNAL LEAN REPRESENTATION

Used internally for verification

pid     := tl.program_id(0)
offsets := pid * $(B) + tl.arange(0, $(B))
x       := tl.load($(x_ptr) + offsets)
y       := tl.load($(y_ptr) + offsets)
output  := x + y
tl.store($(out_ptr) + offsets, output)

Aligned vector addition (no masked tail). Names matched for comparison.

Kernel correctness and equivalence

Verify computation and memory effects against a specification or a reference implementation.

KERNEL → SPECIFICATION

Correctness

A kernel’s output values and memory effects match its specification.

lse_kernel ⊨ lse_spec
Computation
Outputs equal the mathematical function defined by the specification.
Memory
Results occupy the specified output addresses. Memory outside the declared write regions is preserved.

Example · Log-sum-exp

The LSE kernel stores the specified reduction result at each declared output address.

Proof details: Log-sum-exp

out= log Σ exp(xᵢ)

The formula shown omits input scaling and sums over the active lanes of one block.

Arithmetic is modeled over real numbers. For admissible block widths and row lengths, the kernel terminates with the specified value at its output address and leaves every other memory cell unchanged. Active reads and the output store must be in bounds.

View proof: Log-sum-exp ↗

KERNEL ↔ KERNEL

Equivalence

Two kernels agree on output values and observable memory effects.

direct_lse ≡ stable_lse
Computation
Outputs equal those of the other implementation under the same numerical model.
Memory
Both kernels agree at the specified output addresses. Memory outside the declared write regions is preserved.

Example · Direct and stable log-sum-exp

The direct and stable LSE kernels store equal reduction results at each declared output address.

Proof details: Direct and stable log-sum-exp

log Σ exp(xᵢ)= m + log Σ exp(xᵢ − m)

The stable form subtracts m, the maximum input value in a nonempty block.

Intermediate arithmetic is modeled over real numbers. Both kernels round only at the final bf16 store, using the same model R. For a nonempty block with input and output windows in bounds, both terminate from the same state, agree at y[pid], and leave every other memory cell unchanged.

View proof: Direct and stable log-sum-exp ↗

Schematic notation. See Proof details for scope and View proof for the Lean source.

VectorAdd exampleCompare the original kernel and a subtraction variant using recorded Lean checks.

VectorAdd verification example

Switch the operation below. The specification still requires addition.

The values illustrate one input. The addition theorem covers every input allowed by its preconditions.

VectorAdd.lean
triton {
  pid  := tl.program_id(0)
  offs := pid * $(blockSize) + tl.arange(0, $(blockSize))
  x    := tl.load($(xReg) + offs)
  y    := tl.load($(yReg) + offs)
  out  := x + y
  tl.store($(outReg) + offs, out)
}
ONE TILE · FOUR LANES
Values0123
x1234
y4567
Output57911
Expected57911

All four lanes match.

THE CONTRACT
addIO B ⊨ fun xs ys i => xs i + ys i

Proof checked

Lean checks the general addition proof under the stated preconditions.

Lean v4.29.0Recorded Lean checks
Inspect the proof and Lean output

The values illustrate one input. The addition theorem covers every input allowed by its preconditions.

Aligned tiles, mathematical values, and the stated memory preconditions.

  • Each output equals the sum of its inputs.
  • Execution terminates.
  • Memory outside the output window is preserved.
specification add_kernel_correctness (B : Nat) (hB : 0 < B) :
    addIO B ⊨ fun xs ys i => xs i + ys i := by
  refine KernelIO₂.Implements.intro _ ?_ ?_ ?_
  · exact addKernel_flattenOk ⟨"x"⟩ ⟨"y"⟩ ⟨"out"⟩ B
  · intro bounds s h1 h2 h3
    exact addKernel_traceSafe ⟨"x"⟩ ⟨"y"⟩ ⟨"out"⟩ B bounds s h1 h2 h3
  · intro s₀ xs ys hx hy
    exact addKernel_region_run B hB s₀ xs ys hx hy

Lean output · exit 0

VectorAdd-add.lean:139:4: warning: This simp argument is unused:
  evalOp

Hint: Omit it from the simp argument list.
  simp [exec, addKernel, ComputeKernel.toAlgKernel, stepStmts, stepStmt,
  ̵  ̵ ̵ ̵e̵v̵a̵l̵O̵p̵,̵ ̵Tile.bop, NumericDType.add,
  ̲  ̲ ̲ ̲NumericDType.mul] at hExec

Note: This linter can be disabled with `set_option linter.unusedSimpArgs false`
VeriTile.Bench.Examples.VectorAdd.add_kernel_correctness: axiom footprint ⊆ standard base ✓  ([propext,
 Quot.sound,
 Classical.choice])
VeriTile.Bench.Examples.VectorAdd.add_kernel_correctness: statement's project surface ⊆ allowlist ✓

Two fixed variants, checked ahead of time. The four sample values are also proved from the DSL semantics.

Reproduce these checks ↗

Proof generation with agents

The agent uses VeriTile’s reusable lemmas to construct candidate proofs and revises them with feedback from Lean.

Proof generation

Start with a kernel and its specification, then refine a proof with Lean.

Kernel + specification

  1. Proof agent

    Select lemmas and construct a candidate proof.

  2. Lean checker

    Check the proof and report remaining goals.

Lean diagnostics → revise the candidate

WHEN THE PROOF PASSES

Checked Lean proof

FlashAttention and benchmark coverage

The following artifacts include forward FlashAttention contracts and a benchmark corpus with per-kernel audit records.

FEATURED PROOF · FLASHATTENTION

FlashAttention Forward

softmax(QKᵀ · scale) V

  • Streaming softmax
  • Online rescaling
  • Tiling
  • Causal masking
  • Boundary & D-tail

The forward contracts relate tiled, online-softmax attention to a direct attention reference. They cover causal and non-causal variants, with boundary and D-tail conditions.

View example
FlashAttention ForwardKernel, specification, and proof.View exampleClose example

This example follows the non-causal kernel with sequence-boundary and head-dimension masks.

Kernel

Online softmax update

Inside the loop, after loading a tile of K and V.

scores_raw := tl.dot(q, tl.trans(k)) * $(scale)
score_mask := (offs_m[:, None] * $(0) + offs_n[None, :]) < $(S_k)
scores     := tl.where(score_mask, scores_raw, -inf)
m_block    := tl.max(scores, axis = 1)
m_new      := tl.max(m_i, m_block)
alpha      := tl.exp(m_i - m_new)
p          := tl.exp(scores - m_new[:, None])
l_new      := alpha * l_i + tl.sum(p, axis = 1)
o_acc      := alpha[:, None] * o_acc + tl.dot(p, v)
m_i        := m_new
l_i        := l_new
Running maximum
Combine the previous maximum with the current tile’s scores.
Rescaling
Rescale the accumulated sum and output before adding the new tile.
Output
After the loop, divide the accumulator by the normalization sum and store the active lanes.
Complete kernel64 lines

Includes tile loads, the streaming loop, sequence and D-tail masks, and the final store.

def fa1ForwardKernelStridedBoundaryD
    (qReg kReg vReg outReg : RegionName)
    (M Bd Bk numKVBlocks S_q S_k D : Nat)
    -- Q strides (axes [B, H, S_q, D]):
    (stride_qb stride_qh stride_qs stride_qd : Nat)
    -- K strides (axes [B, H, S_k, D]):
    (stride_kb stride_kh stride_kn stride_kd : Nat)
    -- V strides (axes [B, H, S_k, D]):
    (stride_vb stride_vh stride_vn stride_vd : Nat)
    -- Output strides (axes [B, H, S_q, D]):
    (stride_ob stride_oh stride_om stride_od : Nat)
    (scale : ℝ) : ComputeKernel := triton {
  pid_qb := tl.program_id(0)
  pid_h  := tl.program_id(1)
  pid_b  := tl.program_id(2)

  q_base_off := pid_b * $(stride_qb) + pid_h * $(stride_qh)
  k_base_off := pid_b * $(stride_kb) + pid_h * $(stride_kh)
  v_base_off := pid_b * $(stride_vb) + pid_h * $(stride_vh)
  o_base_off := pid_b * $(stride_ob) + pid_h * $(stride_oh)

  offs_m := pid_qb * $(M) + tl.arange(0, $(M))
  offs_d := tl.arange(0, $(Bd))

  q_ptrs := q_base_off + offs_m[:, None] * $(stride_qs) + offs_d[None, :] * $(stride_qd)
  q_seq_mask := (offs_m[:, None] + offs_d[None, :] * $(0)) < $(S_q)
  q_d_mask   := (offs_m[:, None] * $(0) + offs_d[None, :]) < $(D)
  q_mask     := tl.logical_and(q_seq_mask, q_d_mask)
  q          := tl.load($(qReg) + q_ptrs, mask=q_mask, other=0)

  m_i    := tl.full([$(M)], -inf)
  l_i    := tl.zeros([$(M)])
  o_acc  := tl.zeros([$(M), $(Bd)])

  tl.for n in $(numKVBlocks) {
    offs_n  := n * $(Bk) + tl.arange(0, $(Bk))
    k_ptrs  := k_base_off + offs_n[:, None] * $(stride_kn) + offs_d[None, :] * $(stride_kd)
    v_ptrs  := v_base_off + offs_n[:, None] * $(stride_vn) + offs_d[None, :] * $(stride_vd)
    kv_seq_mask := (offs_n[:, None] + offs_d[None, :] * $(0)) < $(S_k)
    kv_d_mask   := (offs_n[:, None] * $(0) + offs_d[None, :]) < $(D)
    kv_mask     := tl.logical_and(kv_seq_mask, kv_d_mask)
    k       := tl.load($(kReg) + k_ptrs, mask=kv_mask, other=0)
    v       := tl.load($(vReg) + v_ptrs, mask=kv_mask, other=0)

    scores_raw := tl.dot(q, tl.trans(k)) * $(scale)
    score_mask := (offs_m[:, None] * $(0) + offs_n[None, :]) < $(S_k)
    scores     := tl.where(score_mask, scores_raw, -inf)
    m_block    := tl.max(scores, axis = 1)
    m_new      := tl.max(m_i, m_block)
    alpha      := tl.exp(m_i - m_new)
    p          := tl.exp(scores - m_new[:, None])
    l_new      := alpha * l_i + tl.sum(p, axis = 1)
    o_acc      := alpha[:, None] * o_acc + tl.dot(p, v)
    m_i        := m_new
    l_i        := l_new
  }

  out    := o_acc / l_i[:, None]
  o_ptrs := o_base_off + offs_m[:, None] * $(stride_om) + offs_d[None, :] * $(stride_od)
  o_seq_mask := (offs_m[:, None] + offs_d[None, :] * $(0)) < $(S_q)
  o_d_mask   := (offs_m[:, None] * $(0) + offs_d[None, :]) < $(D)
  o_mask     := tl.logical_and(o_seq_mask, o_d_mask)
  tl.store($(outReg) + o_ptrs, out, mask=o_mask)
}
Tensor-view binding

The specification’s views.boundaryKernelD selects this kernel with the tensor layout’s strides.

def boundaryKernelD (views : FA1Views4D B H S_q S_k D)
    (M Bd Bk numKVBlocks : Nat) (scale : ℝ) : ComputeKernel :=
  views.layout.boundaryKernelD views.qReg views.kReg views.vReg views.outReg
    M Bd Bk numKVBlocks scale
def boundaryKernelD (layout : FA1Layout4D B H S_q S_k D)
    (qReg kReg vReg outReg : RegionName)
    (M Bd Bk numKVBlocks : Nat) (scale : ℝ) : ComputeKernel :=
  fa1ForwardKernelStridedBoundaryD qReg kReg vReg outReg
    M Bd Bk numKVBlocks S_q S_k D
    layout.qB layout.qH layout.qS layout.qD
    layout.kB layout.kH layout.kS layout.kD
    layout.vB layout.vH layout.vS layout.vD
    layout.oB layout.oH layout.oS layout.oD scale

Specification

Direct attention reference

O = softmax(QKᵀ · scale) V

noncomputable def fa1NaiveReference4D {B H S_q S_k D : Nat}
    (Q : TileIndex [B, H, S_q, D] → ℝ)
    (K V : TileIndex [B, H, S_k, D] → ℝ)
    (scale : ℝ) : TileIndex [B, H, S_q, D] → ℝ :=
  attentionReal4D Q K V scale
Computation
Each active output lane matches the direct attention reference for its batch, head, query, and feature indices.
Memory
The result is checked at its declared output address, determined by the tensor view and the program indices.
Scope
Real arithmetic, under the loaded-input and tile-size conditions stated in the theorem.
Full theorem statement
specification fa1_boundaryD_refines_naive_reference_views
    {B H S_q S_k D Bd Bk numKVBlocks M : Nat}
    (hBk : 0 < Bk) (hSk : 0 < S_k) (hSkLe : S_k ≤ Bk * numKVBlocks)
    (hDLe : D ≤ Bd)
    (views : FA1Views4D B H S_q S_k D)
    (Q4D : TileIndex [B, H, S_q, D] → ℝ)
    (K4D V4D : TileIndex [B, H, S_k, D] → ℝ)
    (scale : ℝ) (s : BlockState)
    (hPidB : s.pids 2 < B) (hPidH : s.pids 1 < H)
    (hQ4D : TensorView.loaded s views.qView Q4D)
    (hK4D : TensorView.loaded s views.kView K4D)
    (hV4D : TensorView.loaded s views.vView V4D) :
    ComputeCorrect.Realizes_without_Rounding
      (kernel := views.boundaryKernelD M Bd Bk numKVBlocks scale)
      (initialState := s)
      (write := fun idx : { idx : TileIndex [M, Bd] //
          s.pids 0 * M + idx.1.val < S_q ∧ idx.2.1.val < D } =>
        some (views.outReg, (views.outBlockOffsetD s M Bd) idx.1))
      (expected := fun idx =>
        fa1NaiveReference4D Q4D K4D V4D scale
          (⟨s.pids 2, hPidB⟩, ⟨s.pids 1, hPidH⟩,
           ⟨s.pids 0 * M + idx.1.1.val, idx.2.1⟩,
           ⟨idx.1.2.1.val, idx.2.2⟩, PUnit.unit))
Reference for each batch and head
noncomputable def attentionReal4D {B H S_q S_k D : Nat}
    (Q : TileIndex [B, H, S_q, D] → ℝ)
    (K V : TileIndex [B, H, S_k, D] → ℝ)
    (scale : ℝ) : TileIndex [B, H, S_q, D] → ℝ :=
  fun (b, h, i, d, _) =>
    attentionReal (sliceBH Q b h) (sliceBH K b h) (sliceBH V b h)
      scale (i, d, PUnit.unit)

Proof

From kernel execution to the specification

The proof uses the forward-correctness theorem to relate each observed output value to direct attention.

Proof body of fa1_boundaryD_refines_naive_reference_views.

by
  apply ComputeKernel.computeCorrect_of_toAlgKernel rfl
  intro s0 s' hExec hs0
  subst s0
  intro idx
  have hview := fa1_boundaryD_refines_naive_reference_exec_views hBk hSk hSkLe hDLe
    views Q4D K4D V4D scale s hPidB hPidH hQ4D hK4D hV4D idx.1 idx.2.1 idx.2.2
  rw [hExec] at hview
  simpa [observeTileAt] using hview

Supporting proofs

Execution agrees with the direct reference
theorem fa1_boundaryD_refines_naive_reference_exec_views
    {B H S_q S_k D Bd Bk numKVBlocks M : Nat}
    (hBk : 0 < Bk) (hSk : 0 < S_k) (hSkLe : S_k ≤ Bk * numKVBlocks)
    (hDLe : D ≤ Bd)
    (views : FA1Views4D B H S_q S_k D)
    (Q4D : TileIndex [B, H, S_q, D] → ℝ)
    (K4D V4D : TileIndex [B, H, S_k, D] → ℝ)
    (scale : ℝ) (s : BlockState)
    (hPidB : s.pids 2 < B) (hPidH : s.pids 1 < H)
    (hQ4D : TensorView.loaded s views.qView Q4D)
    (hK4D : TensorView.loaded s views.kView K4D)
    (hV4D : TensorView.loaded s views.vView V4D) :
    ∀ idx : TileIndex [M, Bd],
      ∀ hLt : s.pids 0 * M + idx.1.val < S_q,
      ∀ hDIdx : idx.2.1.val < D,
      observeTileAt
          (exec (views.boundaryKernelD M Bd Bk numKVBlocks scale) s)
          views.outReg (views.outBlockOffsetD s M Bd) idx
        = some (fa1NaiveReference4D Q4D K4D V4D scale
            (⟨s.pids 2, hPidB⟩, ⟨s.pids 1, hPidH⟩,
             ⟨s.pids 0 * M + idx.1.val, hLt⟩,
             ⟨idx.2.1.val, hDIdx⟩, PUnit.unit)) := by
  intro idx hLt hDIdx
  simpa [fa1NaiveReference4D]
    using fa1_forward_correct_4D_boundaryD_views hBk hSk hSkLe hDLe
      views Q4D K4D V4D scale s hPidB hPidH hQ4D hK4D hV4D idx hLt hDIdx
Forward correctness over tensor views
theorem fa1_forward_correct_4D_boundaryD_views
    {B H S_q S_k D Bd Bk numKVBlocks M : Nat}
    (hBk : 0 < Bk) (hSk : 0 < S_k) (hSkLe : S_k ≤ Bk * numKVBlocks)
    (hDLe : D ≤ Bd)
    (views : FA1Views4D B H S_q S_k D)
    (Q4D : TileIndex [B, H, S_q, D] → ℝ)
    (K4D V4D : TileIndex [B, H, S_k, D] → ℝ)
    (scale : ℝ) (s : BlockState)
    (hPidB : s.pids 2 < B) (hPidH : s.pids 1 < H)
    (hQ4D : TensorView.loaded s views.qView Q4D)
    (hK4D : TensorView.loaded s views.kView K4D)
    (hV4D : TensorView.loaded s views.vView V4D) :
    ∀ idx : TileIndex [M, Bd],
      ∀ hLt : s.pids 0 * M + idx.1.val < S_q,
      ∀ hDIdx : idx.2.1.val < D,
      observeTileAt
          (exec (views.boundaryKernelD M Bd Bk numKVBlocks scale) s)
          views.outReg (views.outBlockOffsetD s M Bd) idx
        = some (attentionReal4D Q4D K4D V4D scale
            (⟨s.pids 2, hPidB⟩, ⟨s.pids 1, hPidH⟩,
             ⟨s.pids 0 * M + idx.1.val, hLt⟩,
             ⟨idx.2.1.val, hDIdx⟩, PUnit.unit)) := by
  intro idx hLt hDIdx
  simpa [FA1Views4D.boundaryKernelD, FA1Views4D.outBlockOffsetD,
         FA1Views4D.qView, FA1Views4D.kView, FA1Views4D.vView]
    using fa1_forward_correct_4D_boundaryD_layout hBk hSk hSkLe hDLe views.layout
      views.qReg views.kReg views.vReg views.outReg
      Q4D K4D V4D scale s hPidB hPidH
      hQ4D hK4D hV4D idx hLt hDIdx
Complete refinement source172 lines

NaiveRefinement.lean includes the non-causal and causal reference contracts and their proofs.

/-
VeriTile.Examples.FlashAttention1.NaiveRefinement

Step 4 of issue #39: make the naive/reference refinement surface explicit.
-/

import VeriTile.Triton.Float
import VeriTile.Examples.FlashAttention1.NaiveKernel
import VeriTile.Meta.Specification

namespace VeriTile.Examples

open VeriTile.Triton

/-! ## Naive FA reference

`attentionReal4D` is the single-block-output, non-online softmax reference:
`softmax(QKᵀ * scale) · V` for each `(batch, head)` slice. The verified
single-pass naive boundary kernels live in `NaiveKernel.lean`; this file
keeps the reference-level refinement aliases.
-/

noncomputable def fa1NaiveReference4D {B H S_q S_k D : Nat}
    (Q : TileIndex [B, H, S_q, D] → ℝ)
    (K V : TileIndex [B, H, S_k, D] → ℝ)
    (scale : ℝ) : TileIndex [B, H, S_q, D] → ℝ :=
  attentionReal4D Q K V scale

noncomputable def fa1NaiveCausalReference4D {B H S_q S_k D : Nat}
    (Q : TileIndex [B, H, S_q, D] → ℝ)
    (K V : TileIndex [B, H, S_k, D] → ℝ)
    (scale : ℝ) : TileIndex [B, H, S_q, D] → ℝ :=
  attentionReal4DCausal Q K V scale

theorem fa1_naive_reference_eq_attentionReal4D {B H S_q S_k D : Nat}
    (Q : TileIndex [B, H, S_q, D] → ℝ)
    (K V : TileIndex [B, H, S_k, D] → ℝ)
    (scale : ℝ) :
    fa1NaiveReference4D Q K V scale = attentionReal4D Q K V scale := rfl

theorem fa1_naive_causal_reference_eq_attentionReal4DCausal {B H S_q S_k D : Nat}
    (Q : TileIndex [B, H, S_q, D] → ℝ)
    (K V : TileIndex [B, H, S_k, D] → ℝ)
    (scale : ℝ) :
    fa1NaiveCausalReference4D Q K V scale =
      attentionReal4DCausal Q K V scale := rfl

/-- FA-1 boundary+D-tail refines the naive direct FA reference. -/
theorem fa1_boundaryD_refines_naive_reference_exec_views
    {B H S_q S_k D Bd Bk numKVBlocks M : Nat}
    (hBk : 0 < Bk) (hSk : 0 < S_k) (hSkLe : S_k ≤ Bk * numKVBlocks)
    (hDLe : D ≤ Bd)
    (views : FA1Views4D B H S_q S_k D)
    (Q4D : TileIndex [B, H, S_q, D] → ℝ)
    (K4D V4D : TileIndex [B, H, S_k, D] → ℝ)
    (scale : ℝ) (s : BlockState)
    (hPidB : s.pids 2 < B) (hPidH : s.pids 1 < H)
    (hQ4D : TensorView.loaded s views.qView Q4D)
    (hK4D : TensorView.loaded s views.kView K4D)
    (hV4D : TensorView.loaded s views.vView V4D) :
    ∀ idx : TileIndex [M, Bd],
      ∀ hLt : s.pids 0 * M + idx.1.val < S_q,
      ∀ hDIdx : idx.2.1.val < D,
      observeTileAt
          (exec (views.boundaryKernelD M Bd Bk numKVBlocks scale) s)
          views.outReg (views.outBlockOffsetD s M Bd) idx
        = some (fa1NaiveReference4D Q4D K4D V4D scale
            (⟨s.pids 2, hPidB⟩, ⟨s.pids 1, hPidH⟩,
             ⟨s.pids 0 * M + idx.1.val, hLt⟩,
             ⟨idx.2.1.val, hDIdx⟩, PUnit.unit)) := by
  intro idx hLt hDIdx
  simpa [fa1NaiveReference4D]
    using fa1_forward_correct_4D_boundaryD_views hBk hSk hSkLe hDLe
      views Q4D K4D V4D scale s hPidB hPidH hQ4D hK4D hV4D idx hLt hDIdx

/-- Compute-facing FA-1 boundary+D-tail correctness against the naive direct
FA reference. -/
specification fa1_boundaryD_refines_naive_reference_views
    {B H S_q S_k D Bd Bk numKVBlocks M : Nat}
    (hBk : 0 < Bk) (hSk : 0 < S_k) (hSkLe : S_k ≤ Bk * numKVBlocks)
    (hDLe : D ≤ Bd)
    (views : FA1Views4D B H S_q S_k D)
    (Q4D : TileIndex [B, H, S_q, D] → ℝ)
    (K4D V4D : TileIndex [B, H, S_k, D] → ℝ)
    (scale : ℝ) (s : BlockState)
    (hPidB : s.pids 2 < B) (hPidH : s.pids 1 < H)
    (hQ4D : TensorView.loaded s views.qView Q4D)
    (hK4D : TensorView.loaded s views.kView K4D)
    (hV4D : TensorView.loaded s views.vView V4D) :
    ComputeCorrect.Realizes_without_Rounding
      (kernel := views.boundaryKernelD M Bd Bk numKVBlocks scale)
      (initialState := s)
      (write := fun idx : { idx : TileIndex [M, Bd] //
          s.pids 0 * M + idx.1.val < S_q ∧ idx.2.1.val < D } =>
        some (views.outReg, (views.outBlockOffsetD s M Bd) idx.1))
      (expected := fun idx =>
        fa1NaiveReference4D Q4D K4D V4D scale
          (⟨s.pids 2, hPidB⟩, ⟨s.pids 1, hPidH⟩,
           ⟨s.pids 0 * M + idx.1.1.val, idx.2.1⟩,
           ⟨idx.1.2.1.val, idx.2.2⟩, PUnit.unit)) := by
  apply ComputeKernel.computeCorrect_of_toAlgKernel rfl
  intro s0 s' hExec hs0
  subst s0
  intro idx
  have hview := fa1_boundaryD_refines_naive_reference_exec_views hBk hSk hSkLe hDLe
    views Q4D K4D V4D scale s hPidB hPidH hQ4D hK4D hV4D idx.1 idx.2.1 idx.2.2
  rw [hExec] at hview
  simpa [observeTileAt] using hview

/-- Causal FA-1 boundary+D-tail refines the naive direct causal FA reference. -/
theorem fa1_causal_boundaryD_refines_naive_reference_exec_views
    {B H S_q S_k D Bd Bk numKVBlocks M : Nat}
    (hBk : 0 < Bk) (hSk : 0 < S_k) (hSkLe : S_k ≤ Bk * numKVBlocks)
    (hDLe : D ≤ Bd)
    (views : FA1Views4D B H S_q S_k D)
    (Q4D : TileIndex [B, H, S_q, D] → ℝ)
    (K4D V4D : TileIndex [B, H, S_k, D] → ℝ)
    (scale : ℝ) (s : BlockState)
    (hPidB : s.pids 2 < B) (hPidH : s.pids 1 < H)
    (hQ4D : TensorView.loaded s views.qView Q4D)
    (hK4D : TensorView.loaded s views.kView K4D)
    (hV4D : TensorView.loaded s views.vView V4D) :
    ∀ idx : TileIndex [M, Bd],
      ∀ hLt : s.pids 0 * M + idx.1.val < S_q,
      ∀ hDIdx : idx.2.1.val < D,
      observeTileAt
          (exec (views.causalBoundaryKernelD M Bd Bk numKVBlocks scale) s)
          views.outReg (views.outBlockOffsetD s M Bd) idx
        = some (fa1NaiveCausalReference4D Q4D K4D V4D scale
            (⟨s.pids 2, hPidB⟩, ⟨s.pids 1, hPidH⟩,
             ⟨s.pids 0 * M + idx.1.val, hLt⟩,
             ⟨idx.2.1.val, hDIdx⟩, PUnit.unit)) := by
  intro idx hLt hDIdx
  simpa [fa1NaiveCausalReference4D]
    using fa1_forward_correct_4D_causal_boundaryD_views hBk hSk hSkLe hDLe
      views Q4D K4D V4D scale s hPidB hPidH hQ4D hK4D hV4D idx hLt hDIdx

/-- Compute-facing causal FA-1 boundary+D-tail correctness against the naive
direct causal FA reference. -/
specification fa1_causal_boundaryD_refines_naive_reference_views
    {B H S_q S_k D Bd Bk numKVBlocks M : Nat}
    (hBk : 0 < Bk) (hSk : 0 < S_k) (hSkLe : S_k ≤ Bk * numKVBlocks)
    (hDLe : D ≤ Bd)
    (views : FA1Views4D B H S_q S_k D)
    (Q4D : TileIndex [B, H, S_q, D] → ℝ)
    (K4D V4D : TileIndex [B, H, S_k, D] → ℝ)
    (scale : ℝ) (s : BlockState)
    (hPidB : s.pids 2 < B) (hPidH : s.pids 1 < H)
    (hQ4D : TensorView.loaded s views.qView Q4D)
    (hK4D : TensorView.loaded s views.kView K4D)
    (hV4D : TensorView.loaded s views.vView V4D) :
    ComputeCorrect.Realizes_without_Rounding
      (kernel := views.causalBoundaryKernelD M Bd Bk numKVBlocks scale)
      (initialState := s)
      (write := fun idx : { idx : TileIndex [M, Bd] //
          s.pids 0 * M + idx.1.val < S_q ∧ idx.2.1.val < D } =>
        some (views.outReg, (views.outBlockOffsetD s M Bd) idx.1))
      (expected := fun idx =>
        fa1NaiveCausalReference4D Q4D K4D V4D scale
          (⟨s.pids 2, hPidB⟩, ⟨s.pids 1, hPidH⟩,
           ⟨s.pids 0 * M + idx.1.1.val, idx.2.1⟩,
           ⟨idx.1.2.1.val, idx.2.2⟩, PUnit.unit)) := by
  apply ComputeKernel.computeCorrect_of_toAlgKernel rfl
  intro s0 s' hExec hs0
  subst s0
  intro idx
  have hview := fa1_causal_boundaryD_refines_naive_reference_exec_views hBk hSk hSkLe hDLe
    views Q4D K4D V4D scale s hPidB hPidH hQ4D hK4D hV4D idx.1 idx.2.1 idx.2.2
  rw [hExec] at hview
  simpa [observeTileAt] using hview

end VeriTile.Examples
Browse all examples ↗

Getting started

Build the library, then check the complete VectorAdd example.

Terminal
git clone https://github.com/Lizn-zn/VeriTile.git
cd VeriTile
lake build
lake env lean bench/examples/VectorAdd.lean

With Git and elan installed; elan uses the pinned Lean toolchain. Setup guide →

Expected: exit 0, with the axiom and statement-surface checks passing.