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
15fn mix_wire_fallback_hashable<T: Hash>(hasher: &mut blake3::Hasher, value: &T) {
16 let mut state = FxHasher::default();
17 value.hash(&mut state);
18 hasher.update(&state.finish().to_le_bytes());
19}
20
21struct FallbackWireHasher<'a>(&'a mut blake3::Hasher);
23
24impl NodeVisitor for FallbackWireHasher<'_> {
25 fn visit_node(&mut self, node: &Node) {
26 let h = &mut *self.0;
27 match node {
28 Node::Let { name, .. } => {
29 h.update(b"n:Let\0");
30 h.update(name.as_bytes());
31 }
32 Node::Assign { name, .. } => {
33 h.update(b"n:Assign\0");
34 h.update(name.as_bytes());
35 }
36 Node::Store { buffer, .. } => {
37 h.update(b"n:Store\0");
38 h.update(buffer.as_bytes());
39 }
40 Node::If { .. } => {
41 h.update(b"n:If\0");
42 }
43 Node::Loop { var, .. } => {
44 h.update(b"n:Loop\0");
45 h.update(var.as_bytes());
46 }
47 Node::IndirectDispatch {
48 count_buffer,
49 count_offset,
50 } => {
51 h.update(b"n:IndirectDispatch\0");
52 h.update(count_buffer.as_bytes());
53 h.update(&count_offset.to_le_bytes());
54 }
55 Node::AsyncLoad {
56 source,
57 destination,
58 tag,
59 ..
60 } => {
61 h.update(b"n:AsyncLoad\0");
62 h.update(source.as_bytes());
63 h.update(destination.as_bytes());
64 h.update(tag.as_bytes());
65 }
66 Node::AsyncStore {
67 source,
68 destination,
69 tag,
70 ..
71 } => {
72 h.update(b"n:AsyncStore\0");
73 h.update(source.as_bytes());
74 h.update(destination.as_bytes());
75 h.update(tag.as_bytes());
76 }
77 Node::AsyncWait { tag } => {
78 h.update(b"n:AsyncWait\0");
79 h.update(tag.as_bytes());
80 }
81 Node::Trap { tag, .. } => {
82 h.update(b"n:Trap\0");
83 h.update(tag.as_bytes());
84 }
85 Node::Resume { tag } => {
86 h.update(b"n:Resume\0");
87 h.update(tag.as_bytes());
88 }
89 Node::Return => {
90 h.update(b"n:Return\0");
91 }
92 Node::Barrier { ordering } => {
93 h.update(b"n:Barrier\0");
94 mix_wire_fallback_hashable(h, ordering);
95 }
96 Node::Block(_) => {
97 h.update(b"n:Block\0");
98 }
99 Node::Region {
100 generator,
101 source_region,
102 ..
103 } => {
104 h.update(b"n:Region\0");
105 h.update(generator.as_bytes());
106 if let Some(gen) = source_region {
107 h.update(gen.name.as_bytes());
108 }
109 }
110 Node::Opaque(ext) => {
111 h.update(b"n:Opaque\0");
112 h.update(ext.extension_kind().as_bytes());
113 }
114 }
115 }
116}
117
118impl ExprVisitor for FallbackWireHasher<'_> {
119 fn visit_expr(&mut self, expr: &Expr) {
120 let h = &mut *self.0;
121 match expr {
122 Expr::LitU32(v) => {
123 h.update(b"e:LitU32\0");
124 h.update(&v.to_le_bytes());
125 }
126 Expr::LitI32(v) => {
127 h.update(b"e:LitI32\0");
128 h.update(&v.to_le_bytes());
129 }
130 Expr::LitF32(v) => {
131 h.update(b"e:LitF32\0");
132 h.update(&v.to_le_bytes());
133 }
134 Expr::LitBool(v) => {
135 h.update(b"e:LitBool\0");
136 h.update(&[u8::from(*v)]);
137 }
138 Expr::Var(name) => {
139 h.update(b"e:Var\0");
140 h.update(name.as_bytes());
141 }
142 Expr::Load { buffer, .. } => {
143 h.update(b"e:Load\0");
144 h.update(buffer.as_bytes());
145 }
146 Expr::BufLen { buffer } => {
147 h.update(b"e:BufLen\0");
148 h.update(buffer.as_bytes());
149 }
150 Expr::InvocationId { axis } => {
151 h.update(b"e:InvocationId\0");
152 h.update(&[*axis]);
153 }
154 Expr::WorkgroupId { axis } => {
155 h.update(b"e:WorkgroupId\0");
156 h.update(&[*axis]);
157 }
158 Expr::LocalId { axis } => {
159 h.update(b"e:LocalId\0");
160 h.update(&[*axis]);
161 }
162 Expr::BinOp { op, .. } => {
163 h.update(b"e:BinOp\0");
164 mix_wire_fallback_hashable(h, op);
165 }
166 Expr::UnOp { op, .. } => {
167 h.update(b"e:UnOp\0");
168 mix_wire_fallback_hashable(h, op);
169 }
170 Expr::Call { op_id, .. } => {
171 h.update(b"e:Call\0");
172 h.update(op_id.as_bytes());
173 }
174 Expr::Select { .. } => {
175 h.update(b"e:Select\0");
176 }
177 Expr::Cast { target, .. } => {
178 h.update(b"e:Cast\0");
179 mix_wire_fallback_hashable(h, target);
180 }
181 Expr::Fma { .. } => {
182 h.update(b"e:Fma\0");
183 }
184 Expr::Atomic {
185 op,
186 buffer,
187 ordering,
188 ..
189 } => {
190 h.update(b"e:Atomic\0");
191 mix_wire_fallback_hashable(h, op);
192 h.update(buffer.as_bytes());
193 mix_wire_fallback_hashable(h, ordering);
194 }
195 Expr::SubgroupBallot { .. } => {
196 h.update(b"e:SubgroupBallot\0");
197 }
198 Expr::SubgroupShuffle { .. } => {
199 h.update(b"e:SubgroupShuffle\0");
200 }
201 Expr::SubgroupAdd { .. } => {
202 h.update(b"e:SubgroupAdd\0");
203 }
204 Expr::SubgroupLocalId => {
205 h.update(b"e:SubgroupLocalId\0");
206 }
207 Expr::SubgroupSize => {
208 h.update(b"e:SubgroupSize\0");
209 }
210 Expr::Opaque(ext) => {
211 h.update(b"e:Opaque\0");
212 h.update(ext.extension_kind().as_bytes());
213 }
214 }
215 }
216}
217
218impl Program {
219 #[must_use]
229 pub fn reconcile_runnable_top_level(self) -> Self {
230 if self.is_top_level_region_wrapped() {
231 return self;
232 }
233 let new_entry = Self::wrap_entry(self.entry().to_vec());
234 self.with_rewritten_entry(new_entry)
235 }
236
237 #[must_use]
239 #[inline]
240 pub fn buffer(&self, name: &str) -> Option<&super::BufferDecl> {
241 self.buffer_index
242 .get(name)
243 .and_then(|&index| self.buffers.get(index))
244 }
245
246 #[must_use]
248 #[inline]
249 pub fn buffers(&self) -> &[super::BufferDecl] {
250 self.buffers.as_ref()
251 }
252
253 #[must_use]
255 #[inline]
256 #[cfg(test)]
257 pub(crate) fn buffers_arc(&self) -> &Arc<[super::BufferDecl]> {
258 &self.buffers
259 }
260
261 #[must_use]
268 #[inline]
269 pub fn structural_eq(&self, other: &Self) -> bool {
270 if std::ptr::eq(self, other)
275 || (Arc::ptr_eq(&self.buffers, &other.buffers)
276 && Arc::ptr_eq(&self.entry, &other.entry)
277 && self.entry_op_id == other.entry_op_id
278 && self.non_composable_with_self == other.non_composable_with_self
279 && self.workgroup_size == other.workgroup_size)
280 {
281 return true;
282 }
283 self.entry_op_id == other.entry_op_id
284 && self.non_composable_with_self == other.non_composable_with_self
285 && buffers_equal_ignoring_declaration_order(&self.buffers, &other.buffers)
286 && self.workgroup_size == other.workgroup_size
287 && self.entry == other.entry
288 }
289
290 #[must_use]
292 #[inline]
293 pub fn workgroup_size(&self) -> [u32; 3] {
294 self.workgroup_size
295 }
296
297 #[must_use]
302 #[inline]
303 pub fn parallel_region_size(&self) -> [u32; 3] {
304 self.workgroup_size
305 }
306
307 #[must_use]
310 #[inline]
311 pub fn is_non_composable_with_self(&self) -> bool {
312 self.non_composable_with_self
313 }
314
315 #[must_use]
317 #[inline]
318 pub fn with_non_composable_with_self(mut self, flag: bool) -> Self {
319 self.non_composable_with_self = flag;
320 self.invalidate_caches();
321 self
322 }
323
324 #[inline]
329 pub fn set_workgroup_size(&mut self, workgroup_size: [u32; 3]) {
330 self.workgroup_size = workgroup_size;
331 self.invalidate_caches();
332 }
333
334 #[inline]
336 pub fn set_parallel_region_size(&mut self, parallel_region_size: [u32; 3]) {
337 self.workgroup_size = parallel_region_size;
338 self.invalidate_caches();
339 }
340
341 #[must_use]
343 #[inline]
344 pub fn entry(&self) -> &[Node] {
345 self.entry.as_ref().as_slice()
346 }
347
348 #[must_use]
350 #[inline]
351 pub fn entry_arc(&self) -> &Arc<Vec<Node>> {
352 &self.entry
353 }
354
355 #[must_use]
358 #[inline]
359 pub fn is_explicit_noop(&self) -> bool {
360 self.buffers().is_empty()
361 && matches!(self.entry(), [Node::Region { body, .. }] if body.is_empty())
362 }
363
364 #[must_use]
368 #[inline]
369 pub fn is_top_level_region_wrapped(&self) -> bool {
370 !self.entry.is_empty()
371 && self
372 .entry()
373 .iter()
374 .all(|node| matches!(node, Node::Region { .. }))
375 }
376
377 #[must_use]
380 pub fn top_level_region_violation(&self) -> Option<String> {
381 if self.entry().is_empty() {
382 return Some(
383 "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."
384 .to_string(),
385 );
386 }
387
388 self.entry()
389 .iter()
390 .enumerate()
391 .find(|(_, node)| !matches!(node, Node::Region { .. }))
392 .map(|(index, node)| {
393 format!(
394 "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.",
395 Self::top_level_node_name(node)
396 )
397 })
398 }
399
400 #[must_use]
402 #[inline]
403 pub fn entry_mut(&mut self) -> &mut Vec<Node> {
404 self.invalidate_caches();
405 Arc::make_mut(&mut self.entry)
406 }
407
408 #[must_use]
410 #[inline]
411 pub fn fingerprint(&self) -> [u8; 32] {
412 *self.fingerprint.get_or_init(|| {
413 let hash = self.compute_wire_hash();
414 let _ = self.hash.set(hash);
415 *hash.as_bytes()
416 })
417 }
418
419 #[must_use]
432 pub fn vsa_fingerprint(&self) -> Vec<u32> {
433 self.fingerprint()
434 .chunks_exact(core::mem::size_of::<u32>())
435 .map(|chunk| u32::from_le_bytes([chunk[0], chunk[1], chunk[2], chunk[3]]))
436 .collect()
437 }
438
439 #[must_use]
441 #[inline]
442 pub fn output_buffer_indices(&self) -> &[u32] {
443 self.output_buffer_index
444 .get_or_init(|| {
445 Arc::new(
446 self.buffers()
447 .iter()
448 .enumerate()
449 .filter_map(|(index, buffer)| {
450 (buffer.access() == BufferAccess::ReadWrite).then_some(index as u32)
451 })
452 .collect(),
453 )
454 })
455 .as_slice()
456 }
457
458 #[must_use]
460 #[inline]
461 pub fn has_indirect_dispatch(&self) -> bool {
462 *self.has_indirect_dispatch.get_or_init(|| {
463 let mut stack: Vec<&Node> = self.entry().iter().rev().collect();
464 while let Some(node) = stack.pop() {
465 match node {
466 Node::IndirectDispatch { .. } => return true,
467 Node::If {
468 then, otherwise, ..
469 } => {
470 stack.extend(otherwise.iter().rev());
471 stack.extend(then.iter().rev());
472 }
473 Node::Loop { body, .. } | Node::Block(body) => {
474 stack.extend(body.iter().rev());
475 }
476 Node::Region { body, .. } => {
477 stack.extend(body.iter().rev());
478 }
479 Node::Let { .. }
480 | Node::Assign { .. }
481 | Node::Store { .. }
482 | Node::Return
483 | Node::Barrier { .. }
484 | Node::AsyncLoad { .. }
485 | Node::AsyncStore { .. }
486 | Node::AsyncWait { .. }
487 | Node::Trap { .. }
488 | Node::Resume { .. }
489 | Node::Opaque(_) => {}
490 }
491 }
492 false
493 })
494 }
495
496 #[must_use]
498 #[inline]
499 pub fn has_buffer(&self, name: &str) -> bool {
500 self.buffer_index.contains_key(name)
501 }
502
503 #[must_use]
505 #[inline]
506 pub fn buffer_count(&self) -> usize {
507 self.buffers.len()
508 }
509
510 #[inline]
511 pub(super) fn build_buffer_index(
512 buffers: &[super::BufferDecl],
513 ) -> rustc_hash::FxHashMap<Arc<str>, usize> {
514 let mut index = rustc_hash::FxHashMap::default();
515 index.reserve(buffers.len());
516 for (buffer_index, buffer) in buffers.iter().enumerate() {
517 index
518 .entry(Arc::clone(&buffer.name))
519 .or_insert(buffer_index);
520 }
521 index
522 }
523
524 #[inline]
526 pub fn mark_structurally_validated(&self) {
527 self.structural_validated.store(true, Ordering::Release);
528 }
529
530 #[must_use]
532 #[inline]
533 pub fn is_structurally_validated(&self) -> bool {
534 self.structural_validated.load(Ordering::Acquire)
535 }
536
537 #[inline]
539 pub fn mark_validated_on(&self, backend_id: &str) {
540 self.validation_set
541 .insert(Arc::from(self.validation_cache_key(backend_id)));
542 }
543
544 #[must_use]
546 #[inline]
547 pub fn is_validated_on(&self, backend_id: &str) -> bool {
548 self.validation_set
549 .contains(self.validation_cache_key(backend_id).as_str())
550 }
551
552 #[deprecated(note = "use is_structurally_validated or is_validated_on")]
554 #[must_use]
555 #[inline]
556 pub fn is_validated(&self) -> bool {
557 self.is_structurally_validated()
558 }
559
560 #[deprecated(note = "use mark_structurally_validated or mark_validated_on")]
562 #[inline]
563 pub fn mark_validated(&self) {
564 self.mark_structurally_validated();
565 }
566
567 pub fn validate(&self) -> crate::error::Result<()> {
574 if self.is_structurally_validated() {
575 return Ok(());
576 }
577 let errors = crate::validate::validate(self);
578 if errors.is_empty() {
579 self.mark_structurally_validated();
580 return Ok(());
581 }
582 let mut message = String::new();
583 for (index, error) in errors.into_iter().enumerate() {
584 if index > 0 {
585 message.push_str("; ");
586 }
587 message.push_str(error.message());
588 }
589 Err(crate::error::Error::WireFormatValidation { message })
590 }
591
592 #[inline]
593 #[must_use]
601 pub fn estimate_peak_vram_bytes(&self) -> u64 {
602 self.buffers
603 .iter()
604 .map(|buffer| {
605 let element_size = buffer.element.size_bytes().unwrap_or(4);
606 (buffer.count as u64) * (element_size as u64)
607 })
608 .sum()
609 }
610
611 #[must_use]
613 pub fn peak_intensity(&self) -> OpIntensity {
614 let mut peak = OpIntensity::Free;
615 for node in self.entry().iter() {
616 peak = peak.max(self.node_intensity(node));
617 }
618 peak
619 }
620
621 fn node_intensity(&self, node: &crate::ir::Node) -> OpIntensity {
622 use crate::ir::Node;
623 match node {
624 Node::Let { value, .. } | Node::Assign { value, .. } => self.expr_intensity(value),
625 Node::Store { index, value, .. } => {
626 self.expr_intensity(index).max(self.expr_intensity(value))
627 }
628 Node::If {
629 cond,
630 then,
631 otherwise,
632 } => {
633 let mut p = self.expr_intensity(cond);
634 for n in then {
635 p = p.max(self.node_intensity(n));
636 }
637 for n in otherwise {
638 p = p.max(self.node_intensity(n));
639 }
640 p
641 }
642 Node::Loop { from, to, body, .. } => {
643 let mut p = self.expr_intensity(from).max(self.expr_intensity(to));
644 for n in body.iter() {
645 p = p.max(self.node_intensity(n));
646 }
647 p
648 }
649 Node::Block(nodes) => {
650 let mut p = OpIntensity::Free;
651 for n in nodes {
652 p = p.max(self.node_intensity(n));
653 }
654 p
655 }
656 Node::Region { body, .. } => {
657 let mut p = OpIntensity::Free;
658 for n in body.iter() {
659 p = p.max(self.node_intensity(n));
660 }
661 p
662 }
663 _ => OpIntensity::Free,
664 }
665 }
666
667 #[allow(clippy::only_used_in_recursion)]
668 fn expr_intensity(&self, expr: &crate::ir::Expr) -> OpIntensity {
669 use crate::ir::Expr;
670 match expr {
671 Expr::BinOp { op, left, right } => op
672 .intensity()
673 .max(self.expr_intensity(left))
674 .max(self.expr_intensity(right)),
675 Expr::UnOp { operand, .. } => self.expr_intensity(operand),
676 Expr::Load { index, .. } => self.expr_intensity(index),
677 Expr::Select {
678 cond,
679 true_val,
680 false_val,
681 } => self
682 .expr_intensity(cond)
683 .max(self.expr_intensity(true_val))
684 .max(self.expr_intensity(false_val)),
685 Expr::Cast { value, .. } => self.expr_intensity(value),
686 Expr::Fma { a, b, c } => self
687 .expr_intensity(a)
688 .max(self.expr_intensity(b))
689 .max(self.expr_intensity(c)),
690 Expr::Atomic {
691 index,
692 value,
693 expected,
694 ..
695 } => {
696 let mut p = self.expr_intensity(index).max(self.expr_intensity(value));
697 if let Some(e) = expected {
698 p = p.max(self.expr_intensity(e));
699 }
700 p.max(OpIntensity::Heavy)
701 }
702 Expr::SubgroupBallot { cond } => self.expr_intensity(cond).max(OpIntensity::Heavy),
703 Expr::SubgroupShuffle { value, lane } => self
704 .expr_intensity(value)
705 .max(self.expr_intensity(lane))
706 .max(OpIntensity::Heavy),
707 Expr::SubgroupAdd { value } => self.expr_intensity(value).max(OpIntensity::Heavy),
708 _ => OpIntensity::Free,
709 }
710 }
711
712 fn compute_wire_hash(&self) -> blake3::Hash {
713 match self.canonical_wire_hash() {
714 Ok(hash) => hash,
715 Err(error) => {
716 let structural = self.structural_fingerprint_fallback();
717 let err_msg = error.to_string();
718 let mut fallback = Vec::with_capacity(96 + err_msg.len() + structural.len());
719 fallback.extend_from_slice(b"VYRE-PROGRAM-CANONICAL-WIRE-HASH-ERROR\0");
720 fallback.extend_from_slice(err_msg.as_bytes());
721 fallback.push(0);
722 fallback.extend_from_slice(structural.as_bytes());
723 blake3::hash(&fallback)
724 }
725 }
726 }
727
728 fn structural_fingerprint_fallback(&self) -> String {
729 let mut hasher = blake3::Hasher::new();
730 hasher.update(b"VYRE-WIRE-FALLBACK-V4\0");
731 if let Some(id) = self.entry_op_id.as_deref() {
732 hasher.update(id.as_bytes());
733 }
734 hasher.update(b"\0");
735 for axis in &self.workgroup_size {
736 hasher.update(&axis.to_le_bytes());
737 }
738 hasher.update(&[u8::from(self.non_composable_with_self)]);
739 let mut keys: Vec<Vec<u8>> = self
740 .buffers()
741 .iter()
742 .map(buffer_decl_canonical_key)
743 .collect();
744 keys.sort_unstable();
745 for key in keys {
746 hasher.update(&key);
747 }
748 let mut visitor = FallbackWireHasher(&mut hasher);
749 walk_nodes_and_exprs(self, &mut visitor);
750 hasher.finalize().to_hex().to_string()
751 }
752
753 fn validation_cache_key(&self, backend_id: &str) -> String {
754 let fingerprint = self.fingerprint();
755 let mut key = String::with_capacity(backend_id.len() + 1 + 64);
756 key.push_str(backend_id);
757 key.push(':');
758 for byte in fingerprint {
759 use std::fmt::Write as _;
760 let _ = write!(&mut key, "{byte:02x}");
761 }
762 key
763 }
764
765 #[inline]
766 pub(super) fn invalidate_caches(&mut self) {
767 self.structural_validated.store(false, Ordering::Release);
768 self.validation_set.clear();
769 let _ = self.hash.take();
770 let _ = self.fingerprint.take();
771 let _ = self.output_buffer_index.take();
772 let _ = self.has_indirect_dispatch.take();
773 let _ = self.stats.take();
774 }
775
776 #[inline]
777 pub(super) fn wrap_entry(entry: Vec<Node>) -> Vec<Node> {
778 if !Self::entry_needs_root_region(&entry) {
779 return entry;
780 }
781 vec![Node::Region {
782 generator: Ident::from(Self::ROOT_REGION_GENERATOR),
783 source_region: None,
784 body: Arc::new(entry),
785 }]
786 }
787
788 #[inline]
789 fn entry_needs_root_region(entry: &[Node]) -> bool {
790 entry.is_empty()
791 || entry
792 .iter()
793 .any(|node| !matches!(node, Node::Region { .. }))
794 }
795
796 #[inline]
797 fn top_level_node_name(node: &Node) -> &'static str {
798 match node {
799 Node::Let { .. } => "Let",
800 Node::Assign { .. } => "Assign",
801 Node::Store { .. } => "Store",
802 Node::If { .. } => "If",
803 Node::Loop { .. } => "Loop",
804 Node::Return => "Return",
805 Node::Block(_) => "Block",
806 Node::Barrier { .. } => "Barrier",
807 Node::Region { .. } => "Region",
808 Node::IndirectDispatch { .. } => "IndirectDispatch",
809 Node::AsyncLoad { .. } => "AsyncLoad",
810 Node::AsyncStore { .. } => "AsyncStore",
811 Node::AsyncWait { .. } => "AsyncWait",
812 Node::Trap { .. } => "Trap",
813 Node::Resume { .. } => "Resume",
814 Node::Opaque(_) => "Opaque",
815 }
816 }
817}
818
819pub(crate) fn buffers_equal_ignoring_declaration_order(
820 left: &[super::BufferDecl],
821 right: &[super::BufferDecl],
822) -> bool {
823 if left.len() != right.len() {
824 return false;
825 }
826
827 if left == right {
834 return true;
835 }
836
837 let mut left_keys = Vec::with_capacity(left.len());
838 left_keys.extend(left.iter().map(buffer_decl_canonical_key));
839 let mut right_keys = Vec::with_capacity(right.len());
840 right_keys.extend(right.iter().map(buffer_decl_canonical_key));
841 left_keys.sort_unstable();
842 right_keys.sort_unstable();
843 left_keys == right_keys
844}
845
846pub(super) fn buffer_decl_canonical_key(buffer: &super::BufferDecl) -> Vec<u8> {
847 use crate::serial::wire::framing::{put_len_u32, put_u32, put_u8};
848 use crate::serial::wire::tags::put_data_type;
849
850 let mut key = Vec::with_capacity(96);
851 if let Err(error) = put_len_u32(&mut key, buffer.name.len(), "buffer name length") {
852 key.extend_from_slice(b"\0name-length-error\0");
853 key.extend_from_slice(error.as_bytes());
854 }
855 key.extend_from_slice(buffer.name.as_bytes());
856 put_u32(&mut key, buffer.binding);
857 match crate::serial::wire::tags::access_tag::access_tag(buffer.access.clone()) {
858 Ok(tag) => put_u8(&mut key, tag),
859 Err(error) => {
860 put_u8(&mut key, u8::MAX);
861 key.extend_from_slice(error.as_bytes());
862 }
863 }
864 put_u8(
865 &mut key,
866 match buffer.kind {
867 super::MemoryKind::Global => 0,
868 super::MemoryKind::Shared => 1,
869 super::MemoryKind::Uniform => 2,
870 super::MemoryKind::Local => 3,
871 super::MemoryKind::Readonly => 4,
872 super::MemoryKind::Persistent => 5,
873 super::MemoryKind::Push => 6,
874 },
875 );
876 if let Err(error) = put_data_type(&mut key, &buffer.element) {
877 key.extend_from_slice(b"\0dtype-error\0");
878 key.extend_from_slice(error.as_bytes());
879 }
880 put_u32(&mut key, buffer.count);
881 put_u8(&mut key, u8::from(buffer.is_output));
882 put_u8(&mut key, u8::from(buffer.pipeline_live_out));
883 match &buffer.output_byte_range {
884 Some(range) => {
885 put_u8(&mut key, 1);
886 match u32::try_from(range.start) {
887 Ok(start) => put_u32(&mut key, start),
888 Err(error) => {
889 put_u32(&mut key, u32::MAX);
890 key.extend_from_slice(error.to_string().as_bytes());
891 }
892 }
893 match u32::try_from(range.end) {
894 Ok(end) => put_u32(&mut key, end),
895 Err(error) => {
896 put_u32(&mut key, u32::MAX);
897 key.extend_from_slice(error.to_string().as_bytes());
898 }
899 }
900 }
901 None => put_u8(&mut key, 0),
902 }
903 match buffer.hints.coalesce_axis {
904 Some(axis) => {
905 put_u8(&mut key, 1);
906 put_u8(&mut key, axis);
907 }
908 None => put_u8(&mut key, 0),
909 }
910 put_u32(&mut key, buffer.hints.preferred_alignment);
911 put_u8(
912 &mut key,
913 match buffer.hints.cache_locality {
914 super::CacheLocality::Streaming => 0,
915 super::CacheLocality::Temporal => 1,
916 super::CacheLocality::Random => 2,
917 },
918 );
919 put_u8(&mut key, u8::from(buffer.bytes_extraction));
920 key
921}