--- /dev/null
+// src/shaders/compute.wgsl
+var compute_default = "struct UBO{blockhash:array<vec4<u32>,2>,seed:vec2<u32>,threshold:u32};@group(0)@binding(0)var<uniform> ubo:UBO;struct WORK{nonce:vec2<u32>,found:atomic<u32>};@group(0)@binding(1)var<storage,read_write>work:WORK;const BLAKE2B_IV_0=vec2(0xF2BDC900u,0x6A09E667u);const Z=vec2(0u);const CARRY=vec4(1u,0u,1u,0u);const ROTATE_1=vec4(1u);const ROTATE_8=vec4(8u);const ROTATE_16=vec4(16u);const ROTATE_24=vec4(24u);const ROTATE_31=vec4(31u);var<workgroup> found:bool;@compute @workgroup_size(32)fn search(@builtin(global_invocation_id)global_id:vec3<u32>,@builtin(local_invocation_id)local_id:vec3<u32>){found=(local_id.x==0u&&atomicLoad(&work.found)!=0u);workgroupBarrier();if(found){return;}main(global_id);}@compute @workgroup_size(1)fn validate(@builtin(global_invocation_id)global_id:vec3<u32>){main(global_id);}fn main(id:vec3<u32>){let m0:vec2<u32>=ubo.seed ^ id.xy;let m1:vec2<u32>=ubo.blockhash[0u].xy;let m2:vec2<u32>=ubo.blockhash[0u].zw;let m3:vec2<u32>=ubo.blockhash[1u].xy;let m4:vec2<u32>=ubo.blockhash[1u].zw;var v01:vec4<u32>=vec4(BLAKE2B_IV_0,0x84CAA73Bu,0xBB67AE85u);var v23:vec4<u32>=vec4(0xFE94F82Bu,0x3C6EF372u,0x5F1D36F1u,0xA54FF53Au);var v45:vec4<u32>=vec4(0xADE682D1u,0x510E527Fu,0x2B3E6C1Fu,0x9B05688Cu);var v67:vec4<u32>=vec4(0xFB41BD6Bu,0x1F83D9ABu,0x137E2179u,0x5BE0CD19u);var v89:vec4<u32>=vec4(0xF3BCC908u,0x6A09E667u,0x84CAA73Bu,0xBB67AE85u);var vAB:vec4<u32>=vec4(0xFE94F82Bu,0x3C6EF372u,0x5F1D36F1u,0xA54FF53Au);var vCD:vec4<u32>=vec4(0xADE682F9u,0x510E527Fu,0x2B3E6C1Fu,0x9B05688Cu);var vEF:vec4<u32>=vec4(0x04BE4294u,0xE07C2654u,0x137E2179u,0x5BE0CD19u);var v56:vec4<u32>;var vFC:vec4<u32>;var v74:vec4<u32>;var vDE:vec4<u32>;var s0:vec4<u32>;var s1:vec4<u32>;s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;s0=v01+vec4(m0,m2);v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;v23+=vec4(m4,Z);v23.y+=u32(v23.x<m4.x);vCD=(vCD ^ v01).yxwz;vEF=(vEF ^ v23).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_24)|(v45<<ROTATE_8).yxwz;v67 ^=vAB;v67=(v67>>ROTATE_24)|(v67<<ROTATE_8).yxwz;s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;s0=v01+vec4(m1,m3);v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;vCD ^=v01;vCD=(vCD>>ROTATE_16)|(vCD<<ROTATE_16).yxwz;vEF ^=v23;vEF=(vEF>>ROTATE_16)|(vEF<<ROTATE_16).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_31).yxwz|(v45<<ROTATE_1);v67 ^=vAB;v67=(v67>>ROTATE_31).yxwz|(v67<<ROTATE_1);v56=vec4(v45.zw,v67.xy);v74=vec4(v67.zw,v45.xy);vFC=vec4(vEF.zw,vCD.xy);vDE=vec4(vCD.zw,vEF.xy);s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;vFC=(vFC ^ v01).yxwz;vDE=(vDE ^ v23).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_24)|(v56<<ROTATE_8).yxwz;v74 ^=v89;v74=(v74>>ROTATE_24)|(v74<<ROTATE_8).yxwz;s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;vFC ^=v01;vFC=(vFC>>ROTATE_16)|(vFC<<ROTATE_16).yxwz;vDE ^=v23;vDE=(vDE>>ROTATE_16)|(vDE<<ROTATE_16).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_31).yxwz|(v56<<ROTATE_1);v74 ^=v89;v74=(v74>>ROTATE_31).yxwz|(v74<<ROTATE_1);v45=vec4(v74.zw,v56.xy);v67=vec4(v56.zw,v74.xy);vCD=vec4(vFC.zw,vDE.xy);vEF=vec4(vDE.zw,vFC.xy);s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v01+=vec4(Z,m4);v01.w+=u32(v01.z<m4.x);vCD=(vCD ^ v01).yxwz;vEF=(vEF ^ v23).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_24)|(v45<<ROTATE_8).yxwz;v67 ^=vAB;v67=(v67>>ROTATE_24)|(v67<<ROTATE_8).yxwz;s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;vCD ^=v01;vCD=(vCD>>ROTATE_16)|(vCD<<ROTATE_16).yxwz;vEF ^=v23;vEF=(vEF>>ROTATE_16)|(vEF<<ROTATE_16).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_31).yxwz|(v45<<ROTATE_1);v67 ^=vAB;v67=(v67>>ROTATE_31).yxwz|(v67<<ROTATE_1);v56=vec4(v45.zw,v67.xy);v74=vec4(v67.zw,v45.xy);vFC=vec4(vEF.zw,vCD.xy);vDE=vec4(vCD.zw,vEF.xy);s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;s0=v01+vec4(m1,m0);v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;vFC=(vFC ^ v01).yxwz;vDE=(vDE ^ v23).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_24)|(v56<<ROTATE_8).yxwz;v74 ^=v89;v74=(v74>>ROTATE_24)|(v74<<ROTATE_8).yxwz;s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v01+=vec4(Z,m2);v01.w+=u32(v01.z<m2.x);v23+=vec4(Z,m3);v23.w+=u32(v23.z<m3.x);vFC ^=v01;vFC=(vFC>>ROTATE_16)|(vFC<<ROTATE_16).yxwz;vDE ^=v23;vDE=(vDE>>ROTATE_16)|(vDE<<ROTATE_16).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_31).yxwz|(v56<<ROTATE_1);v74 ^=v89;v74=(v74>>ROTATE_31).yxwz|(v74<<ROTATE_1);v45=vec4(v74.zw,v56.xy);v67=vec4(v56.zw,v74.xy);vCD=vec4(vFC.zw,vDE.xy);vEF=vec4(vDE.zw,vFC.xy);s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;vCD=(vCD ^ v01).yxwz;vEF=(vEF ^ v23).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_24)|(v45<<ROTATE_8).yxwz;v67 ^=vAB;v67=(v67>>ROTATE_24)|(v67<<ROTATE_8).yxwz;s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;s0=v01+vec4(Z,m0);v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+vec4(m2,Z);v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;vCD ^=v01;vCD=(vCD>>ROTATE_16)|(vCD<<ROTATE_16).yxwz;vEF ^=v23;vEF=(vEF>>ROTATE_16)|(vEF<<ROTATE_16).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_31).yxwz|(v45<<ROTATE_1);v67 ^=vAB;v67=(v67>>ROTATE_31).yxwz|(v67<<ROTATE_1);v56=vec4(v45.zw,v67.xy);v74=vec4(v67.zw,v45.xy);vFC=vec4(vEF.zw,vCD.xy);vDE=vec4(vCD.zw,vEF.xy);s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v01.z+=m3.x;v01.w+=m3.y+u32(v01.z<m3.x);vFC=(vFC ^ v01).yxwz;vDE=(vDE ^ v23).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_24)|(v56<<ROTATE_8).yxwz;v74 ^=v89;v74=(v74>>ROTATE_24)|(v74<<ROTATE_8).yxwz;s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v23+=vec4(m1,m4)+vec4(0u,u32(v23.x+vec4(m1,m4).x<v23.x),0u,u32(v23.z+vec4(m1,m4).z<v23.z));vFC ^=v01;vFC=(vFC>>ROTATE_16)|(vFC<<ROTATE_16).yxwz;vDE ^=v23;vDE=(vDE>>ROTATE_16)|(vDE<<ROTATE_16).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_31).yxwz|(v56<<ROTATE_1);v74 ^=v89;v74=(v74>>ROTATE_31).yxwz|(v74<<ROTATE_1);v45=vec4(v74.zw,v56.xy);v67=vec4(v56.zw,v74.xy);vCD=vec4(vFC.zw,vDE.xy);vEF=vec4(vDE.zw,vFC.xy);s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v01+=vec4(Z,m3)+vec4(Z,0u,u32(v01.z+vec4(Z,m3).z<v01.z));vCD=(vCD ^ v01).yxwz;vEF=(vEF ^ v23).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_24)|(v45<<ROTATE_8).yxwz;v67 ^=vAB;v67=(v67>>ROTATE_24)|(v67<<ROTATE_8).yxwz;s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v01+=vec4(Z,m1)+vec4(Z,0u,u32(v01.z+vec4(Z,m1).z<v01.z));vCD ^=v01;vCD=(vCD>>ROTATE_16)|(vCD<<ROTATE_16).yxwz;vEF ^=v23;vEF=(vEF>>ROTATE_16)|(vEF<<ROTATE_16).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_31).yxwz|(v45<<ROTATE_1);v67 ^=vAB;v67=(v67>>ROTATE_31).yxwz|(v67<<ROTATE_1);v56=vec4(v45.zw,v67.xy);v74=vec4(v67.zw,v45.xy);vFC=vec4(vEF.zw,vCD.xy);vDE=vec4(vCD.zw,vEF.xy);s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v01+=vec4(m2,Z)+vec4(0u,u32(v01.x+vec4(m2,Z).x<v01.x),Z);v23+=vec4(m4,Z)+vec4(0u,u32(v23.x+vec4(m4,Z).x<v23.x),Z);vFC=(vFC ^ v01).yxwz;vDE=(vDE ^ v23).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_24)|(v56<<ROTATE_8).yxwz;v74 ^=v89;v74=(v74>>ROTATE_24)|(v74<<ROTATE_8).yxwz;s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v23+=vec4(m0,Z)+vec4(0u,u32(v23.x+vec4(m0,Z).x<v23.x),Z);vFC ^=v01;vFC=(vFC>>ROTATE_16)|(vFC<<ROTATE_16).yxwz;vDE ^=v23;vDE=(vDE>>ROTATE_16)|(vDE<<ROTATE_16).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_31).yxwz|(v56<<ROTATE_1);v74 ^=v89;v74=(v74>>ROTATE_31).yxwz|(v74<<ROTATE_1);v45=vec4(v74.zw,v56.xy);v67=vec4(v56.zw,v74.xy);vCD=vec4(vFC.zw,vDE.xy);vEF=vec4(vDE.zw,vFC.xy);s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;s1=v23+vec4(m2,Z);v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;vCD=(vCD ^ v01).yxwz;vEF=(vEF ^ v23).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_24)|(v45<<ROTATE_8).yxwz;v67 ^=vAB;v67=(v67>>ROTATE_24)|(v67<<ROTATE_8).yxwz;s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v01+=vec4(m0,Z)+vec4(0u,u32(v01.x+vec4(m0,Z).x<v01.x),Z);v23+=vec4(m4,Z)+vec4(0u,u32(v23.x+vec4(m4,Z).x<v23.x),Z);vCD ^=v01;vCD=(vCD>>ROTATE_16)|(vCD<<ROTATE_16).yxwz;vEF ^=v23;vEF=(vEF>>ROTATE_16)|(vEF<<ROTATE_16).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_31).yxwz|(v45<<ROTATE_1);v67 ^=vAB;v67=(v67>>ROTATE_31).yxwz|(v67<<ROTATE_1);v56=vec4(v45.zw,v67.xy);v74=vec4(v67.zw,v45.xy);vFC=vec4(vEF.zw,vCD.xy);vDE=vec4(vCD.zw,vEF.xy);s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;s1=v23+vec4(Z,m3);v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;vFC=(vFC ^ v01).yxwz;vDE=(vDE ^ v23).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_24)|(v56<<ROTATE_8).yxwz;v74 ^=v89;v74=(v74>>ROTATE_24)|(v74<<ROTATE_8).yxwz;s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v01+=vec4(m1,Z)+vec4(0u,u32(v01.x+vec4(m1,Z).x<v01.x),Z);vFC ^=v01;vFC=(vFC>>ROTATE_16)|(vFC<<ROTATE_16).yxwz;vDE ^=v23;vDE=(vDE>>ROTATE_16)|(vDE<<ROTATE_16).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_31).yxwz|(v56<<ROTATE_1);v74 ^=v89;v74=(v74>>ROTATE_31).yxwz|(v74<<ROTATE_1);v45=vec4(v74.zw,v56.xy);v67=vec4(v56.zw,v74.xy);vCD=vec4(vFC.zw,vDE.xy);vEF=vec4(vDE.zw,vFC.xy);s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v01+=vec4(m2,Z)+vec4(0u,u32(v01.x+vec4(m2,Z).x<v01.x),Z);v23+=vec4(m0,Z)+vec4(0u,u32(v23.x+vec4(m0,Z).x<v23.x),Z);vCD=(vCD ^ v01).yxwz;vEF=(vEF ^ v23).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_24)|(v45<<ROTATE_8).yxwz;v67 ^=vAB;v67=(v67>>ROTATE_24)|(v67<<ROTATE_8).yxwz;s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;s1=v23+vec4(Z,m3);v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;vCD ^=v01;vCD=(vCD>>ROTATE_16)|(vCD<<ROTATE_16).yxwz;vEF ^=v23;vEF=(vEF>>ROTATE_16)|(vEF<<ROTATE_16).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_31).yxwz|(v45<<ROTATE_1);v67 ^=vAB;v67=(v67>>ROTATE_31).yxwz|(v67<<ROTATE_1);v56=vec4(v45.zw,v67.xy);v74=vec4(v67.zw,v45.xy);vFC=vec4(vEF.zw,vCD.xy);vDE=vec4(vCD.zw,vEF.xy);s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v01+=vec4(m4,Z)+vec4(0u,u32(v01.x+vec4(m4,Z).x<v01.x),Z);v23+=vec4(Z,m1)+vec4(Z,0u,u32(v23.z+vec4(Z,m1).z<v23.z));vFC=(vFC ^ v01).yxwz;vDE=(vDE ^ v23).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_24)|(v56<<ROTATE_8).yxwz;v74 ^=v89;v74=(v74>>ROTATE_24)|(v74<<ROTATE_8).yxwz;s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;vFC ^=v01;vFC=(vFC>>ROTATE_16)|(vFC<<ROTATE_16).yxwz;vDE ^=v23;vDE=(vDE>>ROTATE_16)|(vDE<<ROTATE_16).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_31).yxwz|(v56<<ROTATE_1);v74 ^=v89;v74=(v74>>ROTATE_31).yxwz|(v74<<ROTATE_1);v45=vec4(v74.zw,v56.xy);v67=vec4(v56.zw,v74.xy);vCD=vec4(vFC.zw,vDE.xy);vEF=vec4(vDE.zw,vFC.xy);s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v01+=vec4(Z,m1)+vec4(Z,0u,u32(v01.z+vec4(Z,m1).z<v01.z));v23+=vec4(Z,m4)+vec4(Z,0u,u32(v23.z+vec4(Z,m4).z<v23.z));vCD=(vCD ^ v01).yxwz;vEF=(vEF ^ v23).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_24)|(v45<<ROTATE_8).yxwz;v67 ^=vAB;v67=(v67>>ROTATE_24)|(v67<<ROTATE_8).yxwz;s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;vCD ^=v01;vCD=(vCD>>ROTATE_16)|(vCD<<ROTATE_16).yxwz;vEF ^=v23;vEF=(vEF>>ROTATE_16)|(vEF<<ROTATE_16).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_31).yxwz|(v45<<ROTATE_1);v67 ^=vAB;v67=(v67>>ROTATE_31).yxwz|(v67<<ROTATE_1);v56=vec4(v45.zw,v67.xy);v74=vec4(v67.zw,v45.xy);vFC=vec4(vEF.zw,vCD.xy);vDE=vec4(vCD.zw,vEF.xy);s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v01+=vec4(m0,Z)+vec4(0u,u32(v01.x+vec4(m0,Z).x<v01.x),Z);vFC=(vFC ^ v01).yxwz;vDE=(vDE ^ v23).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_24)|(v56<<ROTATE_8).yxwz;v74 ^=v89;v74=(v74>>ROTATE_24)|(v74<<ROTATE_8).yxwz;s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v01+=vec4(Z,m3)+vec4(Z,0u,u32(v01.z+vec4(Z,m3).z<v01.z));s1=v23+vec4(m2,Z);v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;vFC ^=v01;vFC=(vFC>>ROTATE_16)|(vFC<<ROTATE_16).yxwz;vDE ^=v23;vDE=(vDE>>ROTATE_16)|(vDE<<ROTATE_16).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_31).yxwz|(v56<<ROTATE_1);v74 ^=v89;v74=(v74>>ROTATE_31).yxwz|(v74<<ROTATE_1);v45=vec4(v74.zw,v56.xy);v67=vec4(v56.zw,v74.xy);vCD=vec4(vFC.zw,vDE.xy);vEF=vec4(vDE.zw,vFC.xy);s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;s1=v23+vec4(Z,m3);v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;vCD=(vCD ^ v01).yxwz;vEF=(vEF ^ v23).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_24)|(v45<<ROTATE_8).yxwz;v67 ^=vAB;v67=(v67>>ROTATE_24)|(v67<<ROTATE_8).yxwz;s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v23+=vec4(m1,Z)+vec4(0u,u32(v23.x+vec4(m1,Z).x<v23.x),Z);vCD ^=v01;vCD=(vCD>>ROTATE_16)|(vCD<<ROTATE_16).yxwz;vEF ^=v23;vEF=(vEF>>ROTATE_16)|(vEF<<ROTATE_16).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_31).yxwz|(v45<<ROTATE_1);v67 ^=vAB;v67=(v67>>ROTATE_31).yxwz|(v67<<ROTATE_1);v56=vec4(v45.zw,v67.xy);v74=vec4(v67.zw,v45.xy);vFC=vec4(vEF.zw,vCD.xy);vDE=vec4(vCD.zw,vEF.xy);s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v23+=vec4(Z,m2)+vec4(Z,0u,u32(v23.z+vec4(Z,m2).z<v23.z));vFC=(vFC ^ v01).yxwz;vDE=(vDE ^ v23).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_24)|(v56<<ROTATE_8).yxwz;v74 ^=v89;v74=(v74>>ROTATE_24)|(v74<<ROTATE_8).yxwz;s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v01+=vec4(m0,m4)+vec4(0u,u32(v01.x+vec4(m0,m4).x<v01.x),0u,u32(v01.z+vec4(m0,m4).z<v01.z));vFC ^=v01;vFC=(vFC>>ROTATE_16)|(vFC<<ROTATE_16).yxwz;vDE ^=v23;vDE=(vDE>>ROTATE_16)|(vDE<<ROTATE_16).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_31).yxwz|(v56<<ROTATE_1);v74 ^=v89;v74=(v74>>ROTATE_31).yxwz|(v74<<ROTATE_1);v45=vec4(v74.zw,v56.xy);v67=vec4(v56.zw,v74.xy);vCD=vec4(vFC.zw,vDE.xy);vEF=vec4(vDE.zw,vFC.xy);s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v23+=vec4(Z,m0);v23.w+=u32(v23.z<m0.x);vCD=(vCD ^ v01).yxwz;vEF=(vEF ^ v23).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_24)|(v45<<ROTATE_8).yxwz;v67 ^=vAB;v67=(v67>>ROTATE_24)|(v67<<ROTATE_8).yxwz;s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v23+=vec4(m3,Z)+vec4(0u,u32(v23.x+vec4(m3,Z).x<v23.x),Z);vCD ^=v01;vCD=(vCD>>ROTATE_16)|(vCD<<ROTATE_16).yxwz;vEF ^=v23;vEF=(vEF>>ROTATE_16)|(vEF<<ROTATE_16).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_31).yxwz|(v45<<ROTATE_1);v67 ^=vAB;v67=(v67>>ROTATE_31).yxwz|(v67<<ROTATE_1);v56=vec4(v45.zw,v67.xy);v74=vec4(v67.zw,v45.xy);vFC=vec4(vEF.zw,vCD.xy);vDE=vec4(vCD.zw,vEF.xy);s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v23+=vec4(m1,Z)+vec4(0u,u32(v23.x+vec4(m1,Z).x<v23.x),Z);vFC=(vFC ^ v01).yxwz;vDE=(vDE ^ v23).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_24)|(v56<<ROTATE_8).yxwz;v74 ^=v89;v74=(v74>>ROTATE_24)|(v74<<ROTATE_8).yxwz;s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v01+=vec4(m2,Z)+vec4(0u,u32(v01.x+vec4(m2,Z).x<v01.x),Z);v23+=vec4(m4,Z)+vec4(0u,u32(v23.x+vec4(m4,Z).x<v23.x),Z);vFC ^=v01;vFC=(vFC>>ROTATE_16)|(vFC<<ROTATE_16).yxwz;vDE ^=v23;vDE=(vDE>>ROTATE_16)|(vDE<<ROTATE_16).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_31).yxwz|(v56<<ROTATE_1);v74 ^=v89;v74=(v74>>ROTATE_31).yxwz|(v74<<ROTATE_1);v45=vec4(v74.zw,v56.xy);v67=vec4(v56.zw,v74.xy);vCD=vec4(vFC.zw,vDE.xy);vEF=vec4(vDE.zw,vFC.xy);s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v23+=vec4(Z,m1);v23.w+=u32(v23.z<m1.x);vCD=(vCD ^ v01).yxwz;vEF=(vEF ^ v23).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_24)|(v45<<ROTATE_8).yxwz;v67 ^=vAB;v67=(v67>>ROTATE_24)|(v67<<ROTATE_8).yxwz;s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;s0=v01+vec4(m2,m4);v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;vCD ^=v01;vCD=(vCD>>ROTATE_16)|(vCD<<ROTATE_16).yxwz;vEF ^=v23;vEF=(vEF>>ROTATE_16)|(vEF<<ROTATE_16).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_31).yxwz|(v45<<ROTATE_1);v67 ^=vAB;v67=(v67>>ROTATE_31).yxwz|(v67<<ROTATE_1);v56=vec4(v45.zw,v67.xy);v74=vec4(v67.zw,v45.xy);vFC=vec4(vEF.zw,vCD.xy);vDE=vec4(vCD.zw,vEF.xy);s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v23+=vec4(m3,Z);v23.y+=u32(v23.x<m3.x);vFC=(vFC ^ v01).yxwz;vDE=(vDE ^ v23).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_24)|(v56<<ROTATE_8).yxwz;v74 ^=v89;v74=(v74>>ROTATE_24)|(v74<<ROTATE_8).yxwz;s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v23+=vec4(Z,m0);v23.w+=u32(v23.z<m0.x);vFC ^=v01;vFC=(vFC>>ROTATE_16)|(vFC<<ROTATE_16).yxwz;vDE ^=v23;vDE=(vDE>>ROTATE_16)|(vDE<<ROTATE_16).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_31).yxwz|(v56<<ROTATE_1);v74 ^=v89;v74=(v74>>ROTATE_31).yxwz|(v74<<ROTATE_1);v45=vec4(v74.zw,v56.xy);v67=vec4(v56.zw,v74.xy);vCD=vec4(vFC.zw,vDE.xy);vEF=vec4(vDE.zw,vFC.xy);s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;s0=v01+vec4(m0,m2);v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;v23+=vec4(m4,Z);v23.y+=u32(v23.x<m4.x);vCD=(vCD ^ v01).yxwz;vEF=(vEF ^ v23).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_24)|(v45<<ROTATE_8).yxwz;v67 ^=vAB;v67=(v67>>ROTATE_24)|(v67<<ROTATE_8).yxwz;s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;s0=v01+vec4(m1,m3);v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;vCD ^=v01;vCD=(vCD>>ROTATE_16)|(vCD<<ROTATE_16).yxwz;vEF ^=v23;vEF=(vEF>>ROTATE_16)|(vEF<<ROTATE_16).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_31).yxwz|(v45<<ROTATE_1);v67 ^=vAB;v67=(v67>>ROTATE_31).yxwz|(v67<<ROTATE_1);v56=vec4(v45.zw,v67.xy);v74=vec4(v67.zw,v45.xy);vFC=vec4(vEF.zw,vCD.xy);vDE=vec4(vCD.zw,vEF.xy);s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;vFC=(vFC ^ v01).yxwz;vDE=(vDE ^ v23).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_24)|(v56<<ROTATE_8).yxwz;v74 ^=v89;v74=(v74>>ROTATE_24)|(v74<<ROTATE_8).yxwz;s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;vFC ^=v01;vFC=(vFC>>ROTATE_16)|(vFC<<ROTATE_16).yxwz;vDE ^=v23;vDE=(vDE>>ROTATE_16)|(vDE<<ROTATE_16).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_31).yxwz|(v56<<ROTATE_1);v74 ^=v89;v74=(v74>>ROTATE_31).yxwz|(v74<<ROTATE_1);v45=vec4(v74.zw,v56.xy);v67=vec4(v56.zw,v74.xy);vCD=vec4(vFC.zw,vDE.xy);vEF=vec4(vDE.zw,vFC.xy);s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v01+=vec4(Z,m4);v01.w+=u32(v01.z<m4.x);vCD=(vCD ^ v01).yxwz;vEF=(vEF ^ v23).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_24)|(v45<<ROTATE_8).yxwz;v67 ^=vAB;v67=(v67>>ROTATE_24)|(v67<<ROTATE_8).yxwz;s0=v01+v45;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v67;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;vCD ^=v01;vCD=(vCD>>ROTATE_16)|(vCD<<ROTATE_16).yxwz;vEF ^=v23;vEF=(vEF>>ROTATE_16)|(vEF<<ROTATE_16).yxwz;s0=v89+vCD;v89=s0+(vec4<u32>(s0<v89)&CARRY).yxwz;s1=vAB+vEF;vAB=s1+(vec4<u32>(s1<vAB)&CARRY).yxwz;v45 ^=v89;v45=(v45>>ROTATE_31).yxwz|(v45<<ROTATE_1);v67 ^=vAB;v67=(v67>>ROTATE_31).yxwz|(v67<<ROTATE_1);v56=vec4(v45.zw,v67.xy);v74=vec4(v67.zw,v45.xy);vFC=vec4(vEF.zw,vCD.xy);vDE=vec4(vCD.zw,vEF.xy);s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;s0=v01+vec4(m1,m0);v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;vFC=(vFC ^ v01).yxwz;vDE=(vDE ^ v23).yxwz;s0=vAB+vFC;vAB=s0+(vec4<u32>(s0<vAB)&CARRY).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89)&CARRY).yxwz;v56 ^=vAB;v56=(v56>>ROTATE_24)|(v56<<ROTATE_8).yxwz;v74 ^=v89;v74=(v74>>ROTATE_24)|(v74<<ROTATE_8).yxwz;s0=v01+v56;v01=s0+(vec4<u32>(s0<v01)&CARRY).yxwz;s1=v23+v74;v23=s1+(vec4<u32>(s1<v23)&CARRY).yxwz;v01+=vec4(Z,m2);v01.w+=u32(v01.z<m2.x);v23+=vec4(Z,m3);v23.w+=u32(v23.z<m3.x);vDE ^=v23;vDE=(vDE>>ROTATE_16)|(vDE<<ROTATE_16).yxwz;s1=v89+vDE;v89=s1+(vec4<u32>(s1<v89).yxwz&CARRY);if((BLAKE2B_IV_0.y ^ v01.y ^ v89.y)>ubo.threshold&&atomicLoad(&work.found)==0u){atomicStore(&work.found,1u);work.nonce=m0;}return;}";
+
+// src/shaders/gl-downsample.ts
+var NanoPowGlDownsampleShader = `#version 300 es
+#pragma vscode_glsllint_stage: frag
+#ifdef GL_FRAGMENT_PRECISION_HIGH
+precision highp float;
+#else
+precision mediump float;
+#endif
+precision highp int;
+
+out uvec4 nonce;
+
+// source texture to be downsampled
+uniform highp usampler2D src;
+
+void main() {
+ nonce = uvec4(0u);
+ vec2 inputSize = vec2(textureSize(src, 0));
+ vec2 texel = vec2(1.0) / inputSize;
+ vec2 blockCoord = (floor(gl_FragCoord.xy) * 2.0 + vec2(0.5)) / inputSize;
+
+ uvec4 pixel = texture(src, blockCoord);
+ nonce = pixel.x == 0u ? nonce : pixel;
+
+ pixel = texture(src, blockCoord + vec2(texel.x, 0.0));
+ nonce = pixel.x == 0u ? nonce : pixel;
+
+ pixel = texture(src, blockCoord + vec2(0.0, texel.y));
+ nonce = pixel.x == 0u ? nonce : pixel;
+
+ pixel = texture(src, blockCoord + vec2(texel.x, texel.y));
+ nonce = pixel.x == 0u ? nonce : pixel;
+}
+`;
+
+// src/shaders/gl-draw.ts
+var NanoPowGlDrawShader = `#version 300 es
+#pragma vscode_glsllint_stage: frag
+#ifdef GL_FRAGMENT_PRECISION_HIGH
+precision highp float;
+#else
+precision mediump float;
+#endif
+
+out uvec4 nonce;
+
+// blockhash - Array of precalculated block hash components
+// threshold - 0xfffffff8 for send/change blocks, 0xfffffe00 for all else
+// search - Checks all pixels if true, else only checks 1 pixel to validate
+layout(std140) uniform UBO {
+ uint blockhash[8];
+ uint threshold;
+ bool search;
+};
+
+// Random work seed values
+layout(std140) uniform WORK {
+ uvec2 seed;
+};
+
+// Defined separately from uint v[0].y below as the original value is required
+// to calculate the second uint32 of the digest for threshold comparison
+const uint BLAKE2B_IV32_1 = 0x6A09E667u;
+
+// Used during G for vector bit rotations
+const uvec4 ROTATE_1 = uvec4(1u);
+const uvec4 ROTATE_8 = uvec4(8u);
+const uvec4 ROTATE_16 = uvec4(16u);
+const uvec4 ROTATE_24 = uvec4(24u);
+const uvec4 ROTATE_31 = uvec4(31u);
+
+// Both buffers represent 16 uint64s as 32 uint32s
+// because that's what GLSL offers, just like Javascript
+
+// Compression buffer, intialized to 2 instances of the initialization vector
+// The following values have been modified from the BLAKE2B_IV:
+// OUTLEN is constant 8 bytes
+// v[0] ^= 0x01010000u ^ uint(OUTLEN);
+// INLEN is constant 40 bytes: work value (8) + block hash (32)
+// v[12] ^= uint(INLEN);
+// It's always the "last" compression at this INLEN
+// v[14] = ~v[14];
+const uvec2 blake2b_iv[16] = uvec2[16](
+ uvec2(0xF2BDC900u, 0x6A09E667u),
+ uvec2(0x84CAA73Bu, 0xBB67AE85u),
+ uvec2(0xFE94F82Bu, 0x3C6EF372u),
+ uvec2(0x5F1D36F1u, 0xA54FF53Au),
+ uvec2(0xADE682D1u, 0x510E527Fu),
+ uvec2(0x2B3E6C1Fu, 0x9B05688Cu),
+ uvec2(0xFB41BD6Bu, 0x1F83D9ABu),
+ uvec2(0x137E2179u, 0x5BE0CD19u),
+ uvec2(0xF3BCC908u, 0x6A09E667u),
+ uvec2(0x84CAA73Bu, 0xBB67AE85u),
+ uvec2(0xFE94F82Bu, 0x3C6EF372u),
+ uvec2(0x5F1D36F1u, 0xA54FF53Au),
+ uvec2(0xADE682F9u, 0x510E527Fu),
+ uvec2(0x2B3E6C1Fu, 0x9B05688Cu),
+ uvec2(0x04BE4294u, 0xE07C2654u),
+ uvec2(0x137E2179u, 0x5BE0CD19u)
+);
+
+// Iterated initialization vector
+uvec2 v[16];
+
+// Input data buffer
+uvec2 m[16];
+
+// G mixing function, compressing two subprocesses into one
+void G (
+ uint a0, uint b0, uint c0, uint d0, uvec2 x0, uvec2 y0,
+ uint a1, uint b1, uint c1, uint d1, uvec2 x1, uvec2 y1
+) {
+ uvec4 a = uvec4(v[a0], v[a1]);
+ uvec4 b = uvec4(v[b0], v[b1]);
+ uvec4 c = uvec4(v[c0], v[c1]);
+ uvec4 d = uvec4(v[d0], v[d1]);
+ uvec4 mx = uvec4(x0, x1);
+ uvec4 my = uvec4(y0, y1);
+
+ a = a + b + uvec4(0u, uint(a.x + b.x < a.x), 0u, uint(a.z + b.z < a.z));
+ a = a + mx + uvec4(0u, uint(a.x + mx.x < a.x), 0u, uint(a.z + mx.z < a.z));
+ d = (d ^ a).yxwz;
+ c = c + d + uvec4(0u, uint(c.x + d.x < c.x), 0u, uint(c.z + d.z < c.z));
+ b = ((b ^ c) >> ROTATE_24) | ((b ^ c) << ROTATE_8).yxwz;
+ a = a + b + uvec4(0u, uint(a.x + b.x < b.x), 0u, uint(a.z + b.z < b.z));
+ a = a + my + uvec4(0u, uint(a.x + my.x < a.x), 0u, uint(a.z + my.z < a.z));
+ d = ((d ^ a) >> ROTATE_16) | ((d ^ a) << ROTATE_16).yxwz;
+ c = c + d + uvec4(0u, uint(c.x + d.x < c.x), 0u, uint(c.z + d.z < c.z));
+ b = ((b ^ c) >> ROTATE_31).yxwz | ((b ^ c) << ROTATE_1);
+
+ v[a0] = a.xy;
+ v[b0] = b.xy;
+ v[c0] = c.xy;
+ v[d0] = d.xy;
+ v[a1] = a.zw;
+ v[b1] = b.zw;
+ v[c1] = c.zw;
+ v[d1] = d.zw;
+}
+
+void main() {
+ // Initialize fragment output
+ nonce = uvec4(0u);
+
+ // Nonce uniquely differentiated by pixel location
+ m[0u] = seed ^ uvec2(gl_FragCoord);
+
+ // Block hash
+ m[1u] = uvec2(blockhash[0u], blockhash[1u]);
+ m[2u] = uvec2(blockhash[2u], blockhash[3u]);
+ m[3u] = uvec2(blockhash[4u], blockhash[5u]);
+ m[4u] = uvec2(blockhash[6u], blockhash[7u]);
+
+ // Reset v
+ v = blake2b_iv;
+
+ // Twelve rounds of G mixing
+
+ // Round 0
+ G(0u, 4u, 8u, 12u, m[0u], m[1u], 1u, 5u, 9u, 13u, m[2u], m[3u]);
+ G(2u, 6u, 10u, 14u, m[4u], m[5u], 3u, 7u, 11u, 15u, m[6u], m[7u]);
+ G(0u, 5u, 10u, 15u, m[8u], m[9u], 1u, 6u, 11u, 12u, m[10u], m[11u]);
+ G(2u, 7u, 8u, 13u, m[12u], m[13u], 3u, 4u, 9u, 14u, m[14u], m[15u]);
+
+ // Round 1
+ G(0u, 4u, 8u, 12u, m[14u], m[10u], 1u, 5u, 9u, 13u, m[4u], m[8u]);
+ G(2u, 6u, 10u, 14u, m[9u], m[15u], 3u, 7u, 11u, 15u, m[13u], m[6u]);
+ G(0u, 5u, 10u, 15u, m[1u], m[12u], 1u, 6u, 11u, 12u, m[0u], m[2u]);
+ G(2u, 7u, 8u, 13u, m[11u], m[7u], 3u, 4u, 9u, 14u, m[5u], m[3u]);
+
+ // Round 2
+ G(0u, 4u, 8u, 12u, m[11u], m[8u], 1u, 5u, 9u, 13u, m[12u], m[0u]);
+ G(2u, 6u, 10u, 14u, m[5u], m[2u], 3u, 7u, 11u, 15u, m[15u], m[13u]);
+ G(0u, 5u, 10u, 15u, m[10u], m[14u], 1u, 6u, 11u, 12u, m[3u], m[6u]);
+ G(2u, 7u, 8u, 13u, m[7u], m[1u], 3u, 4u, 9u, 14u, m[9u], m[4u]);
+
+ // Round 3
+ G(0u, 4u, 8u, 12u, m[7u], m[9u], 1u, 5u, 9u, 13u, m[3u], m[1u]);
+ G(2u, 6u, 10u, 14u, m[13u], m[12u], 3u, 7u, 11u, 15u, m[11u], m[14u]);
+ G(0u, 5u, 10u, 15u, m[2u], m[6u], 1u, 6u, 11u, 12u, m[5u], m[10u]);
+ G(2u, 7u, 8u, 13u, m[4u], m[0u], 3u, 4u, 9u, 14u, m[15u], m[8u]);
+
+ // Round 4
+ G(0u, 4u, 8u, 12u, m[9u], m[0u], 1u, 5u, 9u, 13u, m[5u], m[7u]);
+ G(2u, 6u, 10u, 14u, m[2u], m[4u], 3u, 7u, 11u, 15u, m[10u], m[15u]);
+ G(0u, 5u, 10u, 15u, m[14u], m[1u], 1u, 6u, 11u, 12u, m[11u], m[12u]);
+ G(2u, 7u, 8u, 13u, m[6u], m[8u], 3u, 4u, 9u, 14u, m[3u], m[13u]);
+
+ // Round 5
+ G(0u, 4u, 8u, 12u, m[2u], m[12u], 1u, 5u, 9u, 13u, m[6u], m[10u]);
+ G(2u, 6u, 10u, 14u, m[0u], m[11u], 3u, 7u, 11u, 15u, m[8u], m[3u]);
+ G(0u, 5u, 10u, 15u, m[4u], m[13u], 1u, 6u, 11u, 12u, m[7u], m[5u]);
+ G(2u, 7u, 8u, 13u, m[15u], m[14u], 3u, 4u, 9u, 14u, m[1u], m[9u]);
+
+ // Round 6
+ G(0u, 4u, 8u, 12u, m[12u], m[5u], 1u, 5u, 9u, 13u, m[1u], m[15u]);
+ G(2u, 6u, 10u, 14u, m[14u], m[13u], 3u, 7u, 11u, 15u, m[4u], m[10u]);
+ G(0u, 5u, 10u, 15u, m[0u], m[7u], 1u, 6u, 11u, 12u, m[6u], m[3u]);
+ G(2u, 7u, 8u, 13u, m[9u], m[2u], 3u, 4u, 9u, 14u, m[8u], m[11u]);
+
+ // Round 7
+ G(0u, 4u, 8u, 12u, m[13u], m[11u], 1u, 5u, 9u, 13u, m[7u], m[14u]);
+ G(2u, 6u, 10u, 14u, m[12u], m[1u], 3u, 7u, 11u, 15u, m[3u], m[9u]);
+ G(0u, 5u, 10u, 15u, m[5u], m[0u], 1u, 6u, 11u, 12u, m[15u], m[4u]);
+ G(2u, 7u, 8u, 13u, m[8u], m[6u], 3u, 4u, 9u, 14u, m[2u], m[10u]);
+
+ // Round 8
+ G(0u, 4u, 8u, 12u, m[6u], m[15u], 1u, 5u, 9u, 13u, m[14u], m[9u]);
+ G(2u, 6u, 10u, 14u, m[11u], m[3u], 3u, 7u, 11u, 15u, m[0u], m[8u]);
+ G(0u, 5u, 10u, 15u, m[12u], m[2u], 1u, 6u, 11u, 12u, m[13u], m[7u]);
+ G(2u, 7u, 8u, 13u, m[1u], m[4u], 3u, 4u, 9u, 14u, m[10u], m[5u]);
+
+ // Round 9
+ G(0u, 4u, 8u, 12u, m[10u], m[2u], 1u, 5u, 9u, 13u, m[8u], m[4u]);
+ G(2u, 6u, 10u, 14u, m[7u], m[6u], 3u, 7u, 11u, 15u, m[1u], m[5u]);
+ G(0u, 5u, 10u, 15u, m[15u], m[11u], 1u, 6u, 11u, 12u, m[9u], m[14u]);
+ G(2u, 7u, 8u, 13u, m[3u], m[12u], 3u, 4u, 9u, 14u, m[13u], m[0u]);
+
+ // Round 10
+ G(0u, 4u, 8u, 12u, m[0u], m[1u], 1u, 5u, 9u, 13u, m[2u], m[3u]);
+ G(2u, 6u, 10u, 14u, m[4u], m[5u], 3u, 7u, 11u, 15u, m[6u], m[7u]);
+ G(0u, 5u, 10u, 15u, m[8u], m[9u], 1u, 6u, 11u, 12u, m[10u], m[11u]);
+ G(2u, 7u, 8u, 13u, m[12u], m[13u], 3u, 4u, 9u, 14u, m[14u], m[15u]);
+
+ // Round 11
+ G(0u, 4u, 8u, 12u, m[14u], m[10u], 1u, 5u, 9u, 13u, m[4u], m[8u]);
+ G(2u, 6u, 10u, 14u, m[9u], m[15u], 3u, 7u, 11u, 15u, m[13u], m[6u]);
+ G(0u, 5u, 10u, 15u, m[1u], m[12u], 1u, 6u, 11u, 12u, m[0u], m[2u]);
+ G(2u, 7u, 8u, 13u, m[11u], m[7u], 3u, 4u, 9u, 14u, m[5u], m[3u]);
+
+ // Pixel data set from work seed values
+ // Finalize digest from high bits, low bits can be safely ignored
+ if ((BLAKE2B_IV32_1 ^ v[0u].y ^ v[8u].y) >= threshold && (search || uvec2(gl_FragCoord) == uvec2(0u))) {
+ nonce = uvec4(1u, m[0u].y, m[0u].x, (uint(gl_FragCoord.x) << 16u) | uint(gl_FragCoord.y));
+ }
+
+ // Valid nonce not found
+ if (nonce.x == 0u) {
+ discard;
+ }
+}
+`;
+
+// src/shaders/gl-vertex.ts
+var NanoPowGlVertexShader = `#version 300 es
+#pragma vscode_glsllint_stage: vert
+#ifdef GL_FRAGMENT_PRECISION_HIGH
+precision highp float;
+#else
+precision mediump float;
+#endif
+
+layout (location=0) in vec4 position;
+
+void main() {
+ gl_Position = position;
+}
+`;
+
+// src/classes/gl.ts
+var NanoPowGl = class _NanoPowGl {
+ static #busy = false;
+ static #debug = false;
+ static #raf = 0;
+ /** Used to set canvas size. */
+ static #cores = Math.max(1, Math.floor(navigator.hardwareConcurrency));
+ static #WORKLOAD = 256 * this.#cores;
+ static #canvas = new OffscreenCanvas(this.#WORKLOAD, this.#WORKLOAD);
+ static get size() {
+ return this.#gl?.drawingBufferWidth;
+ }
+ static #gl;
+ static #drawProgram;
+ static #downsampleProgram;
+ static #vertexShader;
+ static #drawShader;
+ static #downsampleShader;
+ static #positionBuffer;
+ static #drawFbo;
+ static #downsampleFbos = [];
+ static #downsampleSrcLocation;
+ static #uboBuffer;
+ static #uboView = new DataView(new ArrayBuffer(144));
+ static #seedBuffer;
+ static #seed = new BigUint64Array(1);
+ static #query;
+ static #pixels;
+ /**Vertex Positions, 2 triangles */
+ static #positions = new Float32Array([
+ -1,
+ -1,
+ 1,
+ -1,
+ 1,
+ 1,
+ -1,
+ 1
+ ]);
+ /** Compile */
+ static async init() {
+ if (this.#busy) return;
+ this.#busy = true;
+ try {
+ this.#canvas.addEventListener("webglcontextlost", (event) => {
+ event.preventDefault();
+ console.warn("WebGL context lost. Waiting for it to be restored...");
+ cancelAnimationFrame(this.#raf);
+ }, false);
+ this.#canvas.addEventListener("webglcontextrestored", (event) => {
+ console.warn("WebGL context restored. Reinitializing...");
+ _NanoPowGl.init();
+ }, false);
+ this.#gl = this.#canvas.getContext("webgl2");
+ if (this.#gl == null) throw new Error("WebGL 2 is required");
+ this.#drawProgram = this.#gl.createProgram();
+ if (this.#drawProgram == null) throw new Error("Failed to create shader program");
+ this.#vertexShader = this.#gl.createShader(this.#gl.VERTEX_SHADER);
+ if (this.#vertexShader == null) throw new Error("Failed to create vertex shader");
+ this.#gl.shaderSource(this.#vertexShader, NanoPowGlVertexShader);
+ this.#gl.compileShader(this.#vertexShader);
+ if (!this.#gl.getShaderParameter(this.#vertexShader, this.#gl.COMPILE_STATUS))
+ throw new Error(this.#gl.getShaderInfoLog(this.#vertexShader) ?? `Failed to compile vertex shader`);
+ this.#drawShader = this.#gl.createShader(this.#gl.FRAGMENT_SHADER);
+ if (this.#drawShader == null) throw new Error("Failed to create fragment shader");
+ this.#gl.shaderSource(this.#drawShader, NanoPowGlDrawShader);
+ this.#gl.compileShader(this.#drawShader);
+ if (!this.#gl.getShaderParameter(this.#drawShader, this.#gl.COMPILE_STATUS))
+ throw new Error(this.#gl.getShaderInfoLog(this.#drawShader) ?? `Failed to compile fragment shader`);
+ this.#gl.attachShader(this.#drawProgram, this.#vertexShader);
+ this.#gl.attachShader(this.#drawProgram, this.#drawShader);
+ this.#gl.linkProgram(this.#drawProgram);
+ if (!this.#gl.getProgramParameter(this.#drawProgram, this.#gl.LINK_STATUS))
+ throw new Error(this.#gl.getProgramInfoLog(this.#drawProgram) ?? `Failed to link program`);
+ this.#downsampleProgram = this.#gl.createProgram();
+ if (this.#downsampleProgram == null) throw new Error("Failed to create downsample program");
+ this.#downsampleShader = this.#gl.createShader(this.#gl.FRAGMENT_SHADER);
+ if (this.#downsampleShader == null) throw new Error("Failed to create downsample shader");
+ this.#gl.shaderSource(this.#downsampleShader, NanoPowGlDownsampleShader);
+ this.#gl.compileShader(this.#downsampleShader);
+ if (!this.#gl.getShaderParameter(this.#downsampleShader, this.#gl.COMPILE_STATUS))
+ throw new Error(this.#gl.getShaderInfoLog(this.#downsampleShader) ?? `Failed to compile downsample shader`);
+ this.#gl.attachShader(this.#downsampleProgram, this.#vertexShader);
+ this.#gl.attachShader(this.#downsampleProgram, this.#downsampleShader);
+ this.#gl.linkProgram(this.#downsampleProgram);
+ if (!this.#gl.getProgramParameter(this.#downsampleProgram, this.#gl.LINK_STATUS))
+ throw new Error(this.#gl.getProgramInfoLog(this.#downsampleProgram) ?? `Failed to link program`);
+ this.#gl.useProgram(this.#drawProgram);
+ const triangleArray = this.#gl.createVertexArray();
+ this.#gl.bindVertexArray(triangleArray);
+ this.#positionBuffer = this.#gl.createBuffer();
+ this.#gl.bindBuffer(this.#gl.ARRAY_BUFFER, this.#positionBuffer);
+ this.#gl.bufferData(this.#gl.ARRAY_BUFFER, this.#positions, this.#gl.STATIC_DRAW);
+ this.#gl.vertexAttribPointer(0, 2, this.#gl.FLOAT, false, 0, 0);
+ this.#gl.enableVertexAttribArray(0);
+ this.#gl.bindBuffer(this.#gl.ARRAY_BUFFER, null);
+ const texture = this.#gl.createTexture();
+ this.#gl.bindTexture(this.#gl.TEXTURE_2D, texture);
+ this.#gl.texImage2D(this.#gl.TEXTURE_2D, 0, this.#gl.RGBA32UI, this.#gl.drawingBufferWidth, this.#gl.drawingBufferHeight, 0, this.#gl.RGBA_INTEGER, this.#gl.UNSIGNED_INT, null);
+ this.#gl.texParameteri(this.#gl.TEXTURE_2D, this.#gl.TEXTURE_MIN_FILTER, this.#gl.NEAREST);
+ this.#gl.texParameteri(this.#gl.TEXTURE_2D, this.#gl.TEXTURE_MAG_FILTER, this.#gl.NEAREST);
+ const framebuffer = this.#gl.createFramebuffer();
+ this.#gl.bindFramebuffer(this.#gl.FRAMEBUFFER, framebuffer);
+ this.#gl.framebufferTexture2D(this.#gl.FRAMEBUFFER, this.#gl.COLOR_ATTACHMENT0, this.#gl.TEXTURE_2D, texture, 0);
+ if (this.#gl.checkFramebufferStatus(this.#gl.FRAMEBUFFER) !== this.#gl.FRAMEBUFFER_COMPLETE)
+ throw new Error(`Failed to create drawing framebuffer`);
+ this.#drawFbo = { texture, framebuffer, size: { x: this.#gl.drawingBufferWidth, y: this.#gl.drawingBufferHeight } };
+ for (let i = 1; i <= 4; i++) {
+ const width = this.#gl.drawingBufferWidth / 2 ** i;
+ const height = this.#gl.drawingBufferHeight / 2 ** i;
+ const texture2 = this.#gl.createTexture();
+ this.#gl.bindTexture(this.#gl.TEXTURE_2D, texture2);
+ this.#gl.texImage2D(this.#gl.TEXTURE_2D, 0, this.#gl.RGBA32UI, width, height, 0, this.#gl.RGBA_INTEGER, this.#gl.UNSIGNED_INT, null);
+ this.#gl.texParameteri(this.#gl.TEXTURE_2D, this.#gl.TEXTURE_MIN_FILTER, this.#gl.NEAREST);
+ this.#gl.texParameteri(this.#gl.TEXTURE_2D, this.#gl.TEXTURE_MAG_FILTER, this.#gl.NEAREST);
+ const framebuffer2 = this.#gl.createFramebuffer();
+ this.#gl.bindFramebuffer(this.#gl.FRAMEBUFFER, framebuffer2);
+ this.#gl.framebufferTexture2D(this.#gl.FRAMEBUFFER, this.#gl.COLOR_ATTACHMENT0, this.#gl.TEXTURE_2D, texture2, 0);
+ if (this.#gl.checkFramebufferStatus(this.#gl.FRAMEBUFFER) !== this.#gl.FRAMEBUFFER_COMPLETE)
+ throw new Error(`Failed to create downsampling framebuffer ${i}`);
+ this.#downsampleFbos.push({ texture: texture2, framebuffer: framebuffer2, size: { x: width, y: height } });
+ }
+ this.#downsampleSrcLocation = this.#gl.getUniformLocation(this.#downsampleProgram, "src");
+ this.#gl.bindTexture(this.#gl.TEXTURE_2D, null);
+ this.#gl.bindFramebuffer(this.#gl.FRAMEBUFFER, null);
+ this.#uboBuffer = this.#gl.createBuffer();
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, this.#uboBuffer);
+ this.#gl.bufferData(this.#gl.UNIFORM_BUFFER, 144, this.#gl.DYNAMIC_DRAW);
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, null);
+ this.#gl.bindBufferBase(this.#gl.UNIFORM_BUFFER, 0, this.#uboBuffer);
+ this.#gl.uniformBlockBinding(this.#drawProgram, this.#gl.getUniformBlockIndex(this.#drawProgram, "UBO"), 0);
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, null);
+ this.#seedBuffer = this.#gl.createBuffer();
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, this.#seedBuffer);
+ this.#gl.bufferData(this.#gl.UNIFORM_BUFFER, 16, this.#gl.DYNAMIC_DRAW);
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, null);
+ this.#gl.bindBufferBase(this.#gl.UNIFORM_BUFFER, 1, this.#seedBuffer);
+ this.#gl.uniformBlockBinding(this.#drawProgram, this.#gl.getUniformBlockIndex(this.#drawProgram, "WORK"), 1);
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, null);
+ this.#query = this.#gl.createQuery();
+ this.#pixels = new Uint32Array(this.#gl.drawingBufferWidth * this.#gl.drawingBufferHeight * 4);
+ console.log(`NanoPow WebGL initialized at ${this.#gl.drawingBufferWidth}x${this.#gl.drawingBufferHeight}. Maximum nonces checked per frame: ${this.#gl.drawingBufferWidth * this.#gl.drawingBufferHeight}`);
+ } catch (err) {
+ throw new Error("WebGL initialization failed.", { cause: err });
+ } finally {
+ this.#busy = false;
+ }
+ }
+ static reset() {
+ cancelAnimationFrame(_NanoPowGl.#raf);
+ _NanoPowGl.#gl?.deleteQuery(_NanoPowGl.#query);
+ _NanoPowGl.#query = null;
+ _NanoPowGl.#gl?.deleteBuffer(_NanoPowGl.#seedBuffer);
+ _NanoPowGl.#seedBuffer = null;
+ _NanoPowGl.#gl?.deleteBuffer(_NanoPowGl.#uboBuffer);
+ _NanoPowGl.#uboBuffer = null;
+ for (const fbo of _NanoPowGl.#downsampleFbos) {
+ _NanoPowGl.#gl?.deleteFramebuffer(fbo.framebuffer);
+ _NanoPowGl.#gl?.deleteTexture(fbo.texture);
+ }
+ _NanoPowGl.#downsampleFbos = [];
+ _NanoPowGl.#gl?.deleteShader(_NanoPowGl.#downsampleShader);
+ _NanoPowGl.#downsampleShader = null;
+ _NanoPowGl.#gl?.deleteProgram(_NanoPowGl.#downsampleProgram);
+ _NanoPowGl.#downsampleProgram = null;
+ _NanoPowGl.#gl?.deleteFramebuffer(_NanoPowGl.#drawFbo?.framebuffer ?? null);
+ _NanoPowGl.#drawFbo = null;
+ _NanoPowGl.#gl?.deleteTexture(_NanoPowGl.#drawFbo);
+ _NanoPowGl.#drawFbo = null;
+ _NanoPowGl.#gl?.deleteBuffer(_NanoPowGl.#positionBuffer);
+ _NanoPowGl.#positionBuffer = null;
+ _NanoPowGl.#gl?.deleteShader(_NanoPowGl.#drawShader);
+ _NanoPowGl.#drawShader = null;
+ _NanoPowGl.#gl?.deleteShader(_NanoPowGl.#vertexShader);
+ _NanoPowGl.#vertexShader = null;
+ _NanoPowGl.#gl?.deleteProgram(_NanoPowGl.#drawProgram);
+ _NanoPowGl.#drawProgram = null;
+ _NanoPowGl.#gl = null;
+ _NanoPowGl.#busy = false;
+ _NanoPowGl.init();
+ }
+ static #logAverages(times) {
+ let count = times.length, sum = 0, reciprocals = 0, logarithms = 0, truncated = 0, min = 65535, max = 0, rate = 0;
+ times.sort();
+ for (let i = 0; i < count; i++) {
+ sum += times[i];
+ reciprocals += 1 / times[i];
+ logarithms += Math.log(times[i]);
+ min = Math.min(min, times[i]);
+ max = Math.max(max, times[i]);
+ if (count < 3 || i > count * 0.1 && i < count * 0.9) truncated += times[i];
+ }
+ const averages = {
+ "Count (frames)": count,
+ "Total (ms)": sum,
+ "Rate (f/s)": 1e3 * count * 0.8 / (truncated || sum),
+ "Minimum (ms)": min,
+ "Maximum (ms)": max,
+ "Arithmetic Mean (ms)": sum / count,
+ "Truncated Mean (ms)": truncated / count,
+ "Harmonic Mean (ms)": count / reciprocals,
+ "Geometric Mean (ms)": Math.exp(logarithms / count)
+ };
+ console.log(`Averages: ${JSON.stringify(averages)}`);
+ console.table(averages);
+ }
+ static #draw(seed) {
+ if (this.#gl == null || this.#query == null) throw new Error("WebGL 2 is required to draw and query pixels");
+ if (this.#drawFbo == null) throw new Error("FBO is required to draw");
+ if (this.#seed[0] == null || this.#seedBuffer == null) throw new Error("Seed is required to draw");
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, this.#seedBuffer);
+ this.#gl.bufferSubData(this.#gl.UNIFORM_BUFFER, 0, seed);
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, null);
+ this.#gl.useProgram(this.#drawProgram);
+ this.#gl.bindFramebuffer(this.#gl.FRAMEBUFFER, this.#drawFbo.framebuffer);
+ this.#gl.activeTexture(this.#gl.TEXTURE0);
+ this.#gl.bindTexture(this.#gl.TEXTURE_2D, this.#drawFbo.texture);
+ this.#gl.beginQuery(this.#gl.ANY_SAMPLES_PASSED_CONSERVATIVE, this.#query);
+ this.#gl.viewport(0, 0, this.#drawFbo.size.x, this.#drawFbo.size.y);
+ this.#gl.drawArrays(this.#gl.TRIANGLES, 0, 4);
+ this.#gl.bindFramebuffer(this.#gl.FRAMEBUFFER, null);
+ this.#gl.endQuery(this.#gl.ANY_SAMPLES_PASSED_CONSERVATIVE);
+ }
+ static async #checkQueryResult() {
+ return new Promise((resolve, reject) => {
+ function check() {
+ try {
+ if (_NanoPowGl.#gl == null || _NanoPowGl.#query == null) throw new Error("WebGL 2 is required to check query results");
+ if (_NanoPowGl.#gl.getQueryParameter(_NanoPowGl.#query, _NanoPowGl.#gl.QUERY_RESULT_AVAILABLE)) {
+ resolve(!!_NanoPowGl.#gl.getQueryParameter(_NanoPowGl.#query, _NanoPowGl.#gl.QUERY_RESULT));
+ } else {
+ _NanoPowGl.#raf = requestAnimationFrame(check);
+ }
+ } catch (err) {
+ reject(err);
+ }
+ }
+ check();
+ });
+ }
+ /**
+ * Reads pixels into the work buffer, checks every 4th pixel for the 'found'
+ * byte, converts the subsequent 3 pixels with the nonce byte values to a hex
+ * string, and returns the result.
+ *
+ * @param workHex - Original nonce if provided for a validation call
+ * @returns Nonce as an 8-byte (16-char) hexadecimal string
+ */
+ static #readResult(workHex) {
+ if (this.#gl == null) throw new Error("WebGL 2 is required to read pixels");
+ if (this.#drawFbo == null) throw new Error("Source FBO is required to downsample");
+ let source = this.#drawFbo;
+ let pixelCount;
+ const start = performance.now();
+ if (workHex != null) {
+ this.#gl.bindFramebuffer(this.#gl.FRAMEBUFFER, source.framebuffer);
+ this.#gl.readPixels(0, 0, 1, 1, this.#gl.RGBA_INTEGER, this.#gl.UNSIGNED_INT, this.#pixels);
+ pixelCount = 4;
+ } else {
+ this.#gl.useProgram(this.#downsampleProgram);
+ for (const fbo of this.#downsampleFbos) {
+ this.#gl.bindFramebuffer(this.#gl.FRAMEBUFFER, fbo.framebuffer);
+ this.#gl.activeTexture(this.#gl.TEXTURE0);
+ this.#gl.bindTexture(this.#gl.TEXTURE_2D, source.texture);
+ this.#gl.uniform1i(this.#downsampleSrcLocation, 0);
+ this.#gl.viewport(0, 0, fbo.size.x, fbo.size.y);
+ this.#gl.drawArrays(this.#gl.TRIANGLES, 0, 4);
+ source = fbo;
+ }
+ this.#gl.bindFramebuffer(this.#gl.FRAMEBUFFER, source.framebuffer);
+ this.#gl.readPixels(0, 0, source.size.x, source.size.y, this.#gl.RGBA_INTEGER, this.#gl.UNSIGNED_INT, this.#pixels);
+ pixelCount = source.size.x * source.size.y * 4;
+ }
+ this.#gl.bindFramebuffer(this.#gl.FRAMEBUFFER, null);
+ for (let i = 0; i < pixelCount; i += 4) {
+ if (this.#pixels[i] !== 0) {
+ if (this.#debug) console.log(`readResults (${performance.now() - start} ms)`);
+ if (this.#debug) console.log(`Pixel: rgba(${this.#pixels[i]}, ${this.#pixels[i + 1]}, ${this.#pixels[i + 2]}, ${this.#pixels[i + 3].toString(16).padStart(8, "0")})`);
+ const hex = `${this.#pixels[i + 1].toString(16).padStart(8, "0")}${this.#pixels[i + 2].toString(16).padStart(8, "0")}`;
+ if (workHex == null || workHex == hex) return hex;
+ }
+ }
+ throw new Error("Query reported result but nonce value not found");
+ }
+ /**
+ * Finds a nonce that satisfies the Nano proof-of-work requirements.
+ *
+ * @param {string} hash - Hexadecimal hash of previous block, or public key for new accounts
+ * @param {number} [threshold=0xfffffff8] - Difficulty of proof-of-work calculation
+ */
+ static async search(hash, options) {
+ if (this.#busy) {
+ console.log("NanoPowGl is busy. Retrying search...");
+ return new Promise((resolve) => {
+ setTimeout(async () => {
+ const result = this.search(hash, options);
+ resolve(result);
+ }, 100);
+ });
+ }
+ this.#busy = true;
+ if (!/^[A-Fa-f0-9]{64}$/.test(hash)) throw new Error(`Invalid hash ${hash}`);
+ const threshold = typeof options?.threshold !== "number" || options.threshold < 0 || options.threshold > 4294967295 ? 4294967288 : options.threshold;
+ const effort = typeof options?.effort !== "number" || options.effort < 1 || options.effort > 32 ? this.#cores : options.effort;
+ this.#debug = !!options?.debug;
+ if (this.#WORKLOAD !== 256 * effort) {
+ this.#WORKLOAD = 256 * effort;
+ this.#canvas.height = this.#WORKLOAD;
+ this.#canvas.width = this.#WORKLOAD;
+ this.reset();
+ }
+ if (_NanoPowGl.#gl == null) throw new Error("WebGL 2 is required");
+ if (this.#gl == null) throw new Error("WebGL 2 is required");
+ if (this.#drawFbo == null) throw new Error("WebGL framebuffer is required");
+ this.#gl.bindFramebuffer(this.#gl.FRAMEBUFFER, this.#drawFbo.framebuffer);
+ this.#gl.clearBufferuiv(this.#gl.COLOR, 0, [0, 0, 0, 0]);
+ this.#gl.bindFramebuffer(this.#gl.FRAMEBUFFER, null);
+ for (let i = 0; i < this.#uboView.byteLength; i++) this.#uboView.setUint8(i, 0);
+ for (let i = 0; i < 64; i += 8) {
+ const uint32 = hash.slice(i, i + 8);
+ this.#uboView.setUint32(i * 2, parseInt(uint32, 16));
+ }
+ this.#uboView.setUint32(128, threshold, true);
+ this.#uboView.setUint32(132, 1, true);
+ if (this.#debug) console.log("UBO", this.#uboView.buffer.slice(0));
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, this.#uboBuffer);
+ this.#gl.bufferSubData(this.#gl.UNIFORM_BUFFER, 0, this.#uboView);
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, null);
+ let times = [];
+ let start = performance.now();
+ let nonce = null;
+ if (this.#debug) console.groupCollapsed("Seeds (click to view)");
+ while (nonce == null) {
+ start = performance.now();
+ const random0 = Math.floor(Math.random() * 4294967295);
+ const random1 = Math.floor(Math.random() * 4294967295);
+ this.#seed[0] = BigInt(random0) << 32n | BigInt(random1);
+ if (this.#debug) console.log("Seed", this.#seed);
+ this.#draw(this.#seed);
+ const found = await this.#checkQueryResult();
+ times.push(performance.now() - start);
+ if (found) {
+ if (this.#debug) console.groupEnd();
+ nonce = this.#readResult();
+ }
+ }
+ this.#busy = false;
+ if (this.#debug) this.#logAverages(times);
+ return nonce;
+ }
+ /**
+ * Validates that a nonce satisfies Nano proof-of-work requirements.
+ *
+ * @param {string} work - Hexadecimal proof-of-work value to validate
+ * @param {string} hash - Hexadecimal hash of previous block, or public key for new accounts
+ * @param {number} [threshold=0xfffffff8] - Difficulty of proof-of-work calculation
+ */
+ static async validate(work, hash, options) {
+ if (this.#busy) {
+ console.log("NanoPowGl is busy. Retrying validate...");
+ return new Promise((resolve) => {
+ setTimeout(async () => {
+ const result = this.validate(work, hash, options);
+ resolve(result);
+ }, 100);
+ });
+ }
+ this.#busy = true;
+ if (!/^[A-Fa-f0-9]{16}$/.test(work)) throw new Error(`Invalid work ${work}`);
+ if (!/^[A-Fa-f0-9]{64}$/.test(hash)) throw new Error(`Invalid hash ${hash}`);
+ const threshold = typeof options?.threshold !== "number" || options.threshold < 0 || options.threshold > 4294967295 ? 4294967288 : options.threshold;
+ this.#debug = !!options?.debug;
+ if (_NanoPowGl.#gl == null) throw new Error("WebGL 2 is required");
+ if (this.#gl == null) throw new Error("WebGL 2 is required");
+ if (this.#drawFbo == null) throw new Error("WebGL framebuffer is required");
+ this.#gl.bindFramebuffer(this.#gl.FRAMEBUFFER, this.#drawFbo.framebuffer);
+ this.#gl.clearBufferuiv(this.#gl.COLOR, 0, [0, 0, 0, 0]);
+ this.#gl.bindFramebuffer(this.#gl.FRAMEBUFFER, null);
+ for (let i = 0; i < this.#uboView.byteLength; i++) this.#uboView.setUint8(i, 0);
+ for (let i = 0; i < 64; i += 8) {
+ const uint32 = hash.slice(i, i + 8);
+ this.#uboView.setUint32(i * 2, parseInt(uint32, 16));
+ }
+ this.#uboView.setUint32(128, threshold, true);
+ this.#uboView.setUint32(132, 0, true);
+ if (this.#debug) console.log("UBO", this.#uboView.buffer.slice(0));
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, this.#uboBuffer);
+ this.#gl.bufferSubData(this.#gl.UNIFORM_BUFFER, 0, this.#uboView);
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, null);
+ let nonce = null;
+ this.#seed[0] = BigInt(`0x${work}`);
+ if (this.#debug) console.log("Work", this.#seed);
+ this.#draw(this.#seed);
+ let found = await this.#checkQueryResult();
+ if (found) {
+ try {
+ nonce = this.#readResult(work);
+ } catch (err) {
+ found = false;
+ }
+ }
+ this.#busy = false;
+ if (found && nonce !== work) throw new Error(`Nonce found but does not match work`);
+ return found;
+ }
+};
+
+// src/classes/gpu.ts
+var NanoPowGpu = class _NanoPowGpu {
+ // Initialize WebGPU
+ static #busy = false;
+ static #debug = false;
+ static #device = null;
+ static #gpuBufferReset = new BigUint64Array([0n, 0n]);
+ static #gpuBuffer;
+ static #cpuBuffer;
+ static #uboBuffer;
+ static #uboView;
+ static #bindGroupLayout;
+ static #searchPipeline;
+ static #validatePipeline;
+ // Initialize WebGPU
+ static async init() {
+ if (this.#busy) return;
+ this.#busy = true;
+ try {
+ if (navigator.gpu == null) throw new Error("WebGPU is not supported in this browser.");
+ const adapter = await navigator.gpu.requestAdapter();
+ if (adapter == null) throw new Error("WebGPU adapter refused by browser.");
+ const device = await adapter.requestDevice();
+ if (!(device instanceof GPUDevice)) throw new Error("WebGPU device failed to load.");
+ device.lost.then(this.reset);
+ this.#device = device;
+ this.setup();
+ } catch (err) {
+ throw new Error("WebGPU initialization failed.", { cause: err });
+ } finally {
+ this.#busy = false;
+ }
+ }
+ static setup() {
+ if (this.#device == null) throw new Error(`WebGPU device failed to load.`);
+ this.#gpuBuffer = this.#device.createBuffer({
+ size: 16,
+ usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC
+ });
+ this.#cpuBuffer = this.#device.createBuffer({
+ size: 16,
+ usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ
+ });
+ this.#uboBuffer = this.#device.createBuffer({
+ size: 48,
+ usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST
+ });
+ this.#uboView = new DataView(new ArrayBuffer(48));
+ this.#bindGroupLayout = this.#device.createBindGroupLayout({
+ entries: [
+ {
+ binding: 0,
+ visibility: GPUShaderStage.COMPUTE,
+ buffer: { type: "uniform" }
+ },
+ {
+ binding: 1,
+ visibility: GPUShaderStage.COMPUTE,
+ buffer: { type: "storage" }
+ }
+ ]
+ });
+ const shaderModule = this.#device.createShaderModule({
+ code: compute_default
+ });
+ this.#searchPipeline = this.#device.createComputePipeline({
+ layout: this.#device.createPipelineLayout({
+ bindGroupLayouts: [this.#bindGroupLayout]
+ }),
+ compute: {
+ entryPoint: "search",
+ module: shaderModule
+ }
+ });
+ this.#validatePipeline = this.#device.createComputePipeline({
+ layout: this.#device.createPipelineLayout({
+ bindGroupLayouts: [this.#bindGroupLayout]
+ }),
+ compute: {
+ entryPoint: "validate",
+ module: shaderModule
+ }
+ });
+ console.log(`NanoPow WebGPU initialized. Recommended effort: ${Math.max(1, Math.floor(navigator.hardwareConcurrency / 2))}`);
+ }
+ static reset() {
+ console.warn(`GPU device lost. Reinitializing...`);
+ _NanoPowGpu.#cpuBuffer?.destroy();
+ _NanoPowGpu.#gpuBuffer?.destroy();
+ _NanoPowGpu.#uboBuffer?.destroy();
+ _NanoPowGpu.#busy = false;
+ _NanoPowGpu.init();
+ }
+ static #logAverages(times) {
+ let count = times.length, truncatedCount = 0, truncated = 0, sum = 0, reciprocals = 0, logarithms = 0, min = Number.MAX_SAFE_INTEGER, max = 0, median = 0, rate = 0;
+ times.sort();
+ for (let i = 0; i < count; i++) {
+ sum += times[i];
+ reciprocals += 1 / times[i];
+ logarithms += Math.log(times[i]);
+ min = Math.min(min, times[i]);
+ max = Math.max(max, times[i]);
+ if (i === Math.ceil(count / 2)) {
+ median = times[i];
+ }
+ if (count < 3 || i > 0.1 * count && i < 0.9 * (count - 1)) {
+ truncated += times[i];
+ truncatedCount++;
+ }
+ }
+ const averages = {
+ "Count (dispatches)": count,
+ "Total (ms)": sum,
+ "Rate (d/s)": 1e3 * truncatedCount / (truncated || sum),
+ "Minimum (ms)": min,
+ "Maximum (ms)": max,
+ "Median (ms)": median,
+ "Arithmetic Mean (ms)": sum / count,
+ "Truncated Mean (ms)": truncated / truncatedCount,
+ "Harmonic Mean (ms)": count / reciprocals,
+ "Geometric Mean (ms)": Math.exp(logarithms / count)
+ };
+ console.table(averages);
+ }
+ static async #dispatch(pipeline, seed, hash, threshold, passes) {
+ if (this.#device == null) throw new Error(`WebGPU device failed to load.`);
+ for (let i = 0; i < this.#uboView.byteLength; i++) this.#uboView.setUint8(i, 0);
+ for (let i = 0; i < 64; i += 16) {
+ const u64 = hash.slice(i, i + 16);
+ this.#uboView.setBigUint64(i / 2, BigInt(`0x${u64}`));
+ }
+ this.#uboView.setBigUint64(32, seed, true);
+ this.#uboView.setUint32(40, threshold, true);
+ if (this.#debug) console.log("UBO", this.#uboView);
+ this.#device.queue.writeBuffer(this.#uboBuffer, 0, this.#uboView);
+ this.#device.queue.writeBuffer(this.#gpuBuffer, 0, this.#gpuBufferReset);
+ const bindGroup = this.#device.createBindGroup({
+ layout: this.#bindGroupLayout,
+ entries: [
+ {
+ binding: 0,
+ resource: {
+ buffer: this.#uboBuffer
+ }
+ },
+ {
+ binding: 1,
+ resource: {
+ buffer: this.#gpuBuffer
+ }
+ }
+ ]
+ });
+ const commandEncoder = this.#device.createCommandEncoder();
+ const passEncoder = commandEncoder.beginComputePass();
+ passEncoder.setPipeline(pipeline);
+ passEncoder.setBindGroup(0, bindGroup);
+ passEncoder.dispatchWorkgroups(passes, passes);
+ passEncoder.end();
+ commandEncoder.copyBufferToBuffer(this.#gpuBuffer, 0, this.#cpuBuffer, 0, 12);
+ this.#device.queue.submit([commandEncoder.finish()]);
+ let data = null;
+ try {
+ await this.#cpuBuffer.mapAsync(GPUMapMode.READ);
+ await this.#device.queue.onSubmittedWorkDone();
+ data = new DataView(this.#cpuBuffer.getMappedRange().slice(0));
+ this.#cpuBuffer.unmap();
+ } catch (err) {
+ console.warn(`Error getting data from GPU. ${err}`);
+ return this.#dispatch(pipeline, seed, hash, threshold, passes);
+ }
+ if (this.#debug) console.log("gpuBuffer data", data);
+ if (data == null) throw new Error(`Failed to get data from buffer.`);
+ return data;
+ }
+ /**
+ * Finds a nonce that satisfies the Nano proof-of-work requirements.
+ *
+ * @param {string} hash - Hexadecimal hash of previous block, or public key for new accounts
+ * @param {NanoPowOptions} options - Used to configure search execution
+ */
+ static async search(hash, options) {
+ if (!/^[A-Fa-f0-9]{64}$/.test(hash)) throw new TypeError(`Invalid hash ${hash}`);
+ if (this.#busy) {
+ console.log("NanoPowGpu is busy. Retrying search...");
+ return new Promise((resolve) => {
+ setTimeout(async () => {
+ const result = this.search(hash, options);
+ resolve(result);
+ }, 100);
+ });
+ }
+ this.#busy = true;
+ const threshold = typeof options?.threshold !== "number" || options.threshold < 0 || options.threshold > 4294967295 ? 4294967288 : options.threshold;
+ const effort = typeof options?.effort !== "number" || options.effort < 1 || options.effort > 32 ? 2048 : options.effort * 256;
+ this.#debug = !!options?.debug;
+ let loads = 0;
+ while (this.#device == null && loads < 20) {
+ await new Promise((resolve) => {
+ setTimeout(resolve, 500);
+ });
+ }
+ if (this.#device == null) {
+ this.#busy = false;
+ throw new Error(`WebGPU device failed to load.`);
+ }
+ let times = [];
+ let start = performance.now();
+ let nonce = 0n;
+ do {
+ start = performance.now();
+ const random0 = Math.floor(Math.random() * 4294967295);
+ const random1 = Math.floor(Math.random() * 4294967295);
+ const seed = BigInt(random0) << 32n | BigInt(random1);
+ if (this.#debug) console.log(`seed: ${seed}`);
+ const data = await this.#dispatch(this.#searchPipeline, seed, hash, threshold, effort);
+ nonce = data.getBigUint64(0, true);
+ this.#busy = !data.getUint32(8);
+ times.push(performance.now() - start);
+ } while (this.#busy);
+ if (this.#debug) this.#logAverages(times);
+ return nonce.toString(16).padStart(16, "0");
+ }
+ /**
+ * Validates that a nonce satisfies Nano proof-of-work requirements.
+ *
+ * @param {string} work - Hexadecimal proof-of-work value to validate
+ * @param {string} hash - Hexadecimal hash of previous block, or public key for new accounts
+ * @param {NanoPowOptions} options - Options used to configure search execution
+ */
+ static async validate(work, hash, options) {
+ if (!/^[A-Fa-f0-9]{16}$/.test(work)) throw new TypeError(`Invalid work ${work}`);
+ if (!/^[A-Fa-f0-9]{64}$/.test(hash)) throw new TypeError(`Invalid hash ${hash}`);
+ if (this.#busy) {
+ console.log("NanoPowGpu is busy. Retrying validate...");
+ return new Promise((resolve) => {
+ setTimeout(async () => {
+ const result = this.validate(work, hash, options);
+ resolve(result);
+ }, 100);
+ });
+ }
+ this.#busy = true;
+ this.#debug = !!options?.debug;
+ const threshold = typeof options?.threshold !== "number" || options.threshold < 0 || options.threshold > 4294967295 ? 4294967288 : options.threshold;
+ let loads = 0;
+ while (this.#device == null && loads < 20) {
+ await new Promise((resolve) => {
+ setTimeout(resolve, 500);
+ });
+ }
+ if (this.#device == null) {
+ this.#busy = false;
+ throw new Error(`WebGPU device failed to load.`);
+ }
+ const seed = BigInt(`0x${work}`);
+ if (this.#debug) console.log(`work: ${work}`);
+ const data = await this.#dispatch(this.#validatePipeline, seed, hash, threshold, 1);
+ const nonce = data.getBigUint64(0, true).toString(16).padStart(16, "0");
+ if (this.#debug) console.log(`nonce: ${nonce}`);
+ const found = !!data.getUint32(8);
+ this.#busy = false;
+ if (found && work !== nonce) throw new Error(`Nonce (${nonce}) found but does not match work (${work})`);
+ return found;
+ }
+};
+
+// src/classes/index.ts
+var isGlSupported;
+var isGpuSupported = false;
+try {
+ await NanoPowGpu.init();
+ isGpuSupported = true;
+} catch (err) {
+ console.warn("WebGPU is not supported in this environment.\n", err);
+ isGpuSupported = false;
+}
+try {
+ await NanoPowGl.init();
+ isGlSupported = true;
+} catch (err) {
+ console.warn("WebGL is not supported in this environment.\n", err);
+ isGlSupported = false;
+}
+var NanoPow = isGpuSupported ? NanoPowGpu : isGlSupported ? NanoPowGl : null;
+
+// src/main.ts
+var main_default = NanoPow;
+export {
+ NanoPow,
+ NanoPowGl,
+ NanoPowGpu,
+ main_default as default
+};
projects / nano-pow.git / commitdiff