nx_wgsl_front_candidate_t217.nx source
↩ module page · 3423 lines · 200915 B
1// nx_wgsl.nx -- THE SHADER BACKENDS: NishiLang in, GLSL and WGSL out.
2//
3// CONTRACT SYMBOL: wgsl_emit_module (declared _ABSENT_ on the graphics compare board until now).
4//
5// THE SOURCE IS NISHILANG, NOT GLSL (operator 2026-08-23). This organ is NOT a GLSL->WGSL
6// transpiler. A shader is authored ONCE against the nx_shader_ir representation; glsl_emit_module
7// and wgsl_emit_module are two BACKENDS over that one source, exactly as nx_compile_x86 and the
8// WAT lane are two backends over one NishiLang front-end. Neither dialect is authored, so neither
9// can drift from the other -- disagreement is impossible BY CONSTRUCTION rather than by discipline.
10//
11// RUNG 1, SHIPPED HERE: the fullscreen-triangle VERTEX STAGE. It is not a toy -- it is the exact
12// stage that ships TODAY, hand-written TWICE in nx_game_page_emit:
13// GLSL VSH : void main(){vec2 v=vec2((gl_VertexID<<1)&2,gl_VertexID&2);gl_Position=vec4(v*2.-1.,0.,1.);}
14// WGSL vs : @vertex fn vs(@builtin(vertex_index) vi:u32)->@builtin(position) vec4f{...}
15// Those two hand copies are the duplication this rung deletes: below, ONE NishiLang source
16// (shsrc_fullscreen_tri) emits both.
17//
18// WHY THIS STAGE PROVES THE THESIS. The two dialects are not spelling variants of each other:
19// * GLSL ASSIGNS gl_Position; WGSL RETURNS @builtin(position). Different STRUCTURE.
20// * GLSL gl_VertexID is i32; WGSL vertex_index is u32 and must be converted before i32 maths.
21// * WGSL requires a u32 shift amount (1u); GLSL takes a plain int.
22// * WGSL requires 0.0, never 0. ; GLSL accepts either.
23// Every one of those is dialect knowledge the BACKEND owns. A translator would have to rediscover
24// them from GLSL text; a backend simply knows them, once.
25//
26// COVERED SUBSET (rung 1) -- anything outside REFUSES BY NAME via sir_refuse and returns -1.
27// A NAMED REFUSAL IS A CONTRACT; A SILENT MISTRANSLATION IN A SHADER BACKEND IS THE WORST
28// AVAILABLE FAILURE: it compiles, it links, it runs, and it draws the wrong picture forever.
29// decls : uniform (scalar + array); attribute and varying are GLSL-covered and, since S1
30// (2026-09-04), WGSL-REFUSED BY NAME (K_ATTRIB anywhere, K_VARY in a fragment stage) until
31// the entry-parameter derivation that WGSL needs for them lands; a texture-typed uniform
32// refuses in WGSL too -- it was being emitted INSIDE struct U
33// S3 (2026-09-04): a TEXTURE is its own kind (K_TEXTURE / sir_texture) in BOTH backends: GLSL
34// `uniform highp usampler3D`, WGSL a @group/@binding var OUTSIDE struct U with its binding
35// DERIVED from declaration order; only texture_3d<u32> is covered, other texture types refuse
36// S4 (2026-09-04): the FRAGMENT entry signature is DERIVED from the declared parameters (no
37// longer a literal): a sir_param_position parameter is `@builtin(position) name:vec4f` in WGSL and
38// gl_FragCoord in GLSL; a position READ with no such parameter, and any entry parameter a backend
39// cannot bind (plain fragment params, any vertex/compute param), refuse by name
40// stmts : var, assign, return, if/else, discard
41// exprs : literal, identifier, binary op, call, constructor, swizzle, index, cast, builtin,
42// texture load, texture sample
43// types : f32 i32 u32 bool vec2f vec3f vec4f vec3i texture_3d<u32> texture_2d<f32>
44// NOT covered (each REFUSES, and each is a named rung): matrices, structs, arrays of vectors as
45// locals, derivatives. Loops/break landed as rung 2b-general.
46//
47// RUNG GE43 (gameengine board, 2026-09-02) -- THE COMPUTE STAGE, WGSL-DOOR ONLY. Storage buffers
48// (read / read_write, optionally atomic<>), atomic read-modify-write (atomicMin/Max/Add, the
49// backend supplies the `&`), the global_invocation_id builtin, and the @compute entry with its
50// @workgroup_size. This is the first byte of the ONE NishiLang rasterizer kernel the board declares
51// (GE42 rk_kernel_ir -> GE43 wgsl_compute_stage -> GE44 rk_webgpu_door -> GE45 rk_tier_ladder):
52// the kernel is authored HERE, once, and reaches a third-party browser through WebGPU as a pure
53// submission pipe, and NishiOS through the dxg lane. GLSL ES 3.00 (WebGL2) has no compute stage,
54// no storage buffers and no atomics, so the GLSL backend REFUSES every one of these BY NAME --
55// `nx_wgsl compute` prints both outcomes and nx_wgsl_compute_gate asserts the bytes exactly.
56
57import "nx_syscalls.nx"
58import "nx_shader_ir.nx"
59import "nx_cast_shader_src.nx" // S12c (2026-09-05): THE CHARACTER CAST vertex stage written once, transcribed slice by slice and measured against the hand MVS by nx_glsl_tokdiff (verb: cast)
60import "nx_world_shader_src.nx" // GE44 R-G (2026-09-04): THE WORLD SHADER written once; `nx_wgsl world` emits it to both dialects
61// RUNG 2 -- SHARE THE FRONT-END, DO NOT FORK IT. These are exactly the imports nx_compile_wat.nx
62// uses for the WAT backend. A shader is compiled by the SAME lexer, the SAME parser and the SAME
63// optimiser as every other NishiLang program; this organ is one more BACKEND hanging off that
64// front-end, beside nx_compile_x86 and nx_wasm. Two parsers for one language would be the
65// duplicate-ruler defect at the worst possible layer, so there is not a second one anywhere here.
66import "nx_types.nx"
67import "nx_lex_kinds.nx"
68import "nx_ir.nx"
69import "nx_tokenizer.nx"
70import "nx_parse.nx"
71import "nx_opt.nx"
72import "nx_import.nx"
73
74// STAGE_PLAIN is rung 2: an ORDINARY NishiLang function, compiled through the REAL front-end
75// (expand_imports -> lex_source -> parse_module -> opt_run) and emitted as a plain shader
76// function with its own parameters and return type. Rung 1's stages wrap an entry point; this
77// one wraps nothing, because the source is the language itself.
78const STAGE_PLAIN: i64 = 0
79const STAGE_VERTEX: i64 = 1
80const STAGE_FRAGMENT: i64 = 2
81// GE43: the compute stage. WGSL only -- glsl_emit_module refuses it by name before emitting a byte.
82const STAGE_COMPUTE: i64 = 3
83// Invocations per workgroup on the compute entry. 64 is the width both vendor wavefronts divide
84// evenly (AMD wave64, NVIDIA warp32) and one quarter of WebGPU's default
85// maxComputeInvocationsPerWorkgroup of 256 (the limit is banked as a GE42 lesson row on the
86// gameengine board). Named here so a kernel author never spells the geometry by hand.
87const WG_WORKGROUP: i64 = 64
88
89// Buffer bounds, DERIVED from the IR's own cap rather than picked -- ONE OWNER for the number.
90// A trace records at most one entry per IR node, so its ceiling IS the node ceiling.
91const WG_TRACE_CAP: i64 = SIR_MAXNODE
92// The emitted-text buffer. Every node in the covered subset emits a bounded run of characters --
93// the longest single emission is a type name plus punctuation -- so 16 bytes per representable
94// node is a ceiling no module the IR can hold is able to exceed.
95const WG_EMIT_CAP: i64 = SIR_MAXNODE * 16
96// S12c-4: the '!' byte, constructed because the nx_cc lexer refuses a bang inside a string literal
97const WG_C_BANG: i64 = 33
98// Decimal rendering of an i64 needs at most 20 digits plus sign plus terminator.
99const WG_NUMBUF: i64 = 32
100// R-A: the select operands are emitted in IR order and rotated in place by wg_rev -- no scratch buffer, nothing
101// to size (a WG_SEL_SCRATCH constant lived here for one build on 2026-09-04 and was retired before it was used).
102// A REFUSAL MESSAGE IS NOT A NUMBER, AND SIZING IT LIKE ONE TRUNCATES THE ACTIONABLE HALF.
103// Measured on this organ's own first rung-2 run: the refusal printed "opcode outside the covered
104// subs" -- the opcode NUMBER, the only part a reader can act on, fell off the end of a 32-byte
105// buffer named for decimal digits. A named refusal that cannot finish naming is the contract
106// failing in the one place it exists to hold.
107const WG_MSGBUF: i64 = 128
108// S3 (2026-09-04): THE WGSL BINDING LAYOUT, owned here and nowhere else. Every resource lives in ONE bind
109// group. struct U (all uniforms) sits at WG_BIND_UNIFORM whether or not a module declares a uniform, so
110// adding one later never renumbers the resources beside it. Textures take the bindings after struct U in
111// DECLARATION ORDER (first K_TEXTURE = WG_BIND_TEX_FIRST, then +1, ...). A storage buffer carries an
112// EXPLICIT binding (sir_storage); a derived texture binding that lands on one REFUSES by name rather than
113// emitting two vars on one slot.
114const WG_BIND_GROUP: i64 = 0
115// S12 step 2 (2026-09-06): the group a module's resources are spelled in = the base group plus the module's own word
116// (M_BINDGROUP, 0 unless a verb sets it). Every @group( emitter reads THIS, so a module moved to another group moves
117// its struct U, textures, samplers and storage buffers together -- one speller, no site left behind.
118func wg_group(m: *i64) -> i64 { return WG_BIND_GROUP + m[M_BINDGROUP] }
119const WG_BIND_UNIFORM: i64 = 0
120const WG_BIND_TEX_FIRST: i64 = WG_BIND_UNIFORM + 1
121
122// ---- dialect type spelling ----------------------------------------------------------------
123func gl_ty(t: i64) -> *u8 {
124 if t == T_VOID { return "void" as *u8 }
125 if t == T_F32 { return "float" as *u8 }
126 if t == T_I32 { return "int" as *u8 }
127 if t == T_U32 { return "uint" as *u8 }
128 if t == T_BOOL { return "bool" as *u8 }
129 if t == T_V2F { return "vec2" as *u8 }
130 if t == T_V3F { return "vec3" as *u8 }
131 if t == T_V4F { return "vec4" as *u8 }
132 if t == T_V3I { return "ivec3" as *u8 }
133 if t == T_V2I { return "ivec2" as *u8 }
134 if t == T_M3F { return "mat3" as *u8 }
135 if t == T_TEX3U { return "highp usampler3D" as *u8 }
136 if t == T_TEX2F { return "sampler2D" as *u8 }
137 if t == T_TEX2F { return "sampler2D" as *u8 }
138 if t == T_V3U { return "uvec3" as *u8 }
139 return 0 as *u8
140}
141
142func wg_ty(t: i64) -> *u8 {
143 if t == T_VOID { return "void" as *u8 }
144 if t == T_F32 { return "f32" as *u8 }
145 if t == T_I32 { return "i32" as *u8 }
146 if t == T_U32 { return "u32" as *u8 }
147 if t == T_BOOL { return "bool" as *u8 }
148 if t == T_V2F { return "vec2f" as *u8 }
149 if t == T_V3F { return "vec3f" as *u8 }
150 if t == T_V4F { return "vec4f" as *u8 }
151 if t == T_V3I { return "vec3i" as *u8 }
152 if t == T_V2I { return "vec2i" as *u8 }
153 if t == T_M3F { return "mat3x3f" as *u8 }
154 if t == T_TEX3U { return "texture_3d<u32>" as *u8 }
155 if t == T_TEX2F { return "texture_2d<f32>" as *u8 }
156 if t == T_TEX2F { return "texture_2d<f32>" as *u8 }
157 if t == T_V3U { return "vec3u" as *u8 }
158 return 0 as *u8
159}
160
161// R-D (2026-09-04): a STRUCT type is spelled by its declaration's name in BOTH dialects; scalars and vectors keep
162// each dialect's table. ONE lookup (sir_struct_of) decides which declaration a type names, so the two backends
163// cannot disagree about it. A caller that spells a type through these never meets a struct as "outside the
164// covered subset".
165func gl_tyname(m: *i64, t: i64) -> *u8 {
166 let st: i64 = sir_struct_of(m, t)
167 if st != 0 { return sir_cstr(m, sir_get(m, st, SIR_NAME)) }
168 return gl_ty(t)
169}
170func wg_tyname(m: *i64, t: i64) -> *u8 {
171 let st: i64 = sir_struct_of(m, t)
172 if st != 0 { return sir_cstr(m, sir_get(m, st, SIR_NAME)) }
173 return wg_ty(t)
174}
175
176// A texture type is a BOUND RESOURCE in both dialects (a sampler uniform in GLSL, a @binding var in
177// WGSL), never a value type: it cannot be a struct member, a local, a parameter or a cast target.
178// ONE predicate for that class, so the two backends cannot disagree about which types it holds.
179func wg_is_texty(t: i64) -> i64 {
180 if t == T_TEX3U { return 1 }
181 if t == T_TEX2F { return 1 }
182 return 0
183}
184
185// ---- IDENTIFIER RESOLUTION (S2, 2026-09-04) ----------------------------------------------------
186// An E_IDENT carries only a NAME; what it names decides how each dialect spells it. WGSL folds every
187// uniform into `struct U` behind ONE binding, so a uniform read must be spelled `u.<name>` -- before S2
188// the backend emitted the bare name, an undeclared identifier in WGSL, with no refusal (measured on the
189// S1 positive control: `v=(v*scale);` under a `struct U{scale:f32,}`). Resolution order is the
190// language's own: a parameter or local of the CURRENT function (M_CURFN, set by both backends) shadows a
191// module-scope declaration. Names are pool offsets interned per call, so identity is by CONTENT.
192func wg_name_is(m: *i64, id: i64, name: *u8) -> i64 {
193 return wg_streq(sir_cstr(m, sir_get(m, id, SIR_NAME)), name)
194}
195// the K_PARAM of this name on the function, or 0
196func wg_fn_param(m: *i64, fn: i64, name: *u8) -> i64 {
197 if fn == 0 { return 0 }
198 var p: i64 = sir_get(m, fn, SIR_A)
199 while p != 0 {
200 if wg_name_is(m, p, name) == 1 { return p }
201 p = sir_get(m, p, SIR_NEXT)
202 }
203 return 0
204}
205// S4: the K_PARAM bound to @builtin(position) on the function, or 0
206func wg_fn_position_param(m: *i64, fn: i64) -> i64 {
207 if fn == 0 { return 0 }
208 var p: i64 = sir_get(m, fn, SIR_A)
209 while p != 0 {
210 if sir_get(m, p, SIR_A) == P_BUILTIN_POSITION { return p }
211 p = sir_get(m, p, SIR_NEXT)
212 }
213 return 0
214}
215// 1 iff a statement chain -- recursively through if arms and loop bodies -- declares an S_VAR of this name
216func wg_chain_local(m: *i64, head: i64, name: *u8) -> i64 {
217 var s: i64 = head
218 while s != 0 {
219 let k: i64 = sir_get(m, s, SIR_KIND)
220 if k == S_VAR { if wg_name_is(m, s, name) == 1 { return 1 } }
221 if k == S_IF {
222 if wg_chain_local(m, sir_get(m, s, SIR_B), name) == 1 { return 1 }
223 if wg_chain_local(m, sir_get(m, s, SIR_C), name) == 1 { return 1 }
224 }
225 if k == S_LOOP { if wg_chain_local(m, sir_get(m, s, SIR_A), name) == 1 { return 1 } }
226 s = sir_get(m, s, SIR_NEXT)
227 }
228 return 0
229}
230// the module-scope declaration of this KIND with this name, or 0
231func wg_decl_named(m: *i64, kind: i64, name: *u8) -> i64 {
232 var d: i64 = m[M_DECLH]
233 while d != 0 {
234 if sir_get(m, d, SIR_KIND) == kind { if wg_name_is(m, d, name) == 1 { return d } }
235 d = sir_get(m, d, SIR_NEXT)
236 }
237 return 0
238}
239// 1 iff the name, read inside the function being emitted, resolves to a module-scope K_UNIFORM --
240// i.e. no parameter and no local of the current function shadows it.
241func wg_resolves_to_uniform(m: *i64, name: *u8) -> i64 {
242 let fn: i64 = m[M_CURFN]
243 if wg_fn_param(m, fn, name) != 0 { return 0 }
244 if fn != 0 { if wg_chain_local(m, sir_get(m, fn, SIR_B), name) == 1 { return 0 } }
245 if wg_decl_named(m, K_UNIFORM, name) != 0 { return 1 }
246 return 0
247}
248// S7 (2026-09-05): the varying twin of the resolver above -- a bare identifier that no parameter and no
249// local of the current function shadows, resolving to a module-scope K_VARY. The caller spells the
250// qualifier (vo. in the vertex entry, vin. in the fragment entry) from M_VIOMODE.
251func wg_resolves_to_varying(m: *i64, name: *u8) -> i64 {
252 let fn: i64 = m[M_CURFN]
253 if wg_fn_param(m, fn, name) != 0 { return 0 }
254 if fn != 0 { if wg_chain_local(m, sir_get(m, fn, SIR_B), name) == 1 { return 0 } }
255 if wg_decl_named(m, K_VARY, name) != 0 { return 1 }
256 return 0
257}
258
259// ---- S12c-2b (2026-09-05): A float[N] UNIFORM IS CARRIED AS array<vec4f,ceil(N/4)> -----------------------
260// The uniform address space requires a 16-byte array stride (the layout ruler refuses array<f32,N> by name), and the cast
261// declares uOF/uGB/uGT/uGN as float[12] exactly as its GLSL does. ONE rule, read by the struct-U emitter, the E_INDEX arm and
262// the layout ruler: the member is array<vec4f,ceil(N/4)> and the read is u.name[i/4][i%4]. GLSL is untouched (float name[N],
263// name[i]) so the token proof against the hand text is unaffected. The lane count is DERIVED here, never typed at a use site.
264const WG_F32ARR_LANES: i64 = 4
265func wg_f32arr_vec4n(n: i64) -> i64 { return (n + WG_F32ARR_LANES - 1) / WG_F32ARR_LANES }
266// N when `name` resolves (unshadowed) to a module-scope K_UNIFORM float[N] with N>0; 0 otherwise
267func wg_f32arr_len(m: *i64, name: *u8) -> i64 {
268 if wg_resolves_to_uniform(m, name) == 0 { return 0 }
269 let d: i64 = wg_decl_named(m, K_UNIFORM, name)
270 if d == 0 { return 0 }
271 if sir_get(m, d, SIR_TY) != T_F32 { return 0 }
272 return sir_get(m, d, SIR_A)
273}
274
275// ---- S5 (2026-09-04): THE FRAGMENT OUTPUT ---------------------------------------------------------
276// the FIRST declaration of this kind, or 0
277func wg_decl_kind_first(m: *i64, kind: i64) -> i64 {
278 var d: i64 = m[M_DECLH]
279 while d != 0 {
280 if sir_get(m, d, SIR_KIND) == kind { return d }
281 d = sir_get(m, d, SIR_NEXT)
282 }
283 return 0
284}
285// how many declarations of this kind the module carries
286func wg_decl_kind_count(m: *i64, kind: i64) -> i64 {
287 var n: i64 = 0
288 var d: i64 = m[M_DECLH]
289 while d != 0 {
290 if sir_get(m, d, SIR_KIND) == kind { n = n + 1 }
291 d = sir_get(m, d, SIR_NEXT)
292 }
293 return n
294}
295// the module's declared fragment output, or 0. ONE lookup, so the two backends cannot disagree
296// about which declaration is the output -- the duplicate-ruler defect at the worst possible layer.
297func wg_fragout(m: *i64) -> i64 { return wg_decl_kind_first(m, K_FRAGOUT) }
298// 1 iff a statement chain -- recursively through if arms and loop bodies -- carries an S_RETURN that
299// returns a VALUE. The output requirement binds to what the body DOES, not to the stage: a fragment
300// that only discards is a legal program with no colour output, and a guard that refused it would be a
301// false-positive generator. This is the cheapest exact test, and it reuses wg_chain_local's shape.
302func wg_chain_returns_value(m: *i64, head: i64) -> i64 {
303 var s: i64 = head
304 while s != 0 {
305 let k: i64 = sir_get(m, s, SIR_KIND)
306 if k == S_RETURN { if sir_get(m, s, SIR_A) != 0 { return 1 } }
307 if k == S_IF {
308 if wg_chain_returns_value(m, sir_get(m, s, SIR_B)) == 1 { return 1 }
309 if wg_chain_returns_value(m, sir_get(m, s, SIR_C)) == 1 { return 1 }
310 }
311 if k == S_LOOP { if wg_chain_returns_value(m, sir_get(m, s, SIR_A)) == 1 { return 1 } }
312 s = sir_get(m, s, SIR_NEXT)
313 }
314 return 0
315}
316
317// S12c-5 (2026-09-05): 1 iff the chain (recursing into if/else and loop bodies) ASSIGNS an identifier named `name` --
318// the test that decides whether a fragment entry's declared output is carried as a WGSL local (see M_WGFRAGOUT).
319func wg_chain_assigns(m: *i64, head: i64, name: *u8) -> i64 {
320 var s: i64 = head
321 while s != 0 {
322 let k: i64 = sir_get(m, s, SIR_KIND)
323 if k == S_ASSIGN {
324 let l: i64 = sir_get(m, s, SIR_A)
325 if sir_get(m, l, SIR_KIND) == E_IDENT { if wg_name_is(m, l, name) == 1 { return 1 } }
326 }
327 if k == S_IF {
328 if wg_chain_assigns(m, sir_get(m, s, SIR_B), name) == 1 { return 1 }
329 if wg_chain_assigns(m, sir_get(m, s, SIR_C), name) == 1 { return 1 }
330 }
331 if k == S_LOOP { if wg_chain_assigns(m, sir_get(m, s, SIR_A), name) == 1 { return 1 } }
332 s = sir_get(m, s, SIR_NEXT)
333 }
334 return 0
335}
336// S12c-5: 1 iff the LAST top-level statement of the chain is a return (then no trailing `return fc;` is appended)
337func wg_chain_last_is_return(m: *i64, head: i64) -> i64 {
338 if head == 0 { return 0 }
339 var s: i64 = head
340 while sir_get(m, s, SIR_NEXT) != 0 { s = sir_get(m, s, SIR_NEXT) }
341 if sir_get(m, s, SIR_KIND) == S_RETURN { return 1 }
342 return 0
343}
344
345// ---- BINDING DERIVATION (S3) ----------------------------------------------------------------------
346// the binding a K_TEXTURE gets: WG_BIND_TEX_FIRST plus its ordinal among the K_TEXTURE declarations
347func wg_tex_binding(m: *i64, tex: i64) -> i64 {
348 var ord: i64 = 0
349 var d: i64 = m[M_DECLH]
350 while d != 0 {
351 if d == tex { return WG_BIND_TEX_FIRST + ord }
352 if sir_get(m, d, SIR_KIND) == K_TEXTURE { if sir_get(m, d, SIR_TY) == T_TEX2F { ord = ord + 2 } else { ord = ord + 1 } }
353 d = sir_get(m, d, SIR_NEXT)
354 }
355 return WG_BIND_TEX_FIRST + ord
356}
357// the K_STORAGE declared at this explicit binding, or 0
358func wg_storage_at(m: *i64, binding: i64) -> i64 {
359 var d: i64 = m[M_DECLH]
360 while d != 0 {
361 if sir_get(m, d, SIR_KIND) == K_STORAGE { if sir_get(m, d, SIR_A) == binding { return d } }
362 d = sir_get(m, d, SIR_NEXT)
363 }
364 return 0
365}
366
367// ---- GLSL backend ---------------------------------------------------------------------------
368// A rasterizer-facing binding is legal only on the fragment entry and is boolean.
369// Reject shadowing because GLSL reads map to its implicit builtin rather than a local.
370func wg_front_contract(m:*i64,stage:i64)->i64{
371 var f:i64=m[M_FUNCH]
372 while f!=0{
373 var p:i64=sir_get(m,f,SIR_A);var seen:i64=0
374 while p!=0{
375 if sir_get(m,p,SIR_A)==P_BUILTIN_FRONT_FACING{
376 if stage!=STAGE_FRAGMENT||wg_fn_is_plain(m,f,stage)==1||sir_get(m,p,SIR_TY)!=T_BOOL||seen!=0{sir_refuse(m,"shader: front_facing requires one boolean fragment entry input");return 0}
377 seen=1;let name:*u8=sir_cstr(m,sir_get(m,p,SIR_NAME))
378 if wg_chain_local(m,sir_get(m,f,SIR_B),name)==1{sir_refuse(m,"shader: front_facing input name shadowed by local");return 0}
379 if wg_decl_kind_count(m,K_VARY)>0{if wg_streq(name,"vin")==1{sir_refuse(m,"shader: front_facing input collides with varying interface");return 0}}
380 var q:i64=sir_get(m,f,SIR_A);while q!=0{if q!=p{if wg_name_is(m,q,name)==1{sir_refuse(m,"shader: duplicate front_facing input name");return 0}};q=sir_get(m,q,SIR_NEXT)}
381 };p=sir_get(m,p,SIR_NEXT)
382 };f=sir_get(m,f,SIR_NEXT)
383 };return 1
384}
385
386func gl_expr(m: *i64, id: i64, out: *u8, pos: i64, cap: i64, trace: *i64, tcap: i64, tn: *i64) -> i64 {
387 if id == 0 { return pos }
388 var p: i64 = pos
389 let k: i64 = sir_get(m, id, SIR_KIND)
390 let ty: i64 = sir_get(m, id, SIR_TY)
391 tr_add(trace, tcap, tn, id)
392 if k == E_LIT {
393 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, id, SIR_NAME)))
394 // S12c-5 (2026-09-05): a literal built with sir_lit_u carries its `u` suffix in GLSL too -- the spelling a uint compare
395 // needs in GLSL ES 3.00 (`b!=0u`), and the spelling the hand cast fragment carries. A plain T_U32 literal stays bare
396 // here on purpose: the shipping fullscreen triangle's shift amount is bare in GLSL and suffixed only in WGSL.
397 if sir_get(m, id, SIR_A) == SIR_LIT_USUFFIX { p = eb_put(out, p, cap, "u" as *u8) }
398 return p
399 }
400 if k == E_IDENT {
401 // S4: a parameter bound to @builtin(position) has no name of its own in GLSL -- it IS gl_FragCoord.
402 // Every other identifier stays bare (GLSL uniforms are free-standing; see wg_expr for the WGSL twin).
403 let gp: i64 = wg_fn_param(m, m[M_CURFN], sir_cstr(m, sir_get(m, id, SIR_NAME)))
404 if gp != 0 { if sir_get(m, gp, SIR_A) == P_BUILTIN_POSITION { return eb_put(out, p, cap, "gl_FragCoord" as *u8) }; if sir_get(m,gp,SIR_A)==P_BUILTIN_FRONT_FACING{return eb_put(out,p,cap,"gl_FrontFacing")} }
405 return eb_put(out, p, cap, sir_cstr(m, sir_get(m, id, SIR_NAME)))
406 }
407 if k == E_BUILTIN {
408 let b: *u8 = sir_cstr(m, sir_get(m, id, SIR_NAME))
409 // GE43: WebGL2 has no compute stage, so it has no invocation id. Refuse by name -- an
410 // emitter that improvised gl_FragCoord here would compile and draw the wrong thing forever.
411 if b[0] == (103 as u8) { sir_refuse(m, "glsl: global_invocation_id -- GLSL ES 3.00 (WebGL2) has no compute stage; this kernel reaches a third-party browser only through the WGSL door (GE43)" as *u8); return p }
412 if b[0] == (118 as u8) { return eb_put(out, p, cap, "gl_VertexID" as *u8) }
413 // S12a (2026-09-05): the instance index needs no declaration in GLSL ES 3.00.
414 if b[0] == (105 as u8) { return eb_put(out, p, cap, "gl_InstanceID" as *u8) }
415 // S4: a declared position INPUT (sir_param_position, fragment stage) reads as gl_FragCoord; with no
416 // such parameter the builtin keeps its vertex-stage meaning, gl_Position (the output being assigned).
417 if wg_fn_position_param(m, m[M_CURFN]) != 0 { return eb_put(out, p, cap, "gl_FragCoord" as *u8) }
418 return eb_put(out, p, cap, "gl_Position" as *u8)
419 }
420 if k == E_ATOMIC { sir_refuse(m, "glsl: atomic read-modify-write -- GLSL ES 3.00 (WebGL2) has no atomics and no storage buffers; WGSL-door only (GE43)" as *u8); return p }
421 if k == E_CAST {
422 let tn2: *u8 = gl_ty(ty)
423 if (tn2 as i64) == 0 { sir_refuse(m, "glsl: cast to a type outside the covered subset" as *u8); return p }
424 p = eb_put(out, p, cap, tn2)
425 p = eb_put(out, p, cap, "(" as *u8)
426 p = gl_expr(m, sir_get(m, id, SIR_A), out, p, cap, trace, tcap, tn)
427 return eb_put(out, p, cap, ")" as *u8)
428 }
429 if k == E_BIN {
430 p = eb_put(out, p, cap, "(" as *u8)
431 p = gl_expr(m, sir_get(m, id, SIR_A), out, p, cap, trace, tcap, tn)
432 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, id, SIR_NAME)))
433 p = gl_expr(m, sir_get(m, id, SIR_B), out, p, cap, trace, tcap, tn)
434 return eb_put(out, p, cap, ")" as *u8)
435 }
436 if k == E_NOT {
437 // S12c-4 (2026-09-05): logical not, (!x) in both dialects. The bang is CONSTRUCTED (WG_C_BANG) because the nx_cc lexer
438 // refuses '!' inside a string literal; the shape mirrors E_NEG exactly.
439 p = eb_put(out, p, cap, "(" as *u8)
440 if p + 1 < cap { out[p] = WG_C_BANG as u8; p = p + 1 }
441 p = gl_expr(m, sir_get(m, id, SIR_A), out, p, cap, trace, tcap, tn)
442 return eb_put(out, p, cap, ")" as *u8)
443 }
444 if k == E_NEG {
445 // R-G0 (2026-09-04): unary negation, (-x) in both dialects; the parentheses keep `a- -b` unambiguous.
446 p = eb_put(out, p, cap, "(-" as *u8)
447 p = gl_expr(m, sir_get(m, id, SIR_A), out, p, cap, trace, tcap, tn)
448 return eb_put(out, p, cap, ")" as *u8)
449 }
450 if k == E_SELECT {
451 // R-A (2026-09-04): GLSL spells a select as the ternary. The IR's condition is a scalar bool by
452 // construction (no vector-bool type exists), so the ternary's scalar-condition rule always holds.
453 p = eb_put(out, p, cap, "(" as *u8)
454 p = gl_expr(m, sir_get(m, id, SIR_A), out, p, cap, trace, tcap, tn)
455 p = eb_put(out, p, cap, "?" as *u8)
456 p = gl_expr(m, sir_get(m, id, SIR_B), out, p, cap, trace, tcap, tn)
457 p = eb_put(out, p, cap, ":" as *u8)
458 p = gl_expr(m, sir_get(m, id, SIR_C), out, p, cap, trace, tcap, tn)
459 return eb_put(out, p, cap, ")" as *u8)
460 }
461 if k == E_SWZ {
462 p = gl_expr(m, sir_get(m, id, SIR_A), out, p, cap, trace, tcap, tn)
463 p = eb_put(out, p, cap, "." as *u8)
464 return eb_put(out, p, cap, sir_cstr(m, sir_get(m, id, SIR_NAME)))
465 }
466 if k == E_INDEX {
467 p = gl_expr(m, sir_get(m, id, SIR_A), out, p, cap, trace, tcap, tn)
468 p = eb_put(out, p, cap, "[" as *u8)
469 p = gl_expr(m, sir_get(m, id, SIR_B), out, p, cap, trace, tcap, tn)
470 return eb_put(out, p, cap, "]" as *u8)
471 }
472 if k == E_TEXLOAD {
473 p = eb_put(out, p, cap, "texelFetch(" as *u8)
474 p = gl_expr(m, sir_get(m, id, SIR_A), out, p, cap, trace, tcap, tn)
475 p = eb_put(out, p, cap, "," as *u8)
476 p = gl_expr(m, sir_get(m, id, SIR_B), out, p, cap, trace, tcap, tn)
477 p = eb_put(out, p, cap, "," as *u8)
478 p = gl_expr(m, sir_get(m, id, SIR_C), out, p, cap, trace, tcap, tn)
479 return eb_put(out, p, cap, ")" as *u8)
480 }
481 if k == E_TEXSAMPLE {
482 p = eb_put(out, p, cap, "texture(" as *u8)
483 p = gl_expr(m, sir_get(m, id, SIR_A), out, p, cap, trace, tcap, tn)
484 p = eb_put(out, p, cap, "," as *u8)
485 p = gl_expr(m, sir_get(m, id, SIR_B), out, p, cap, trace, tcap, tn)
486 return eb_put(out, p, cap, ")" as *u8)
487 }
488 if k == E_CALL {
489 // S9: the callee is re-spelled for GLSL through the ONE map (dpdx/dpdy/inverseSqrt); the trace keeps the IR name.
490 p = eb_put(out, p, cap, gl_call_name(sir_cstr(m, sir_get(m, id, SIR_NAME))))
491 p = eb_put(out, p, cap, "(" as *u8)
492 var a: i64 = sir_get(m, id, SIR_A)
493 var first: i64 = 1
494 while a != 0 {
495 if first == 0 { p = eb_put(out, p, cap, "," as *u8) }
496 first = 0
497 p = gl_expr(m, a, out, p, cap, trace, tcap, tn)
498 a = sir_get(m, a, SIR_NEXT)
499 }
500 return eb_put(out, p, cap, ")" as *u8)
501 }
502 if k == E_CTOR {
503 let tn3: *u8 = gl_ty(ty)
504 if (tn3 as i64) == 0 { sir_refuse(m, "glsl: constructor for a type outside the covered subset" as *u8); return p }
505 p = eb_put(out, p, cap, tn3)
506 p = eb_put(out, p, cap, "(" as *u8)
507 var a2: i64 = sir_get(m, id, SIR_A)
508 var f2: i64 = 1
509 while a2 != 0 {
510 if f2 == 0 { p = eb_put(out, p, cap, "," as *u8) }
511 f2 = 0
512 p = gl_expr(m, a2, out, p, cap, trace, tcap, tn)
513 a2 = sir_get(m, a2, SIR_NEXT)
514 }
515 return eb_put(out, p, cap, ")" as *u8)
516 }
517 sir_refuse(m, "glsl: expression kind outside the covered subset" as *u8)
518 return p
519}
520
521func gl_stmts(m: *i64, head: i64, out: *u8, pos: i64, cap: i64, trace: *i64, tcap: i64, tn: *i64) -> i64 {
522 var p: i64 = pos
523 var s: i64 = head
524 while s != 0 {
525 let k: i64 = sir_get(m, s, SIR_KIND)
526 tr_add(trace, tcap, tn, s)
527 if k == S_VAR {
528 // R-D (2026-09-04): a local may be struct-typed; gl_tyname spells a declared struct by name.
529 let t: *u8 = gl_tyname(m, sir_get(m, s, SIR_TY))
530 if (t as i64) == 0 { sir_refuse(m, "glsl: local of a type outside the covered subset" as *u8); return p }
531 // S12c-5 (2026-09-05): a constant local (sir_const, SIR_D=2) is `const T n=...;` -- the hand cast fragment's SHD_CTR.
532 if sir_get(m, s, SIR_D) == SIR_VAR_CONST { p = eb_put(out, p, cap, "const " as *u8) }
533 p = eb_put(out, p, cap, t)
534 p = eb_put(out, p, cap, " " as *u8)
535 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, s, SIR_NAME)))
536 // R-G0 (2026-09-04): a local declared WITHOUT an initialiser (init 0) is `float v;` -- the world shader
537 // declares `var col:vec3f;` and fills it per branch; an emitted `=` with nothing after it would not compile.
538 if sir_get(m, s, SIR_A) != 0 {
539 p = eb_put(out, p, cap, "=" as *u8)
540 p = gl_expr(m, sir_get(m, s, SIR_A), out, p, cap, trace, tcap, tn)
541 }
542 p = eb_put(out, p, cap, ";\n" as *u8)
543 }
544 if k == S_ASSIGN {
545 p = gl_expr(m, sir_get(m, s, SIR_A), out, p, cap, trace, tcap, tn)
546 p = eb_put(out, p, cap, "=" as *u8)
547 p = gl_expr(m, sir_get(m, s, SIR_B), out, p, cap, trace, tcap, tn)
548 p = eb_put(out, p, cap, ";\n" as *u8)
549 }
550 if k == S_RETURN {
551 // S5 (2026-09-04): GLSL ES 3.00 HAS NO VALUE-RETURNING ENTRY POINT. A fragment entry ASSIGNS
552 // its declared `out` variable; WGSL RETURNS one. One IR node, two shapes -- the same structural
553 // split the vertex stage already has (GLSL assigns gl_Position, WGSL returns it), and the exact
554 // pair this backend exists to own. M_FRAGOUT is non-zero ONLY inside the fragment entry's body,
555 // so an ordinary helper in the same module still emits a `return`.
556 if m[M_FRAGOUT] != 0 {
557 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, m[M_FRAGOUT], SIR_NAME)))
558 p = eb_put(out, p, cap, "=" as *u8)
559 p = gl_expr(m, sir_get(m, s, SIR_A), out, p, cap, trace, tcap, tn)
560 p = eb_put(out, p, cap, ";\n" as *u8)
561 } else {
562 // S12c-1 (2026-09-05): a VOID return is `return;` -- never `return ;` with a dangling expression slot.
563 if sir_get(m, s, SIR_A) == 0 { p = eb_put(out, p, cap, "return;\n" as *u8) } else {
564 p = eb_put(out, p, cap, "return " as *u8)
565 p = gl_expr(m, sir_get(m, s, SIR_A), out, p, cap, trace, tcap, tn)
566 p = eb_put(out, p, cap, ";\n" as *u8)
567 }
568 }
569 }
570 if k == S_DISCARD { p = eb_put(out, p, cap, "discard;\n" as *u8) }
571 if k == S_IF {
572 p = eb_put(out, p, cap, "if(" as *u8)
573 p = gl_expr(m, sir_get(m, s, SIR_A), out, p, cap, trace, tcap, tn)
574 p = eb_put(out, p, cap, "){\n" as *u8)
575 p = gl_stmts(m, sir_get(m, s, SIR_B), out, p, cap, trace, tcap, tn)
576 p = eb_put(out, p, cap, "}\n" as *u8)
577 if sir_get(m, s, SIR_C) != 0 {
578 p = eb_put(out, p, cap, "else{\n" as *u8)
579 p = gl_stmts(m, sir_get(m, s, SIR_C), out, p, cap, trace, tcap, tn)
580 p = eb_put(out, p, cap, "}\n" as *u8)
581 }
582 }
583 if k == S_LOOP {
584 p = eb_put(out, p, cap, "while(true){\n" as *u8)
585 p = gl_stmts(m, sir_get(m, s, SIR_A), out, p, cap, trace, tcap, tn)
586 p = eb_put(out, p, cap, "}\n" as *u8)
587 }
588 if k == S_BREAK { p = eb_put(out, p, cap, "break;\n" as *u8) }
589 if k == S_EXPR {
590 p = gl_expr(m, sir_get(m, s, SIR_A), out, p, cap, trace, tcap, tn)
591 p = eb_put(out, p, cap, ";\n" as *u8)
592 }
593 s = sir_get(m, s, SIR_NEXT)
594 }
595 return p
596}
597
598func glsl_emit_module(m: *i64, stage: i64, out: *u8, cap: i64, trace: *i64, tcap: i64, tn: *i64) -> i64 {
599 if wg_front_contract(m,stage)!=1{return -1}
600 if m[M_CLIP_DEPTH] != SIR_CLIP_NATIVE && m[M_CLIP_DEPTH] != SIR_CLIP_NEGATIVE_ONE_TO_ONE { sir_refuse(m, "shader: unsupported source clip-depth convention" as *u8); return 0 - 1 }
601 var p: i64 = 0
602 // GE43: WebGL2 has no compute stage at all. Refuse the STAGE by name before emitting a byte, so a
603 // kernel can never be mis-shaped into a fragment program that compiles, links and draws nothing.
604 if stage == STAGE_COMPUTE { sir_refuse(m, "glsl: compute stage -- GLSL ES 3.00 (WebGL2) has no compute; a kernel reaches a third-party browser only through the WGSL door (GE43/GE44), NishiOS through the dxg lane" as *u8); return 0 - 1 }
605 p = eb_put(out, p, cap, "#version 300 es\n" as *u8)
606 if stage == STAGE_FRAGMENT { p = eb_put(out, p, cap, "precision highp float;precision highp int;\n" as *u8) }
607 var d: i64 = m[M_DECLH]
608 while d != 0 {
609 let k: i64 = sir_get(m, d, SIR_KIND)
610 tr_add(trace, tcap, tn, d)
611 // GE43: named refusal FIRST -- a u32 storage element has a GLSL type name, so without this
612 // the declaration would fall through every branch below and vanish SILENTLY from the output.
613 if k == K_STORAGE { sir_refuse(m, "glsl: storage buffer declaration -- GLSL ES 3.00 (WebGL2) has no SSBO; WGSL-door only (GE43)" as *u8); return 0 - 1 }
614 // R-D (2026-09-04): a struct declaration -- `struct Hit{vec3 pos;float d;};`. Members are the K_PARAM
615 // chain under SIR_A, each traced and spelled through gl_tyname (a struct-typed member resolves by name;
616 // an uncovered or texture-typed member refuses by name, never vanishes). Handled BEFORE the generic
617 // type check for the same reason K_TEXTURE is.
618 if k == K_STRUCT {
619 p = eb_put(out, p, cap, "struct " as *u8)
620 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, d, SIR_NAME)))
621 p = eb_put(out, p, cap, "{" as *u8)
622 var mb: i64 = sir_get(m, d, SIR_A)
623 while mb != 0 {
624 tr_add(trace, tcap, tn, mb)
625 if wg_is_texty(sir_get(m, mb, SIR_TY)) == 1 { sir_refuse(m, "glsl: texture-typed struct member -- a texture is a bound resource, never a value; refused in both dialects" as *u8); return 0 - 1 }
626 let mt: *u8 = gl_tyname(m, sir_get(m, mb, SIR_TY))
627 if (mt as i64) == 0 { sir_refuse(m, "glsl: struct member of a type outside the covered subset" as *u8); return 0 - 1 }
628 p = eb_put(out, p, cap, mt)
629 p = eb_put(out, p, cap, " " as *u8)
630 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, mb, SIR_NAME)))
631 p = eb_put(out, p, cap, ";" as *u8)
632 mb = sir_get(m, mb, SIR_NEXT)
633 }
634 p = eb_put(out, p, cap, "};\n" as *u8)
635 }
636 // S3 (2026-09-04): a texture is a sampler uniform in GLSL. Handled BEFORE the generic type check for
637 // the same reason K_STORAGE is: gl_ty names the type, so the declaration would otherwise fall through
638 // every branch below and vanish SILENTLY. Only the covered texture type is lowered.
639 if k == K_TEXTURE {
640 let gtk: i64 = sir_get(m, d, SIR_TY)
641 if gtk != T_TEX3U { if gtk != T_TEX2F { sir_refuse(m, "glsl: texture type outside the covered subset -- usampler3D (T_TEX3U, texelFetch) and sampler2D (T_TEX2F, texture()) are lowered; refused rather than guessed" as *u8); return 0 - 1 } }
642 p = eb_put(out, p, cap, "uniform " as *u8)
643 p = eb_put(out, p, cap, gl_ty(gtk))
644 p = eb_put(out, p, cap, " " as *u8)
645 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, d, SIR_NAME)))
646 p = eb_put(out, p, cap, ";\n" as *u8)
647 }
648 // S3: the deprecated form refuses here as well, so BOTH backends cover exactly the same declarations
649 // -- a source that emits in one dialect and refuses in the other is a fork wearing one name.
650 if k == K_UNIFORM { if wg_is_texty(sir_get(m, d, SIR_TY)) == 1 { sir_refuse(m, "glsl: texture-typed uniform -- a texture is its own declaration kind (sir_texture / K_TEXTURE); refused so both backends cover the same source" as *u8); return 0 - 1 } }
651 let t: *u8 = gl_ty(sir_get(m, d, SIR_TY))
652 if (t as i64) == 0 { sir_refuse(m, "glsl: declaration of a type outside the covered subset" as *u8); return 0 - 1 }
653 if k == K_PRIVATE {
654 // GE55: module-scope private state is a bare global in GLSL (`T name;` / `T name[N];`), emitted in
655 // declaration order. `t` is this declaration's GLSL type, already null-guarded above; the trace entry
656 // was added at the loop head like every other declaration. A texture-typed private refuses by name in
657 // BOTH backends, so the two cover exactly the same source.
658 if wg_is_texty(sir_get(m, d, SIR_TY)) == 1 { sir_refuse(m, "glsl: texture-typed private variable -- a texture is a bound resource, never module state; refused so both backends cover the same source" as *u8); return 0 - 1 }
659 p = eb_put(out, p, cap, t)
660 p = eb_put(out, p, cap, " " as *u8)
661 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, d, SIR_NAME)))
662 if sir_get(m, d, SIR_A) > 0 {
663 p = eb_put(out, p, cap, "[" as *u8)
664 p = eb_num(out, p, cap, sir_get(m, d, SIR_A))
665 p = eb_put(out, p, cap, "]" as *u8)
666 }
667 p = eb_put(out, p, cap, ";\n" as *u8)
668 }
669 if k == K_UNIFORM {
670 p = eb_put(out, p, cap, "uniform " as *u8)
671 p = eb_put(out, p, cap, t)
672 p = eb_put(out, p, cap, " " as *u8)
673 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, d, SIR_NAME)))
674 if sir_get(m, d, SIR_A) > 0 {
675 p = eb_put(out, p, cap, "[" as *u8)
676 p = eb_num(out, p, cap, sir_get(m, d, SIR_A))
677 p = eb_put(out, p, cap, "]" as *u8)
678 }
679 p = eb_put(out, p, cap, ";\n" as *u8)
680 }
681 if k == K_ATTRIB {
682 p = eb_put(out, p, cap, "layout(location=" as *u8)
683 p = eb_num(out, p, cap, sir_get(m, d, SIR_A))
684 p = eb_put(out, p, cap, ") in " as *u8)
685 p = eb_put(out, p, cap, t)
686 p = eb_put(out, p, cap, " " as *u8)
687 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, d, SIR_NAME)))
688 p = eb_put(out, p, cap, ";\n" as *u8)
689 }
690 // S5 (2026-09-04): THE FRAGMENT OUTPUT IS A DECLARATION IN GLSL. `layout(location=N)` is emitted
691 // only for N>0: location 0 is the default the shipping FSH spells as a bare `out vec4 fc;`, so the
692 // derived bytes reproduce the hand-written form exactly instead of merely being equivalent to it.
693 if k == K_FRAGOUT {
694 if stage != STAGE_FRAGMENT { sir_refuse(m, "glsl: fragment output declaration (K_FRAGOUT) outside a fragment stage -- only the fragment stage has a colour output; refused rather than declared where nothing can write it" as *u8); return 0 - 1 }
695 if wg_decl_kind_count(m, K_FRAGOUT) > 1 { sir_refuse(m, "glsl: more than one fragment output declared -- multiple render targets are their own rung; refused rather than emitting the first and silently dropping the rest" as *u8); return 0 - 1 }
696 if sir_get(m, d, SIR_A) > 0 {
697 p = eb_put(out, p, cap, "layout(location=" as *u8)
698 p = eb_num(out, p, cap, sir_get(m, d, SIR_A))
699 p = eb_put(out, p, cap, ") " as *u8)
700 }
701 p = eb_put(out, p, cap, "out " as *u8)
702 p = eb_put(out, p, cap, t)
703 p = eb_put(out, p, cap, " " as *u8)
704 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, d, SIR_NAME)))
705 p = eb_put(out, p, cap, ";\n" as *u8)
706 }
707 if k == K_VARY {
708 if sir_get(m, d, SIR_B) == 1 { p = eb_put(out, p, cap, "flat " as *u8) }
709 if stage == STAGE_VERTEX { p = eb_put(out, p, cap, "out " as *u8) } else { p = eb_put(out, p, cap, "in " as *u8) }
710 p = eb_put(out, p, cap, t)
711 p = eb_put(out, p, cap, " " as *u8)
712 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, d, SIR_NAME)))
713 p = eb_put(out, p, cap, ";\n" as *u8)
714 }
715 d = sir_get(m, d, SIR_NEXT)
716 }
717 var f: i64 = m[M_FUNCH]
718 while f != 0 {
719 tr_add(trace, tcap, tn, f)
720 m[M_CURFN] = f
721 if wg_fn_is_plain(m, f, stage) == 1 {
722 // R-E (2026-09-04): a function may RETURN a struct in GLSL ES 3.00 exactly as in WGSL, so the canonical
723 // shape is the struct return in both dialects and no out-parameter twin exists; gl_tyname spells it.
724 let rt: *u8 = gl_tyname(m, sir_get(m, f, SIR_TY))
725 if (rt as i64) == 0 { sir_refuse(m, "glsl: return type outside the covered subset" as *u8); return 0 - 1 }
726 p = eb_put(out, p, cap, rt)
727 p = eb_put(out, p, cap, " " as *u8)
728 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, f, SIR_NAME)))
729 p = eb_put(out, p, cap, "(" as *u8)
730 var pa: i64 = sir_get(m, f, SIR_A)
731 var pf: i64 = 1
732 while pa != 0 {
733 if pf == 0 { p = eb_put(out, p, cap, "," as *u8) }
734 pf = 0
735 tr_add(trace, tcap, tn, pa)
736 let pt: *u8 = gl_tyname(m, sir_get(m, pa, SIR_TY))
737 if (pt as i64) == 0 { sir_refuse(m, "glsl: parameter type outside the covered subset" as *u8); return 0 - 1 }
738 p = eb_put(out, p, cap, pt)
739 p = eb_put(out, p, cap, " " as *u8)
740 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, pa, SIR_NAME)))
741 pa = sir_get(m, pa, SIR_NEXT)
742 }
743 p = eb_put(out, p, cap, "){\n" as *u8)
744 } else {
745 // S4 (2026-09-04): GLSL main() takes no parameters. A parameter bound to @builtin(position) on a
746 // FRAGMENT entry is gl_FragCoord -- implicit, so nothing is emitted for it here and its reads spell
747 // it (gl_expr). Any other entry-point parameter is an input GLSL binds by module-scope declaration
748 // (K_ATTRIB / K_VARY), never by parameter: refused by name rather than dropped from the program.
749 var gp: i64 = sir_get(m, f, SIR_A)
750 while gp != 0 {
751 tr_add(trace, tcap, tn, gp)
752 if sir_get(m, gp, SIR_A) != P_BUILTIN_POSITION && sir_get(m,gp,SIR_A)!=P_BUILTIN_FRONT_FACING { sir_refuse(m, "glsl: entry-point parameter that is not the @builtin(position) input -- GLSL main() takes no parameters; inputs are module-scope declarations (K_ATTRIB / K_VARY); refused rather than silently dropped" as *u8); return 0 - 1 }
753 if stage != STAGE_FRAGMENT { sir_refuse(m, "glsl: @builtin(position) declared as an input outside a fragment stage -- position is a fragment-stage input (gl_FragCoord) and a vertex-stage output (gl_Position); refused rather than guessed" as *u8); return 0 - 1 }
754 gp = sir_get(m, gp, SIR_NEXT)
755 }
756 // S5: a fragment entry that RETURNS A VALUE must have somewhere to put it. Before S5 this emitted
757 // `void main(){ return vec4(...); }` -- a value returned from a void function, against no declared
758 // output -- invalid GLSL ES 3.00 on BOTH counts, shipped with no refusal, and invisible because
759 // nothing pinned the GLSL fragment bytes. A shader that only discards declares no output and is
760 // untouched by this rule.
761 if stage == STAGE_FRAGMENT { if wg_chain_returns_value(m, sir_get(m, f, SIR_B)) == 1 {
762 let fo: i64 = wg_fragout(m)
763 if fo == 0 { sir_refuse(m, "glsl: fragment entry returns a value with no declared fragment output -- GLSL ES 3.00 has no value-returning entry point; declare the output with sir_fragout (K_FRAGOUT) and the return lowers to an assignment" as *u8); return 0 - 1 }
764 m[M_FRAGOUT] = fo
765 } }
766 p = eb_put(out, p, cap, "void " as *u8)
767 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, f, SIR_NAME)))
768 p = eb_put(out, p, cap, "(){\n" as *u8)
769 }
770 p = gl_stmts(m, sir_get(m, f, SIR_B), out, p, cap, trace, tcap, tn)
771 // S5: the assignment shape belongs to the ENTRY body only -- clear it before the next function so a
772 // helper declared after the entry cannot inherit it and silently assign the colour output.
773 m[M_FRAGOUT] = 0
774 p = eb_put(out, p, cap, "}\n" as *u8)
775 f = sir_get(m, f, SIR_NEXT)
776 }
777 m[M_CURFN] = 0
778 if sir_refused(m) != 0 { return 0 - 1 }
779 if p + 1 < cap { out[p] = 0 as u8 }
780 return p
781}
782
783// ---- WGSL backend ----------------------------------------------------------------------------
784// R-G1 (2026-09-04): WHICH FUNCTION IS THE ENTRY. Until the world shader arrived every staged module held ONE
785// function, so the backends treated every function of a vertex/fragment/compute module as the entry and refused
786// helpers by name ("entry-point parameter that is not the @builtin(position) input"). The rule, ONE for both
787// dialects: in a staged module the function named `main` is the entry (GLSL requires that name; the WGSL entry
788// keeps it and the page names its entryPoint); every other function is a plain function in any stage.
789func wg_fn_is_plain(m: *i64, f: i64, stage: i64) -> i64 {
790 if stage == STAGE_PLAIN { return 1 }
791 if wg_name_is(m, f, "main" as *u8) == 1 { return 0 }
792 return 1
793}
794// R-A TRACE IDENTITY (2026-09-04): reverse out[a..b) in place. The E_SELECT arm emits its three operands in IR
795// order (the trace nx_wgsl_gate compares across dialects) and then rotates the spans into WGSL argument order
796// with three reversals -- no scratch buffer, no second visitation order, nothing to size.
797func wg_rev(out: *u8, a: i64, b: i64) -> i64 {
798 var i: i64 = a
799 var j: i64 = b - 1
800 while i < j {
801 let t: i64 = out[i] as i64
802 out[i] = out[j]
803 out[j] = t as u8
804 i = i + 1
805 j = j - 1
806 }
807 return 0
808}
809
810// S12c-6 (2026-09-05): WGSL's clamp/min/max/smoothstep/step take ONE type T for every argument (f32 or vecNf) -- there
811// is no scalar-bound overload -- while GLSL ES accepts clamp(vec3,float,float), and the hand cast fragment uses exactly
812// that form (the skin tail: clamp((vec3(ndl)+w3)/(1.+w3),0.,1.)). When any argument of one of these callees is a float
813// vector, every f32 argument is SPLATTED as vecNf(<arg>) in the WGSL dialect only; the GLSL emitter is a different
814// backend and never sees this, so the token ruler's IDENTICAL cannot move. A user function or a builtin with all-scalar
815// arguments (max(0.,-n.y), clamp(fl9,0.,.3)) passes through untouched -- the neg-control in nx_wgsl_gate pins that.
816func wg_splat_callee(n: *u8) -> i64 {
817 if wg_streq(n, "clamp" as *u8) == 1 { return 1 }
818 if wg_streq(n, "min" as *u8) == 1 { return 1 }
819 if wg_streq(n, "max" as *u8) == 1 { return 1 }
820 if wg_streq(n, "smoothstep" as *u8) == 1 { return 1 }
821 if wg_streq(n, "step" as *u8) == 1 { return 1 }
822 return 0
823}
824// the float-vector type among a call's arguments (the type every f32 argument must be splatted to), 0 when none
825func wg_vec_arg_ty(m: *i64, id: i64) -> i64 {
826 var a: i64 = sir_get(m, id, SIR_A)
827 while a != 0 {
828 let t: i64 = sir_get(m, a, SIR_TY)
829 if t == T_V2F { return t }
830 if t == T_V3F { return t }
831 if t == T_V4F { return t }
832 a = sir_get(m, a, SIR_NEXT)
833 }
834 return 0
835}
836
837func wg_expr(m: *i64, id: i64, out: *u8, pos: i64, cap: i64, trace: *i64, tcap: i64, tn: *i64) -> i64 {
838 if id == 0 { return pos }
839 var p: i64 = pos
840 let k: i64 = sir_get(m, id, SIR_KIND)
841 let ty: i64 = sir_get(m, id, SIR_TY)
842 tr_add(trace, tcap, tn, id)
843 if k == E_LIT {
844 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, id, SIR_NAME)))
845 // WGSL DIALECT RULE: an unsigned literal carries its u suffix. A shift amount MUST be u32
846 // in WGSL and must NOT be in GLSL -- the backend owns that, the author never sees it.
847 if ty == T_U32 { p = eb_put(out, p, cap, "u" as *u8) }
848 return p
849 }
850 if k == E_IDENT {
851 // S2: a uniform lives in struct U behind binding 0, so its read is `u.<name>`; a parameter or a
852 // local of the current function keeps the bare name (it shadows the uniform, as in the language).
853 // GLSL keeps every identifier bare -- its uniforms are free-standing -- so only THIS backend qualifies.
854 if wg_resolves_to_uniform(m, sir_cstr(m, sir_get(m, id, SIR_NAME))) == 1 { p = eb_put(out, p, cap, "u." as *u8) }
855 // S7 (2026-09-05): a varying lives in the interstage struct -- written as vo.<name> in the vertex entry,
856 // read as vin.<name> in the fragment entry (M_VIOMODE 1 / 2). Outside an entry (a plain helper) the
857 // struct is not in scope, so the reference refuses by name rather than emitting a bare undeclared
858 // identifier. GLSL keeps every varying bare -- its varyings are module-scope -- so only THIS backend qualifies.
859 if wg_resolves_to_varying(m, sir_cstr(m, sir_get(m, id, SIR_NAME))) == 1 {
860 if m[M_VIOMODE] == 1 { p = eb_put(out, p, cap, "vo." as *u8) }
861 if m[M_VIOMODE] == 2 { p = eb_put(out, p, cap, "vin." as *u8) }
862 if m[M_VIOMODE] == 0 { sir_refuse(m, "wgsl: varying referenced outside a vertex or fragment entry -- the interstage struct is an entry parameter/return, not module scope; pass the value to the helper explicitly" as *u8); return p }
863 }
864 return eb_put(out, p, cap, sir_cstr(m, sir_get(m, id, SIR_NAME)))
865 }
866 if k == E_BUILTIN {
867 let b: *u8 = sir_cstr(m, sir_get(m, id, SIR_NAME))
868 // vertex_index arrives as u32; the IR types the expression i32, so the backend converts.
869 if b[0] == (118 as u8) { return eb_put(out, p, cap, "i32(vi)" as *u8) }
870 // S12a (2026-09-05): instance_index arrives as u32 like vertex_index and is typed i32 by the IR; the entry binds
871 // it as `ii` only when the module uses it (M_USESII), so unrelated modules stay byte-identical.
872 if b[0] == (105 as u8) { return eb_put(out, p, cap, "i32(iidx)" as *u8) }
873 // GE43: the compute entry binds @builtin(global_invocation_id) to the parameter `gid`.
874 if b[0] == (103 as u8) { return eb_put(out, p, cap, "gid" as *u8) }
875 // S4 (2026-09-04): a READ of @builtin(position) resolves to the entry parameter the function declared
876 // for it (sir_param_position); with no such parameter it is an UNDECLARED INPUT and refuses by name.
877 // Before S4 this emitted the placeholder identifier POSITION with no refusal. (An ASSIGNMENT to the
878 // builtin never reaches here -- wg_stmts turns it into the vertex stage's return.)
879 // S7 (2026-09-05): inside a fragment entry that takes the interstage struct (M_VIOMODE=2) the position
880 // arrives as vin.bpos -- the declared parameter (if any) was folded into it by the entry emitter.
881 if m[M_VIOMODE] == 2 { return eb_put(out, p, cap, "vin.bpos" as *u8) }
882 let pp: i64 = wg_fn_position_param(m, m[M_CURFN])
883 if pp == 0 { sir_refuse(m, "wgsl: @builtin(position) read without a declared entry parameter -- declare it with sir_param_position; an undeclared input is refused rather than emitted as a placeholder identifier" as *u8); return p }
884 return eb_put(out, p, cap, sir_cstr(m, sir_get(m, pp, SIR_NAME)))
885 }
886 if k == E_ATOMIC {
887 // WGSL DIALECT RULE: an atomic builtin takes a POINTER to its target -- the `&` is the
888 // backend's, never the author's. The op spelling is the IR's (atomicMin/atomicMax/atomicAdd).
889 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, id, SIR_NAME)))
890 p = eb_put(out, p, cap, "(&" as *u8)
891 p = wg_expr(m, sir_get(m, id, SIR_A), out, p, cap, trace, tcap, tn)
892 p = eb_put(out, p, cap, "," as *u8)
893 p = wg_expr(m, sir_get(m, id, SIR_B), out, p, cap, trace, tcap, tn)
894 return eb_put(out, p, cap, ")" as *u8)
895 }
896 if k == E_CAST {
897 let tn2: *u8 = wg_ty(ty)
898 if (tn2 as i64) == 0 { sir_refuse(m, "wgsl: cast to a type outside the covered subset" as *u8); return p }
899 p = eb_put(out, p, cap, tn2)
900 p = eb_put(out, p, cap, "(" as *u8)
901 p = wg_expr(m, sir_get(m, id, SIR_A), out, p, cap, trace, tcap, tn)
902 return eb_put(out, p, cap, ")" as *u8)
903 }
904 if k == E_BIN {
905 p = eb_put(out, p, cap, "(" as *u8)
906 p = wg_expr(m, sir_get(m, id, SIR_A), out, p, cap, trace, tcap, tn)
907 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, id, SIR_NAME)))
908 p = wg_expr(m, sir_get(m, id, SIR_B), out, p, cap, trace, tcap, tn)
909 return eb_put(out, p, cap, ")" as *u8)
910 }
911 if k == E_NOT {
912 // S12c-4 (2026-09-05): logical not, (!x) -- the WGSL twin of the GLSL arm, bang constructed (WG_C_BANG).
913 p = eb_put(out, p, cap, "(" as *u8)
914 if p + 1 < cap { out[p] = WG_C_BANG as u8; p = p + 1 }
915 p = wg_expr(m, sir_get(m, id, SIR_A), out, p, cap, trace, tcap, tn)
916 return eb_put(out, p, cap, ")" as *u8)
917 }
918 if k == E_NEG {
919 // R-G0 (2026-09-04): unary negation, (-x) in both dialects (see the GLSL twin).
920 p = eb_put(out, p, cap, "(-" as *u8)
921 p = wg_expr(m, sir_get(m, id, SIR_A), out, p, cap, trace, tcap, tn)
922 return eb_put(out, p, cap, ")" as *u8)
923 }
924 if k == E_SELECT {
925 // R-A (2026-09-04): WGSL DIALECT RULE -- select(f, t, cond) takes the FALSE value FIRST. The IR carries
926 // (cond, then, else); the reordering is the backend's, the author never writes it.
927 // TRACE IDENTITY (same day, caught by nx_wgsl_gate's selectlet_trace_equal tooth): the operands are VISITED
928 // in IR order (A, B, C) exactly as gl_expr visits them, so the node trace is dialect-independent; the WGSL
929 // argument order is produced afterwards by an in-place rotation (reverse the span, reverse each operand
930 // back) and only when every separator landed inside cap -- a truncated span is left as emitted.
931 p = eb_put(out, p, cap, "select(" as *u8)
932 let s0: i64 = p
933 p = wg_expr(m, sir_get(m, id, SIR_A), out, p, cap, trace, tcap, tn)
934 let la: i64 = p - s0
935 p = eb_put(out, p, cap, "," as *u8)
936 let s1: i64 = p
937 p = wg_expr(m, sir_get(m, id, SIR_B), out, p, cap, trace, tcap, tn)
938 let lb: i64 = p - s1
939 p = eb_put(out, p, cap, "," as *u8)
940 let s2: i64 = p
941 p = wg_expr(m, sir_get(m, id, SIR_C), out, p, cap, trace, tcap, tn)
942 let lc: i64 = p - s2
943 if p - s0 == la + lb + lc + 2 {
944 wg_rev(out, s0, p)
945 wg_rev(out, s0, s0 + lc)
946 wg_rev(out, s0 + lc + 1, s0 + lc + 1 + lb)
947 wg_rev(out, s0 + lc + 1 + lb + 1, p)
948 }
949 return eb_put(out, p, cap, ")" as *u8)
950 }
951 if k == E_SWZ {
952 p = wg_expr(m, sir_get(m, id, SIR_A), out, p, cap, trace, tcap, tn)
953 p = eb_put(out, p, cap, "." as *u8)
954 return eb_put(out, p, cap, sir_cstr(m, sir_get(m, id, SIR_NAME)))
955 }
956 if k == E_INDEX {
957 // S12c-2b: a float[N] uniform is carried as array<vec4f,ceil(N/4)> in struct U, so its read is u.name[i/4][i%4] --
958 // the ONE rule the struct emitter and the layout ruler apply (wg_f32arr_len). The index expression is emitted TWICE
959 // here (once per bracket), so the WGSL trace visits its nodes twice: the trace oracle is DECLARED inapplicable to
960 // modules that read a float[N] uniform (the cast is one), and `cast_trace_equal` is printed, never asserted, for them.
961 let bx: i64 = sir_get(m, id, SIR_A)
962 if sir_get(m, bx, SIR_KIND) == E_IDENT { if wg_f32arr_len(m, sir_cstr(m, sir_get(m, bx, SIR_NAME))) > 0 {
963 p = wg_expr(m, bx, out, p, cap, trace, tcap, tn)
964 p = eb_put(out, p, cap, "[(" as *u8)
965 p = wg_expr(m, sir_get(m, id, SIR_B), out, p, cap, trace, tcap, tn)
966 p = eb_put(out, p, cap, ")/" as *u8)
967 p = eb_num(out, p, cap, WG_F32ARR_LANES)
968 p = eb_put(out, p, cap, "][(" as *u8)
969 p = wg_expr(m, sir_get(m, id, SIR_B), out, p, cap, trace, tcap, tn)
970 p = eb_put(out, p, cap, ")%" as *u8)
971 p = eb_num(out, p, cap, WG_F32ARR_LANES)
972 return eb_put(out, p, cap, "]" as *u8)
973 } }
974 p = wg_expr(m, sir_get(m, id, SIR_A), out, p, cap, trace, tcap, tn)
975 p = eb_put(out, p, cap, "[" as *u8)
976 p = wg_expr(m, sir_get(m, id, SIR_B), out, p, cap, trace, tcap, tn)
977 return eb_put(out, p, cap, "]" as *u8)
978 }
979 if k == E_TEXLOAD {
980 p = eb_put(out, p, cap, "textureLoad(" as *u8)
981 p = wg_expr(m, sir_get(m, id, SIR_A), out, p, cap, trace, tcap, tn)
982 p = eb_put(out, p, cap, "," as *u8)
983 p = wg_expr(m, sir_get(m, id, SIR_B), out, p, cap, trace, tcap, tn)
984 p = eb_put(out, p, cap, "," as *u8)
985 p = wg_expr(m, sir_get(m, id, SIR_C), out, p, cap, trace, tcap, tn)
986 return eb_put(out, p, cap, ")" as *u8)
987 }
988 if k == E_TEXSAMPLE {
989 // S5 (2026-09-04): THIS EMITTED `textureSample(t,samp,c)` AGAINST AN UNDECLARED `samp` -- WGSL that
990 // does not compile, produced with NO refusal, while the DECLARATION pass in wgsl_emit_module already
991 // refused a sampled texture_2d<f32> for precisely the sampler binding this backend does not declare.
992 // The declaration side refused and the expression side guessed: half a law, and the half left undone
993 // was the half that ships. MEASURED on the shipping world shader before choosing between building a
994 // sampler kind and refusing: texelFetch 1 / textureLoad 1 / texture( 0 / textureSample 0 over both
995 // whole files -- it performs ZERO sampled reads, because a texture_3d<u32> is read with textureLoad
996 // and textureLoad takes NO sampler. A sampler declaration kind would therefore be built for no
997 // measured caller. When a sampled read IS needed, the rung is a sampler declaration + its binding,
998 // and this refusal is the contract that will name it.
999 // GE54 (2026-09-05): textureSample(<tex>, <tex>_s, <coord>) -- the companion sampler declared beside
1000 // the texture_2d<f32> binding. The texture node (SIR_A) is an identifier; its name plus `_s` is the
1001 // sampler, so the two agree by construction. GLSL emits texture(<tex>, <coord>) -- one combined sampler2D.
1002 p = eb_put(out, p, cap, "textureSample(" as *u8)
1003 p = wg_expr(m, sir_get(m, id, SIR_A), out, p, cap, trace, tcap, tn)
1004 p = eb_put(out, p, cap, "," as *u8)
1005 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, sir_get(m, id, SIR_A), SIR_NAME)))
1006 p = eb_put(out, p, cap, "_s," as *u8)
1007 p = wg_expr(m, sir_get(m, id, SIR_B), out, p, cap, trace, tcap, tn)
1008 return eb_put(out, p, cap, ")" as *u8)
1009 }
1010 if k == E_CALL {
1011 let cn: *u8 = sir_cstr(m, sir_get(m, id, SIR_NAME))
1012 p = eb_put(out, p, cap, cn)
1013 p = eb_put(out, p, cap, "(" as *u8)
1014 // S12c-6: every f32 argument of clamp/min/max/smoothstep/step is splatted to the call's float-vector type
1015 // (wg_splat_callee / wg_vec_arg_ty above): WGSL has no scalar-bound overload, GLSL ES does, ONE source serves both.
1016 var splat: i64 = 0
1017 if wg_splat_callee(cn) == 1 { splat = wg_vec_arg_ty(m, id) }
1018 var a: i64 = sir_get(m, id, SIR_A)
1019 var first: i64 = 1
1020 while a != 0 {
1021 if first == 0 { p = eb_put(out, p, cap, "," as *u8) }
1022 first = 0
1023 var wrap: i64 = 0
1024 if splat != 0 { if sir_get(m, a, SIR_TY) == T_F32 { wrap = 1 } }
1025 if wrap == 1 { p = eb_put(out, p, cap, wg_ty(splat)); p = eb_put(out, p, cap, "(" as *u8) }
1026 p = wg_expr(m, a, out, p, cap, trace, tcap, tn)
1027 if wrap == 1 { p = eb_put(out, p, cap, ")" as *u8) }
1028 a = sir_get(m, a, SIR_NEXT)
1029 }
1030 return eb_put(out, p, cap, ")" as *u8)
1031 }
1032 if k == E_CTOR {
1033 let tn3: *u8 = wg_ty(ty)
1034 if (tn3 as i64) == 0 { sir_refuse(m, "wgsl: constructor for a type outside the covered subset" as *u8); return p }
1035 p = eb_put(out, p, cap, tn3)
1036 p = eb_put(out, p, cap, "(" as *u8)
1037 var a2: i64 = sir_get(m, id, SIR_A)
1038 var f2: i64 = 1
1039 while a2 != 0 {
1040 if f2 == 0 { p = eb_put(out, p, cap, "," as *u8) }
1041 f2 = 0
1042 p = wg_expr(m, a2, out, p, cap, trace, tcap, tn)
1043 a2 = sir_get(m, a2, SIR_NEXT)
1044 }
1045 return eb_put(out, p, cap, ")" as *u8)
1046 }
1047 sir_refuse(m, "wgsl: expression kind outside the covered subset" as *u8)
1048 return p
1049}
1050
1051func wg_is_position(m: *i64, id: i64) -> i64 {
1052 if id == 0 { return 0 }
1053 if sir_get(m, id, SIR_KIND) != E_BUILTIN { return 0 }
1054 let b: *u8 = sir_cstr(m, sir_get(m, id, SIR_NAME))
1055 if b[0] == (112 as u8) { return 1 }
1056 return 0
1057}
1058
1059func wg_stmts(m: *i64, head: i64, out: *u8, pos: i64, cap: i64, trace: *i64, tcap: i64, tn: *i64) -> i64 {
1060 var p: i64 = pos
1061 var s: i64 = head
1062 while s != 0 {
1063 let k: i64 = sir_get(m, s, SIR_KIND)
1064 tr_add(trace, tcap, tn, s)
1065 if k == S_VAR {
1066 // R-D (2026-09-04): a local may be struct-typed; wg_tyname spells a declared struct by name.
1067 let t: *u8 = wg_tyname(m, sir_get(m, s, SIR_TY))
1068 if (t as i64) == 0 { sir_refuse(m, "wgsl: local of a type outside the covered subset" as *u8); return p }
1069 // R-C (2026-09-04): an immutable local (sir_let, SIR_D=1) is `let` in WGSL; the mutable one stays `var`.
1070 // S12c-5 (2026-09-05): a constant local (sir_const, SIR_D=2) is `const n:T=...;` (see the GLSL twin).
1071 if sir_get(m, s, SIR_D) == SIR_VAR_CONST { p = eb_put(out, p, cap, "const " as *u8) } else {
1072 if sir_get(m, s, SIR_D) == SIR_VAR_LET { p = eb_put(out, p, cap, "let " as *u8) } else { p = eb_put(out, p, cap, "var " as *u8) }
1073 }
1074 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, s, SIR_NAME)))
1075 p = eb_put(out, p, cap, ":" as *u8)
1076 p = eb_put(out, p, cap, t)
1077 // R-G0 (2026-09-04): a local declared WITHOUT an initialiser (init 0) is `var v:f32;` (see the GLSL twin).
1078 if sir_get(m, s, SIR_A) != 0 {
1079 p = eb_put(out, p, cap, "=" as *u8)
1080 p = wg_expr(m, sir_get(m, s, SIR_A), out, p, cap, trace, tcap, tn)
1081 }
1082 p = eb_put(out, p, cap, ";\n" as *u8)
1083 }
1084 if k == S_ASSIGN {
1085 // STRUCTURAL DIALECT DIFFERENCE, OWNED BY THE BACKEND: GLSL ASSIGNS gl_Position;
1086 // a WGSL vertex entry point RETURNS @builtin(position). One IR node, two shapes.
1087 if wg_is_position(m, sir_get(m, s, SIR_A)) == 1 {
1088 tr_add(trace, tcap, tn, sir_get(m, s, SIR_A))
1089 // S7 (2026-09-05): inside a vertex entry that returns the interstage struct (M_VIOMODE=1) the
1090 // position lands in vo.bpos and the entry returns vo after its body; with no varyings it is
1091 // the return itself, exactly as before -- one IR node, the same two shapes, plus the struct.
1092 let remap_depth: i64 = m[M_CLIP_DEPTH] == SIR_CLIP_NEGATIVE_ONE_TO_ONE
1093 if m[M_VIOMODE] == 1 { p = eb_put(out, p, cap, "vo.bpos=" as *u8) } else {
1094 if remap_depth { p = eb_put(out, p, cap, "{let nishiClipPosition=" as *u8) } else { p = eb_put(out, p, cap, "return " as *u8) }
1095 }
1096 p = wg_expr(m, sir_get(m, s, SIR_B), out, p, cap, trace, tcap, tn)
1097 p = eb_put(out, p, cap, ";\n" as *u8)
1098 if remap_depth {
1099 if m[M_VIOMODE] == 1 {
1100 p = eb_put(out, p, cap, "vo.bpos.z=(vo.bpos.z+vo.bpos.w)*0.5;\n" as *u8)
1101 } else {
1102 p = eb_put(out, p, cap, "return vec4f(nishiClipPosition.xy,(nishiClipPosition.z+nishiClipPosition.w)*0.5,nishiClipPosition.w);}\n" as *u8)
1103 }
1104 }
1105 } else {
1106 p = wg_expr(m, sir_get(m, s, SIR_A), out, p, cap, trace, tcap, tn)
1107 p = eb_put(out, p, cap, "=" as *u8)
1108 p = wg_expr(m, sir_get(m, s, SIR_B), out, p, cap, trace, tcap, tn)
1109 p = eb_put(out, p, cap, ";\n" as *u8)
1110 }
1111 }
1112 if k == S_RETURN {
1113 // S12c-1 (2026-09-05): a VOID return (no expression) inside the interstage-struct vertex entry (M_VIOMODE=1)
1114 // hands back the struct built so far -- `return vo;` -- the early exit the cast vertex shader uses to cull a
1115 // hair lock after writing its outputs. Anywhere else a void return is `return;`. GLSL says `return;` in both.
1116 if sir_get(m, s, SIR_A) == 0 {
1117 // S12c-5 (2026-09-05): inside a fragment entry whose body ASSIGNS the declared output (M_WGFRAGOUT), a void
1118 // return hands back that variable -- `return fc;` -- the WGSL spelling of the hand's `fc=...;return;` early exits.
1119 if m[M_VIOMODE] == 1 { p = eb_put(out, p, cap, "return vo;\n" as *u8) } else {
1120 if m[M_WGFRAGOUT] != 0 {
1121 p = eb_put(out, p, cap, "return " as *u8)
1122 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, m[M_WGFRAGOUT], SIR_NAME)))
1123 p = eb_put(out, p, cap, ";\n" as *u8)
1124 } else { p = eb_put(out, p, cap, "return;\n" as *u8) }
1125 }
1126 } else {
1127 p = eb_put(out, p, cap, "return " as *u8)
1128 p = wg_expr(m, sir_get(m, s, SIR_A), out, p, cap, trace, tcap, tn)
1129 p = eb_put(out, p, cap, ";\n" as *u8)
1130 }
1131 }
1132 if k == S_DISCARD { p = eb_put(out, p, cap, "discard;\n" as *u8) }
1133 if k == S_IF {
1134 p = eb_put(out, p, cap, "if(" as *u8)
1135 p = wg_expr(m, sir_get(m, s, SIR_A), out, p, cap, trace, tcap, tn)
1136 p = eb_put(out, p, cap, "){\n" as *u8)
1137 p = wg_stmts(m, sir_get(m, s, SIR_B), out, p, cap, trace, tcap, tn)
1138 p = eb_put(out, p, cap, "}\n" as *u8)
1139 if sir_get(m, s, SIR_C) != 0 {
1140 p = eb_put(out, p, cap, "else{\n" as *u8)
1141 p = wg_stmts(m, sir_get(m, s, SIR_C), out, p, cap, trace, tcap, tn)
1142 p = eb_put(out, p, cap, "}\n" as *u8)
1143 }
1144 }
1145 if k == S_LOOP {
1146 p = eb_put(out, p, cap, "loop{\n" as *u8)
1147 p = wg_stmts(m, sir_get(m, s, SIR_A), out, p, cap, trace, tcap, tn)
1148 p = eb_put(out, p, cap, "}\n" as *u8)
1149 }
1150 if k == S_BREAK { p = eb_put(out, p, cap, "break;\n" as *u8) }
1151 if k == S_EXPR {
1152 p = wg_expr(m, sir_get(m, s, SIR_A), out, p, cap, trace, tcap, tn)
1153 p = eb_put(out, p, cap, ";\n" as *u8)
1154 }
1155 s = sir_get(m, s, SIR_NEXT)
1156 }
1157 return p
1158}
1159
1160// GE43 CONTRACT SYMBOL: the compute entry point. WGSL spells the dispatch geometry on the entry
1161// (@workgroup_size) and hands each invocation its global id as a builtin bound to `gid`, which is
1162// what wg_expr emits for the global_invocation_id builtin -- one name, owned in two places on
1163// purpose so a rename here breaks the gate rather than the kernel.
1164func wgsl_compute_stage(m: *i64, f: i64, out: *u8, pos: i64, cap: i64) -> i64 {
1165 var p: i64 = pos
1166 p = eb_put(out, p, cap, "@compute @workgroup_size(" as *u8)
1167 p = eb_num(out, p, cap, WG_WORKGROUP)
1168 p = eb_put(out, p, cap, ") fn " as *u8)
1169 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, f, SIR_NAME)))
1170 p = eb_put(out, p, cap, "(@builtin(global_invocation_id) gid:vec3u){\n" as *u8)
1171 return p
1172}
1173
1174// THE CONTRACT SYMBOL. Returns bytes written, or -1 with the refusal named in the module.
1175func wgsl_emit_module(m: *i64, stage: i64, out: *u8, cap: i64, trace: *i64, tcap: i64, tn: *i64) -> i64 {
1176 if wg_front_contract(m,stage)!=1{return -1}
1177 if m[M_CLIP_DEPTH] != SIR_CLIP_NATIVE && m[M_CLIP_DEPTH] != SIR_CLIP_NEGATIVE_ONE_TO_ONE { sir_refuse(m, "shader: unsupported source clip-depth convention" as *u8); return 0 - 1 }
1178 var p: i64 = 0
1179 var d: i64 = m[M_DECLH]
1180 var nu: i64 = 0
1181 while d != 0 {
1182 if sir_get(m, d, SIR_KIND) == K_UNIFORM { nu = nu + 1 }
1183 d = sir_get(m, d, SIR_NEXT)
1184 }
1185 // uniforms become ONE uniform struct + binding -- WGSL has no free-standing uniforms.
1186 if nu > 0 {
1187 p = eb_put(out, p, cap, "struct U{\n" as *u8)
1188 d = m[M_DECLH]
1189 while d != 0 {
1190 if sir_get(m, d, SIR_KIND) == K_UNIFORM {
1191 tr_add(trace, tcap, tn, d)
1192 // S1 (2026-09-04): A TEXTURE IS A BOUND RESOURCE, NEVER A STRUCT MEMBER. wg_ty returns a
1193 // non-null name for a texture type, so before this guard a texture-typed K_UNIFORM passed
1194 // the type check and was emitted INSIDE struct U -- WGSL that does not compile, produced
1195 // with no refusal. The remedy is the dedicated kind (sir_texture / K_TEXTURE).
1196 if wg_is_texty(sir_get(m, d, SIR_TY)) == 1 { sir_refuse(m, "wgsl: texture-typed uniform -- a texture is a bound resource, never a struct U member; declare it with sir_texture (K_TEXTURE), refused rather than emitted inside struct U" as *u8); return 0 - 1 }
1197 let t: *u8 = wg_ty(sir_get(m, d, SIR_TY))
1198 if (t as i64) == 0 { sir_refuse(m, "wgsl: uniform of a type outside the covered subset" as *u8); return 0 - 1 }
1199 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, d, SIR_NAME)))
1200 p = eb_put(out, p, cap, ":" as *u8)
1201 if sir_get(m, d, SIR_A) > 0 {
1202 // S12c-2b: float[N] lowers to array<vec4f,ceil(N/4)> (wg_f32arr_vec4n -- the same rule the layout ruler
1203 // and the E_INDEX arm apply); every other element type keeps its own spelling.
1204 if sir_get(m, d, SIR_TY) == T_F32 {
1205 p = eb_put(out, p, cap, "array<" as *u8)
1206 p = eb_put(out, p, cap, wg_ty(T_V4F))
1207 p = eb_put(out, p, cap, "," as *u8)
1208 p = eb_num(out, p, cap, wg_f32arr_vec4n(sir_get(m, d, SIR_A)))
1209 p = eb_put(out, p, cap, ">" as *u8)
1210 } else {
1211 p = eb_put(out, p, cap, "array<" as *u8)
1212 p = eb_put(out, p, cap, t)
1213 p = eb_put(out, p, cap, "," as *u8)
1214 p = eb_num(out, p, cap, sir_get(m, d, SIR_A))
1215 p = eb_put(out, p, cap, ">" as *u8)
1216 }
1217 } else { p = eb_put(out, p, cap, t) }
1218 p = eb_put(out, p, cap, ",\n" as *u8)
1219 }
1220 d = sir_get(m, d, SIR_NEXT)
1221 }
1222 p = eb_put(out, p, cap, "}\n@group(" as *u8)
1223 p = eb_num(out, p, cap, wg_group(m))
1224 p = eb_put(out, p, cap, ")@binding(" as *u8)
1225 p = eb_num(out, p, cap, WG_BIND_UNIFORM)
1226 p = eb_put(out, p, cap, ")var<uniform> u:U;\n" as *u8)
1227 }
1228 d = m[M_DECLH]
1229 while d != 0 {
1230 let k: i64 = sir_get(m, d, SIR_KIND)
1231 if k == K_PRIVATE {
1232 // GE55: module-scope private state -- `var<private> name:T;` -- emitted in declaration order beside the
1233 // other module declarations, never inside struct U (so the layout never sees it) and never qualified
1234 // (wg_resolves_to_uniform matches K_UNIFORM only, so every read and write of it stays bare).
1235 tr_add(trace, tcap, tn, d)
1236 if wg_is_texty(sir_get(m, d, SIR_TY)) == 1 { sir_refuse(m, "wgsl: texture-typed private variable -- a texture is a bound resource, never module state; refused so both backends cover the same source" as *u8); return 0 - 1 }
1237 let tp: *u8 = wg_ty(sir_get(m, d, SIR_TY))
1238 if (tp as i64) == 0 { sir_refuse(m, "wgsl: private variable of a type outside the covered subset" as *u8); return 0 - 1 }
1239 p = eb_put(out, p, cap, "var<private> " as *u8)
1240 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, d, SIR_NAME)))
1241 p = eb_put(out, p, cap, ":" as *u8)
1242 if sir_get(m, d, SIR_A) > 0 {
1243 p = eb_put(out, p, cap, "array<" as *u8)
1244 p = eb_put(out, p, cap, tp)
1245 p = eb_put(out, p, cap, "," as *u8)
1246 p = eb_num(out, p, cap, sir_get(m, d, SIR_A))
1247 p = eb_put(out, p, cap, ">" as *u8)
1248 } else { p = eb_put(out, p, cap, tp) }
1249 p = eb_put(out, p, cap, ";\n" as *u8)
1250 }
1251 if k == K_VARY {
1252 tr_add(trace, tcap, tn, d)
1253 // S7 (2026-09-05): varyings are ONE interstage struct in WGSL -- struct VIO{@builtin(position)
1254 // bpos:vec4f,@location(N) [@interpolate(flat) ]name:type,...} -- RETURNED by the vertex entry and
1255 // TAKEN as the fragment entry's parameter (see the entry emitters; reads/writes qualify vo./vin.
1256 // through M_VIOMODE). Emitted ONCE, at the first K_VARY the loop reaches, by walking every K_VARY;
1257 // each node is traced by ITS OWN arm visit (never inside the walk), so the WGSL trace matches the
1258 // GLSL loop-head trace exactly. Before S7 a vertex-stage varying became a COMMENT and a
1259 // fragment-stage one REFUSED; GLSL keeps its module-scope [flat ]out|in declarations unchanged.
1260 let t2: *u8 = wg_ty(sir_get(m, d, SIR_TY))
1261 if (t2 as i64) == 0 { sir_refuse(m, "wgsl: varying of a type outside the covered subset" as *u8); return 0 - 1 }
1262 if wg_name_is(m, d, "bpos" as *u8) == 1 { sir_refuse(m, "wgsl: a varying named bpos collides with the interstage struct's reserved @builtin(position) member; rename it" as *u8); return 0 - 1 }
1263 if wg_decl_kind_first(m, K_VARY) == d {
1264 p = eb_put(out, p, cap, "struct VIO{@builtin(position) bpos:vec4f," as *u8)
1265 var vd: i64 = m[M_DECLH]
1266 while vd != 0 {
1267 if sir_get(m, vd, SIR_KIND) == K_VARY {
1268 let vt: *u8 = wg_ty(sir_get(m, vd, SIR_TY))
1269 if (vt as i64) == 0 { sir_refuse(m, "wgsl: varying of a type outside the covered subset" as *u8); return 0 - 1 }
1270 p = eb_put(out, p, cap, "@location(" as *u8)
1271 p = eb_num(out, p, cap, sir_get(m, vd, SIR_A))
1272 p = eb_put(out, p, cap, ") " as *u8)
1273 if sir_get(m, vd, SIR_B) == 1 { p = eb_put(out, p, cap, "@interpolate(flat) " as *u8) }
1274 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, vd, SIR_NAME)))
1275 p = eb_put(out, p, cap, ":" as *u8)
1276 p = eb_put(out, p, cap, vt)
1277 p = eb_put(out, p, cap, "," as *u8)
1278 }
1279 vd = sir_get(m, vd, SIR_NEXT)
1280 }
1281 p = eb_put(out, p, cap, "}\n" as *u8)
1282 }
1283 }
1284 // R-D (2026-09-04): a struct declaration -- `struct Hit{pos:vec3f,d:f32,}` (the WGSL member separator is
1285 // the comma and a trailing comma is legal, so the emitter has one shape). Traced HERE because this
1286 // backend traces per arm; members through wg_tyname; the same refusals as the GLSL arm, so both
1287 // backends cover exactly the same source.
1288 if k == K_STRUCT {
1289 tr_add(trace, tcap, tn, d)
1290 p = eb_put(out, p, cap, "struct " as *u8)
1291 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, d, SIR_NAME)))
1292 p = eb_put(out, p, cap, "{" as *u8)
1293 var mb: i64 = sir_get(m, d, SIR_A)
1294 while mb != 0 {
1295 tr_add(trace, tcap, tn, mb)
1296 if wg_is_texty(sir_get(m, mb, SIR_TY)) == 1 { sir_refuse(m, "wgsl: texture-typed struct member -- a texture is a bound resource, never a value; refused in both dialects" as *u8); return 0 - 1 }
1297 let mt: *u8 = wg_tyname(m, sir_get(m, mb, SIR_TY))
1298 if (mt as i64) == 0 { sir_refuse(m, "wgsl: struct member of a type outside the covered subset" as *u8); return 0 - 1 }
1299 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, mb, SIR_NAME)))
1300 p = eb_put(out, p, cap, ":" as *u8)
1301 p = eb_put(out, p, cap, mt)
1302 p = eb_put(out, p, cap, "," as *u8)
1303 mb = sir_get(m, mb, SIR_NEXT)
1304 }
1305 p = eb_put(out, p, cap, "}\n" as *u8)
1306 }
1307 // S1 (2026-09-04): before this line a K_ATTRIB was traced and then DROPPED -- zero bytes, no
1308 // refusal, and a trace that still matched GLSL's. Vertex inputs are @location entry parameters
1309 // in WGSL and this backend does not derive them yet, so the honest answer is a named refusal.
1310 if k == K_ATTRIB {
1311 tr_add(trace, tcap, tn, d)
1312 // S6 (2026-09-05): a vertex attribute is a @location entry PARAMETER in WGSL (emitted in the
1313 // vertex entry signature below), never a module-scope declaration -- nothing is emitted here,
1314 // exactly as K_FRAGOUT emits nothing and becomes the fragment entry's return type. Traced here so
1315 // the WGSL trace matches the GLSL arm, which DOES emit it module-scope (layout(location=N) in ...).
1316 if stage != STAGE_VERTEX { sir_refuse(m, "wgsl: vertex attribute declaration (K_ATTRIB) outside a vertex stage -- attributes are vertex-stage inputs; refused rather than emitted where nothing feeds them" as *u8); return 0 - 1 }
1317 if (wg_ty(sir_get(m, d, SIR_TY)) as i64) == 0 { sir_refuse(m, "wgsl: vertex attribute of a type outside the covered subset" as *u8); return 0 - 1 }
1318 }
1319 // S5 (2026-09-04): the fragment output is the entry's RETURN TYPE in WGSL, so NO declaration is
1320 // emitted for it -- but it IS traced, because both backends CONSUME the node and trace identity is
1321 // the equivalence oracle on a box with no GPU. A backend that silently skipped it would produce a
1322 // shorter trace, which is exactly what that oracle exists to catch.
1323 if k == K_FRAGOUT {
1324 tr_add(trace, tcap, tn, d)
1325 if stage != STAGE_FRAGMENT { sir_refuse(m, "wgsl: fragment output declaration (K_FRAGOUT) outside a fragment stage -- only the fragment stage has a colour output; refused rather than accepted where nothing can return it" as *u8); return 0 - 1 }
1326 if wg_decl_kind_count(m, K_FRAGOUT) > 1 { sir_refuse(m, "wgsl: more than one fragment output declared -- multiple render targets are their own rung; refused rather than returning the first and silently dropping the rest" as *u8); return 0 - 1 }
1327 if (wg_ty(sir_get(m, d, SIR_TY)) as i64) == 0 { sir_refuse(m, "wgsl: fragment output of a type outside the covered subset" as *u8); return 0 - 1 }
1328 }
1329 // S3 (2026-09-04): a texture is a module-scope binding var, OUTSIDE struct U, at a binding DERIVED
1330 // from declaration order (see WG_BIND_TEX_FIRST). Only texture_3d<u32> is lowered: a sampled
1331 // texture_2d<f32> needs a sampler binding this backend does not declare, so it refuses by name.
1332 if k == K_TEXTURE {
1333 tr_add(trace, tcap, tn, d)
1334 let ttk: i64 = sir_get(m, d, SIR_TY)
1335 if ttk != T_TEX3U { if ttk != T_TEX2F { sir_refuse(m, "wgsl: texture type outside the covered subset -- texture_3d<u32> (textureLoad) and texture_2d<f32> (textureSample with a derived companion sampler) are lowered; refused rather than guessed" as *u8); return 0 - 1 } }
1336 let tb: i64 = wg_tex_binding(m, d)
1337 if wg_storage_at(m, tb) != 0 { sir_refuse(m, "wgsl: derived texture binding collides with a storage buffer declared at the same binding -- declare the storage buffer (sir_storage) above the texture bindings" as *u8); return 0 - 1 }
1338 p = eb_put(out, p, cap, "@group(" as *u8)
1339 p = eb_num(out, p, cap, wg_group(m))
1340 p = eb_put(out, p, cap, ")@binding(" as *u8)
1341 p = eb_num(out, p, cap, tb)
1342 p = eb_put(out, p, cap, ")var " as *u8)
1343 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, d, SIR_NAME)))
1344 p = eb_put(out, p, cap, ":" as *u8)
1345 p = eb_put(out, p, cap, wg_ty(ttk))
1346 p = eb_put(out, p, cap, ";\n" as *u8)
1347 // GE54 (2026-09-05): a sampled texture_2d<f32> carries a DERIVED companion sampler `<name>_s` at the
1348 // next binding (wg_tex_binding counts a 2D texture as two slots, so tb+1 is reserved and no later
1349 // binding collides). GLSL folds both into one combined sampler2D; WGSL keeps texture and sampler apart.
1350 if ttk == T_TEX2F {
1351 p = eb_put(out, p, cap, "@group(" as *u8)
1352 p = eb_num(out, p, cap, wg_group(m))
1353 p = eb_put(out, p, cap, ")@binding(" as *u8)
1354 p = eb_num(out, p, cap, tb + 1)
1355 p = eb_put(out, p, cap, ")var " as *u8)
1356 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, d, SIR_NAME)))
1357 p = eb_put(out, p, cap, "_s:sampler;\n" as *u8)
1358 }
1359 }
1360 d = sir_get(m, d, SIR_NEXT)
1361 }
1362 // GE43: storage buffers, one @binding each. WGSL has them; GLSL ES 3.00 does not (its backend
1363 // refuses the kind by name), so this pass has no twin on the GLSL side BY DESIGN.
1364 d = m[M_DECLH]
1365 while d != 0 {
1366 if sir_get(m, d, SIR_KIND) == K_STORAGE {
1367 tr_add(trace, tcap, tn, d)
1368 let t4: *u8 = wg_ty(sir_get(m, d, SIR_TY))
1369 if (t4 as i64) == 0 { sir_refuse(m, "wgsl: storage element type outside the covered subset" as *u8); return 0 - 1 }
1370 p = eb_put(out, p, cap, "@group(" as *u8)
1371 p = eb_num(out, p, cap, wg_group(m))
1372 p = eb_put(out, p, cap, ")@binding(" as *u8)
1373 p = eb_num(out, p, cap, sir_get(m, d, SIR_A))
1374 if sir_get(m, d, SIR_B) == 1 { p = eb_put(out, p, cap, ")var<storage,read_write> " as *u8) } else { p = eb_put(out, p, cap, ")var<storage,read> " as *u8) }
1375 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, d, SIR_NAME)))
1376 p = eb_put(out, p, cap, ":array<" as *u8)
1377 if sir_get(m, d, SIR_C) == 1 {
1378 p = eb_put(out, p, cap, "atomic<" as *u8)
1379 p = eb_put(out, p, cap, t4)
1380 p = eb_put(out, p, cap, ">" as *u8)
1381 } else { p = eb_put(out, p, cap, t4) }
1382 p = eb_put(out, p, cap, ">;\n" as *u8)
1383 }
1384 d = sir_get(m, d, SIR_NEXT)
1385 }
1386 var f: i64 = m[M_FUNCH]
1387 while f != 0 {
1388 tr_add(trace, tcap, tn, f)
1389 m[M_CURFN] = f
1390 if wg_fn_is_plain(m, f, stage) == 1 {
1391 // R-E (2026-09-04): struct returns and struct parameters spell through wg_tyname (see the GLSL twin).
1392 let rt2: *u8 = wg_tyname(m, sir_get(m, f, SIR_TY))
1393 if (rt2 as i64) == 0 { sir_refuse(m, "wgsl: return type outside the covered subset" as *u8); return 0 - 1 }
1394 p = eb_put(out, p, cap, "fn " as *u8)
1395 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, f, SIR_NAME)))
1396 p = eb_put(out, p, cap, "(" as *u8)
1397 var pb: i64 = sir_get(m, f, SIR_A)
1398 var pg: i64 = 1
1399 while pb != 0 {
1400 if pg == 0 { p = eb_put(out, p, cap, "," as *u8) }
1401 pg = 0
1402 tr_add(trace, tcap, tn, pb)
1403 let pt2: *u8 = wg_tyname(m, sir_get(m, pb, SIR_TY))
1404 if (pt2 as i64) == 0 { sir_refuse(m, "wgsl: parameter type outside the covered subset" as *u8); return 0 - 1 }
1405 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, pb, SIR_NAME)))
1406 p = eb_put(out, p, cap, ":" as *u8)
1407 p = eb_put(out, p, cap, pt2)
1408 pb = sir_get(m, pb, SIR_NEXT)
1409 }
1410 // GE55 (2026-09-05): a function that returns nothing has NO arrow in WGSL -- `fn upk(){` -- because
1411 // `->void` is not WGSL and a browser refuses the whole module. Measured on the first void non-entry
1412 // function this backend ever emitted (the world pass upk); GLSL already spelled `void upk(){` correctly.
1413 if sir_get(m, f, SIR_TY) == T_VOID { p = eb_put(out, p, cap, "){\n" as *u8) } else {
1414 p = eb_put(out, p, cap, ")->" as *u8)
1415 p = eb_put(out, p, cap, rt2)
1416 p = eb_put(out, p, cap, "{\n" as *u8)
1417 }
1418 } else { if stage == STAGE_VERTEX {
1419 // S4: the vertex entry binds vertex_index itself; a declared parameter is an input this backend
1420 // does not derive yet (K_ATTRIB refuses for the same reason) -- refused rather than dropped.
1421 if sir_get(m, f, SIR_A) != 0 { sir_refuse(m, "wgsl: vertex entry parameter -- vertex inputs are not derived from the declaration yet (the vertex_index builtin is bound by the backend); refused rather than silently dropped" as *u8); return 0 - 1 }
1422 p = eb_put(out, p, cap, "@vertex fn " as *u8)
1423 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, f, SIR_NAME)))
1424 p = eb_put(out, p, cap, "(@builtin(vertex_index) vi:u32" as *u8)
1425 // S12a (2026-09-05): the instance index is bound ONLY when the module read it (sir_builtin set M_USESII),
1426 // so every module that never asks keeps a byte-identical signature. GLSL needs no declaration.
1427 if m[M_USESII] == 1 { p = eb_put(out, p, cap, ",@builtin(instance_index) iidx:u32" as *u8) }
1428 // S6 (2026-09-05): declared vertex attributes are @location entry parameters, appended after the
1429 // backend-bound vertex_index in declaration order. GLSL binds them module-scope; WGSL binds them
1430 // here. The location is the attribute's own SIR_A, so a mesh declaring aP/aJ/aW/aN/aUV/aHS derives
1431 // its own signature, and the fullscreen triangle (no attributes) is byte-identical to before.
1432 // Not traced here: the K_ATTRIB was traced once in the declarations loop, matching the GLSL arm.
1433 var vsa: i64 = m[M_DECLH]
1434 while vsa != 0 {
1435 if sir_get(m, vsa, SIR_KIND) == K_ATTRIB {
1436 p = eb_put(out, p, cap, ",@location(" as *u8)
1437 p = eb_num(out, p, cap, sir_get(m, vsa, SIR_A))
1438 p = eb_put(out, p, cap, ") " as *u8)
1439 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, vsa, SIR_NAME)))
1440 p = eb_put(out, p, cap, ":" as *u8)
1441 p = eb_put(out, p, cap, wg_ty(sir_get(m, vsa, SIR_TY)))
1442 }
1443 vsa = sir_get(m, vsa, SIR_NEXT)
1444 }
1445 // S7 (2026-09-05): with varyings declared the vertex entry RETURNS the interstage struct VIO --
1446 // `var vo:VIO;` opens the body, the position assignment and every varying write land in vo.<name>
1447 // (wg_stmts / wg_expr read M_VIOMODE=1), and `return vo;` closes it after the body. With none
1448 // declared the entry keeps ->@builtin(position) vec4f and `return <expr>;` -- byte-identical.
1449 if wg_decl_kind_count(m, K_VARY) > 0 {
1450 p = eb_put(out, p, cap, ")->VIO{\nvar vo:VIO;\n" as *u8)
1451 m[M_VIOMODE] = 1
1452 } else {
1453 p = eb_put(out, p, cap, ")->@builtin(position) vec4f{\n" as *u8)
1454 }
1455 } else { if stage == STAGE_COMPUTE {
1456 // S4: the compute entry binds global_invocation_id itself (wgsl_compute_stage); same rule.
1457 if sir_get(m, f, SIR_A) != 0 { sir_refuse(m, "wgsl: compute entry parameter -- the compute entry binds global_invocation_id itself; a declared parameter is not derived yet, refused rather than silently dropped" as *u8); return 0 - 1 }
1458 p = wgsl_compute_stage(m, f, out, p, cap)
1459 } else {
1460 p = eb_put(out, p, cap, "@fragment fn " as *u8)
1461 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, f, SIR_NAME)))
1462 p = eb_put(out, p, cap, "(" as *u8)
1463 var fp: i64 = sir_get(m, f, SIR_A)
1464 var ff: i64 = 1
1465 // S7 (2026-09-05): with varyings declared the fragment entry TAKES the interstage struct -- `vin:VIO`
1466 // -- and every varying read qualifies vin.<name> (M_VIOMODE=2). @builtin(position) then arrives as
1467 // vin.bpos, so a declared position parameter (S4) FOLDS into it: traced here (identity with the GLSL
1468 // arm, which traces it in its own loop) and not emitted, because a second @builtin(position)
1469 // binding is invalid WGSL. Any other parameter still refuses by name.
1470 if wg_decl_kind_count(m,K_VARY)>0{
1471 p=eb_put(out,p,cap,"vin:VIO");ff=0;m[M_VIOMODE]=2
1472 }
1473 while fp!=0{
1474 tr_add(trace,tcap,tn,fp)
1475 let binding:i64=sir_get(m,fp,SIR_A)
1476 if binding!=P_BUILTIN_POSITION&&binding!=P_BUILTIN_FRONT_FACING{sir_refuse(m,"wgsl: unsupported fragment entry parameter binding");return -1}
1477 if binding!=P_BUILTIN_POSITION||m[M_VIOMODE]!=2{
1478 if ff==0{p=eb_put(out,p,cap,",")};ff=0
1479 if binding==P_BUILTIN_FRONT_FACING{p=eb_put(out,p,cap,"@builtin(front_facing) ")}else{p=eb_put(out,p,cap,"@builtin(position) ")}
1480 p=eb_put(out,p,cap,sir_cstr(m,sir_get(m,fp,SIR_NAME)));p=eb_put(out,p,cap,":")
1481 if binding==P_BUILTIN_FRONT_FACING{p=eb_put(out,p,cap,wg_ty(T_BOOL))}else{p=eb_put(out,p,cap,wg_ty(T_V4F))}
1482 }
1483 fp=sir_get(m,fp,SIR_NEXT)
1484 }
1485 // S5 (2026-09-04): THE RETURN HALF OF THE SIGNATURE IS DERIVED TOO. It was the literal
1486 // ")->@location(0) vec4f{" -- so the location and the type were unreadable constants rather than
1487 // data, the same defect S4 had just removed from the PARAMETER half of this very line. Both
1488 // halves now come from declarations. A fragment that only discards declares no output and gets
1489 // no return type; one that returns a value with none declared refuses rather than assuming.
1490 let fo2: i64 = wg_fragout(m)
1491 if fo2 == 0 {
1492 if wg_chain_returns_value(m, sir_get(m, f, SIR_B)) == 1 { sir_refuse(m, "wgsl: fragment entry returns a value with no declared fragment output -- declare it with sir_fragout (K_FRAGOUT) so the @location and the type are data; refused rather than assuming @location(0) vec4f" as *u8); return 0 - 1 }
1493 p = eb_put(out, p, cap, "){\n" as *u8)
1494 } else {
1495 p = eb_put(out, p, cap, ")->@location(" as *u8)
1496 p = eb_num(out, p, cap, sir_get(m, fo2, SIR_A))
1497 p = eb_put(out, p, cap, ") " as *u8)
1498 p = eb_put(out, p, cap, wg_ty(sir_get(m, fo2, SIR_TY)))
1499 p = eb_put(out, p, cap, "{\n" as *u8)
1500 // S12c-5 (2026-09-05): a body that ASSIGNS the declared output (the hand cast fragment's `fc=...;return;`
1501 // shape) gets the output as a local variable; every void return inside it returns that variable (wg_stmts
1502 // reads M_WGFRAGOUT) and one trailing `return fc;` closes the body when its last statement is not a return.
1503 // A body that RETURNS its value (the world pass) is untouched: no local, no trailing return, byte-identical.
1504 if wg_chain_assigns(m, sir_get(m, f, SIR_B), sir_cstr(m, sir_get(m, fo2, SIR_NAME))) == 1 {
1505 p = eb_put(out, p, cap, "var " as *u8)
1506 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, fo2, SIR_NAME)))
1507 p = eb_put(out, p, cap, ":" as *u8)
1508 p = eb_put(out, p, cap, wg_ty(sir_get(m, fo2, SIR_TY)))
1509 p = eb_put(out, p, cap, ";\n" as *u8)
1510 m[M_WGFRAGOUT] = fo2
1511 }
1512 }
1513 } } }
1514 p = wg_stmts(m, sir_get(m, f, SIR_B), out, p, cap, trace, tcap, tn)
1515 // S7: a vertex entry that returns the interstage struct closes with `return vo;`; the mode is cleared
1516 // after EVERY function so a helper declared after an entry cannot inherit the vo./vin. qualification.
1517 if m[M_VIOMODE] == 1 { p = eb_put(out, p, cap, "return vo;\n" as *u8) }
1518 m[M_VIOMODE] = 0
1519 // S12c-5: a fragment entry carrying its output as a local closes with `return fc;` unless its last statement already
1520 // returns; the word is cleared after EVERY function so a helper declared after the entry cannot inherit it.
1521 if m[M_WGFRAGOUT] != 0 {
1522 if wg_chain_last_is_return(m, sir_get(m, f, SIR_B)) == 0 {
1523 p = eb_put(out, p, cap, "return " as *u8)
1524 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, m[M_WGFRAGOUT], SIR_NAME)))
1525 p = eb_put(out, p, cap, ";\n" as *u8)
1526 }
1527 m[M_WGFRAGOUT] = 0
1528 }
1529 p = eb_put(out, p, cap, "}\n" as *u8)
1530 f = sir_get(m, f, SIR_NEXT)
1531 }
1532 m[M_CURFN] = 0
1533 if sir_refused(m) != 0 { return 0 - 1 }
1534 if p + 1 < cap { out[p] = 0 as u8 }
1535 return p
1536}
1537
1538// ---- THE SHADER, AUTHORED IN NISHILANG --------------------------------------------------------
1539// This function IS the shader source. No GLSL and no WGSL is written anywhere in this estate for
1540// this stage any more -- both dialects above are emitted from exactly these calls.
1541//
1542// vertex stage: fullscreen triangle
1543// var v: vec2 = vec2(float((vertex_index << 1) & 2), float(vertex_index & 2))
1544// position = vec4(v * 2.0 - 1.0, 0.0, 1.0)
1545func shsrc_fullscreen_tri(m: *i64) -> i64 {
1546 let fn: i64 = sir_func(m, "main" as *u8, T_VOID)
1547 let vi1: i64 = sir_builtin(m, "vertex_index" as *u8, T_I32)
1548 let sh: i64 = sir_bin(m, "<<" as *u8, vi1, sir_lit(m, "1" as *u8, T_U32), T_I32)
1549 let m1: i64 = sir_bin(m, "&" as *u8, sh, sir_lit(m, "2" as *u8, T_I32), T_I32)
1550 let vi2: i64 = sir_builtin(m, "vertex_index" as *u8, T_I32)
1551 let m2: i64 = sir_bin(m, "&" as *u8, vi2, sir_lit(m, "2" as *u8, T_I32), T_I32)
1552 let c1: i64 = sir_ctor(m, T_V2F)
1553 sir_arg(m, c1, sir_cast(m, m1, T_F32))
1554 sir_arg(m, c1, sir_cast(m, m2, T_F32))
1555 sir_stmt(m, fn, sir_var(m, "v" as *u8, T_V2F, c1))
1556 let mul: i64 = sir_bin(m, "*" as *u8, sir_ident(m, "v" as *u8, T_V2F), sir_lit(m, "2.0" as *u8, T_F32), T_V2F)
1557 let sub: i64 = sir_bin(m, "-" as *u8, mul, sir_lit(m, "1.0" as *u8, T_F32), T_V2F)
1558 let c2: i64 = sir_ctor(m, T_V4F)
1559 sir_arg(m, c2, sub)
1560 sir_arg(m, c2, sir_lit(m, "0.0" as *u8, T_F32))
1561 sir_arg(m, c2, sir_lit(m, "1.0" as *u8, T_F32))
1562 sir_stmt(m, fn, sir_assign(m, sir_builtin(m, "position" as *u8, T_V4F), c2))
1563 return fn
1564}
1565
1566// GE43 -- THE FIRST KERNEL PRIMITIVE, AUTHORED IN NISHILANG: a depth test resolved by atomicMin over
1567// a u32 depth buffer. This is the byte a software rasterizer cannot skip -- every tile, every
1568// covered fragment, one atomicMin -- and it is INTEGER end to end, which is the determinism exceed
1569// the gameengine board claims over the float fixed-function pipelines of WebGL2 and its ancestors.
1570// compute stage, WG_WORKGROUP invocations per workgroup:
1571// storage depth: array<atomic<u32>> @binding(1) read_write
1572// storage zin: array<u32> @binding(2) read
1573// var i: u32 = global_invocation_id.x
1574// atomicMin(&depth[i], zin[i])
1575func shsrc_depth_min(m: *i64) -> i64 {
1576 sir_storage(m, "depth" as *u8, T_U32, 1, 1, 1)
1577 sir_storage(m, "zin" as *u8, T_U32, 2, 0, 0)
1578 let fn: i64 = sir_func(m, "main" as *u8, T_VOID)
1579 let gx: i64 = sir_swz(m, sir_builtin(m, "global_invocation_id" as *u8, T_V3U), "x" as *u8, T_U32)
1580 sir_stmt(m, fn, sir_var(m, "i" as *u8, T_U32, gx))
1581 let tgt: i64 = sir_index(m, sir_ident(m, "depth" as *u8, T_U32), sir_ident(m, "i" as *u8, T_U32), T_U32)
1582 let src: i64 = sir_index(m, sir_ident(m, "zin" as *u8, T_U32), sir_ident(m, "i" as *u8, T_U32), T_U32)
1583 sir_stmt(m, fn, sir_expr(m, sir_atomic(m, "atomicMin" as *u8, tgt, src, T_U32)))
1584 return fn
1585}
1586
1587func wg_puts(s: *u8) -> i64 { var n: i64 = 0; while s[n] != (0 as u8) { n = n + 1 } sys_write(1, s, n); return 0 }
1588
1589func wg_streq(a: *u8, b: *u8) -> i64 {
1590 var i: i64 = 0
1591 while a[i] != (0 as u8) { if a[i] != b[i] { return 0 } i = i + 1 }
1592 if b[i] != (0 as u8) { return 0 }
1593 return 1
1594}
1595
1596// `nx_wgsl compute` -- the GE43 kernel through BOTH backends: WGSL emits, GLSL refuses BY NAME.
1597// nx_wgsl_compute_gate asserts these bytes exactly, so this output IS the contract.
1598func wg_compute_demo() -> i64 {
1599 let m: *i64 = sir_new()
1600 shsrc_depth_min(m)
1601 let wb: *u8 = sys_mmap(WG_EMIT_CAP)
1602 let wt: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
1603 let wn: *i64 = sys_mmap(WG_NUMBUF) as *i64
1604 wn[0] = 0
1605 let wl: i64 = wgsl_emit_module(m, STAGE_COMPUTE, wb, WG_EMIT_CAP, wt, WG_TRACE_CAP, wn)
1606 wg_puts("=== NX-WGSL GE43: COMPUTE STAGE -- storage buffers, atomics, global_invocation_id, ONE NishiLang source ===\n\n--- WGSL ---\n" as *u8)
1607 if wl > 0 { sys_write(1, wb, wl) }
1608 let nb: *u8 = sys_mmap(WG_NUMBUF)
1609 var q: i64 = 0
1610 wg_puts("\nwgsl_bytes=" as *u8)
1611 q = eb_num(nb, 0, WG_NUMBUF, wl); nb[q] = 0 as u8; wg_puts(nb)
1612 wg_puts(" wgsl_nodes=" as *u8)
1613 q = eb_num(nb, 0, WG_NUMBUF, wn[0]); nb[q] = 0 as u8; wg_puts(nb)
1614 wg_puts(" wgsl_refused=" as *u8)
1615 if sir_refused(m) != 0 { wg_puts(sir_cstr(m, sir_refused(m))) } else { wg_puts("NONE" as *u8) }
1616 // THE GLSL SIDE MUST REFUSE BY NAME. A WebGL2 program that silently dropped the storage buffers
1617 // would compile, link, run and draw nothing, forever -- the failure this backend exists to make impossible.
1618 let m2: *i64 = sir_new()
1619 shsrc_depth_min(m2)
1620 let gb: *u8 = sys_mmap(WG_EMIT_CAP)
1621 let gt: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
1622 let gn: *i64 = sys_mmap(WG_NUMBUF) as *i64
1623 gn[0] = 0
1624 let gl: i64 = glsl_emit_module(m2, STAGE_COMPUTE, gb, WG_EMIT_CAP, gt, WG_TRACE_CAP, gn)
1625 wg_puts(" glsl_compute_rc=" as *u8)
1626 q = eb_num(nb, 0, WG_NUMBUF, gl); nb[q] = 0 as u8; wg_puts(nb)
1627 wg_puts(" glsl_compute_named=" as *u8)
1628 if sir_refused(m2) != 0 { wg_puts(sir_cstr(m2, sir_refused(m2))) } else { wg_puts("NONE" as *u8) }
1629 // NEG-CONTROL -- the GLSL refusal must be the STAGE refusal, not a later accident: emit the same
1630 // kernel as a GLSL FRAGMENT stage and the storage DECLARATION must refuse by its own name.
1631 let m3: *i64 = sir_new()
1632 shsrc_depth_min(m3)
1633 let g3: *u8 = sys_mmap(WG_EMIT_CAP)
1634 let t3: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
1635 let n3: *i64 = sys_mmap(WG_NUMBUF) as *i64
1636 n3[0] = 0
1637 let r3: i64 = glsl_emit_module(m3, STAGE_FRAGMENT, g3, WG_EMIT_CAP, t3, WG_TRACE_CAP, n3)
1638 wg_puts(" glsl_fragment_of_kernel_rc=" as *u8)
1639 q = eb_num(nb, 0, WG_NUMBUF, r3); nb[q] = 0 as u8; wg_puts(nb)
1640 wg_puts(" glsl_fragment_of_kernel_named=" as *u8)
1641 if sir_refused(m3) != 0 { wg_puts(sir_cstr(m3, sir_refused(m3))) } else { wg_puts("NONE" as *u8) }
1642 wg_puts("\n" as *u8)
1643 return 0
1644}
1645
1646// ---- `nx_wgsl decls` -- THE DECLARATION PROBES (S1..S4, 2026-09-04) --------------------------------
1647// One module through one backend, printed the way every other verb prints: the emitted text (when it
1648// emitted), then `<label>_rc=` and `<label>_named=` on ONE line for nx_wgsl_gate to anchor on. The
1649// refusal that fired is printed VERBATIM so the gate asserts WHICH rule fired, never merely that one did.
1650func wg_probe(m: *i64, stage: i64, label: *u8, dialect_wgsl: i64, trace: *i64, tn: *i64) -> i64 {
1651 let buf: *u8 = sys_mmap(WG_EMIT_CAP)
1652 let nb: *u8 = sys_mmap(WG_NUMBUF)
1653 tn[0] = 0
1654 var rc: i64 = 0
1655 if dialect_wgsl == 1 { rc = wgsl_emit_module(m, stage, buf, WG_EMIT_CAP, trace, WG_TRACE_CAP, tn) } else { rc = glsl_emit_module(m, stage, buf, WG_EMIT_CAP, trace, WG_TRACE_CAP, tn) }
1656 wg_puts("--- " as *u8); wg_puts(label)
1657 if dialect_wgsl == 1 { wg_puts(" WGSL ---\n" as *u8) } else { wg_puts(" GLSL ---\n" as *u8) }
1658 if rc > 0 { sys_write(1, buf, rc) }
1659 wg_puts(label); wg_puts("_rc=" as *u8)
1660 var q: i64 = eb_num(nb, 0, WG_NUMBUF, rc); nb[q] = 0 as u8; wg_puts(nb)
1661 wg_puts(" " as *u8); wg_puts(label); wg_puts("_named=" as *u8)
1662 if sir_refused(m) != 0 { wg_puts(sir_cstr(m, sir_refused(m))) } else { wg_puts("NONE" as *u8) }
1663 wg_puts("\n" as *u8)
1664 return rc
1665}
1666
1667// Element-wise trace identity: same count AND same ids in the same order. The oracle main() applies
1668// inline, extracted once so the probes cannot drift from it.
1669func wg_trace_eq(a: *i64, an: i64, b: *i64, bn: i64) -> i64 {
1670 if an != bn { return 0 }
1671 var i: i64 = 0
1672 while i < an { if a[i] != b[i] { return 0 } i = i + 1 }
1673 return 1
1674}
1675
1676// A minimal fragment stage: one constant colour. The declaration probes need a fragment FUNCTION to
1677// reach the fragment-stage decl pass; this is the smallest one that is a real fragment program.
1678// fragment stage: return vec4(1.0, 0.0, 0.0, 1.0)
1679func shsrc_frag_const(m: *i64) -> i64 {
1680 sir_fragout(m, "fc" as *u8, T_V4F, 0)
1681 let fn: i64 = sir_func(m, "main" as *u8, T_VOID)
1682 let c: i64 = sir_ctor(m, T_V4F)
1683 sir_arg(m, c, sir_lit(m, "1.0" as *u8, T_F32))
1684 sir_arg(m, c, sir_lit(m, "0.0" as *u8, T_F32))
1685 sir_arg(m, c, sir_lit(m, "0.0" as *u8, T_F32))
1686 sir_arg(m, c, sir_lit(m, "1.0" as *u8, T_F32))
1687 sir_stmt(m, fn, sir_return(m, c))
1688 return fn
1689}
1690
1691// POSITIVE CONTROL for the declaration guards -- the fullscreen triangle with ONE scalar uniform folded
1692// into the clip transform. A deny-guard that has only ever refused is unverified (a guard that refuses
1693// everything passes every negative test), so the same declaration pass must be shown to ADMIT a scalar
1694// uniform; the identifier is then read through it in both dialects.
1695// uniform scale: f32
1696// vertex stage: var v = vec2(...); v = v * scale; position = vec4(v*2.0-1.0, 0.0, 1.0)
1697func shsrc_fullscreen_tri_scaled(m: *i64) -> i64 {
1698 sir_uniform(m, "scale" as *u8, T_F32, 0)
1699 let fn: i64 = shsrc_fullscreen_tri(m)
1700 // splice `v = v * scale` between the two statements the triangle already carries
1701 let s1: i64 = sir_get(m, fn, SIR_B)
1702 let s2: i64 = sir_get(m, s1, SIR_NEXT)
1703 let mul: i64 = sir_bin(m, "*" as *u8, sir_ident(m, "v" as *u8, T_V2F), sir_ident(m, "scale" as *u8, T_F32), T_V2F)
1704 let asg: i64 = sir_assign(m, sir_ident(m, "v" as *u8, T_V2F), mul)
1705 sir_set(m, s1, SIR_NEXT, asg)
1706 sir_set(m, asg, SIR_NEXT, s2)
1707 return fn
1708}
1709
1710// S2 NEG-CONTROL -- the triangle with a uniform `k` AND a local `k`: the local shadows the uniform, so the
1711// read is the bare `k`, never `u.k`. A resolver that qualified on name alone would fail this.
1712// uniform k: f32
1713// vertex stage: var v = vec2(...); var k = 2.0; v = v * k; position = ...
1714func shsrc_fullscreen_tri_shadowlocal(m: *i64) -> i64 {
1715 sir_uniform(m, "k" as *u8, T_F32, 0)
1716 let fn: i64 = shsrc_fullscreen_tri(m)
1717 let s1: i64 = sir_get(m, fn, SIR_B)
1718 let s2: i64 = sir_get(m, s1, SIR_NEXT)
1719 let lk: i64 = sir_var(m, "k" as *u8, T_F32, sir_lit(m, "2.0" as *u8, T_F32))
1720 let mul: i64 = sir_bin(m, "*" as *u8, sir_ident(m, "v" as *u8, T_V2F), sir_ident(m, "k" as *u8, T_F32), T_V2F)
1721 let asg: i64 = sir_assign(m, sir_ident(m, "v" as *u8, T_V2F), mul)
1722 sir_set(m, s1, SIR_NEXT, lk)
1723 sir_set(m, lk, SIR_NEXT, asg)
1724 sir_set(m, asg, SIR_NEXT, s2)
1725 return fn
1726}
1727
1728// S2 NEG-CONTROL -- a plain function whose PARAMETER is named like a uniform: `fn f(k:f32)` reads its own
1729// `k`, never `u.k`.
1730// uniform k: f32
1731// func f(k: f32) -> f32 { return k * 2.0 }
1732func shsrc_plain_shadowparam(m: *i64) -> i64 {
1733 sir_uniform(m, "k" as *u8, T_F32, 0)
1734 let fn: i64 = sir_func(m, "f" as *u8, T_F32)
1735 sir_param(m, fn, "k" as *u8, T_F32)
1736 sir_stmt(m, fn, sir_return(m, sir_bin(m, "*" as *u8, sir_ident(m, "k" as *u8, T_F32), sir_lit(m, "2.0" as *u8, T_F32), T_F32)))
1737 return fn
1738}
1739
1740// S3 COVERAGE -- a plain function reading a 3D u32 texture beside one scalar uniform: the texture must come
1741// out as a binding var OUTSIDE struct U in WGSL (struct U carries only the scalar) and as a sampler
1742// uniform in GLSL, and the texture identifier is read BARE in both (it is not a uniform).
1743// uniform scale: f32 ; texture vox: texture_3d<u32>
1744// func vox_at(x: i32, y: i32, z: i32) -> u32 { return texelFetch(vox, ivec3(x, y, z), 0).x }
1745// K_PRIVATE fixture (GE55): module-scope private state -- `var<private>` in WGSL, a bare global in GLSL -- a scalar
1746// and an array, written by one function and read bare by another, beside a uniform of a DIFFERENT name so the
1747// qualifier rule is exercised (u.q is qualified, gk and pv never are). STAGE_PLAIN: no entry point.
1748func shsrc_private_demo(m: *i64) -> i64 {
1749 sir_uniform(m, "q" as *u8, T_F32, 0)
1750 sir_private(m, "gk" as *u8, T_F32, 0)
1751 sir_private(m, "pv" as *u8, T_V4F, 2)
1752 let f: i64 = sir_func(m, "upk" as *u8, T_VOID)
1753 sir_stmt(m, f, sir_assign(m, sir_ident(m, "gk" as *u8, T_F32), sir_bin(m, "*" as *u8, sir_ident(m, "q" as *u8, T_F32), sir_lit(m, "2.0" as *u8, T_F32), T_F32)))
1754 let ctor: i64 = sir_ctor(m, T_V4F)
1755 sir_arg(m, ctor, sir_ident(m, "gk" as *u8, T_F32))
1756 sir_stmt(m, f, sir_assign(m, sir_index(m, sir_ident(m, "pv" as *u8, T_V4F), sir_lit(m, "0" as *u8, T_I32), T_V4F), ctor))
1757 let g: i64 = sir_func(m, "rd" as *u8, T_F32)
1758 sir_stmt(m, g, sir_return(m, sir_bin(m, "+" as *u8, sir_ident(m, "gk" as *u8, T_F32), sir_swz(m, sir_index(m, sir_ident(m, "pv" as *u8, T_V4F), sir_lit(m, "0" as *u8, T_I32), T_V4F), "x" as *u8, T_F32), T_F32)))
1759 return 0
1760}
1761
1762// GE54 fixture: a sampled texture_2d<f32> -- the cast fragment's need (wc_mob_tex). Declares `atlas` (T_TEX2F) and a
1763// plain function that samples it at a uv param; GLSL lowers to sampler2D + texture(), WGSL to a texture + a derived
1764// companion sampler + textureSample(). STAGE_PLAIN (no entry point) -- a decl+function module, valid to compile.
1765func shsrc_tex2d_sample(m: *i64) -> i64 {
1766 sir_texture(m, "atlas" as *u8, T_TEX2F)
1767 let f: i64 = sir_func(m, "smp" as *u8, T_V4F)
1768 sir_param(m, f, "uv" as *u8, T_V2F)
1769 sir_stmt(m, f, sir_return(m, sir_texsample(m, sir_ident(m, "atlas" as *u8, T_TEX2F), sir_ident(m, "uv" as *u8, T_V2F))))
1770 return 0
1771}
1772
1773func shsrc_tex_plain(m: *i64) -> i64 {
1774 sir_uniform(m, "scale" as *u8, T_F32, 0)
1775 sir_texture(m, "vox" as *u8, T_TEX3U)
1776 let fn: i64 = sir_func(m, "vox_at" as *u8, T_U32)
1777 sir_param(m, fn, "x" as *u8, T_I32)
1778 sir_param(m, fn, "y" as *u8, T_I32)
1779 sir_param(m, fn, "z" as *u8, T_I32)
1780 let c: i64 = sir_ctor(m, T_V3I)
1781 sir_arg(m, c, sir_ident(m, "x" as *u8, T_I32))
1782 sir_arg(m, c, sir_ident(m, "y" as *u8, T_I32))
1783 sir_arg(m, c, sir_ident(m, "z" as *u8, T_I32))
1784 let ld: i64 = sir_texload(m, sir_ident(m, "vox" as *u8, T_TEX3U), c, sir_lit(m, "0" as *u8, T_I32))
1785 sir_stmt(m, fn, sir_return(m, sir_swz(m, ld, "x" as *u8, T_U32)))
1786 return fn
1787}
1788
1789// S4 COVERAGE -- a fragment that READS its position. The WGSL entry signature is derived from the declared
1790// parameter (`@builtin(position) pos:vec4f`) and the body reads `pos`; GLSL reads it as gl_FragCoord.
1791// fragment stage: param pos = @builtin(position)
1792// return vec4(pos.x, pos.y, 0.0, 1.0)
1793func shsrc_frag_pos(m: *i64) -> i64 {
1794 sir_fragout(m, "fc" as *u8, T_V4F, 0)
1795 let fn: i64 = sir_func(m, "main" as *u8, T_VOID)
1796 sir_param_position(m, fn, "pos" as *u8)
1797 let c: i64 = sir_ctor(m, T_V4F)
1798 sir_arg(m, c, sir_swz(m, sir_ident(m, "pos" as *u8, T_V4F), "x" as *u8, T_F32))
1799 sir_arg(m, c, sir_swz(m, sir_ident(m, "pos" as *u8, T_V4F), "y" as *u8, T_F32))
1800 sir_arg(m, c, sir_lit(m, "0.0" as *u8, T_F32))
1801 sir_arg(m, c, sir_lit(m, "1.0" as *u8, T_F32))
1802 sir_stmt(m, fn, sir_return(m, c))
1803 return fn
1804}
1805
1806// S4 NEG-CONTROL -- a fragment that reads the position BUILTIN without declaring the parameter: an undeclared
1807// input, refused by name (before S4 it emitted the placeholder identifier POSITION with no refusal).
1808// fragment stage: return vec4(position.x, 0.0, 0.0, 1.0) -- no parameter declared
1809func shsrc_frag_pos_undeclared(m: *i64) -> i64 {
1810 sir_fragout(m, "fc" as *u8, T_V4F, 0)
1811 let fn: i64 = sir_func(m, "main" as *u8, T_VOID)
1812 let c: i64 = sir_ctor(m, T_V4F)
1813 sir_arg(m, c, sir_swz(m, sir_builtin(m, "position" as *u8, T_V4F), "x" as *u8, T_F32))
1814 sir_arg(m, c, sir_lit(m, "0.0" as *u8, T_F32))
1815 sir_arg(m, c, sir_lit(m, "0.0" as *u8, T_F32))
1816 sir_arg(m, c, sir_lit(m, "1.0" as *u8, T_F32))
1817 sir_stmt(m, fn, sir_return(m, c))
1818 return fn
1819}
1820
1821// S4 NEG-CONTROL -- a fragment entry with a PLAIN parameter (no builtin binding): neither backend can bind it
1822// yet (it would be a @location varying), so both refuse by name rather than dropping it from the signature.
1823// fragment stage: param uv: vec2 (plain) return vec4(1.0, 0.0, 0.0, 1.0)
1824func shsrc_frag_plainparam(m: *i64) -> i64 {
1825 sir_fragout(m, "fc" as *u8, T_V4F, 0)
1826 let fn: i64 = sir_func(m, "main" as *u8, T_VOID)
1827 sir_param(m, fn, "uv" as *u8, T_V2F)
1828 let c: i64 = sir_ctor(m, T_V4F)
1829 sir_arg(m, c, sir_lit(m, "1.0" as *u8, T_F32))
1830 sir_arg(m, c, sir_lit(m, "0.0" as *u8, T_F32))
1831 sir_arg(m, c, sir_lit(m, "0.0" as *u8, T_F32))
1832 sir_arg(m, c, sir_lit(m, "1.0" as *u8, T_F32))
1833 sir_stmt(m, fn, sir_return(m, c))
1834 return fn
1835}
1836
1837// S5 COVERAGE -- a fragment that DISCARDS and returns nothing. It declares NO output, and neither backend
1838// requires one: the output rule binds to what the body does, not to the stage. This is the POSITIVE control
1839// for that precision -- without it, "a fragment must declare an output" would be a guard that refuses a
1840// legal program, and a deny-guard whose false-positive rate nobody measured is worse than none.
1841// fragment stage: discard
1842func shsrc_frag_discard(m: *i64) -> i64 {
1843 let fn: i64 = sir_func(m, "main" as *u8, T_VOID)
1844 sir_stmt(m, fn, sir_discard(m))
1845 return fn
1846}
1847
1848// S5 NEG-CONTROL -- a fragment that RETURNS A VALUE with no declared output. This is the exact source that
1849// used to emit `void main(){ return vec4(1.0,0.0,0.0,1.0); }` -- invalid GLSL ES 3.00 twice over, shipped
1850// with no refusal because nothing pinned the GLSL fragment bytes. Both backends now refuse it by name.
1851// fragment stage: return vec4(1.0, 0.0, 0.0, 1.0) -- no output declared
1852func shsrc_frag_noout(m: *i64) -> i64 {
1853 let fn: i64 = sir_func(m, "main" as *u8, T_VOID)
1854 let c: i64 = sir_ctor(m, T_V4F)
1855 sir_arg(m, c, sir_lit(m, "1.0" as *u8, T_F32))
1856 sir_arg(m, c, sir_lit(m, "0.0" as *u8, T_F32))
1857 sir_arg(m, c, sir_lit(m, "0.0" as *u8, T_F32))
1858 sir_arg(m, c, sir_lit(m, "1.0" as *u8, T_F32))
1859 sir_stmt(m, fn, sir_return(m, c))
1860 return fn
1861}
1862
1863// S5 COVERAGE -- the output at a NON-ZERO location, proving the @location and the layout() qualifier are
1864// DATA and not the constant 0 they used to be. GLSL spells it `layout(location=2) out vec4 fc;` and WGSL
1865// `->@location(2) vec4f`; at location 0 both collapse to the bare form the shipping FSH already uses.
1866func shsrc_frag_loc2(m: *i64) -> i64 {
1867 sir_fragout(m, "fc" as *u8, T_V4F, 2)
1868 let fn: i64 = sir_func(m, "main" as *u8, T_VOID)
1869 let c: i64 = sir_ctor(m, T_V4F)
1870 sir_arg(m, c, sir_lit(m, "1.0" as *u8, T_F32))
1871 sir_arg(m, c, sir_lit(m, "0.0" as *u8, T_F32))
1872 sir_arg(m, c, sir_lit(m, "0.0" as *u8, T_F32))
1873 sir_arg(m, c, sir_lit(m, "1.0" as *u8, T_F32))
1874 sir_stmt(m, fn, sir_return(m, c))
1875 return fn
1876}
1877
1878// S5 NEG-CONTROL -- a fragment output declared on a VERTEX stage. Only the fragment stage has a colour
1879// output; a vertex stage that declared one would emit a variable nothing can ever write.
1880func shsrc_vert_fragout(m: *i64) -> i64 {
1881 sir_fragout(m, "fc" as *u8, T_V4F, 0)
1882 return shsrc_fullscreen_tri(m)
1883}
1884
1885// S5 NEG-CONTROL -- TWO fragment outputs. Multiple render targets are a real feature and their own rung;
1886// emitting the first and dropping the second would be the silent mistranslation this backend forbids.
1887func shsrc_frag_twoout(m: *i64) -> i64 {
1888 sir_fragout(m, "fc" as *u8, T_V4F, 0)
1889 sir_fragout(m, "fc2" as *u8, T_V4F, 1)
1890 return shsrc_frag_const_body(m)
1891}
1892
1893// the shared body the two-output control needs (shsrc_frag_const declares an output of its own)
1894func shsrc_frag_const_body(m: *i64) -> i64 {
1895 let fn: i64 = sir_func(m, "main" as *u8, T_VOID)
1896 let c: i64 = sir_ctor(m, T_V4F)
1897 sir_arg(m, c, sir_lit(m, "1.0" as *u8, T_F32))
1898 sir_arg(m, c, sir_lit(m, "0.0" as *u8, T_F32))
1899 sir_arg(m, c, sir_lit(m, "0.0" as *u8, T_F32))
1900 sir_arg(m, c, sir_lit(m, "1.0" as *u8, T_F32))
1901 sir_stmt(m, fn, sir_return(m, c))
1902 return fn
1903}
1904
1905// S5 NEG-CONTROL -- a SAMPLED texture read. wg_expr used to emit textureSample(t,samp,c) against an
1906// undeclared `samp`; the declaration pass had refused the same missing binding for two weeks. GLSL has no
1907// such gap (its texture() needs no separate sampler object), so this is a one-sided refusal BY DESIGN --
1908// the mirror image of K_STORAGE, which GLSL refuses and WGSL emits.
1909// texture vox: texture_3d<u32> ; func s(uv: vec2) -> vec4 { return textureSample(vox, uv) }
1910func shsrc_texsample(m: *i64) -> i64 {
1911 sir_texture(m, "vox" as *u8, T_TEX3U)
1912 let fn: i64 = sir_func(m, "s" as *u8, T_V4F)
1913 sir_param(m, fn, "uv" as *u8, T_V2F)
1914 let sm: i64 = sir_texsample(m, sir_ident(m, "vox" as *u8, T_TEX3U), sir_ident(m, "uv" as *u8, T_V2F))
1915 sir_stmt(m, fn, sir_return(m, sm))
1916 return fn
1917}
1918
1919// R-A / R-C fixture (2026-09-04): one plain function using a select and an immutable local. Expected bytes:
1920// GLSL float t=((c>0)?a:b); WGSL let t:f32=select(b,a,(c>0));
1921func shsrc_select_let(m: *i64) -> i64 {
1922 let f: i64 = sir_func(m, "pick" as *u8, T_F32)
1923 sir_param(m, f, "a" as *u8, T_F32)
1924 sir_param(m, f, "b" as *u8, T_F32)
1925 sir_param(m, f, "c" as *u8, T_I32)
1926 let cond: i64 = sir_bin(m, ">" as *u8, sir_ident(m, "c" as *u8, T_I32), sir_lit(m, "0" as *u8, T_I32), T_BOOL)
1927 let sel: i64 = sir_select(m, cond, sir_ident(m, "a" as *u8, T_F32), sir_ident(m, "b" as *u8, T_F32), T_F32)
1928 sir_stmt(m, f, sir_let(m, "t" as *u8, T_F32, sel))
1929 sir_stmt(m, f, sir_return(m, sir_ident(m, "t" as *u8, T_F32)))
1930 return f
1931}
1932
1933// R-B fixture (2026-09-04): a counted loop written ONCE as builder sugar (sir_for_until) -- the counter is a
1934// mutable local, the exit test and the step ride the shared S_IF/S_LOOP/S_ASSIGN arms, so the two dialects
1935// differ only in the loop keyword and the declaration spelling. Expected bytes are asserted by the gate.
1936func shsrc_for_until(m: *i64) -> i64 {
1937 let f: i64 = sir_func(m, "sum" as *u8, T_I32)
1938 sir_param(m, f, "n" as *u8, T_I32)
1939 sir_stmt(m, f, sir_var(m, "s" as *u8, T_I32, sir_lit(m, "0" as *u8, T_I32)))
1940 let init: i64 = sir_var(m, "i" as *u8, T_I32, sir_lit(m, "0" as *u8, T_I32))
1941 let until: i64 = sir_bin(m, ">=" as *u8, sir_ident(m, "i" as *u8, T_I32), sir_ident(m, "n" as *u8, T_I32), T_BOOL)
1942 let step: i64 = sir_assign(m, sir_ident(m, "i" as *u8, T_I32), sir_bin(m, "+" as *u8, sir_ident(m, "i" as *u8, T_I32), sir_lit(m, "1" as *u8, T_I32), T_I32))
1943 let body: i64 = sir_assign(m, sir_ident(m, "s" as *u8, T_I32), sir_bin(m, "+" as *u8, sir_ident(m, "s" as *u8, T_I32), sir_ident(m, "i" as *u8, T_I32), T_I32))
1944 sir_for_until(m, f, init, until, step, body)
1945 sir_stmt(m, f, sir_return(m, sir_ident(m, "s" as *u8, T_I32)))
1946 return f
1947}
1948
1949// R-D fixture (2026-09-04): one struct, a struct-typed local built by the struct's own constructor (an E_CALL
1950// named after the struct), a member write, and two member reads -- one through a nested swizzle. Expected bytes
1951// are asserted by the gate; the two dialects differ only in the declaration spelling and the local's keyword.
1952func shsrc_struct(m: *i64) -> i64 {
1953 let st: i64 = sir_struct(m, "Hit" as *u8)
1954 sir_member(m, st, "pos" as *u8, T_V3F)
1955 sir_member(m, st, "d" as *u8, T_F32)
1956 let hty: i64 = sir_struct_ty(st)
1957 let f: i64 = sir_func(m, "hitd" as *u8, T_F32)
1958 sir_param(m, f, "p" as *u8, T_V3F)
1959 sir_param(m, f, "k" as *u8, T_F32)
1960 let ctor: i64 = sir_call(m, "Hit" as *u8, hty)
1961 sir_arg(m, ctor, sir_ident(m, "p" as *u8, T_V3F))
1962 sir_arg(m, ctor, sir_ident(m, "k" as *u8, T_F32))
1963 sir_stmt(m, f, sir_var(m, "h" as *u8, hty, ctor))
1964 let hd: i64 = sir_member_of(m, sir_ident(m, "h" as *u8, hty), "d" as *u8, T_F32)
1965 sir_stmt(m, f, sir_assign(m, hd, sir_bin(m, "*" as *u8, sir_member_of(m, sir_ident(m, "h" as *u8, hty), "d" as *u8, T_F32), sir_lit(m, "2.0" as *u8, T_F32), T_F32)))
1966 let px: i64 = sir_swz(m, sir_member_of(m, sir_ident(m, "h" as *u8, hty), "pos" as *u8, T_V3F), "x" as *u8, T_F32)
1967 sir_stmt(m, f, sir_return(m, sir_bin(m, "+" as *u8, sir_member_of(m, sir_ident(m, "h" as *u8, hty), "d" as *u8, T_F32), px, T_F32)))
1968 return f
1969}
1970
1971// R-E fixture (2026-09-04): ONE function shape for the struct-returning march the two hand shaders spell two ways
1972// (GLSL out-params vs WGSL struct return). GLSL ES 3.00 returns structs, so the canonical shape is the struct
1973// return in BOTH dialects: mk builds and returns a Hit, useh takes a Hit PARAMETER and reads a member off the
1974// call's result directly. Expected bytes are asserted by the gate.
1975func shsrc_struct_ret(m: *i64) -> i64 {
1976 let st: i64 = sir_struct(m, "Hit" as *u8)
1977 sir_member(m, st, "pos" as *u8, T_V3F)
1978 sir_member(m, st, "d" as *u8, T_F32)
1979 let hty: i64 = sir_struct_ty(st)
1980 let f1: i64 = sir_func(m, "mk" as *u8, hty)
1981 sir_param(m, f1, "p" as *u8, T_V3F)
1982 sir_param(m, f1, "k" as *u8, T_F32)
1983 let ctor: i64 = sir_call(m, "Hit" as *u8, hty)
1984 sir_arg(m, ctor, sir_ident(m, "p" as *u8, T_V3F))
1985 sir_arg(m, ctor, sir_ident(m, "k" as *u8, T_F32))
1986 sir_stmt(m, f1, sir_var(m, "h" as *u8, hty, ctor))
1987 sir_stmt(m, f1, sir_assign(m, sir_member_of(m, sir_ident(m, "h" as *u8, hty), "d" as *u8, T_F32), sir_bin(m, "*" as *u8, sir_ident(m, "k" as *u8, T_F32), sir_lit(m, "2.0" as *u8, T_F32), T_F32)))
1988 sir_stmt(m, f1, sir_return(m, sir_ident(m, "h" as *u8, hty)))
1989 let f2: i64 = sir_func(m, "useh" as *u8, T_F32)
1990 sir_param(m, f2, "h" as *u8, hty)
1991 sir_param(m, f2, "k" as *u8, T_F32)
1992 let call: i64 = sir_call(m, "mk" as *u8, hty)
1993 sir_arg(m, call, sir_member_of(m, sir_ident(m, "h" as *u8, hty), "pos" as *u8, T_V3F))
1994 sir_arg(m, call, sir_ident(m, "k" as *u8, T_F32))
1995 sir_stmt(m, f2, sir_return(m, sir_bin(m, "+" as *u8, sir_member_of(m, call, "d" as *u8, T_F32), sir_member_of(m, sir_ident(m, "h" as *u8, hty), "d" as *u8, T_F32), T_F32)))
1996 return f2
1997}
1998
1999// ---- R-F (2026-09-04, GE44): THE UNIFORM LAYOUT IS DERIVED, NEVER HAND-COUNTED --------------------------
2000// The shipping page packs struct U by hand-counted float indices beside a hand-written declaration: two copies
2001// of one shape. WGSL defines the layout mechanically (spec, memory layouts, mirrored as gameengine.refs
2002// ge-wgsl-spec): AlignOf/SizeOf per type, array stride = roundUp(AlignOf(E), SizeOf(E)), struct size rounded up
2003// to its alignment, and in the uniform address space an array stride that is a multiple of 16 -- an array<f32,N>
2004// member is INVALID there, so it is refused by name here rather than laid out into a buffer the GPU rejects.
2005// `nx_wgsl layout` prints one row per uniform in declaration order plus the struct size; nx_wgsl_gate holds the
2006// shipping packer's own numbers (pal at 112, palf at 304, 96 floats = 384 bytes) as the KAT.
2007const WG_UNIFORM_STRIDE_MULTIPLE: i64 = 16
2008func wg_round_up(al: i64, n: i64) -> i64 { return ((n + al - 1) / al) * al }
2009func wg_align_of(t: i64) -> i64 {
2010 if t == T_F32 { return 4 }
2011 if t == T_I32 { return 4 }
2012 if t == T_U32 { return 4 }
2013 if t == T_BOOL { return 4 }
2014 if t == T_V2F { return 8 }
2015 if t == T_V3F { return 16 }
2016 if t == T_V4F { return 16 }
2017 if t == T_V3I { return 16 }
2018 if t == T_V3U { return 16 }
2019 return 0
2020}
2021func wg_size_of(t: i64) -> i64 {
2022 if t == T_F32 { return 4 }
2023 if t == T_I32 { return 4 }
2024 if t == T_U32 { return 4 }
2025 if t == T_BOOL { return 4 }
2026 if t == T_V2F { return 8 }
2027 if t == T_V3F { return 12 }
2028 if t == T_V4F { return 16 }
2029 if t == T_V3I { return 12 }
2030 if t == T_V3U { return 12 }
2031 return 0
2032}
2033// Writes `uniform|<name>|offset=|size=|align=|stride=|count=` rows for every K_UNIFORM in declaration order and a
2034// final `struct_u_size=` row; returns the struct size, or -1 with the refusal named in the module.
2035// The summary key is a PARAMETER (GE55, 2026-09-05): a capture that lays out several modules must give each its own
2036// key, or a reader anchored on `struct_u_size=` takes whichever module printed last -- measured on the live gate: the
2037// world block's 384 read as a fixture's 4 the moment a second layout shared the key. One writer per key.
2038func wg_uniform_layout(m: *i64, out: *u8, pos: i64, cap: i64) -> i64 { return wg_uniform_layout_k(m, out, pos, cap, "struct_u_size=" as *u8) }
2039func wg_uniform_layout_k(m: *i64, out: *u8, pos: i64, cap: i64, key: *u8) -> i64 {
2040 var p: i64 = pos
2041 var off: i64 = 0
2042 var maxal: i64 = 0
2043 var d: i64 = m[M_DECLH]
2044 while d != 0 {
2045 if sir_get(m, d, SIR_KIND) == K_UNIFORM {
2046 var ty: i64 = sir_get(m, d, SIR_TY)
2047 var n: i64 = sir_get(m, d, SIR_A)
2048 // S12c-2b: float[N] is laid out as the array<vec4f,ceil(N/4)> the struct emitter declares -- ONE rule
2049 // (wg_f32arr_vec4n), so the ruler and the emitter cannot disagree about the same declaration; the row
2050 // carries `f32n=` so a reader can see the lowering happened and from what.
2051 var f32n: i64 = 0
2052 if ty == T_F32 { if n > 0 { f32n = n; ty = T_V4F; n = wg_f32arr_vec4n(n) } }
2053 let al: i64 = wg_align_of(ty)
2054 let sz: i64 = wg_size_of(ty)
2055 if al == 0 { sir_refuse(m, "wgsl: uniform of a type with no layout rule in this backend (a texture is a bound resource, a struct-typed uniform is its own rung) -- refused rather than laid out by guess" as *u8); return 0 - 1 }
2056 var stride: i64 = 0
2057 var total: i64 = sz
2058 if n > 0 {
2059 stride = wg_round_up(al, sz)
2060 if (stride - (stride / WG_UNIFORM_STRIDE_MULTIPLE) * WG_UNIFORM_STRIDE_MULTIPLE) != 0 { sir_refuse(m, "wgsl: uniform array element stride is not a multiple of 16 -- the uniform address space requires it (spec: an array<f32,N> member of a uniform struct is invalid); use a vec4f element as the shipping pal/palf do. Refused rather than laid out into a buffer the GPU rejects" as *u8); return 0 - 1 }
2061 total = stride * n
2062 }
2063 off = wg_round_up(al, off)
2064 p = eb_put(out, p, cap, "uniform|" as *u8)
2065 p = eb_put(out, p, cap, sir_cstr(m, sir_get(m, d, SIR_NAME)))
2066 p = eb_put(out, p, cap, "|offset=" as *u8)
2067 p = eb_num(out, p, cap, off)
2068 p = eb_put(out, p, cap, "|size=" as *u8)
2069 p = eb_num(out, p, cap, total)
2070 p = eb_put(out, p, cap, "|align=" as *u8)
2071 p = eb_num(out, p, cap, al)
2072 p = eb_put(out, p, cap, "|stride=" as *u8)
2073 p = eb_num(out, p, cap, stride)
2074 p = eb_put(out, p, cap, "|count=" as *u8)
2075 p = eb_num(out, p, cap, n)
2076 if f32n > 0 { p = eb_put(out, p, cap, "|f32n=" as *u8); p = eb_num(out, p, cap, f32n) }
2077 p = eb_put(out, p, cap, "\n" as *u8)
2078 off = off + total
2079 if al > maxal { maxal = al }
2080 }
2081 d = sir_get(m, d, SIR_NEXT)
2082 }
2083 var size: i64 = 0
2084 if maxal > 0 { size = wg_round_up(maxal, off) }
2085 p = eb_put(out, p, cap, key)
2086 p = eb_num(out, p, cap, size)
2087 p = eb_put(out, p, cap, "\n" as *u8)
2088 if p + 1 < cap { out[p] = 0 as u8 }
2089 return size
2090}
2091// R-F fixture: the shipping world pass's uniform block, in its declaration order, plus one uniform read.
2092func shsrc_world_uniforms(m: *i64) -> i64 {
2093 // R-G: the uniform block has ONE owner -- the world shader source (ws_world_uniforms); this fixture reads it.
2094 ws_world_uniforms(m)
2095 let f: i64 = sir_func(m, "tk" as *u8, T_F32)
2096 sir_param(m, f, "k" as *u8, T_F32)
2097 sir_stmt(m, f, sir_return(m, sir_bin(m, "*" as *u8, sir_ident(m, "t" as *u8, T_F32), sir_ident(m, "k" as *u8, T_F32), T_F32)))
2098 return f
2099}
2100// `nx_wgsl layout`: the world uniform block through both backends (the emitted struct U is the WGSL spelling the
2101// hand shader carries), then the derived layout rows, then the stride neg-control on an array<f32,8> member.
2102// R-G0 fixture (2026-09-04): an uninitialised local and a unary negation -- `var v:f32; v=(-a); return v;` in both
2103// spellings. Expected bytes are asserted by the gate.
2104func shsrc_neg_uninit(m: *i64) -> i64 {
2105 let f: i64 = sir_func(m, "nu" as *u8, T_F32)
2106 sir_param(m, f, "a" as *u8, T_F32)
2107 sir_stmt(m, f, sir_var(m, "v" as *u8, T_F32, 0))
2108 sir_stmt(m, f, sir_assign(m, sir_ident(m, "v" as *u8, T_F32), sir_neg(m, sir_ident(m, "a" as *u8, T_F32), T_F32)))
2109 sir_stmt(m, f, sir_return(m, sir_ident(m, "v" as *u8, T_F32)))
2110 return f
2111}
2112// `nx_wgsl world` (R-G): THE WORLD SHADER (nx_world_shader_src) through both backends -- the fragment module
2113// (uniforms, the voxel texture, struct Hit, every function, the entry) and the backend's own fullscreen-triangle
2114// vertex stage -- printed the way every verb prints, for nx_game_page_emit to capture and for the gate to pin.
2115func wg_world_demo() -> i64 {
2116 wg_puts("=== NX-WGSL world: the world pass written ONCE, emitted to both dialects ===\n" as *u8)
2117 let tr: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2118 let tn: *i64 = sys_mmap(WG_NUMBUF) as *i64
2119 let nb: *u8 = sys_mmap(WG_NUMBUF)
2120 var q: i64 = 0
2121 let m1: *i64 = sir_new()
2122 shsrc_world(m1)
2123 let gt1: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2124 let gn1: *i64 = sys_mmap(WG_NUMBUF) as *i64
2125 wg_probe(m1, STAGE_FRAGMENT, "world_glsl" as *u8, 0, gt1, gn1)
2126 wg_probe(m1, STAGE_FRAGMENT, "world_wgsl" as *u8, 1, tr, tn)
2127 wg_puts("world_trace_equal=" as *u8)
2128 q = eb_num(nb, 0, WG_NUMBUF, wg_trace_eq(gt1, gn1[0], tr, tn[0])); nb[q] = 0 as u8; wg_puts(nb)
2129 wg_puts("\n" as *u8)
2130 let m2: *i64 = sir_new()
2131 shsrc_fullscreen_tri(m2)
2132 wg_probe(m2, STAGE_VERTEX, "worldvs_glsl" as *u8, 0, gt1, gn1)
2133 wg_probe(m2, STAGE_VERTEX, "worldvs_wgsl" as *u8, 1, tr, tn)
2134 return 0
2135}
2136// S12c (2026-09-05): `nx_wgsl cast [<glsl-out>]` -- THE CHARACTER CAST vertex stage (nx_cast_shader_src, transcribed slice by
2137// slice from the hand MVS) through both backends, printed the way every verb prints; when a path is given the emitted GLSL is
2138// ALSO written there VERBATIM, so nx_glsl_tokdiff can measure it against buildroot/knowledge/compare/cast_mvs_hand.glsl (the
2139// pre-declared accept rule: canonical-token prefix per slice, IDENTICAL for the whole stage). The WGSL probe is expected to
2140// REFUSE BY NAME until S12c-2b lowers the hand's float[12] uniform arrays (uniform-space stride) -- printed, never hidden.
2141const WG_MODE_0644: i64 = 420
2142// emit the GLSL of module m for `stage` and write it VERBATIM to outp, printing `<label>_written=<bytes> fd_ok= path=`
2143func wg_glsl_to_path(m: *i64, stage: i64, outp: *u8, label: *u8) -> i64 {
2144 let nb: *u8 = sys_mmap(WG_NUMBUF)
2145 var q: i64 = 0
2146 let buf: *u8 = sys_mmap(WG_EMIT_CAP)
2147 let t2: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2148 let n2: *i64 = sys_mmap(WG_NUMBUF) as *i64
2149 n2[0] = 0
2150 let rc: i64 = glsl_emit_module(m, stage, buf, WG_EMIT_CAP, t2, WG_TRACE_CAP, n2)
2151 wg_puts(label)
2152 if rc > 0 {
2153 let fd: i64 = sys_openat_wr(outp, WG_MODE_0644)
2154 if fd >= 0 { sys_write(fd, buf, rc); sys_close(fd) }
2155 wg_puts("_written=" as *u8); q = eb_num(nb, 0, WG_NUMBUF, rc); nb[q] = 0 as u8; wg_puts(nb)
2156 wg_puts(" fd_ok=" as *u8); if fd >= 0 { wg_puts("1" as *u8) } else { wg_puts("0" as *u8) }
2157 wg_puts(" path=" as *u8); wg_puts(outp); wg_puts("\n" as *u8)
2158 } else { wg_puts("_written=0 REFUSED\n" as *u8) }
2159 return rc
2160}
2161func wg_cast_demo(outp: *u8, outf: *u8) -> i64 {
2162 wg_puts("=== NX-WGSL cast: the character cast vertex stage written ONCE, emitted to both dialects ===\n" as *u8)
2163 let tr: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2164 let tn: *i64 = sys_mmap(WG_NUMBUF) as *i64
2165 let nb: *u8 = sys_mmap(WG_NUMBUF)
2166 var q: i64 = 0
2167 let m1: *i64 = sir_new()
2168 shsrc_cast_vs(m1)
2169 let gt1: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2170 let gn1: *i64 = sys_mmap(WG_NUMBUF) as *i64
2171 wg_probe(m1, STAGE_VERTEX, "cast_glsl" as *u8, 0, gt1, gn1)
2172 let m1w: *i64 = sir_new()
2173 shsrc_cast_vs(m1w)
2174 wg_probe(m1w, STAGE_VERTEX, "cast_wgsl" as *u8, 1, tr, tn)
2175 wg_puts("cast_trace_equal=" as *u8)
2176 q = eb_num(nb, 0, WG_NUMBUF, wg_trace_eq(gt1, gn1[0], tr, tn[0])); nb[q] = 0 as u8; wg_puts(nb)
2177 wg_puts(" cast_nodes=" as *u8)
2178 q = eb_num(nb, 0, WG_NUMBUF, m1[M_NCOUNT]); nb[q] = 0 as u8; wg_puts(nb)
2179 wg_puts("\n" as *u8)
2180 // S12c-6: the node arena overflows SILENTLY (sir_node returns 0, sets M_OVER) -- printed so a truncated module cannot pass as a small one
2181 wg_puts("cast_over=" as *u8)
2182 q = eb_num(nb, 0, WG_NUMBUF, m1[M_OVER]); nb[q] = 0 as u8; wg_puts(nb)
2183 wg_puts("\n" as *u8)
2184 if (outp as i64) != 0 {
2185 let m2: *i64 = sir_new()
2186 shsrc_cast_vs(m2)
2187 wg_glsl_to_path(m2, STAGE_VERTEX, outp, "cast_glsl" as *u8)
2188 }
2189 // S12c-6 (2026-09-05): THE FRAGMENT STAGE (shsrc_cast_fs) through both backends -- its own probe pair, trace line and file,
2190 // measured by nx_glsl_tokdiff against buildroot/knowledge/compare/cast_mfs_hand.glsl exactly as the vertex stage was.
2191 let m3: *i64 = sir_new()
2192 shsrc_cast_fs(m3)
2193 let gt3: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2194 let gn3: *i64 = sys_mmap(WG_NUMBUF) as *i64
2195 wg_probe(m3, STAGE_FRAGMENT, "castfs_glsl" as *u8, 0, gt3, gn3)
2196 let m3w: *i64 = sir_new()
2197 shsrc_cast_fs(m3w)
2198 // S12 step 2: the fragment stage's resources live in bind group 1 (see M_BINDGROUP) -- the vertex stage keeps group 0
2199 m3w[M_BINDGROUP] = 1
2200 wg_probe(m3w, STAGE_FRAGMENT, "castfs_wgsl" as *u8, 1, tr, tn)
2201 wg_puts("castfs_trace_equal=" as *u8)
2202 q = eb_num(nb, 0, WG_NUMBUF, wg_trace_eq(gt3, gn3[0], tr, tn[0])); nb[q] = 0 as u8; wg_puts(nb)
2203 wg_puts(" castfs_nodes=" as *u8)
2204 q = eb_num(nb, 0, WG_NUMBUF, m3[M_NCOUNT]); nb[q] = 0 as u8; wg_puts(nb)
2205 wg_puts(" castfs_over=" as *u8)
2206 q = eb_num(nb, 0, WG_NUMBUF, m3[M_OVER]); nb[q] = 0 as u8; wg_puts(nb)
2207 wg_puts("\n" as *u8)
2208 if (outf as i64) != 0 {
2209 let m4: *i64 = sir_new()
2210 shsrc_cast_fs(m4)
2211 wg_glsl_to_path(m4, STAGE_FRAGMENT, outf, "castfs_glsl" as *u8)
2212 }
2213 // S12 step 2 (2026-09-06): the cast structs' DERIVED byte layouts, one key per module (one writer per key -- the world
2214 // block measured what a shared key does), so the page packs each stage's u:U at the ruler's offsets, never a hand count.
2215 let lbv: *u8 = sys_mmap(WG_EMIT_CAP)
2216 let lsv: i64 = wg_uniform_layout_k(m1w, lbv, 0, WG_EMIT_CAP, "castv_struct_size=" as *u8)
2217 wg_puts("castv_layout_rc=" as *u8)
2218 q = eb_num(nb, 0, WG_NUMBUF, lsv); nb[q] = 0 as u8; wg_puts(nb)
2219 wg_puts(" rows:\n" as *u8)
2220 wg_puts(lbv)
2221 let lbf: *u8 = sys_mmap(WG_EMIT_CAP)
2222 let lsf: i64 = wg_uniform_layout_k(m3w, lbf, 0, WG_EMIT_CAP, "castf_struct_size=" as *u8)
2223 wg_puts("castf_layout_rc=" as *u8)
2224 q = eb_num(nb, 0, WG_NUMBUF, lsf); nb[q] = 0 as u8; wg_puts(nb)
2225 wg_puts(" rows:\n" as *u8)
2226 wg_puts(lbf)
2227 wg_puts("castf_bind_group=" as *u8)
2228 q = eb_num(nb, 0, WG_NUMBUF, m3w[M_BINDGROUP]); nb[q] = 0 as u8; wg_puts(nb)
2229 wg_puts("\n" as *u8)
2230 return 0
2231}
2232func wg_layout_demo() -> i64 {
2233 wg_puts("=== NX-WGSL layout: the uniform block's memory layout is DERIVED from the IR by the WGSL rules ===\n" as *u8)
2234 let tr: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2235 let tn: *i64 = sys_mmap(WG_NUMBUF) as *i64
2236 let nb: *u8 = sys_mmap(WG_NUMBUF)
2237 var q: i64 = 0
2238 let m1: *i64 = sir_new()
2239 shsrc_world_uniforms(m1)
2240 wg_probe(m1, STAGE_PLAIN, "worldu_glsl" as *u8, 0, tr, tn)
2241 wg_probe(m1, STAGE_PLAIN, "worldu_wgsl" as *u8, 1, tr, tn)
2242 let lb: *u8 = sys_mmap(WG_EMIT_CAP)
2243 let ls: i64 = wg_uniform_layout(m1, lb, 0, WG_EMIT_CAP)
2244 wg_puts(lb)
2245 wg_puts("layout_rc=" as *u8)
2246 q = eb_num(nb, 0, WG_NUMBUF, ls); nb[q] = 0 as u8; wg_puts(nb)
2247 wg_puts(" layout_named=" as *u8)
2248 if sir_refused(m1) != 0 { wg_puts(sir_cstr(m1, sir_refused(m1))) } else { wg_puts("NONE" as *u8) }
2249 wg_puts("\n" as *u8)
2250 // GE55 K_PRIVATE is INVISIBLE to the uniform layout: a module with one uniform beside private state lays out the
2251 // uniform alone (struct size 4) and never lists the private name.
2252 let m3: *i64 = sir_new()
2253 shsrc_private_demo(m3)
2254 let lb3: *u8 = sys_mmap(WG_EMIT_CAP)
2255 let ls3: i64 = wg_uniform_layout_k(m3, lb3, 0, WG_EMIT_CAP, "priv_struct_size=" as *u8)
2256 wg_puts("layoutpriv_rc=" as *u8)
2257 q = eb_num(nb, 0, WG_NUMBUF, ls3); nb[q] = 0 as u8; wg_puts(nb)
2258 wg_puts(" rows:\n" as *u8)
2259 wg_puts(lb3)
2260 // GE55 the RAW world block: the one ruler both the sim's UR_* word table and the page's NXU_SIZE are pinned to.
2261 // 40 four-byte scalars, so every offset is index*4 and the struct is exactly 160 B; the gate reads these rows.
2262 let m4: *i64 = sir_new()
2263 ws_world_uniforms_raw(m4)
2264 let lb4: *u8 = sys_mmap(WG_EMIT_CAP)
2265 let ls4: i64 = wg_uniform_layout_k(m4, lb4, 0, WG_EMIT_CAP, "raw_struct_size=" as *u8)
2266 wg_puts("layoutraw_rc=" as *u8)
2267 q = eb_num(nb, 0, WG_NUMBUF, ls4); nb[q] = 0 as u8; wg_puts(nb)
2268 wg_puts(" rows:\n" as *u8)
2269 wg_puts(lb4)
2270 // S12c-2b POSITIVE CONTROL: an array<f32,8> uniform member is LOWERED to array<vec4f,2> -- stride 16, size 32, f32n=8 --
2271 // by the same rule the struct emitter applies (wg_f32arr_vec4n), so the layout ruler and the emitted struct U agree.
2272 let m5: *i64 = sir_new()
2273 sir_uniform(m5, "w" as *u8, T_F32, 8)
2274 let lb5: *u8 = sys_mmap(WG_EMIT_CAP)
2275 let ls5: i64 = wg_uniform_layout_k(m5, lb5, 0, WG_EMIT_CAP, "f32_struct_size=" as *u8)
2276 wg_puts("layoutf32_rc=" as *u8)
2277 q = eb_num(nb, 0, WG_NUMBUF, ls5); nb[q] = 0 as u8; wg_puts(nb)
2278 wg_puts(" rows:\n" as *u8)
2279 wg_puts(lb5)
2280 // NEG-CONTROL: an array<vec2f,8> uniform member has stride 8, invalid in the uniform address space and with no lowering:
2281 // refused by name (this control moved off array<f32,8> when S12c-2b gave that case a lowering)
2282 let m2: *i64 = sir_new()
2283 sir_uniform(m2, "w" as *u8, T_V2F, 8)
2284 let lb2: *u8 = sys_mmap(WG_EMIT_CAP)
2285 let ls2: i64 = wg_uniform_layout(m2, lb2, 0, WG_EMIT_CAP)
2286 wg_puts("layoutbad_rc=" as *u8)
2287 q = eb_num(nb, 0, WG_NUMBUF, ls2); nb[q] = 0 as u8; wg_puts(nb)
2288 wg_puts(" layoutbad_named=" as *u8)
2289 if sir_refused(m2) != 0 { wg_puts(sir_cstr(m2, sir_refused(m2))) } else { wg_puts("NONE" as *u8) }
2290 wg_puts("\n" as *u8)
2291 return 0
2292}
2293
2294// S6 2026-09-05 closure bump: fixture attribute apos->aP (substring-collision fix); a prior /api/build
2295// served a stale cached artifact despite a changed src, so this comment forces a genuine recompile.
2296// S9 (2026-09-05): the ONE place a callee is re-spelled per dialect. IR call names are the WGSL spellings; GLSL ES 3.00
2297// spells the fragment derivatives dFdx/dFdy and inverseSqrt as inversesqrt (fwidth is the same word in both). Any other
2298// name passes through unchanged, and the trace carries the IR name, so trace identity across dialects is untouched.
2299// Stage refusal (derivatives are fragment-only in both dialects) is NOT decided here: the expression arms have no stage
2300// in scope; both compilers refuse a derivative outside a fragment entry at pipeline creation, which is the oracle.
2301func gl_call_name(name: *u8) -> *u8 {
2302 if wg_streq(name, "dpdx" as *u8) == 1 { return "dFdx" as *u8 }
2303 if wg_streq(name, "dpdy" as *u8) == 1 { return "dFdy" as *u8 }
2304 if wg_streq(name, "inverseSqrt" as *u8) == 1 { return "inversesqrt" as *u8 }
2305 return name
2306}
2307
2308// S8 provenance 2026-09-05 (seat fable): samplers as texture_2d+sampler pairs for the named cast set uAtl uNrm uJT uGD,
2309// texload typed by its texture, T_V2I lowered to vec2i/ivec2. This comment also bumps the build closure: the /api/build
2310// cache answered the changed source with the RELEASE twin e646b12c built by the sov lane, and the live binary is the
2311// DEBUG flavour, so a flavour-matched rebuild needs a key the cache has never seen.
2312func wg_decls_demo() -> i64 {
2313 wg_puts("=== NX-WGSL decls: declaration lowering is NAMED -- refusals, qualification, resources, entry params ===\n" as *u8)
2314 let tr: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2315 let tn: *i64 = sys_mmap(WG_NUMBUF) as *i64
2316 // NEG-CONTROL S1a -- a texture-typed K_UNIFORM must refuse, never land inside struct U
2317 let m1: *i64 = sir_new()
2318 sir_uniform(m1, "vox" as *u8, T_TEX3U, 0)
2319 shsrc_fullscreen_tri(m1)
2320 wg_probe(m1, STAGE_VERTEX, "texuniform" as *u8, 1, tr, tn)
2321 // S6 (2026-09-05) COVERAGE -- a vertex attribute lowers to a @location entry parameter in WGSL and a
2322 // module-scope `layout(location=N) in` in GLSL; emitted through BOTH backends, trace-identical (the
2323 // no-GPU equivalence oracle). It used to REFUSE (the S1b neg-control); the refusal path is now the
2324 // out-of-vertex-stage case, covered by the S6 neg-control below.
2325 let m2: *i64 = sir_new()
2326 sir_attrib(m2, "aP" as *u8, T_V3F, 0)
2327 shsrc_fullscreen_tri(m2)
2328 let gt2: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2329 let gn2: *i64 = sys_mmap(WG_NUMBUF) as *i64
2330 wg_probe(m2, STAGE_VERTEX, "attrib_glsl" as *u8, 0, gt2, gn2)
2331 wg_probe(m2, STAGE_VERTEX, "attrib" as *u8, 1, tr, tn)
2332 let nb2: *u8 = sys_mmap(WG_NUMBUF)
2333 wg_puts("attrib_trace_equal=" as *u8)
2334 var q2: i64 = eb_num(nb2, 0, WG_NUMBUF, wg_trace_eq(gt2, gn2[0], tr, tn[0])); nb2[q2] = 0 as u8; wg_puts(nb2)
2335 wg_puts("\n" as *u8)
2336 // S6 NEG-CONTROL -- a K_ATTRIB reached in a FRAGMENT stage must refuse BY NAME (attributes are vertex-only)
2337 let m2b: *i64 = sir_new()
2338 sir_attrib(m2b, "aP" as *u8, T_V3F, 0)
2339 shsrc_frag_const(m2b)
2340 wg_probe(m2b, STAGE_FRAGMENT, "attribfrag" as *u8, 1, tr, tn)
2341 // S7 (2026-09-05) COVERAGE -- varyings are ONE interstage struct in WGSL (returned by the vertex entry,
2342 // taken by the fragment entry, a flat int carried with @interpolate(flat)) and plain [flat ]out|in in
2343 // GLSL; BOTH dialects, BOTH stages, trace-identical. The old S1c neg-control (a fragment K_VARY must
2344 // refuse) is now the fragvary POSITIVE: declared-but-unread is admitted with the struct parameter in place.
2345 let m3: *i64 = sir_new()
2346 sir_vary(m3, "vcol" as *u8, T_V4F, 0, 0)
2347 shsrc_frag_const(m3)
2348 wg_probe(m3, STAGE_FRAGMENT, "fragvary" as *u8, 1, tr, tn)
2349 // vertex fixture: vcol = vec4(1,0,0,1); vii = 7 (flat i32); position = vec4(0,0,0,1)
2350 let m3v: *i64 = sir_new()
2351 sir_vary(m3v, "vcol" as *u8, T_V4F, 0, 0)
2352 sir_vary(m3v, "vii" as *u8, T_I32, 1, 1)
2353 let f3v: i64 = sir_func(m3v, "main" as *u8, T_VOID)
2354 let c3a: i64 = sir_ctor(m3v, T_V4F)
2355 sir_arg(m3v, c3a, sir_lit(m3v, "1.0" as *u8, T_F32))
2356 sir_arg(m3v, c3a, sir_lit(m3v, "0.0" as *u8, T_F32))
2357 sir_arg(m3v, c3a, sir_lit(m3v, "0.0" as *u8, T_F32))
2358 sir_arg(m3v, c3a, sir_lit(m3v, "1.0" as *u8, T_F32))
2359 sir_stmt(m3v, f3v, sir_assign(m3v, sir_ident(m3v, "vcol" as *u8, T_V4F), c3a))
2360 sir_stmt(m3v, f3v, sir_assign(m3v, sir_ident(m3v, "vii" as *u8, T_I32), sir_lit(m3v, "7" as *u8, T_I32)))
2361 let c3b: i64 = sir_ctor(m3v, T_V4F)
2362 sir_arg(m3v, c3b, sir_lit(m3v, "0.0" as *u8, T_F32))
2363 sir_arg(m3v, c3b, sir_lit(m3v, "0.0" as *u8, T_F32))
2364 sir_arg(m3v, c3b, sir_lit(m3v, "0.0" as *u8, T_F32))
2365 sir_arg(m3v, c3b, sir_lit(m3v, "1.0" as *u8, T_F32))
2366 sir_stmt(m3v, f3v, sir_assign(m3v, sir_builtin(m3v, "position" as *u8, T_V4F), c3b))
2367 let gt3v: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2368 let gn3v: *i64 = sys_mmap(WG_NUMBUF) as *i64
2369 wg_probe(m3v, STAGE_VERTEX, "varyvs_glsl" as *u8, 0, gt3v, gn3v)
2370 wg_probe(m3v, STAGE_VERTEX, "varyvs" as *u8, 1, tr, tn)
2371 let nb3: *u8 = sys_mmap(WG_NUMBUF)
2372 wg_puts("varyvs_trace_equal=" as *u8)
2373 var q3: i64 = eb_num(nb3, 0, WG_NUMBUF, wg_trace_eq(gt3v, gn3v[0], tr, tn[0])); nb3[q3] = 0 as u8; wg_puts(nb3)
2374 wg_puts("\n" as *u8)
2375 // fragment fixture: fc = vcol * float(vii) -- reads both varyings, one of them flat
2376 let m3f: *i64 = sir_new()
2377 sir_vary(m3f, "vcol" as *u8, T_V4F, 0, 0)
2378 sir_vary(m3f, "vii" as *u8, T_I32, 1, 1)
2379 sir_fragout(m3f, "fc" as *u8, T_V4F, 0)
2380 let f3f: i64 = sir_func(m3f, "main" as *u8, T_VOID)
2381 let e3: i64 = sir_bin(m3f, "*" as *u8, sir_ident(m3f, "vcol" as *u8, T_V4F), sir_cast(m3f, sir_ident(m3f, "vii" as *u8, T_I32), T_F32), T_V4F)
2382 sir_stmt(m3f, f3f, sir_return(m3f, e3))
2383 let gt3f: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2384 let gn3f: *i64 = sys_mmap(WG_NUMBUF) as *i64
2385 wg_probe(m3f, STAGE_FRAGMENT, "varyfs_glsl" as *u8, 0, gt3f, gn3f)
2386 wg_probe(m3f, STAGE_FRAGMENT, "varyfs" as *u8, 1, tr, tn)
2387 wg_puts("varyfs_trace_equal=" as *u8)
2388 q3 = eb_num(nb3, 0, WG_NUMBUF, wg_trace_eq(gt3f, gn3f[0], tr, tn[0])); nb3[q3] = 0 as u8; wg_puts(nb3)
2389 wg_puts("\n" as *u8)
2390 // S7 NEG-CONTROL -- a varying named bpos collides with the reserved position member and must refuse BY NAME
2391 let m3n: *i64 = sir_new()
2392 // (bpos is the reserved @builtin(position) member of struct VIO)
2393 sir_vary(m3n, "bpos" as *u8, T_V4F, 0, 0)
2394 shsrc_frag_const(m3n)
2395 wg_probe(m3n, STAGE_FRAGMENT, "varybpos" as *u8, 1, tr, tn)
2396 // S8 (2026-09-05) COVERAGE -- THE NAMED CAST SAMPLER SET uAtl uNrm uJT uGD: four texture_2d<f32> at DERIVED
2397 // two-slot bindings (1,3,5,7 + companion samplers 2,4,6,8); the joint table (uJT) and the garment table (uGD)
2398 // are read by INTEGER coordinate in the VERTEX stage (textureLoad / texelFetch over the new T_V2I, the load
2399 // typed vec4f by its texture), the atlas and normal map are SAMPLED in the fragment stage through the
2400 // derived companion samplers. Both dialects, both stages, trace-identical.
2401 let m8v: *i64 = sir_new()
2402 sir_texture(m8v, "uAtl" as *u8, T_TEX2F)
2403 sir_texture(m8v, "uNrm" as *u8, T_TEX2F)
2404 sir_texture(m8v, "uJT" as *u8, T_TEX2F)
2405 sir_texture(m8v, "uGD" as *u8, T_TEX2F)
2406 sir_attrib(m8v, "aJ" as *u8, T_V4F, 1)
2407 sir_vary(m8v, "vcol" as *u8, T_V4F, 0, 0)
2408 let f8v: i64 = sir_func(m8v, "main" as *u8, T_VOID)
2409 let c8j: i64 = sir_ctor(m8v, T_V2I)
2410 sir_arg(m8v, c8j, sir_cast(m8v, sir_swz(m8v, sir_ident(m8v, "aJ" as *u8, T_V4F), "x" as *u8, T_F32), T_I32))
2411 sir_arg(m8v, c8j, sir_lit(m8v, "0" as *u8, T_I32))
2412 sir_stmt(m8v, f8v, sir_assign(m8v, sir_ident(m8v, "vcol" as *u8, T_V4F), sir_texload(m8v, sir_ident(m8v, "uJT" as *u8, T_TEX2F), c8j, sir_lit(m8v, "0" as *u8, T_I32))))
2413 let c8g: i64 = sir_ctor(m8v, T_V2I)
2414 sir_arg(m8v, c8g, sir_lit(m8v, "0" as *u8, T_I32))
2415 sir_arg(m8v, c8g, sir_lit(m8v, "0" as *u8, T_I32))
2416 sir_stmt(m8v, f8v, sir_assign(m8v, sir_builtin(m8v, "position" as *u8, T_V4F), sir_texload(m8v, sir_ident(m8v, "uGD" as *u8, T_TEX2F), c8g, sir_lit(m8v, "0" as *u8, T_I32))))
2417 let gt8v: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2418 let gn8v: *i64 = sys_mmap(WG_NUMBUF) as *i64
2419 wg_probe(m8v, STAGE_VERTEX, "cast8vs_glsl" as *u8, 0, gt8v, gn8v)
2420 wg_probe(m8v, STAGE_VERTEX, "cast8vs" as *u8, 1, tr, tn)
2421 let nb8: *u8 = sys_mmap(WG_NUMBUF)
2422 wg_puts("cast8vs_trace_equal=" as *u8)
2423 var q8: i64 = eb_num(nb8, 0, WG_NUMBUF, wg_trace_eq(gt8v, gn8v[0], tr, tn[0])); nb8[q8] = 0 as u8; wg_puts(nb8)
2424 wg_puts("\n" as *u8)
2425 // S9 (2026-09-05) COVERAGE -- mat3 + derivatives + inverseSqrt. cast9vs: M = mat3(aA,aB,aC); position = vec4(M*aP, 1.0).
2426 // cast9fs: vec4(dpdx(v.x), dpdy(v.y), fwidth(v.z), inverseSqrt(v.w)) -- IR names are the WGSL spellings; the GLSL arm
2427 // re-spells dFdx/dFdy/inversesqrt through gl_call_name. Both stages, both dialects, trace-identical.
2428 let m9v: *i64 = sir_new()
2429 sir_attrib(m9v, "aP" as *u8, T_V3F, 0)
2430 sir_attrib(m9v, "aA" as *u8, T_V3F, 1)
2431 sir_attrib(m9v, "aB" as *u8, T_V3F, 2)
2432 sir_attrib(m9v, "aC" as *u8, T_V3F, 3)
2433 let f9v: i64 = sir_func(m9v, "main" as *u8, T_VOID)
2434 let c9m: i64 = sir_ctor(m9v, T_M3F)
2435 sir_arg(m9v, c9m, sir_ident(m9v, "aA" as *u8, T_V3F))
2436 sir_arg(m9v, c9m, sir_ident(m9v, "aB" as *u8, T_V3F))
2437 sir_arg(m9v, c9m, sir_ident(m9v, "aC" as *u8, T_V3F))
2438 sir_stmt(m9v, f9v, sir_let(m9v, "M" as *u8, T_M3F, c9m))
2439 let c9p: i64 = sir_ctor(m9v, T_V4F)
2440 sir_arg(m9v, c9p, sir_bin(m9v, "*" as *u8, sir_ident(m9v, "M" as *u8, T_M3F), sir_ident(m9v, "aP" as *u8, T_V3F), T_V3F))
2441 sir_arg(m9v, c9p, sir_lit(m9v, "1.0" as *u8, T_F32))
2442 sir_stmt(m9v, f9v, sir_assign(m9v, sir_builtin(m9v, "position" as *u8, T_V4F), c9p))
2443 let gt9v: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2444 let gn9v: *i64 = sys_mmap(WG_NUMBUF) as *i64
2445 wg_probe(m9v, STAGE_VERTEX, "cast9vs_glsl" as *u8, 0, gt9v, gn9v)
2446 wg_probe(m9v, STAGE_VERTEX, "cast9vs" as *u8, 1, tr, tn)
2447 let nb9c: *u8 = sys_mmap(WG_NUMBUF)
2448 wg_puts("cast9vs_trace_equal=" as *u8)
2449 var q9c: i64 = eb_num(nb9c, 0, WG_NUMBUF, wg_trace_eq(gt9v, gn9v[0], tr, tn[0])); nb9c[q9c] = 0 as u8; wg_puts(nb9c)
2450 wg_puts("\n" as *u8)
2451 let m9f: *i64 = sir_new()
2452 sir_vary(m9f, "vcol" as *u8, T_V4F, 0, 0)
2453 sir_fragout(m9f, "fc" as *u8, T_V4F, 0)
2454 let f9f: i64 = sir_func(m9f, "main" as *u8, T_VOID)
2455 let c9o: i64 = sir_ctor(m9f, T_V4F)
2456 let k9a: i64 = sir_call(m9f, "dpdx" as *u8, T_F32)
2457 sir_arg(m9f, k9a, sir_swz(m9f, sir_ident(m9f, "vcol" as *u8, T_V4F), "x" as *u8, T_F32))
2458 sir_arg(m9f, c9o, k9a)
2459 let k9b: i64 = sir_call(m9f, "dpdy" as *u8, T_F32)
2460 sir_arg(m9f, k9b, sir_swz(m9f, sir_ident(m9f, "vcol" as *u8, T_V4F), "y" as *u8, T_F32))
2461 sir_arg(m9f, c9o, k9b)
2462 let k9c: i64 = sir_call(m9f, "fwidth" as *u8, T_F32)
2463 sir_arg(m9f, k9c, sir_swz(m9f, sir_ident(m9f, "vcol" as *u8, T_V4F), "z" as *u8, T_F32))
2464 sir_arg(m9f, c9o, k9c)
2465 let k9d: i64 = sir_call(m9f, "inverseSqrt" as *u8, T_F32)
2466 sir_arg(m9f, k9d, sir_swz(m9f, sir_ident(m9f, "vcol" as *u8, T_V4F), "w" as *u8, T_F32))
2467 sir_arg(m9f, c9o, k9d)
2468 sir_stmt(m9f, f9f, sir_return(m9f, c9o))
2469 let gt9f: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2470 let gn9f: *i64 = sys_mmap(WG_NUMBUF) as *i64
2471 wg_probe(m9f, STAGE_FRAGMENT, "cast9fs_glsl" as *u8, 0, gt9f, gn9f)
2472 wg_probe(m9f, STAGE_FRAGMENT, "cast9fs" as *u8, 1, tr, tn)
2473 wg_puts("cast9fs_trace_equal=" as *u8)
2474 q9c = eb_num(nb9c, 0, WG_NUMBUF, wg_trace_eq(gt9f, gn9f[0], tr, tn[0])); nb9c[q9c] = 0 as u8; wg_puts(nb9c)
2475 wg_puts("\n" as *u8)
2476 // S12a (2026-09-05) COVERAGE -- the instance index: position = vec4(f32(instance_index), aP.y, 0, 1). WGSL binds
2477 // @builtin(instance_index) ii:u32 only in THIS module (M_USESII); GLSL reads gl_InstanceID undeclared. Trace-identical.
2478 let m10v: *i64 = sir_new()
2479 sir_attrib(m10v, "aP" as *u8, T_V3F, 0)
2480 let f10v: i64 = sir_func(m10v, "main" as *u8, T_VOID)
2481 let c10p: i64 = sir_ctor(m10v, T_V4F)
2482 sir_arg(m10v, c10p, sir_cast(m10v, sir_builtin(m10v, "instance_index" as *u8, T_I32), T_F32))
2483 sir_arg(m10v, c10p, sir_swz(m10v, sir_ident(m10v, "aP" as *u8, T_V3F), "y" as *u8, T_F32))
2484 sir_arg(m10v, c10p, sir_lit(m10v, "0.0" as *u8, T_F32))
2485 sir_arg(m10v, c10p, sir_lit(m10v, "1.0" as *u8, T_F32))
2486 sir_stmt(m10v, f10v, sir_assign(m10v, sir_builtin(m10v, "position" as *u8, T_V4F), c10p))
2487 let gt10v: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2488 let gn10v: *i64 = sys_mmap(WG_NUMBUF) as *i64
2489 wg_probe(m10v, STAGE_VERTEX, "cast10vs_glsl" as *u8, 0, gt10v, gn10v)
2490 wg_probe(m10v, STAGE_VERTEX, "cast10vs" as *u8, 1, tr, tn)
2491 let nb10: *u8 = sys_mmap(WG_NUMBUF)
2492 wg_puts("cast10vs_trace_equal=" as *u8)
2493 var q10: i64 = eb_num(nb10, 0, WG_NUMBUF, wg_trace_eq(gt10v, gn10v[0], tr, tn[0])); nb10[q10] = 0 as u8; wg_puts(nb10)
2494 wg_puts("\n" as *u8)
2495 // S12c-1 (2026-09-05) COVERAGE -- a VOID return inside the interstage-struct vertex entry: cast11vs writes vo.vcol and
2496 // returns early under an if (the hair-lock cull shape). WGSL hands back the struct (`return vo;`), GLSL says `return;`.
2497 let m11v: *i64 = sir_new()
2498 sir_attrib(m11v, "aP" as *u8, T_V3F, 0)
2499 sir_vary(m11v, "vcol" as *u8, T_V4F, 0, 0)
2500 let f11v: i64 = sir_func(m11v, "main" as *u8, T_VOID)
2501 let c11z: i64 = sir_ctor(m11v, T_V4F)
2502 sir_arg(m11v, c11z, sir_lit(m11v, "0.0" as *u8, T_F32))
2503 var tc11: i64 = 0
2504 tc11 = sir_seq(m11v, tc11, sir_assign(m11v, sir_ident(m11v, "vcol" as *u8, T_V4F), c11z))
2505 tc11 = sir_seq(m11v, tc11, sir_return(m11v, 0))
2506 sir_stmt(m11v, f11v, sir_if(m11v, sir_bin(m11v, "<" as *u8, sir_swz(m11v, sir_ident(m11v, "aP" as *u8, T_V3F), "y" as *u8, T_F32), sir_lit(m11v, "0.0" as *u8, T_F32), T_BOOL), tc11, 0))
2507 let c11c: i64 = sir_ctor(m11v, T_V4F)
2508 sir_arg(m11v, c11c, sir_ident(m11v, "aP" as *u8, T_V3F))
2509 sir_arg(m11v, c11c, sir_lit(m11v, "1.0" as *u8, T_F32))
2510 sir_stmt(m11v, f11v, sir_assign(m11v, sir_ident(m11v, "vcol" as *u8, T_V4F), c11c))
2511 let c11p: i64 = sir_ctor(m11v, T_V4F)
2512 sir_arg(m11v, c11p, sir_ident(m11v, "aP" as *u8, T_V3F))
2513 sir_arg(m11v, c11p, sir_lit(m11v, "1.0" as *u8, T_F32))
2514 sir_stmt(m11v, f11v, sir_assign(m11v, sir_builtin(m11v, "position" as *u8, T_V4F), c11p))
2515 let gt11v: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2516 let gn11v: *i64 = sys_mmap(WG_NUMBUF) as *i64
2517 wg_probe(m11v, STAGE_VERTEX, "cast11vs_glsl" as *u8, 0, gt11v, gn11v)
2518 wg_probe(m11v, STAGE_VERTEX, "cast11vs" as *u8, 1, tr, tn)
2519 let nb11: *u8 = sys_mmap(WG_NUMBUF)
2520 wg_puts("cast11vs_trace_equal=" as *u8)
2521 var q11: i64 = eb_num(nb11, 0, WG_NUMBUF, wg_trace_eq(gt11v, gn11v[0], tr, tn[0])); nb11[q11] = 0 as u8; wg_puts(nb11)
2522 wg_puts("\n" as *u8)
2523 // S12c-5 fixture (2026-09-05): a fragment that ASSIGNS its declared output and exits early with a void return -- the hand
2524 // cast fragment's shape. GLSL: `fc=...;` and `return;` (fc is an out variable). WGSL: `var fc:vec4f;` opens the body, each
2525 // void return becomes `return fc;`, a trailing `return fc;` closes it. It also carries a const local and a u-suffixed unsigned
2526 // literal -- the two other spellings the cast fragment needs in BOTH dialects. Trace-identical (the lowering emits no node).
2527 let m12f: *i64 = sir_new()
2528 sir_uniform(m12f, "a" as *u8, T_F32, 0)
2529 sir_fragout(m12f, "fc" as *u8, T_V4F, 0)
2530 let f12f: i64 = sir_func(m12f, "main" as *u8, T_VOID)
2531 let c12a: i64 = sir_ctor(m12f, T_V4F)
2532 sir_arg(m12f, c12a, sir_lit(m12f, "1.0" as *u8, T_F32))
2533 var tc12: i64 = 0
2534 tc12 = sir_seq(m12f, tc12, sir_assign(m12f, sir_ident(m12f, "fc" as *u8, T_V4F), c12a))
2535 tc12 = sir_seq(m12f, tc12, sir_return(m12f, 0))
2536 sir_stmt(m12f, f12f, sir_if(m12f, sir_bin(m12f, "<" as *u8, sir_ident(m12f, "a" as *u8, T_F32), sir_lit(m12f, "0.5" as *u8, T_F32), T_BOOL), tc12, 0))
2537 sir_stmt(m12f, f12f, sir_const(m12f, "K" as *u8, T_F32, sir_lit(m12f, "1.2" as *u8, T_F32)))
2538 sir_stmt(m12f, f12f, sir_var(m12f, "z" as *u8, T_U32, sir_lit_u(m12f, "0" as *u8)))
2539 let c12b: i64 = sir_ctor(m12f, T_V4F)
2540 sir_arg(m12f, c12b, sir_ident(m12f, "K" as *u8, T_F32))
2541 sir_stmt(m12f, f12f, sir_assign(m12f, sir_ident(m12f, "fc" as *u8, T_V4F), c12b))
2542 let gt12f: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2543 let gn12f: *i64 = sys_mmap(WG_NUMBUF) as *i64
2544 wg_probe(m12f, STAGE_FRAGMENT, "fragassign_glsl" as *u8, 0, gt12f, gn12f)
2545 wg_probe(m12f, STAGE_FRAGMENT, "fragassign_wgsl" as *u8, 1, tr, tn)
2546 let nb12: *u8 = sys_mmap(WG_NUMBUF)
2547 wg_puts("fragassign_trace_equal=" as *u8)
2548 var q12: i64 = eb_num(nb12, 0, WG_NUMBUF, wg_trace_eq(gt12f, gn12f[0], tr, tn[0])); nb12[q12] = 0 as u8; wg_puts(nb12)
2549 wg_puts("\n" as *u8)
2550 let m8f: *i64 = sir_new()
2551 sir_texture(m8f, "uAtl" as *u8, T_TEX2F)
2552 sir_texture(m8f, "uNrm" as *u8, T_TEX2F)
2553 sir_texture(m8f, "uJT" as *u8, T_TEX2F)
2554 sir_texture(m8f, "uGD" as *u8, T_TEX2F)
2555 sir_vary(m8f, "vcol" as *u8, T_V4F, 0, 0)
2556 sir_fragout(m8f, "fc" as *u8, T_V4F, 0)
2557 let f8f: i64 = sir_func(m8f, "main" as *u8, T_VOID)
2558 let s8a: i64 = sir_texsample(m8f, sir_ident(m8f, "uAtl" as *u8, T_TEX2F), sir_swz(m8f, sir_ident(m8f, "vcol" as *u8, T_V4F), "xy" as *u8, T_V2F))
2559 let s8n: i64 = sir_texsample(m8f, sir_ident(m8f, "uNrm" as *u8, T_TEX2F), sir_swz(m8f, sir_ident(m8f, "vcol" as *u8, T_V4F), "xy" as *u8, T_V2F))
2560 sir_stmt(m8f, f8f, sir_return(m8f, sir_bin(m8f, "*" as *u8, s8a, s8n, T_V4F)))
2561 let gt8f: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2562 let gn8f: *i64 = sys_mmap(WG_NUMBUF) as *i64
2563 wg_probe(m8f, STAGE_FRAGMENT, "cast8fs_glsl" as *u8, 0, gt8f, gn8f)
2564 wg_probe(m8f, STAGE_FRAGMENT, "cast8fs" as *u8, 1, tr, tn)
2565 wg_puts("cast8fs_trace_equal=" as *u8)
2566 q8 = eb_num(nb8, 0, WG_NUMBUF, wg_trace_eq(gt8f, gn8f[0], tr, tn[0])); nb8[q8] = 0 as u8; wg_puts(nb8)
2567 wg_puts("\n" as *u8)
2568 // POSITIVE CONTROL -- the same passes ADMIT a scalar uniform, in both dialects, trace-identical
2569 let m4: *i64 = sir_new()
2570 shsrc_fullscreen_tri_scaled(m4)
2571 let gt: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2572 let gn: *i64 = sys_mmap(WG_NUMBUF) as *i64
2573 wg_probe(m4, STAGE_VERTEX, "posctl_glsl" as *u8, 0, gt, gn)
2574 wg_probe(m4, STAGE_VERTEX, "posctl_wgsl" as *u8, 1, tr, tn)
2575 let nb: *u8 = sys_mmap(WG_NUMBUF)
2576 wg_puts("posctl_trace_equal=" as *u8)
2577 var q: i64 = eb_num(nb, 0, WG_NUMBUF, wg_trace_eq(gt, gn[0], tr, tn[0])); nb[q] = 0 as u8; wg_puts(nb)
2578 wg_puts("\n" as *u8)
2579 // NEG-CONTROL S2a -- a LOCAL named like a uniform is read bare (it shadows the uniform)
2580 let m5: *i64 = sir_new()
2581 shsrc_fullscreen_tri_shadowlocal(m5)
2582 wg_probe(m5, STAGE_VERTEX, "shadowlocal" as *u8, 1, tr, tn)
2583 // NEG-CONTROL S2b -- a PARAMETER named like a uniform is read bare
2584 let m6: *i64 = sir_new()
2585 shsrc_plain_shadowparam(m6)
2586 wg_probe(m6, STAGE_PLAIN, "shadowparam" as *u8, 1, tr, tn)
2587 // S3 COVERAGE -- K_TEXTURE through BOTH backends: binding var OUTSIDE struct U / sampler uniform, trace-identical
2588 let m7: *i64 = sir_new()
2589 shsrc_tex_plain(m7)
2590 let gt7: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2591 let gn7: *i64 = sys_mmap(WG_NUMBUF) as *i64
2592 wg_probe(m7, STAGE_PLAIN, "tex_glsl" as *u8, 0, gt7, gn7)
2593 wg_probe(m7, STAGE_PLAIN, "tex_wgsl" as *u8, 1, tr, tn)
2594 wg_puts("tex_trace_equal=" as *u8)
2595 q = eb_num(nb, 0, WG_NUMBUF, wg_trace_eq(gt7, gn7[0], tr, tn[0])); nb[q] = 0 as u8; wg_puts(nb)
2596 wg_puts("\n" as *u8)
2597 // GE54 tex2d_sample: a sampled texture_2d<f32> emits in BOTH dialects (GLSL sampler2D, WGSL texture + companion sampler)
2598 let m11: *i64 = sir_new()
2599 shsrc_tex2d_sample(m11)
2600 let gt11: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2601 let gn11: *i64 = sys_mmap(WG_NUMBUF) as *i64
2602 wg_probe(m11, STAGE_PLAIN, "tex2dsmp_glsl" as *u8, 0, gt11, gn11)
2603 wg_probe(m11, STAGE_PLAIN, "tex2dsmp_wgsl" as *u8, 1, tr, tn)
2604 wg_puts("tex2dsmp_trace_equal=" as *u8)
2605 q = eb_num(nb, 0, WG_NUMBUF, wg_trace_eq(gt11, gn11[0], tr, tn[0])); nb[q] = 0 as u8; wg_puts(nb)
2606 wg_puts("\n" as *u8)
2607 // GE55 K_PRIVATE -- module-scope private state through BOTH backends: var<private> / bare global, written by one
2608 // function and read bare by another, NEVER qualified as u.<name>, trace-identical across dialects.
2609 let m8: *i64 = sir_new()
2610 shsrc_private_demo(m8)
2611 let gt8: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2612 let gn8: *i64 = sys_mmap(WG_NUMBUF) as *i64
2613 wg_probe(m8, STAGE_PLAIN, "priv_glsl" as *u8, 0, gt8, gn8)
2614 wg_probe(m8, STAGE_PLAIN, "priv_wgsl" as *u8, 1, tr, tn)
2615 wg_puts("priv_trace_equal=" as *u8)
2616 q = eb_num(nb, 0, WG_NUMBUF, wg_trace_eq(gt8, gn8[0], tr, tn[0])); nb[q] = 0 as u8; wg_puts(nb)
2617 wg_puts("\n" as *u8)
2618 // NEG-CONTROL -- a texture-typed private must refuse BY NAME in each dialect (module state is never a resource)
2619 let m9: *i64 = sir_new()
2620 sir_private(m9, "tp" as *u8, T_TEX3U, 0)
2621 shsrc_fullscreen_tri(m9)
2622 wg_probe(m9, STAGE_VERTEX, "privtex_glsl" as *u8, 0, gt8, gn8)
2623 let m10: *i64 = sir_new()
2624 sir_private(m10, "tp" as *u8, T_TEX3U, 0)
2625 shsrc_fullscreen_tri(m10)
2626 wg_probe(m10, STAGE_VERTEX, "privtex_wgsl" as *u8, 1, tr, tn)
2627 // NEG-CONTROL S3a -- a texture type outside the covered subset refuses BY NAME in BOTH dialects
2628 let m8: *i64 = sir_new()
2629 sir_texture(m8, "img" as *u8, T_TEX2F)
2630 shsrc_frag_const(m8)
2631 wg_probe(m8, STAGE_FRAGMENT, "tex2d_wgsl" as *u8, 1, tr, tn)
2632 let m9: *i64 = sir_new()
2633 sir_texture(m9, "img" as *u8, T_TEX2F)
2634 shsrc_frag_const(m9)
2635 wg_probe(m9, STAGE_FRAGMENT, "tex2d_glsl" as *u8, 0, tr, tn)
2636 // NEG-CONTROL S3b -- the deprecated form (a texture-typed K_UNIFORM) refuses in GLSL too, naming K_TEXTURE
2637 let m10: *i64 = sir_new()
2638 sir_uniform(m10, "vox" as *u8, T_TEX3U, 0)
2639 shsrc_fullscreen_tri(m10)
2640 wg_probe(m10, STAGE_VERTEX, "texuniform_glsl" as *u8, 0, tr, tn)
2641 // NEG-CONTROL S3c -- a derived texture binding landing on an explicit storage binding refuses by name
2642 let m11: *i64 = sir_new()
2643 shsrc_depth_min(m11)
2644 sir_texture(m11, "vox" as *u8, T_TEX3U)
2645 wg_probe(m11, STAGE_COMPUTE, "texcollide" as *u8, 1, tr, tn)
2646 // S4 COVERAGE -- the fragment entry signature is DERIVED from the declared position parameter, both dialects
2647 let m12: *i64 = sir_new()
2648 shsrc_frag_pos(m12)
2649 let gt12: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2650 let gn12: *i64 = sys_mmap(WG_NUMBUF) as *i64
2651 wg_probe(m12, STAGE_FRAGMENT, "fragpos_glsl" as *u8, 0, gt12, gn12)
2652 wg_probe(m12, STAGE_FRAGMENT, "fragpos_wgsl" as *u8, 1, tr, tn)
2653 wg_puts("fragpos_trace_equal=" as *u8)
2654 q = eb_num(nb, 0, WG_NUMBUF, wg_trace_eq(gt12, gn12[0], tr, tn[0])); nb[q] = 0 as u8; wg_puts(nb)
2655 wg_puts("\n" as *u8)
2656 // CONTROL S4 -- the zero-parameter fragment derives to exactly the signature the literal used to carry
2657 let m13: *i64 = sir_new()
2658 shsrc_frag_const(m13)
2659 wg_probe(m13, STAGE_FRAGMENT, "fragconst" as *u8, 1, tr, tn)
2660 // NEG-CONTROL S4a -- reading position with NO declared parameter is an undeclared input: refused by name
2661 let m14: *i64 = sir_new()
2662 shsrc_frag_pos_undeclared(m14)
2663 wg_probe(m14, STAGE_FRAGMENT, "fragundecl" as *u8, 1, tr, tn)
2664 // NEG-CONTROL S4b -- a plain (unbound) fragment parameter refuses by name in BOTH dialects
2665 let m15: *i64 = sir_new()
2666 shsrc_frag_plainparam(m15)
2667 wg_probe(m15, STAGE_FRAGMENT, "fragplain_wgsl" as *u8, 1, tr, tn)
2668 let m16: *i64 = sir_new()
2669 shsrc_frag_plainparam(m16)
2670 wg_probe(m16, STAGE_FRAGMENT, "fragplain_glsl" as *u8, 0, tr, tn)
2671 // S5 COVERAGE -- the declared fragment output through BOTH backends, trace-identical: GLSL declares
2672 // `out vec4 fc;` and ASSIGNS it, WGSL declares nothing and DERIVES `->@location(0) vec4f`.
2673 let m17: *i64 = sir_new()
2674 shsrc_frag_const(m17)
2675 let gt17: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2676 let gn17: *i64 = sys_mmap(WG_NUMBUF) as *i64
2677 wg_probe(m17, STAGE_FRAGMENT, "fragout_glsl" as *u8, 0, gt17, gn17)
2678 wg_probe(m17, STAGE_FRAGMENT, "fragout_wgsl" as *u8, 1, tr, tn)
2679 wg_puts("fragout_trace_equal=" as *u8)
2680 q = eb_num(nb, 0, WG_NUMBUF, wg_trace_eq(gt17, gn17[0], tr, tn[0])); nb[q] = 0 as u8; wg_puts(nb)
2681 wg_puts("\n" as *u8)
2682 // S5 POSITIVE CONTROL -- a discard-only fragment declares NO output and BOTH backends admit it. Without
2683 // this the output rule would be a guard that refuses a legal program and nobody would have measured it.
2684 let m18: *i64 = sir_new()
2685 shsrc_frag_discard(m18)
2686 wg_probe(m18, STAGE_FRAGMENT, "fragdiscard_glsl" as *u8, 0, tr, tn)
2687 let m19: *i64 = sir_new()
2688 shsrc_frag_discard(m19)
2689 wg_probe(m19, STAGE_FRAGMENT, "fragdiscard_wgsl" as *u8, 1, tr, tn)
2690 // S5 NEG-CONTROL -- returning a value with no declared output refuses BY NAME in both dialects
2691 let m20: *i64 = sir_new()
2692 shsrc_frag_noout(m20)
2693 wg_probe(m20, STAGE_FRAGMENT, "fragnoout_glsl" as *u8, 0, tr, tn)
2694 let m21: *i64 = sir_new()
2695 shsrc_frag_noout(m21)
2696 wg_probe(m21, STAGE_FRAGMENT, "fragnoout_wgsl" as *u8, 1, tr, tn)
2697 // S5 COVERAGE -- a NON-ZERO @location proves the location is data, not the constant 0
2698 let m22: *i64 = sir_new()
2699 shsrc_frag_loc2(m22)
2700 wg_probe(m22, STAGE_FRAGMENT, "fragloc2_glsl" as *u8, 0, tr, tn)
2701 let m23: *i64 = sir_new()
2702 shsrc_frag_loc2(m23)
2703 wg_probe(m23, STAGE_FRAGMENT, "fragloc2_wgsl" as *u8, 1, tr, tn)
2704 // S5 NEG-CONTROL -- a fragment output on a VERTEX stage refuses by name in both dialects
2705 let m24: *i64 = sir_new()
2706 shsrc_vert_fragout(m24)
2707 wg_probe(m24, STAGE_VERTEX, "vertfragout_glsl" as *u8, 0, tr, tn)
2708 let m25: *i64 = sir_new()
2709 shsrc_vert_fragout(m25)
2710 wg_probe(m25, STAGE_VERTEX, "vertfragout_wgsl" as *u8, 1, tr, tn)
2711 // S5 NEG-CONTROL -- two declared outputs (MRT) refuse by name rather than dropping one
2712 let m26: *i64 = sir_new()
2713 shsrc_frag_twoout(m26)
2714 wg_probe(m26, STAGE_FRAGMENT, "fragtwoout_glsl" as *u8, 0, tr, tn)
2715 let m27: *i64 = sir_new()
2716 shsrc_frag_twoout(m27)
2717 wg_probe(m27, STAGE_FRAGMENT, "fragtwoout_wgsl" as *u8, 1, tr, tn)
2718 // S5 NEG-CONTROL -- a sampled texture read refuses in WGSL (undeclared sampler) and EMITS in GLSL.
2719 // Both halves are asserted: a refusal with no admitting twin would not prove the rule is the sampler.
2720 let m28: *i64 = sir_new()
2721 shsrc_texsample(m28)
2722 wg_probe(m28, STAGE_PLAIN, "texsample_wgsl" as *u8, 1, tr, tn)
2723 let m29: *i64 = sir_new()
2724 shsrc_texsample(m29)
2725 wg_probe(m29, STAGE_PLAIN, "texsample_glsl" as *u8, 0, tr, tn)
2726 // R-A / R-C COVERAGE (2026-09-04, GE44 R-A and R-C): a select and an immutable local through BOTH backends,
2727 // trace-identical -- the ternary in GLSL, select(f,t,cond) and `let` in WGSL.
2728 let m30: *i64 = sir_new()
2729 shsrc_select_let(m30)
2730 let gt30: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2731 let gn30: *i64 = sys_mmap(WG_NUMBUF) as *i64
2732 wg_probe(m30, STAGE_PLAIN, "selectlet_glsl" as *u8, 0, gt30, gn30)
2733 wg_probe(m30, STAGE_PLAIN, "selectlet_wgsl" as *u8, 1, tr, tn)
2734 wg_puts("selectlet_trace_equal=" as *u8)
2735 q = eb_num(nb, 0, WG_NUMBUF, wg_trace_eq(gt30, gn30[0], tr, tn[0])); nb[q] = 0 as u8; wg_puts(nb)
2736 wg_puts("\n" as *u8)
2737 // R-B COVERAGE (2026-09-04, GE44 R-B): a counted loop written ONCE as builder sugar (sir_for_until), through
2738 // BOTH backends, trace-identical -- while(true) in GLSL, loop in WGSL, the exit test and the step in the
2739 // shared S_IF/S_LOOP/S_ASSIGN arms, the counter a mutable local in each dialect's own spelling.
2740 let m31: *i64 = sir_new()
2741 shsrc_for_until(m31)
2742 let gt31: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2743 let gn31: *i64 = sys_mmap(WG_NUMBUF) as *i64
2744 wg_probe(m31, STAGE_PLAIN, "foruntil_glsl" as *u8, 0, gt31, gn31)
2745 wg_probe(m31, STAGE_PLAIN, "foruntil_wgsl" as *u8, 1, tr, tn)
2746 wg_puts("foruntil_trace_equal=" as *u8)
2747 q = eb_num(nb, 0, WG_NUMBUF, wg_trace_eq(gt31, gn31[0], tr, tn[0])); nb[q] = 0 as u8; wg_puts(nb)
2748 wg_puts("\n" as *u8)
2749 // R-D COVERAGE (2026-09-04, GE44 R-D): a struct declared once, a struct-typed local built by its constructor,
2750 // a member write and member reads, through BOTH backends, trace-identical -- `struct Hit{vec3 pos;float d;};`
2751 // in GLSL, `struct Hit{pos:vec3f,d:f32,}` in WGSL, `h.d` and `h.pos.x` spelled the same in both.
2752 let m32: *i64 = sir_new()
2753 shsrc_struct(m32)
2754 let gt32: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2755 let gn32: *i64 = sys_mmap(WG_NUMBUF) as *i64
2756 wg_probe(m32, STAGE_PLAIN, "struct_glsl" as *u8, 0, gt32, gn32)
2757 wg_probe(m32, STAGE_PLAIN, "struct_wgsl" as *u8, 1, tr, tn)
2758 wg_puts("struct_trace_equal=" as *u8)
2759 q = eb_num(nb, 0, WG_NUMBUF, wg_trace_eq(gt32, gn32[0], tr, tn[0])); nb[q] = 0 as u8; wg_puts(nb)
2760 wg_puts("\n" as *u8)
2761 // R-E COVERAGE (2026-09-04, GE44 R-E): a struct RETURN and a struct PARAMETER through BOTH backends, one
2762 // canonical shape, trace-identical -- `Hit mk(vec3 p,float k)` / `fn mk(p:vec3f,k:f32)->Hit`, and a member
2763 // read straight off the call (`mk(h.pos,k).d`) spelled the same in both.
2764 let m33: *i64 = sir_new()
2765 shsrc_struct_ret(m33)
2766 let gt33: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2767 let gn33: *i64 = sys_mmap(WG_NUMBUF) as *i64
2768 wg_probe(m33, STAGE_PLAIN, "structret_glsl" as *u8, 0, gt33, gn33)
2769 wg_probe(m33, STAGE_PLAIN, "structret_wgsl" as *u8, 1, tr, tn)
2770 wg_puts("structret_trace_equal=" as *u8)
2771 q = eb_num(nb, 0, WG_NUMBUF, wg_trace_eq(gt33, gn33[0], tr, tn[0])); nb[q] = 0 as u8; wg_puts(nb)
2772 wg_puts("\n" as *u8)
2773 // R-G0 COVERAGE (2026-09-04, GE44 R-G0): an uninitialised local and a unary negation through BOTH backends,
2774 // trace-identical -- `float v;` / `var v:f32;` and `v=(-a);` spelled the same in both.
2775 let m34: *i64 = sir_new()
2776 shsrc_neg_uninit(m34)
2777 let gt34: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
2778 let gn34: *i64 = sys_mmap(WG_NUMBUF) as *i64
2779 wg_probe(m34, STAGE_PLAIN, "neguninit_glsl" as *u8, 0, gt34, gn34)
2780 wg_probe(m34, STAGE_PLAIN, "neguninit_wgsl" as *u8, 1, tr, tn)
2781 wg_puts("neguninit_trace_equal=" as *u8)
2782 q = eb_num(nb, 0, WG_NUMBUF, wg_trace_eq(gt34, gn34[0], tr, tn[0])); nb[q] = 0 as u8; wg_puts(nb)
2783 wg_puts("\n" as *u8)
2784 return 0
2785}
2786
2787// ============================ RUNG 2: THE FRONT HALF ============================
2788// Lower an nx_ir SSA Function -- produced by the REAL lexer/parser/optimiser -- into the shader
2789// IR, so a shader is an ORDINARY NISHILANG FUNCTION rather than builder calls. Straight-line
2790// functions only: GLSL and WGSL have no goto, so a multi-block CFG needs structured control-flow
2791// recovery (the problem WebAssembly had to solve). That is REFUSED BY NAME here, not guessed.
2792
2793// Driver buffers, named for purpose (the magic ratchet refused inline literals and was right).
2794const SHL_EXPAND_CAP: i64 = 4194304
2795const SHL_LEX_CAP: i64 = 262144
2796
2797// nx_ir Type kind -> shader type. i64/f64/pointer/struct have NO GLSL ES 3.00 counterpart, so
2798// they REFUSE rather than narrow: silently turning an i64 into an int is precisely the
2799// mistranslation this backend exists to make impossible.
2800func shl_ty(t: *Type) -> i64 {
2801 if (t as i64) == 0 { return T_VOID }
2802 let k: i64 = t.kind
2803 if k == 0 { return T_VOID }
2804 if k == 1 { return T_BOOL }
2805 if k == 4 { return T_I32 }
2806 if k == 9 { return T_F32 }
2807 return 0 - 1
2808}
2809
2810func shl_op(op: i64) -> *u8 {
2811 if op == OP_ADD { return "+" as *u8 }
2812 if op == OP_SUB { return "-" as *u8 }
2813 if op == OP_MUL { return "*" as *u8 }
2814 if op == OP_DIV_S { return "/" as *u8 }
2815 if op == OP_AND { return "&" as *u8 }
2816 if op == OP_OR { return "|" as *u8 }
2817 if op == OP_XOR { return "^" as *u8 }
2818 if op == OP_SHL { return "<<" as *u8 }
2819 if op == OP_SHR_S { return ">>" as *u8 }
2820 if op == OP_EQ { return "==" as *u8 }
2821 if op == OP_NE { return "!=" as *u8 }
2822 if op == OP_LT_S { return "<" as *u8 }
2823 if op == OP_LE_S { return "<=" as *u8 }
2824 if op == OP_GT_S { return ">" as *u8 }
2825 if op == OP_GE_S { return ">=" as *u8 }
2826 return 0 as *u8
2827}
2828
2829// An SSA operand becomes: a literal (CONST_INT), a parameter name (PARAM), or the named
2830// temporary emitted for the instruction that defined it (INSTR). No expression-tree rebuild is
2831// needed -- one named temporary per instruction is the classic SSA-to-text lowering, and it is
2832// exactly why the straight-line case needs no relooper.
2833func shl_val(m2: *i64, f: *Function, vid: i64, nb: *u8) -> i64 {
2834 let v: *Value = val_at(f, vid)
2835 if (v as i64) == 0 { return 0 }
2836 let ty: i64 = shl_ty(v.ty)
2837 var q: i64 = 0
2838 if v.kind == 0 {
2839 q = eb_num(nb, 0, WG_NUMBUF, v.const_int)
2840 nb[q] = 0 as u8
2841 return sir_lit(m2, nb, ty)
2842 }
2843 if v.kind == 1 {
2844 q = eb_put(nb, 0, WG_NUMBUF, "p" as *u8)
2845 q = eb_num(nb, q, WG_NUMBUF, v.param_index)
2846 nb[q] = 0 as u8
2847 return sir_ident(m2, nb, ty)
2848 }
2849 q = eb_put(nb, 0, WG_NUMBUF, "t" as *u8)
2850 q = eb_num(nb, q, WG_NUMBUF, vid)
2851 nb[q] = 0 as u8
2852 return sir_ident(m2, nb, ty)
2853}
2854
2855// Lower ONE basic block. Statements are appended to fn2 and then DETACHED as a standalone chain,
2856// so the caller can place them inside an if-branch instead of at function top level. Terminators
2857// (BR / BR_COND) emit NOTHING here: control flow is the SHAPE recogniser's job, and a terminator
2858// that leaked through as a statement would be a goto in a language that has none.
2859func shl_blk(m2: *i64, f: *Function, fn2: i64, bb: *BasicBlock, nb: *u8, msg: *u8, seen: *i64) -> i64 {
2860 if (bb as i64) == 0 { sir_refuse(m2, "basic block is absent" as *u8); return 0 - 1 }
2861 let t0: i64 = sir_get(m2, fn2, SIR_D)
2862 var ins: *Instr = bb.head
2863 while (ins as i64) != 0 {
2864 let op: i64 = ins.op
2865 if op == OP_BR { } else { if op == OP_BR_COND { } else {
2866 if op == OP_RETURN {
2867 var e: i64 = 0
2868 if ins.n_operands > 0 { e = shl_val(m2, f, ins.op0, nb) }
2869 sir_stmt(m2, fn2, sir_return(m2, e))
2870 } else { if op == OP_ALLOCA {
2871 // DELIBERATELY NO EMISSION -- see shl_lower.
2872 } else { if op == OP_STORE {
2873 let ad: i64 = ins.op0
2874 let vx: i64 = shl_val(m2, f, ins.op1, nb)
2875 let vv: *Value = val_at(f, ins.op1)
2876 var vt: i64 = 0 - 1
2877 if (vv as i64) != 0 { vt = shl_ty(vv.ty) }
2878 if vt < 0 { sir_refuse(m2, "stored value type has no GLSL ES 3.00 counterpart -- refused rather than narrowed" as *u8); return 0 - 1 }
2879 var qa: i64 = eb_put(nb, 0, WG_NUMBUF, "a" as *u8)
2880 qa = eb_num(nb, qa, WG_NUMBUF, ad)
2881 nb[qa] = 0 as u8
2882 if seen[ad] == 0 {
2883 seen[ad] = 1
2884 sir_stmt(m2, fn2, sir_var(m2, nb, vt, vx))
2885 } else {
2886 sir_stmt(m2, fn2, sir_assign(m2, sir_ident(m2, nb, vt), vx))
2887 }
2888 } else { if op == OP_LOAD {
2889 let ad2: i64 = ins.op0
2890 let lt2: i64 = shl_ty(ins.ty)
2891 if lt2 < 0 { sir_refuse(m2, "loaded value type has no GLSL ES 3.00 counterpart -- refused rather than narrowed" as *u8); return 0 - 1 }
2892 var qb: i64 = eb_put(nb, 0, WG_NUMBUF, "a" as *u8)
2893 qb = eb_num(nb, qb, WG_NUMBUF, ad2)
2894 nb[qb] = 0 as u8
2895 let srcv: i64 = sir_ident(m2, nb, lt2)
2896 var qc: i64 = eb_put(nb, 0, WG_NUMBUF, "t" as *u8)
2897 qc = eb_num(nb, qc, WG_NUMBUF, ins.result)
2898 nb[qc] = 0 as u8
2899 sir_stmt(m2, fn2, sir_var(m2, nb, lt2, srcv))
2900 } else {
2901 let ops: *u8 = shl_op(op)
2902 if (ops as i64) == 0 {
2903 var qm: i64 = eb_put(msg, 0, WG_MSGBUF, "opcode outside the covered subset: op=" as *u8)
2904 qm = eb_num(msg, qm, WG_MSGBUF, op)
2905 msg[qm] = 0 as u8
2906 sir_refuse(m2, msg)
2907 return 0 - 1
2908 }
2909 let lt: i64 = shl_ty(ins.ty)
2910 if lt < 0 { sir_refuse(m2, "instruction result type has no GLSL ES 3.00 counterpart -- refused rather than narrowed" as *u8); return 0 - 1 }
2911 let l: i64 = shl_val(m2, f, ins.op0, nb)
2912 let r: i64 = shl_val(m2, f, ins.op1, nb)
2913 let bx: i64 = sir_bin(m2, ops, l, r, lt)
2914 var qt: i64 = eb_put(nb, 0, WG_NUMBUF, "t" as *u8)
2915 qt = eb_num(nb, qt, WG_NUMBUF, ins.result)
2916 nb[qt] = 0 as u8
2917 sir_stmt(m2, fn2, sir_var(m2, nb, lt, bx))
2918 } } } } } }
2919 ins = ins.next
2920 }
2921 var head: i64 = 0
2922 if t0 == 0 {
2923 head = sir_get(m2, fn2, SIR_B)
2924 sir_set(m2, fn2, SIR_B, 0)
2925 sir_set(m2, fn2, SIR_D, 0)
2926 } else {
2927 head = sir_get(m2, t0, SIR_NEXT)
2928 sir_set(m2, t0, SIR_NEXT, 0)
2929 sir_set(m2, fn2, SIR_D, t0)
2930 }
2931 return head
2932}
2933
2934// ---- RUNG 2b-GENERAL: recursive region reduction ------------------------------------------
2935// block_at() takes an INDEX; the CFG talks in IDs. Resolve rather than assume they coincide.
2936func shl_bid(f: *Function, id: i64) -> *BasicBlock {
2937 var i: i64 = 0
2938 while i < f.n_blocks {
2939 let b: *BasicBlock = block_at(f, i)
2940 if (b as i64) != 0 { if b.id == id { return b } }
2941 i = i + 1
2942 }
2943 return 0 as *BasicBlock
2944}
2945
2946// Can control reach `to` starting at `from`? Breadth-first over successors. This is the ONLY
2947// back-edge test used below: a heuristic on block ORDER would silently mis-shape any CFG whose
2948// numbering is not the traversal order, and the whole point of this rung is to stop guessing.
2949func shl_reaches(f: *Function, from: i64, to: i64, mark: *i64, wl: *i64) -> i64 {
2950 if from < 0 { return 0 }
2951 var i: i64 = 0
2952 while i < f.n_blocks { mark[i] = 0; i = i + 1 }
2953 var wn: i64 = 0
2954 var wp: i64 = 0
2955 if from < f.n_blocks { mark[from] = 1 }
2956 wl[wn] = from
2957 wn = wn + 1
2958 while wp < wn {
2959 let cid: i64 = wl[wp]
2960 wp = wp + 1
2961 if cid == to { return 1 }
2962 let b: *BasicBlock = shl_bid(f, cid)
2963 if (b as i64) != 0 {
2964 if b.n_succs > 0 { if (b.succ0 as i64) != 0 {
2965 let s: i64 = b.succ0.id
2966 if s >= 0 { if s < f.n_blocks { if mark[s] == 0 { mark[s] = 1; wl[wn] = s; wn = wn + 1 } } }
2967 } }
2968 if b.n_succs > 1 { if (b.succ1 as i64) != 0 {
2969 let s2: i64 = b.succ1.id
2970 if s2 >= 0 { if s2 < f.n_blocks { if mark[s2] == 0 { mark[s2] = 1; wl[wn] = s2; wn = wn + 1 } } }
2971 } }
2972 }
2973 }
2974 return 0
2975}
2976
2977// Reachability that must not pass THROUGH `avoid`. This is the whole dominance test: B dominates
2978// P exactly when P cannot be reached from the entry without going through B.
2979func shl_reaches_avoid(f: *Function, from: i64, to: i64, avoid: i64, mark: *i64, wl: *i64) -> i64 {
2980 if from < 0 { return 0 }
2981 if from == avoid { return 0 }
2982 var i: i64 = 0
2983 while i < f.n_blocks { mark[i] = 0; i = i + 1 }
2984 var wn: i64 = 0
2985 var wp: i64 = 0
2986 if from < f.n_blocks { mark[from] = 1 }
2987 wl[wn] = from
2988 wn = wn + 1
2989 while wp < wn {
2990 let cid: i64 = wl[wp]
2991 wp = wp + 1
2992 if cid == to { return 1 }
2993 let b: *BasicBlock = shl_bid(f, cid)
2994 if (b as i64) != 0 {
2995 if b.n_succs > 0 { if (b.succ0 as i64) != 0 {
2996 let s: i64 = b.succ0.id
2997 if s != avoid { if s >= 0 { if s < f.n_blocks { if mark[s] == 0 { mark[s] = 1; wl[wn] = s; wn = wn + 1 } } } }
2998 } }
2999 if b.n_succs > 1 { if (b.succ1 as i64) != 0 {
3000 let s2: i64 = b.succ1.id
3001 if s2 != avoid { if s2 >= 0 { if s2 < f.n_blocks { if mark[s2] == 0 { mark[s2] = 1; wl[wn] = s2; wn = wn + 1 } } } }
3002 } }
3003 }
3004 }
3005 return 0
3006}
3007
3008func shl_dominates(f: *Function, b: i64, p: i64, mark: *i64, wl: *i64) -> i64 {
3009 if b == p { return 1 }
3010 if (f.entry as i64) == 0 { return 0 }
3011 if shl_reaches_avoid(f, f.entry.id, p, b, mark, wl) == 1 { return 0 }
3012 return 1
3013}
3014
3015// A block is a loop HEADER iff some edge P -> bb is a BACK EDGE, i.e. bb DOMINATES P.
3016//
3017// THE OBVIOUS PREDICATE IS WRONG AND THIS LANE SHIPPED IT ONCE: "can one of my successors reach me
3018// again" is true of EVERY block in a cycle, not just its header, so it flagged the latch as a
3019// header and refused a textbook while-loop for having n_preds=1. It was caught because the refusal
3020// CONTRADICTED THE CFG DUMP PRINTED DIRECTLY ABOVE IT -- the message said n_preds != 2 about a
3021// block the dump showed with preds=2, which can only mean the predicate had selected a different
3022// block. A refusal that names its numbers is what made a wrong predicate visible in one run.
3023func shl_is_header(f: *Function, bb: *BasicBlock, mark: *i64, wl: *i64) -> i64 {
3024 var i: i64 = 0
3025 while i < f.n_blocks {
3026 let p: *BasicBlock = block_at(f, i)
3027 if (p as i64) != 0 {
3028 var edge: i64 = 0
3029 if p.n_succs > 0 { if (p.succ0 as i64) != 0 { if p.succ0.id == bb.id { edge = 1 } } }
3030 if p.n_succs > 1 { if (p.succ1 as i64) != 0 { if p.succ1.id == bb.id { edge = 1 } } }
3031 if edge == 1 { if shl_dominates(f, bb.id, p.id, mark, wl) == 1 { return 1 } }
3032 }
3033 i = i + 1
3034 }
3035 return 0
3036}
3037
3038// Lower the region that begins at block `eid` and ends when control reaches `stop` (a block id,
3039// or -1 for "run to the end of the function"). Returns a statement chain, or -1 having NAMED the
3040// refusal. This is the Relooper's shape trichotomy -- Simple / Loop / Multiple -- restricted to
3041// the single-entry reducible case that MEASUREMENT showed our front-end actually emits, and
3042// refusing everything else by name instead of guessing at it.
3043//
3044// WHY NO LABEL VARIABLE, AND WHY NO NODE DUPLICATION. The Relooper needs both because it accepts
3045// arbitrary CFGs: its label variable exists to re-enter the right block of a MULTI-ENTRY region.
3046// Every loop header in our measured corpus has exactly TWO predecessors -- one preheader, one
3047// latch -- so there is no multi-entry region for a label variable to disambiguate. LLVM's
3048// WebAssembly Stackifier states the same precondition in its own header comment: its marker
3049// insertion "is sufficient to convert arbitrary CFGs into a form that works on WebAssembly,
3050// provided that all loops are single-entry" (read from the mirrored source, key llvm-cfgstackify).
3051// The moment a header appears with n_preds != 2 this refuses BY NAME rather than mis-shaping it.
3052func shl_region(m2: *i64, f: *Function, fn2: i64, eid: i64, stop: i64, nb: *u8, msg: *u8, seen: *i64, mark: *i64, wl: *i64, depth: i64) -> i64 {
3053 if depth > f.n_blocks { sir_refuse(m2, "region nesting deeper than the block count -- the CFG is not the reducible shape this reducer accepts" as *u8); return 0 - 1 }
3054 var chain: i64 = 0
3055 var cur: i64 = eid
3056 var step: i64 = 0
3057 let maxstep: i64 = f.n_blocks + 1
3058 while step < maxstep {
3059 step = step + 1
3060 if cur < 0 { return chain }
3061 if cur == stop { return chain }
3062 let bb: *BasicBlock = shl_bid(f, cur)
3063 if (bb as i64) == 0 { sir_refuse(m2, "region walk reached a block id with no block" as *u8); return 0 - 1 }
3064
3065 if shl_is_header(f, bb, mark, wl) == 1 {
3066 if bb.n_preds != 2 { sir_refuse(m2, "loop header with n_preds != 2: a MULTI-ENTRY or multi-latch loop needs the Relooper's label variable or node duplication -- rung 2b-irreducible, refused rather than guessed" as *u8); return 0 - 1 }
3067 if bb.n_succs != 2 { sir_refuse(m2, "loop header without a two-way exit test -- a loop with no break has no terminating shader meaning" as *u8); return 0 - 1 }
3068 if (bb.succ0 as i64) == 0 { sir_refuse(m2, "loop header succ0 missing" as *u8); return 0 - 1 }
3069 if (bb.succ1 as i64) == 0 { sir_refuse(m2, "loop header succ1 missing" as *u8); return 0 - 1 }
3070 let s0: i64 = bb.succ0.id
3071 let s1: i64 = bb.succ1.id
3072 var body0: i64 = 0 - 1
3073 var exitb: i64 = 0 - 1
3074 var thenisbody: i64 = 0
3075 if shl_reaches(f, s0, bb.id, mark, wl) == 1 { body0 = s0; exitb = s1; thenisbody = 1 }
3076 if body0 < 0 { if shl_reaches(f, s1, bb.id, mark, wl) == 1 { body0 = s1; exitb = s0; thenisbody = 0 } }
3077 if body0 < 0 { sir_refuse(m2, "loop header has a back edge but neither successor re-reaches it" as *u8); return 0 - 1 }
3078 var cnd: i64 = 0
3079 var tm: *Instr = bb.head
3080 while (tm as i64) != 0 {
3081 if tm.op == OP_BR_COND { cnd = shl_val(m2, f, tm.op0, nb) }
3082 tm = tm.next
3083 }
3084 if cnd == 0 { sir_refuse(m2, "loop header carries no BR_COND condition operand" as *u8); return 0 - 1 }
3085 // The header's own straight-line code is the loop TEST and must be re-evaluated every
3086 // iteration, so it is emitted INSIDE the loop, above the break.
3087 let hc: i64 = shl_blk(m2, f, fn2, bb, nb, msg, seen)
3088 if hc == (0 - 1) { return 0 - 1 }
3089 let bc: i64 = shl_region(m2, f, fn2, body0, bb.id, nb, msg, seen, mark, wl, depth + 1)
3090 if bc == (0 - 1) { return 0 - 1 }
3091 let brk: i64 = sir_break(m2)
3092 // No NOT node is invented: the two polarities are expressed by which ARM carries the
3093 // break, so an inverted test can never be silently emitted as an un-inverted one.
3094 var ifs: i64 = 0
3095 if thenisbody == 1 { ifs = sir_if(m2, cnd, bc, brk) }
3096 if thenisbody == 0 { ifs = sir_if(m2, cnd, brk, bc) }
3097 let lbody: i64 = sir_seq(m2, hc, ifs)
3098 let lp: i64 = sir_loop(m2, lbody)
3099 chain = sir_seq(m2, chain, lp)
3100 cur = exitb
3101 }
3102 if shl_is_header(f, bb, mark, wl) == 0 {
3103 if bb.n_succs == 0 {
3104 let tc0: i64 = shl_blk(m2, f, fn2, bb, nb, msg, seen)
3105 if tc0 == (0 - 1) { return 0 - 1 }
3106 return sir_seq(m2, chain, tc0)
3107 }
3108 if bb.n_succs == 1 {
3109 let sc: i64 = shl_blk(m2, f, fn2, bb, nb, msg, seen)
3110 if sc == (0 - 1) { return 0 - 1 }
3111 chain = sir_seq(m2, chain, sc)
3112 if (bb.succ0 as i64) == 0 { sir_refuse(m2, "single-successor block with a null successor" as *u8); return 0 - 1 }
3113 cur = bb.succ0.id
3114 }
3115 if bb.n_succs == 2 {
3116 if (bb.succ0 as i64) == 0 { sir_refuse(m2, "two-way block succ0 missing" as *u8); return 0 - 1 }
3117 if (bb.succ1 as i64) == 0 { sir_refuse(m2, "two-way block succ1 missing" as *u8); return 0 - 1 }
3118 let tb2: *BasicBlock = bb.succ0
3119 let jb2: *BasicBlock = bb.succ1
3120 // SHAPE A -- if WITHOUT else: the then-arm falls through to the OTHER successor,
3121 // so that other successor IS the join and there are three blocks in play.
3122 var ok2: i64 = 0
3123 // SHAPE B -- if/ELSE: BOTH arms fall through to a common THIRD block. Four blocks.
3124 // This shape was found by its own negative control (knowledge/nxsh_ifelse.nx),
3125 // which refused by name and thereby named the gap: shaders use else constantly.
3126 var okelse: i64 = 0
3127 var joinid: i64 = 0 - 1
3128 if tb2.n_succs == 1 { if (tb2.succ0 as i64) != 0 {
3129 if tb2.succ0.id == jb2.id { if jb2.n_preds == 2 { ok2 = 1; joinid = jb2.id } }
3130 } }
3131 if ok2 == 0 { if tb2.n_succs == 1 { if jb2.n_succs == 1 {
3132 if (tb2.succ0 as i64) != 0 { if (jb2.succ0 as i64) != 0 {
3133 if tb2.succ0.id == jb2.succ0.id {
3134 if tb2.succ0.n_preds == 2 { okelse = 1; joinid = tb2.succ0.id }
3135 }
3136 } }
3137 } } }
3138 if ok2 == 0 { if okelse == 0 { sir_refuse(m2, "two-way branch that is neither an if-diamond (then-arm falls through to the other successor) nor an if/else (both arms fall through to one common join whose only predecessors are those two arms) -- rung 2b-irreducible" as *u8); return 0 - 1 } }
3139 var cnd2: i64 = 0
3140 var tm2: *Instr = bb.head
3141 while (tm2 as i64) != 0 {
3142 if tm2.op == OP_BR_COND { cnd2 = shl_val(m2, f, tm2.op0, nb) }
3143 tm2 = tm2.next
3144 }
3145 if cnd2 == 0 { sir_refuse(m2, "if-diamond header carries no BR_COND condition operand" as *u8); return 0 - 1 }
3146 let hc2: i64 = shl_blk(m2, f, fn2, bb, nb, msg, seen)
3147 if hc2 == (0 - 1) { return 0 - 1 }
3148 let tc2: i64 = shl_region(m2, f, fn2, tb2.id, joinid, nb, msg, seen, mark, wl, depth + 1)
3149 if tc2 == (0 - 1) { return 0 - 1 }
3150 // The else arm is a REGION too, so an else containing a loop, or a nested if, or
3151 // anything else the reducer knows, composes for free. sir_if already carries an
3152 // else slot (SIR_C) and both emitters already render it: 0 means no else arm.
3153 var ec2: i64 = 0
3154 if okelse == 1 {
3155 ec2 = shl_region(m2, f, fn2, jb2.id, joinid, nb, msg, seen, mark, wl, depth + 1)
3156 if ec2 == (0 - 1) { return 0 - 1 }
3157 }
3158 chain = sir_seq(m2, chain, hc2)
3159 chain = sir_seq(m2, chain, sir_if(m2, cnd2, tc2, ec2))
3160 cur = joinid
3161 }
3162 if bb.n_succs > 2 { sir_refuse(m2, "block with more than two successors (a switch) -- the Relooper's Multiple block is rung 2b-multiple, not implemented, refused rather than guessed" as *u8); return 0 - 1 }
3163 }
3164 }
3165 sir_refuse(m2, "region walk exceeded its derived step budget (one visit per block) -- the CFG is not the reducible shape this reducer accepts" as *u8)
3166 return 0 - 1
3167}
3168
3169func shl_lower(m2: *i64, f: *Function, fname: *u8) -> i64 {
3170 let nb: *u8 = sys_mmap(WG_NUMBUF)
3171 let msg: *u8 = sys_mmap(WG_MSGBUF)
3172 // One flag per SSA value: has this alloca been stored to yet? The FIRST store declares the
3173 // shader local (with the stored value's type, which is where the type actually comes from --
3174 // an alloca's own type is a pointer, and a shader has no pointers); later stores assign to it.
3175 // sys_mmap zero-fills, so "not yet stored" needs no initialisation pass.
3176 let seen: *i64 = sys_mmap(f.n_values*8 + 8) as *i64
3177 let rt: i64 = shl_ty(f.ret_ty)
3178 if rt < 0 { sir_refuse(m2, "return type has no GLSL ES 3.00 counterpart (i64/f64/pointer) -- refused rather than narrowed" as *u8); return 0 - 1 }
3179 let fn2: i64 = sir_func(m2, fname, rt)
3180 var vi: i64 = 0
3181 while vi < f.n_values {
3182 let v: *Value = val_at(f, vi)
3183 if (v as i64) != 0 { if v.kind == 1 {
3184 let pt: i64 = shl_ty(v.ty)
3185 if pt < 0 { sir_refuse(m2, "parameter type has no GLSL ES 3.00 counterpart -- refused rather than narrowed" as *u8); return 0 - 1 }
3186 var q: i64 = eb_put(nb, 0, WG_NUMBUF, "p" as *u8)
3187 q = eb_num(nb, q, WG_NUMBUF, v.param_index)
3188 nb[q] = 0 as u8
3189 sir_param(m2, fn2, nb, pt)
3190 } }
3191 vi = vi + 1
3192 }
3193 let e0: *BasicBlock = f.entry
3194 if (e0 as i64) == 0 { sir_refuse(m2, "function has no entry block" as *u8); return 0 - 1 }
3195
3196 // ---- RUNG 2b-GENERAL: STRUCTURED CONTROL-FLOW RECOVERY, INCLUDING LOOPS --------------
3197 // GLSL and WGSL have no goto, so a CFG must be rebuilt into nested statements. This is the
3198 // problem WebAssembly had to solve for the same reason, and the field's answers are the
3199 // Relooper (Zakai, Emscripten, OOPSLA 2011) and LLVM's WebAssembly Stackifier, which reduce
3200 // a GENERAL CFG by duplicating nodes or dispatching on a state variable.
3201 //
3202 // THE REFS DEBT THIS COMMENT USED TO DECLARE IS PAID. Both are now fetched, mirrored and
3203 // hash-pinned in graphics.refs (keys relooper-zakai11, llvm-cfgstackify), and every phrase
3204 // quoted here came from bytes that were read, not recalled. Ramsey's ICFP 2022 "Beyond
3205 // Relooper" is deliberately NOT cited: dl.acm.org answered 403 with a 5,470-byte bot
3206 // challenge page on four attempts, so it was never read and must not be leaned on.
3207 //
3208 // WHY A SHAPE REDUCER IS SUFFICIENT HERE -- MEASURED, NOT ASSUMED. nx_parse keeps locals
3209 // ALLOCA-BACKED (proven when this backend first refused op=43) and the opcode table carries
3210 // NO PHI, so values crossing blocks -- including across a BACK EDGE -- travel through
3211 // memory. There are no phi nodes to place, and phi placement is the hard half of relooping.
3212 // The measured corpus knowledge/nxsh_{clip,branch,loop,skin,layers}.nx is 1/3/4/6/7 blocks,
3213 // covering straight-line, if-diamond, while, loop-containing-if and nested loops, and EVERY
3214 // loop header in it has exactly TWO predecessors: one preheader, one latch. Single-entry and
3215 // reducible throughout, with no multi-entry region anywhere. That is precisely the
3216 // Stackifier's own stated precondition -- its markers are "sufficient to convert arbitrary
3217 // CFGs into a form that works on WebAssembly, provided that all loops are single-entry".
3218 //
3219 // So the Relooper's label variable (which exists to re-enter the right block of a MULTI-entry
3220 // region) is not needed here, and neither is node duplication. Both remain the right answer
3221 // for the general case, and the moment a header appears with n_preds != 2, shl_region REFUSES
3222 // IT BY NAME rather than mis-shaping it into something that compiles and computes the wrong
3223 // thing. The straight-line and if-diamond special cases that used to sit here are gone: they
3224 // are shapes of the reducer now, not rivals to it.
3225 // ONE recursive region reducer owns every shape decision AND every refusal. There is no
3226 // second ruler here to drift out of agreement with it.
3227 let mark: *i64 = sys_mmap(f.n_blocks*8 + 64) as *i64
3228 let wlq: *i64 = sys_mmap(f.n_blocks*8 + 64) as *i64
3229 let rbody: i64 = shl_region(m2, f, fn2, e0.id, 0 - 1, nb, msg, seen, mark, wlq, 0)
3230 if rbody == (0 - 1) { return 0 - 1 }
3231 sir_set(m2, fn2, SIR_B, rbody)
3232 return fn2
3233}
3234
3235func main(argc: i64, argv: *i64) -> i64 {
3236 // GE43 PATH: `nx_wgsl compute` -- the compute kernel through both backends (see wg_compute_demo).
3237 if argc > 1 { if wg_streq(argv[1] as *u8, "compute" as *u8) == 1 { return wg_compute_demo() } }
3238 // R-G PATH: `nx_wgsl world` -- the world shader through both backends, see wg_world_demo.
3239 if argc > 1 { if wg_streq(argv[1] as *u8, "world" as *u8) == 1 { return wg_world_demo() } }
3240 // R-F PATH: `nx_wgsl layout` -- the derived uniform layout rows and the stride neg-control, see wg_layout_demo.
3241 if argc > 1 { if wg_streq(argv[1] as *u8, "layout" as *u8) == 1 { return wg_layout_demo() } }
3242 // S12c PATH: `nx_wgsl cast [<glsl-out>]` -- the character cast vertex stage through both backends, see wg_cast_demo.
3243 if argc > 1 { if wg_streq(argv[1] as *u8, "cast" as *u8) == 1 { var co: *u8 = 0 as *u8; var cf: *u8 = 0 as *u8; if argc > 2 { co = argv[2] as *u8 } if argc > 3 { cf = argv[3] as *u8 } return wg_cast_demo(co, cf) } }
3244 // S1..S4 PATH: `nx_wgsl decls` -- the declaration probes (named refusals + positive control), see wg_decls_demo.
3245 if argc > 1 { if wg_streq(argv[1] as *u8, "decls" as *u8) == 1 { return wg_decls_demo() } }
3246 // RUNG 2 PATH: a real .nx source file, through the real front-end, out as both dialects.
3247 if argc > 1 {
3248 let path: *u8 = argv[1] as *u8
3249 let xb: *u8 = sys_mmap(SHL_EXPAND_CAP)
3250 let cx: *ExpandCtx = expand_ctx_new(xb, SHL_EXPAND_CAP)
3251 if expand_imports(cx, path) < 0 { wg_puts("nx_wgsl: expand_imports failed\n" as *u8); return 10 }
3252 let opz: *i64 = cx.out_pos
3253 let endp: i64 = *opz
3254 xb[endp] = 0 as u8
3255 let toks: *Tok = lex_source(xb, SHL_LEX_CAP)
3256 if toks == (0 as *Tok) { wg_puts("nx_wgsl: lex_source failed\n" as *u8); return 2 }
3257 let mm: *Module = parse_module(toks, 0 as *Module)
3258 if mm == (0 as *Module) { wg_puts("nx_wgsl: parse_module failed\n" as *u8); return 3 }
3259 if mm.n_functions <= 0 { wg_puts("nx_wgsl: no functions\n" as *u8); return 4 }
3260 var fi: i64 = 0
3261 while fi < mm.n_functions {
3262 let fb: i64 = mm.functions as i64
3263 let ff: *Function = (fb + fi * NX_MODULE_FN_STRIDE) as *Function
3264 opt_run(ff)
3265 fi = fi + 1
3266 }
3267 let f0: *Function = (mm.functions as i64) as *Function
3268 let m2: *i64 = sir_new()
3269 let fnid: i64 = shl_lower(m2, f0, "shader_main" as *u8)
3270 let nb2: *u8 = sys_mmap(WG_NUMBUF)
3271 var q2: i64 = 0
3272 wg_puts("=== NX-WGSL rung 2: an ORDINARY NishiLang function, through the REAL front-end ===\n" as *u8)
3273 wg_puts("ir_functions=" as *u8)
3274 q2 = eb_num(nb2, 0, WG_NUMBUF, mm.n_functions); nb2[q2] = 0 as u8; wg_puts(nb2)
3275 wg_puts(" ir_blocks=" as *u8)
3276 q2 = eb_num(nb2, 0, WG_NUMBUF, f0.n_blocks); nb2[q2] = 0 as u8; wg_puts(nb2)
3277 wg_puts(" ir_values=" as *u8)
3278 q2 = eb_num(nb2, 0, WG_NUMBUF, f0.n_values); nb2[q2] = 0 as u8; wg_puts(nb2)
3279 wg_puts(" ir_instrs=" as *u8)
3280 q2 = eb_num(nb2, 0, WG_NUMBUF, f0.n_instrs); nb2[q2] = 0 as u8; wg_puts(nb2)
3281 wg_puts("\n" as *u8)
3282 // CFG SHAPE, PRINTED BEFORE ANY LOWERING DECISION. Structured control-flow recovery is a
3283 // CFG-shape problem (rung 2b), and the shape must be MEASURED before an algorithm is
3284 // chosen for it -- picking a reducer for a graph nobody looked at is how a backend ends up
3285 // with a relooper it did not need or a pattern-matcher that cannot see the real cases.
3286 var bi: i64 = 0
3287 while bi < f0.n_blocks {
3288 let bbx: *BasicBlock = block_at(f0, bi)
3289 if (bbx as i64) != 0 {
3290 wg_puts("blk" as *u8)
3291 q2 = eb_num(nb2, 0, WG_NUMBUF, bbx.id); nb2[q2] = 0 as u8; wg_puts(nb2)
3292 wg_puts(" preds=" as *u8)
3293 q2 = eb_num(nb2, 0, WG_NUMBUF, bbx.n_preds); nb2[q2] = 0 as u8; wg_puts(nb2)
3294 wg_puts(" succs=" as *u8)
3295 q2 = eb_num(nb2, 0, WG_NUMBUF, bbx.n_succs); nb2[q2] = 0 as u8; wg_puts(nb2)
3296 if bbx.n_succs > 0 { if (bbx.succ0 as i64) != 0 {
3297 wg_puts(" s0=" as *u8)
3298 q2 = eb_num(nb2, 0, WG_NUMBUF, bbx.succ0.id); nb2[q2] = 0 as u8; wg_puts(nb2)
3299 } }
3300 if bbx.n_succs > 1 { if (bbx.succ1 as i64) != 0 {
3301 wg_puts(" s1=" as *u8)
3302 q2 = eb_num(nb2, 0, WG_NUMBUF, bbx.succ1.id); nb2[q2] = 0 as u8; wg_puts(nb2)
3303 } }
3304 var nin: i64 = 0
3305 var term: i64 = 0
3306 var ix: *Instr = bbx.head
3307 while (ix as i64) != 0 { nin = nin + 1; term = ix.op; ix = ix.next }
3308 wg_puts(" n=" as *u8)
3309 q2 = eb_num(nb2, 0, WG_NUMBUF, nin); nb2[q2] = 0 as u8; wg_puts(nb2)
3310 wg_puts(" term=" as *u8)
3311 q2 = eb_num(nb2, 0, WG_NUMBUF, term); nb2[q2] = 0 as u8; wg_puts(nb2)
3312 wg_puts("\n" as *u8)
3313 }
3314 bi = bi + 1
3315 }
3316 if fnid < 0 {
3317 wg_puts("lower_rc=-1 REFUSED: " as *u8)
3318 wg_puts(sir_cstr(m2, sir_refused(m2)))
3319 wg_puts("\n" as *u8)
3320 return 0
3321 }
3322 let gb2: *u8 = sys_mmap(WG_EMIT_CAP)
3323 let wb2: *u8 = sys_mmap(WG_EMIT_CAP)
3324 let gt2: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
3325 let wt2: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
3326 let gn2: *i64 = sys_mmap(WG_NUMBUF) as *i64
3327 let wn2: *i64 = sys_mmap(WG_NUMBUF) as *i64
3328 gn2[0] = 0
3329 wn2[0] = 0
3330 let gl2: i64 = glsl_emit_module(m2, STAGE_PLAIN, gb2, WG_EMIT_CAP, gt2, WG_TRACE_CAP, gn2)
3331 let wl2: i64 = wgsl_emit_module(m2, STAGE_PLAIN, wb2, WG_EMIT_CAP, wt2, WG_TRACE_CAP, wn2)
3332 wg_puts("\n--- GLSL (from NishiLang) ---\n" as *u8)
3333 if gl2 > 0 { sys_write(1, gb2, gl2) }
3334 wg_puts("\n--- WGSL (from the SAME NishiLang) ---\n" as *u8)
3335 if wl2 > 0 { sys_write(1, wb2, wl2) }
3336 var eq2: i64 = 1
3337 if gn2[0] != wn2[0] { eq2 = 0 }
3338 var i2: i64 = 0
3339 while i2 < gn2[0] {
3340 if i2 < wn2[0] { if gt2[i2] != wt2[i2] { eq2 = 0 } }
3341 i2 = i2 + 1
3342 }
3343 wg_puts("\nlower_rc=0 glsl_nodes=" as *u8)
3344 q2 = eb_num(nb2, 0, WG_NUMBUF, gn2[0]); nb2[q2] = 0 as u8; wg_puts(nb2)
3345 wg_puts(" wgsl_nodes=" as *u8)
3346 q2 = eb_num(nb2, 0, WG_NUMBUF, wn2[0]); nb2[q2] = 0 as u8; wg_puts(nb2)
3347 wg_puts(" trace_equal=" as *u8)
3348 q2 = eb_num(nb2, 0, WG_NUMBUF, eq2); nb2[q2] = 0 as u8; wg_puts(nb2)
3349 wg_puts("\n" as *u8)
3350 return 0
3351 }
3352 let m: *i64 = sir_new()
3353 shsrc_fullscreen_tri(m)
3354 let gb: *u8 = sys_mmap(WG_EMIT_CAP)
3355 let wb: *u8 = sys_mmap(WG_EMIT_CAP)
3356 let gt: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
3357 let wt: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
3358 let gn: *i64 = sys_mmap(WG_NUMBUF) as *i64
3359 let wn: *i64 = sys_mmap(WG_NUMBUF) as *i64
3360 gn[0] = 0
3361 wn[0] = 0
3362 let gl: i64 = glsl_emit_module(m, STAGE_VERTEX, gb, WG_EMIT_CAP, gt, WG_TRACE_CAP, gn)
3363 let wl: i64 = wgsl_emit_module(m, STAGE_VERTEX, wb, WG_EMIT_CAP, wt, WG_TRACE_CAP, wn)
3364 wg_puts("=== NX-WGSL rung 1: ONE NishiLang source, TWO dialects ===\n\n--- GLSL ---\n" as *u8)
3365 if gl > 0 { sys_write(1, gb, gl) }
3366 wg_puts("\n--- WGSL ---\n" as *u8)
3367 if wl > 0 { sys_write(1, wb, wl) }
3368 wg_puts("\nglsl_bytes=" as *u8)
3369 let nb: *u8 = sys_mmap(WG_NUMBUF)
3370 var q: i64 = 0
3371 q = eb_num(nb, 0, WG_NUMBUF, gl); nb[q] = 0 as u8; wg_puts(nb)
3372 wg_puts(" wgsl_bytes=" as *u8)
3373 q = eb_num(nb, 0, WG_NUMBUF, wl); nb[q] = 0 as u8; wg_puts(nb)
3374 wg_puts("\nglsl_nodes=" as *u8)
3375 q = eb_num(nb, 0, WG_NUMBUF, gn[0]); nb[q] = 0 as u8; wg_puts(nb)
3376 wg_puts("\nwgsl_nodes=" as *u8)
3377 q = eb_num(nb, 0, WG_NUMBUF, wn[0]); nb[q] = 0 as u8; wg_puts(nb)
3378 wg_puts("\n" as *u8)
3379 if sir_refused(m) != 0 { wg_puts("REFUSED: " as *u8); wg_puts(sir_cstr(m, sir_refused(m))); wg_puts("\n" as *u8) }
3380 // TRACE IDENTITY, ELEMENT-WISE. Two backends that consumed the SAME source COMPLETELY emit
3381 // the same node ids in the same order. This is the equivalence oracle available on a box with
3382 // no GPU, and it is NOT vacuous: corrupt_nodes below proves a dropped statement changes it.
3383 var eq: i64 = 1
3384 if gn[0] != wn[0] { eq = 0 }
3385 var i9: i64 = 0
3386 while i9 < gn[0] {
3387 if i9 < wn[0] { if gt[i9] != wt[i9] { eq = 0 } }
3388 i9 = i9 + 1
3389 }
3390 wg_puts("trace_equal=" as *u8)
3391 q = eb_num(nb, 0, WG_NUMBUF, eq); nb[q] = 0 as u8; wg_puts(nb)
3392 // NEG-CONTROL 1 -- A CONSTRUCT OUTSIDE THE COVERED SUBSET MUST REFUSE BY NAME. A backend that
3393 // guesses here compiles, links, runs, and draws the wrong picture forever.
3394 let m2: *i64 = sir_new()
3395 let f2: i64 = sir_func(m2, "main" as *u8, T_VOID)
3396 let bad: i64 = sir_ctor(m2, T_MAX + 89)
3397 sir_arg(m2, bad, sir_lit(m2, "0.0" as *u8, T_F32))
3398 sir_stmt(m2, f2, sir_var(m2, "x" as *u8, T_V2F, bad))
3399 let g2: *u8 = sys_mmap(WG_EMIT_CAP)
3400 let t2: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
3401 let n2: *i64 = sys_mmap(WG_NUMBUF) as *i64
3402 n2[0] = 0
3403 let r2: i64 = glsl_emit_module(m2, STAGE_VERTEX, g2, WG_EMIT_CAP, t2, WG_TRACE_CAP, n2)
3404 wg_puts("\nuncovered_rc=" as *u8)
3405 q = eb_num(nb, 0, WG_NUMBUF, r2); nb[q] = 0 as u8; wg_puts(nb)
3406 wg_puts(" uncovered_named=" as *u8)
3407 if sir_refused(m2) != 0 { wg_puts(sir_cstr(m2, sir_refused(m2))) } else { wg_puts("NONE" as *u8) }
3408 // NEG-CONTROL 2 -- A MODULE MISSING ONE STATEMENT MUST NOT TRACE THE SAME. This is what makes
3409 // trace_equal load-bearing rather than decorative.
3410 let m3: *i64 = sir_new()
3411 shsrc_fullscreen_tri(m3)
3412 let f3: i64 = m3[M_FUNCH]
3413 sir_set(m3, sir_get(m3, f3, SIR_B), SIR_NEXT, 0)
3414 let g3: *u8 = sys_mmap(WG_EMIT_CAP)
3415 let t3: *i64 = sys_mmap(WG_TRACE_CAP*8) as *i64
3416 let n3: *i64 = sys_mmap(WG_NUMBUF) as *i64
3417 n3[0] = 0
3418 glsl_emit_module(m3, STAGE_VERTEX, g3, WG_EMIT_CAP, t3, WG_TRACE_CAP, n3)
3419 wg_puts("\ncorrupt_nodes=" as *u8)
3420 q = eb_num(nb, 0, WG_NUMBUF, n3[0]); nb[q] = 0 as u8; wg_puts(nb)
3421 wg_puts("\n" as *u8)
3422 return 0
3423}