vyre_primitives/text/
char_class.rs

1//! Tier 2.5 byte classifier  -  the canonical char-class primitive.
2//!
3//! Each invocation classifies one source byte by loading a host-supplied
4//! 256-entry lookup table from the `table` buffer. The table stays in data
5//! rather than code so alternate classifier sets can be swapped in without
6//! rebuilding the crate.
7//!
8//! Tier 3 dialects call this builder and may register wrapper ops
9//! with their own ids. This primitive keeps its own Tier 2.5 id so
10//! op coverage and composition audits can distinguish the reusable
11//! substrate from user-facing library wrappers.
12
13use vyre_foundation::ir::{BufferAccess, BufferDecl, DataType, Expr, Node, Program};
14
15/// `\0`
16pub const C_EOF: u32 = 0;
17/// Space or tab.
18pub const C_WS: u32 = 1;
19/// `\n` or `\r`
20pub const C_NEWLINE: u32 = 2;
21/// `A-Z`, `a-z`, `_`
22pub const C_ALPHA: u32 = 3;
23/// `0-9`
24pub const C_DIGIT: u32 = 4;
25/// `(`
26pub const C_OPEN_PAREN: u32 = 5;
27/// `)`
28pub const C_CLOSE_PAREN: u32 = 6;
29/// `{`
30pub const C_OPEN_BRACE: u32 = 7;
31/// `}`
32pub const C_CLOSE_BRACE: u32 = 8;
33/// `;`
34pub const C_SEMICOLON: u32 = 9;
35/// `,`
36pub const C_COMMA: u32 = 10;
37/// `.`
38pub const C_DOT: u32 = 11;
39/// `*`
40pub const C_STAR: u32 = 12;
41/// `+`
42pub const C_PLUS: u32 = 13;
43/// `-`
44pub const C_MINUS: u32 = 14;
45/// `/`
46pub const C_SLASH: u32 = 15;
47/// `#`
48pub const C_HASH: u32 = 16;
49/// `'`
50pub const C_QUOTE: u32 = 17;
51/// `"`
52pub const C_DQUOTE: u32 = 18;
53/// `=`
54pub const C_EQUALS: u32 = 19;
55/// `<`
56pub const C_LT: u32 = 20;
57/// `>`
58pub const C_GT: u32 = 21;
59/// `!`
60pub const C_BANG: u32 = 22;
61/// `&`
62pub const C_AMP: u32 = 23;
63/// `|`
64pub const C_PIPE: u32 = 24;
65/// `^`
66pub const C_CARET: u32 = 25;
67/// `~`
68pub const C_TILDE: u32 = 26;
69/// `%`
70pub const C_PERCENT: u32 = 27;
71/// `\`
72pub const C_BACKSLASH: u32 = 28;
73/// `[`
74pub const C_OPEN_BRACKET: u32 = 29;
75/// `]`
76pub const C_CLOSE_BRACKET: u32 = 30;
77/// Anything else.
78pub const C_OTHER: u32 = 31;
79
80/// Stable op id for the registered Tier 2.5 primitive.
81pub const CHAR_CLASS_OP_ID: &str = "vyre-primitives::text::char_class";
82/// Byte-lane workgroup used by the table-driven classifier.
83pub const CHAR_CLASS_WORKGROUP_SIZE: [u32; 3] = [256, 1, 1];
84
85/// Dispatch grid for classifying `n` byte lanes.
86#[must_use]
87pub const fn char_class_dispatch_grid(n: u32) -> [u32; 3] {
88    let blocks = n.div_ceil(CHAR_CLASS_WORKGROUP_SIZE[0]);
89    if blocks == 0 {
90        [1, 1, 1]
91    } else {
92        [blocks, 1, 1]
93    }
94}
95
96/// Build the default ASCII byte-classification table.
97#[must_use]
98pub fn build_char_class_table() -> [u32; 256] {
99    let mut table = [C_OTHER; 256];
100
101    table[0] = C_EOF;
102    table[usize::from(b' ')] = C_WS;
103    table[usize::from(b'\t')] = C_WS;
104    table[usize::from(b'\n')] = C_NEWLINE;
105    table[usize::from(b'\r')] = C_NEWLINE;
106    table[usize::from(b'(')] = C_OPEN_PAREN;
107    table[usize::from(b')')] = C_CLOSE_PAREN;
108    table[usize::from(b'{')] = C_OPEN_BRACE;
109    table[usize::from(b'}')] = C_CLOSE_BRACE;
110    table[usize::from(b';')] = C_SEMICOLON;
111    table[usize::from(b',')] = C_COMMA;
112    table[usize::from(b'.')] = C_DOT;
113    table[usize::from(b'*')] = C_STAR;
114    table[usize::from(b'+')] = C_PLUS;
115    table[usize::from(b'-')] = C_MINUS;
116    table[usize::from(b'/')] = C_SLASH;
117    table[usize::from(b'#')] = C_HASH;
118    table[usize::from(b'\'')] = C_QUOTE;
119    table[usize::from(b'"')] = C_DQUOTE;
120    table[usize::from(b'=')] = C_EQUALS;
121    table[usize::from(b'<')] = C_LT;
122    table[usize::from(b'>')] = C_GT;
123    table[usize::from(b'!')] = C_BANG;
124    table[usize::from(b'&')] = C_AMP;
125    table[usize::from(b'|')] = C_PIPE;
126    table[usize::from(b'^')] = C_CARET;
127    table[usize::from(b'~')] = C_TILDE;
128    table[usize::from(b'%')] = C_PERCENT;
129    table[usize::from(b'\\')] = C_BACKSLASH;
130    table[usize::from(b'[')] = C_OPEN_BRACKET;
131    table[usize::from(b']')] = C_CLOSE_BRACKET;
132    table[usize::from(b'_')] = C_ALPHA;
133
134    for byte in b'0'..=b'9' {
135        table[usize::from(byte)] = C_DIGIT;
136    }
137    for byte in b'A'..=b'Z' {
138        table[usize::from(byte)] = C_ALPHA;
139    }
140    for byte in b'a'..=b'z' {
141        table[usize::from(byte)] = C_ALPHA;
142    }
143
144    table
145}
146
147fn char_class_body(source: &str, classified: &str, n: u32) -> Vec<Node> {
148    vec![Node::Region {
149        generator: vyre_foundation::ir::model::expr::Ident::from(CHAR_CLASS_OP_ID),
150        source_region: None,
151        body: std::sync::Arc::new(vec![
152            Node::let_bind("idx", Expr::InvocationId { axis: 0 }),
153            Node::if_then(
154                Expr::lt(Expr::var("idx"), Expr::u32(n)),
155                vec![Node::store(
156                    classified,
157                    Expr::var("idx"),
158                    // Canonical masked source-byte → 256-table lookup (ONE-PLACE:
159                    // crate::ir_safe), widens the source byte to u32 and masks the
160                    // table index with `& 0xFF` so a >255 element can't read past it.
161                    crate::ir_safe::source_byte_table_lookup("table", source, Expr::var("idx")),
162                )],
163            ),
164        ]),
165    }]
166}
167
168/// Build a Program that writes one character-class code per source byte.
169///
170/// This compatibility entry point expects one `DataType::U32` element per
171/// source byte and reads the low byte of each word. Use [`char_class_u8`] when
172/// the source is packed as one byte per element. `table` is loaded from a
173/// host-provided buffer named `"table"`.
174#[must_use]
175pub fn char_class(source: &str, classified: &str, n: u32) -> Program {
176    char_class_with_source_type(source, classified, n, DataType::U32)
177}
178
179/// Build a Program that writes one character-class code per packed source byte.
180///
181/// It emits the same class stream as [`char_class`] while cutting source input
182/// bandwidth from four bytes per logical byte to one.
183#[must_use]
184pub fn char_class_u8(source: &str, classified: &str, n: u32) -> Program {
185    char_class_with_source_type(source, classified, n, DataType::U8)
186}
187
188fn char_class_with_source_type(
189    source: &str,
190    classified: &str,
191    n: u32,
192    source_type: DataType,
193) -> Program {
194    let output_byte_len = usize::try_from(n).unwrap_or(usize::MAX).saturating_mul(4);
195    Program::wrapped(
196        vec![
197            BufferDecl::storage(source, 0, BufferAccess::ReadOnly, source_type).with_count(n),
198            BufferDecl::storage("table", 1, BufferAccess::ReadOnly, DataType::U32).with_count(256),
199            BufferDecl::output(classified, 2, DataType::U32)
200                .with_count(n.max(1))
201                .with_output_byte_range(0..output_byte_len),
202        ],
203        CHAR_CLASS_WORKGROUP_SIZE,
204        char_class_body(source, classified, n),
205    )
206}
207
208/// Reference oracle: classify each source byte through the lookup table.
209///
210/// Pure function, exposed for fixture generation + harness oracles.
211#[must_use]
212#[cfg(any(test, feature = "cpu-parity", feature = "text"))]
213pub fn reference_char_class(source: &[u8], table: &[u32; 256]) -> Vec<u32> {
214    source
215        .iter()
216        .map(|byte| table[usize::from(*byte)])
217        .collect()
218}
219
220/// Pack a `[u32]` slice into the LE-byte layout the harness uses.
221#[must_use]
222pub fn pack_u32(words: &[u32]) -> Vec<u8> {
223    crate::wire::pack_u32_slice(words)
224}
225
226/// Pack a `[u8]` source slice into the per-element u32 layout the GPU
227/// kernel expects (each byte in the low 8 bits of a u32 lane).
228#[must_use]
229pub fn pack_bytes_as_u32(bytes: &[u8]) -> Vec<u8> {
230    crate::wire::pack_bytes_as_u32_slice(bytes)
231}
232
233#[cfg(feature = "inventory-registry")]
234inventory::submit! {
235    vyre_foundation::operation::OperationRegistration::primitive(
236        CHAR_CLASS_OP_ID,
237        || char_class("source", "classified", 3),
238        Some(|| {
239            let table = build_char_class_table();
240            vec![vec![
241                pack_bytes_as_u32(b"A1 "),
242                pack_u32(&table),
243                vec![0u8; 3 * 4],
244            ]]
245        }),
246        Some(|| {
247            vec![vec![pack_u32(&[C_ALPHA, C_DIGIT, C_WS])]]
248        }),
249    )
250}
251
252#[cfg(test)]
253mod tests {
254    use super::*;
255
256    #[test]
257    fn table_classifies_ascii_letter_as_alpha() {
258        let table = build_char_class_table();
259        assert_eq!(table[usize::from(b'A')], C_ALPHA);
260        assert_eq!(table[usize::from(b'z')], C_ALPHA);
261        assert_eq!(table[usize::from(b'_')], C_ALPHA);
262    }
263
264    #[test]
265    fn table_classifies_digits() {
266        let table = build_char_class_table();
267        for byte in b'0'..=b'9' {
268            assert_eq!(table[usize::from(byte)], C_DIGIT);
269        }
270    }
271
272    #[test]
273    fn reference_walks_table() {
274        let table = build_char_class_table();
275        assert_eq!(
276            reference_char_class(b"A1 ", &table),
277            vec![C_ALPHA, C_DIGIT, C_WS]
278        );
279    }
280
281    #[test]
282    fn reference_covers_every_byte_value() {
283        let table = build_char_class_table();
284        let source: Vec<u8> = (0u8..=255).collect();
285        assert_eq!(reference_char_class(&source, &table), table.to_vec());
286    }
287
288    #[test]
289    fn program_uses_block_sized_workgroup() {
290        let program = char_class("source", "classified", 513);
291        assert_eq!(program.workgroup_size(), CHAR_CLASS_WORKGROUP_SIZE);
292    }
293
294    #[test]
295    fn packed_u8_program_declares_one_source_byte_per_element() {
296        let program = char_class_u8("source", "classified", 513);
297        let source = program
298            .buffers()
299            .iter()
300            .find(|buffer| buffer.name() == "source")
301            .expect("Fix: packed-u8 char_class source buffer must be declared");
302        let classified = program
303            .buffers()
304            .iter()
305            .find(|buffer| buffer.name() == "classified")
306            .expect("Fix: char_class output buffer must be declared");
307
308        assert_eq!(source.element(), DataType::U8);
309        assert_eq!(source.count(), 513);
310        assert_eq!(classified.element(), DataType::U32);
311        assert_eq!(classified.count(), 513);
312        assert_eq!(classified.output_byte_range(), Some(0..513 * 4));
313        assert_eq!(program.workgroup_size(), CHAR_CLASS_WORKGROUP_SIZE);
314    }
315
316    #[test]
317    fn empty_program_declares_empty_output_range() {
318        let program = char_class_u8("source", "classified", 0);
319        let classified = program
320            .buffers()
321            .iter()
322            .find(|buffer| buffer.name() == "classified")
323            .expect("Fix: char_class output buffer must be declared");
324
325        assert_eq!(classified.count(), 1);
326        assert_eq!(classified.output_byte_range(), Some(0..0));
327    }
328
329    #[test]
330    fn dispatch_grid_packs_byte_lanes_into_blocks() {
331        assert_eq!(char_class_dispatch_grid(0), [1, 1, 1]);
332        assert_eq!(char_class_dispatch_grid(1), [1, 1, 1]);
333        assert_eq!(char_class_dispatch_grid(256), [1, 1, 1]);
334        assert_eq!(char_class_dispatch_grid(257), [2, 1, 1]);
335        assert_eq!(char_class_dispatch_grid(513), [3, 1, 1]);
336    }
337
338    #[test]
339    fn primitive_id_names_the_primitive_tier() {
340        assert_eq!(CHAR_CLASS_OP_ID, "vyre-primitives::text::char_class");
341    }
342}
vyre_primitives/text/char_class.rs

vyre_primitives/text/
char_class.rs
