nx_wgsl_sky_review_t154.nx source

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