nx_game_surface_prepare_t132.nx source
↩ module page · 22 lines · 17740 B
1func gpe_surface_sources_t132(cast_path:*u8,cast_sha:*u8,adapter_path:*u8,adapter_sha:*u8)->i64{
2 return gsb_artifact_matches(cast_path,cast_sha,128275)&gsb_artifact_matches(adapter_path,adapter_sha,7728)
3}
4// Generated capture closure: changes to either deformation source refuse emission until requalified.
5func gpe_surface_runtime_t132(out:*u8,pos:i64)->i64{
6 if gpe_surface_sources_t132("buildroot/runtime/nx_cast_shader_src.nx","1f11de1c8fe2c2f0a5d7c6c2a705873081754ab48d739c6d5f69215cda565ba9","buildroot/runtime/nx_cast_surface_normal_t133.nx","b93fc53053eadbee8d6bee38ded29532e5dae9471bd1fbbe695357bf4c9494f9")!=1{gsb_failed=1;return 0}
7 let compute:*u8="struct U{\nuSc:f32,\nuIP:array<vec3f,12>,\nuIH:array<vec4f,12>,\nuSB:array<vec3f,72>,\nuOF:array<vec4f,3>,\nuGB:array<vec4f,3>,\nuGT:array<vec4f,3>,\nuGK:i32,\nuGN:array<vec4f,3>,\nuTm:f32,\nuWnd:f32,\nuZMN:f32,\nuH:f32,\nuNIP:f32,\nuDNL:vec4f,\nuDNR:vec4f,\nuDNB:vec4f,\nuDNG:vec4f,\nuDNTl:vec4f,\nuDNTr:vec4f,\nuCamM:vec3f,\nuYPM:vec4f,\nuResM:vec2f,\nuSunM:vec3f,\nuHF2:array<vec3f,84>,\nuHM:i32,\nuGar:i32,\nuAuth:i32,\nuShd:i32,\nuShdP:vec4f,\nuSurfaceVertices:i32,\nuSurfaceActors:i32,\n}\n@group(0)@binding(0)var<uniform> u:U;\nvar<private> aP:vec3f;\nvar<private> aJ:vec4f;\nvar<private> aW:vec4f;\nvar<private> aN:vec3f;\nvar<private> aUV:vec2f;\nvar<private> aHS:f32;\n@group(0)@binding(1)var uJT:texture_2d<f32>;\n@group(0)@binding(2)var uJT_s:sampler;\n@group(0)@binding(3)var uGD:texture_2d<f32>;\n@group(0)@binding(4)var uGD_s:sampler;\nvar<private> vB:vec3f;\nvar<private> vW:vec3f;\nvar<private> vN:vec3f;\nvar<private> vII:i32;\nvar<private> vSh:f32;\nvar<private> vUV:vec2f;\nvar<private> vSt:f32;\n@group(0)@binding(5)var<storage,read> surfaceIn:array<vec4f>;\n@group(0)@binding(6)var<storage,read_write> surfaceOut:array<vec4f>;\nfn hsh(s:f32)->f32{\nreturn fract((sin((s*127.1))*43758.5453));\n}\n@compute @workgroup_size(64) fn main(@builtin(global_invocation_id) gid:vec3u){\nvar svIndex:i32=i32(gid.x);\nvar ii:i32=i32(gid.y);\nif((((svIndex>=u.uSurfaceVertices)||(ii>=u.uSurfaceActors))||((u.uGar!=0)||(u.uHM!=1)))){\nreturn;\n}\naP=surfaceIn[((svIndex*4)+0)].xyz;\naN=surfaceIn[((svIndex*4)+1)].xyz;\naJ=surfaceIn[((svIndex*4)+2)];\naW=surfaceIn[((svIndex*4)+3)];\naHS=0.0;\nvSh=1.0;\nvar jit:vec3f=vec3f(0.0);\nvar rb:i32=i32((u.uIH[ii].w+0.5));\nvar rd2:i32=(rb+1);\nvar r0:vec4f=textureLoad(uJT,vec2i(i32((aJ.x+0.5)),rb),0);\nvar d0:vec4f=textureLoad(uJT,vec2i(i32((aJ.x+0.5)),rd2),0);\nvar r1:vec4f=textureLoad(uJT,vec2i(i32((aJ.y+0.5)),rb),0);\nvar d1:vec4f=textureLoad(uJT,vec2i(i32((aJ.y+0.5)),rd2),0);\nvar r2:vec4f=textureLoad(uJT,vec2i(i32((aJ.z+0.5)),rb),0);\nvar d2:vec4f=textureLoad(uJT,vec2i(i32((aJ.z+0.5)),rd2),0);\nvar r3:vec4f=textureLoad(uJT,vec2i(i32((aJ.w+0.5)),rb),0);\nvar d3:vec4f=textureLoad(uJT,vec2i(i32((aJ.w+0.5)),rd2),0);\nvar s1:f32=select(1.0,(-1.0),(dot(r0,r1)<0.0));\nvar s2:f32=select(1.0,(-1.0),(dot(r0,r2)<0.0));\nvar s3:f32=select(1.0,(-1.0),(dot(r0,r3)<0.0));\nvar br:vec4f=((((aW.x*r0)+((aW.y*s1)*r1))+((aW.z*s2)*r2))+((aW.w*s3)*r3));\nvar bd:vec4f=((((aW.x*d0)+((aW.y*s1)*d1))+((aW.z*s2)*d2))+((aW.w*s3)*d3));\nvar nr:f32=max(length(br),1e-6);\nbr=(br/nr);\nbd=(bd/nr);\nvar aP2:vec3f=(aP+jit);\nif((u.uGar==1)){\naP2=(aP2+textureLoad(uGD,vec2i(svIndex,ii),0).xyz);\n}\nif((u.uGar==2)){\nvar hf9:vec3f=u.uHF2[((ii*7)+i32((aHS+0.5)))];\nvar hw9:f32=clamp((aHS/6.0),0.0,1.0);\nhw9=(hw9*hw9);\naP2=(aP2+(vec3f(hf9.x,hf9.y,hf9.z)*hw9));\n}\nif(((u.uHM==1)&&(u.uAuth==0))){\nvar hb1:f32=((aP.z-u.uZMN)/u.uH);\nvar w1:f32=(smoothstep(0.885,0.905,hb1)*(1.0-smoothstep(0.930,0.9448,hb1)));\nif((aP.y>(-1200.0))){\nw1=0.0;\n}\naP2.z=(aP2.z+(704.0*w1));\nvar w2:f32=(((1.0-smoothstep(0.915,0.925,hb1))*smoothstep(0.885,0.895,hb1))*smoothstep((u.uH*0.015),(u.uH*0.03),abs(aP.x)));\nif((aP.y>(-1200.0))){\nw2=0.0;\n}\naP2.x=(aP2.x*(1.0+(0.08*w2)));\n}\nif(((u.uAuth==0)&&(u.uGar==0))){\nvar fh9:f32=((aP.z-u.uZMN)/u.uH);\nif((((fh9>0.885)&&(fh9<0.985))&&(aP.y<(-1200.0)))){\nvar HH7:f32=(0.115*u.uH);\nvar hc7:vec3f=vec3f(0.0,0.0,(u.uZMN+(0.94*u.uH)));\nvar no7:vec3f=(aN*sign(dot(aN,(aP-hc7))));\nvar ez7:f32=((fh9-0.9420)*u.uH);\nvar bz7:f32=((fh9-0.954)*u.uH);\nvar mz7:f32=((fh9-0.908)*u.uH);\nvar ax7:f32=abs(aP.x);\nvar bo7:f32=(ax7/HH7);\nvar bt7:f32=clamp(((bo7-0.085)/0.16),0.0,1.0);\nvar bc7:f32=(((0.024*sin((3.14159*pow(bt7,1.45))))-(0.006*bt7))*HH7);\nvar br7:f32=(((1.0-smoothstep(0.0,(0.020*HH7),abs((bz7-bc7))))*smoothstep(0.0,0.10,bt7))*(1.0-smoothstep(0.90,1.0,bt7)));\nvar ex7:f32=(ax7-(0.16*HH7));\nvar ed7:f32=sqrt(((ex7*ex7)+((ez7*ez7)*1.4)));\nvar so7:f32=(1.0-smoothstep(0.0,(0.085*HH7),ed7));\nvar lp7:f32=((1.0-smoothstep(0.0,(0.018*HH7),abs(mz7)))*(1.0-smoothstep((0.030*HH7),(0.062*HH7),ax7)));\nvar d7:f32=((((br7*0.0100)*HH7)-((so7*0.0250)*HH7))+((lp7*0.0100)*HH7));\naP2=(aP2+(no7*d7));\n}\n}\nvar sp:vec3f=((aP2+(2.0*cross(br.xyz,(cross(br.xyz,aP2)+(br.w*aP2)))))+(2.0*(((br.w*bd.xyz)-(bd.w*br.xyz))+cross(br.xyz,bd.xyz))));\nvar sn:vec3f=(aN+(2.0*cross(br.xyz,(cross(br.xyz,aN)+(br.w*aN)))));\nvar nyup:vec3f=vec3f(sn.x,sn.z,(-sn.y));\nvN=vec3f(((nyup.x*u.uIH[ii].y)-(nyup.z*u.uIH[ii].x)),nyup.y,((nyup.x*u.uIH[ii].x)+(nyup.z*u.uIH[ii].y)));\nvB=aP;\nvar yup:vec3f=vec3f(sp.x,sp.z,(-sp.y));\nvar rot:vec3f=vec3f(((yup.x*u.uIH[ii].y)-(yup.z*u.uIH[ii].x)),yup.y,((yup.x*u.uIH[ii].x)+(yup.z*u.uIH[ii].y)));\nvar wp:vec3f=(u.uIP[ii]+(rot*(u.uSc*u.uIH[ii].z)));\nvar hb2:f32=((aP.z-u.uZMN)/u.uH);\nvar sbo:vec3f=vec3f(0.0);\nif(((u.uHM>1)||(u.uGar==2))){\nvar tw9:f32=clamp(((1.02-hb2)*1.7),0.0,1.0);\nsbo=(u.uSB[(ii*6)]*tw9);\nvar wp9:f32=((u.uTm*3.1)+(vSh*7.9));\nsbo=(sbo+(vec3f(0.7,(0.15*sin((wp9*1.7))),0.7)*((sin(wp9)*u.uWnd)*tw9)));\n}\nelse{\nif((u.uGar==1)){\nsbo=(u.uSB[((ii*6)+2)]*(1.0-smoothstep(0.30,0.55,hb2)));\nif((u.uDNL.w>0.5)){\nvar wLg:f32=(1.0-smoothstep(0.0,(u.uDNL.w*1.4),length((aP-u.uDNL.xyz))));\nvar wRg:f32=(1.0-smoothstep(0.0,(u.uDNR.w*1.4),length((aP-u.uDNR.xyz))));\nsbo=(sbo+((u.uSB[((ii*6)+1)]*max(wLg,wRg))*(1.0-smoothstep((-1000.0),(-14000.0),aP.y))));\n}\nif((u.uDNB.w>0.5)){\nsbo=(sbo+(u.uSB[((ii*6)+3)]*(1.0-smoothstep(0.0,(u.uDNB.w*1.6),length((aP-u.uDNB.xyz))))));\n}\nif((u.uDNG.w>0.5)){\nsbo=(sbo+(u.uSB[((ii*6)+4)]*(1.0-smoothstep(0.0,(u.uDNG.w*1.6),length((aP-u.uDNG.xyz))))));\n}\nif((u.uDNTl.w>0.5)){\nvar wT9:f32=max((1.0-smoothstep(0.0,(u.uDNTl.w*1.6),length((aP-u.uDNTl.xyz)))),(1.0-smoothstep(0.0,(u.uDNTr.w*1.6),length((aP-u.uDNTr.xyz)))));\nsbo=(sbo+(u.uSB[((ii*6)+5)]*wT9));\n}\n}\nelse{\nvar wB:f32;\nif((u.uDNL.w>0.5)){\nvar wL:f32=(1.0-smoothstep(0.0,(u.uDNL.w*1.4),length((aP-u.uDNL.xyz))));\nvar wR:f32=(1.0-smoothstep(0.0,(u.uDNR.w*1.4),length((aP-u.uDNR.xyz))));\nwB=(max(wL,wR)*(1.0-smoothstep((-1000.0),(-14000.0),aP.y)));\n}\nelse{\nwB=(((1.0-smoothstep(0.0,0.085,abs((hb2-u.uNIP))))*(1.0-smoothstep(6000.0,26000.0,abs(aP.x))))*(1.0-smoothstep((-1000.0),(-14000.0),aP.y)));\n}\nsbo=(u.uSB[((ii*6)+1)]*wB);\nif((u.uDNB.w>0.5)){\nsbo=(sbo+(u.uSB[((ii*6)+3)]*(1.0-smoothstep(0.0,(u.uDNB.w*1.6),length((aP-u.uDNB.xyz))))));\n}\nif((u.uDNG.w>0.5)){\nsbo=(sbo+(u.uSB[((ii*6)+4)]*(1.0-smoothstep(0.0,(u.uDNG.w*1.6),length((aP-u.uDNG.xyz))))));\n}\nif((u.uDNTl.w>0.5)){\nvar wT8:f32=max((1.0-smoothstep(0.0,(u.uDNTl.w*1.6),length((aP-u.uDNTl.xyz)))),(1.0-smoothstep(0.0,(u.uDNTr.w*1.6),length((aP-u.uDNTr.xyz)))));\nsbo=(sbo+(u.uSB[((ii*6)+5)]*wT8));\n}\n}\n}\nwp=(wp+sbo);\nvar surfaceIndex:i32=(((ii*u.uSurfaceVertices)+svIndex)*2);\nsurfaceOut[surfaceIndex]=vec4f(wp.x,wp.y,wp.z,1.0);\nsurfaceOut[(surfaceIndex+1)]=vec4f(vN.x,vN.y,vN.z,0.0);\n}\n" as *u8
8 let vertex:*u8="struct U{\nuSc:f32,\nuIP:array<vec3f,12>,\nuIH:array<vec4f,12>,\nuSB:array<vec3f,72>,\nuOF:array<vec4f,3>,\nuGB:array<vec4f,3>,\nuGT:array<vec4f,3>,\nuGK:i32,\nuGN:array<vec4f,3>,\nuTm:f32,\nuWnd:f32,\nuZMN:f32,\nuH:f32,\nuNIP:f32,\nuDNL:vec4f,\nuDNR:vec4f,\nuDNB:vec4f,\nuDNG:vec4f,\nuDNTl:vec4f,\nuDNTr:vec4f,\nuCamM:vec3f,\nuYPM:vec4f,\nuResM:vec2f,\nuSunM:vec3f,\nuHF2:array<vec3f,84>,\nuHM:i32,\nuGar:i32,\nuAuth:i32,\nuShd:i32,\nuShdP:vec4f,\nuSurfaceVertices:i32,\nuSurfaceActors:i32,\n}\n@group(0)@binding(0)var<uniform> u:U;\n@group(0)@binding(1)var uJT:texture_2d<f32>;\n@group(0)@binding(2)var uJT_s:sampler;\n@group(0)@binding(3)var uGD:texture_2d<f32>;\n@group(0)@binding(4)var uGD_s:sampler;\nstruct VIO{@builtin(position) bpos:vec4f,@location(0) vB:vec3f,@location(1) vW:vec3f,@location(2) vN:vec3f,@location(3) @interpolate(flat) vII:i32,@location(4) @interpolate(flat) vSh:f32,@location(5) vUV:vec2f,@location(6) vSt:f32,}\n@group(0)@binding(6)var<storage,read> surfaceOut:array<vec4f>;\n@vertex fn main(@builtin(vertex_index) vi:u32,@builtin(instance_index) iidx:u32,@location(0) aP:vec3f,@location(1) aJ:vec4f,@location(2) aW:vec4f,@location(3) aN:vec3f,@location(4) aUV:vec2f,@location(5) aHS:f32)->VIO{\nvar vo:VIO;\nvar ii:i32=i32(iidx);\nvo.vII=ii;\nvo.vB=aP;\nvo.vUV=aUV;\nvo.vSh=1.0;\nvar surfaceNormal:vec3f=(((surfaceOut[((((ii*u.uSurfaceVertices)+i32(aJ.x))*2)+1)].xyz*aW.x)+(surfaceOut[((((ii*u.uSurfaceVertices)+i32(aJ.y))*2)+1)].xyz*aW.y))+(surfaceOut[((((ii*u.uSurfaceVertices)+i32(aJ.z))*2)+1)].xyz*aW.z));\nvo.vN=vec3f(0.0);\nif((dot(surfaceNormal,surfaceNormal)>0.0)){\nvo.vN=normalize(surfaceNormal);\n}\nelse{\nvar surfaceGeometric:vec3f=cross((surfaceOut[((((ii*u.uSurfaceVertices)+i32(aJ.y))*2)+0)].xyz-surfaceOut[((((ii*u.uSurfaceVertices)+i32(aJ.x))*2)+0)].xyz),(surfaceOut[((((ii*u.uSurfaceVertices)+i32(aJ.z))*2)+0)].xyz-surfaceOut[((((ii*u.uSurfaceVertices)+i32(aJ.x))*2)+0)].xyz));\nif((dot(surfaceGeometric,surfaceGeometric)>0.0)){\nvo.vN=normalize(surfaceGeometric);\n}\n}\nvar wp:vec3f=((((surfaceOut[((((ii*u.uSurfaceVertices)+i32(aJ.x))*2)+0)].xyz*aW.x)+(surfaceOut[((((ii*u.uSurfaceVertices)+i32(aJ.y))*2)+0)].xyz*aW.y))+(surfaceOut[((((ii*u.uSurfaceVertices)+i32(aJ.z))*2)+0)].xyz*aW.z))+(vo.vN*(aW.w*(u.uSc*u.uIH[ii].z))));\nvar gbm:i32=i32((u.uGB[(ii)/4][(ii)%4]+0.5));\nif(((u.uGar==1)&&(!((((gbm==4)||(gbm==5))||(gbm==8))||(gbm==9))))){\nwp=vec3f(0.0,(-999.0),0.0);\n}\nif((u.uGar==3)){\nvar gk9:i32=u.uGK;\nvar hv9:i32=select(i32((u.uGB[(ii)/4][(ii)%4]+0.5)),i32((u.uGT[(ii)/4][(ii)%4]+0.5)),((gk9==1)||(gk9==6)));\nvar wr9:bool=(hv9==gk9);\nif((gk9==8)){\nwr9=((((hv9==4)||(hv9==5))||(hv9==8))||(hv9==9));\n}\nif((!wr9)){\nwp=vec3f(0.0,(-999.0),0.0);\n}\n}\nvo.vSt=0.0;\nif(((u.uShd==1)&&(wp.y>(-900.0)))){\nvar gY9:f32=(u.uIP[ii].y+u.uShdP.x);\nvar sd9:vec3f=normalize(u.uSunM);\nvar sy9:f32=max(sd9.y,u.uShdP.y);\nvo.vSt=((wp.y-gY9)/sy9);\nwp=vec3f((wp.x-(sd9.x*vo.vSt)),gY9,(wp.z-(sd9.z*vo.vSt)));\n}\nvo.vW=wp;\nvar rel:vec3f=(wp-u.uCamM);\nvar fw:vec3f=vec3f((u.uYPM.x*u.uYPM.w),u.uYPM.z,(u.uYPM.y*u.uYPM.w));\nvar rt:vec3f=vec3f(u.uYPM.y,0.0,(-u.uYPM.x));\nvar up:vec3f=vec3f(((-u.uYPM.x)*u.uYPM.z),u.uYPM.w,((-u.uYPM.y)*u.uYPM.z));\nvar zv:f32=max(dot(rel,fw),0.081);\nvo.bpos=vec4f(dot(rel,rt),(dot(rel,up)*(u.uResM.x/u.uResM.y)),((zv*(300.08/299.92))-(48.0/299.92)),zv);\nvo.bpos.z=(vo.bpos.z+vo.bpos.w)*0.5;\nreturn vo;\n}\n" as *u8
9 let digest:*u8=sys_mmap(65)
10 gpa_sha(compute,gpe_slen(compute),digest)
11 if fsx_seq(digest,"8c2ed6f1fb5b1c02a807531c5affd4ae76791c80cf8c3b4edb637c5c53e96543" as *u8)!=1{gsb_failed=1;sys_munmap(digest,65);return 0}
12 gpa_sha(vertex,gpe_slen(vertex),digest)
13 if fsx_seq(digest,"11268cd64433ea762cdcd6d1091627eb4ea8d72caf491cca218de011df4edbb1" as *u8)!=1{gsb_failed=1;sys_munmap(digest,65);return 0}
14 sys_munmap(digest,65)
15 var p:i64=gpe_ap(out,pos,"const NX_SURFACE_COMPUTE=" as *u8)
16 p=rk_js_str(out,p,compute,gpe_slen(compute))
17 p=gpe_ap(out,p,";\nconst NX_ATTACHMENT_VERTEX=" as *u8)
18 p=rk_js_str(out,p,vertex,gpe_slen(vertex))
19 p=gpe_ap(out,p,";\n" as *u8)
20 p=gpe_ap(out,p,"\nfunction nxGarmentFamilyWorn(kind,id){return id===kind||(kind===8&&(id===4||id===5||id===8||id===9));}\nconst NX_SURFACE_LAYOUT={uSc:0,uIP:16,uIH:208,uSB:400,uOF:1552,uGB:1600,uGT:1648,uGK:1696,uGN:1712,uTm:1760,uWnd:1764,uZMN:1768,uH:1772,uNIP:1776,uDNL:1792,uDNR:1808,uDNB:1824,uDNG:1840,uDNTl:1856,uDNTr:1872,uCamM:1888,uYPM:1904,uResM:1920,uSunM:1936,uHF2:1952,uHM:3296,uGar:3300,uAuth:3304,uShd:3308,uShdP:3312,};\nasync function nxGpuSurfaceCreate(dev,G,R,P,mkg,owned){\n if(NXCV_SIZE!==3328||Object.keys(NXCV).length!==Object.keys(NX_SURFACE_LAYOUT).length||Object.keys(NX_SURFACE_LAYOUT).some(k=>NXCV[k]!==NX_SURFACE_LAYOUT[k]))throw Error('surface-source-layout-identity');\n const packet=P.attachment,h=new Uint32Array(packet,0,16),nr=h[2],nt=h[3],nv=h[4],kind=h[5];\n if(h[0]!==827343175||h[1]!==1||nv!==P.nv||kind<1||kind>9||h[10]!==packet.byteLength||h[6]!==64||h[7]!==64+nr*12||h[8]!==h[7]+nr*16||h[9]!==h[8]+nr*16||h[9]+nt*12!==packet.byteLength)throw Error('native-attachment-packet-contract');\n const rows=12,inputBytes=nv*64,outputBytes=nv*32*rows,uniformBytes=3344,L=dev.limits;\n if(!Number.isSafeInteger(outputBytes)||nv>Math.floor(2147483647/(2*rows))||inputBytes>L.maxStorageBufferBindingSize||inputBytes>L.maxBufferSize||outputBytes>L.maxStorageBufferBindingSize||outputBytes>L.maxBufferSize||Math.ceil(nv/64)>L.maxComputeWorkgroupsPerDimension||rows>L.maxComputeWorkgroupsPerDimension)throw Error('surface-device-budget');\n const input=new Float32Array(nv*16);for(let i=0;i<nv;i++){input.set(P.pos.subarray(i*3,i*3+3),i*16);input.set(P.nrm.subarray(i*3,i*3+3),i*16+4);input.set(P.aj.subarray(i*4,i*4+4),i*16+8);input.set(P.aw3.subarray(i*4,i*4+4),i*16+12);}\n const mk=(size,usage)=>{const b=dev.createBuffer({size,usage:usage|GPUBufferUsage.COPY_DST});owned.push(b);return b;};\n const source=mkg(input,GPUBufferUsage.STORAGE),surface=mk(outputBytes,GPUBufferUsage.STORAGE),cu=mk(uniformBytes,GPUBufferUsage.UNIFORM),vu=mk(uniformBytes,GPUBufferUsage.UNIFORM),fu=mk(NXCF_SIZE,GPUBufferUsage.UNIFORM);\n const dummy=dev.createTexture({size:[1,1],format:'rgba8unorm',usage:GPUTextureUsage.TEXTURE_BINDING|GPUTextureUsage.COPY_DST});owned.push(dummy);dev.queue.writeTexture({texture:dummy},new Uint8Array([128,128,255,255]),{bytesPerRow:4},[1,1]);\n const module=async code=>{const m=dev.createShaderModule({code}),info=await m.getCompilationInfo();const error=info.messages.find(x=>x.type==='error');if(error)throw Error(error.message);return m;};\n const cm=await module(NX_SURFACE_COMPUTE),vm=await module(NX_ATTACHMENT_VERTEX),fm=await module(WGSL_CASTF);\n const cp=await dev.createComputePipelineAsync({layout:'auto',compute:{module:cm,entryPoint:'main'}}),vb=(arrayStride,shaderLocation,format)=>({arrayStride,attributes:[{shaderLocation,offset:0,format}]});\n const pl=await dev.createRenderPipelineAsync({layout:'auto',depthStencil:{format:'depth24plus',depthWriteEnabled:true,depthCompare:'less'},vertex:{module:vm,entryPoint:'main',buffers:[vb(12,0,'float32x3'),vb(16,1,'float32x4'),vb(16,2,'float32x4'),vb(12,3,'float32x3'),vb(8,4,'float32x2'),vb(4,5,'float32')]},fragment:{module:fm,entryPoint:'main',targets:[{format:navigator.gpu.getPreferredCanvasFormat()}]},primitive:{topology:'triangle-list'}});\n const cb=dev.createBindGroup({layout:cp.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:cu}},{binding:1,resource:R.jt.createView()},{binding:3,resource:dummy.createView()},{binding:5,resource:{buffer:source}},{binding:6,resource:{buffer:surface}}]});\n const vg=dev.createBindGroup({layout:pl.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:vu}},{binding:6,resource:{buffer:surface}}]});\n const fg=dev.createBindGroup({layout:pl.getBindGroupLayout(1),entries:[{binding:0,resource:{buffer:fu}},{binding:1,resource:G.t3.createView()},{binding:2,resource:(R.atl||dummy).createView()},{binding:3,resource:R.samp},{binding:4,resource:(R.nrm||dummy).createView()},{binding:5,resource:R.samp}]});\n const vbs=[new Float32Array(packet,h[6],nr*3),new Float32Array(packet,h[7],nr*4),new Float32Array(packet,h[8],nr*4),new Float32Array(nr*3),new Float32Array(nr*2),new Float32Array(nr)].map(x=>mkg(x,GPUBufferUsage.VERTEX));const ib=mkg(new Uint32Array(packet,h[9],nt*3),GPUBufferUsage.INDEX);\n const cv=new ArrayBuffer(uniformBytes),cf=new ArrayBuffer(NXCF_SIZE),cvBytes=new Uint8Array(cv),cfBytes=new Uint8Array(cf),ci=new Int32Array(cv),fi=new Int32Array(cf),ids=new Int32Array(rows);\n const report=window.__nx_surface={state:'ready',kind,source_vertices:nv,attachment_vertices:nr,triangles:nt,input_bytes:inputBytes,output_bytes:outputBytes,dispatches:0,draws:0,rows:0,ids,reason:null};\n const state={kind,rows:0,prepare(enc,G,N,count,nm){state.rows=0;report.rows=0;if(!Number.isInteger(count)||count<1||count>rows||!N.nativeIds)throw Error('surface-actor-rows');let worn=false;\n for(let i=0;i<count;i++){const id=N.nativeIds[i];if(!Number.isInteger(id)||id<0||id>=nm)throw Error('surface-actor-id');for(let j=0;j<i;j++)if(ids[j]===id)throw Error('surface-actor-duplicate');ids[i]=id;const wardrobe=(kind===1||kind===6)?N.gt:N.gb;if(wardrobe&&nxGarmentFamilyWorn(kind,Math.round(wardrobe[i])))worn=true;}\n if(!worn)return;cvBytes.set(G.castCVBytes);ci[3328/4]=nv;ci[3332/4]=count;if(ci[NXCV.uGar/4]!==0||ci[NXCV.uHM/4]!==1)throw Error('surface-mode');dev.queue.writeBuffer(cu,0,cv);const pass=enc.beginComputePass();pass.setPipeline(cp);pass.setBindGroup(0,cb);pass.dispatchWorkgroups(Math.ceil(nv/64),count);pass.end();\n ci[NXCV.uGar/4]=3;ci[NXCV.uGK/4]=kind;dev.queue.writeBuffer(vu,0,cv);cfBytes.set(G.castCFBytes);fi[NXCF.uGar/4]=3;fi[NXCF.uGK/4]=kind;dev.queue.writeBuffer(fu,0,cf);state.rows=count;report.rows=count;report.dispatches++;report.state='prepared';\n },draw(rp,mesh){fi[NXCF.uGM/4]=mesh.m;dev.queue.writeBuffer(fu,0,cf);rp.setPipeline(pl);rp.setBindGroup(0,vg);rp.setBindGroup(1,fg);for(let k=0;k<6;k++)rp.setVertexBuffer(k,vbs[k]);rp.setIndexBuffer(ib,'uint32');rp.drawIndexed(nt*3,state.rows);rp.setPipeline(G.cpl);report.draws++;report.state='submitted';}};return state;\n}\n" as *u8)
21 return p
22}