nx_wgsl_compute_gate.nx
buildroot/runtime/nx_wgsl_compute_gate.nx
about
nx_wgsl_compute_gate.nx -- GE43 REFEREE: the compute stage of the ONE NishiLang shader backend
emits EXACTLY the WGSL the kernel contract names, and the GLSL side REFUSES it BY NAME.
SUBJECT: the DEPLOYED nx_wgsl.elf (end-to-end -- this gate forks the artifact a browser page
would be built from, not the library). `nx_wgsl compute` prints the depth-min kernel through both
backends; every tooth below is a KAT over those bytes. A KAT is the right shape here: a shader
backend that drifts by one character produces a program that compiles and draws the wrong
picture, so "contains the right idea" is not a bar -- the bytes are.
WHY THE GLSL REFUSAL IS LOAD-BEARING: WebGL2 has no compute stage, no storage buffers and no
atomics. A backend that silently dropped the storage declarations would emit a fragment program
that compiles, links, runs and draws nothing forever. Two teeth pin the refusal (the STAGE refusal
on a compute emit, the DECLARATION refusal when the same kernel is asked for as a fragment stage),
so the refusal cannot degrade into a later accident and still read GREEN.
CONTROLS: the rung-1 contract (fullscreen triangle, trace_equal=1, uncovered_rc=-1) must still
hold on the same binary -- a compute edit that broke rung 1 must read RED here, not on some other
board weeks later. neg-control-* teeth assert refusals and absences.
nx_wgsl_compute_gate [subject-elf] default ./nx_wgsl.elf, then _offc/nx_wgsl.elf
license_tier: ORIGINAL No hw writes (Rule 26).
dependencies 3 imports · 0 importers
imports: nx_syscalls.nxnx_gate_verdict.nxnx_tool_run.nx
imported by: nobody (leaf or entry point)
structs
| none |
consts
| 26 | const WCG_CAP: i64 = 65536 |
| 27 | const WCG_TMO_MS: i64 = 20000 |
| 28 | const WCG_ARGV: i64 = 4 |
| 29 | const WCG_SUBJECT: *u8 = "./nx_wgsl.elf" |
| 30 | const WCG_SUBJECT_OFFC: *u8 = "_offc/nx_wgsl.elf" |
| 32 | const WCG_L_DEPTH: *u8 = "@group(0)@binding(1)var<storage,read_write> depth:array<atomic<u32>>;\n" |
| 33 | const WCG_L_ZIN: *u8 = "@group(0)@binding(2)var<storage,read> zin:array<u32>;\n" |
| 34 | const WCG_L_ENTRY: *u8 = "@compute @workgroup_size(64) fn main(@builtin(global_invocation_id) gid:vec3u){\n" |
| 35 | const WCG_L_GID: *u8 = "var i:u32=gid.x;\n" |
| 36 | const WCG_L_ATOMIC: *u8 = "atomicMin(&depth[i],zin[i]);\n" |
| 37 | const WCG_L_CLOSE: *u8 = "}\n" |
| 38 | const WCG_GL_RC: *u8 = "glsl_compute_rc=-1" |
| 39 | const WCG_GL_NAMED: *u8 = "glsl_compute_named=glsl: compute stage" |
| 40 | const WCG_GL_FRAG_RC: *u8 = "glsl_fragment_of_kernel_rc=-1" |
| 41 | const WCG_GL_FRAG_NAMED: *u8 = "glsl_fragment_of_kernel_named=glsl: storage buffer declaration" |
| 42 | const WCG_WG_CLEAN: *u8 = "wgsl_refused=NONE" |
| 43 | const WCG_R1_TRACE: *u8 = "trace_equal=1" |
| 44 | const WCG_R1_UNCOV: *u8 = "uncovered_rc=-1" |
| 45 | const WCG_LEAK_GL: *u8 = "gl_" |
| 46 | const WCG_LEAK_VERSION: *u8 = "version 300" |
functions
| 54 | func wcg_slen(s: *u8) -> i64 { var n: i64 = 0; while s[n] != (0 as u8) { n = n + 1 } return n } |
| 55 | func wcg_exists(path: *u8) -> i64 { let fd: i64 = sys_openat_rd(path); if fd < 0 { return 0 } sys_close(fd); return 1 } |
| 58 | func wcg_find(buf: *u8, n: i64, needle: *u8) -> i64 |
| 71 | func wcg_has(buf: *u8, n: i64, needle: *u8) -> i64 { if wcg_find(buf, n, needle) >= 0 { return 1 } return 0 } |
| 72 | func wcg_contract_len() -> i64 |
| 77 | func wcg_run(subject: *u8, verb: *u8, out: *u8, ol: *i64) -> i64 |
| 92 | func main(argc: i64, argv: *i64) -> i64 |