+++ /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