summaryrefslogtreecommitdiff
path: root/cnn_v3/tools/shaders.js
diff options
context:
space:
mode:
authorskal <pascal.massimino@gmail.com>2026-03-26 07:03:01 +0100
committerskal <pascal.massimino@gmail.com>2026-03-26 07:03:01 +0100
commit8f14bdd66cb002b2f89265b2a578ad93249089c9 (patch)
tree2ccdb3939b673ebc3a5df429160631240239cee2 /cnn_v3/tools/shaders.js
parent4ca498277b033ae10134045dae9c8c249a8d2b2b (diff)
feat(cnn_v3): upgrade architecture to enc_channels=[8,16]
Double encoder capacity: enc0 4→8ch, enc1 8→16ch, bottleneck 16→16ch, dec1 32→8ch, dec0 16→4ch. Total weights 2476→7828 f16 (~15.3 KB). FiLM MLP output 40→72 params (L1: 16×40→16×72). 16-ch textures split into _lo/_hi rgba32uint pairs (enc1, bottleneck). enc0 and dec1 textures changed from rgba16float to rgba32uint (8ch). GBUF_RGBA32UINT node gains CopySrc for parity test readback. - WGSL shaders: all 5 passes rewritten for new channel counts - C++ CNNv3Effect: new weight offsets/sizes, 8ch uniform structs - Web tool (shaders.js + tester.js): matching texture formats and bindings - Parity test: readback_rgba32uint_8ch helper, updated vector counts - Training scripts: default enc_channels=[8,16], updated docstrings - Docs + architecture PNG regenerated handoff(Gemini): CNN v3 [8,16] upgrade complete. All code, tests, web tool, training scripts, and docs updated. Next: run training pass.
Diffstat (limited to 'cnn_v3/tools/shaders.js')
-rw-r--r--cnn_v3/tools/shaders.js186
1 files changed, 106 insertions, 80 deletions
diff --git a/cnn_v3/tools/shaders.js b/cnn_v3/tools/shaders.js
index 36f53c8..6f2176d 100644
--- a/cnn_v3/tools/shaders.js
+++ b/cnn_v3/tools/shaders.js
@@ -1,10 +1,10 @@
'use strict';
// CNN v3 WGSL shaders — matches cnn_v3/shaders/*.wgsl exactly.
-// Weight offsets (f16 index): enc0=0, enc1=724, bn=1020, dec1=1604, dec0=2184, total=2476
-// BN is now Conv(8→8, 3×3, dilation=2): 8*8*9+8=584 weights (was 72 for 1×1)
+// Architecture: enc_channels=[8,16]
+// Weight offsets (f16 index): enc0=0, enc1=1448, bn=2616, dec1=4936, dec0=7248, total=7828
-const ENC0_OFF=0, ENC1_OFF=724, BN_OFF=1020, DEC1_OFF=1604, DEC0_OFF=2184;
-const TOTAL_F16=2476, TOTAL_U32=1238;
+const ENC0_OFF=0, ENC1_OFF=1448, BN_OFF=2616, DEC1_OFF=4936, DEC0_OFF=7248;
+const TOTAL_F16=7828, TOTAL_U32=3914;
// Inlined helpers — prepended to shaders that need them.
const H = `
@@ -41,20 +41,24 @@ fn main(@builtin(global_invocation_id) id:vec3u){
pack4x8unorm(vec4f(m2.g,m2.b,1.,tr)),0u));
}`;
-// Enc0: Conv(20→4, 3×3, zero-pad) + FiLM + ReLU → rgba16float
-// Params (48 bytes): weight_offset u32 _pad×3 gamma vec4f beta vec4f
+// Enc0: Conv(20→8, 3×3, zero-pad) + FiLM + ReLU → rgba32uint (pack2x16float, 8ch)
+// Params (80 bytes): wo u32 _pad×3 gl gh bl bh (vec4f×4)
const ENC0_SHADER=H+`
-struct P{wo:u32,_a:u32,_b:u32,_c:u32,g:vec4f,b:vec4f}
+struct P{wo:u32,_a:u32,_b:u32,_c:u32,gl:vec4f,gh:vec4f,bl:vec4f,bh:vec4f}
@group(0) @binding(0) var t0:texture_2d<u32>;
@group(0) @binding(1) var t1:texture_2d<u32>;
@group(0) @binding(2) var<storage,read> weights:array<u32>;
@group(0) @binding(3) var<uniform> p:P;
-@group(0) @binding(4) var out:texture_storage_2d<rgba16float,write>;
+@group(0) @binding(4) var out:texture_storage_2d<rgba32uint,write>;
+fn fg(o:u32)->f32{if(o<4u){return p.gl[o];}return p.gh[o-4u];}
+fn fb(o:u32)->f32{if(o<4u){return p.bl[o];}return p.bh[o-4u];}
fn feat(c:vec2i,d:vec2i)->array<f32,20>{
if(c.x<0||c.y<0||c.x>=d.x||c.y>=d.y){return array<f32,20>(0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.);}
- let a=unpack2x16float(textureLoad(t0,c,0).x); let b=unpack2x16float(textureLoad(t0,c,0).y);
- let cc=unpack2x16float(textureLoad(t0,c,0).z);let dd=unpack2x16float(textureLoad(t0,c,0).w);
- let e=unpack4x8unorm(textureLoad(t1,c,0).x); let f=unpack4x8unorm(textureLoad(t1,c,0).y);
+ let t=textureLoad(t0,c,0);
+ let a=unpack2x16float(t.x);let b=unpack2x16float(t.y);
+ let cc=unpack2x16float(t.z);let dd=unpack2x16float(t.w);
+ let e=unpack4x8unorm(textureLoad(t1,c,0).x);
+ let f=unpack4x8unorm(textureLoad(t1,c,0).y);
let g=unpack4x8unorm(textureLoad(t1,c,0).z);
return array<f32,20>(a.x,a.y,b.x,b.y,cc.x,cc.y,dd.x,dd.y,e.x,e.y,e.z,e.w,f.x,f.y,f.z,f.w,g.x,g.y,g.z,g.w);
}
@@ -62,41 +66,50 @@ fn feat(c:vec2i,d:vec2i)->array<f32,20>{
fn main(@builtin(global_invocation_id) id:vec3u){
let c=vec2i(id.xy); let d=vec2i(textureDimensions(t0));
if(c.x>=d.x||c.y>=d.y){return;}
- const IN:u32=20u; const OUT:u32=4u;
- var o:array<f32,4>;
+ const IN:u32=20u; const OUT:u32=8u;
+ var o:array<f32,8>;
for(var oc:u32=0u;oc<OUT;oc++){
var s=get_w(p.wo,OUT*IN*9u+oc);
for(var ky:i32=-1;ky<=1;ky++){for(var kx:i32=-1;kx<=1;kx++){
let ft=feat(c+vec2i(kx,ky),d); let ki=u32(ky+1)*3u+u32(kx+1);
for(var i:u32=0u;i<IN;i++){s+=get_w(p.wo,oc*IN*9u+i*9u+ki)*ft[i];}
}}
- o[oc]=max(0.,p.g[oc]*s+p.b[oc]);
+ o[oc]=max(0.,fg(oc)*s+fb(oc));
}
- textureStore(out,c,vec4f(o[0],o[1],o[2],o[3]));
+ textureStore(out,c,vec4u(pack2x16float(vec2f(o[0],o[1])),pack2x16float(vec2f(o[2],o[3])),
+ pack2x16float(vec2f(o[4],o[5])),pack2x16float(vec2f(o[6],o[7]))));
}`;
-// Enc1: AvgPool(enc0) + Conv(4→8, 3×3) + FiLM + ReLU → rgba32uint half-res
-// Params (80 bytes): wo u32 _pad×3 glo ghi blo bhi vec4f×4
+// Enc1: AvgPool(enc0) + Conv(8→16, 3×3) + FiLM + ReLU → 2× rgba32uint half-res (lo ch0-7, hi ch8-15)
+// Params (144 bytes): wo u32 _pad×3 g0 g1 g2 g3 b0 b1 b2 b3 (vec4f×8)
const ENC1_SHADER=H+`
-struct P{wo:u32,_a:u32,_b:u32,_c:u32,gl:vec4f,gh:vec4f,bl:vec4f,bh:vec4f}
-@group(0) @binding(0) var e0:texture_2d<f32>;
+struct P{wo:u32,_a:u32,_b:u32,_c:u32,g0:vec4f,g1:vec4f,g2:vec4f,g3:vec4f,b0:vec4f,b1:vec4f,b2:vec4f,b3:vec4f}
+@group(0) @binding(0) var e0:texture_2d<u32>;
@group(0) @binding(1) var<storage,read> weights:array<u32>;
@group(0) @binding(2) var<uniform> p:P;
-@group(0) @binding(3) var out:texture_storage_2d<rgba32uint,write>;
-fn fg(o:u32)->f32{if(o<4u){return p.gl[o];}return p.gh[o-4u];}
-fn fb(o:u32)->f32{if(o<4u){return p.bl[o];}return p.bh[o-4u];}
-fn avg(hc:vec2i,fd:vec2i)->array<f32,4>{
- let hd=fd/2; if(hc.x<0||hc.y<0||hc.x>=hd.x||hc.y>=hd.y){return array<f32,4>(0.,0.,0.,0.);}
- var s=vec4f(0.);
- for(var y:i32=0;y<2;y++){for(var x:i32=0;x<2;x++){s+=textureLoad(e0,clamp(hc*2+vec2i(x,y),vec2i(0),fd-vec2i(1)),0);}}
- let a=s*.25; return array<f32,4>(a.x,a.y,a.z,a.w);
+@group(0) @binding(3) var olo:texture_storage_2d<rgba32uint,write>;
+@group(0) @binding(4) var ohi:texture_storage_2d<rgba32uint,write>;
+fn fg(o:u32)->f32{
+ if(o<4u){return p.g0[o];}if(o<8u){return p.g1[o-4u];}
+ if(o<12u){return p.g2[o-8u];}return p.g3[o-12u];}
+fn fb(o:u32)->f32{
+ if(o<4u){return p.b0[o];}if(o<8u){return p.b1[o-4u];}
+ if(o<12u){return p.b2[o-8u];}return p.b3[o-12u];}
+fn avg(hc:vec2i,fd:vec2i)->array<f32,8>{
+ let hd=fd/2; if(hc.x<0||hc.y<0||hc.x>=hd.x||hc.y>=hd.y){return array<f32,8>(0.,0.,0.,0.,0.,0.,0.,0.);}
+ var s:array<f32,8>;
+ for(var y:i32=0;y<2;y++){for(var x:i32=0;x<2;x++){
+ let f=unpack8(e0,clamp(hc*2+vec2i(x,y),vec2i(0),fd-vec2i(1)));
+ for(var i:u32=0u;i<8u;i++){s[i]+=f[i];}
+ }}
+ for(var i:u32=0u;i<8u;i++){s[i]*=.25;} return s;
}
@compute @workgroup_size(8,8)
fn main(@builtin(global_invocation_id) id:vec3u){
let fd=vec2i(textureDimensions(e0)); let hd=fd/2; let c=vec2i(id.xy);
if(c.x>=hd.x||c.y>=hd.y){return;}
- const IN:u32=4u; const OUT:u32=8u;
- var o:array<f32,8>;
+ const IN:u32=8u; const OUT:u32=16u;
+ var o:array<f32,16>;
for(var oc:u32=0u;oc<OUT;oc++){
var s=get_w(p.wo,OUT*IN*9u+oc);
for(var ky:i32=-1;ky<=1;ky++){for(var kx:i32=-1;kx<=1;kx++){
@@ -105,97 +118,114 @@ fn main(@builtin(global_invocation_id) id:vec3u){
}}
o[oc]=max(0.,fg(oc)*s+fb(oc));
}
- textureStore(out,c,vec4u(pack2x16float(vec2f(o[0],o[1])),pack2x16float(vec2f(o[2],o[3])),
+ textureStore(olo,c,vec4u(pack2x16float(vec2f(o[0],o[1])),pack2x16float(vec2f(o[2],o[3])),
pack2x16float(vec2f(o[4],o[5])),pack2x16float(vec2f(o[6],o[7]))));
+ textureStore(ohi,c,vec4u(pack2x16float(vec2f(o[8],o[9])),pack2x16float(vec2f(o[10],o[11])),
+ pack2x16float(vec2f(o[12],o[13])),pack2x16float(vec2f(o[14],o[15]))));
}`;
-// Bottleneck: AvgPool(enc1) + Conv(8→8, 3×3, dilation=2) + ReLU → rgba32uint quarter-res (no FiLM)
+// Bottleneck: AvgPool(enc1) + Conv(16→16, 3×3, dil=2) + ReLU → 2× rgba32uint quarter-res (no FiLM)
// Params (16 bytes): wo u32 _pad×3
const BN_SHADER=H+`
struct P{wo:u32,_a:u32,_b:u32,_c:u32}
-@group(0) @binding(0) var e1:texture_2d<u32>;
-@group(0) @binding(1) var<storage,read> weights:array<u32>;
-@group(0) @binding(2) var<uniform> p:P;
-@group(0) @binding(3) var out:texture_storage_2d<rgba32uint,write>;
-fn avg(qc:vec2i,hd:vec2i)->array<f32,8>{
- let qd=hd/2; if(qc.x<0||qc.y<0||qc.x>=qd.x||qc.y>=qd.y){return array<f32,8>(0.,0.,0.,0.,0.,0.,0.,0.);}
- var s:array<f32,8>;
+@group(0) @binding(0) var elo:texture_2d<u32>;
+@group(0) @binding(1) var ehi:texture_2d<u32>;
+@group(0) @binding(2) var<storage,read> weights:array<u32>;
+@group(0) @binding(3) var<uniform> p:P;
+@group(0) @binding(4) var olo:texture_storage_2d<rgba32uint,write>;
+@group(0) @binding(5) var ohi:texture_storage_2d<rgba32uint,write>;
+fn avg(qc:vec2i,hd:vec2i)->array<f32,16>{
+ let qd=hd/2; if(qc.x<0||qc.y<0||qc.x>=qd.x||qc.y>=qd.y){return array<f32,16>(0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.);}
+ var s:array<f32,16>;
for(var y:i32=0;y<2;y++){for(var x:i32=0;x<2;x++){
- let f=unpack8(e1,clamp(qc*2+vec2i(x,y),vec2i(0),hd-vec2i(1)));
- for(var i:u32=0u;i<8u;i++){s[i]+=f[i];}
+ let c2=clamp(qc*2+vec2i(x,y),vec2i(0),hd-vec2i(1));
+ let lo=unpack8(elo,c2);let hi=unpack8(ehi,c2);
+ for(var i:u32=0u;i<8u;i++){s[i]+=lo[i];s[i+8u]+=hi[i];}
}}
- for(var i:u32=0u;i<8u;i++){s[i]*=.25;} return s;
+ for(var i:u32=0u;i<16u;i++){s[i]*=.25;} return s;
}
@compute @workgroup_size(8,8)
fn main(@builtin(global_invocation_id) id:vec3u){
- let hd=vec2i(textureDimensions(e1)); let qd=hd/2; let c=vec2i(id.xy);
+ let hd=vec2i(textureDimensions(elo)); let qd=hd/2; let c=vec2i(id.xy);
if(c.x>=qd.x||c.y>=qd.y){return;}
- var o:array<f32,8>;
- for(var oc:u32=0u;oc<8u;oc++){
- var s=get_w(p.wo,576u+oc);
+ var o:array<f32,16>;
+ for(var oc:u32=0u;oc<16u;oc++){
+ var s=get_w(p.wo,2304u+oc);
for(var ky:i32=-1;ky<=1;ky++){for(var kx:i32=-1;kx<=1;kx++){
let ft=avg(c+vec2i(kx,ky)*2,hd); let ki=u32(ky+1)*3u+u32(kx+1);
- for(var i:u32=0u;i<8u;i++){s+=get_w(p.wo,oc*72u+i*9u+ki)*ft[i];}
+ for(var i:u32=0u;i<16u;i++){s+=get_w(p.wo,oc*144u+i*9u+ki)*ft[i];}
}}
o[oc]=max(0.,s);
}
- textureStore(out,c,vec4u(pack2x16float(vec2f(o[0],o[1])),pack2x16float(vec2f(o[2],o[3])),
+ textureStore(olo,c,vec4u(pack2x16float(vec2f(o[0],o[1])),pack2x16float(vec2f(o[2],o[3])),
pack2x16float(vec2f(o[4],o[5])),pack2x16float(vec2f(o[6],o[7]))));
+ textureStore(ohi,c,vec4u(pack2x16float(vec2f(o[8],o[9])),pack2x16float(vec2f(o[10],o[11])),
+ pack2x16float(vec2f(o[12],o[13])),pack2x16float(vec2f(o[14],o[15]))));
}`;
-// Dec1: NearestUp(bn)+cat(enc1_skip) → Conv(16→4,3×3) + FiLM + ReLU → rgba16float half-res
-// Params (48 bytes): same layout as enc0
+// Dec1: NearestUp(bn_lo/hi)+cat(enc1_lo/hi) → Conv(32→8,3×3) + FiLM + ReLU → rgba32uint half-res
+// Params (80 bytes): wo u32 _pad×3 gl gh bl bh (vec4f×4)
const DEC1_SHADER=H+`
-struct P{wo:u32,_a:u32,_b:u32,_c:u32,g:vec4f,b:vec4f}
-@group(0) @binding(0) var bn:texture_2d<u32>;
-@group(0) @binding(1) var e1:texture_2d<u32>;
-@group(0) @binding(2) var<storage,read> weights:array<u32>;
-@group(0) @binding(3) var<uniform> p:P;
-@group(0) @binding(4) var out:texture_storage_2d<rgba16float,write>;
-fn cat(hc:vec2i,hd:vec2i)->array<f32,16>{
- if(hc.x<0||hc.y<0||hc.x>=hd.x||hc.y>=hd.y){return array<f32,16>(0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.);}
- let qd=hd/2; let b=unpack8(bn,clamp(hc/2,vec2i(0),qd-vec2i(1)));
- let s=unpack8(e1,hc);
- return array<f32,16>(b[0],b[1],b[2],b[3],b[4],b[5],b[6],b[7],s[0],s[1],s[2],s[3],s[4],s[5],s[6],s[7]);
+struct P{wo:u32,_a:u32,_b:u32,_c:u32,gl:vec4f,gh:vec4f,bl:vec4f,bh:vec4f}
+@group(0) @binding(0) var bnlo:texture_2d<u32>;
+@group(0) @binding(1) var bnhi:texture_2d<u32>;
+@group(0) @binding(2) var e1lo:texture_2d<u32>;
+@group(0) @binding(3) var e1hi:texture_2d<u32>;
+@group(0) @binding(4) var<storage,read> weights:array<u32>;
+@group(0) @binding(5) var<uniform> p:P;
+@group(0) @binding(6) var out:texture_storage_2d<rgba32uint,write>;
+fn fg(o:u32)->f32{if(o<4u){return p.gl[o];}return p.gh[o-4u];}
+fn fb(o:u32)->f32{if(o<4u){return p.bl[o];}return p.bh[o-4u];}
+fn cat(hc:vec2i,hd:vec2i)->array<f32,32>{
+ if(hc.x<0||hc.y<0||hc.x>=hd.x||hc.y>=hd.y){return array<f32,32>(0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.);}
+ let qd=hd/2; let q=clamp(hc/2,vec2i(0),qd-vec2i(1));
+ let b0=unpack8(bnlo,q);let b1=unpack8(bnhi,q);
+ let s0=unpack8(e1lo,hc);let s1=unpack8(e1hi,hc);
+ return array<f32,32>(b0[0],b0[1],b0[2],b0[3],b0[4],b0[5],b0[6],b0[7],
+ b1[0],b1[1],b1[2],b1[3],b1[4],b1[5],b1[6],b1[7],
+ s0[0],s0[1],s0[2],s0[3],s0[4],s0[5],s0[6],s0[7],
+ s1[0],s1[1],s1[2],s1[3],s1[4],s1[5],s1[6],s1[7]);
}
@compute @workgroup_size(8,8)
fn main(@builtin(global_invocation_id) id:vec3u){
- let hd=vec2i(textureDimensions(e1)); let c=vec2i(id.xy);
+ let hd=vec2i(textureDimensions(e1lo)); let c=vec2i(id.xy);
if(c.x>=hd.x||c.y>=hd.y){return;}
- const IN:u32=16u; const OUT:u32=4u;
- var o:array<f32,4>;
+ const IN:u32=32u; const OUT:u32=8u;
+ var o:array<f32,8>;
for(var oc:u32=0u;oc<OUT;oc++){
var s=get_w(p.wo,OUT*IN*9u+oc);
for(var ky:i32=-1;ky<=1;ky++){for(var kx:i32=-1;kx<=1;kx++){
let ft=cat(c+vec2i(kx,ky),hd); let ki=u32(ky+1)*3u+u32(kx+1);
for(var i:u32=0u;i<IN;i++){s+=get_w(p.wo,oc*IN*9u+i*9u+ki)*ft[i];}
}}
- o[oc]=max(0.,p.g[oc]*s+p.b[oc]);
+ o[oc]=max(0.,fg(oc)*s+fb(oc));
}
- textureStore(out,c,vec4f(o[0],o[1],o[2],o[3]));
+ textureStore(out,c,vec4u(pack2x16float(vec2f(o[0],o[1])),pack2x16float(vec2f(o[2],o[3])),
+ pack2x16float(vec2f(o[4],o[5])),pack2x16float(vec2f(o[6],o[7]))));
}`;
-// Dec0: NearestUp(dec1)+cat(enc0_skip) → Conv(8→4,3×3) + FiLM + ReLU + Sigmoid → rgba16float
-// Params (48 bytes): same layout as enc0
+// Dec0: NearestUp(dec1)+cat(enc0_skip) → Conv(16→4,3×3) + FiLM + ReLU + Sigmoid → rgba16float
+// Params (48 bytes): wo u32 _pad×3 g vec4f b vec4f
const DEC0_SHADER=H+`
struct P{wo:u32,_a:u32,_b:u32,_c:u32,g:vec4f,b:vec4f}
-@group(0) @binding(0) var d1:texture_2d<f32>;
-@group(0) @binding(1) var e0:texture_2d<f32>;
+@group(0) @binding(0) var d1:texture_2d<u32>;
+@group(0) @binding(1) var e0:texture_2d<u32>;
@group(0) @binding(2) var<storage,read> weights:array<u32>;
@group(0) @binding(3) var<uniform> p:P;
@group(0) @binding(4) var out:texture_storage_2d<rgba16float,write>;
-fn cat(c:vec2i,fd:vec2i)->array<f32,8>{
- if(c.x<0||c.y<0||c.x>=fd.x||c.y>=fd.y){return array<f32,8>(0.,0.,0.,0.,0.,0.,0.,0.);}
+fn cat(c:vec2i,fd:vec2i)->array<f32,16>{
+ if(c.x<0||c.y<0||c.x>=fd.x||c.y>=fd.y){return array<f32,16>(0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.,0.);}
let hd=vec2i(textureDimensions(d1));
- let a=textureLoad(d1,clamp(c/2,vec2i(0),hd-vec2i(1)),0);
- let b=textureLoad(e0,c,0);
- return array<f32,8>(a.x,a.y,a.z,a.w,b.x,b.y,b.z,b.w);
+ let a=unpack8(d1,clamp(c/2,vec2i(0),hd-vec2i(1)));
+ let b=unpack8(e0,c);
+ return array<f32,16>(a[0],a[1],a[2],a[3],a[4],a[5],a[6],a[7],
+ b[0],b[1],b[2],b[3],b[4],b[5],b[6],b[7]);
}
@compute @workgroup_size(8,8)
fn main(@builtin(global_invocation_id) id:vec3u){
let fd=vec2i(textureDimensions(e0)); let c=vec2i(id.xy);
if(c.x>=fd.x||c.y>=fd.y){return;}
- const IN:u32=8u; const OUT:u32=4u;
+ const IN:u32=16u; const OUT:u32=4u;
var o:array<f32,4>;
for(var oc:u32=0u;oc<OUT;oc++){
var s=get_w(p.wo,OUT*IN*9u+oc);
@@ -227,8 +257,6 @@ const DISP_SHADER=`
}`;
// Viz f32: show one channel of rgba16float layer
-// Uniform layout: ch(u32) _p(u32) ox(i32) oy(i32) — 16 bytes
-// ox/oy = texel offset (top-left of view); 0,0 for full-texture vignettes.
const VIZ_F32=`
struct Vu{ch:u32,_p:u32,ox:i32,oy:i32}
@group(0) @binding(0) var t:texture_2d<f32>;
@@ -265,8 +293,6 @@ struct Vu{ch:u32,_p:u32,ox:i32,oy:i32}
// Full G-buffer pack: assembles feat_tex0/feat_tex1 from individual G-buffer images.
// Bindings: albedo(0) normal(1) depth(2) matid(3) shadow(4) transp(5) f0(6) f1(7)
-// All source textures are rgba8unorm (browser-loaded images, R channel for depth/matid/shadow/transp).
-// Uses textureLoad() only (no sampler needed). Matches gbuf_pack.wgsl packing exactly.
const FULL_PACK_SHADER=`
@group(0) @binding(0) var albedo: texture_2d<f32>;
@group(0) @binding(1) var normal: texture_2d<f32>;
@@ -295,7 +321,7 @@ fn main(@builtin(global_invocation_id) id:vec3u){
if(c.x>=d.x||c.y>=d.y){return;}
let alb=textureLoad(albedo,c,0).rgb;
let nrm=textureLoad(normal,c,0).rg;
- let oct=nrm*2.-vec2f(1.); // [0,1] -> [-1,1]
+ let oct=nrm*2.-vec2f(1.);
let dv=ld(c,d);
let dzdx=(ld(c+vec2i(1,0),d)-ld(c-vec2i(1,0),d))*.5;
let dzdy=(ld(c+vec2i(0,1),d)-ld(c-vec2i(0,1),d))*.5;