nx_cast_affine_compute_t38.nx source
↩ module page · 4 lines · 2945 B
1// Private additive compute implementation for the existing cast shader owner.
2// Full AFM1 affine palettes + SKV1 variable influences; no DQ projection and no page-side deformation.
3const CS_AFFINE_COMPUTE_T38:*u8="struct Params{nv:u32,nj:u32,f0:u32,f1:u32,blend:f32,axis:u32,pad0:u32,pad1:u32}\nstruct Influence{joint:u32,weight:f32}\n@group(0) @binding(0) var<uniform> u:Params;\n@group(0) @binding(1) var<storage,read> source:array<f32>;\n@group(0) @binding(2) var<storage,read> offsets:array<u32>;\n@group(0) @binding(3) var<storage,read> influences:array<Influence>;\n@group(0) @binding(4) var<storage,read> animation:array<f32>;\n@group(0) @binding(5) var<storage,read_write> positions:array<f32>;\n@group(0) @binding(6) var<storage,read_write> normals:array<f32>;\n@group(0) @binding(7) var<storage,read_write> faults:array<atomic<u32>>;\nfn from_render(v:vec3<f32>)->vec3<f32>{\n if(u.axis==1u){return vec3<f32>(v.x,v.z,-v.y);}\n if(u.axis==2u){return vec3<f32>(v.z,v.y,-v.x);}\n return v;\n}\nfn to_render(v:vec3<f32>)->vec3<f32>{\n if(u.axis==1u){return vec3<f32>(v.x,-v.z,v.y);}\n if(u.axis==2u){return vec3<f32>(-v.z,v.y,v.x);}\n return v;\n}\n@compute @workgroup_size(64)\nfn main(@builtin(global_invocation_id) id:vec3<u32>){\n let v=id.x;if(v>=u.nv){return;}\n let start=offsets[v];let end=offsets[v+1u];var sum=0.0;\n for(var k=start;k<end;k=k+1u){sum=sum+influences[k].weight;}\n if(!(sum>0.0)){atomicAdd(&faults[0],1u);return;}\n var m:array<f32,12>;\n for(var k=start;k<end;k=k+1u){\n let influence=influences[k];let weight=influence.weight/sum;\n let a=(u.f0*u.nj+influence.joint)*12u;let b=(u.f1*u.nj+influence.joint)*12u;\n for(var c=0u;c<12u;c=c+1u){m[c]=m[c]+weight*mix(animation[a+c],animation[b+c],u.blend);}\n }\n let A=m[4]*m[8]-m[7]*m[5];let B=m[7]*m[2]-m[1]*m[8];let C=m[1]*m[5]-m[4]*m[2];\n let D=m[6]*m[5]-m[3]*m[8];let E=m[0]*m[8]-m[6]*m[2];let F=m[3]*m[2]-m[0]*m[5];\n let G=m[3]*m[7]-m[6]*m[4];let H=m[6]*m[1]-m[0]*m[7];let I=m[0]*m[4]-m[3]*m[1];\n let determinant=m[0]*A+m[3]*B+m[6]*C;\n if(determinant==0.0){atomicAdd(&faults[1],1u);return;}\n let at=v*6u;let p=from_render(vec3<f32>(source[at],source[at+1u],source[at+2u]));\n let n=from_render(vec3<f32>(source[at+3u],source[at+4u],source[at+5u]));\n let moved=to_render(vec3<f32>(m[0]*p.x+m[3]*p.y+m[6]*p.z+m[9],m[1]*p.x+m[4]*p.y+m[7]*p.z+m[10],m[2]*p.x+m[5]*p.y+m[8]*p.z+m[11]));\n let inverse_normal=vec3<f32>(A*n.x+B*n.y+C*n.z,D*n.x+E*n.y+F*n.z,G*n.x+H*n.y+I*n.z)/determinant;\n let magnitude=length(inverse_normal);\n if(!(magnitude>0.0)||!(magnitude<=3.402823466e38)||!all(abs(moved)<=vec3<f32>(3.402823466e38))){atomicAdd(&faults[2],1u);return;}\n let transformed=to_render(inverse_normal/magnitude);\n positions[v*3u]=moved.x;positions[v*3u+1u]=moved.y;positions[v*3u+2u]=moved.z;\n normals[v*3u]=transformed.x;normals[v*3u+1u]=transformed.y;normals[v*3u+2u]=transformed.z;\n}\n"
4func cs_affine_compute_src()->*u8{return CS_AFFINE_COMPUTE_T38}