1use std::sync::Arc;
17
18use vyre_foundation::ir::model::expr::Ident;
19use vyre_foundation::ir::{BufferAccess, BufferDecl, DataType, Expr, Node, Program};
20
21pub const OP_ID: &str = "vyre-primitives::graph::path_reconstruct";
23
24pub const BATCHED_OP_ID: &str = "vyre-primitives::graph::batched_path_reconstruct";
26
27pub const BATCHED_WORKGROUP_SIZE: u32 = 256;
29pub const PATH_PARENT_BUFFER: &str = "path_reconstruct parent";
31pub const PATH_TARGET_BUFFER: &str = "path_reconstruct target";
33pub const PATH_OUT_BUFFER: &str = "path_reconstruct path_out";
35pub const PATH_LEN_BUFFER: &str = "path_reconstruct path_len";
37pub const BATCHED_PATH_TARGETS_BUFFER: &str = "batched_path_reconstruct targets";
39pub const BATCHED_PATHS_BUFFER: &str = "batched_path_reconstruct paths";
41pub const BATCHED_LENS_BUFFER: &str = "batched_path_reconstruct lens";
43pub const PATH_RECONSTRUCT_DISPATCH_GRID: [u32; 3] = [1, 1, 1];
45
46#[derive(Clone, Copy, Debug, Eq, PartialEq)]
48pub struct BatchedPathReconstructLayout {
49 pub target_count: u32,
51 pub path_words: usize,
53 pub path_words_u32: u32,
55}
56
57#[derive(Clone, Debug, Eq, PartialEq)]
59pub struct PathReconstructDispatchPlan {
60 pub parent_words: usize,
62 pub target_words: usize,
64 pub path_words: usize,
66 pub len_words: usize,
68 pub max_depth: u32,
70 pub grid: [u32; 3],
72}
73
74impl PathReconstructDispatchPlan {
75 #[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 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#[derive(Clone, Debug, Eq, PartialEq)]
116pub struct BatchedPathReconstructDispatchPlan {
117 pub layout: BatchedPathReconstructLayout,
119 pub parent_words: usize,
121 pub target_words: usize,
123 pub path_words: usize,
125 pub len_words: usize,
127 pub max_depth: u32,
129 pub grid: [u32; 3],
131}
132
133#[derive(Clone, Copy, Debug, Eq, PartialEq)]
139pub struct PathReconstructStaticInputKey {
140 pub parent_words: usize,
142 pub parent_hash: u64,
144 pub target_count: u32,
146 pub max_depth: u32,
148 pub batched: bool,
150}
151
152impl BatchedPathReconstructDispatchPlan {
153 #[must_use]
155 pub fn program(&self) -> Program {
156 batched_path_reconstruct(self.layout.target_count, self.max_depth)
157 }
158
159 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#[must_use]
192pub 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 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 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#[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
383pub 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
417pub 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
435pub 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
465pub 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
485pub 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#[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#[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#[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#[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 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 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 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 #[test]
799 fn parent_self_loops_terminate_immediately() {
800 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 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 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 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}