weirflow 0.1.0

GPU-first dataflow analysis primitives for Vyre and Santh compiler pipelines.
Documentation 
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143

use super::*;

/// PHASE6_DATAFLOW HIGH regression: backward_slice now requires
/// a real ProgramGraphShape; pre-fix the 5-arg entry hardcoded
/// (1, 1), which produced a 1-thread dispatch grid that silently
/// lost every dependency edge in the supergraph.
#[test]
fn backward_slice_requires_caller_supplied_shape() {
    let shape = ProgramGraphShape::new(64, 128);
    let p = backward_slice(shape, "sink_in", "slice_out");
    let buffer_count = |name: &str| {
        p.buffers
            .iter()
            .find(|buffer| buffer.name() == name)
            .map(|buffer| buffer.count)
            .unwrap_or_else(|| panic!("{name} buffer must be declared"))
    };

    assert_eq!(
        buffer_count("sink_in"),
        2,
        "Fix: backward slice input frontier must be sized to the caller's full node domain."
    );
    assert_eq!(
        buffer_count("slice_out"),
        2,
        "Fix: backward slice output frontier must be sized to the caller's full node domain."
    );
    assert_eq!(
        buffer_count("pg_edge_targets"),
        128,
        "Fix: backward slice CSR target buffer must preserve the caller's full edge domain."
    );
    assert_eq!(
        buffer_count("pg_edge_offsets"),
        65,
        "Fix: backward slice CSR offset buffer must contain node_count + 1 offsets."
    );
}

#[test]
fn slice_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 & 0b10 != 0 {
                frontier | 0b01
            } else {
                frontier
            };
            if outputs.is_empty() {
                outputs.push(Vec::new());
            }
            outputs[0].clear();
            outputs[0].extend_from_slice(&next.to_le_bytes());
            Ok(())
        };

    let mut scratch = crate::fixed_point_scratch::FixedPointScratch::default();
    let result = slice_closure_borrowed_into_with_scratch_via(
        &dispatch,
        2,
        &[0, 1, 1],
        &[1],
        &[vyre_primitives::predicate::edge_kind::CONTROL],
        &[0b10],
        4,
        &mut scratch,
    )
    .expect("slice closure into dispatch must converge");

    assert_eq!(result, vec![0b11]);
    assert_eq!(calls.get(), 2);
    assert!(saw_reused_slot.get());
    assert!(scratch.frontier_word_capacity() >= 1);
}

#[test]
fn slice_closure_into_result_reuses_caller_result_storage() {
    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 frontier = u32::from_le_bytes(inputs[5][..4].try_into().unwrap());
            let next = if frontier & 0b10 != 0 {
                frontier | 0b01
            } else {
                frontier
            };
            if outputs.is_empty() {
                outputs.push(Vec::new());
            }
            outputs[0].clear();
            outputs[0].extend_from_slice(&next.to_le_bytes());
            Ok(())
        };

    let mut scratch = crate::fixed_point_scratch::FixedPointScratch::default();
    let mut result = Vec::with_capacity(1);
    let result_ptr = result.as_ptr();
    let result_capacity = result.capacity();

    slice_closure_borrowed_into_result_with_scratch_via(
        &dispatch,
        2,
        &[0, 1, 1],
        &[1],
        &[vyre_primitives::predicate::edge_kind::CONTROL],
        &[0b10],
        4,
        &mut scratch,
        &mut result,
    )
    .expect("slice closure must fill caller-owned result");

    assert_eq!(result, vec![0b11]);
    assert_eq!(result.as_ptr(), result_ptr);
    assert_eq!(result.capacity(), result_capacity);
}

/// PHASE6_DATAFLOW HIGH regression: the deprecated alias still
/// emits the same Program shape so back-compat callers continue
/// to compile.
#[test]
#[allow(deprecated)]
fn deprecated_alias_emits_same_program_shape() {
    let shape = ProgramGraphShape::new(32, 64);
    let canonical = backward_slice(shape, "fin", "fout");
    let alias = backward_slice_with_shape(shape, "fin", "fout");
    assert_eq!(canonical.workgroup_size, alias.workgroup_size);
    assert_eq!(canonical.buffers.len(), alias.buffers.len());
}