vyre_primitives/graph/
dominator_frontier.rs

1//! `dominator_frontier`  -  query the dominance frontier of a node
2//! set, packed as a per-node bitset.
3//!
4//! The dominance frontier of node `n` is the set of nodes `m` such
5//! that `n` dominates a predecessor of `m` but does NOT dominate `m`
6//! itself. SSA phi placement uses this directly; rule pipelines can
7//! reach for it via the `vyre.graph.dominator_frontier.v1` ExternCall.
8//!
9//! Soundness: exact when the supplied dominator-tree CSR is
10//! correctly computed (the caller is responsible for that  -  usually
11//! via `vyre-libs::dataflow::ssa::compute_dominators`).
12
13use std::sync::Arc;
14
15use vyre_foundation::ir::model::expr::Ident;
16use vyre_foundation::ir::{BufferAccess, BufferDecl, DataType, Expr, Node, Program};
17
18use crate::graph::csr_forward_traverse::bitset_words;
19
20/// Canonical op id.
21pub const OP_ID: &str = "vyre-primitives::graph::dominator_frontier";
22
23/// Dominance-closure CSR offsets input binding.
24pub const DOMINATOR_FRONTIER_DOM_OFFSETS_BUFFER: u32 = 0;
25/// Dominance-closure CSR targets input binding.
26pub const DOMINATOR_FRONTIER_DOM_TARGETS_BUFFER: u32 = 1;
27/// Predecessor CSR offsets input binding.
28pub const DOMINATOR_FRONTIER_PRED_OFFSETS_BUFFER: u32 = 2;
29/// Predecessor CSR targets input binding.
30pub const DOMINATOR_FRONTIER_PRED_TARGETS_BUFFER: u32 = 3;
31/// Seed bitset input binding.
32pub const DOMINATOR_FRONTIER_SEED_BUFFER: u32 = 4;
33/// Frontier bitset output binding.
34pub const DOMINATOR_FRONTIER_OUT_BUFFER: u32 = 5;
35/// Candidate-node workgroup for dominance-frontier queries.
36pub const DOMINATOR_FRONTIER_WORKGROUP_SIZE: [u32; 3] = [256, 1, 1];
37
38/// Dispatch grid for one dominance-frontier query over candidate nodes.
39#[must_use]
40pub const fn dominator_frontier_dispatch_grid(node_count: u32) -> [u32; 3] {
41    if node_count == 0 {
42        [0, 1, 1]
43    } else {
44        [
45            node_count.div_ceil(DOMINATOR_FRONTIER_WORKGROUP_SIZE[0]),
46            1,
47            1,
48        ]
49    }
50}
51
52/// Validated dominance-frontier dispatch layout.
53#[derive(Clone, Copy, Debug, Eq, PartialEq)]
54pub struct DominatorFrontierLayout {
55    /// Number of u32 words in the frontier/seed bitset.
56    pub words: usize,
57    /// Number of dominance-closure CSR edges.
58    pub dom_edge_count: u32,
59    /// Number of predecessor CSR edges.
60    pub pred_edge_count: u32,
61}
62
63/// Program-shape key for dominance-frontier IR materialization.
64#[derive(Clone, Copy, Debug, Eq, PartialEq)]
65pub struct DominatorFrontierProgramShape {
66    /// Number of candidate nodes.
67    pub node_count: u32,
68    /// Number of dominance-closure CSR edges.
69    pub dom_edge_count: u32,
70    /// Number of predecessor CSR edges.
71    pub pred_edge_count: u32,
72}
73
74/// Content fingerprint for one immutable dominance-frontier input slice.
75#[derive(Clone, Copy, Debug, Eq, PartialEq)]
76pub struct DominatorFrontierSliceFingerprint {
77    len: usize,
78    first: u32,
79    last: u32,
80    xor: u32,
81    sum: u64,
82}
83
84/// Primitive-owned identity for immutable dominance-frontier dispatch inputs.
85///
86/// Dynamic seed/frontier buffers are intentionally excluded: wrappers refresh
87/// those every dispatch. This key covers only graph shape and graph content
88/// that determine whether static device inputs can be reused safely.
89#[derive(Clone, Copy, Debug, Eq, PartialEq)]
90pub struct DominatorFrontierStaticInputKey {
91    shape: DominatorFrontierProgramShape,
92    layout: DominatorFrontierLayout,
93    dom_target_words: usize,
94    pred_target_words: usize,
95    frontier_words: usize,
96    dom_offsets: DominatorFrontierSliceFingerprint,
97    dom_targets: DominatorFrontierSliceFingerprint,
98    pred_offsets: DominatorFrontierSliceFingerprint,
99    pred_targets: DominatorFrontierSliceFingerprint,
100}
101
102/// Compute the primitive-owned fingerprint used for immutable dispatch inputs.
103#[must_use]
104pub fn dominator_frontier_slice_fingerprint(words: &[u32]) -> DominatorFrontierSliceFingerprint {
105    let mut xor = 0u32;
106    let mut sum = 0u64;
107    for &word in words {
108        xor ^= word;
109        sum = sum.wrapping_add(u64::from(word));
110    }
111    DominatorFrontierSliceFingerprint {
112        len: words.len(),
113        first: words.first().copied().unwrap_or(0),
114        last: words.last().copied().unwrap_or(0),
115        xor,
116        sum,
117    }
118}
119
120#[cfg(test)]
121mod static_input_key_tests {
122    use super::*;
123
124    #[test]
125    fn slice_fingerprint_tracks_interior_content_not_only_len_edges() {
126        let baseline = dominator_frontier_slice_fingerprint(&[7, 11, 13, 17]);
127        let changed = dominator_frontier_slice_fingerprint(&[7, 11, 19, 17]);
128
129        assert_ne!(baseline, changed);
130    }
131
132    #[test]
133    fn static_input_key_tracks_graph_content_but_not_dynamic_seed_bits() {
134        let plan_a = plan_dominator_frontier_launch(
135            4,
136            &[0, 4, 5, 6, 7],
137            &[0, 1, 2, 3, 1, 2, 3],
138            &[0, 0, 1, 2, 4],
139            &[0, 0, 1, 2],
140            &[0b0010],
141        )
142        .expect("Fix: valid dominator-frontier launch plan should build");
143        let plan_b = plan_dominator_frontier_launch(
144            4,
145            &[0, 4, 5, 6, 7],
146            &[0, 1, 2, 3, 1, 2, 3],
147            &[0, 0, 1, 2, 4],
148            &[0, 0, 1, 2],
149            &[0b0100],
150        )
151        .expect("Fix: seed-only changes should keep the same static launch shape");
152
153        let baseline = plan_a.static_input_key(
154            &[0, 4, 5, 6, 7],
155            &[0, 1, 2, 3, 1, 2, 3],
156            &[0, 0, 1, 2, 4],
157            &[0, 0, 1, 2],
158        );
159        let seed_only_change = plan_b.static_input_key(
160            &[0, 4, 5, 6, 7],
161            &[0, 1, 2, 3, 1, 2, 3],
162            &[0, 0, 1, 2, 4],
163            &[0, 0, 1, 2],
164        );
165        let graph_content_change = plan_a.static_input_key(
166            &[0, 4, 5, 6, 7],
167            &[0, 1, 2, 2, 1, 2, 3],
168            &[0, 0, 1, 2, 4],
169            &[0, 0, 1, 2],
170        );
171
172        assert_eq!(baseline, seed_only_change);
173        assert_ne!(baseline, graph_content_change);
174    }
175}
176
177/// Primitive-owned dominance-frontier launch plan without eager IR materialization.
178#[derive(Clone, Copy, Debug, Eq, PartialEq)]
179pub struct DominatorFrontierLaunchPlan {
180    layout: DominatorFrontierLayout,
181    shape: DominatorFrontierProgramShape,
182    dispatch_grid: [u32; 3],
183}
184
185impl DominatorFrontierLaunchPlan {
186    /// Validated CSR and bitset layout.
187    #[must_use]
188    pub const fn layout(&self) -> DominatorFrontierLayout {
189        self.layout
190    }
191
192    /// Program-shape key for cache lookups.
193    #[must_use]
194    pub const fn shape(&self) -> DominatorFrontierProgramShape {
195        self.shape
196    }
197
198    /// Exact GPU dispatch grid for this query.
199    #[must_use]
200    pub const fn dispatch_grid(&self) -> [u32; 3] {
201        self.dispatch_grid
202    }
203
204    /// Number of u32 words in the seed/frontier bitsets.
205    #[must_use]
206    pub const fn frontier_words(&self) -> usize {
207        self.layout.words
208    }
209
210    /// Number of u32 target words required by the dominance-closure input.
211    #[must_use]
212    pub const fn dom_target_words(&self) -> usize {
213        if self.layout.dom_edge_count == 0 {
214            1
215        } else {
216            self.layout.dom_edge_count as usize
217        }
218    }
219
220    /// Number of u32 target words required by the predecessor input.
221    #[must_use]
222    pub const fn pred_target_words(&self) -> usize {
223        if self.layout.pred_edge_count == 0 {
224            1
225        } else {
226            self.layout.pred_edge_count as usize
227        }
228    }
229
230    /// Stable identity for immutable graph inputs associated with this plan.
231    #[must_use]
232    pub fn static_input_key(
233        &self,
234        dom_offsets: &[u32],
235        dom_targets: &[u32],
236        pred_offsets: &[u32],
237        pred_targets: &[u32],
238    ) -> DominatorFrontierStaticInputKey {
239        DominatorFrontierStaticInputKey {
240            shape: self.shape,
241            layout: self.layout,
242            dom_target_words: self.dom_target_words(),
243            pred_target_words: self.pred_target_words(),
244            frontier_words: self.frontier_words(),
245            dom_offsets: dominator_frontier_slice_fingerprint(dom_offsets),
246            dom_targets: dominator_frontier_slice_fingerprint(dom_targets),
247            pred_offsets: dominator_frontier_slice_fingerprint(pred_offsets),
248            pred_targets: dominator_frontier_slice_fingerprint(pred_targets),
249        }
250    }
251
252    /// Build the dominance-frontier Program for this launch plan.
253    pub fn program(&self, seed_buffer: &str, out_buffer: &str) -> Result<Program, String> {
254        try_dominator_frontier(
255            self.shape.node_count,
256            self.shape.dom_edge_count,
257            self.shape.pred_edge_count,
258            seed_buffer,
259            out_buffer,
260        )
261    }
262}
263
264/// Primitive-owned dominance-frontier dispatch plan with eager IR materialization.
265pub struct DominatorFrontierDispatchPlan {
266    launch: DominatorFrontierLaunchPlan,
267    program: Program,
268}
269
270impl DominatorFrontierDispatchPlan {
271    /// Validated CSR and bitset layout.
272    #[must_use]
273    pub const fn layout(&self) -> DominatorFrontierLayout {
274        self.launch.layout()
275    }
276
277    /// Program-shape key for cache lookups.
278    #[must_use]
279    pub const fn shape(&self) -> DominatorFrontierProgramShape {
280        self.launch.shape()
281    }
282
283    /// Program wired to the canonical primitive buffer layout.
284    #[must_use]
285    pub const fn program(&self) -> &Program {
286        &self.program
287    }
288
289    /// Exact GPU dispatch grid for this query.
290    #[must_use]
291    pub const fn dispatch_grid(&self) -> [u32; 3] {
292        self.launch.dispatch_grid()
293    }
294
295    /// Number of u32 words in the seed/frontier bitsets.
296    #[must_use]
297    pub const fn frontier_words(&self) -> usize {
298        self.launch.frontier_words()
299    }
300
301    /// Number of u32 target words required by the dominance-closure input.
302    #[must_use]
303    pub const fn dom_target_words(&self) -> usize {
304        self.launch.dom_target_words()
305    }
306
307    /// Number of u32 target words required by the predecessor input.
308    #[must_use]
309    pub const fn pred_target_words(&self) -> usize {
310        self.launch.pred_target_words()
311    }
312}
313
314/// Validate inputs and build a dominance-frontier launch plan without
315/// materializing IR.
316///
317/// # Errors
318///
319/// Returns an actionable diagnostic when either CSR is malformed, the seed
320/// bitset is not exactly shaped for `node_count`, or the dispatch shape would
321/// overflow.
322pub fn plan_dominator_frontier_launch(
323    node_count: u32,
324    dom_offsets: &[u32],
325    dom_targets: &[u32],
326    pred_offsets: &[u32],
327    pred_targets: &[u32],
328    seed: &[u32],
329) -> Result<DominatorFrontierLaunchPlan, String> {
330    let layout = validate_dominator_frontier_inputs(
331        node_count,
332        dom_offsets,
333        dom_targets,
334        pred_offsets,
335        pred_targets,
336        seed,
337    )?;
338    let _offset_count = node_count.checked_add(1).ok_or_else(|| {
339        format!(
340            "dominator_frontier node_count={node_count} overflows CSR offset buffer count. Fix: shard the graph before GPU dispatch."
341        )
342    })?;
343
344    Ok(DominatorFrontierLaunchPlan {
345        layout,
346        shape: DominatorFrontierProgramShape {
347            node_count,
348            dom_edge_count: layout.dom_edge_count,
349            pred_edge_count: layout.pred_edge_count,
350        },
351        dispatch_grid: dominator_frontier_dispatch_grid(node_count),
352    })
353}
354
355/// Validate inputs and build the canonical dominance-frontier dispatch plan.
356///
357/// # Errors
358///
359/// Returns an actionable diagnostic when either CSR is malformed, the seed
360/// bitset is not exactly shaped for `node_count`, or the generated dispatch
361/// program would overflow its CSR launch shape.
362pub fn plan_dominator_frontier_dispatch(
363    node_count: u32,
364    dom_offsets: &[u32],
365    dom_targets: &[u32],
366    pred_offsets: &[u32],
367    pred_targets: &[u32],
368    seed: &[u32],
369    seed_buffer: &str,
370    out_buffer: &str,
371) -> Result<DominatorFrontierDispatchPlan, String> {
372    let launch = plan_dominator_frontier_launch(
373        node_count,
374        dom_offsets,
375        dom_targets,
376        pred_offsets,
377        pred_targets,
378        seed,
379    )?;
380    let program = launch.program(seed_buffer, out_buffer)?;
381
382    Ok(DominatorFrontierDispatchPlan { launch, program })
383}
384
385/// Build a Program that evaluates the exact dominance-frontier
386/// predicate:
387///
388/// `m ∈ DF(seed)` iff some seeded node `n` dominates at least one
389/// predecessor of `m`, and `n` does not strictly dominate `m`.
390///
391/// `dom_offsets`/`dom_targets` encode dominance closure by dominator.
392/// `pred_offsets`/`pred_targets` encode CFG predecessors by candidate
393/// node.
394///
395/// # Panics
396/// Panics on an invalid CSR launch shape. Callers that must recover use the checked
397/// twin below.
398#[must_use]
399pub fn dominator_frontier(
400    node_count: u32,
401    dom_edge_count: u32,
402    pred_edge_count: u32,
403    seed: &str,
404    out: &str,
405) -> Program {
406    // Fail fast on an overflowing CSR launch shape rather than silently
407    // degrading to an inert empty kernel (silent recall loss). Use
408    // `try_dominator_frontier` for structured handling.
409    try_dominator_frontier(node_count, dom_edge_count, pred_edge_count, seed, out)
410        .unwrap_or_else(|error| panic!("{error}"))
411}
412
413/// Build a dominance-frontier Program with checked CSR launch-shape
414/// validation.
415pub fn try_dominator_frontier(
416    node_count: u32,
417    dom_edge_count: u32,
418    pred_edge_count: u32,
419    seed: &str,
420    out: &str,
421) -> Result<Program, String> {
422    let words = bitset_words(node_count).max(1);
423    let offset_count = node_count.checked_add(1).ok_or_else(|| {
424        format!(
425            "dominator_frontier node_count={node_count} overflows CSR offset buffer count. Fix: shard the graph before GPU dispatch."
426        )
427    });
428    let offset_count = offset_count?;
429    let t = Expr::InvocationId { axis: 0 };
430    let dominator_is_seed = vec![
431        Node::let_bind(
432            "seed_word",
433            Expr::load(seed, Expr::shr(Expr::var("n"), Expr::u32(5))),
434        ),
435        Node::let_bind(
436            "seed_bit",
437            Expr::shl(Expr::u32(1), Expr::bitand(Expr::var("n"), Expr::u32(31))),
438        ),
439        Node::if_then(
440            Expr::ne(
441                Expr::bitand(Expr::var("seed_word"), Expr::var("seed_bit")),
442                Expr::u32(0),
443            ),
444            vec![
445                Node::let_bind(
446                    "pred_start",
447                    Expr::load("pred_offsets", Expr::var("candidate")),
448                ),
449                Node::let_bind(
450                    "pred_end",
451                    Expr::load(
452                        "pred_offsets",
453                        Expr::add(Expr::var("candidate"), Expr::u32(1)),
454                    ),
455                ),
456                Node::let_bind("dominates_a_predecessor", Expr::u32(0)),
457                Node::loop_for(
458                    "p",
459                    Expr::var("pred_start"),
460                    Expr::var("pred_end"),
461                    vec![Node::if_then(
462                        Expr::eq(Expr::var("dominates_a_predecessor"), Expr::u32(0)),
463                        vec![
464                            Node::let_bind("pred", Expr::load("pred_targets", Expr::var("p"))),
465                            Node::let_bind(
466                                "dom_start_pred",
467                                Expr::load("dom_offsets", Expr::var("n")),
468                            ),
469                            Node::let_bind(
470                                "dom_end_pred",
471                                Expr::load("dom_offsets", Expr::add(Expr::var("n"), Expr::u32(1))),
472                            ),
473                            Node::loop_for(
474                                "d_pred",
475                                Expr::var("dom_start_pred"),
476                                Expr::var("dom_end_pred"),
477                                vec![Node::if_then(
478                                    Expr::eq(
479                                        Expr::load("dom_targets", Expr::var("d_pred")),
480                                        Expr::var("pred"),
481                                    ),
482                                    vec![Node::assign("dominates_a_predecessor", Expr::u32(1))],
483                                )],
484                            ),
485                        ],
486                    )],
487                ),
488                Node::let_bind("dominates_candidate", Expr::u32(0)),
489                Node::let_bind(
490                    "dom_start_candidate",
491                    Expr::load("dom_offsets", Expr::var("n")),
492                ),
493                Node::let_bind(
494                    "dom_end_candidate",
495                    Expr::load("dom_offsets", Expr::add(Expr::var("n"), Expr::u32(1))),
496                ),
497                Node::loop_for(
498                    "d_candidate",
499                    Expr::var("dom_start_candidate"),
500                    Expr::var("dom_end_candidate"),
501                    vec![Node::if_then(
502                        Expr::eq(
503                            Expr::load("dom_targets", Expr::var("d_candidate")),
504                            Expr::var("candidate"),
505                        ),
506                        vec![Node::assign("dominates_candidate", Expr::u32(1))],
507                    )],
508                ),
509                Node::let_bind("strictly_dominates_candidate", Expr::u32(0)),
510                Node::if_then(
511                    Expr::and(
512                        Expr::eq(Expr::var("dominates_candidate"), Expr::u32(1)),
513                        Expr::ne(Expr::var("n"), Expr::var("candidate")),
514                    ),
515                    vec![Node::assign("strictly_dominates_candidate", Expr::u32(1))],
516                ),
517                Node::if_then(
518                    Expr::and(
519                        Expr::eq(Expr::var("dominates_a_predecessor"), Expr::u32(1)),
520                        Expr::eq(Expr::var("strictly_dominates_candidate"), Expr::u32(0)),
521                    ),
522                    vec![
523                        Node::let_bind(
524                            "candidate_word",
525                            Expr::shr(Expr::var("candidate"), Expr::u32(5)),
526                        ),
527                        Node::let_bind(
528                            "candidate_bit",
529                            Expr::shl(
530                                Expr::u32(1),
531                                Expr::bitand(Expr::var("candidate"), Expr::u32(31)),
532                            ),
533                        ),
534                        Node::let_bind(
535                            "_prev",
536                            Expr::atomic_or(
537                                out,
538                                Expr::var("candidate_word"),
539                                Expr::var("candidate_bit"),
540                            ),
541                        ),
542                    ],
543                ),
544            ],
545        ),
546    ];
547    Ok(Program::wrapped(
548        vec![
549            BufferDecl::storage(
550                "dom_offsets",
551                DOMINATOR_FRONTIER_DOM_OFFSETS_BUFFER,
552                BufferAccess::ReadOnly,
553                DataType::U32,
554            )
555            .with_count(offset_count),
556            BufferDecl::storage(
557                "dom_targets",
558                DOMINATOR_FRONTIER_DOM_TARGETS_BUFFER,
559                BufferAccess::ReadOnly,
560                DataType::U32,
561            )
562            .with_count(dom_edge_count.max(1)),
563            BufferDecl::storage(
564                "pred_offsets",
565                DOMINATOR_FRONTIER_PRED_OFFSETS_BUFFER,
566                BufferAccess::ReadOnly,
567                DataType::U32,
568            )
569            .with_count(offset_count),
570            BufferDecl::storage(
571                "pred_targets",
572                DOMINATOR_FRONTIER_PRED_TARGETS_BUFFER,
573                BufferAccess::ReadOnly,
574                DataType::U32,
575            )
576            .with_count(pred_edge_count.max(1)),
577            BufferDecl::storage(
578                seed,
579                DOMINATOR_FRONTIER_SEED_BUFFER,
580                BufferAccess::ReadOnly,
581                DataType::U32,
582            )
583            .with_count(words),
584            BufferDecl::storage(
585                out,
586                DOMINATOR_FRONTIER_OUT_BUFFER,
587                BufferAccess::ReadWrite,
588                DataType::U32,
589            )
590            .with_count(words),
591        ],
592        DOMINATOR_FRONTIER_WORKGROUP_SIZE,
593        vec![Node::Region {
594            generator: Ident::from(OP_ID),
595            source_region: None,
596            body: Arc::new(vec![Node::if_then(
597                Expr::lt(t.clone(), Expr::u32(node_count)),
598                vec![
599                    Node::let_bind("candidate", t),
600                    Node::loop_for("n", Expr::u32(0), Expr::u32(node_count), dominator_is_seed),
601                ],
602            )]),
603        }],
604    ))
605}
606
607/// CPU oracle: returns the dominance-frontier bitset for the seed set.
608///
609/// `dom_offsets` / `dom_targets` encode the dominance closure by dominator:
610/// row `n` contains every node dominated by `n`, including `n`.
611#[must_use]
612#[cfg(any(test, feature = "cpu-parity"))]
613pub fn cpu_ref(
614    node_count: u32,
615    dom_offsets: &[u32],
616    dom_targets: &[u32],
617    pred_offsets: &[u32],
618    pred_targets: &[u32],
619    seed: &[u32],
620) -> Vec<u32> {
621    try_cpu_ref(
622        node_count,
623        dom_offsets,
624        dom_targets,
625        pred_offsets,
626        pred_targets,
627        seed,
628    )
629    .expect("Fix: reject malformed oracle input via try_* APIs; do not call panicking wrappers on hostile data - dominator_frontier CPU oracle received malformed input or could not reserve output")
630}
631
632/// Fallible CPU oracle: returns the dominance-frontier bitset for the seed set.
633///
634/// This is the primitive-owned entry point for wrappers and generated tests that
635/// must reject hostile CSR/seed inputs without panicking.
636#[cfg(any(test, feature = "cpu-parity"))]
637pub fn try_cpu_ref(
638    node_count: u32,
639    dom_offsets: &[u32],
640    dom_targets: &[u32],
641    pred_offsets: &[u32],
642    pred_targets: &[u32],
643    seed: &[u32],
644) -> Result<Vec<u32>, String> {
645    let mut frontier = Vec::new();
646    try_cpu_ref_into(
647        node_count,
648        dom_offsets,
649        dom_targets,
650        pred_offsets,
651        pred_targets,
652        seed,
653        &mut frontier,
654    )?;
655    Ok(frontier)
656}
657
658/// CPU oracle into caller-owned output storage.
659///
660/// `dom_offsets` / `dom_targets` encode the dominance closure by dominator:
661/// row `n` contains every node dominated by `n`, including `n`.
662#[cfg(any(test, feature = "cpu-parity"))]
663pub fn cpu_ref_into(
664    node_count: u32,
665    dom_offsets: &[u32],
666    dom_targets: &[u32],
667    pred_offsets: &[u32],
668    pred_targets: &[u32],
669    seed: &[u32],
670    frontier: &mut Vec<u32>,
671) {
672    try_cpu_ref_into(
673        node_count,
674        dom_offsets,
675        dom_targets,
676        pred_offsets,
677        pred_targets,
678        seed,
679        frontier,
680    )
681    .expect("Fix: reject malformed oracle input via try_* APIs; do not call panicking wrappers on hostile data - dominator_frontier CPU oracle received malformed input or could not reserve output")
682}
683
684/// Fallible CPU oracle into caller-owned output storage.
685///
686/// On error, `frontier` is left unchanged so dispatch wrappers and parity tests
687/// can surface malformed input as a typed finding instead of losing the last
688/// useful diagnostic output.
689#[cfg(any(test, feature = "cpu-parity"))]
690pub fn try_cpu_ref_into(
691    node_count: u32,
692    dom_offsets: &[u32],
693    dom_targets: &[u32],
694    pred_offsets: &[u32],
695    pred_targets: &[u32],
696    seed: &[u32],
697    frontier: &mut Vec<u32>,
698) -> Result<(), String> {
699    let layout = validate_dominator_frontier_inputs(
700        node_count,
701        dom_offsets,
702        dom_targets,
703        pred_offsets,
704        pred_targets,
705        seed,
706    )?;
707    let words = layout.words;
708    crate::graph::scratch::reserve_graph_items(
709        frontier,
710        words,
711        "dominator frontier CPU oracle",
712        "frontier output",
713    )?;
714    frontier.clear();
715    frontier.resize(words, 0);
716    for n in 0..node_count {
717        let n_word = (n / 32) as usize;
718        let n_bit = 1u32 << (n % 32);
719        if seed[n_word] & n_bit == 0 {
720            continue;
721        }
722        for m in 0..node_count {
723            let pred_start = pred_offsets[m as usize] as usize;
724            let pred_end = pred_offsets[m as usize + 1] as usize;
725            let dominates_a_predecessor = pred_targets[pred_start..pred_end]
726                .iter()
727                .copied()
728                .any(|pred| dominates(dom_offsets, dom_targets, n, pred));
729            let strictly_dominates_m = n != m && dominates(dom_offsets, dom_targets, n, m);
730            if dominates_a_predecessor && !strictly_dominates_m {
731                let m_word = (m / 32) as usize;
732                let m_bit = 1u32 << (m % 32);
733                frontier[m_word] |= m_bit;
734            }
735        }
736    }
737    Ok(())
738}
739
740/// Number of nodes flagged in a dominance-frontier bitset.
741#[must_use]
742pub fn frontier_size(frontier: &[u32]) -> u32 {
743    frontier.iter().map(|word| word.count_ones()).sum()
744}
745
746/// Validate a CSR buffer pair for `node_count` rows.
747///
748/// # Errors
749///
750/// Returns an actionable diagnostic when offsets are the wrong length,
751/// non-monotonic, inconsistent with target count, or targets point outside
752/// `0..node_count`.
753pub fn validate_csr_shape(
754    label: &str,
755    node_count: u32,
756    offsets: &[u32],
757    targets: &[u32],
758) -> Result<u32, String> {
759    let expected_offsets = (node_count as usize).checked_add(1).ok_or_else(|| {
760        format!(
761            "Fix: dominator_frontier {label} node_count + 1 overflows usize for node_count={node_count}."
762        )
763    })?;
764    if offsets.len() != expected_offsets {
765        return Err(format!(
766            "Fix: dominator_frontier {label} offsets length must be {expected_offsets}, got {}.",
767            offsets.len()
768        ));
769    }
770    let mut previous = 0u32;
771    for (idx, &offset) in offsets.iter().enumerate() {
772        if idx > 0 && offset < previous {
773            return Err(format!(
774                "Fix: dominator_frontier {label} offsets must be monotonic; offsets[{idx}]={offset} after {previous}."
775            ));
776        }
777        previous = offset;
778    }
779    if offsets.last().copied().unwrap_or(0) as usize != targets.len() {
780        return Err(format!(
781            "Fix: dominator_frontier {label} final offset must equal target count {}, got {}.",
782            targets.len(),
783            offsets.last().copied().unwrap_or(0)
784        ));
785    }
786    for (idx, &target) in targets.iter().enumerate() {
787        if target >= node_count {
788            return Err(format!(
789                "Fix: dominator_frontier {label} target[{idx}]={target} is outside node_count {node_count}."
790            ));
791        }
792    }
793    u32::try_from(targets.len()).map_err(|_| {
794        format!(
795            "Fix: dominator_frontier {label} target count {} exceeds u32 index space.",
796            targets.len()
797        )
798    })
799}
800
801/// Validate the full dominance-frontier CPU/dispatch input contract.
802///
803/// # Errors
804///
805/// Returns an actionable diagnostic when either CSR is malformed or when the
806/// seed bitset does not contain exactly the required number of words.
807pub fn validate_dominator_frontier_inputs(
808    node_count: u32,
809    dom_offsets: &[u32],
810    dom_targets: &[u32],
811    pred_offsets: &[u32],
812    pred_targets: &[u32],
813    seed: &[u32],
814) -> Result<DominatorFrontierLayout, String> {
815    let words = bitset_words(node_count) as usize;
816    if seed.len() != words {
817        return Err(format!(
818            "Fix: dominator_frontier expected seed length {words} words for {node_count} nodes, got {}.",
819            seed.len()
820        ));
821    }
822    let dom_edge_count = validate_csr_shape("dominance", node_count, dom_offsets, dom_targets)?;
823    let pred_edge_count =
824        validate_csr_shape("predecessor", node_count, pred_offsets, pred_targets)?;
825    Ok(DominatorFrontierLayout {
826        words,
827        dom_edge_count,
828        pred_edge_count,
829    })
830}
831
832#[cfg(any(test, feature = "cpu-parity"))]
833fn dominates(dom_offsets: &[u32], dom_targets: &[u32], dominator: u32, node: u32) -> bool {
834    let start = dom_offsets[dominator as usize] as usize;
835    let end = dom_offsets[dominator as usize + 1] as usize;
836    dom_targets[start..end].contains(&node)
837}
838
839#[cfg(feature = "inventory-registry")]
840inventory::submit! {
841    vyre_foundation::operation::OperationRegistration::primitive(
842        OP_ID,
843        || dominator_frontier(4, 4, 4, "idom", "df"),
844        Some(|| {
845            vec![vec![
846                crate::wire::pack_u32_slice(&[0, 1, 2, 3, 4]),
847                crate::wire::pack_u32_slice(&[0, 1, 2, 3]),
848                crate::wire::pack_u32_slice(&[0, 0, 1, 2, 3]),
849                crate::wire::pack_u32_slice(&[0, 1, 2, 0]),
850                crate::wire::pack_u32_slice(&[0]),
851                crate::wire::pack_u32_slice(&[0]),
852            ]]
853        }),
854        Some(|| {
855            vec![vec![crate::wire::pack_u32_slice(&[0])]]
856        }),
857    )
858}
859
860#[cfg(test)]
861mod tests {
862    use super::*;
863
864    #[test]
865    fn empty_seed_yields_empty_frontier() {
866        let out = cpu_ref(4, &[0, 0, 0, 0, 0], &[], &[0, 0, 0, 0, 0], &[], &[0]);
867        assert_eq!(out, vec![0]);
868    }
869
870    #[test]
871    fn single_node_with_no_predecessors_has_empty_frontier() {
872        // node 0 with no predecessors → df(0) = {}.
873        let out = cpu_ref(2, &[0, 0, 0], &[], &[0, 0, 0], &[], &[0b01]);
874        assert_eq!(out, vec![0]);
875    }
876
877    #[test]
878    fn dom_frontier_picks_up_join_node() {
879        // CFG: 0 -> 1, 0 -> 2, 1 -> 3, 2 -> 3.
880        // Predecessors of 3: [1, 2]. 1 dominates itself only, 2 same.
881        // df(1) includes 3 because 1 dominates predecessor 1 of 3,
882        // but 1 does not dominate 3.
883        let pred_offsets = vec![0u32, 0, 1, 2, 4];
884        let pred_targets = vec![0u32, 0, 1, 2];
885        // Dominator sets: 0 dominates {0,1,2,3}; 1 dominates {1};
886        // 2 dominates {2}; 3 dominates {3}.
887        let dom_offsets = vec![0u32, 4, 5, 6, 7];
888        let dom_targets = vec![0u32, 1, 2, 3, 1, 2, 3];
889        let out = cpu_ref(
890            4,
891            &dom_offsets,
892            &dom_targets,
893            &pred_offsets,
894            &pred_targets,
895            &[0b0010],
896        );
897        assert_eq!(out, vec![0b1000]);
898    }
899
900    #[test]
901    fn cpu_ref_into_reuses_frontier_storage() {
902        let mut out = Vec::with_capacity(8);
903        let dom_offsets = vec![0u32, 4, 5, 6, 7];
904        let dom_targets = vec![0u32, 1, 2, 3, 1, 2, 3];
905        let pred_offsets = vec![0u32, 0, 1, 2, 4];
906        let pred_targets = vec![0u32, 0, 1, 2];
907        cpu_ref_into(
908            4,
909            &dom_offsets,
910            &dom_targets,
911            &pred_offsets,
912            &pred_targets,
913            &[0b0010],
914            &mut out,
915        );
916        let capacity = out.capacity();
917        assert_eq!(out, vec![0b1000]);
918
919        cpu_ref_into(
920            4,
921            &dom_offsets,
922            &dom_targets,
923            &pred_offsets,
924            &pred_targets,
925            &[0],
926            &mut out,
927        );
928        assert_eq!(out.capacity(), capacity);
929        assert_eq!(out, vec![0]);
930    }
931
932    #[test]
933    fn cpu_ref_into_uses_shared_validation_before_output_mutation() {
934        let mut out = vec![0xCAFE_BABEu32];
935        let ptr = out.as_ptr();
936        let err = try_cpu_ref_into(
937            4,
938            &[0, 4, 5, 6, 7],
939            &[0, 1, 2, 3, 1, 2, 3],
940            &[0, 0, 1, 2, 4],
941            &[0, 1, 2, 3],
942            &[0b0010, 0],
943            &mut out,
944        );
945
946        assert!(err.is_err(), "extra seed word must be rejected");
947        assert_eq!(
948            out,
949            vec![0xCAFE_BABEu32],
950            "Fix: shared validation must reject malformed input before clearing caller output."
951        );
952        assert_eq!(out.as_ptr(), ptr);
953    }
954
955    #[test]
956    fn fallible_cpu_ref_matches_compatibility_oracle_on_generated_diamonds() {
957        for diamond_count in [1_u32, 2, 7, 16, 33, 64] {
958            let node_count = diamond_count
959                .checked_mul(4)
960                .expect("Fix: generated diamond node count should fit u32");
961            let words = bitset_words(node_count) as usize;
962            let mut dom_offsets = Vec::with_capacity(node_count as usize + 1);
963            let mut dom_targets = Vec::new();
964            let mut pred_offsets = Vec::with_capacity(node_count as usize + 1);
965            let mut pred_targets = Vec::new();
966            dom_offsets.push(0);
967            pred_offsets.push(0);
968            for diamond in 0..diamond_count {
969                let base = diamond * 4;
970                dom_targets.extend_from_slice(&[base, base + 1, base + 2, base + 3]);
971                dom_offsets.push(dom_targets.len() as u32);
972                dom_targets.push(base + 1);
973                dom_offsets.push(dom_targets.len() as u32);
974                dom_targets.push(base + 2);
975                dom_offsets.push(dom_targets.len() as u32);
976                dom_targets.push(base + 3);
977                dom_offsets.push(dom_targets.len() as u32);
978
979                pred_offsets.push(pred_targets.len() as u32);
980                pred_targets.push(base);
981                pred_offsets.push(pred_targets.len() as u32);
982                pred_targets.push(base);
983                pred_offsets.push(pred_targets.len() as u32);
984                pred_targets.extend_from_slice(&[base + 1, base + 2]);
985                pred_offsets.push(pred_targets.len() as u32);
986            }
987            let mut seed = vec![0; words];
988            for diamond in 0..diamond_count {
989                let node = diamond * 4 + 1;
990                seed[(node / 32) as usize] |= 1_u32 << (node % 32);
991            }
992
993            let expected = cpu_ref(
994                node_count,
995                &dom_offsets,
996                &dom_targets,
997                &pred_offsets,
998                &pred_targets,
999                &seed,
1000            );
1001            let actual = try_cpu_ref(
1002                node_count,
1003                &dom_offsets,
1004                &dom_targets,
1005                &pred_offsets,
1006                &pred_targets,
1007                &seed,
1008            )
1009            .expect("Fix: generated dominance-frontier diamonds should run fallibly");
1010            assert_eq!(actual, expected, "diamond_count={diamond_count}");
1011        }
1012    }
1013
1014    #[test]
1015    fn reusable_validation_rejects_bad_csr_and_seed() {
1016        let err = validate_dominator_frontier_inputs(2, &[0, 1, 1], &[1], &[0, 1, 0], &[0], &[1])
1017            .unwrap_err();
1018        assert!(err.contains("predecessor offsets must be monotonic"));
1019
1020        let err =
1021            validate_dominator_frontier_inputs(33, &[0; 34], &[], &[0; 34], &[], &[1]).unwrap_err();
1022        assert!(err.contains("expected seed length 2 words"));
1023    }
1024
1025    #[test]
1026    fn reusable_validation_returns_dispatch_layout() {
1027        let layout = validate_dominator_frontier_inputs(
1028            4,
1029            &[0, 4, 5, 6, 7],
1030            &[0, 1, 2, 3, 1, 2, 3],
1031            &[0, 0, 1, 2, 4],
1032            &[0, 1, 2, 3],
1033            &[0b0010],
1034        )
1035        .expect("Fix: canonical dominance-frontier input should validate");
1036
1037        assert_eq!(
1038            layout,
1039            DominatorFrontierLayout {
1040                words: 1,
1041                dom_edge_count: 7,
1042                pred_edge_count: 4,
1043            }
1044        );
1045    }
1046
1047    #[test]
1048    fn launch_plan_validates_layout_without_eager_program_materialization() {
1049        let plan = plan_dominator_frontier_launch(
1050            4,
1051            &[0, 4, 5, 6, 7],
1052            &[0, 1, 2, 3, 1, 2, 3],
1053            &[0, 0, 1, 2, 4],
1054            &[0, 1, 2, 3],
1055            &[0b0010],
1056        )
1057        .expect("Fix: canonical dominance-frontier launch plan should validate");
1058
1059        assert_eq!(plan.dispatch_grid(), [1, 1, 1]);
1060        assert_eq!(plan.frontier_words(), 1);
1061        assert_eq!(plan.dom_target_words(), 7);
1062        assert_eq!(plan.pred_target_words(), 4);
1063        assert_eq!(
1064            plan.shape(),
1065            DominatorFrontierProgramShape {
1066                node_count: 4,
1067                dom_edge_count: 7,
1068                pred_edge_count: 4,
1069            }
1070        );
1071        assert_eq!(
1072            plan.program("seed", "frontier_out")
1073                .expect("Fix: validated launch plan should materialize IR")
1074                .workgroup_size,
1075            DOMINATOR_FRONTIER_WORKGROUP_SIZE
1076        );
1077    }
1078
1079    #[test]
1080    fn launch_plan_packs_candidate_lanes_into_blocks() {
1081        assert_eq!(dominator_frontier_dispatch_grid(0), [0, 1, 1]);
1082        assert_eq!(dominator_frontier_dispatch_grid(1), [1, 1, 1]);
1083        assert_eq!(dominator_frontier_dispatch_grid(256), [1, 1, 1]);
1084        assert_eq!(dominator_frontier_dispatch_grid(257), [2, 1, 1]);
1085        assert_eq!(dominator_frontier_dispatch_grid(513), [3, 1, 1]);
1086    }
1087
1088    #[test]
1089    fn generated_launch_grid_covers_candidate_shapes_to_8192() {
1090        for node_count in 1..=8_192 {
1091            let grid = dominator_frontier_dispatch_grid(node_count);
1092            assert_eq!(
1093                grid[1], 1,
1094                "Fix: dominator-frontier grid y dimension drifted at node_count={node_count}"
1095            );
1096            assert_eq!(
1097                grid[2], 1,
1098                "Fix: dominator-frontier grid z dimension drifted at node_count={node_count}"
1099            );
1100            assert!(
1101                grid[0] * DOMINATOR_FRONTIER_WORKGROUP_SIZE[0] >= node_count,
1102                "Fix: dominator-frontier grid under-covers node_count={node_count}"
1103            );
1104            assert!(
1105                grid[0] == 1 || (grid[0] - 1) * DOMINATOR_FRONTIER_WORKGROUP_SIZE[0] < node_count,
1106                "Fix: dominator-frontier grid over-launches an avoidable extra block at node_count={node_count}"
1107            );
1108        }
1109    }
1110
1111    #[test]
1112    fn dispatch_plan_owns_buffer_slots_grid_and_readback_shape() {
1113        let plan = plan_dominator_frontier_dispatch(
1114            4,
1115            &[0, 4, 5, 6, 7],
1116            &[0, 1, 2, 3, 1, 2, 3],
1117            &[0, 0, 1, 2, 4],
1118            &[0, 1, 2, 3],
1119            &[0b0010],
1120            "seed",
1121            "frontier_out",
1122        )
1123        .expect("Fix: canonical dominance-frontier plan should validate");
1124
1125        assert_eq!(plan.dispatch_grid(), [1, 1, 1]);
1126        assert_eq!(plan.frontier_words(), 1);
1127        assert_eq!(plan.dom_target_words(), 7);
1128        assert_eq!(plan.pred_target_words(), 4);
1129        assert_eq!(
1130            plan.program().workgroup_size,
1131            DOMINATOR_FRONTIER_WORKGROUP_SIZE
1132        );
1133        let bindings = plan
1134            .program()
1135            .buffers
1136            .iter()
1137            .map(|buffer| buffer.binding)
1138            .collect::<Vec<_>>();
1139        assert_eq!(
1140            bindings,
1141            vec![
1142                DOMINATOR_FRONTIER_DOM_OFFSETS_BUFFER,
1143                DOMINATOR_FRONTIER_DOM_TARGETS_BUFFER,
1144                DOMINATOR_FRONTIER_PRED_OFFSETS_BUFFER,
1145                DOMINATOR_FRONTIER_PRED_TARGETS_BUFFER,
1146                DOMINATOR_FRONTIER_SEED_BUFFER,
1147                DOMINATOR_FRONTIER_OUT_BUFFER,
1148            ]
1149        );
1150    }
1151
1152    #[test]
1153    fn dispatch_plan_pads_empty_target_buffers_without_hiding_empty_offsets() {
1154        let plan = plan_dominator_frontier_dispatch(
1155            1,
1156            &[0, 0],
1157            &[],
1158            &[0, 0],
1159            &[],
1160            &[1],
1161            "seed",
1162            "frontier_out",
1163        )
1164        .expect("Fix: empty edge sets are valid CSR inputs");
1165
1166        assert_eq!(plan.layout().dom_edge_count, 0);
1167        assert_eq!(plan.layout().pred_edge_count, 0);
1168        assert_eq!(plan.dom_target_words(), 1);
1169        assert_eq!(plan.pred_target_words(), 1);
1170    }
1171
1172    #[test]
1173    fn frontier_size_counts_set_bits() {
1174        assert_eq!(frontier_size(&[0]), 0);
1175        assert_eq!(frontier_size(&[0b1011]), 3);
1176        assert_eq!(frontier_size(&[u32::MAX, 1]), 33);
1177    }
1178
1179    #[test]
1180    fn checked_builder_rejects_offset_count_overflow() {
1181        let error = try_dominator_frontier(u32::MAX, 0, 0, "seed", "out")
1182            .expect_err("checked dominator-frontier builder must reject CSR offset overflow");
1183
1184        assert!(
1185            error.contains("overflows CSR offset buffer count"),
1186            "error should describe the CSR offset overflow: {error}"
1187        );
1188    }
1189
1190    #[test]
1191    fn legacy_builder_fails_fast_on_offset_count_overflow() {
1192        let panic = std::panic::catch_unwind(|| {
1193            let _ = dominator_frontier(u32::MAX, 0, 0, "seed", "out");
1194        })
1195        .expect_err("legacy dominator-frontier builder must fail fast on CSR offset overflow");
1196
1197        let message = panic_payload_message(panic);
1198        assert!(
1199            message.contains("overflows CSR offset buffer count"),
1200            "error should describe the CSR offset overflow: {message}"
1201        );
1202    }
1203
1204    fn panic_payload_message(payload: Box<dyn std::any::Any + Send>) -> String {
1205        if let Some(message) = payload.downcast_ref::<&str>() {
1206            message.to_string()
1207        } else if let Some(message) = payload.downcast_ref::<String>() {
1208            message.clone()
1209        } else {
1210            format!("{payload:?}")
1211        }
1212    }
1213
1214    #[test]
1215    fn missing_seed_word_fails_loudly() {
1216        let previous_hook = std::panic::take_hook();
1217        std::panic::set_hook(Box::new(|_| {}));
1218        let err = std::panic::catch_unwind(|| {
1219            let _ = cpu_ref(2, &[0, 0, 0], &[], &[0, 0, 0], &[], &[]);
1220        });
1221        std::panic::set_hook(previous_hook);
1222
1223        let payload = err.expect_err("missing seed word must fail loudly");
1224        let message = payload
1225            .downcast_ref::<String>()
1226            .map(String::as_str)
1227            .or_else(|| payload.downcast_ref::<&str>().copied())
1228            .unwrap_or("<non-string panic>");
1229        assert!(
1230            message.contains("expected seed length"),
1231            "Fix: missing seed panic should explain the exact seed length mismatch, got: {message}"
1232        );
1233    }
1234}
vyre_primitives/graph/dominator_frontier.rs

vyre_primitives/graph/
dominator_frontier.rs
