1use vyre_foundation::ir::model::expr::GeneratorRef;
21use vyre_foundation::ir::{BufferDecl, DataType, Expr, Node, Program};
22
23use crate::tensor_ref::{TensorRef, TensorRefError};
24
25pub(crate) const INDEXED_MAP_OP_ID: &str = "vyre-libs::substrate::indexed_map";
31pub(crate) const STRIDED_ACCUMULATE_OP_ID: &str = "vyre-libs::substrate::strided_accumulate";
33pub(crate) const STRIDED_WRITEBACK_OP_ID: &str =
35 "anonymous::vyre-libs::substrate::strided_writeback";
36
37#[derive(Debug, Clone, Default)]
40#[non_exhaustive]
41pub struct BuildOptions {
42 pub workgroup_size: Option<[u32; 3]>,
44 pub region_generator: Option<&'static str>,
48 pub tenant_id: Option<u32>,
52}
53
54impl BuildOptions {
55 #[must_use]
57 pub fn new() -> Self {
58 Self::default()
59 }
60
61 #[must_use]
63 pub fn with_workgroup_size(mut self, size: [u32; 3]) -> Self {
64 self.workgroup_size = Some(size);
65 self
66 }
67
68 #[must_use]
70 pub fn with_region_generator(mut self, name: &'static str) -> Self {
71 self.region_generator = Some(name);
72 self
73 }
74
75 #[must_use]
77 pub fn with_tenant_id(mut self, tenant_id: u32) -> Self {
78 self.tenant_id = Some(tenant_id);
79 self
80 }
81}
82
83macro_rules! impl_cat_a_builder_options {
84 ($builder:ident) => {
85 impl $builder {
86 #[must_use]
88 pub fn with_workgroup_size(mut self, size: [u32; 3]) -> Self {
89 self.options = self.options.with_workgroup_size(size);
90 self
91 }
92
93 #[must_use]
95 pub fn with_region_generator(mut self, name: &'static str) -> Self {
96 self.options = self.options.with_region_generator(name);
97 self
98 }
99
100 #[must_use]
102 pub fn with_tenant_id(mut self, tenant_id: u32) -> Self {
103 self.options = self.options.with_tenant_id(tenant_id);
104 self
105 }
106 }
107 };
108}
109
110pub(crate) use impl_cat_a_builder_options;
111
112pub fn check_tensors(
116 op: &'static str,
117 tensors: &[(&TensorRef, DataType)],
118) -> Result<(), TensorRefError> {
119 for (r, expected) in tensors {
121 crate::tensor_ref::check_dtype(r, expected.clone(), op)?;
122 if r.element_count().is_none() {
123 return Err(TensorRefError::ElementCountOverflow {
124 name: r.name.as_str().to_string(),
125 shape: r.shape.to_vec(),
126 });
127 }
128 }
129 for (idx, (left, _)) in tensors.iter().enumerate() {
130 for (right, _) in &tensors[idx + 1..] {
131 if left.name_str() == right.name_str() {
132 return Err(TensorRefError::NameCollision {
133 name: left.name.as_str().to_string(),
134 op,
135 });
136 }
137 }
138 }
139 Ok(())
140}
141
142#[cfg(test)]
143mod cat_a_builder_option_macro_tests {
144 #![allow(unreachable_pub)]
145
146 use super::BuildOptions;
147
148 #[derive(Debug, Clone)]
149 struct DemoBuilder {
150 options: BuildOptions,
151 }
152
153 impl DemoBuilder {
154 fn new() -> Self {
155 Self {
156 options: BuildOptions::default(),
157 }
158 }
159 }
160
161 super::impl_cat_a_builder_options!(DemoBuilder);
162
163 #[test]
164 fn generated_option_surface_threads_every_shared_knob() {
165 let builder = DemoBuilder::new()
166 .with_workgroup_size([8, 4, 2])
167 .with_region_generator("custom::generator")
168 .with_tenant_id(17);
169
170 assert_eq!(builder.options.workgroup_size, Some([8, 4, 2]));
171 assert_eq!(builder.options.region_generator, Some("custom::generator"));
172 assert_eq!(builder.options.tenant_id, Some(17));
173 }
174}
175
176pub(crate) fn build_indexed_map<F>(
182 op_id: &'static str,
183 buffers: Vec<BufferDecl>,
184 output: &str,
185 count: u32,
186 workgroup_size: [u32; 3],
187 f: F,
188) -> Program
189where
190 F: FnOnce(Expr) -> (Expr, Expr),
191{
192 let i = Expr::var("i");
193 let (dst_index, value) = f(i.clone());
194 let child_body = vec![
195 Node::let_bind("i", Expr::InvocationId { axis: 0 }),
196 Node::if_then(
197 Expr::lt(i, Expr::u32(count)),
198 vec![Node::store(output, dst_index, value)],
199 ),
200 ];
201 let parent = GeneratorRef {
202 name: op_id.to_string(),
203 };
204
205 Program::wrapped(
206 buffers,
207 workgroup_size,
208 vec![crate::region::wrap_anonymous(
209 op_id,
210 vec![crate::region::wrap_child(
211 INDEXED_MAP_OP_ID,
212 parent,
213 child_body,
214 )],
215 )],
216 )
217}
218
219pub(crate) fn strided_accumulate_child<F>(
225 parent_op_id: &'static str,
226 tile: u32,
227 chunks: u32,
228 n: u32,
229 acc_name: &'static str,
230 initial: Expr,
231 scratch: &'static str,
232 step: F,
233) -> Node
234where
235 F: Fn(Expr, Expr) -> Expr,
236{
237 let local = Expr::var("local");
238 let idx = Expr::var("idx");
239 let acc = Expr::var(acc_name);
240 let child_body = vec![Node::if_then(
241 Expr::is_first_workgroup(),
242 vec![
243 Node::let_bind(acc_name, initial),
244 strided_loop(
245 tile,
246 chunks,
247 n,
248 vec![Node::assign(acc_name, step(idx, acc))],
249 ),
250 Node::store(scratch, local, Expr::var(acc_name)),
251 ],
252 )];
253
254 child_region(parent_op_id, STRIDED_ACCUMULATE_OP_ID, child_body)
255}
256
257#[allow(dead_code)]
262pub(crate) fn strided_accumulate2_child<F1, F2>(
263 parent_op_id: &'static str,
264 tile: u32,
265 chunks: u32,
266 n: u32,
267 first: (&'static str, Expr, &'static str, F1),
268 second: (&'static str, Expr, &'static str, F2),
269) -> Node
270where
271 F1: Fn(Expr, Expr) -> Expr,
272 F2: Fn(Expr, Expr) -> Expr,
273{
274 let (first_name, first_initial, first_scratch, first_step) = first;
275 let (second_name, second_initial, second_scratch, second_step) = second;
276 let local = Expr::var("local");
277 let idx = Expr::var("idx");
278 let child_body = vec![Node::if_then(
279 Expr::is_first_workgroup(),
280 vec![
281 Node::let_bind(first_name, first_initial),
282 Node::let_bind(second_name, second_initial),
283 strided_loop(
284 tile,
285 chunks,
286 n,
287 vec![
288 Node::assign(first_name, first_step(idx.clone(), Expr::var(first_name))),
289 Node::assign(second_name, second_step(idx, Expr::var(second_name))),
290 ],
291 ),
292 Node::store(first_scratch, local.clone(), Expr::var(first_name)),
293 Node::store(second_scratch, local, Expr::var(second_name)),
294 ],
295 )];
296
297 child_region(parent_op_id, STRIDED_ACCUMULATE_OP_ID, child_body)
298}
299
300pub(crate) fn strided_writeback_child<F>(
306 parent_op_id: &'static str,
307 tile: u32,
308 chunks: u32,
309 n: u32,
310 output: &str,
311 prelude: Vec<Node>,
312 value: F,
313) -> Node
314where
315 F: Fn(Expr) -> Expr,
316{
317 let idx = Expr::var("idx");
318 let mut guarded = prelude;
319 guarded.push(strided_loop(
320 tile,
321 chunks,
322 n,
323 vec![Node::store(output, idx.clone(), value(idx))],
324 ));
325 child_region(
326 parent_op_id,
327 STRIDED_WRITEBACK_OP_ID,
328 vec![Node::if_then(Expr::is_first_workgroup(), guarded)],
329 )
330}
331
332fn strided_loop(tile: u32, chunks: u32, n: u32, guarded_body: Vec<Node>) -> Node {
333 Node::loop_for(
334 "chunk",
335 Expr::u32(0),
336 Expr::u32(chunks),
337 vec![
338 Node::let_bind(
339 "idx",
340 Expr::add(
341 Expr::mul(Expr::var("chunk"), Expr::u32(tile)),
342 Expr::var("local"),
343 ),
344 ),
345 Node::if_then(Expr::lt(Expr::var("idx"), Expr::u32(n)), guarded_body),
346 ],
347 )
348}
349
350fn child_region(parent_op_id: &'static str, child_op_id: &'static str, body: Vec<Node>) -> Node {
351 crate::region::wrap_child(
352 child_op_id,
353 GeneratorRef {
354 name: parent_op_id.to_string(),
355 },
356 body,
357 )
358}
359
360pub(crate) fn invalid_builder_trap_program(
366 op_id: &'static str,
367 output: &str,
368 data_type: DataType,
369 message: String,
370) -> Program {
371 Program::wrapped(
372 vec![BufferDecl::output(output, 0, data_type).with_count(1)],
373 [1, 1, 1],
374 vec![crate::region::wrap_anonymous(
375 op_id,
376 vec![Node::trap(Expr::u32(0), message)],
377 )],
378 )
379}
380
381#[allow(dead_code)]
384pub(crate) fn build_elementwise_binary<F>(
385 op_id: &'static str,
386 a: crate::tensor_ref::TensorRef,
387 b: crate::tensor_ref::TensorRef,
388 out: crate::tensor_ref::TensorRef,
389 options: BuildOptions,
390 f: F,
391) -> Result<vyre_foundation::ir::Program, crate::tensor_ref::TensorRefError>
392where
393 F: Fn(vyre_foundation::ir::Expr, vyre_foundation::ir::Expr) -> vyre_foundation::ir::Expr,
394{
395 check_tensors(
396 op_id,
397 &[
398 (&a, vyre_foundation::ir::DataType::U32),
399 (&b, vyre_foundation::ir::DataType::U32),
400 (&out, vyre_foundation::ir::DataType::U32),
401 ],
402 )?;
403
404 if a.shape != b.shape || a.shape != out.shape {
405 return Err(crate::tensor_ref::TensorRefError::ShapeMismatch {
406 name: "elementwise_binary".into(),
407 found: vec![],
408 expected: vec![],
409 op: op_id,
410 });
411 }
412
413 let a_count = a.element_count().ok_or_else(|| {
414 crate::tensor_ref::TensorRefError::ElementCountOverflow {
415 name: a.name_str().to_string(),
416 shape: a.shape.to_vec(),
417 }
418 })?;
419 let out_count = out.element_count().ok_or_else(|| {
420 crate::tensor_ref::TensorRefError::ElementCountOverflow {
421 name: out.name_str().to_string(),
422 shape: out.shape.to_vec(),
423 }
424 })?;
425 if out_count < a_count {
426 return Err(crate::tensor_ref::TensorRefError::ShapeMismatch {
427 name: out.name_str().to_string(),
428 found: out.shape.to_vec(),
429 expected: a.shape.to_vec(),
430 op: op_id,
431 });
432 }
433
434 let n = a_count;
435 let body = vec![
436 vyre_foundation::ir::Node::let_bind(
437 "idx",
438 vyre_foundation::ir::Expr::InvocationId { axis: 0 },
439 ),
440 vyre_foundation::ir::Node::if_then(
441 vyre_foundation::ir::Expr::lt(
442 vyre_foundation::ir::Expr::var("idx"),
443 vyre_foundation::ir::Expr::u32(n),
444 ),
445 vec![vyre_foundation::ir::Node::store(
446 out.name_str(),
447 vyre_foundation::ir::Expr::var("idx"),
448 f(
449 vyre_foundation::ir::Expr::load(
450 a.name_str(),
451 vyre_foundation::ir::Expr::var("idx"),
452 ),
453 vyre_foundation::ir::Expr::load(
454 b.name_str(),
455 vyre_foundation::ir::Expr::var("idx"),
456 ),
457 ),
458 )],
459 ),
460 ];
461
462 let group = options.workgroup_size.unwrap_or([64, 1, 1]);
463
464 Ok(vyre_foundation::ir::Program::wrapped(
465 vec![
466 vyre_foundation::ir::BufferDecl::storage(
467 a.name_str(),
468 0,
469 vyre_foundation::ir::BufferAccess::ReadOnly,
470 vyre_foundation::ir::DataType::U32,
471 )
472 .with_count(n),
473 vyre_foundation::ir::BufferDecl::storage(
474 b.name_str(),
475 1,
476 vyre_foundation::ir::BufferAccess::ReadOnly,
477 vyre_foundation::ir::DataType::U32,
478 )
479 .with_count(n),
480 vyre_foundation::ir::BufferDecl::output(
481 out.name_str(),
482 2,
483 vyre_foundation::ir::DataType::U32,
484 )
485 .with_count(n),
486 ],
487 group,
488 vec![crate::region::wrap_anonymous(op_id, body)],
489 ))
490}
491
492#[allow(dead_code)]
493pub(crate) fn build_elementwise_unary<F>(
494 op_id: &'static str,
495 a: crate::tensor_ref::TensorRef,
496 out: crate::tensor_ref::TensorRef,
497 options: BuildOptions,
498 f: F,
499) -> Result<vyre_foundation::ir::Program, crate::tensor_ref::TensorRefError>
500where
501 F: Fn(vyre_foundation::ir::Expr) -> vyre_foundation::ir::Expr,
502{
503 check_tensors(
504 op_id,
505 &[
506 (&a, vyre_foundation::ir::DataType::U32),
507 (&out, vyre_foundation::ir::DataType::U32),
508 ],
509 )?;
510
511 if a.shape != out.shape {
512 return Err(crate::tensor_ref::TensorRefError::ShapeMismatch {
513 name: "elementwise_unary".into(),
514 found: vec![],
515 expected: vec![],
516 op: op_id,
517 });
518 }
519
520 let n = a.element_count().ok_or_else(|| {
521 crate::tensor_ref::TensorRefError::ElementCountOverflow {
522 name: a.name_str().to_string(),
523 shape: a.shape.to_vec(),
524 }
525 })?;
526 let body = vec![
527 vyre_foundation::ir::Node::let_bind(
528 "idx",
529 vyre_foundation::ir::Expr::InvocationId { axis: 0 },
530 ),
531 vyre_foundation::ir::Node::if_then(
532 vyre_foundation::ir::Expr::lt(
533 vyre_foundation::ir::Expr::var("idx"),
534 vyre_foundation::ir::Expr::u32(n),
535 ),
536 vec![vyre_foundation::ir::Node::store(
537 out.name_str(),
538 vyre_foundation::ir::Expr::var("idx"),
539 f(vyre_foundation::ir::Expr::load(
540 a.name_str(),
541 vyre_foundation::ir::Expr::var("idx"),
542 )),
543 )],
544 ),
545 ];
546
547 let group = options.workgroup_size.unwrap_or([64, 1, 1]);
548
549 Ok(vyre_foundation::ir::Program::wrapped(
550 vec![
551 vyre_foundation::ir::BufferDecl::storage(
552 a.name_str(),
553 0,
554 vyre_foundation::ir::BufferAccess::ReadOnly,
555 vyre_foundation::ir::DataType::U32,
556 )
557 .with_count(n),
558 vyre_foundation::ir::BufferDecl::output(
559 out.name_str(),
560 1,
561 vyre_foundation::ir::DataType::U32,
562 )
563 .with_count(n),
564 ],
565 group,
566 vec![crate::region::wrap_anonymous(op_id, body)],
567 ))
568}
569
570#[cfg(test)]
571mod tests {
572 use super::*;
573
574 #[test]
575 fn build_options_defaults_are_all_none() {
576 let o = BuildOptions::default();
577 assert!(o.workgroup_size.is_none());
578 assert!(o.region_generator.is_none());
579 assert!(o.tenant_id.is_none());
580 }
581
582 #[test]
583 fn build_options_chain_preserves_earlier_setters() {
584 let o = BuildOptions::new()
585 .with_workgroup_size([128, 1, 1])
586 .with_region_generator("test::op")
587 .with_tenant_id(7);
588 assert_eq!(o.workgroup_size, Some([128, 1, 1]));
589 assert_eq!(o.region_generator, Some("test::op"));
590 assert_eq!(o.tenant_id, Some(7));
591 }
592
593 #[test]
594 fn check_tensors_passes_on_clean_inputs() {
595 let a = TensorRef::u32_1d("a", 4);
596 let b = TensorRef::u32_1d("b", 4);
597 assert!(matches!(
598 check_tensors("op", &[(&a, DataType::U32), (&b, DataType::U32)]),
599 Ok(())
600 ));
601 }
602
603 #[test]
604 fn check_tensors_catches_dtype_mismatch() {
605 let a = TensorRef::u32_1d("a", 4);
606 let err = check_tensors("op", &[(&a, DataType::F32)]).unwrap_err();
607 assert!(matches!(err, TensorRefError::DtypeMismatch { .. }));
608 }
609
610 #[test]
611 fn check_tensors_catches_overflow() {
612 let a = TensorRef::new("big", DataType::U32, vec![1u32 << 20, 1u32 << 20]);
613 let err = check_tensors("op", &[(&a, DataType::U32)]).unwrap_err();
614 assert!(matches!(err, TensorRefError::ElementCountOverflow { .. }));
615 }
616
617 #[test]
618 fn check_tensors_catches_name_collision() {
619 let a = TensorRef::u32_1d("x", 4);
620 let b = TensorRef::u32_1d("x", 4);
621 let err = check_tensors("op", &[(&a, DataType::U32), (&b, DataType::U32)]).unwrap_err();
622 assert!(matches!(err, TensorRefError::NameCollision { .. }));
623 }
624
625 #[test]
626 fn indexed_map_builder_emits_shared_child_region() {
627 let program = build_indexed_map(
628 "vyre-libs::test::indexed_map_user",
629 vec![
630 BufferDecl::storage(
631 "input",
632 0,
633 vyre_foundation::ir::BufferAccess::ReadOnly,
634 DataType::U32,
635 )
636 .with_count(4),
637 BufferDecl::output("output", 1, DataType::U32).with_count(4),
638 ],
639 "output",
640 4,
641 [64, 1, 1],
642 |i| (i.clone(), Expr::load("input", i)),
643 );
644 let rendered = format!("{:?}", program.entry());
645 assert!(
646 rendered.contains(INDEXED_MAP_OP_ID),
647 "Fix: indexed-map users must share the same child region instead of copying loop skeletons: {rendered}"
648 );
649 }
650
651 #[test]
652 fn strided_writeback_builder_emits_shared_child_region() {
653 let program = Program::wrapped(
654 vec![BufferDecl::output("out", 0, DataType::F32).with_count(4)],
655 [4, 1, 1],
656 vec![crate::region::wrap_anonymous(
657 "vyre-libs::test::row_reduction_user",
658 vec![
659 Node::let_bind("local", Expr::LocalId { axis: 0 }),
660 strided_writeback_child(
661 "vyre-libs::test::row_reduction_user",
662 4,
663 1,
664 4,
665 "out",
666 vec![Node::let_bind("scale", Expr::f32(0.5))],
667 |_idx| Expr::var("scale"),
668 ),
669 ],
670 )],
671 );
672 let rendered = format!("{:?}", program.entry());
673 assert!(
674 rendered.contains(STRIDED_WRITEBACK_OP_ID),
675 "Fix: row-reduction writeback users must share the same child region instead of copying loop skeletons: {rendered}"
676 );
677 }
678}