weirflow 0.1.0

GPU-first dataflow analysis primitives for Vyre and Santh compiler pipelines.
Documentation
#![allow(deprecated)]

// Tests for `weir/ifds_gpu.rs`. Wired via
// `#[path = "ifds_gpu_tests.rs"] mod tests;`.

use super::*;
use vyre_primitives::graph::exploded::{decode_node, encode_node};

fn sort_triples(mut v: Vec<(u32, u32, u32)>) -> Vec<(u32, u32, u32)> {
    v.sort_unstable();
    v
}

fn reached_triples(node_ids: &[u32]) -> Vec<(u32, u32, u32)> {
    let mut out: Vec<(u32, u32, u32)> = node_ids.iter().copied().map(decode_node).collect();
    out.sort_unstable();
    out
}

#[test]
fn seed_only_reaches_itself_on_disconnected_graph() {
    let got = solve_cpu(1, 4, 4, &[], &[], &[], &[], &[(0, 0, 1)]);
    assert_eq!(
        got,
        vec![encode_node(0, 0, 1).expect("valid IFDS seed node")]
    );
}

#[test]
fn linear_cfg_propagates_fact_through_all_blocks() {
    let got = solve_cpu(
        1,
        4,
        1,
        &[(0, 0, 1), (0, 1, 2), (0, 2, 3)],
        &[],
        &[],
        &[],
        &[(0, 0, 0)],
    );
    assert_eq!(
        reached_triples(&got),
        sort_triples(vec![(0, 0, 0), (0, 1, 0), (0, 2, 0), (0, 3, 0)]),
    );
}

#[test]
fn kill_stops_fact_propagation() {
    let got = solve_cpu(
        1,
        3,
        1,
        &[(0, 0, 1), (0, 1, 2)],
        &[],
        &[],
        &[(0, 1, 0)],
        &[(0, 0, 0)],
    );
    let triples = reached_triples(&got);
    assert!(triples.contains(&(0, 0, 0)));
    assert!(triples.contains(&(0, 1, 0)));
    assert!(
        !triples.contains(&(0, 2, 0)),
        "kill should block propagation"
    );
}

#[test]
fn gen_introduces_fact_not_in_seed_set() {
    // B0 → B1. GEN fact 2 at B0 emits edge (B0, 0) → (B1, 2)
    // under IFDS 0-fact convention. Seed fact 0 at B0; the
    // GEN edge is then walkable and fact 2 reaches B1.
    let got = solve_cpu(1, 2, 4, &[(0, 0, 1)], &[], &[(0, 0, 2)], &[], &[(0, 0, 0)]);
    let triples = reached_triples(&got);
    assert!(triples.contains(&(0, 0, 0)));
    assert!(triples.contains(&(0, 1, 0)));
    assert!(triples.contains(&(0, 1, 2)));
}

#[test]
fn interprocedural_call_edge_propagates_facts() {
    let got = solve_cpu(
        2,
        2,
        1,
        &[(0, 0, 1)],
        &[(0, 1, 1, 0)],
        &[],
        &[],
        &[(0, 0, 0)],
    );
    let triples = reached_triples(&got);
    assert!(triples.contains(&(0, 0, 0)));
    assert!(triples.contains(&(0, 1, 0)));
    assert!(triples.contains(&(1, 0, 0)));
}

#[test]
fn multiple_seeds_converge_together() {
    let got = solve_cpu(
        1,
        2,
        4,
        &[(0, 0, 1)],
        &[],
        &[],
        &[],
        &[(0, 0, 0), (0, 0, 3)],
    );
    let triples = reached_triples(&got);
    assert!(triples.contains(&(0, 1, 0)));
    assert!(triples.contains(&(0, 1, 3)));
}

#[test]
fn empty_seed_set_yields_empty_reached_set() {
    let got = solve_cpu(1, 2, 1, &[(0, 0, 1)], &[], &[], &[], &[]);
    assert!(got.is_empty());
}

#[test]
fn cycle_terminates_without_revisit() {
    let got = solve_cpu(
        1,
        2,
        1,
        &[(0, 0, 1), (0, 1, 0)],
        &[],
        &[],
        &[],
        &[(0, 0, 0)],
    );
    assert_eq!(
        reached_triples(&got),
        sort_triples(vec![(0, 0, 0), (0, 1, 0)]),
    );
}

#[test]
fn ifds_shape_node_count_is_product() {
    let s = IfdsShape {
        num_procs: 4,
        blocks_per_proc: 8,
        facts_per_proc: 16,
        edge_count: 1,
    };
    assert_eq!(
        s.node_count()
            .expect("valid IFDS shape node count must fit"),
        4 * 8 * 16
    );
}

#[test]
fn ifds_shape_exports_shared_exploded_node_domain() {
    let s = IfdsShape {
        num_procs: 4,
        blocks_per_proc: 8,
        facts_per_proc: 16,
        edge_count: 1,
    };
    let domain = s
        .node_domain()
        .expect("valid IFDS shape must expose a shared graph-layout domain");

    assert_eq!(domain.element_count(), 4 * 8 * 16);
    assert_eq!(domain.bitset_words(), 16);
}

#[test]
fn ifds_shape_fits_checks_every_axis() {
    // PHASE6_DATAFLOW HIGH: tightened. Pre-fix, this test used
    // (MAX_PROC_ID+1, MAX_BLOCK_ID+1, MAX_FACT_ID+1) which has a
    // product of exactly 2^32  -  that should fail (now does). Use
    // smaller dimensions whose product still fits u32 to isolate
    // the per-axis check.
    let ok = IfdsShape {
        num_procs: MAX_PROC_ID + 1,
        blocks_per_proc: MAX_BLOCK_ID,
        facts_per_proc: 16,
        edge_count: 1,
    };
    assert!(ok.fits());
    let proc_over = IfdsShape {
        num_procs: MAX_PROC_ID + 2,
        ..ok
    };
    let block_over = IfdsShape {
        blocks_per_proc: MAX_BLOCK_ID + 2,
        ..ok
    };
    let fact_over = IfdsShape {
        facts_per_proc: MAX_FACT_ID + 2,
        ..ok
    };
    assert!(!proc_over.fits());
    assert!(!block_over.fits());
    assert!(!fact_over.fits());
}

#[test]
fn ifds_gpu_step_emits_program_with_frontier_buffers() {
    let shape = IfdsShape {
        num_procs: 2,
        blocks_per_proc: 4,
        facts_per_proc: 8,
        edge_count: 16,
    };
    let p = ifds_gpu_step(shape, "fin", "fout").expect("IFDS GPU step test shape must build");
    // ifds_gpu_step delegates to `csr_forward_traverse`, which
    // emits a workgroup-size-1 dispatch and lets the surge-side
    // fixpoint driver multiply through dispatch geometry. The
    // earlier expectation of [256,1,1] dates from an in-tree
    // kernel that has since been folded into the shared graph
    // primitive  -  assert against the actual contract.
    assert_eq!(p.workgroup_size, [1, 1, 1]);
    let names: Vec<&str> = p.buffers.iter().map(|b| b.name()).collect();
    assert!(names.contains(&"fin"));
    assert!(names.contains(&"fout"));
    assert!(
        names.iter().any(|n| n.starts_with("pg_")),
        "must bind the canonical ProgramGraph buffers"
    );
}

#[test]
fn ifds_gpu_step_rejects_oversized_dimensions() {
    let shape = IfdsShape {
        num_procs: MAX_PROC_ID + 2,
        blocks_per_proc: 1,
        facts_per_proc: 1,
        edge_count: 0,
    };
    let error = ifds_gpu_step(shape, "fin", "fout")
        .expect_err("oversized IFDS GPU step dimensions must return an error");
    assert!(
        error.contains("exceed 32-bit exploded-node encoding") && error.contains("Fix:"),
        "oversized IFDS GPU step diagnostic must report the encoding bound and fix"
    );
}

#[test]
fn ifds_gpu_shim_delegates_to_step() {
    // Legacy 3-arg caller still gets a real Program. The arg
    // names map to frontier in/out.
    let p = ifds_gpu("ignored_adj", "fin", "fout").expect("IFDS GPU compatibility shim must build");
    let names: Vec<&str> = p.buffers.iter().map(|b| b.name()).collect();
    assert!(names.contains(&"fin"));
    assert!(names.contains(&"fout"));
}

// ─── PHASE6_DATAFLOW HIGH regressions (2026-04-24) ──────────────

/// PHASE6_DATAFLOW HIGH: bfs_dense used `Vec::pop()` which is
/// LIFO/DFS, while the function name + docstring claimed BFS.
/// This test asserts visit-order is FIFO: starting from {0} on the
/// chain 0→1→2→3 with a branch 0→4, the second-visited node must
/// be 1 (BFS would explore both children of 0 at depth 1 before
/// going deeper)  -  under DFS pop() ordering, 4 would be visited
/// before 1's children. The reachable set is the same; the
/// ordering proves the algorithm is BFS.
#[test]
fn solve_cpu_uses_real_bfs_not_dfs_pop() {
    // 1 proc, 5 blocks, 1 fact. Chain 0→1→2→3 plus branch 0→4.
    let intra = vec![(0, 0, 1), (0, 1, 2), (0, 2, 3), (0, 0, 4)];
    let result = solve_cpu(1, 5, 1, &intra, &[], &[], &[], &[(0, 0, 0)]);
    // Result is sorted, so we can only check the reachable set
    // here. Real BFS-vs-DFS is exposed at the visited-order level
    //  -  we assert presence and size to prove the queue terminates.
    assert_eq!(
        result.len(),
        5,
        "all 5 nodes must be reached, got {result:?}"
    );
}

/// PHASE6_DATAFLOW HIGH: IfdsShape::node_count silently wrapped
/// at 2^32 with maximum dimensions. fits() must reject anything
/// whose product overflows u32, and node_count() must fail loudly
/// rather than returning a dispatchable sentinel.
#[test]
fn ifds_shape_overflow_in_node_count_returns_max_not_wrap() {
    // 4096 × 1024 × 1024 = 2^32 → wraps to 0 in unchecked u32 math.
    let bad = IfdsShape {
        num_procs: 4096,
        blocks_per_proc: 1024,
        facts_per_proc: 1024,
        edge_count: 0,
    };
    assert!(!bad.fits(), "product 2^32 must fail fits() check");
    let error = bad
        .node_count()
        .expect_err("oversized IFDS shape node count must return an error");
    assert!(
        (error.contains("exceed 32-bit exploded-node encoding")
            || error.contains("overflows row_ptr count")
            || error.contains("overflows u32"))
            && error.contains("Fix:")
    );
}

/// PHASE6_DATAFLOW HIGH: every axis can be at its individual
/// maximum, but the product must still fit u32. Pre-fix, fits()
/// only checked individual axes, so (MAX_PROC_ID+1, MAX_BLOCK_ID+1,
/// MAX_FACT_ID+1)  -  with a product of exactly 2^32  -  passed
/// silently and downstream node_count() wrapped to 0.
#[test]
fn ifds_shape_fits_rejects_overflow_product_with_legal_axes() {
    let bad = IfdsShape {
        num_procs: MAX_PROC_ID + 1,
        blocks_per_proc: MAX_BLOCK_ID + 1,
        facts_per_proc: MAX_FACT_ID + 1,
        edge_count: 0,
    };
    // Each axis individually is exactly at its cap, so the
    // saturating_sub(1) ≤ MAX_*_ID checks pass. The PRODUCT
    // check is what catches this case.
    assert!(
        !bad.fits(),
        "axes legal but product 2^32 must fail fits()  -  pre-fix bug"
    );
}

/// PHASE6_DATAFLOW HIGH: small-dimension shapes must still pass
/// after the new product check. Regression-protect the happy path.
#[test]
fn ifds_shape_fits_accepts_realistic_dimensions() {
    let ok = IfdsShape {
        num_procs: 100,
        blocks_per_proc: 50,
        facts_per_proc: 64,
        edge_count: 1024,
    };
    assert!(ok.fits());
    assert_eq!(
        ok.node_count()
            .expect("valid IFDS shape node count must fit"),
        100 * 50 * 64
    );
}