Skip to main content

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}