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):

1. strength_reduce: Mul/Div/Mod by power-of-2 → shift/and.
2. 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).
3. 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).
4. branch_collapse: If(Lit_bool, ...) → selected arm inlined.
5. loop_unroll: small constant-bound loops (≤ 4 iterations).
6. licm: currently a no-op (see module docs).
7. load_forwarding: store-to-load and load-to-load forwarding with per-slot aliasing rules.
8. dce: drops result-producing ops with no users.
9. dead_store: drops stores whose value is overwritten before any observation.
10. dce (again): cleans up what dead_store orphaned.
11. canonicalize: sorts commutative-op operands so CSE catches Add(a, b) == Add(b, a).
12. cse: merges structurally-equivalent ops.
13. drop_unused_bindings: strips binding slots no surviving op references.
14. drop_unused_literals: strips pool entries no Literal op references (const_fold + identity_elim leave plenty of orphans).
15. 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.

See also

• vyre-emit-naga / vyre-emit-ptx / vyre-emit-spirv: substrate emitters that consume KernelDescriptor. Each exposes emit and emit_optimized.
• vyre-foundation: IR primitives (BinOp, UnOp, DataType, MemoryOrdering) that KernelOpKind embeds.

License

MIT OR Apache-2.0.