Skip to main content

vyre_libs/
builder.rs

1//! Shared helpers used by the per-op Cat-A builders.
2//!
3//! Each op in `vyre-libs` ships a chainable builder that:
4//!
5//! 1. Accepts [`TensorRef`]s instead of bare `&str` buffer names, so
6//!    dtype + shape mismatches fail at `build()` time.
7//! 2. Checks every pair of buffer names is unique.
8//! 3. Verifies every [`TensorRef`]'s dtype against the op's expected dtype.
9//! 4. Verifies element-count overflow.
10//! 5. Allows chained overrides (workgroup size, region generator,
11//!    tenant id) without churning the function signature  -  extension
12//!    fields live inside a `#[non_exhaustive]` options struct so new
13//!    knobs never break existing call sites.
14//!
15//! `BuildOptions` is intentionally small at launch; fields are added
16//! rather than removed (the `#[non_exhaustive]` attribute enforces
17//! this). Every Cat-A op exposes its builder as `<Op>Builder::new(...)`
18//! and delegates defaults through `BuildOptions::default()`.
19
20use vyre_foundation::ir::model::expr::GeneratorRef;
21use vyre_foundation::ir::{BufferDecl, DataType, Expr, Node, Program};
22
23use crate::tensor_ref::{TensorRef, TensorRefError};
24
25/// Shared child region for one-output indexed maps.
26///
27/// This is the kernel skeleton behind embedding lookup, byte shuffles,
28/// quant pack/unpack, and similar data-layout transforms:
29/// `for i in 0..n { out[dst(i)] = value(i) }`.
30pub(crate) const INDEXED_MAP_OP_ID: &str = "vyre-libs::substrate::indexed_map";
31/// Shared child region for strided per-lane workgroup accumulators.
32pub(crate) const STRIDED_ACCUMULATE_OP_ID: &str = "vyre-libs::substrate::strided_accumulate";
33/// Shared child region for strided writeback after a tiled row reduction.
34pub(crate) const STRIDED_WRITEBACK_OP_ID: &str =
35    "anonymous::vyre-libs::substrate::strided_writeback";
36
37/// Shared options every Cat-A builder threads through. Lives here so
38/// every op agrees on the same surface.
39#[derive(Debug, Clone, Default)]
40#[non_exhaustive]
41pub struct BuildOptions {
42    /// Workgroup size override. `None` = op's canonical default.
43    pub workgroup_size: Option<[u32; 3]>,
44    /// Region generator override. `None` = op's canonical `"vyre-libs::…"`
45    /// identifier. Used when a downstream crate wraps a Cat-A op and
46    /// wants its own generator id in conformance certificates.
47    pub region_generator: Option<&'static str>,
48    /// Tenant id baked into the region metadata for multi-tenant
49    /// deployments. Routed through the megakernel's tenant-mask table
50    /// when the Program runs inside `vyre-runtime`.
51    pub tenant_id: Option<u32>,
52}
53
54impl BuildOptions {
55    /// Fluent constructor  -  start with defaults and chain overrides.
56    #[must_use]
57    pub fn new() -> Self {
58        Self::default()
59    }
60
61    /// Override the workgroup size.
62    #[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    /// Override the region generator name (must be `&'static str`).
69    #[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    /// Stamp a tenant id into the Cat-A op's region metadata.
76    #[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            /// Override the generated Program workgroup size.
87            #[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            /// Override the Region generator id.
94            #[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            /// Stamp the Region metadata with a tenant id.
101            #[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
112/// Validate a slice of `TensorRef`s against an expected `DataType`
113/// for each position, plus name-uniqueness across the whole slice.
114/// Used by every op's `build()` to consolidate the fanout of checks.
115pub fn check_tensors(
116    op: &'static str,
117    tensors: &[(&TensorRef, DataType)],
118) -> Result<(), TensorRefError> {
119    // Dtype check per tensor.
120    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
176/// Build the canonical one-output indexed-map skeleton.
177///
178/// Callers provide buffer declarations plus the semantic mapping from logical
179/// element `i` to `(dst_index, value)`. The loop, bounds guard, invocation id,
180/// workgroup default, and composition region stay centralized.
181pub(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
219/// Build a shared strided single-accumulator child region.
220///
221/// The parent must bind `local = LocalId(0)` before this child. The child
222/// accumulates `i = chunk * tile + local` for `chunk in 0..chunks`, guards
223/// `i < n`, and stores the lane-local accumulator into `scratch[local]`.
224pub(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/// Build a shared strided dual-accumulator child region.
258///
259/// This keeps paired reductions such as `(sum, sum_sq)` in one memory pass
260/// instead of forcing two separate scans over the input.
261#[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
300/// Build a shared strided writeback child region.
301///
302/// The parent must bind `local = LocalId(0)` before this child. Optional
303/// `prelude` nodes run once in workgroup zero before the strided write loop,
304/// which lets row reductions load reduced scalars exactly once per lane.
305pub(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
360/// Build an explicit trap program for an invalid infallible builder input.
361///
362/// The trap preserves the builder's `Program` contract while making the invalid
363/// input observable at execution. Fallible builders return their validation
364/// error before this boundary.
365pub(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/// Tensor-ref elementwise binary builder, used by `math::avg_floor`,
382/// `math::algebra`, and other binary-arithmetic primitives.
383#[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}