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