vyre_foundation/ir_inner/model/program/
meta.rs

1use std::hash::{Hash, Hasher as _};
2use std::sync::atomic::Ordering;
3use std::sync::Arc;
4
5use rustc_hash::FxHasher;
6use vyre_spec::bin_op::OpIntensity;
7
8use crate::ir::{Expr, Node};
9use crate::ir_inner::model::expr::Ident;
10use crate::ir_inner::model::types::BufferAccess;
11use crate::transform::visit::{walk_nodes_and_exprs, ExprVisitor, NodeVisitor};
12
13use super::Program;
14
15/// Provenance for mutations that invalidate Program validation/cache state.
16#[derive(Clone, Copy, Debug, Eq, PartialEq)]
17#[repr(u8)]
18pub enum ProgramMutationProvenance {
19    /// Program has not been mutated since construction or successful validation.
20    Clean = 0,
21    /// The non-composable dispatch flag changed.
22    NonComposableFlag = 1,
23    /// The target workgroup dimensions changed.
24    WorkgroupSize = 2,
25    /// The substrate-neutral parallel-region dimensions changed.
26    ParallelRegionSize = 3,
27    /// A caller borrowed the mutable entry vector.
28    EntryMutation = 4,
29    /// Internal builder or decode path rewrote Program shape.
30    InternalShapeMutation = 5,
31    /// Mutation provenance is unknown, so validation must fail closed.
32    Unknown = 255,
33}
34
35impl ProgramMutationProvenance {
36    #[inline]
37    const fn from_code(code: u8) -> Self {
38        match code {
39            0 => Self::Clean,
40            1 => Self::NonComposableFlag,
41            2 => Self::WorkgroupSize,
42            3 => Self::ParallelRegionSize,
43            4 => Self::EntryMutation,
44            5 => Self::InternalShapeMutation,
45            _ => Self::Unknown,
46        }
47    }
48}
49
50fn mix_wire_fallback_hashable<T: Hash>(hasher: &mut blake3::Hasher, value: &T) {
51    let mut state = FxHasher::default();
52    value.hash(&mut state);
53    hasher.update(&state.finish().to_le_bytes());
54}
55
56/// Bounded IR structure digest for wire-hash fallback (never formats full IR via `Debug`).
57struct FallbackWireHasher<'a>(&'a mut blake3::Hasher);
58
59impl NodeVisitor for FallbackWireHasher<'_> {
60    fn visit_node(&mut self, node: &Node) {
61        let h = &mut *self.0;
62        match node {
63            Node::Let { name, .. } => {
64                h.update(b"n:Let\0");
65                h.update(name.as_bytes());
66            }
67            Node::Assign { name, .. } => {
68                h.update(b"n:Assign\0");
69                h.update(name.as_bytes());
70            }
71            Node::Store { buffer, .. } => {
72                h.update(b"n:Store\0");
73                h.update(buffer.as_bytes());
74            }
75            Node::If { .. } => {
76                h.update(b"n:If\0");
77            }
78            Node::Loop { var, .. } => {
79                h.update(b"n:Loop\0");
80                h.update(var.as_bytes());
81            }
82            Node::IndirectDispatch {
83                count_buffer,
84                count_offset,
85            } => {
86                h.update(b"n:IndirectDispatch\0");
87                h.update(count_buffer.as_bytes());
88                h.update(&count_offset.to_le_bytes());
89            }
90            Node::AsyncLoad {
91                source,
92                destination,
93                tag,
94                ..
95            } => {
96                h.update(b"n:AsyncLoad\0");
97                h.update(source.as_bytes());
98                h.update(destination.as_bytes());
99                h.update(tag.as_bytes());
100            }
101            Node::AsyncStore {
102                source,
103                destination,
104                tag,
105                ..
106            } => {
107                h.update(b"n:AsyncStore\0");
108                h.update(source.as_bytes());
109                h.update(destination.as_bytes());
110                h.update(tag.as_bytes());
111            }
112            Node::AsyncWait { tag } => {
113                h.update(b"n:AsyncWait\0");
114                h.update(tag.as_bytes());
115            }
116            Node::Trap { tag, .. } => {
117                h.update(b"n:Trap\0");
118                h.update(tag.as_bytes());
119            }
120            Node::Resume { tag } => {
121                h.update(b"n:Resume\0");
122                h.update(tag.as_bytes());
123            }
124            Node::AllReduce { buffer, op, group } => {
125                h.update(b"n:AllReduce\0");
126                h.update(buffer.as_bytes());
127                h.update(&op.builtin_wire_tag().to_le_bytes());
128                h.update(&group.as_u32().to_le_bytes());
129            }
130            Node::AllGather {
131                input,
132                output,
133                group,
134            } => {
135                h.update(b"n:AllGather\0");
136                h.update(input.as_bytes());
137                h.update(output.as_bytes());
138                h.update(&group.as_u32().to_le_bytes());
139            }
140            Node::ReduceScatter {
141                input,
142                output,
143                op,
144                group,
145            } => {
146                h.update(b"n:ReduceScatter\0");
147                h.update(input.as_bytes());
148                h.update(output.as_bytes());
149                h.update(&op.builtin_wire_tag().to_le_bytes());
150                h.update(&group.as_u32().to_le_bytes());
151            }
152            Node::Broadcast {
153                buffer,
154                root,
155                group,
156            } => {
157                h.update(b"n:Broadcast\0");
158                h.update(buffer.as_bytes());
159                h.update(&root.to_le_bytes());
160                h.update(&group.as_u32().to_le_bytes());
161            }
162            Node::Return => {
163                h.update(b"n:Return\0");
164            }
165            Node::Barrier { ordering } => {
166                h.update(b"n:Barrier\0");
167                mix_wire_fallback_hashable(h, ordering);
168            }
169            Node::Block(_) => {
170                h.update(b"n:Block\0");
171            }
172            Node::Region {
173                generator,
174                source_region,
175                ..
176            } => {
177                h.update(b"n:Region\0");
178                h.update(generator.as_bytes());
179                if let Some(source_gen) = source_region {
180                    h.update(source_gen.name.as_bytes());
181                }
182            }
183            Node::Opaque(ext) => {
184                h.update(b"n:Opaque\0");
185                h.update(ext.extension_kind().as_bytes());
186            }
187        }
188    }
189}
190
191impl ExprVisitor for FallbackWireHasher<'_> {
192    fn visit_expr(&mut self, expr: &Expr) {
193        let h = &mut *self.0;
194        match expr {
195            Expr::LitU32(v) => {
196                h.update(b"e:LitU32\0");
197                h.update(&v.to_le_bytes());
198            }
199            Expr::LitI32(v) => {
200                h.update(b"e:LitI32\0");
201                h.update(&v.to_le_bytes());
202            }
203            Expr::LitF32(v) => {
204                h.update(b"e:LitF32\0");
205                h.update(&v.to_le_bytes());
206            }
207            Expr::LitBool(v) => {
208                h.update(b"e:LitBool\0");
209                h.update(&[u8::from(*v)]);
210            }
211            Expr::Var(name) => {
212                h.update(b"e:Var\0");
213                h.update(name.as_bytes());
214            }
215            Expr::Load { buffer, .. } => {
216                h.update(b"e:Load\0");
217                h.update(buffer.as_bytes());
218            }
219            Expr::BufLen { buffer } => {
220                h.update(b"e:BufLen\0");
221                h.update(buffer.as_bytes());
222            }
223            Expr::BufferRef { buffer } => {
224                h.update(b"e:BufferRef\0");
225                h.update(buffer.as_bytes());
226            }
227            Expr::InvocationId { axis } => {
228                h.update(b"e:InvocationId\0");
229                h.update(&[*axis]);
230            }
231            Expr::WorkgroupId { axis } => {
232                h.update(b"e:WorkgroupId\0");
233                h.update(&[*axis]);
234            }
235            Expr::LocalId { axis } => {
236                h.update(b"e:LocalId\0");
237                h.update(&[*axis]);
238            }
239            Expr::BinOp { op, .. } => {
240                h.update(b"e:BinOp\0");
241                mix_wire_fallback_hashable(h, op);
242            }
243            Expr::UnOp { op, .. } => {
244                h.update(b"e:UnOp\0");
245                mix_wire_fallback_hashable(h, op);
246            }
247            Expr::Call { op_id, .. } => {
248                h.update(b"e:Call\0");
249                h.update(op_id.as_bytes());
250            }
251            Expr::Select { .. } => {
252                h.update(b"e:Select\0");
253            }
254            Expr::Cast { target, .. } => {
255                h.update(b"e:Cast\0");
256                mix_wire_fallback_hashable(h, target);
257            }
258            Expr::Fma { .. } => {
259                h.update(b"e:Fma\0");
260            }
261            Expr::Atomic {
262                op,
263                buffer,
264                ordering,
265                ..
266            } => {
267                h.update(b"e:Atomic\0");
268                mix_wire_fallback_hashable(h, op);
269                h.update(buffer.as_bytes());
270                mix_wire_fallback_hashable(h, ordering);
271            }
272            Expr::SubgroupBallot { .. } => {
273                h.update(b"e:SubgroupBallot\0");
274            }
275            Expr::SubgroupShuffle { .. } => {
276                h.update(b"e:SubgroupShuffle\0");
277            }
278            Expr::SubgroupReduce { op, .. } => {
279                h.update(b"e:SubgroupReduce\0");
280                h.update(&[op.builtin_wire_tag()]);
281            }
282            Expr::SubgroupLocalId => {
283                h.update(b"e:SubgroupLocalId\0");
284            }
285            Expr::SubgroupSize => {
286                h.update(b"e:SubgroupSize\0");
287            }
288            Expr::Opaque(ext) => {
289                h.update(b"e:Opaque\0");
290                h.update(ext.extension_kind().as_bytes());
291            }
292        }
293    }
294}
295
296impl Program {
297    /// Re-apply the same top-level `Node::Region` contract as
298    /// [`Program::wrapped`].
299    ///
300    /// The [`region_inline_engine`](crate::optimizer::passes::cleanup::region_inline_engine)
301    /// pass flattens small Category-A regions so CSE/DCE can see a single
302    /// function-shaped body, which can leave a statement-shaped entry list. The
303    /// standard optimizer run ends with this helper so the program remains in
304    /// a runnable, validator/reference-interpreter–compatible form while
305    /// still benefiting from the inline pass.
306    #[must_use]
307    pub fn reconcile_runnable_top_level(self) -> Self {
308        if self.is_top_level_region_wrapped() {
309            return self;
310        }
311        // Move the entry Vec out via map_entry's Arc-aware path; one
312        // Program rebuild instead of two scaffold rebuilds.
313        self.map_entry(Self::wrap_entry)
314    }
315
316    /// Look up a buffer declaration by name.
317    #[must_use]
318    #[inline]
319    pub fn buffer(&self, name: &str) -> Option<&super::BufferDecl> {
320        self.buffer_index
321            .get(name)
322            .and_then(|&index| self.buffers.get(index))
323    }
324
325    /// Declared buffers.
326    #[must_use]
327    #[inline]
328    pub fn buffers(&self) -> &[super::BufferDecl] {
329        self.buffers.as_ref()
330    }
331
332    /// Access the buffer declaration Arc directly for identity checks.
333    #[must_use]
334    #[inline]
335    #[cfg(test)]
336    pub(crate) fn buffers_arc(&self) -> &Arc<[super::BufferDecl]> {
337        &self.buffers
338    }
339
340    /// Compare two programs by observable IR structure.
341    ///
342    /// This walk intentionally ignores buffer declaration order and never
343    /// consults arena-local allocation identity. Two programs are structurally
344    /// equal when they declare the same buffers, workgroup size, optional entry
345    /// op id, and entry body semantics.
346    #[must_use]
347    #[inline]
348    pub fn structural_eq(&self, other: &Self) -> bool {
349        // Identity short-circuit: Program::clone shares all the
350        // inner Arcs, so comparing a cloned program against its
351        // source (the common optimizer-pipeline pattern) is pure
352        // refcount comparison.
353        if std::ptr::eq(self, other)
354            || (Arc::ptr_eq(&self.buffers, &other.buffers)
355                && Arc::ptr_eq(&self.entry, &other.entry)
356                && self.entry_op_id == other.entry_op_id
357                && self.non_composable_with_self == other.non_composable_with_self
358                && self.workgroup_size == other.workgroup_size)
359        {
360            return true;
361        }
362        self.entry_op_id == other.entry_op_id
363            && self.non_composable_with_self == other.non_composable_with_self
364            && buffers_equal_ignoring_declaration_order(&self.buffers, &other.buffers)
365            && self.workgroup_size == other.workgroup_size
366            && self.entry == other.entry
367    }
368
369    /// Workgroup dimensions.
370    #[must_use]
371    #[inline]
372    pub fn workgroup_size(&self) -> [u32; 3] {
373        self.workgroup_size
374    }
375
376    /// Substrate-neutral alias for [`workgroup_size`](Self::workgroup_size).
377    ///
378    /// Naming: "parallel region" avoids picking a single target substrate's
379    /// word for one dispatch invocation grouping.
380    #[must_use]
381    #[inline]
382    pub fn parallel_region_size(&self) -> [u32; 3] {
383        self.workgroup_size
384    }
385
386    /// Return true when this program must not be fused with another copy
387    /// of itself in the same megakernel.
388    #[must_use]
389    #[inline]
390    pub fn is_non_composable_with_self(&self) -> bool {
391        self.non_composable_with_self
392    }
393
394    /// Mark this program as non-composable with itself.
395    #[must_use]
396    #[inline]
397    pub fn with_non_composable_with_self(mut self, flag: bool) -> Self {
398        self.non_composable_with_self = flag;
399        self.invalidate_caches_for(ProgramMutationProvenance::NonComposableFlag);
400        self
401    }
402
403    /// Set the workgroup dimensions in place. Used by harnesses that
404    /// need to clone-and-rewrite a program's workgroup size for fallback
405    /// dispatch  -  the alternative was to reconstruct the entire Program,
406    /// which is unnecessarily expensive when only one field changes.
407    #[inline]
408    pub fn set_workgroup_size(&mut self, workgroup_size: [u32; 3]) {
409        self.workgroup_size = workgroup_size;
410        self.invalidate_caches_for(ProgramMutationProvenance::WorkgroupSize);
411    }
412
413    /// Substrate-neutral alias for [`set_workgroup_size`](Self::set_workgroup_size).
414    #[inline]
415    pub fn set_parallel_region_size(&mut self, parallel_region_size: [u32; 3]) {
416        self.workgroup_size = parallel_region_size;
417        self.invalidate_caches_for(ProgramMutationProvenance::ParallelRegionSize);
418    }
419
420    /// Entry-point nodes.
421    #[must_use]
422    #[inline]
423    pub fn entry(&self) -> &[Node] {
424        self.entry.as_ref().as_slice()
425    }
426
427    /// Shared entry-point body Arc for identity checks.
428    #[must_use]
429    #[inline]
430    pub fn entry_arc(&self) -> &Arc<Vec<Node>> {
431        &self.entry
432    }
433
434    /// Return true when this Program is the canonical no-op shape produced by
435    /// [`Program::empty`]: no buffers and a single empty root Region.
436    #[must_use]
437    #[inline]
438    pub fn is_explicit_noop(&self) -> bool {
439        self.buffers().is_empty()
440            && matches!(self.entry(), [Node::Region { body, .. }] if body.is_empty())
441    }
442
443    /// Return true when the program satisfies the top-level region-chain
444    /// invariant: at least one top-level node, and every top-level node is a
445    /// `Node::Region`.
446    #[must_use]
447    #[inline]
448    pub fn is_top_level_region_wrapped(&self) -> bool {
449        !self.entry.is_empty()
450            && self
451                .entry()
452                .iter()
453                .all(|node| matches!(node, Node::Region { .. }))
454    }
455
456    /// Actionable error text describing why the top-level region invariant
457    /// failed, or `None` when the entry is valid.
458    #[must_use]
459    pub fn top_level_region_violation(&self) -> Option<String> {
460        if self.entry().is_empty() {
461            return Some(
462                "program entry has no top-level Region. Fix: construct runnable programs with Program::wrapped(...) or wrap the body in Node::Region before validation, interpretation, or dispatch."
463                    .to_string(),
464            );
465        }
466
467        self.entry()
468            .iter()
469            .enumerate()
470            .find(|(_, node)| !matches!(node, Node::Region { .. }))
471            .map(|(index, node)| {
472                format!(
473                    "program entry node {index} is `{}` instead of `Node::Region`. Fix: construct runnable programs with Program::wrapped(...) or wrap the top-level body in Node::Region; raw Program::new is reserved for wire decode and negative tests.",
474                    crate::ir_inner::model::node::node_op_id(node)
475                )
476            })
477    }
478
479    /// Mutable entry-point nodes for transformation passes.
480    #[must_use]
481    #[inline]
482    pub fn entry_mut(&mut self) -> &mut Vec<Node> {
483        self.invalidate_caches_for(ProgramMutationProvenance::EntryMutation);
484        Arc::make_mut(&mut self.entry)
485    }
486
487    /// Stable BLAKE3 fingerprint of the canonical wire-format bytes.
488    #[must_use]
489    #[inline]
490    pub fn fingerprint(&self) -> [u8; 32] {
491        *self.fingerprint.get_or_init(|| {
492            let hash = self.compute_wire_hash();
493            let _ = self.hash.set(hash);
494            *hash.as_bytes()
495        })
496    }
497
498    /// VSA-style hypervector fingerprint of the canonical wire-format
499    /// bytes. Each `u32` lane is one segment of the program's blake3
500    /// hash; together they form an 8-lane hypervector suitable for
501    /// approximate similarity search via hamming distance.
502    ///
503    /// Use as the canonical cache key for approximate-match caches
504    /// (e.g. validation cache, AOT artifact dedup); use
505    /// [`Self::fingerprint`] for exact-match lookups.
506    ///
507    /// Wires the substrate's #29 hypervector primitive into Program
508    /// itself  -  every Program now carries its own VSA fingerprint
509    /// without callers having to reach into the substrate explicitly.
510    #[must_use]
511    pub fn vsa_fingerprint(&self) -> Vec<u32> {
512        self.fingerprint()
513            .chunks_exact(core::mem::size_of::<u32>())
514            .map(|chunk| u32::from_le_bytes([chunk[0], chunk[1], chunk[2], chunk[3]]))
515            .collect()
516    }
517
518    /// Indices of writable storage outputs in `buffers()` order.
519    #[must_use]
520    #[inline]
521    pub fn output_buffer_indices(&self) -> &[u32] {
522        self.output_buffer_index
523            .get_or_init(|| {
524                Arc::new(
525                    self.buffers()
526                        .iter()
527                        .enumerate()
528                        .filter_map(|(index, buffer)| {
529                            matches!(
530                                buffer.access(),
531                                BufferAccess::ReadWrite | BufferAccess::WriteOnly
532                            )
533                            .then(|| u32::try_from(index).ok())
534                            .flatten()
535                        })
536                        .collect(),
537                )
538            })
539            .as_slice()
540    }
541
542    /// True when the entry walk discovers any indirect dispatch node.
543    #[must_use]
544    #[inline]
545    pub fn has_indirect_dispatch(&self) -> bool {
546        *self.has_indirect_dispatch.get_or_init(|| {
547            // Fast-path: ProgramStats records every node kind seen during
548            // its single-pass walk. If the IndirectDispatch bit is unset,
549            // the tree definitely contains no IndirectDispatch nodes and
550            // the explicit traversal below would redundantly visit every
551            // node only to return false. Reading the bit is O(1).
552            if !self
553                .stats()
554                .has_any_node_kind(super::stats::NODE_KIND_INDIRECT_DISPATCH)
555            {
556                return false;
557            }
558            let mut stack: smallvec::SmallVec<[&Node; 32]> = self.entry().iter().rev().collect();
559            while let Some(node) = stack.pop() {
560                match node {
561                    Node::IndirectDispatch { .. } => return true,
562                    Node::If {
563                        then, otherwise, ..
564                    } => {
565                        stack.extend(otherwise.iter().rev());
566                        stack.extend(then.iter().rev());
567                    }
568                    Node::Loop { body, .. } | Node::Block(body) => {
569                        stack.extend(body.iter().rev());
570                    }
571                    Node::Region { body, .. } => {
572                        stack.extend(body.iter().rev());
573                    }
574                    Node::Let { .. }
575                    | Node::Assign { .. }
576                    | Node::Store { .. }
577                    | Node::AllReduce { .. }
578                    | Node::AllGather { .. }
579                    | Node::ReduceScatter { .. }
580                    | Node::Broadcast { .. }
581                    | Node::Return
582                    | Node::Barrier { .. }
583                    | Node::AsyncLoad { .. }
584                    | Node::AsyncStore { .. }
585                    | Node::AsyncWait { .. }
586                    | Node::Trap { .. }
587                    | Node::Resume { .. }
588                    | Node::Opaque(_) => {}
589                }
590            }
591            false
592        })
593    }
594
595    /// Check whether a named buffer exists.
596    #[must_use]
597    #[inline]
598    pub fn has_buffer(&self, name: &str) -> bool {
599        self.buffer_index.contains_key(name)
600    }
601
602    /// Number of declared buffers.
603    #[must_use]
604    #[inline]
605    pub fn buffer_count(&self) -> usize {
606        self.buffers.len()
607    }
608
609    #[inline]
610    pub(super) fn build_buffer_index(
611        buffers: &[super::BufferDecl],
612    ) -> rustc_hash::FxHashMap<Arc<str>, usize> {
613        let mut index = rustc_hash::FxHashMap::default();
614        index.reserve(buffers.len());
615        for (buffer_index, buffer) in buffers.iter().enumerate() {
616            index
617                .entry(Arc::clone(&buffer.name))
618                .or_insert(buffer_index);
619        }
620        index
621    }
622
623    /// Mark this program as successfully validated structurally.
624    #[inline]
625    pub fn mark_structurally_validated(&self) {
626        self.structural_validation_fingerprint.store(
627            self.current_validation_fingerprint_token(),
628            Ordering::Release,
629        );
630        self.mutation_provenance
631            .store(ProgramMutationProvenance::Clean as u8, Ordering::Release);
632        self.structural_validated.store(true, Ordering::Release);
633    }
634
635    /// Return true once structural validation has succeeded for this program shape.
636    #[must_use]
637    #[inline]
638    pub fn is_structurally_validated(&self) -> bool {
639        if !self.structural_validated.load(Ordering::Acquire) {
640            return false;
641        }
642        if self.validation_mutation_provenance() == ProgramMutationProvenance::Unknown {
643            self.structural_validated.store(false, Ordering::Release);
644            return false;
645        }
646        let recorded = self
647            .structural_validation_fingerprint
648            .load(Ordering::Acquire);
649        if recorded == 0 || recorded != self.current_validation_fingerprint_token() {
650            self.structural_validated.store(false, Ordering::Release);
651            return false;
652        }
653        true
654    }
655
656    /// Last mutation provenance recorded for validation/cache invalidation.
657    #[must_use]
658    #[inline]
659    pub fn validation_mutation_provenance(&self) -> ProgramMutationProvenance {
660        ProgramMutationProvenance::from_code(self.mutation_provenance.load(Ordering::Acquire))
661    }
662
663    /// Mark the Program as having been mutated by a boundary that cannot name a
664    /// concrete provenance. Validation fails closed until the Program is rebuilt
665    /// through a known constructor or known mutation API.
666    #[inline]
667    pub fn mark_unknown_mutation_provenance(&mut self) {
668        self.invalidate_caches_for(ProgramMutationProvenance::Unknown);
669    }
670
671    /// Mark this program as successfully validated for a specific backend.
672    #[inline]
673    pub fn mark_validated_on(&self, backend_id: &str) {
674        if self.validation_mutation_provenance() == ProgramMutationProvenance::Unknown {
675            return;
676        }
677        self.validation_set
678            .get_or_init(|| Arc::new(dashmap::DashSet::new()))
679            .insert(Arc::from(self.validation_cache_key(backend_id)));
680    }
681
682    /// Return true if this program has been validated for the given backend.
683    #[must_use]
684    #[inline]
685    pub fn is_validated_on(&self, backend_id: &str) -> bool {
686        self.validation_set
687            .get()
688            .is_some_and(|set| set.contains(self.validation_cache_key(backend_id).as_str()))
689    }
690
691    /// Deprecated: use `is_structurally_validated` or `is_validated_on`.
692    #[deprecated(note = "use is_structurally_validated or is_validated_on")]
693    #[must_use]
694    #[inline]
695    pub fn is_validated(&self) -> bool {
696        self.is_structurally_validated()
697    }
698
699    /// Deprecated: use `mark_structurally_validated` or `mark_validated_on`.
700    #[deprecated(note = "use mark_structurally_validated or mark_validated_on")]
701    #[inline]
702    pub fn mark_validated(&self) {
703        self.mark_structurally_validated();
704    }
705
706    /// Validate the program and cache the successful result on the program.
707    ///
708    /// # Errors
709    ///
710    /// Returns [`crate::Error::WireFormatValidation`] with every validation
711    /// message joined when the structural validator rejects the program.
712    pub fn validate(&self) -> crate::error::Result<()> {
713        if self.validation_mutation_provenance() == ProgramMutationProvenance::Unknown {
714            return Err(crate::error::Error::WireFormatValidation {
715                message: "program validation cache was invalidated by unknown mutation provenance. Fix: rebuild the Program through Program::wrapped/from_wire or use a named Program mutation API before validating.".into(),
716            });
717        }
718        if self.is_structurally_validated() {
719            return Ok(());
720        }
721        let errors = crate::validate::validate(self);
722        if errors.is_empty() {
723            self.mark_structurally_validated();
724            return Ok(());
725        }
726        let mut message = String::new();
727        for (index, error) in errors.into_iter().enumerate() {
728            if index > 0 {
729                message.push_str("; ");
730            }
731            message.push_str(error.message());
732        }
733        Err(crate::error::Error::WireFormatValidation { message })
734    }
735
736    #[inline]
737    /// Estimate the peak VRAM byte size of this Program.
738    ///
739    /// Innovation I.11: Static VRAM Pressure Analysis.
740    /// Returns the total bytes required by all storage and uniform buffers
741    /// declared in the Program. Optimizer passes use this to automatically
742    /// partition workloads if they would exceed a backend-specific safety
743    /// margin.
744    #[must_use]
745    pub fn estimate_peak_vram_bytes(&self) -> u64 {
746        self.buffers
747            .iter()
748            .map(|buffer| {
749                let Some(element_size) = buffer.element.size_bytes() else {
750                    return u64::MAX;
751                };
752                u64::from(buffer.count)
753                    .saturating_mul(u64::try_from(element_size).unwrap_or(u64::MAX))
754            })
755            .fold(0u64, u64::saturating_add)
756    }
757
758    /// Return the peak computational intensity found in any instruction.
759    #[must_use]
760    pub fn peak_intensity(&self) -> OpIntensity {
761        let mut peak = OpIntensity::Free;
762        for node in self.entry() {
763            peak = peak.max(Self::node_intensity(node));
764        }
765        peak
766    }
767
768    fn node_intensity(node: &crate::ir::Node) -> OpIntensity {
769        use crate::ir::Node;
770        match node {
771            Node::Let { value, .. } | Node::Assign { value, .. } => Self::expr_intensity(value),
772            Node::Store { index, value, .. } => {
773                Self::expr_intensity(index).max(Self::expr_intensity(value))
774            }
775            Node::If {
776                cond,
777                then,
778                otherwise,
779            } => {
780                let mut p = Self::expr_intensity(cond);
781                for n in then {
782                    p = p.max(Self::node_intensity(n));
783                }
784                for n in otherwise {
785                    p = p.max(Self::node_intensity(n));
786                }
787                p
788            }
789            Node::Loop { from, to, body, .. } => {
790                let mut p = Self::expr_intensity(from).max(Self::expr_intensity(to));
791                for n in body {
792                    p = p.max(Self::node_intensity(n));
793                }
794                p
795            }
796            Node::Block(nodes) => {
797                let mut p = OpIntensity::Free;
798                for n in nodes {
799                    p = p.max(Self::node_intensity(n));
800                }
801                p
802            }
803            Node::Region { body, .. } => {
804                let mut p = OpIntensity::Free;
805                for n in body.iter() {
806                    p = p.max(Self::node_intensity(n));
807                }
808                p
809            }
810            _ => OpIntensity::Free,
811        }
812    }
813
814    fn expr_intensity(expr: &crate::ir::Expr) -> OpIntensity {
815        use crate::ir::Expr;
816        match expr {
817            Expr::BinOp { op, left, right } => op
818                .intensity()
819                .max(Self::expr_intensity(left))
820                .max(Self::expr_intensity(right)),
821            Expr::UnOp { operand, .. } => Self::expr_intensity(operand),
822            Expr::Load { index, .. } => Self::expr_intensity(index),
823            Expr::Select {
824                cond,
825                true_val,
826                false_val,
827            } => Self::expr_intensity(cond)
828                .max(Self::expr_intensity(true_val))
829                .max(Self::expr_intensity(false_val)),
830            Expr::Cast { value, .. } => Self::expr_intensity(value),
831            Expr::Fma { a, b, c } => Self::expr_intensity(a)
832                .max(Self::expr_intensity(b))
833                .max(Self::expr_intensity(c)),
834            Expr::Atomic {
835                index,
836                value,
837                expected,
838                ..
839            } => {
840                let mut p = Self::expr_intensity(index).max(Self::expr_intensity(value));
841                if let Some(e) = expected {
842                    p = p.max(Self::expr_intensity(e));
843                }
844                p.max(OpIntensity::Heavy)
845            }
846            Expr::SubgroupBallot { cond } => Self::expr_intensity(cond).max(OpIntensity::Heavy),
847            Expr::SubgroupShuffle { value, lane } => Self::expr_intensity(value)
848                .max(Self::expr_intensity(lane))
849                .max(OpIntensity::Heavy),
850            Expr::SubgroupReduce { value, .. } => {
851                Self::expr_intensity(value).max(OpIntensity::Heavy)
852            }
853            _ => OpIntensity::Free,
854        }
855    }
856
857    fn compute_wire_hash(&self) -> blake3::Hash {
858        match self.canonical_wire_hash() {
859            Ok(hash) => hash,
860            Err(error) => {
861                let structural = self.structural_fingerprint_fallback();
862                let err_msg = error.to_string();
863                let mut fallback = Vec::with_capacity(96 + err_msg.len() + structural.len());
864                fallback.extend_from_slice(b"VYRE-PROGRAM-CANONICAL-WIRE-HASH-ERROR\0");
865                fallback.extend_from_slice(err_msg.as_bytes());
866                fallback.push(0);
867                fallback.extend_from_slice(structural.as_bytes());
868                blake3::hash(&fallback)
869            }
870        }
871    }
872
873    fn structural_fingerprint_fallback(&self) -> String {
874        let mut hasher = blake3::Hasher::new();
875        hasher.update(b"VYRE-WIRE-FALLBACK-V4\0");
876        if let Some(id) = self.entry_op_id.as_deref() {
877            hasher.update(id.as_bytes());
878        }
879        hasher.update(b"\0");
880        for axis in &self.workgroup_size {
881            hasher.update(&axis.to_le_bytes());
882        }
883        hasher.update(&[u8::from(self.non_composable_with_self)]);
884        let mut keys: Vec<Vec<u8>> = self
885            .buffers()
886            .iter()
887            .map(buffer_decl_canonical_key)
888            .collect();
889        keys.sort_unstable();
890        for key in keys {
891            hasher.update(&key);
892        }
893        let mut visitor = FallbackWireHasher(&mut hasher);
894        walk_nodes_and_exprs(self, &mut visitor);
895        hasher.finalize().to_hex().to_string()
896    }
897
898    fn validation_cache_key(&self, backend_id: &str) -> String {
899        const HEX: &[u8; 16] = b"0123456789abcdef";
900        let fingerprint = self.current_validation_fingerprint();
901        let mut key = String::with_capacity(backend_id.len() + 1 + 64);
902        key.push_str(backend_id);
903        key.push(':');
904        for &byte in &fingerprint {
905            key.push(HEX[(byte >> 4) as usize] as char);
906            key.push(HEX[(byte & 0x0f) as usize] as char);
907        }
908        key
909    }
910
911    #[inline]
912    pub(super) fn invalidate_caches(&mut self) {
913        self.invalidate_caches_for(ProgramMutationProvenance::InternalShapeMutation);
914    }
915
916    #[inline]
917    pub(super) fn invalidate_caches_for(&mut self, provenance: ProgramMutationProvenance) {
918        self.structural_validated.store(false, Ordering::Release);
919        self.structural_validation_fingerprint
920            .store(0, Ordering::Release);
921        self.mutation_provenance
922            .store(provenance as u8, Ordering::Release);
923        if let Some(set) = self.validation_set.get() {
924            set.clear();
925        }
926        let _ = self.hash.take();
927        let _ = self.fingerprint.take();
928        let _ = self.normalized_cache_digest.take();
929        drop(self.output_buffer_index.take());
930        let _ = self.has_indirect_dispatch.take();
931        drop(self.stats.take());
932    }
933
934    fn current_validation_fingerprint(&self) -> [u8; 32] {
935        *self.compute_wire_hash().as_bytes()
936    }
937
938    fn current_validation_fingerprint_token(&self) -> u64 {
939        let bytes = self.current_validation_fingerprint();
940        let token = u64::from_le_bytes([
941            bytes[0], bytes[1], bytes[2], bytes[3], bytes[4], bytes[5], bytes[6], bytes[7],
942        ]);
943        token.max(1)
944    }
945
946    #[inline]
947    pub(super) fn wrap_entry(entry: Vec<Node>) -> Vec<Node> {
948        if !Self::entry_needs_root_region(&entry) {
949            return entry;
950        }
951        vec![Node::Region {
952            generator: Ident::from(Self::ROOT_REGION_GENERATOR),
953            source_region: None,
954            body: Arc::new(entry),
955        }]
956    }
957
958    #[inline]
959    fn entry_needs_root_region(entry: &[Node]) -> bool {
960        entry.is_empty()
961            || entry
962                .iter()
963                .any(|node| !matches!(node, Node::Region { .. }))
964    }
965}
966
967pub(crate) fn buffers_equal_ignoring_declaration_order(
968    left: &[super::BufferDecl],
969    right: &[super::BufferDecl],
970) -> bool {
971    if left.len() != right.len() {
972        return false;
973    }
974
975    // VYRE_IR_HOTSPOTS HIGH (meta.rs:360-379): previous impl allocated
976    // two Vec<Vec<u8>> then sorted on every equality call. Fast-path:
977    // if the slices compare equal in-place (declaration orders match)
978    // we skip the key-materialization entirely. This catches every
979    // Program::clone(prog) == prog and every `Arc::clone`-equivalent
980    // comparison, which dominate the call distribution.
981    if left == right {
982        return true;
983    }
984
985    let mut left_keys = Vec::with_capacity(left.len());
986    left_keys.extend(left.iter().map(buffer_decl_canonical_key));
987    let mut right_keys = Vec::with_capacity(right.len());
988    right_keys.extend(right.iter().map(buffer_decl_canonical_key));
989    left_keys.sort_unstable();
990    right_keys.sort_unstable();
991    left_keys == right_keys
992}
993
994/// Canonical bytes for one buffer declaration, used to compare buffer sets
995/// independently of declaration order.
996///
997/// THIS KEY DECIDES PROGRAM EQUALITY. `Program::eq` delegates to
998/// `structural_eq`, which compares buffers through
999/// `buffers_equal_ignoring_declaration_order`, which keys each declaration
1000/// here. So any field this function omits is a field two programs may differ in
1001/// while comparing EQUAL.
1002///
1003/// It previously omitted `linear_type` and `shape_predicate` while including
1004/// `bytes_extraction`: one of three fields wired in, the other two forgotten.
1005/// Both omitted fields DECIDE VALIDATION VERDICTS (`validate::linear_type` and
1006/// `validate::shape_predicate::check_shape_predicates`), so a VALID program and
1007/// an INVALID one compared equal, and any caller shaped as "a == b, therefore
1008/// same program, therefore safe to reuse" was wrong for them.
1009///
1010/// COMPLETENESS IS ENFORCED BY THE COMPILER, not by review. The destructuring
1011/// below binds every field by name, so adding a field to `BufferDecl` FAILS TO
1012/// COMPILE here until somebody decides whether program equality depends on it.
1013/// That is the only mechanism that catches the next omission at the moment it is
1014/// introduced rather than years later through a wrong result. Do not replace it
1015/// with `buffer.field` accesses, and do not silence a new field with `..`.
1016///
1017/// `linear_type` and `shape_predicate` are encoded by the SAME functions the
1018/// wire encoder uses, so program equality and `Program::fingerprint` cannot
1019/// drift apart on these fields.
1020pub(super) fn buffer_decl_canonical_key(buffer: &super::BufferDecl) -> Vec<u8> {
1021    use crate::serial::wire::encode::to_wire::{linear_type_tag, put_shape_predicate};
1022    use crate::serial::wire::framing::{put_len_u32, put_u32, put_u8};
1023    use crate::serial::wire::tags::put_data_type;
1024
1025    // Exhaustive by field name: see the completeness note above. A new
1026    // `BufferDecl` field breaks this pattern and forces a decision.
1027    let super::BufferDecl {
1028        name,
1029        binding,
1030        access,
1031        kind,
1032        element,
1033        count,
1034        is_output,
1035        pipeline_live_out,
1036        output_byte_range,
1037        hints,
1038        bytes_extraction,
1039        linear_type,
1040        shape_predicate,
1041    } = buffer;
1042
1043    let mut key = Vec::with_capacity(96);
1044    if let Err(error) = put_len_u32(&mut key, name.len(), "buffer name length") {
1045        key.extend_from_slice(b"\0name-length-error\0");
1046        key.extend_from_slice(error.as_bytes());
1047    }
1048    key.extend_from_slice(name.as_bytes());
1049    put_u32(&mut key, *binding);
1050    match crate::serial::wire::tags::access_tag::access_tag(access) {
1051        Ok(tag) => put_u8(&mut key, tag),
1052        Err(error) => {
1053            put_u8(&mut key, u8::MAX);
1054            key.extend_from_slice(error.as_bytes());
1055        }
1056    }
1057    put_u8(&mut key, super::cache_digest::memory_kind_cache_tag(*kind));
1058    if let Err(error) = put_data_type(&mut key, element) {
1059        key.extend_from_slice(b"\0dtype-error\0");
1060        key.extend_from_slice(error.as_bytes());
1061    }
1062    put_u32(&mut key, *count);
1063    put_u8(&mut key, u8::from(*is_output));
1064    put_u8(&mut key, u8::from(*pipeline_live_out));
1065    match output_byte_range {
1066        Some(range) => {
1067            put_u8(&mut key, 1);
1068            match u32::try_from(range.start) {
1069                Ok(start) => put_u32(&mut key, start),
1070                Err(error) => {
1071                    put_u32(&mut key, u32::MAX);
1072                    key.extend_from_slice(error.to_string().as_bytes());
1073                }
1074            }
1075            match u32::try_from(range.end) {
1076                Ok(end) => put_u32(&mut key, end),
1077                Err(error) => {
1078                    put_u32(&mut key, u32::MAX);
1079                    key.extend_from_slice(error.to_string().as_bytes());
1080                }
1081            }
1082        }
1083        None => put_u8(&mut key, 0),
1084    }
1085    match hints.coalesce_axis {
1086        Some(axis) => {
1087            put_u8(&mut key, 1);
1088            put_u8(&mut key, axis);
1089        }
1090        None => put_u8(&mut key, 0),
1091    }
1092    put_u32(&mut key, hints.preferred_alignment);
1093    put_u8(
1094        &mut key,
1095        match hints.cache_locality {
1096            super::CacheLocality::Streaming => 0,
1097            super::CacheLocality::Temporal => 1,
1098            super::CacheLocality::Random => 2,
1099        },
1100    );
1101    put_u8(&mut key, u8::from(*bytes_extraction));
1102    put_u8(&mut key, linear_type_tag(*linear_type));
1103    if let Err(error) = put_shape_predicate(&mut key, shape_predicate.as_ref(), 0) {
1104        // The encoder refuses a predicate nested past the wire depth bound, and
1105        // it may already have emitted a prefix before refusing. A bare error
1106        // marker would let two DIFFERENT over-deep predicates share a key,
1107        // reintroducing the exact collision this field was added to close, so
1108        // fall back to the Debug form, which stays injective at any depth.
1109        key.extend_from_slice(b"\0shape-predicate-error\0");
1110        key.extend_from_slice(error.as_bytes());
1111        key.extend_from_slice(format!("{shape_predicate:?}").as_bytes());
1112    }
1113    key
1114}
vyre_foundation/ir_inner/model/program/meta.rs

vyre_foundation/ir_inner/model/program/
meta.rs
