vyre_foundation/ir_inner/model/program/cache_digest.rs
1//! Normalized compiled-pipeline cache digest, and its per-`Program` memo.
2//!
3//! # What this digest is for
4//!
5//! Backend pipeline caches key generated primary text or binary on this digest.
6//! It therefore has exactly one correctness obligation:
7//! two programs that compile to different backend code MUST get different
8//! digests. Over-keying (two programs that compile identically getting
9//! different digests) only costs a redundant compile; under-keying serves code
10//! generated for a different program, which surfaces as wrong output.
11//!
12//! # Why these inputs and no others
13//!
14//! The input set is derived from the single Program-to-emitter boundary rather
15//! than guessed. `vyre_lower::lower_for_emit` is the only
16//! Program-to-descriptor lowering, and every emitter reads only the resulting
17//! `KernelDescriptor`. `vyre-lower/src/lower.rs` reads exactly these program
18//! inputs into that descriptor:
19//!
20//! - `Program::workgroup_size` into `Dispatch`
21//! - `Program::entry` into the descriptor body
22//! - per `BufferDecl`: `name`, `binding`, `access`, `kind`, `element`, `count`
23//!
24//! Nothing else on `Program` or `BufferDecl` can reach an emitter, so nothing
25//! else belongs in the digest. `entry_op_id` is additionally included: it does
26//! not reach the descriptor, but it is cheap and keeps distinct certified
27//! operations in distinct cache lanes.
28//!
29//! `count` participates only where it is a static array length, which is
30//! [`BufferDecl::has_static_element_count`]. Runtime-sized storage and uniform
31//! buffers keep their count erased so that resizing a buffer does not force a
32//! shader recompile, an invariance pinned by the driver-owned disk-cache
33//! contracts.
34//!
35//! That erasure is safe only because the two emission paths differ. The primary
36//! text emitter reads `element_count` into generated text only under
37//! `MemoryClass::Shared`; every other class remains dynamically sized, so a
38//! runtime storage length cannot reach primary text. The primary binary emitter
39//! can bake a binding count as an immediate for an asynchronous copy. Its
40//! owning driver therefore carries a second full-wire-hash cache-key lane
41//! beside this digest rather than relying on this digest alone.
42//!
43//! # Staleness boundary
44//!
45//! `Program`'s IR fields are `pub`, so writing one directly through a `&mut
46//! Program` leaves this memo stale, exactly as it leaves `hash` and
47//! `fingerprint` stale. No production path does that: every mutator that can
48//! change these inputs (`entry_mut`, `set_workgroup_size`,
49//! `set_parallel_region_size`, `with_entry_op_id`,
50//! `with_non_composable_with_self`) routes through `invalidate_caches_for`.
51//!
52//! # Memoization
53//!
54//! The digest is a pure function of the program value, so it is memoized on
55//! the value itself in a `OnceLock`, exactly as [`Program::fingerprint`] is.
56//! A memo that lives on the value is keyed by nothing, so it structurally
57//! cannot serve another program's digest, and every sanctioned mutator clears
58//! it through `invalidate_caches_for`.
59//!
60//! # Three keys over `BufferDecl`, and why this one is the narrow one
61//!
62//! There are three independent keys derived from `BufferDecl`, and their
63//! coverage differs ON PURPOSE. `to_wire`/`fingerprint` and
64//! `buffer_decl_canonical_key` (program equality) both cover
65//! `bytes_extraction`, `linear_type` and `shape_predicate`; this digest covers
66//! NONE of the three. That is a decision, not an omission: those fields are
67//! declaration-level disciplines that feed validation verdicts, and none of
68//! them reaches an emitter, so a backend artifact cache must not vary on them.
69//! Do not "fix" this to match the other two keys. A field added to
70//! `BufferDecl` belongs here only if `vyre_lower::lower` reads it into the
71//! descriptor.
72
73use super::{BufferDecl, MemoryKind, Program};
74
75/// Version label for the normalized `Program` cache digest.
76///
77/// Single source of truth for both the digest's own domain separator and the
78/// label recorded in dispatch evidence, so the label can never describe an
79/// algorithm the digest no longer implements.
80///
81/// `v3` dropped `Program::is_structurally_validated` (validation state is
82/// provably not a codegen input) and added `BufferDecl::binding` plus the
83/// static-array `count`, which are.
84pub const NORMALIZED_PROGRAM_CACHE_DIGEST_VERSION: &str = "vyre-pipeline-cache-norm-v3";
85
86impl Program {
87 /// Normalized digest used by backend compiled-pipeline caches.
88 ///
89 /// Computed at most once per `Program` value: the result is memoized on the
90 /// program and cleared by every cache-invalidating mutation.
91 ///
92 /// # Errors
93 ///
94 /// Returns when the program contains an IR type or node shape that cannot
95 /// be serialized into stable cache identity. Dispatch admission surfaces
96 /// the error rather than generating a lossy cache key. Failures are not
97 /// memoized, so a caller that repairs the program sees the repair.
98 pub fn try_normalized_cache_digest(&self) -> Result<[u8; 32], String> {
99 if let Some(digest) = self.normalized_cache_digest.get() {
100 return Ok(*digest);
101 }
102 let digest = self.compute_normalized_cache_digest()?;
103 let _ = self.normalized_cache_digest.set(digest);
104 Ok(digest)
105 }
106
107 /// Uncached digest computation.
108 ///
109 /// `pub(super)` rather than private so the memo-soundness test can compare
110 /// a memoized read against a fresh recompute; a memo that returns a wrong
111 /// value consistently is invisible to any test that only calls the cached
112 /// path.
113 pub(super) fn compute_normalized_cache_digest(&self) -> Result<[u8; 32], String> {
114 super::record_digest_computation();
115
116 thread_local! {
117 static SCRATCH: std::cell::RefCell<Vec<u8>> =
118 std::cell::RefCell::new(Vec::with_capacity(1024));
119 }
120 SCRATCH.with(|cell| {
121 let mut scratch = cell.borrow_mut();
122 scratch.clear();
123 scratch.extend_from_slice(NORMALIZED_PROGRAM_CACHE_DIGEST_VERSION.as_bytes());
124 scratch.extend_from_slice(b"\0wg\0");
125 for axis in self.workgroup_size {
126 scratch.extend_from_slice(&axis.to_le_bytes());
127 }
128 scratch.extend_from_slice(b"\0op\0");
129 match self.entry_op_id.as_deref() {
130 Some(op) => {
131 // Length-prefixed: a raw name plus a NUL terminator lets an
132 // op id containing an interior NUL impersonate a different
133 // id followed by the next field.
134 scratch.extend_from_slice(&op_len_bytes(op.len())?);
135 scratch.extend_from_slice(op.as_bytes());
136 }
137 None => scratch.extend_from_slice(&[0u8; 4]),
138 }
139 scratch.extend_from_slice(b"\0bufs\0");
140 for buffer in self.buffers.iter() {
141 append_buffer_cache_key(&mut scratch, buffer)?;
142 }
143 scratch.extend_from_slice(b"\0body\0");
144 crate::serial::wire::append_node_list_fingerprint(&mut scratch, self.entry()).map_err(
145 |message| {
146 format!(
147 "failed to fingerprint pipeline-cache Program body: {message}. Fix: validate and normalize the Program before computing a compiled-pipeline cache key; invalid IR must not enter cache identity."
148 )
149 },
150 )?;
151 Ok(*blake3::hash(&scratch).as_bytes())
152 })
153 }
154}
155
156fn op_len_bytes(len: usize) -> Result<[u8; 4], String> {
157 u32::try_from(len)
158 .map(u32::to_le_bytes)
159 .map_err(|_| {
160 format!(
161 "pipeline-cache Program entry op id length {len} exceeds u32. Fix: shorten the certified operation id before computing a compiled-pipeline cache key."
162 )
163 })
164}
165
166fn append_buffer_cache_key(scratch: &mut Vec<u8>, buffer: &BufferDecl) -> Result<(), String> {
167 // Length-prefixed name. The v2 encoding wrote the raw name followed by a
168 // NUL, so a buffer literally named "a\0<tag bytes>" could produce the same
169 // byte stream as a buffer named "a" with different tags.
170 let name = buffer.name();
171 let name_len = u32::try_from(name.len()).map_err(|_| {
172 format!(
173 "pipeline-cache buffer name length {} exceeds u32. Fix: shorten the buffer name before computing a compiled-pipeline cache key.",
174 name.len()
175 )
176 })?;
177 scratch.extend_from_slice(&name_len.to_le_bytes());
178 scratch.extend_from_slice(name.as_bytes());
179
180 // Stable tags, never `enum as u8`: both enums are `#[non_exhaustive]`, so a
181 // variant inserted mid-list would silently remap discriminants and could
182 // alias a persisted entry recorded under the same version label.
183 scratch.push(memory_kind_cache_tag(buffer.kind()));
184 let access_tag = crate::serial::wire::tags::access_tag::access_tag(&buffer.access).map_err(
185 |message| {
186 format!(
187 "failed to tag pipeline-cache buffer access for `{name}`: {message}. Fix: validate and normalize the Program before computing a compiled-pipeline cache key; invalid IR must not enter cache identity."
188 )
189 },
190 )?;
191 scratch.push(access_tag);
192
193 // Read by vyre_lower::lower as the requested host binding slot, then emitted
194 // as the corresponding primary-text binding. Absent from v2, which let two
195 // programs with different binding layouts share one compiled-pipeline entry.
196 scratch.extend_from_slice(&buffer.binding().to_le_bytes());
197
198 crate::serial::wire::append_data_type_fingerprint(scratch, &buffer.element()).map_err(
199 |message| {
200 format!(
201 "failed to fingerprint pipeline-cache buffer data type `{name}`: {message}. Fix: validate and normalize the Program before computing a compiled-pipeline cache key; invalid IR must not enter cache identity."
202 )
203 },
204 )?;
205
206 // Fixed width so the count lane can never shift the following bytes:
207 // zero means "no static element count", matching how vyre_lower maps a
208 // zero count to `element_count: None`.
209 let static_count = if buffer.has_static_element_count() {
210 buffer.count()
211 } else {
212 0
213 };
214 scratch.extend_from_slice(&static_count.to_le_bytes());
215 Ok(())
216}
217
218/// Stable cache-identity tag for a memory tier.
219///
220/// Single owner for both this digest and `buffer_decl_canonical_key`. The match
221/// is exhaustive on purpose: a new `MemoryKind` must fail to compile here so the
222/// author decides its cache-identity tag instead of inheriting a wildcard.
223pub(super) const fn memory_kind_cache_tag(kind: MemoryKind) -> u8 {
224 match kind {
225 MemoryKind::Global => 0,
226 MemoryKind::Shared => 1,
227 MemoryKind::Uniform => 2,
228 MemoryKind::Local => 3,
229 MemoryKind::Readonly => 4,
230 MemoryKind::Persistent => 5,
231 MemoryKind::Push => 6,
232 }
233}