diff options
| author | skal <pascal.massimino@gmail.com> | 2026-03-26 07:03:01 +0100 |
|---|---|---|
| committer | skal <pascal.massimino@gmail.com> | 2026-03-26 07:03:01 +0100 |
| commit | 8f14bdd66cb002b2f89265b2a578ad93249089c9 (patch) | |
| tree | 2ccdb3939b673ebc3a5df429160631240239cee2 /cnn_v3/tools/shaders.js | |
| parent | 4ca498277b033ae10134045dae9c8c249a8d2b2b (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.js | 186 |
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; |
