weirflow 0.1.0

GPU-first dataflow analysis primitives for Vyre and Santh compiler pipelines.
Documentation
use super::*;
use proptest::prelude::*;

fn csr_from_edges(node_count: u32, edges: &[(u32, u32)]) -> (Vec<u32>, Vec<u32>, Vec<u32>) {
    let mut outgoing = vec![Vec::<u32>::new(); node_count as usize];
    for &(src, dst) in edges {
        outgoing[src as usize].push(dst);
    }
    let mut offsets = Vec::with_capacity(node_count as usize + 1);
    let mut targets = Vec::new();
    offsets.push(0);
    for targets_for_node in outgoing {
        targets.extend(targets_for_node);
        offsets.push(targets.len() as u32);
    }
    let masks = vec![1; targets.len()];
    (offsets, targets, masks)
}

fn seed_bitset(node_count: u32, nodes: &[u32]) -> Vec<u32> {
    let mut bits = vec![0; bitset_words(node_count) as usize];
    for &node in nodes {
        let node = node % node_count;
        bits[(node / 32) as usize] |= 1u32 << (node % 32);
    }
    bits
}

fn assert_subset(lhs: &[u32], rhs: &[u32]) {
    for (word_index, (left, right)) in lhs.iter().zip(rhs.iter()).enumerate() {
        assert_eq!(
                left & !right,
                0,
                "Fix: points-to closure must be monotone at word {word_index}; left={left:#034b}, right={right:#034b}"
            );
    }
}

fn bit_is_set(bits: &[u32], node: u32) -> bool {
    (bits[(node / 32) as usize] & (1u32 << (node % 32))) != 0
}

/// PHASE6_DATAFLOW HIGH regression: pre-fix, `andersen_points_to`
/// was a 2-arg entry point that hardcoded `ProgramGraphShape::new(1, 1)`,
/// silently producing a 1-node grid. Now the function REQUIRES an
/// explicit shape, so the emitted Program reflects the caller's
/// real constraint-graph dimensions.
#[test]
fn andersen_points_to_uses_caller_supplied_shape() {
    let shape = ProgramGraphShape::new(64, 128);
    let program = andersen_points_to(shape, "constraints_in", "pts_out");
    let buffer_count = |name: &str| {
        program
            .buffers
            .iter()
            .find(|buffer| buffer.name() == name)
            .map(|buffer| buffer.count)
            .unwrap_or_else(|| panic!("{name} buffer must be declared"))
    };

    assert_eq!(
        buffer_count("constraints_in"),
        bitset_words(64),
        "Fix: points-to input frontier must be sized to the caller's full node domain."
    );
    assert_eq!(
        buffer_count("pts_out"),
        bitset_words(64),
        "Fix: points-to output frontier must be sized to the caller's full node domain."
    );
    assert_eq!(
        buffer_count("pg_edge_targets"),
        128,
        "Fix: points-to CSR target buffer must preserve the caller's full edge domain."
    );
    assert_eq!(
        buffer_count("pg_edge_offsets"),
        65,
        "Fix: points-to CSR offset buffer must contain node_count + 1 offsets."
    );
}

/// PHASE6_DATAFLOW HIGH regression: the deprecated alias still
/// works for back-compat callers but emits the same Program as the
/// canonical entry. Drop in a future major version.
#[test]
#[allow(deprecated)]
fn deprecated_alias_emits_same_program_shape() {
    let shape = ProgramGraphShape::new(32, 64);
    let canonical = andersen_points_to(shape, "ci", "po");
    let alias = andersen_points_to_with_shape(shape, "ci", "po");
    assert_eq!(
        canonical.workgroup_size, alias.workgroup_size,
        "deprecated alias must delegate to canonical entry"
    );
    assert_eq!(canonical.buffers.len(), alias.buffers.len());
}

#[test]
fn andersen_subset_closure_soundness() {
    let node_count = 4;
    let edges = [(0, 1), (1, 2), (2, 3)];
    let (offsets, targets, masks) = csr_from_edges(node_count, &edges);
    let seed = seed_bitset(node_count, &[0]);

    let closure = cpu_subset_closure(node_count, &offsets, &targets, &masks, &seed);

    for node in 0..node_count {
        assert!(
                bit_is_set(&closure, node),
                "Fix: subset closure must propagate addr-of seed across transitive edge chain; missing node {node}"
            );
    }
}

#[test]
fn subset_closure_rejects_short_seed() {
    let node_count = 64;
    let edges = [(0, 63)];
    let (offsets, targets, masks) = csr_from_edges(node_count, &edges);

    let panic = std::panic::catch_unwind(|| {
        let _ = cpu_subset_closure(node_count, &offsets, &targets, &masks, &[1]);
    })
    .expect_err("short seed bitset must fail loudly");
    let message = panic
        .downcast_ref::<String>()
        .map(String::as_str)
        .or_else(|| panic.downcast_ref::<&str>().copied())
        .unwrap_or("<non-string panic>");
    assert!(
        message.contains("malformed seed bitset"),
        "unexpected diagnostic: {message}"
    );
}

#[test]
fn subset_closure_gpu_boundary_rejects_out_of_domain_seed_tail_bits() {
    let err = subset_closure_borrowed_via(
        &|_, _, _| unreachable!("validation must reject before dispatch"),
        3,
        &[0, 0, 0, 0],
        &[],
        &[],
        &[0b1000],
        1,
    )
    .expect_err("tail bit outside node_count must be rejected");
    assert!(
        err.contains("outside the declared domain"),
        "unexpected diagnostic: {err}"
    );
}

#[test]
fn subset_closure_into_reuses_output_storage_across_iterations() {
    let calls = std::cell::Cell::new(0_u32);
    let saw_reused_slot = std::cell::Cell::new(false);
    let dispatch =
        |_: &Program, inputs: &[&[u8]], _: Option<[u32; 3]>, outputs: &mut Vec<Vec<u8>>| {
            if inputs.len() < 6 {
                return crate::fixed_point_closure::fallback_bitset_equal_dispatch(inputs, outputs);
            }
            let call = calls.get();
            if call > 0 && outputs.len() == 1 {
                saw_reused_slot.set(true);
            }
            calls.set(call + 1);
            let frontier = u32::from_le_bytes(inputs[5][..4].try_into().unwrap());
            let next = if frontier & 1 != 0 {
                frontier | 0b10
            } else {
                frontier
            };
            if outputs.is_empty() {
                outputs.push(Vec::new());
            }
            outputs[0].clear();
            outputs[0].extend_from_slice(&next.to_le_bytes());
            Ok(())
        };

    let result = subset_closure_borrowed_into_via(
        &dispatch,
        2,
        &[0, 1, 1],
        &[1],
        &[vyre_primitives::predicate::edge_kind::ASSIGNMENT],
        &[0b01],
        4,
    )
    .expect("points-to subset closure into dispatch must converge");

    assert_eq!(result, vec![0b11]);
    assert_eq!(calls.get(), 2);
    assert!(saw_reused_slot.get());
}

proptest! {
    #![proptest_config(ProptestConfig::with_cases(10_000))]
    #[test]
    fn andersen_monotonicity_proptest(
        node_count in 1u32..96,
        raw_edges in prop::collection::vec((0u32..256, 0u32..256), 0..256),
        raw_seed in prop::collection::vec(0u32..256, 0..96),
        raw_extra_seed in prop::collection::vec(0u32..256, 0..96),
    ) {
        let edges = raw_edges
            .into_iter()
            .map(|(src, dst)| (src % node_count, dst % node_count))
            .collect::<Vec<_>>();
        let (offsets, targets, masks) = csr_from_edges(node_count, &edges);
        let seed_a = seed_bitset(node_count, &raw_seed);
        let mut seed_b_nodes = raw_seed;
        seed_b_nodes.extend(raw_extra_seed);
        let seed_b = seed_bitset(node_count, &seed_b_nodes);

        let closure_a = cpu_subset_closure(node_count, &offsets, &targets, &masks, &seed_a);
        let closure_b = cpu_subset_closure(node_count, &offsets, &targets, &masks, &seed_b);

        assert_subset(&seed_a, &seed_b);
        assert_subset(&closure_a, &closure_b);
        assert_subset(&seed_a, &closure_a);
        assert_subset(&seed_b, &closure_b);
    }
}