code wiki / (root) / nx_wgsl_compute_gate.nx

nx_wgsl_compute_gate.nx

buildroot/runtime/nx_wgsl_compute_gate.nx

10042 B170 linesdepth 5pulls 6 transitivereach 0 importersview sourcekind gate/prooftopic wgsl
docsdependenciesstructsconstsfunctions

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

nx_syscalls.nx nx_gate_verdict.nx nx_tool_run.nx nx_wgsl_compute_gate.nx

imports: nx_syscalls.nxnx_gate_verdict.nxnx_tool_run.nx

imported by: nobody (leaf or entry point)

structs

none

consts

26const WCG_CAP: i64 = 65536
27const WCG_TMO_MS: i64 = 20000
28const WCG_ARGV: i64 = 4
29const WCG_SUBJECT: *u8 = "./nx_wgsl.elf"
30const WCG_SUBJECT_OFFC: *u8 = "_offc/nx_wgsl.elf"
32const WCG_L_DEPTH: *u8 = "@group(0)@binding(1)var<storage,read_write> depth:array<atomic<u32>>;\n"
33const WCG_L_ZIN: *u8 = "@group(0)@binding(2)var<storage,read> zin:array<u32>;\n"
34const WCG_L_ENTRY: *u8 = "@compute @workgroup_size(64) fn main(@builtin(global_invocation_id) gid:vec3u){\n"
35const WCG_L_GID: *u8 = "var i:u32=gid.x;\n"
36const WCG_L_ATOMIC: *u8 = "atomicMin(&depth[i],zin[i]);\n"
37const WCG_L_CLOSE: *u8 = "}\n"
38const WCG_GL_RC: *u8 = "glsl_compute_rc=-1"
39const WCG_GL_NAMED: *u8 = "glsl_compute_named=glsl: compute stage"
40const WCG_GL_FRAG_RC: *u8 = "glsl_fragment_of_kernel_rc=-1"
41const WCG_GL_FRAG_NAMED: *u8 = "glsl_fragment_of_kernel_named=glsl: storage buffer declaration"
42const WCG_WG_CLEAN: *u8 = "wgsl_refused=NONE"
43const WCG_R1_TRACE: *u8 = "trace_equal=1"
44const WCG_R1_UNCOV: *u8 = "uncovered_rc=-1"
45const WCG_LEAK_GL: *u8 = "gl_"
46const WCG_LEAK_VERSION: *u8 = "version 300"

functions

54func wcg_slen(s: *u8) -> i64 { var n: i64 = 0; while s[n] != (0 as u8) { n = n + 1 } return n }
55func wcg_exists(path: *u8) -> i64 { let fd: i64 = sys_openat_rd(path); if fd < 0 { return 0 } sys_close(fd); return 1 }
58func wcg_find(buf: *u8, n: i64, needle: *u8) -> i64
71func wcg_has(buf: *u8, n: i64, needle: *u8) -> i64 { if wcg_find(buf, n, needle) >= 0 { return 1 } return 0 }
72func wcg_contract_len() -> i64
77func wcg_run(subject: *u8, verb: *u8, out: *u8, ol: *i64) -> i64
92func main(argc: i64, argv: *i64) -> i64