vyre_libs/parsing/c/preprocess/
gpu_include_parse.rs

1//! GPU `#include` row parser.
2//!
3//! Phase 17b.7: per `TOK_PREPROC` token classified as
4//! `TOK_PP_INCLUDE` or `TOK_PP_INCLUDE_NEXT`, extract the include
5//! path's byte span and whether it was the `<…>` (system) form or
6//! `"…"` (local) form. Per-thread, fully parallel.
7//!
8//! ## Output columns (one row per token)
9//!
10//! - `path_start`, `path_len`         -  byte span between the
11//!                                     delimiters (`<`/`>` or `"`/`"`).
12//! - `is_system`                      -  `1` for `<…>`, `0` for `"…"`.
13//!
14//! Non-INCLUDE rows get all-zero output. `path_len == 0` after this
15//! kernel means "not a parsed `#include` row"  -  equivalent to the CPU
16//! `parse_include_literal` returning `None`/error.
17//!
18//! ## Real-GPU lowering note
19//!
20//! Two real-GPU lowering pitfalls (both shared with
21//! `gpu_directive_metadata`):
22//!
23//! 1. `DataType::U8` storage buffers are emitted by vyre-emit-naga as
24//!    `array<u32>` (WGSL has no u8 storage). `Expr::load("source",
25//!    addr)` therefore returns the u32 word at index `addr`, not the
26//!    byte at byte-address `addr`. The kernel does the byte
27//!    extraction inline so it produces the correct value on both
28//!    backends.
29//! 2. Whitespace skipping uses fixed-depth chained Selects because C
30//!    directive separators are short in practice. Path extraction is
31//!    bounded by the directive row length, so Linux-scale include paths
32//!    are not truncated by a compile-time probe cap.
33//!
34//! ## Wire layout
35//!
36//! Inputs:
37//!   - `tok_starts` (U32), `tok_lens` (U32),
38//!     `directive_kinds` (U32)  -  output of 17a.
39//!   - `source` (U8).
40//!
41//! Outputs (all U32, one element per token):
42//!   - `path_start_out`, `path_len_out`, `is_system_out`.
43
44use super::gpu_directive_parse_shared::{
45    directive_program_from_parse_with_source_layout, push_bounded_byte_scan_until,
46    push_directive_row_bounds, push_hash_and_keyword_start, push_keyword_end,
47    push_ws_skip_from_expr, safe_source_byte_expr, DirectiveOutputColumn, DirectiveSourceLayout,
48    DirectiveThreadLayout,
49};
50use crate::parsing::c::lex::tokens::{TOK_PP_INCLUDE, TOK_PP_INCLUDE_NEXT};
51use vyre::ir::{Expr, Node, Program};
52
53/// Canonical op id.
54pub const OP_ID: &str = "vyre-libs::parsing::c::preprocess::gpu_include_parse_v2";
55
56/// Canonical binding for the input per-token start-offset buffer.
57pub const BINDING_TOK_STARTS: u32 = 0;
58/// Canonical binding for the input per-token length buffer.
59pub const BINDING_TOK_LENS: u32 = 1;
60/// Canonical binding for the input directive-kinds buffer.
61pub const BINDING_DIRECTIVE_KINDS: u32 = 2;
62/// Canonical binding for the input source bytes.
63pub const BINDING_SOURCE: u32 = 3;
64/// Canonical binding for the output `path_start` column.
65pub const BINDING_PATH_START_OUT: u32 = 4;
66/// Canonical binding for the output `path_len` column.
67pub const BINDING_PATH_LEN_OUT: u32 = 5;
68/// Canonical binding for the output `is_system` column.
69pub const BINDING_IS_SYSTEM_OUT: u32 = 6;
70
71const OUTPUT_COLUMNS: [DirectiveOutputColumn; 3] = [
72    DirectiveOutputColumn {
73        name: "path_start_out",
74        binding: BINDING_PATH_START_OUT,
75    },
76    DirectiveOutputColumn {
77        name: "path_len_out",
78        binding: BINDING_PATH_LEN_OUT,
79    },
80    DirectiveOutputColumn {
81        name: "is_system_out",
82        binding: BINDING_IS_SYSTEM_OUT,
83    },
84];
85
86/// Build the 17b.7 `#include` row parser `Program`.
87///
88/// Hybrid runtime/static-bound: kernel BODY uses `Expr::buf_len()` for
89/// every per-thread bound (so program AST is constant across files),
90/// `num_tokens` is kept ONLY for output buffer sizing (CUDA backend
91/// requires static byte length for readback), `source_len` is unused.
92#[must_use]
93pub fn gpu_include_parse(num_tokens: u32, source_len: u32) -> Program {
94    gpu_include_parse_with_source_layout(num_tokens, source_len, DirectiveSourceLayout::PackedU32)
95}
96
97/// Build the 17b.7 `#include` row parser over raw `DataType::U8` source bytes.
98#[must_use]
99pub fn gpu_include_parse_u8(num_tokens: u32, source_len: u32) -> Program {
100    gpu_include_parse_with_source_layout(num_tokens, source_len, DirectiveSourceLayout::RawU8)
101}
102
103fn gpu_include_parse_with_source_layout(
104    num_tokens: u32,
105    source_len: u32,
106    source_layout: DirectiveSourceLayout,
107) -> Program {
108    let t = Expr::var("t");
109    let safe_load = |addr: Expr| safe_source_byte_expr(source_layout, addr);
110
111    let mut parse: Vec<Node> = Vec::new();
112    push_directive_row_bounds(&mut parse);
113    push_hash_and_keyword_start(&mut parse, source_layout);
114
115    // ---- step 3: skip past keyword. kw_len = 7 (`include`) or 12
116    // (`include_next`). Decided by `kind`. ----
117    parse.push(Node::let_bind(
118        "kw_len_skip",
119        Expr::select(
120            Expr::eq(Expr::var("kind"), Expr::u32(TOK_PP_INCLUDE_NEXT)),
121            Expr::u32(12),
122            Expr::u32(7),
123        ),
124    ));
125    push_keyword_end(&mut parse, Expr::var("kw_len_skip"));
126
127    // ---- step 4: skip WS between keyword and delimiter. ----
128    push_ws_skip_from_expr(
129        &mut parse,
130        source_layout,
131        "dp",
132        Expr::var("post_kw"),
133        "delim_skip",
134        "delim_pos",
135    );
136
137    // ---- step 5: classify delimiter. ----
138    parse.push(Node::let_bind(
139        "delim_byte",
140        safe_load(Expr::var("delim_pos")),
141    ));
142    parse.push(Node::let_bind(
143        "is_angle",
144        Expr::select(
145            Expr::eq(Expr::var("delim_byte"), Expr::u32(b'<' as u32)),
146            Expr::u32(1),
147            Expr::u32(0),
148        ),
149    ));
150    parse.push(Node::let_bind(
151        "is_quote",
152        Expr::select(
153            Expr::eq(Expr::var("delim_byte"), Expr::u32(b'"' as u32)),
154            Expr::u32(1),
155            Expr::u32(0),
156        ),
157    ));
158    parse.push(Node::let_bind(
159        "valid_delim",
160        Expr::select(
161            Expr::or(
162                Expr::eq(Expr::var("is_angle"), Expr::u32(1)),
163                Expr::eq(Expr::var("is_quote"), Expr::u32(1)),
164            ),
165            Expr::u32(1),
166            Expr::u32(0),
167        ),
168    ));
169    parse.push(Node::let_bind(
170        "close_byte",
171        Expr::select(
172            Expr::eq(Expr::var("is_angle"), Expr::u32(1)),
173            Expr::u32(b'>' as u32),
174            Expr::u32(b'"' as u32),
175        ),
176    ));
177    parse.push(Node::let_bind(
178        "path_start_val",
179        Expr::add(Expr::var("delim_pos"), Expr::u32(1)),
180    ));
181
182    // ---- step 6: scan path bytes to the directive row end for the
183    // closing delimiter. This used to be a fixed 48-byte unrolled
184    // probe, which silently rejected long Linux/generated include
185    // paths. The row-length loop keeps the program shape constant but
186    // removes the semantic cap.
187    push_bounded_byte_scan_until(
188        &mut parse,
189        source_layout,
190        "path_i",
191        "path_start_val",
192        "path_scan_limit",
193        "path_byte",
194        "path_len_val",
195        "path_done",
196        Expr::var("close_byte"),
197        Expr::eq(Expr::u32(1), Expr::u32(1)),
198    );
199
200    // ---- step 7: commit if found_hash AND valid_delim ----
201    // Both are u32 0/1; bitand stays u32; convert to bool for if_then.
202    parse.push(Node::if_then(
203        Expr::eq(
204            Expr::bitand(
205                Expr::bitand(Expr::var("found_hash"), Expr::var("valid_delim")),
206                Expr::var("path_done"),
207            ),
208            Expr::u32(1),
209        ),
210        vec![
211            Node::store("path_start_out", t.clone(), Expr::var("path_start_val")),
212            Node::store("path_len_out", t.clone(), Expr::var("path_len_val")),
213            Node::store("is_system_out", t.clone(), Expr::var("is_angle")),
214        ],
215    ));
216
217    directive_program_from_parse_with_source_layout(
218        OP_ID,
219        num_tokens,
220        source_len,
221        source_layout,
222        &OUTPUT_COLUMNS,
223        DirectiveThreadLayout::WorkgroupLinear,
224        Expr::or(
225            Expr::eq(Expr::var("kind"), Expr::u32(TOK_PP_INCLUDE)),
226            Expr::eq(Expr::var("kind"), Expr::u32(TOK_PP_INCLUDE_NEXT)),
227        ),
228        parse,
229    )
230}
231
232#[cfg(test)]
233mod tests {
234    use super::*;
235    use vyre::ir::DataType;
236
237    #[test]
238    fn op_id_is_canonical_and_stable() {
239        assert_eq!(
240            OP_ID,
241            "vyre-libs::parsing::c::preprocess::gpu_include_parse_v2"
242        );
243    }
244
245    #[test]
246    fn build_program_returns_well_formed_program() {
247        let p = gpu_include_parse(8, 64);
248        assert_eq!(p.buffers().len(), 7);
249        assert_eq!(p.workgroup_size(), [256, 1, 1]);
250    }
251
252    #[test]
253    fn source_buffer_layouts_preserve_packed_abi_and_raw_u8_variant() {
254        let packed = gpu_include_parse(8, 64);
255        let raw_u8 = gpu_include_parse_u8(8, 64);
256        let packed_source = packed
257            .buffers()
258            .iter()
259            .find(|buffer| buffer.name() == "source")
260            .expect("Fix: packed include parser source buffer must exist");
261        let raw_u8_source = raw_u8
262            .buffers()
263            .iter()
264            .find(|buffer| buffer.name() == "source")
265            .expect("Fix: raw-U8 include parser source buffer must exist");
266
267        assert_eq!(packed_source.element(), DataType::U32);
268        assert_ne!(packed_source.count(), 0);
269        assert_eq!(raw_u8_source.element(), DataType::U8);
270        assert_eq!(raw_u8_source.count(), 0);
271    }
272}
vyre_libs/parsing/c/preprocess/gpu_include_parse.rs

vyre_libs/parsing/c/preprocess/
gpu_include_parse.rs
