vyre-lower

Substrate-neutral lowering, optimization, analysis, and verification for vyre's KernelDescriptor IR.

This crate sits between vyre's frontend (which produces a high-level vyre::Program) and the substrate-specific emitters (vyre-emit-naga, vyre-emit-ptx, vyre-emit-spirv). It owns:

The KernelDescriptor IR: a flat, SSA-shaped, structured-control- flow program that every emitter consumes verbatim.
12 rewrite passes that simplify the IR before lowering.
11 analyses that report on the IR (coalescing, bank conflicts, shared-mem promotion candidates, def-use chains, etc.).
A structural verifier (verify) that catches dangling refs, duplicate result-ids, and out-of-range pool/child-body indices.
Performance instrumentation (OptimizationStats).

If you're emitting GPU code, you want vyre_emit_*::emit_optimized: those wrappers call vyre_lower::rewrites::run_all for you.

Quick start

use vyre_lower::{
    BindingLayout, BindingSlot, BindingVisibility, Dispatch, KernelBody,
    KernelDescriptor, KernelOp, KernelOpKind, LiteralValue, MemoryClass,
    rewrites::{run_all, run_all_with_stats},
    verify,
};
use vyre_foundation::ir::DataType;

let desc = KernelDescriptor {
    id: "store_seven".into(),
    bindings: BindingLayout {
        slots: vec![BindingSlot {
            slot: 0,
            element_type: DataType::U32,
            element_count: None,
            memory_class: MemoryClass::Global,
            visibility: BindingVisibility::ReadWrite,
            name: "out".into(),
        }],
    },
    dispatch: Dispatch::new(64, 1, 1),
    body: KernelBody {
        ops: vec![
            KernelOp { kind: KernelOpKind::Literal, operands: vec![0], result: Some(0) },
            KernelOp { kind: KernelOpKind::Literal, operands: vec![1], result: Some(1) },
            KernelOp {
                kind: KernelOpKind::StoreGlobal,
                operands: vec![0, 0, 1],
                result: None,
            },
        ],
        child_bodies: vec![],
        literals: vec![LiteralValue::U32(0), LiteralValue::U32(7)],
    },
};

// Run the optimization pipeline.
let (optimized, stats) = run_all_with_stats(&desc);
println!(
    "{} ops -> {} ops in {} iterations",
    stats.ops_before, stats.ops_after, stats.iterations,
);

// Sanity: the optimized form is structurally valid.
assert!(verify(&optimized).is_ok());

The IR

A KernelDescriptor is:

id: String: diagnostic.
bindings: BindingLayout { slots: Vec<BindingSlot> }: buffers bound at the kernel boundary, looked up by BindingSlot.slot field.
dispatch: Dispatch { workgroup_size }: thread-group geometry.
body: KernelBody: the program.

A KernelBody is:

ops: Vec<KernelOp>: flat op stream, walked linearly.
child_bodies: Vec<KernelBody>: referenced by If/ForLoop/ Block/Region ops via a child-body index in their operands.
literals: Vec<LiteralValue>: pool, referenced by Literal ops.

A KernelOp is { kind: KernelOpKind, operands: Vec<u32>, result: Option<u32> }. Operands are typed by position per KernelOpKind: some positions are SSA result-id refs, some are literal-pool indices, some are binding slot ids, some are child-body indices.

Per-body id space. Each KernelBody has its own SSA id space. Result-ids in a child body do NOT exist in the parent body's id space. Rewrites that move ops across bodies must respect this: see the LICM module for the consequence of getting it wrong.

Rewrite pipeline

run_all applies these passes in order, then iterates to fixed point (up to 4 iterations):

strength_reduce: Mul/Div/Mod by power-of-2 → shift/and.
const_fold: folds compile-time-constant arithmetic. Coverage:
- BinOp(Lit, Lit) → single Lit (full BinOp×Type matrix: U32/I32/F32/Bool × all 22 BinOp variants including comparisons, bitwise, wrapping arithmetic).
- UnOp(Lit) → single Lit (10 unary ops: BitNot, Negate, LogicalNot, Popcount, Clz, Ctz, ReverseBits, Abs, Floor, Ceil, Round, Trunc, Sqrt, Cos, Sin).
- Cast(Lit) → typed Lit (int↔int, int↔float, bool→int, same-type; float→int only when finite + in range).
- Fma(Lit, Lit, Lit) → single Lit (F32 only, finite-result).
identity_elim: Add(x, 0), Mul(x, 1), etc. → x; Mul(x, 0), BitAnd(x, 0) → 0; Select(Lit_bool, then, else) → then or else (bool-cond folding).
branch_collapse: If(Lit_bool, ...) → selected arm inlined.
loop_unroll: small constant-bound loops (≤ 4 iterations).
licm: currently a no-op (see module docs).
load_forwarding: store-to-load and load-to-load forwarding with per-slot aliasing rules.
dce: drops result-producing ops with no users.
dead_store: drops stores whose value is overwritten before any observation.
dce (again): cleans up what dead_store orphaned.
canonicalize: sorts commutative-op operands so CSE catches Add(a, b) == Add(b, a).
cse: merges structurally-equivalent ops.
drop_unused_bindings: strips binding slots no surviving op references.
drop_unused_literals: strips pool entries no Literal op references (const_fold + identity_elim leave plenty of orphans).
drop_unused_child_bodies: strips child bodies orphaned by branch_collapse / loop_unroll inlining.

Each pass is total (no Result, returns input on no-op), preserves semantic equivalence, and is individually idempotent. The wrapping fixed-point loop catches inter-pass dependencies (e.g., CSE merging two index ops exposes a dead_store opportunity that dead_store missed in the first pass).

Analyses

11 substrate-neutral analyses in vyre_lower::analyses:

coalesce: memory-access coalescence per warp/workgroup.
shared_mem_promote: global → shared-memory tile candidates.
bank_conflict: shared-memory bank conflict detection.
vec_pack: adjacent-load vectorization candidates (companion to vyre_emit_naga::patterns::vec_pack).
workgroup_uniform: values uniform across a workgroup.
texture_promote: read-mostly buffer → texture candidates.
layout_aos_to_soa: AoS-to-SoA layout transform candidates.
const_buffer_promote: uniform-buffer promotion candidates.
dead_op: result-producing ops with no users (a less efficient cousin of def_use::dead_by_no_use).
common_subexpr: equivalence groups for CSE.
def_use: full def-use chains with per-body UseSites.

Each returns a serializable report. Run audit::audit(desc) for a unified PerfAuditReport with prioritized recommendations, or audit::audit_optimized(desc) to audit the post-run_all form (answers "what perf issues remain after the standard pipeline?"). The same audit + audit_optimized pair is mirrored in vyre_emit_naga::patterns, vyre_emit_ptx::patterns, and vyre_emit_spirv::patterns for substrate-specific concerns.

Verifier

verify(desc) -> Result<(), Vec<VerifyError>> checks:

Result-id uniqueness within each body.
No dangling result-id refs.
Literal-pool indices in range.
Child-body indices in range.
Literal ops have ≥1 operand.
Per-kind minimum operand counts.

Errors are collected (not short-circuited) so a single call surfaces every violation. Both vyre_emit_*::emit_optimized functions debug_assert!(verify(optimized).is_ok()): production builds skip the check; debug/test builds catch any rewrite bug at the boundary with a clean panic message.

The 1000-descriptor fuzz harness at tests/rewrite_soundness_fuzz.rs is the regression gate: every shape in the corpus must produce a descriptor that verifies after run_all.

License

MIT OR Apache-2.0.

vyre-lower 0.6.1