Skip to main content

vyre_primitives/graph/
path_reconstruct.rs

1//! `path_reconstruct`  -  walk a parent-pointer array back from a
2//! target node, emitting the materialized path into an output
3//! buffer.
4//!
5//! Given:
6//! - `parent`: u32 buffer where `parent[v] == u` means `u → v` is
7//!   the chosen predecessor edge (and `parent[root] == root` marks
8//!   termination).
9//! - `target`: u32 buffer whose slot 0 names the node to walk back
10//!   from.
11//!
12//! Emits `path_out[0..len]` = `[target, parent[target], parent[parent[target]], …, root]`
13//! and writes the path length into `path_len[0]`. Bounded by
14//! `max_depth` so a corrupt parent array cannot hang the GPU.
15
16use std::sync::Arc;
17
18use vyre_foundation::ir::model::expr::Ident;
19use vyre_foundation::ir::{BufferAccess, BufferDecl, DataType, Expr, Node, Program};
20
21/// Canonical op id.
22pub const OP_ID: &str = "vyre-primitives::graph::path_reconstruct";
23
24/// Canonical batched op id.
25pub const BATCHED_OP_ID: &str = "vyre-primitives::graph::batched_path_reconstruct";
26
27/// Workgroup size used by the batched path-reconstruction primitive.
28pub const BATCHED_WORKGROUP_SIZE: u32 = 256;
29/// Canonical parent input buffer.
30pub const PATH_PARENT_BUFFER: &str = "path_reconstruct parent";
31/// Canonical single-target input buffer.
32pub const PATH_TARGET_BUFFER: &str = "path_reconstruct target";
33/// Canonical single-path output buffer.
34pub const PATH_OUT_BUFFER: &str = "path_reconstruct path_out";
35/// Canonical single-path length output buffer.
36pub const PATH_LEN_BUFFER: &str = "path_reconstruct path_len";
37/// Canonical batched-target input buffer.
38pub const BATCHED_PATH_TARGETS_BUFFER: &str = "batched_path_reconstruct targets";
39/// Canonical batched-path output buffer.
40pub const BATCHED_PATHS_BUFFER: &str = "batched_path_reconstruct paths";
41/// Canonical batched-length output buffer.
42pub const BATCHED_LENS_BUFFER: &str = "batched_path_reconstruct lens";
43/// Single-lane path-reconstruction dispatch grid.
44pub const PATH_RECONSTRUCT_DISPATCH_GRID: [u32; 3] = [1, 1, 1];
45
46/// Validated batched path-reconstruction buffer layout.
47#[derive(Clone, Copy, Debug, Eq, PartialEq)]
48pub struct BatchedPathReconstructLayout {
49    /// Number of target nodes in the batch, narrowed for GPU dispatch lanes.
50    pub target_count: u32,
51    /// Total number of u32 path output words.
52    pub path_words: usize,
53    /// Total number of u32 path output words, narrowed for primitive buffer metadata.
54    pub path_words_u32: u32,
55}
56
57/// Primitive-owned dispatch plan for one path reconstruction.
58#[derive(Clone, Debug, Eq, PartialEq)]
59pub struct PathReconstructDispatchPlan {
60    /// Number of parent entries supplied by the caller.
61    pub parent_words: usize,
62    /// Number of words in the target buffer.
63    pub target_words: usize,
64    /// Number of words in the padded path output.
65    pub path_words: usize,
66    /// Number of words in the length output.
67    pub len_words: usize,
68    /// Maximum path depth accepted by this dispatch.
69    pub max_depth: u32,
70    /// Dispatch grid override.
71    pub grid: [u32; 3],
72}
73
74impl PathReconstructDispatchPlan {
75    /// Build the single-target path-reconstruction program for this plan.
76    #[must_use]
77    pub fn program(&self) -> Program {
78        path_reconstruct(
79            PATH_PARENT_BUFFER,
80            PATH_TARGET_BUFFER,
81            PATH_OUT_BUFFER,
82            PATH_LEN_BUFFER,
83            self.max_depth,
84        )
85    }
86
87    /// Return the primitive-owned cache identity for this plan's static parent input.
88    ///
89    /// # Errors
90    ///
91    /// Returns an actionable diagnostic when the parent slice no longer matches
92    /// the validated single-target dispatch plan.
93    pub fn static_input_key(
94        &self,
95        parent: &[u32],
96    ) -> Result<PathReconstructStaticInputKey, String> {
97        if parent.len() != self.parent_words {
98            return Err(format!(
99                "Fix: path_reconstruct static key expected {} parent word(s), got {}.",
100                self.parent_words,
101                parent.len()
102            ));
103        }
104        Ok(PathReconstructStaticInputKey {
105            parent_words: self.parent_words,
106            parent_hash: path_reconstruct_u32_slice_fingerprint(parent),
107            target_count: 1,
108            max_depth: self.max_depth,
109            batched: false,
110        })
111    }
112}
113
114/// Primitive-owned dispatch plan for batched path reconstruction.
115#[derive(Clone, Debug, Eq, PartialEq)]
116pub struct BatchedPathReconstructDispatchPlan {
117    /// Validated batched output layout.
118    pub layout: BatchedPathReconstructLayout,
119    /// Number of parent entries supplied by the caller.
120    pub parent_words: usize,
121    /// Number of words in the targets input.
122    pub target_words: usize,
123    /// Number of words in the padded paths output.
124    pub path_words: usize,
125    /// Number of words in the lengths output.
126    pub len_words: usize,
127    /// Maximum path depth accepted by this dispatch.
128    pub max_depth: u32,
129    /// Dispatch grid override.
130    pub grid: [u32; 3],
131}
132
133/// Primitive-owned identity for reusable path-reconstruction static inputs.
134///
135/// Dispatch wrappers use this key to decide whether the parent-pointer array
136/// can remain staged across calls. Keeping the fingerprint beside the primitive
137/// plan prevents every wrapper from forking parent-vector cache semantics.
138#[derive(Clone, Copy, Debug, Eq, PartialEq)]
139pub struct PathReconstructStaticInputKey {
140    /// Number of parent entries supplied by the caller.
141    pub parent_words: usize,
142    /// Stable content fingerprint for the parent-pointer array.
143    pub parent_hash: u64,
144    /// Number of target nodes in this dispatch.
145    pub target_count: u32,
146    /// Maximum path depth accepted by this dispatch.
147    pub max_depth: u32,
148    /// True when this key describes a batched dispatch.
149    pub batched: bool,
150}
151
152impl BatchedPathReconstructDispatchPlan {
153    /// Build the batched path-reconstruction program for this plan.
154    #[must_use]
155    pub fn program(&self) -> Program {
156        batched_path_reconstruct(self.layout.target_count, self.max_depth)
157    }
158
159    /// Return the primitive-owned cache identity for this plan's static parent input.
160    ///
161    /// # Errors
162    ///
163    /// Returns an actionable diagnostic when the parent slice no longer matches
164    /// the validated batched dispatch plan.
165    pub fn static_input_key(
166        &self,
167        parent: &[u32],
168    ) -> Result<PathReconstructStaticInputKey, String> {
169        if parent.len() != self.parent_words {
170            return Err(format!(
171                "Fix: batched_path_reconstruct static key expected {} parent word(s), got {}.",
172                self.parent_words,
173                parent.len()
174            ));
175        }
176        Ok(PathReconstructStaticInputKey {
177            parent_words: self.parent_words,
178            parent_hash: path_reconstruct_u32_slice_fingerprint(parent),
179            target_count: self.layout.target_count,
180            max_depth: self.max_depth,
181            batched: true,
182        })
183    }
184}
185
186fn path_reconstruct_u32_slice_fingerprint(values: &[u32]) -> u64 {
187    super::u32_slice_fingerprint(values)
188}
189
190/// Build the IR `Program` for path reconstruction.
191#[must_use]
192///
193pub fn path_reconstruct(
194    parent: &str,
195    target: &str,
196    path_out: &str,
197    path_len: &str,
198    max_depth: u32,
199) -> Program {
200    if max_depth == 0 {
201        return crate::invalid_output_program(
202            OP_ID,
203            path_out,
204            DataType::U32,
205            "Fix: path_reconstruct max_depth must be >= 1.".to_string(),
206        );
207    }
208    // Single-threaded walk (invocation 0 owns the chain). The work
209    // is O(path_length) which is typically small (stack trace length,
210    // tiny CFG path), so parallelism is not meaningful here.
211    //
212    // AUDIT_2026-04-24 F-PR-01: two divergences from cpu_ref fixed
213    // here.
214    //   (1) Prior code overloaded `len` as both the path-length
215    //       counter and the loop-termination signal (setting
216    //       `len = max_depth` on root-hit), so the stored
217    //       `path_len[0]` reported `max_depth` instead of the true
218    //       path length whenever a root was reached before the cap.
219    //       Now uses a separate `done` flag; `len` stays truthful.
220    //   (2) Prior code left `path_out[len..max_depth]` uninitialized
221    //       while cpu_ref explicitly pads that tail with zeros, so
222    //       harness byte-compare diverged unless the dispatcher
223    //       zeroed the buffer between runs. IR now writes 0 into
224    //       the unused tail slots on early termination.
225    let body = vec![
226        Node::let_bind("current", Expr::load(target, Expr::u32(0))),
227        Node::let_bind("len", Expr::u32(0)),
228        Node::let_bind("done", Expr::u32(0)),
229        Node::loop_for(
230            "step",
231            Expr::u32(0),
232            Expr::u32(max_depth),
233            vec![Node::if_then(
234                Expr::eq(Expr::var("done"), Expr::u32(0)),
235                vec![
236                    Node::store(path_out, Expr::var("len"), Expr::var("current")),
237                    Node::assign("len", Expr::add(Expr::var("len"), Expr::u32(1))),
238                    Node::let_bind(
239                        "next",
240                        Expr::select(
241                            Expr::lt(Expr::var("current"), Expr::buf_len(parent)),
242                            Expr::load(parent, Expr::var("current")),
243                            Expr::var("current"),
244                        ),
245                    ),
246                    Node::if_then(
247                        Expr::eq(Expr::var("next"), Expr::var("current")),
248                        vec![Node::assign("done", Expr::u32(1))],
249                    ),
250                    Node::assign("current", Expr::var("next")),
251                ],
252            )],
253        ),
254        // Zero-fill path_out[len..max_depth] so harness byte-compare
255        // matches cpu_ref tail-padding convention.
256        Node::loop_for(
257            "pad",
258            Expr::var("len"),
259            Expr::u32(max_depth),
260            vec![Node::store(path_out, Expr::var("pad"), Expr::u32(0))],
261        ),
262        Node::store(path_len, Expr::u32(0), Expr::var("len")),
263    ];
264
265    Program::wrapped(
266        vec![
267            BufferDecl::storage(parent, 0, BufferAccess::ReadOnly, DataType::U32),
268            BufferDecl::storage(target, 1, BufferAccess::ReadOnly, DataType::U32).with_count(1),
269            BufferDecl::storage(path_out, 2, BufferAccess::ReadWrite, DataType::U32)
270                .with_count(max_depth),
271            BufferDecl::storage(path_len, 3, BufferAccess::ReadWrite, DataType::U32).with_count(1),
272        ],
273        [1, 1, 1],
274        vec![Node::Region {
275            generator: Ident::from(OP_ID),
276            source_region: None,
277            body: Arc::new(vec![Node::if_then(
278                Expr::eq(Expr::InvocationId { axis: 0 }, Expr::u32(0)),
279                body,
280            )]),
281        }],
282    )
283}
284
285/// Build the IR `Program` for batched path reconstruction.
286///
287/// Each invocation reconstructs one target's parent chain and writes
288/// a `max_depth`-padded slice at `paths[target_index * max_depth..]`.
289/// `lens[target_index]` receives the valid path length before padding.
290///
291/// # Contract
292///
293/// The caller supplies:
294/// - `parent`: parent-pointer array.
295/// - `targets`: `target_count` target nodes.
296/// - `paths`: `target_count * max_depth` u32 slots.
297/// - `lens`: `target_count` u32 slots.
298///
299/// `max_depth == 0` or `target_count * max_depth` overflow produces a trap
300/// program rather than silently emitting malformed buffer metadata.
301#[must_use]
302pub fn batched_path_reconstruct(target_count: u32, max_depth: u32) -> Program {
303    let layout = match validate_batched_path_reconstruct_layout(target_count as usize, max_depth) {
304        Ok(layout) => layout,
305        Err(error) => {
306            return crate::invalid_output_program(BATCHED_OP_ID, "paths", DataType::U32, error);
307        }
308    };
309    let path_words = layout.path_words_u32;
310
311    let body = vec![
312        Node::let_bind("idx", Expr::InvocationId { axis: 0 }),
313        Node::if_then(
314            Expr::lt(Expr::var("idx"), Expr::u32(target_count)),
315            vec![
316                Node::let_bind("base", Expr::mul(Expr::var("idx"), Expr::u32(max_depth))),
317                Node::let_bind("current", Expr::load("targets", Expr::var("idx"))),
318                Node::let_bind("len", Expr::u32(0)),
319                Node::let_bind("done", Expr::u32(0)),
320                Node::loop_for(
321                    "step",
322                    Expr::u32(0),
323                    Expr::u32(max_depth),
324                    vec![Node::if_then(
325                        Expr::eq(Expr::var("done"), Expr::u32(0)),
326                        vec![
327                            Node::store(
328                                "paths",
329                                Expr::add(Expr::var("base"), Expr::var("len")),
330                                Expr::var("current"),
331                            ),
332                            Node::assign("len", Expr::add(Expr::var("len"), Expr::u32(1))),
333                            Node::let_bind(
334                                "next",
335                                Expr::select(
336                                    Expr::lt(Expr::var("current"), Expr::buf_len("parent")),
337                                    Expr::load("parent", Expr::var("current")),
338                                    Expr::var("current"),
339                                ),
340                            ),
341                            Node::if_then(
342                                Expr::eq(Expr::var("next"), Expr::var("current")),
343                                vec![Node::assign("done", Expr::u32(1))],
344                            ),
345                            Node::assign("current", Expr::var("next")),
346                        ],
347                    )],
348                ),
349                Node::loop_for(
350                    "pad",
351                    Expr::var("len"),
352                    Expr::u32(max_depth),
353                    vec![Node::store(
354                        "paths",
355                        Expr::add(Expr::var("base"), Expr::var("pad")),
356                        Expr::u32(0),
357                    )],
358                ),
359                Node::store("lens", Expr::var("idx"), Expr::var("len")),
360            ],
361        ),
362    ];
363
364    Program::wrapped(
365        vec![
366            BufferDecl::storage("parent", 0, BufferAccess::ReadOnly, DataType::U32),
367            BufferDecl::storage("targets", 1, BufferAccess::ReadOnly, DataType::U32)
368                .with_count(target_count),
369            BufferDecl::storage("paths", 2, BufferAccess::ReadWrite, DataType::U32)
370                .with_count(path_words),
371            BufferDecl::storage("lens", 3, BufferAccess::ReadWrite, DataType::U32)
372                .with_count(target_count),
373        ],
374        [BATCHED_WORKGROUP_SIZE, 1, 1],
375        vec![Node::Region {
376            generator: Ident::from(BATCHED_OP_ID),
377            source_region: None,
378            body: Arc::new(body),
379        }],
380    )
381}
382
383/// Validate batched path-reconstruction output layout.
384///
385/// # Errors
386///
387/// Returns an actionable diagnostic when `max_depth` is zero, target count
388/// exceeds the primitive's u32 lane space, or `target_count * max_depth`
389/// cannot be represented by primitive buffer metadata.
390pub fn validate_batched_path_reconstruct_layout(
391    target_len: usize,
392    max_depth: u32,
393) -> Result<BatchedPathReconstructLayout, String> {
394    if max_depth == 0 {
395        return Err("Fix: batched_path_reconstruct max_depth must be >= 1.".to_string());
396    }
397    let target_count = u32::try_from(target_len).map_err(|_| {
398        format!(
399            "Fix: batched_path_reconstruct target count {target_len} exceeds the primitive u32 lane limit."
400        )
401    })?;
402    let path_words_u32 = target_count.checked_mul(max_depth).ok_or_else(|| {
403        format!(
404            "Fix: batched_path_reconstruct target_count*max_depth overflows u32 for target_count={target_count}, max_depth={max_depth}."
405        )
406    })?;
407    let path_words = usize::try_from(path_words_u32).map_err(|_| {
408        format!("Fix: batched_path_reconstruct path word count {path_words_u32} exceeds usize.")
409    })?;
410    Ok(BatchedPathReconstructLayout {
411        target_count,
412        path_words,
413        path_words_u32,
414    })
415}
416
417/// Validate a single-target reconstruction dispatch and return its buffer plan.
418pub fn plan_path_reconstruct_dispatch(
419    parent_len: usize,
420    max_depth: u32,
421) -> Result<PathReconstructDispatchPlan, String> {
422    if max_depth == 0 {
423        return Err("Fix: path_reconstruct max_depth must be >= 1.".to_string());
424    }
425    Ok(PathReconstructDispatchPlan {
426        parent_words: parent_len,
427        target_words: 1,
428        path_words: max_depth as usize,
429        len_words: 1,
430        max_depth,
431        grid: PATH_RECONSTRUCT_DISPATCH_GRID,
432    })
433}
434
435/// Validate a batched path-reconstruction dispatch and return its buffer plan.
436pub fn plan_batched_path_reconstruct_dispatch(
437    parent_len: usize,
438    target_len: usize,
439    max_depth: u32,
440) -> Result<BatchedPathReconstructDispatchPlan, String> {
441    let layout = validate_batched_path_reconstruct_layout(target_len, max_depth)?;
442    Ok(BatchedPathReconstructDispatchPlan {
443        parent_words: parent_len,
444        target_words: target_len,
445        path_words: layout.path_words,
446        len_words: target_len,
447        max_depth,
448        grid: [
449            ceil_div_u32(layout.target_count, BATCHED_WORKGROUP_SIZE),
450            1,
451            1,
452        ],
453        layout,
454    })
455}
456
457fn ceil_div_u32(value: u32, divisor: u32) -> u32 {
458    if value == 0 {
459        0
460    } else {
461        ((value - 1) / divisor) + 1
462    }
463}
464
465/// Validate the length word returned by a single-target path-reconstruction dispatch.
466///
467/// Backend wrappers use this after readback so malformed GPU output is checked
468/// by the primitive owner rather than by each consumer.
469pub fn validate_path_reconstruct_readback(
470    plan: &PathReconstructDispatchPlan,
471    len: u32,
472) -> Result<usize, String> {
473    let len_usize = usize::try_from(len).map_err(|_| {
474        format!("Fix: path_reconstruct returned length {len}, which cannot fit this host usize.")
475    })?;
476    if len_usize > plan.path_words {
477        return Err(format!(
478            "Fix: path_reconstruct returned length {len}, exceeding max_depth {}. Treat this as malformed GPU readback or a backend bug.",
479            plan.max_depth
480        ));
481    }
482    Ok(len_usize)
483}
484
485/// Validate output-buffer shapes and length words returned by a batched
486/// path-reconstruction dispatch.
487///
488/// This is deliberately primitive-owned: dispatch wrappers may stage buffers
489/// differently, but the post-dispatch correctness envelope is part of the
490/// graph primitive contract.
491pub fn validate_batched_path_reconstruct_readback(
492    plan: &BatchedPathReconstructDispatchPlan,
493    path_words: usize,
494    len_words: usize,
495    lens: &[u32],
496) -> Result<(), String> {
497    if path_words != plan.path_words || len_words != plan.len_words {
498        return Err(format!(
499            "Fix: batched_path_reconstruct returned {path_words} path word(s) and {len_words} len word(s), expected {} and {}.",
500            plan.path_words, plan.len_words
501        ));
502    }
503    for (target_index, &len) in lens.iter().enumerate() {
504        let len_usize = usize::try_from(len).map_err(|_| {
505            format!(
506                "Fix: batched_path_reconstruct target {target_index} returned length {len}, which cannot fit this host usize."
507            )
508        })?;
509        if len_usize > plan.max_depth as usize {
510            return Err(format!(
511                "Fix: batched_path_reconstruct target {target_index} returned length {len}, exceeding max_depth {}. Treat this as malformed GPU readback or a backend bug.",
512                plan.max_depth
513            ));
514        }
515    }
516    Ok(())
517}
518
519#[cfg(test)]
520mod dispatch_plan_tests {
521    use super::*;
522
523    #[test]
524    fn single_path_dispatch_plan_owns_outputs_and_grid() {
525        let plan = plan_path_reconstruct_dispatch(4, 8)
526            .expect("Fix: nonzero max_depth should plan single reconstruction");
527
528        assert_eq!(plan.parent_words, 4);
529        assert_eq!(plan.target_words, 1);
530        assert_eq!(plan.path_words, 8);
531        assert_eq!(plan.len_words, 1);
532        assert_eq!(plan.grid, PATH_RECONSTRUCT_DISPATCH_GRID);
533    }
534
535    #[test]
536    fn single_path_dispatch_plan_rejects_zero_depth() {
537        let err = plan_path_reconstruct_dispatch(4, 0).unwrap_err();
538        assert!(err.contains("max_depth"));
539    }
540
541    #[test]
542    fn batched_path_dispatch_plan_owns_layout_and_grid() {
543        let plan = plan_batched_path_reconstruct_dispatch(4, 513, 3)
544            .expect("Fix: valid batched reconstruction should plan");
545
546        assert_eq!(plan.parent_words, 4);
547        assert_eq!(plan.target_words, 513);
548        assert_eq!(plan.path_words, 1539);
549        assert_eq!(plan.len_words, 513);
550        assert_eq!(plan.grid, [3, 1, 1]);
551        assert_eq!(plan.layout.target_count, 513);
552    }
553
554    #[test]
555    fn static_input_key_tracks_parent_content_and_dispatch_shape() {
556        let single = plan_path_reconstruct_dispatch(4, 8)
557            .expect("Fix: nonzero max_depth should plan single reconstruction");
558        let batched = plan_batched_path_reconstruct_dispatch(4, 2, 8)
559            .expect("Fix: valid batched reconstruction should plan");
560
561        let first = single
562            .static_input_key(&[0, 0, 1, 2])
563            .expect("Fix: matching parent slice should key");
564        let same = single
565            .static_input_key(&[0, 0, 1, 2])
566            .expect("Fix: matching parent slice should key");
567        let changed = single
568            .static_input_key(&[0, 0, 0, 2])
569            .expect("Fix: same-shape parent content should key");
570        let batched_key = batched
571            .static_input_key(&[0, 0, 1, 2])
572            .expect("Fix: matching batched parent slice should key");
573
574        assert_eq!(first, same);
575        assert_ne!(first, changed);
576        assert_ne!(first, batched_key);
577        assert_eq!(first.parent_words, 4);
578        assert_eq!(first.target_count, 1);
579        assert!(!first.batched);
580        assert_eq!(batched_key.target_count, 2);
581        assert!(batched_key.batched);
582    }
583
584    #[test]
585    fn static_input_key_rejects_parent_length_drift() {
586        let single = plan_path_reconstruct_dispatch(4, 8)
587            .expect("Fix: nonzero max_depth should plan single reconstruction");
588        let batched = plan_batched_path_reconstruct_dispatch(4, 2, 8)
589            .expect("Fix: valid batched reconstruction should plan");
590
591        let err = single.static_input_key(&[0, 0, 1]).unwrap_err();
592        assert!(err.contains("expected 4 parent word"));
593
594        let err = batched.static_input_key(&[0, 0, 1]).unwrap_err();
595        assert!(err.contains("expected 4 parent word"));
596    }
597
598    #[test]
599    fn single_path_readback_validation_rejects_impossible_len() {
600        let plan = plan_path_reconstruct_dispatch(4, 4)
601            .expect("Fix: nonzero max_depth should plan single reconstruction");
602
603        assert_eq!(validate_path_reconstruct_readback(&plan, 4), Ok(4));
604
605        let err = validate_path_reconstruct_readback(&plan, 5).unwrap_err();
606        assert!(err.contains("exceeding max_depth 4"));
607    }
608
609    #[test]
610    fn batched_path_readback_validation_rejects_shape_and_len_drift() {
611        let plan = plan_batched_path_reconstruct_dispatch(4, 2, 4)
612            .expect("Fix: valid batched reconstruction should plan");
613
614        validate_batched_path_reconstruct_readback(&plan, 8, 2, &[4, 1]).unwrap();
615
616        let err = validate_batched_path_reconstruct_readback(&plan, 7, 2, &[4, 1]).unwrap_err();
617        assert!(err.contains("expected 8 and 2"));
618
619        let err = validate_batched_path_reconstruct_readback(&plan, 8, 2, &[4, 5]).unwrap_err();
620        assert!(err.contains("target 1"));
621        assert!(err.contains("exceeding max_depth 4"));
622    }
623}
624
625/// CPU reference: walks parent pointers up to `max_depth`, writing
626/// the materialized path into `scratch` and returning its length.
627/// Early-terminates when a node's parent points at itself (root
628/// convention).
629///
630/// # Performance
631///
632/// Callers doing many reconstructions (e.g. one per node in a deep
633/// call graph) should reuse one scratch vector across calls to avoid
634/// an allocation per walk.
635#[must_use]
636#[cfg(any(test, feature = "cpu-parity"))]
637pub fn cpu_ref(parent: &[u32], target: u32, max_depth: u32, scratch: &mut Vec<u32>) -> u32 {
638    scratch.clear();
639    let mut current = target;
640    let mut len = 0u32;
641    let cap = max_depth as usize;
642    while (len as usize) < cap {
643        scratch.push(current);
644        len += 1;
645        let next = parent.get(current as usize).copied().unwrap_or(current);
646        if next == current {
647            break;
648        }
649        current = next;
650    }
651    while scratch.len() < cap {
652        scratch.push(0);
653    }
654    len
655}
656
657/// Fallible CPU reference for the batched path-reconstruction contract.
658///
659/// This is the allocation-safe entry point for wrappers that run CPU parity on
660/// hostile batch dimensions. The legacy `cpu_ref_batched` wrapper delegates here
661/// to preserve the existing call surface while avoiding infallible reservations.
662#[cfg(any(test, feature = "cpu-parity"))]
663pub fn try_cpu_ref_batched(
664    parent: &[u32],
665    targets: &[u32],
666    max_depth: u32,
667    paths: &mut Vec<u32>,
668    lens: &mut Vec<u32>,
669) -> Result<(), String> {
670    let mut scratch = Vec::new();
671    try_cpu_ref_batched_with_scratch(parent, targets, max_depth, paths, lens, &mut scratch)
672}
673
674/// Fallible batched CPU reference using caller-owned output and per-target scratch storage.
675#[cfg(any(test, feature = "cpu-parity"))]
676pub fn try_cpu_ref_batched_with_scratch(
677    parent: &[u32],
678    targets: &[u32],
679    max_depth: u32,
680    paths: &mut Vec<u32>,
681    lens: &mut Vec<u32>,
682    scratch: &mut Vec<u32>,
683) -> Result<(), String> {
684    let layout = validate_batched_path_reconstruct_layout(targets.len(), max_depth)?;
685    let depth = max_depth as usize;
686    scratch.clear();
687    crate::graph::scratch::reserve_graph_items(
688        paths,
689        layout.path_words,
690        "path reconstruction CPU oracle",
691        "batched path output",
692    )?;
693    crate::graph::scratch::reserve_graph_items(
694        lens,
695        layout.target_count as usize,
696        "path reconstruction CPU oracle",
697        "batched length output",
698    )?;
699    crate::graph::scratch::reserve_graph_items(
700        scratch,
701        depth,
702        "path reconstruction CPU oracle",
703        "per-target path scratch",
704    )?;
705    paths.clear();
706    lens.clear();
707    for &target in targets {
708        let len = cpu_ref(parent, target, max_depth, scratch);
709        paths.extend_from_slice(&scratch);
710        lens.push(len);
711    }
712    Ok(())
713}
714
715/// CPU reference for the batched path-reconstruction contract.
716///
717/// `paths` is rewritten to `targets.len() * max_depth` entries, where each
718/// target owns one zero-padded segment. `lens` is rewritten to one valid
719/// length per target.
720#[cfg(any(test, feature = "cpu-parity"))]
721pub fn cpu_ref_batched(
722    parent: &[u32],
723    targets: &[u32],
724    max_depth: u32,
725    paths: &mut Vec<u32>,
726    lens: &mut Vec<u32>,
727) {
728    try_cpu_ref_batched(parent, targets, max_depth, paths, lens)
729        .expect("Fix: batched path reconstruction CPU oracle allocation failed");
730}
731
732#[cfg(feature = "inventory-registry")]
733inventory::submit! {
734    vyre_foundation::operation::OperationRegistration::primitive(
735        OP_ID,
736        || path_reconstruct("parent", "target", "path_out", "path_len", 4),
737        Some(|| {
738            let to_bytes = |w: &[u32]| crate::wire::pack_u32_slice(w);
739            // parent: [0, 0, 1, 2]  (0 is root; 1→0, 2→1, 3→2)
740            // target = 3
741            // expected path = [3, 2, 1, 0], len = 4
742            vec![vec![
743                to_bytes(&[0, 0, 1, 2]),
744                to_bytes(&[3]),
745                to_bytes(&[0, 0, 0, 0]),
746                to_bytes(&[0]),
747            ]]
748        }),
749        Some(|| {
750            let to_bytes = |w: &[u32]| crate::wire::pack_u32_slice(w);
751            vec![vec![
752                to_bytes(&[3, 2, 1, 0]),
753                to_bytes(&[4]),
754            ]]
755        }),
756    )
757}
758
759#[cfg(test)]
760mod tests {
761    use super::*;
762
763    #[test]
764    fn walks_parent_chain_to_root() {
765        let mut scratch = Vec::with_capacity(4);
766        let len = cpu_ref(&[0, 0, 1, 2], 3, 4, &mut scratch);
767        assert_eq!(len, 4);
768        assert_eq!(&scratch[0..4], &[3, 2, 1, 0]);
769    }
770
771    #[test]
772    fn terminates_on_max_depth() {
773        // Cycle: 0 → 1 → 0. Without max_depth we'd loop forever.
774        // AUDIT_2026-04-24 F-PR-02: also assert path contents so a
775        // silent buffer corruption cannot slip past the test.
776        let mut scratch = Vec::with_capacity(8);
777        let len = cpu_ref(&[1, 0], 0, 8, &mut scratch);
778        assert_eq!(len, 8);
779        assert_eq!(&scratch[..], &[0, 1, 0, 1, 0, 1, 0, 1]);
780    }
781
782    #[test]
783    fn tail_is_zero_padded_when_root_reached_before_cap() {
784        // AUDIT_2026-04-24 F-PR-01: cpu_ref must zero-fill the tail
785        // beyond the materialized path so harness byte-compare with
786        // the IR builder stays stable.
787        let mut scratch = Vec::with_capacity(8);
788        let len = cpu_ref(&[0, 0, 1, 2], 3, 8, &mut scratch);
789        assert_eq!(len, 4);
790        assert_eq!(&scratch[..4], &[3, 2, 1, 0]);
791        assert_eq!(&scratch[4..], &[0, 0, 0, 0]);
792    }
793
794    // ------------------------------------------------------------------
795    // Adversarial fixtures  -  self-loops, deep chain, OOB target, max_depth.
796    // ------------------------------------------------------------------
797
798    #[test]
799    fn parent_self_loops_terminate_immediately() {
800        // parent[0]=0, parent[1]=1. target=1 → path [1].
801        let mut scratch = Vec::with_capacity(4);
802        let len = cpu_ref(&[0, 1], 1, 4, &mut scratch);
803        assert_eq!(len, 1);
804        assert_eq!(scratch[0], 1);
805        assert_eq!(&scratch[1..], &[0, 0, 0]);
806    }
807
808    #[test]
809    fn deep_chain_within_max_depth() {
810        // 0 ← 1 ← 2 ← 3 ← 4
811        let parent = &[0, 0, 1, 2, 3];
812        let mut scratch = Vec::with_capacity(8);
813        let len = cpu_ref(parent, 4, 8, &mut scratch);
814        assert_eq!(len, 5);
815        assert_eq!(&scratch[..5], &[4, 3, 2, 1, 0]);
816        assert_eq!(&scratch[5..], &[0, 0, 0]);
817    }
818
819    #[test]
820    fn target_not_in_parent_array_terminates_at_target() {
821        // parent has 3 entries. target=5 is OOB.
822        let mut scratch = Vec::with_capacity(4);
823        let len = cpu_ref(&[0, 0, 1], 5, 4, &mut scratch);
824        assert_eq!(len, 1);
825        assert_eq!(scratch[0], 5);
826        assert_eq!(&scratch[1..], &[0, 0, 0]);
827    }
828
829    #[test]
830    fn max_depth_zero_returns_empty_path() {
831        let mut scratch = Vec::with_capacity(4);
832        let len = cpu_ref(&[0, 0, 1, 2], 3, 0, &mut scratch);
833        assert_eq!(len, 0);
834        assert!(scratch.is_empty());
835    }
836
837    #[test]
838    fn max_depth_one_returns_only_target() {
839        let mut scratch = Vec::with_capacity(4);
840        let len = cpu_ref(&[0, 0, 1, 2], 3, 1, &mut scratch);
841        assert_eq!(len, 1);
842        assert_eq!(scratch[0], 3);
843        assert_eq!(scratch.len(), 1, "cap == max_depth == 1, no padding needed");
844    }
845
846    #[test]
847    fn program_builder_max_depth_zero_emits_trap() {
848        let p = path_reconstruct("parent", "target", "out", "len", 0);
849        // Trap programs contain a Node::Trap in the entry region body.
850        let entry = p.entry();
851        let has_trap = entry.iter().any(|n| {
852            if let Node::Region { body, .. } = n {
853                body.iter().any(|inner| matches!(inner, Node::Trap { .. }))
854            } else {
855                matches!(n, Node::Trap { .. })
856            }
857        });
858        assert!(
859            has_trap,
860            "max_depth == 0 must produce a trap program, not panic"
861        );
862    }
863
864    #[test]
865    fn batched_program_has_expected_buffers_and_workgroup() {
866        let p = batched_path_reconstruct(3, 4);
867        assert_eq!(p.workgroup_size, [BATCHED_WORKGROUP_SIZE, 1, 1]);
868        let names: Vec<&str> = p.buffers.iter().map(|b| b.name()).collect();
869        assert_eq!(names, vec!["parent", "targets", "paths", "lens"]);
870        assert_eq!(p.buffers[1].count(), 3);
871        assert_eq!(p.buffers[2].count(), 12);
872        assert_eq!(p.buffers[3].count(), 3);
873    }
874
875    #[test]
876    fn batched_layout_validator_accepts_empty_and_canonical_batches() {
877        assert_eq!(
878            validate_batched_path_reconstruct_layout(0, 4).unwrap(),
879            BatchedPathReconstructLayout {
880                target_count: 0,
881                path_words: 0,
882                path_words_u32: 0,
883            }
884        );
885        assert_eq!(
886            validate_batched_path_reconstruct_layout(3, 4).unwrap(),
887            BatchedPathReconstructLayout {
888                target_count: 3,
889                path_words: 12,
890                path_words_u32: 12,
891            }
892        );
893    }
894
895    #[test]
896    fn batched_layout_validator_rejects_zero_depth_and_overflow() {
897        let err = validate_batched_path_reconstruct_layout(3, 0).unwrap_err();
898        assert!(err.contains("max_depth must be >= 1"));
899
900        let err = validate_batched_path_reconstruct_layout(u32::MAX as usize + 1, 1).unwrap_err();
901        assert!(err.contains("target count"));
902
903        let err = validate_batched_path_reconstruct_layout(u32::MAX as usize, 2).unwrap_err();
904        assert!(err.contains("target_count*max_depth"));
905    }
906
907    #[test]
908    fn batched_cpu_ref_matches_single_target_segments() {
909        let mut paths = Vec::new();
910        let mut lens = Vec::new();
911        cpu_ref_batched(&[0, 0, 1, 2], &[3, 0, 2], 4, &mut paths, &mut lens);
912        assert_eq!(lens, vec![4, 1, 3]);
913        assert_eq!(&paths[0..4], &[3, 2, 1, 0]);
914        assert_eq!(&paths[4..8], &[0, 0, 0, 0]);
915        assert_eq!(&paths[8..12], &[2, 1, 0, 0]);
916    }
917
918    #[test]
919    fn batched_cpu_ref_with_scratch_reuses_all_storage() {
920        let mut paths = Vec::with_capacity(32);
921        let mut lens = Vec::with_capacity(8);
922        let mut scratch = Vec::with_capacity(8);
923        paths.extend_from_slice(&[0xDEAD_BEEF; 5]);
924        lens.extend_from_slice(&[0xCAFE_BABE; 3]);
925        scratch.extend_from_slice(&[0xFEED_FACE; 6]);
926        let paths_capacity = paths.capacity();
927        let lens_capacity = lens.capacity();
928        let scratch_capacity = scratch.capacity();
929
930        try_cpu_ref_batched_with_scratch(
931            &[0, 0, 1, 2],
932            &[3, 0, 2],
933            4,
934            &mut paths,
935            &mut lens,
936            &mut scratch,
937        )
938        .expect("Fix: valid batched path reconstruction must evaluate.");
939
940        assert_eq!(lens, vec![4, 1, 3]);
941        assert_eq!(&paths[0..4], &[3, 2, 1, 0]);
942        assert_eq!(&paths[4..8], &[0, 0, 0, 0]);
943        assert_eq!(&paths[8..12], &[2, 1, 0, 0]);
944        assert_eq!(scratch, vec![2, 1, 0, 0]);
945        assert_eq!(paths.capacity(), paths_capacity);
946        assert_eq!(lens.capacity(), lens_capacity);
947        assert_eq!(scratch.capacity(), scratch_capacity);
948
949        try_cpu_ref_batched_with_scratch(
950            &[0, 0, 1, 2],
951            &[1],
952            2,
953            &mut paths,
954            &mut lens,
955            &mut scratch,
956        )
957        .expect("Fix: second valid batch must reuse and truncate buffers.");
958
959        assert_eq!(lens, vec![2]);
960        assert_eq!(paths, vec![1, 0]);
961        assert_eq!(scratch, vec![1, 0]);
962        assert_eq!(paths.capacity(), paths_capacity);
963        assert_eq!(lens.capacity(), lens_capacity);
964        assert_eq!(scratch.capacity(), scratch_capacity);
965    }
966
967    #[test]
968    fn batched_cpu_ref_rejects_zero_depth_like_dispatch_planner() {
969        let mut paths = vec![0xDEAD_BEEF];
970        let mut lens = vec![0xCAFE_BABE];
971        let err = try_cpu_ref_batched(&[0], &[0], 0, &mut paths, &mut lens).unwrap_err();
972
973        assert!(err.contains("max_depth must be >= 1"));
974        assert_eq!(paths, vec![0xDEAD_BEEF]);
975        assert_eq!(lens, vec![0xCAFE_BABE]);
976    }
977
978    #[test]
979    fn generated_batched_cpu_ref_matches_single_target_oracle_shapes() {
980        for target_count in 0usize..64 {
981            for depth in 1u32..65 {
982                let parent: Vec<u32> = (0..128u32)
983                    .map(|node| if node == 0 { 0 } else { node - 1 })
984                    .collect();
985                let targets: Vec<u32> = (0..target_count)
986                    .map(|index| ((index * 17 + depth as usize * 3) % parent.len()) as u32)
987                    .collect();
988                let mut paths = vec![0xDEAD_BEEFu32; 3];
989                let mut lens = vec![0xCAFE_BABEu32; 2];
990                try_cpu_ref_batched(&parent, &targets, depth, &mut paths, &mut lens).unwrap();
991                assert_eq!(lens.len(), targets.len());
992                assert_eq!(paths.len(), targets.len() * depth as usize);
993                let mut single = Vec::new();
994                for (target_index, &target) in targets.iter().enumerate() {
995                    let expected_len = cpu_ref(&parent, target, depth, &mut single);
996                    assert_eq!(lens[target_index], expected_len);
997                    let start = target_index * depth as usize;
998                    let end = start + depth as usize;
999                    assert_eq!(&paths[start..end], &single[..]);
1000                }
1001            }
1002        }
1003    }
1004
1005    #[test]
1006    fn batched_program_zero_depth_emits_trap() {
1007        let p = batched_path_reconstruct(3, 0);
1008        let entry = p.entry();
1009        let has_trap = entry.iter().any(|n| {
1010            if let Node::Region { body, .. } = n {
1011                body.iter().any(|inner| matches!(inner, Node::Trap { .. }))
1012            } else {
1013                matches!(n, Node::Trap { .. })
1014            }
1015        });
1016        assert!(has_trap, "zero-depth batched path reconstruction must trap");
1017    }
1018}