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