--- /dev/null
+// src/shaders/compute.wgsl
+var compute_default = "struct UBO{blockhash:array<vec4<u32>,2>,random: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_IV32_1:u32=0x6A09E667u;const ROTATE_1=vec2(1u,1u);const ROTATE_8=vec2(8u,8u);const ROTATE_16=vec2(16u,16u);const ROTATE_24=vec2(24u,24u);const ROTATE_31=vec2(31u,31u);@compute @workgroup_size(8,8)fn search(@builtin(global_invocation_id)global_id:vec3<u32>){if(atomicLoad(&work.found)!=0u){return;}main(global_id);}@compute @workgroup_size(1)fn validate(@builtin(global_invocation_id)global_id:vec3<u32>){if(atomicLoad(&work.found)!=0u){return;}main(global_id);}fn main(id:vec3<u32>){let threshold:u32=ubo.threshold;var m0:vec2<u32>=ubo.random ^ vec2(id.x,id.y);var m1:vec2<u32>=vec2(ubo.blockhash[0u].x,ubo.blockhash[0u].y);var m2:vec2<u32>=vec2(ubo.blockhash[0u].z,ubo.blockhash[0u].w);var m3:vec2<u32>=vec2(ubo.blockhash[1u].x,ubo.blockhash[1u].y);var m4:vec2<u32>=vec2(ubo.blockhash[1u].z,ubo.blockhash[1u].w);var v0:vec2<u32>=vec2(0xF2BDC900u,0x6A09E667u);var v1:vec2<u32>=vec2(0x84CAA73Bu,0xBB67AE85u);var v2:vec2<u32>=vec2(0xFE94F82Bu,0x3C6EF372u);var v3:vec2<u32>=vec2(0x5F1D36F1u,0xA54FF53Au);var v4:vec2<u32>=vec2(0xADE682D1u,0x510E527Fu);var v5:vec2<u32>=vec2(0x2B3E6C1Fu,0x9B05688Cu);var v6:vec2<u32>=vec2(0xFB41BD6Bu,0x1F83D9ABu);var v7:vec2<u32>=vec2(0x137E2179u,0x5BE0CD19u);var v8:vec2<u32>=vec2(0xF3BCC908u,0x6A09E667u);var v9:vec2<u32>=vec2(0x84CAA73Bu,0xBB67AE85u);var v10:vec2<u32>=vec2(0xFE94F82Bu,0x3C6EF372u);var v11:vec2<u32>=vec2(0x5F1D36F1u,0xA54FF53Au);var v12:vec2<u32>=vec2(0xADE682F9u,0x510E527Fu);var v13:vec2<u32>=vec2(0x2B3E6C1Fu,0x9B05688Cu);var v14:vec2<u32>=vec2(0x04BE4294u,0xE07C2654u);var v15:vec2<u32>=vec2(0x137E2179u,0x5BE0CD19u);v0=v0+v4+m0+vec2(0u,u32(v0.x+v4.x<v0.x))+vec2(0u,u32(v0.x+v4.x+m0.x<v0.x+v4.x));v12=v12.yx ^ v0.yx;v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy>>ROTATE_24)|((v4 ^ v8).yx<<ROTATE_8);v0=v0+v4+m1+vec2(0u,u32(v0.x+v4.x<v0.x))+vec2(0u,u32(v0.x+v4.x+m1.x<v0.x+v4.x));v12=((v12 ^ v0).xy>>ROTATE_16)|((v12 ^ v0).yx<<ROTATE_16);v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy<<ROTATE_1)|((v4 ^ v8).yx>>ROTATE_31);v1=v1+v5+m2+vec2(0u,u32(v1.x+v5.x<v1.x))+vec2(0u,u32(v1.x+v5.x+m2.x<v1.x+v5.x));v13=v13.yx ^ v1.yx;v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy>>ROTATE_24)|((v5 ^ v9).yx<<ROTATE_8);v1=v1+v5+m3+vec2(0u,u32(v1.x+v5.x<v1.x))+vec2(0u,u32(v1.x+v5.x+m3.x<v1.x+v5.x));v13=((v13 ^ v1).xy>>ROTATE_16)|((v13 ^ v1).yx<<ROTATE_16);v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy<<ROTATE_1)|((v5 ^ v9).yx>>ROTATE_31);v2=v2+v6+m4+vec2(0u,u32(v2.x+v6.x<v2.x))+vec2(0u,u32(v2.x+v6.x+m4.x<v2.x+v6.x));v14=v14.yx ^ v2.yx;v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy>>ROTATE_24)|((v6 ^ v10).yx<<ROTATE_8);v2=v2+v6+vec2(0u,u32(v2.x+v6.x<v2.x));v14=((v14 ^ v2).xy>>ROTATE_16)|((v14 ^ v2).yx<<ROTATE_16);v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy<<ROTATE_1)|((v6 ^ v10).yx>>ROTATE_31);v3=v3+v7+vec2(0u,u32(v3.x+v7.x<v3.x));v15=v15.yx ^ v3.yx;v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy>>ROTATE_24)|((v7 ^ v11).yx<<ROTATE_8);v3=v3+v7+vec2(0u,u32(v3.x+v7.x<v3.x));v15=((v15 ^ v3).xy>>ROTATE_16)|((v15 ^ v3).yx<<ROTATE_16);v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy<<ROTATE_1)|((v7 ^ v11).yx>>ROTATE_31);v0=v0+v5+vec2(0u,u32(v0.x+v5.x<v0.x));v15=v15.yx ^ v0.yx;v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy>>ROTATE_24)|((v5 ^ v10).yx<<ROTATE_8);v0=v0+v5+vec2(0u,u32(v0.x+v5.x<v0.x));v15=((v15 ^ v0).xy>>ROTATE_16)|((v15 ^ v0).yx<<ROTATE_16);v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy<<ROTATE_1)|((v5 ^ v10).yx>>ROTATE_31);v1=v1+v6+vec2(0u,u32(v1.x+v6.x<v1.x));v12=v12.yx ^ v1.yx;v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy>>ROTATE_24)|((v6 ^ v11).yx<<ROTATE_8);v1=v1+v6+vec2(0u,u32(v1.x+v6.x<v1.x));v12=((v12 ^ v1).xy>>ROTATE_16)|((v12 ^ v1).yx<<ROTATE_16);v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy<<ROTATE_1)|((v6 ^ v11).yx>>ROTATE_31);v2=v2+v7+vec2(0u,u32(v2.x+v7.x<v2.x));v13=v13.yx ^ v2.yx;v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy>>ROTATE_24)|((v7 ^ v8).yx<<ROTATE_8);v2=v2+v7+vec2(0u,u32(v2.x+v7.x<v2.x));v13=((v13 ^ v2).xy>>ROTATE_16)|((v13 ^ v2).yx<<ROTATE_16);v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy<<ROTATE_1)|((v7 ^ v8).yx>>ROTATE_31);v3=v3+v4+vec2(0u,u32(v3.x+v4.x<v3.x));v14=v14.yx ^ v3.yx;v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy>>ROTATE_24)|((v4 ^ v9).yx<<ROTATE_8);v3=v3+v4+vec2(0u,u32(v3.x+v4.x<v3.x));v14=((v14 ^ v3).xy>>ROTATE_16)|((v14 ^ v3).yx<<ROTATE_16);v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy<<ROTATE_1)|((v4 ^ v9).yx>>ROTATE_31);v0=v0+v4+vec2(0u,u32(v0.x+v4.x<v0.x));v12=v12.yx ^ v0.yx;v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy>>ROTATE_24)|((v4 ^ v8).yx<<ROTATE_8);v0=v0+v4+vec2(0u,u32(v0.x+v4.x<v0.x));v12=((v12 ^ v0).xy>>ROTATE_16)|((v12 ^ v0).yx<<ROTATE_16);v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy<<ROTATE_1)|((v4 ^ v8).yx>>ROTATE_31);v1=v1+v5+m4+vec2(0u,u32(v1.x+v5.x<v1.x))+vec2(0u,u32(v1.x+v5.x+m4.x<v1.x+v5.x));v13=v13.yx ^ v1.yx;v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy>>ROTATE_24)|((v5 ^ v9).yx<<ROTATE_8);v1=v1+v5+vec2(0u,u32(v1.x+v5.x<v1.x));v13=((v13 ^ v1).xy>>ROTATE_16)|((v13 ^ v1).yx<<ROTATE_16);v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy<<ROTATE_1)|((v5 ^ v9).yx>>ROTATE_31);v2=v2+v6+vec2(0u,u32(v2.x+v6.x<v2.x));v14=v14.yx ^ v2.yx;v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy>>ROTATE_24)|((v6 ^ v10).yx<<ROTATE_8);v2=v2+v6+vec2(0u,u32(v2.x+v6.x<v2.x));v14=((v14 ^ v2).xy>>ROTATE_16)|((v14 ^ v2).yx<<ROTATE_16);v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy<<ROTATE_1)|((v6 ^ v10).yx>>ROTATE_31);v3=v3+v7+vec2(0u,u32(v3.x+v7.x<v3.x));v15=v15.yx ^ v3.yx;v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy>>ROTATE_24)|((v7 ^ v11).yx<<ROTATE_8);v3=v3+v7+vec2(0u,u32(v3.x+v7.x<v3.x));v15=((v15 ^ v3).xy>>ROTATE_16)|((v15 ^ v3).yx<<ROTATE_16);v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy<<ROTATE_1)|((v7 ^ v11).yx>>ROTATE_31);v0=v0+v5+m1+vec2(0u,u32(v0.x+v5.x<v0.x))+vec2(0u,u32(v0.x+v5.x+m1.x<v0.x+v5.x));v15=v15.yx ^ v0.yx;v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy>>ROTATE_24)|((v5 ^ v10).yx<<ROTATE_8);v0=v0+v5+vec2(0u,u32(v0.x+v5.x<v0.x));v15=((v15 ^ v0).xy>>ROTATE_16)|((v15 ^ v0).yx<<ROTATE_16);v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy<<ROTATE_1)|((v5 ^ v10).yx>>ROTATE_31);v1=v1+v6+m0+vec2(0u,u32(v1.x+v6.x<v1.x))+vec2(0u,u32(v1.x+v6.x+m0.x<v1.x+v6.x));v12=v12.yx ^ v1.yx;v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy>>ROTATE_24)|((v6 ^ v11).yx<<ROTATE_8);v1=v1+v6+m2+vec2(0u,u32(v1.x+v6.x<v1.x))+vec2(0u,u32(v1.x+v6.x+m2.x<v1.x+v6.x));v12=((v12 ^ v1).xy>>ROTATE_16)|((v12 ^ v1).yx<<ROTATE_16);v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy<<ROTATE_1)|((v6 ^ v11).yx>>ROTATE_31);v2=v2+v7+vec2(0u,u32(v2.x+v7.x<v2.x));v13=v13.yx ^ v2.yx;v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy>>ROTATE_24)|((v7 ^ v8).yx<<ROTATE_8);v2=v2+v7+vec2(0u,u32(v2.x+v7.x<v2.x));v13=((v13 ^ v2).xy>>ROTATE_16)|((v13 ^ v2).yx<<ROTATE_16);v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy<<ROTATE_1)|((v7 ^ v8).yx>>ROTATE_31);v3=v3+v4+vec2(0u,u32(v3.x+v4.x<v3.x));v14=v14.yx ^ v3.yx;v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy>>ROTATE_24)|((v4 ^ v9).yx<<ROTATE_8);v3=v3+v4+m3+vec2(0u,u32(v3.x+v4.x<v3.x))+vec2(0u,u32(v3.x+v4.x+m3.x<v3.x+v4.x));v14=((v14 ^ v3).xy>>ROTATE_16)|((v14 ^ v3).yx<<ROTATE_16);v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy<<ROTATE_1)|((v4 ^ v9).yx>>ROTATE_31);v0=v0+v4+vec2(0u,u32(v0.x+v4.x<v0.x));v12=v12.yx ^ v0.yx;v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy>>ROTATE_24)|((v4 ^ v8).yx<<ROTATE_8);v0=v0+v4+vec2(0u,u32(v0.x+v4.x<v0.x));v12=((v12 ^ v0).xy>>ROTATE_16)|((v12 ^ v0).yx<<ROTATE_16);v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy<<ROTATE_1)|((v4 ^ v8).yx>>ROTATE_31);v1=v1+v5+vec2(0u,u32(v1.x+v5.x<v1.x));v13=v13.yx ^ v1.yx;v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy>>ROTATE_24)|((v5 ^ v9).yx<<ROTATE_8);v1=v1+v5+m0+vec2(0u,u32(v1.x+v5.x<v1.x))+vec2(0u,u32(v1.x+v5.x+m0.x<v1.x+v5.x));v13=((v13 ^ v1).xy>>ROTATE_16)|((v13 ^ v1).yx<<ROTATE_16);v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy<<ROTATE_1)|((v5 ^ v9).yx>>ROTATE_31);v2=v2+v6+vec2(0u,u32(v2.x+v6.x<v2.x));v14=v14.yx ^ v2.yx;v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy>>ROTATE_24)|((v6 ^ v10).yx<<ROTATE_8);v2=v2+v6+m2+vec2(0u,u32(v2.x+v6.x<v2.x))+vec2(0u,u32(v2.x+v6.x+m2.x<v2.x+v6.x));v14=((v14 ^ v2).xy>>ROTATE_16)|((v14 ^ v2).yx<<ROTATE_16);v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy<<ROTATE_1)|((v6 ^ v10).yx>>ROTATE_31);v3=v3+v7+vec2(0u,u32(v3.x+v7.x<v3.x));v15=v15.yx ^ v3.yx;v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy>>ROTATE_24)|((v7 ^ v11).yx<<ROTATE_8);v3=v3+v7+vec2(0u,u32(v3.x+v7.x<v3.x));v15=((v15 ^ v3).xy>>ROTATE_16)|((v15 ^ v3).yx<<ROTATE_16);v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy<<ROTATE_1)|((v7 ^ v11).yx>>ROTATE_31);v0=v0+v5+vec2(0u,u32(v0.x+v5.x<v0.x));v15=v15.yx ^ v0.yx;v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy>>ROTATE_24)|((v5 ^ v10).yx<<ROTATE_8);v0=v0+v5+vec2(0u,u32(v0.x+v5.x<v0.x));v15=((v15 ^ v0).xy>>ROTATE_16)|((v15 ^ v0).yx<<ROTATE_16);v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy<<ROTATE_1)|((v5 ^ v10).yx>>ROTATE_31);v1=v1+v6+m3+vec2(0u,u32(v1.x+v6.x<v1.x))+vec2(0u,u32(v1.x+v6.x+m3.x<v1.x+v6.x));v12=v12.yx ^ v1.yx;v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy>>ROTATE_24)|((v6 ^ v11).yx<<ROTATE_8);v1=v1+v6+vec2(0u,u32(v1.x+v6.x<v1.x));v12=((v12 ^ v1).xy>>ROTATE_16)|((v12 ^ v1).yx<<ROTATE_16);v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy<<ROTATE_1)|((v6 ^ v11).yx>>ROTATE_31);v2=v2+v7+vec2(0u,u32(v2.x+v7.x<v2.x));v13=v13.yx ^ v2.yx;v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy>>ROTATE_24)|((v7 ^ v8).yx<<ROTATE_8);v2=v2+v7+m1+vec2(0u,u32(v2.x+v7.x<v2.x))+vec2(0u,u32(v2.x+v7.x+m1.x<v2.x+v7.x));v13=((v13 ^ v2).xy>>ROTATE_16)|((v13 ^ v2).yx<<ROTATE_16);v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy<<ROTATE_1)|((v7 ^ v8).yx>>ROTATE_31);v3=v3+v4+vec2(0u,u32(v3.x+v4.x<v3.x));v14=v14.yx ^ v3.yx;v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy>>ROTATE_24)|((v4 ^ v9).yx<<ROTATE_8);v3=v3+v4+m4+vec2(0u,u32(v3.x+v4.x<v3.x))+vec2(0u,u32(v3.x+v4.x+m4.x<v3.x+v4.x));v14=((v14 ^ v3).xy>>ROTATE_16)|((v14 ^ v3).yx<<ROTATE_16);v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy<<ROTATE_1)|((v4 ^ v9).yx>>ROTATE_31);v0=v0+v4+vec2(0u,u32(v0.x+v4.x<v0.x));v12=v12.yx ^ v0.yx;v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy>>ROTATE_24)|((v4 ^ v8).yx<<ROTATE_8);v0=v0+v4+vec2(0u,u32(v0.x+v4.x<v0.x));v12=((v12 ^ v0).xy>>ROTATE_16)|((v12 ^ v0).yx<<ROTATE_16);v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy<<ROTATE_1)|((v4 ^ v8).yx>>ROTATE_31);v1=v1+v5+m3+vec2(0u,u32(v1.x+v5.x<v1.x))+vec2(0u,u32(v1.x+v5.x+m3.x<v1.x+v5.x));v13=v13.yx ^ v1.yx;v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy>>ROTATE_24)|((v5 ^ v9).yx<<ROTATE_8);v1=v1+v5+m1+vec2(0u,u32(v1.x+v5.x<v1.x))+vec2(0u,u32(v1.x+v5.x+m1.x<v1.x+v5.x));v13=((v13 ^ v1).xy>>ROTATE_16)|((v13 ^ v1).yx<<ROTATE_16);v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy<<ROTATE_1)|((v5 ^ v9).yx>>ROTATE_31);v2=v2+v6+vec2(0u,u32(v2.x+v6.x<v2.x));v14=v14.yx ^ v2.yx;v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy>>ROTATE_24)|((v6 ^ v10).yx<<ROTATE_8);v2=v2+v6+vec2(0u,u32(v2.x+v6.x<v2.x));v14=((v14 ^ v2).xy>>ROTATE_16)|((v14 ^ v2).yx<<ROTATE_16);v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy<<ROTATE_1)|((v6 ^ v10).yx>>ROTATE_31);v3=v3+v7+vec2(0u,u32(v3.x+v7.x<v3.x));v15=v15.yx ^ v3.yx;v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy>>ROTATE_24)|((v7 ^ v11).yx<<ROTATE_8);v3=v3+v7+vec2(0u,u32(v3.x+v7.x<v3.x));v15=((v15 ^ v3).xy>>ROTATE_16)|((v15 ^ v3).yx<<ROTATE_16);v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy<<ROTATE_1)|((v7 ^ v11).yx>>ROTATE_31);v0=v0+v5+m2+vec2(0u,u32(v0.x+v5.x<v0.x))+vec2(0u,u32(v0.x+v5.x+m2.x<v0.x+v5.x));v15=v15.yx ^ v0.yx;v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy>>ROTATE_24)|((v5 ^ v10).yx<<ROTATE_8);v0=v0+v5+vec2(0u,u32(v0.x+v5.x<v0.x));v15=((v15 ^ v0).xy>>ROTATE_16)|((v15 ^ v0).yx<<ROTATE_16);v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy<<ROTATE_1)|((v5 ^ v10).yx>>ROTATE_31);v1=v1+v6+vec2(0u,u32(v1.x+v6.x<v1.x));v12=v12.yx ^ v1.yx;v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy>>ROTATE_24)|((v6 ^ v11).yx<<ROTATE_8);v1=v1+v6+vec2(0u,u32(v1.x+v6.x<v1.x));v12=((v12 ^ v1).xy>>ROTATE_16)|((v12 ^ v1).yx<<ROTATE_16);v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy<<ROTATE_1)|((v6 ^ v11).yx>>ROTATE_31);v2=v2+v7+m4+vec2(0u,u32(v2.x+v7.x<v2.x))+vec2(0u,u32(v2.x+v7.x+m4.x<v2.x+v7.x));v13=v13.yx ^ v2.yx;v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy>>ROTATE_24)|((v7 ^ v8).yx<<ROTATE_8);v2=v2+v7+m0+vec2(0u,u32(v2.x+v7.x<v2.x))+vec2(0u,u32(v2.x+v7.x+m0.x<v2.x+v7.x));v13=((v13 ^ v2).xy>>ROTATE_16)|((v13 ^ v2).yx<<ROTATE_16);v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy<<ROTATE_1)|((v7 ^ v8).yx>>ROTATE_31);v3=v3+v4+vec2(0u,u32(v3.x+v4.x<v3.x));v14=v14.yx ^ v3.yx;v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy>>ROTATE_24)|((v4 ^ v9).yx<<ROTATE_8);v3=v3+v4+vec2(0u,u32(v3.x+v4.x<v3.x));v14=((v14 ^ v3).xy>>ROTATE_16)|((v14 ^ v3).yx<<ROTATE_16);v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy<<ROTATE_1)|((v4 ^ v9).yx>>ROTATE_31);v0=v0+v4+vec2(0u,u32(v0.x+v4.x<v0.x));v12=v12.yx ^ v0.yx;v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy>>ROTATE_24)|((v4 ^ v8).yx<<ROTATE_8);v0=v0+v4+m0+vec2(0u,u32(v0.x+v4.x<v0.x))+vec2(0u,u32(v0.x+v4.x+m0.x<v0.x+v4.x));v12=((v12 ^ v0).xy>>ROTATE_16)|((v12 ^ v0).yx<<ROTATE_16);v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy<<ROTATE_1)|((v4 ^ v8).yx>>ROTATE_31);v1=v1+v5+vec2(0u,u32(v1.x+v5.x<v1.x));v13=v13.yx ^ v1.yx;v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy>>ROTATE_24)|((v5 ^ v9).yx<<ROTATE_8);v1=v1+v5+vec2(0u,u32(v1.x+v5.x<v1.x));v13=((v13 ^ v1).xy>>ROTATE_16)|((v13 ^ v1).yx<<ROTATE_16);v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy<<ROTATE_1)|((v5 ^ v9).yx>>ROTATE_31);v2=v2+v6+m2+vec2(0u,u32(v2.x+v6.x<v2.x))+vec2(0u,u32(v2.x+v6.x+m2.x<v2.x+v6.x));v14=v14.yx ^ v2.yx;v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy>>ROTATE_24)|((v6 ^ v10).yx<<ROTATE_8);v2=v2+v6+m4+vec2(0u,u32(v2.x+v6.x<v2.x))+vec2(0u,u32(v2.x+v6.x+m4.x<v2.x+v6.x));v14=((v14 ^ v2).xy>>ROTATE_16)|((v14 ^ v2).yx<<ROTATE_16);v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy<<ROTATE_1)|((v6 ^ v10).yx>>ROTATE_31);v3=v3+v7+vec2(0u,u32(v3.x+v7.x<v3.x));v15=v15.yx ^ v3.yx;v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy>>ROTATE_24)|((v7 ^ v11).yx<<ROTATE_8);v3=v3+v7+vec2(0u,u32(v3.x+v7.x<v3.x));v15=((v15 ^ v3).xy>>ROTATE_16)|((v15 ^ v3).yx<<ROTATE_16);v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy<<ROTATE_1)|((v7 ^ v11).yx>>ROTATE_31);v0=v0+v5+vec2(0u,u32(v0.x+v5.x<v0.x));v15=v15.yx ^ v0.yx;v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy>>ROTATE_24)|((v5 ^ v10).yx<<ROTATE_8);v0=v0+v5+m1+vec2(0u,u32(v0.x+v5.x<v0.x))+vec2(0u,u32(v0.x+v5.x+m1.x<v0.x+v5.x));v15=((v15 ^ v0).xy>>ROTATE_16)|((v15 ^ v0).yx<<ROTATE_16);v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy<<ROTATE_1)|((v5 ^ v10).yx>>ROTATE_31);v1=v1+v6+vec2(0u,u32(v1.x+v6.x<v1.x));v12=v12.yx ^ v1.yx;v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy>>ROTATE_24)|((v6 ^ v11).yx<<ROTATE_8);v1=v1+v6+vec2(0u,u32(v1.x+v6.x<v1.x));v12=((v12 ^ v1).xy>>ROTATE_16)|((v12 ^ v1).yx<<ROTATE_16);v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy<<ROTATE_1)|((v6 ^ v11).yx>>ROTATE_31);v2=v2+v7+vec2(0u,u32(v2.x+v7.x<v2.x));v13=v13.yx ^ v2.yx;v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy>>ROTATE_24)|((v7 ^ v8).yx<<ROTATE_8);v2=v2+v7+vec2(0u,u32(v2.x+v7.x<v2.x));v13=((v13 ^ v2).xy>>ROTATE_16)|((v13 ^ v2).yx<<ROTATE_16);v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy<<ROTATE_1)|((v7 ^ v8).yx>>ROTATE_31);v3=v3+v4+m3+vec2(0u,u32(v3.x+v4.x<v3.x))+vec2(0u,u32(v3.x+v4.x+m3.x<v3.x+v4.x));v14=v14.yx ^ v3.yx;v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy>>ROTATE_24)|((v4 ^ v9).yx<<ROTATE_8);v3=v3+v4+vec2(0u,u32(v3.x+v4.x<v3.x));v14=((v14 ^ v3).xy>>ROTATE_16)|((v14 ^ v3).yx<<ROTATE_16);v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy<<ROTATE_1)|((v4 ^ v9).yx>>ROTATE_31);v0=v0+v4+m2+vec2(0u,u32(v0.x+v4.x<v0.x))+vec2(0u,u32(v0.x+v4.x+m2.x<v0.x+v4.x));v12=v12.yx ^ v0.yx;v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy>>ROTATE_24)|((v4 ^ v8).yx<<ROTATE_8);v0=v0+v4+vec2(0u,u32(v0.x+v4.x<v0.x));v12=((v12 ^ v0).xy>>ROTATE_16)|((v12 ^ v0).yx<<ROTATE_16);v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy<<ROTATE_1)|((v4 ^ v8).yx>>ROTATE_31);v1=v1+v5+vec2(0u,u32(v1.x+v5.x<v1.x));v13=v13.yx ^ v1.yx;v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy>>ROTATE_24)|((v5 ^ v9).yx<<ROTATE_8);v1=v1+v5+vec2(0u,u32(v1.x+v5.x<v1.x));v13=((v13 ^ v1).xy>>ROTATE_16)|((v13 ^ v1).yx<<ROTATE_16);v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy<<ROTATE_1)|((v5 ^ v9).yx>>ROTATE_31);v2=v2+v6+m0+vec2(0u,u32(v2.x+v6.x<v2.x))+vec2(0u,u32(v2.x+v6.x+m0.x<v2.x+v6.x));v14=v14.yx ^ v2.yx;v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy>>ROTATE_24)|((v6 ^ v10).yx<<ROTATE_8);v2=v2+v6+vec2(0u,u32(v2.x+v6.x<v2.x));v14=((v14 ^ v2).xy>>ROTATE_16)|((v14 ^ v2).yx<<ROTATE_16);v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy<<ROTATE_1)|((v6 ^ v10).yx>>ROTATE_31);v3=v3+v7+vec2(0u,u32(v3.x+v7.x<v3.x));v15=v15.yx ^ v3.yx;v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy>>ROTATE_24)|((v7 ^ v11).yx<<ROTATE_8);v3=v3+v7+m3+vec2(0u,u32(v3.x+v7.x<v3.x))+vec2(0u,u32(v3.x+v7.x+m3.x<v3.x+v7.x));v15=((v15 ^ v3).xy>>ROTATE_16)|((v15 ^ v3).yx<<ROTATE_16);v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy<<ROTATE_1)|((v7 ^ v11).yx>>ROTATE_31);v0=v0+v5+m4+vec2(0u,u32(v0.x+v5.x<v0.x))+vec2(0u,u32(v0.x+v5.x+m4.x<v0.x+v5.x));v15=v15.yx ^ v0.yx;v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy>>ROTATE_24)|((v5 ^ v10).yx<<ROTATE_8);v0=v0+v5+vec2(0u,u32(v0.x+v5.x<v0.x));v15=((v15 ^ v0).xy>>ROTATE_16)|((v15 ^ v0).yx<<ROTATE_16);v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy<<ROTATE_1)|((v5 ^ v10).yx>>ROTATE_31);v1=v1+v6+vec2(0u,u32(v1.x+v6.x<v1.x));v12=v12.yx ^ v1.yx;v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy>>ROTATE_24)|((v6 ^ v11).yx<<ROTATE_8);v1=v1+v6+vec2(0u,u32(v1.x+v6.x<v1.x));v12=((v12 ^ v1).xy>>ROTATE_16)|((v12 ^ v1).yx<<ROTATE_16);v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy<<ROTATE_1)|((v6 ^ v11).yx>>ROTATE_31);v2=v2+v7+vec2(0u,u32(v2.x+v7.x<v2.x));v13=v13.yx ^ v2.yx;v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy>>ROTATE_24)|((v7 ^ v8).yx<<ROTATE_8);v2=v2+v7+vec2(0u,u32(v2.x+v7.x<v2.x));v13=((v13 ^ v2).xy>>ROTATE_16)|((v13 ^ v2).yx<<ROTATE_16);v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy<<ROTATE_1)|((v7 ^ v8).yx>>ROTATE_31);v3=v3+v4+m1+vec2(0u,u32(v3.x+v4.x<v3.x))+vec2(0u,u32(v3.x+v4.x+m1.x<v3.x+v4.x));v14=v14.yx ^ v3.yx;v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy>>ROTATE_24)|((v4 ^ v9).yx<<ROTATE_8);v3=v3+v4+vec2(0u,u32(v3.x+v4.x<v3.x));v14=((v14 ^ v3).xy>>ROTATE_16)|((v14 ^ v3).yx<<ROTATE_16);v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy<<ROTATE_1)|((v4 ^ v9).yx>>ROTATE_31);v0=v0+v4+vec2(0u,u32(v0.x+v4.x<v0.x));v12=v12.yx ^ v0.yx;v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy>>ROTATE_24)|((v4 ^ v8).yx<<ROTATE_8);v0=v0+v4+vec2(0u,u32(v0.x+v4.x<v0.x));v12=((v12 ^ v0).xy>>ROTATE_16)|((v12 ^ v0).yx<<ROTATE_16);v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy<<ROTATE_1)|((v4 ^ v8).yx>>ROTATE_31);v1=v1+v5+m1+vec2(0u,u32(v1.x+v5.x<v1.x))+vec2(0u,u32(v1.x+v5.x+m1.x<v1.x+v5.x));v13=v13.yx ^ v1.yx;v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy>>ROTATE_24)|((v5 ^ v9).yx<<ROTATE_8);v1=v1+v5+vec2(0u,u32(v1.x+v5.x<v1.x));v13=((v13 ^ v1).xy>>ROTATE_16)|((v13 ^ v1).yx<<ROTATE_16);v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy<<ROTATE_1)|((v5 ^ v9).yx>>ROTATE_31);v2=v2+v6+vec2(0u,u32(v2.x+v6.x<v2.x));v14=v14.yx ^ v2.yx;v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy>>ROTATE_24)|((v6 ^ v10).yx<<ROTATE_8);v2=v2+v6+vec2(0u,u32(v2.x+v6.x<v2.x));v14=((v14 ^ v2).xy>>ROTATE_16)|((v14 ^ v2).yx<<ROTATE_16);v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy<<ROTATE_1)|((v6 ^ v10).yx>>ROTATE_31);v3=v3+v7+m4+vec2(0u,u32(v3.x+v7.x<v3.x))+vec2(0u,u32(v3.x+v7.x+m4.x<v3.x+v7.x));v15=v15.yx ^ v3.yx;v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy>>ROTATE_24)|((v7 ^ v11).yx<<ROTATE_8);v3=v3+v7+vec2(0u,u32(v3.x+v7.x<v3.x));v15=((v15 ^ v3).xy>>ROTATE_16)|((v15 ^ v3).yx<<ROTATE_16);v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy<<ROTATE_1)|((v7 ^ v11).yx>>ROTATE_31);v0=v0+v5+m0+vec2(0u,u32(v0.x+v5.x<v0.x))+vec2(0u,u32(v0.x+v5.x+m0.x<v0.x+v5.x));v15=v15.yx ^ v0.yx;v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy>>ROTATE_24)|((v5 ^ v10).yx<<ROTATE_8);v0=v0+v5+vec2(0u,u32(v0.x+v5.x<v0.x));v15=((v15 ^ v0).xy>>ROTATE_16)|((v15 ^ v0).yx<<ROTATE_16);v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy<<ROTATE_1)|((v5 ^ v10).yx>>ROTATE_31);v1=v1+v6+vec2(0u,u32(v1.x+v6.x<v1.x));v12=v12.yx ^ v1.yx;v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy>>ROTATE_24)|((v6 ^ v11).yx<<ROTATE_8);v1=v1+v6+m3+vec2(0u,u32(v1.x+v6.x<v1.x))+vec2(0u,u32(v1.x+v6.x+m3.x<v1.x+v6.x));v12=((v12 ^ v1).xy>>ROTATE_16)|((v12 ^ v1).yx<<ROTATE_16);v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy<<ROTATE_1)|((v6 ^ v11).yx>>ROTATE_31);v2=v2+v7+vec2(0u,u32(v2.x+v7.x<v2.x));v13=v13.yx ^ v2.yx;v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy>>ROTATE_24)|((v7 ^ v8).yx<<ROTATE_8);v2=v2+v7+m2+vec2(0u,u32(v2.x+v7.x<v2.x))+vec2(0u,u32(v2.x+v7.x+m2.x<v2.x+v7.x));v13=((v13 ^ v2).xy>>ROTATE_16)|((v13 ^ v2).yx<<ROTATE_16);v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy<<ROTATE_1)|((v7 ^ v8).yx>>ROTATE_31);v3=v3+v4+vec2(0u,u32(v3.x+v4.x<v3.x));v14=v14.yx ^ v3.yx;v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy>>ROTATE_24)|((v4 ^ v9).yx<<ROTATE_8);v3=v3+v4+vec2(0u,u32(v3.x+v4.x<v3.x));v14=((v14 ^ v3).xy>>ROTATE_16)|((v14 ^ v3).yx<<ROTATE_16);v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy<<ROTATE_1)|((v4 ^ v9).yx>>ROTATE_31);v0=v0+v4+vec2(0u,u32(v0.x+v4.x<v0.x));v12=v12.yx ^ v0.yx;v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy>>ROTATE_24)|((v4 ^ v8).yx<<ROTATE_8);v0=v0+v4+vec2(0u,u32(v0.x+v4.x<v0.x));v12=((v12 ^ v0).xy>>ROTATE_16)|((v12 ^ v0).yx<<ROTATE_16);v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy<<ROTATE_1)|((v4 ^ v8).yx>>ROTATE_31);v1=v1+v5+vec2(0u,u32(v1.x+v5.x<v1.x));v13=v13.yx ^ v1.yx;v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy>>ROTATE_24)|((v5 ^ v9).yx<<ROTATE_8);v1=v1+v5+vec2(0u,u32(v1.x+v5.x<v1.x));v13=((v13 ^ v1).xy>>ROTATE_16)|((v13 ^ v1).yx<<ROTATE_16);v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy<<ROTATE_1)|((v5 ^ v9).yx>>ROTATE_31);v2=v2+v6+vec2(0u,u32(v2.x+v6.x<v2.x));v14=v14.yx ^ v2.yx;v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy>>ROTATE_24)|((v6 ^ v10).yx<<ROTATE_8);v2=v2+v6+m1+vec2(0u,u32(v2.x+v6.x<v2.x))+vec2(0u,u32(v2.x+v6.x+m1.x<v2.x+v6.x));v14=((v14 ^ v2).xy>>ROTATE_16)|((v14 ^ v2).yx<<ROTATE_16);v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy<<ROTATE_1)|((v6 ^ v10).yx>>ROTATE_31);v3=v3+v7+m3+vec2(0u,u32(v3.x+v7.x<v3.x))+vec2(0u,u32(v3.x+v7.x+m3.x<v3.x+v7.x));v15=v15.yx ^ v3.yx;v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy>>ROTATE_24)|((v7 ^ v11).yx<<ROTATE_8);v3=v3+v7+vec2(0u,u32(v3.x+v7.x<v3.x));v15=((v15 ^ v3).xy>>ROTATE_16)|((v15 ^ v3).yx<<ROTATE_16);v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy<<ROTATE_1)|((v7 ^ v11).yx>>ROTATE_31);v0=v0+v5+vec2(0u,u32(v0.x+v5.x<v0.x));v15=v15.yx ^ v0.yx;v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy>>ROTATE_24)|((v5 ^ v10).yx<<ROTATE_8);v0=v0+v5+m0+vec2(0u,u32(v0.x+v5.x<v0.x))+vec2(0u,u32(v0.x+v5.x+m0.x<v0.x+v5.x));v15=((v15 ^ v0).xy>>ROTATE_16)|((v15 ^ v0).yx<<ROTATE_16);v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy<<ROTATE_1)|((v5 ^ v10).yx>>ROTATE_31);v1=v1+v6+vec2(0u,u32(v1.x+v6.x<v1.x));v12=v12.yx ^ v1.yx;v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy>>ROTATE_24)|((v6 ^ v11).yx<<ROTATE_8);v1=v1+v6+m4+vec2(0u,u32(v1.x+v6.x<v1.x))+vec2(0u,u32(v1.x+v6.x+m4.x<v1.x+v6.x));v12=((v12 ^ v1).xy>>ROTATE_16)|((v12 ^ v1).yx<<ROTATE_16);v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy<<ROTATE_1)|((v6 ^ v11).yx>>ROTATE_31);v2=v2+v7+vec2(0u,u32(v2.x+v7.x<v2.x));v13=v13.yx ^ v2.yx;v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy>>ROTATE_24)|((v7 ^ v8).yx<<ROTATE_8);v2=v2+v7+vec2(0u,u32(v2.x+v7.x<v2.x));v13=((v13 ^ v2).xy>>ROTATE_16)|((v13 ^ v2).yx<<ROTATE_16);v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy<<ROTATE_1)|((v7 ^ v8).yx>>ROTATE_31);v3=v3+v4+m2+vec2(0u,u32(v3.x+v4.x<v3.x))+vec2(0u,u32(v3.x+v4.x+m2.x<v3.x+v4.x));v14=v14.yx ^ v3.yx;v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy>>ROTATE_24)|((v4 ^ v9).yx<<ROTATE_8);v3=v3+v4+vec2(0u,u32(v3.x+v4.x<v3.x));v14=((v14 ^ v3).xy>>ROTATE_16)|((v14 ^ v3).yx<<ROTATE_16);v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy<<ROTATE_1)|((v4 ^ v9).yx>>ROTATE_31);v0=v0+v4+vec2(0u,u32(v0.x+v4.x<v0.x));v12=v12.yx ^ v0.yx;v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy>>ROTATE_24)|((v4 ^ v8).yx<<ROTATE_8);v0=v0+v4+vec2(0u,u32(v0.x+v4.x<v0.x));v12=((v12 ^ v0).xy>>ROTATE_16)|((v12 ^ v0).yx<<ROTATE_16);v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy<<ROTATE_1)|((v4 ^ v8).yx>>ROTATE_31);v1=v1+v5+vec2(0u,u32(v1.x+v5.x<v1.x));v13=v13.yx ^ v1.yx;v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy>>ROTATE_24)|((v5 ^ v9).yx<<ROTATE_8);v1=v1+v5+vec2(0u,u32(v1.x+v5.x<v1.x));v13=((v13 ^ v1).xy>>ROTATE_16)|((v13 ^ v1).yx<<ROTATE_16);v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy<<ROTATE_1)|((v5 ^ v9).yx>>ROTATE_31);v2=v2+v6+vec2(0u,u32(v2.x+v6.x<v2.x));v14=v14.yx ^ v2.yx;v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy>>ROTATE_24)|((v6 ^ v10).yx<<ROTATE_8);v2=v2+v6+m3+vec2(0u,u32(v2.x+v6.x<v2.x))+vec2(0u,u32(v2.x+v6.x+m3.x<v2.x+v6.x));v14=((v14 ^ v2).xy>>ROTATE_16)|((v14 ^ v2).yx<<ROTATE_16);v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy<<ROTATE_1)|((v6 ^ v10).yx>>ROTATE_31);v3=v3+v7+m0+vec2(0u,u32(v3.x+v7.x<v3.x))+vec2(0u,u32(v3.x+v7.x+m0.x<v3.x+v7.x));v15=v15.yx ^ v3.yx;v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy>>ROTATE_24)|((v7 ^ v11).yx<<ROTATE_8);v3=v3+v7+vec2(0u,u32(v3.x+v7.x<v3.x));v15=((v15 ^ v3).xy>>ROTATE_16)|((v15 ^ v3).yx<<ROTATE_16);v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy<<ROTATE_1)|((v7 ^ v11).yx>>ROTATE_31);v0=v0+v5+vec2(0u,u32(v0.x+v5.x<v0.x));v15=v15.yx ^ v0.yx;v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy>>ROTATE_24)|((v5 ^ v10).yx<<ROTATE_8);v0=v0+v5+m2+vec2(0u,u32(v0.x+v5.x<v0.x))+vec2(0u,u32(v0.x+v5.x+m2.x<v0.x+v5.x));v15=((v15 ^ v0).xy>>ROTATE_16)|((v15 ^ v0).yx<<ROTATE_16);v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy<<ROTATE_1)|((v5 ^ v10).yx>>ROTATE_31);v1=v1+v6+vec2(0u,u32(v1.x+v6.x<v1.x));v12=v12.yx ^ v1.yx;v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy>>ROTATE_24)|((v6 ^ v11).yx<<ROTATE_8);v1=v1+v6+vec2(0u,u32(v1.x+v6.x<v1.x));v12=((v12 ^ v1).xy>>ROTATE_16)|((v12 ^ v1).yx<<ROTATE_16);v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy<<ROTATE_1)|((v6 ^ v11).yx>>ROTATE_31);v2=v2+v7+m1+vec2(0u,u32(v2.x+v7.x<v2.x))+vec2(0u,u32(v2.x+v7.x+m1.x<v2.x+v7.x));v13=v13.yx ^ v2.yx;v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy>>ROTATE_24)|((v7 ^ v8).yx<<ROTATE_8);v2=v2+v7+m4+vec2(0u,u32(v2.x+v7.x<v2.x))+vec2(0u,u32(v2.x+v7.x+m4.x<v2.x+v7.x));v13=((v13 ^ v2).xy>>ROTATE_16)|((v13 ^ v2).yx<<ROTATE_16);v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy<<ROTATE_1)|((v7 ^ v8).yx>>ROTATE_31);v3=v3+v4+vec2(0u,u32(v3.x+v4.x<v3.x));v14=v14.yx ^ v3.yx;v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy>>ROTATE_24)|((v4 ^ v9).yx<<ROTATE_8);v3=v3+v4+vec2(0u,u32(v3.x+v4.x<v3.x));v14=((v14 ^ v3).xy>>ROTATE_16)|((v14 ^ v3).yx<<ROTATE_16);v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy<<ROTATE_1)|((v4 ^ v9).yx>>ROTATE_31);v0=v0+v4+vec2(0u,u32(v0.x+v4.x<v0.x));v12=v12.yx ^ v0.yx;v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy>>ROTATE_24)|((v4 ^ v8).yx<<ROTATE_8);v0=v0+v4+m2+vec2(0u,u32(v0.x+v4.x<v0.x))+vec2(0u,u32(v0.x+v4.x+m2.x<v0.x+v4.x));v12=((v12 ^ v0).xy>>ROTATE_16)|((v12 ^ v0).yx<<ROTATE_16);v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy<<ROTATE_1)|((v4 ^ v8).yx>>ROTATE_31);v1=v1+v5+vec2(0u,u32(v1.x+v5.x<v1.x));v13=v13.yx ^ v1.yx;v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy>>ROTATE_24)|((v5 ^ v9).yx<<ROTATE_8);v1=v1+v5+m4+vec2(0u,u32(v1.x+v5.x<v1.x))+vec2(0u,u32(v1.x+v5.x+m4.x<v1.x+v5.x));v13=((v13 ^ v1).xy>>ROTATE_16)|((v13 ^ v1).yx<<ROTATE_16);v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy<<ROTATE_1)|((v5 ^ v9).yx>>ROTATE_31);v2=v2+v6+vec2(0u,u32(v2.x+v6.x<v2.x));v14=v14.yx ^ v2.yx;v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy>>ROTATE_24)|((v6 ^ v10).yx<<ROTATE_8);v2=v2+v6+vec2(0u,u32(v2.x+v6.x<v2.x));v14=((v14 ^ v2).xy>>ROTATE_16)|((v14 ^ v2).yx<<ROTATE_16);v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy<<ROTATE_1)|((v6 ^ v10).yx>>ROTATE_31);v3=v3+v7+m1+vec2(0u,u32(v3.x+v7.x<v3.x))+vec2(0u,u32(v3.x+v7.x+m1.x<v3.x+v7.x));v15=v15.yx ^ v3.yx;v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy>>ROTATE_24)|((v7 ^ v11).yx<<ROTATE_8);v3=v3+v7+vec2(0u,u32(v3.x+v7.x<v3.x));v15=((v15 ^ v3).xy>>ROTATE_16)|((v15 ^ v3).yx<<ROTATE_16);v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy<<ROTATE_1)|((v7 ^ v11).yx>>ROTATE_31);v0=v0+v5+vec2(0u,u32(v0.x+v5.x<v0.x));v15=v15.yx ^ v0.yx;v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy>>ROTATE_24)|((v5 ^ v10).yx<<ROTATE_8);v0=v0+v5+vec2(0u,u32(v0.x+v5.x<v0.x));v15=((v15 ^ v0).xy>>ROTATE_16)|((v15 ^ v0).yx<<ROTATE_16);v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy<<ROTATE_1)|((v5 ^ v10).yx>>ROTATE_31);v1=v1+v6+vec2(0u,u32(v1.x+v6.x<v1.x));v12=v12.yx ^ v1.yx;v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy>>ROTATE_24)|((v6 ^ v11).yx<<ROTATE_8);v1=v1+v6+vec2(0u,u32(v1.x+v6.x<v1.x));v12=((v12 ^ v1).xy>>ROTATE_16)|((v12 ^ v1).yx<<ROTATE_16);v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy<<ROTATE_1)|((v6 ^ v11).yx>>ROTATE_31);v2=v2+v7+m3+vec2(0u,u32(v2.x+v7.x<v2.x))+vec2(0u,u32(v2.x+v7.x+m3.x<v2.x+v7.x));v13=v13.yx ^ v2.yx;v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy>>ROTATE_24)|((v7 ^ v8).yx<<ROTATE_8);v2=v2+v7+vec2(0u,u32(v2.x+v7.x<v2.x));v13=((v13 ^ v2).xy>>ROTATE_16)|((v13 ^ v2).yx<<ROTATE_16);v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy<<ROTATE_1)|((v7 ^ v8).yx>>ROTATE_31);v3=v3+v4+vec2(0u,u32(v3.x+v4.x<v3.x));v14=v14.yx ^ v3.yx;v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy>>ROTATE_24)|((v4 ^ v9).yx<<ROTATE_8);v3=v3+v4+m0+vec2(0u,u32(v3.x+v4.x<v3.x))+vec2(0u,u32(v3.x+v4.x+m0.x<v3.x+v4.x));v14=((v14 ^ v3).xy>>ROTATE_16)|((v14 ^ v3).yx<<ROTATE_16);v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy<<ROTATE_1)|((v4 ^ v9).yx>>ROTATE_31);v0=v0+v4+m0+vec2(0u,u32(v0.x+v4.x<v0.x))+vec2(0u,u32(v0.x+v4.x+m0.x<v0.x+v4.x));v12=v12.yx ^ v0.yx;v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy>>ROTATE_24)|((v4 ^ v8).yx<<ROTATE_8);v0=v0+v4+m1+vec2(0u,u32(v0.x+v4.x<v0.x))+vec2(0u,u32(v0.x+v4.x+m1.x<v0.x+v4.x));v12=((v12 ^ v0).xy>>ROTATE_16)|((v12 ^ v0).yx<<ROTATE_16);v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy<<ROTATE_1)|((v4 ^ v8).yx>>ROTATE_31);v1=v1+v5+m2+vec2(0u,u32(v1.x+v5.x<v1.x))+vec2(0u,u32(v1.x+v5.x+m2.x<v1.x+v5.x));v13=v13.yx ^ v1.yx;v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy>>ROTATE_24)|((v5 ^ v9).yx<<ROTATE_8);v1=v1+v5+m3+vec2(0u,u32(v1.x+v5.x<v1.x))+vec2(0u,u32(v1.x+v5.x+m3.x<v1.x+v5.x));v13=((v13 ^ v1).xy>>ROTATE_16)|((v13 ^ v1).yx<<ROTATE_16);v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy<<ROTATE_1)|((v5 ^ v9).yx>>ROTATE_31);v2=v2+v6+m4+vec2(0u,u32(v2.x+v6.x<v2.x))+vec2(0u,u32(v2.x+v6.x+m4.x<v2.x+v6.x));v14=v14.yx ^ v2.yx;v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy>>ROTATE_24)|((v6 ^ v10).yx<<ROTATE_8);v2=v2+v6+vec2(0u,u32(v2.x+v6.x<v2.x));v14=((v14 ^ v2).xy>>ROTATE_16)|((v14 ^ v2).yx<<ROTATE_16);v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy<<ROTATE_1)|((v6 ^ v10).yx>>ROTATE_31);v3=v3+v7+vec2(0u,u32(v3.x+v7.x<v3.x));v15=v15.yx ^ v3.yx;v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy>>ROTATE_24)|((v7 ^ v11).yx<<ROTATE_8);v3=v3+v7+vec2(0u,u32(v3.x+v7.x<v3.x));v15=((v15 ^ v3).xy>>ROTATE_16)|((v15 ^ v3).yx<<ROTATE_16);v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy<<ROTATE_1)|((v7 ^ v11).yx>>ROTATE_31);v0=v0+v5+vec2(0u,u32(v0.x+v5.x<v0.x));v15=v15.yx ^ v0.yx;v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy>>ROTATE_24)|((v5 ^ v10).yx<<ROTATE_8);v0=v0+v5+vec2(0u,u32(v0.x+v5.x<v0.x));v15=((v15 ^ v0).xy>>ROTATE_16)|((v15 ^ v0).yx<<ROTATE_16);v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy<<ROTATE_1)|((v5 ^ v10).yx>>ROTATE_31);v1=v1+v6+vec2(0u,u32(v1.x+v6.x<v1.x));v12=v12.yx ^ v1.yx;v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy>>ROTATE_24)|((v6 ^ v11).yx<<ROTATE_8);v1=v1+v6+vec2(0u,u32(v1.x+v6.x<v1.x));v12=((v12 ^ v1).xy>>ROTATE_16)|((v12 ^ v1).yx<<ROTATE_16);v11=v11+v12+vec2(0u,u32(v11.x+v12.x<v11.x));v6=((v6 ^ v11).xy<<ROTATE_1)|((v6 ^ v11).yx>>ROTATE_31);v2=v2+v7+vec2(0u,u32(v2.x+v7.x<v2.x));v13=v13.yx ^ v2.yx;v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy>>ROTATE_24)|((v7 ^ v8).yx<<ROTATE_8);v2=v2+v7+vec2(0u,u32(v2.x+v7.x<v2.x));v13=((v13 ^ v2).xy>>ROTATE_16)|((v13 ^ v2).yx<<ROTATE_16);v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy<<ROTATE_1)|((v7 ^ v8).yx>>ROTATE_31);v3=v3+v4+vec2(0u,u32(v3.x+v4.x<v3.x));v14=v14.yx ^ v3.yx;v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy>>ROTATE_24)|((v4 ^ v9).yx<<ROTATE_8);v3=v3+v4+vec2(0u,u32(v3.x+v4.x<v3.x));v14=((v14 ^ v3).xy>>ROTATE_16)|((v14 ^ v3).yx<<ROTATE_16);v9=v9+v14+vec2(0u,u32(v9.x+v14.x<v9.x));v4=((v4 ^ v9).xy<<ROTATE_1)|((v4 ^ v9).yx>>ROTATE_31);v0=v0+v4+vec2(0u,u32(v0.x+v4.x<v0.x));v12=v12.yx ^ v0.yx;v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v4=((v4 ^ v8).xy>>ROTATE_24)|((v4 ^ v8).yx<<ROTATE_8);v0=v0+v4+vec2(0u,u32(v0.x+v4.x<v0.x));v12=((v12 ^ v0).xy>>ROTATE_16)|((v12 ^ v0).yx<<ROTATE_16);v8=v8+v12+vec2(0u,u32(v8.x+v12.x<v8.x));v1=v1+v5+m4+vec2(0u,u32(v1.x+v5.x<v1.x))+vec2(0u,u32(v1.x+v5.x+m4.x<v1.x+v5.x));v13=v13.yx ^ v1.yx;v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy>>ROTATE_24)|((v5 ^ v9).yx<<ROTATE_8);v1=v1+v5+vec2(0u,u32(v1.x+v5.x<v1.x));v13=((v13 ^ v1).xy>>ROTATE_16)|((v13 ^ v1).yx<<ROTATE_16);v9=v9+v13+vec2(0u,u32(v9.x+v13.x<v9.x));v5=((v5 ^ v9).xy<<ROTATE_1)|((v5 ^ v9).yx>>ROTATE_31);v2=v2+v6+vec2(0u,u32(v2.x+v6.x<v2.x));v14=v14.yx ^ v2.yx;v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v6=((v6 ^ v10).xy>>ROTATE_24)|((v6 ^ v10).yx<<ROTATE_8);v2=v2+v6+vec2(0u,u32(v2.x+v6.x<v2.x));v14=((v14 ^ v2).xy>>ROTATE_16)|((v14 ^ v2).yx<<ROTATE_16);v10=v10+v14+vec2(0u,u32(v10.x+v14.x<v10.x));v3=v3+v7+vec2(0u,u32(v3.x+v7.x<v3.x));v15=v15.yx ^ v3.yx;v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy>>ROTATE_24)|((v7 ^ v11).yx<<ROTATE_8);v3=v3+v7+vec2(0u,u32(v3.x+v7.x<v3.x));v15=((v15 ^ v3).xy>>ROTATE_16)|((v15 ^ v3).yx<<ROTATE_16);v11=v11+v15+vec2(0u,u32(v11.x+v15.x<v11.x));v7=((v7 ^ v11).xy<<ROTATE_1)|((v7 ^ v11).yx>>ROTATE_31);v0=v0+v5+m1+vec2(0u,u32(v0.x+v5.x<v0.x))+vec2(0u,u32(v0.x+v5.x+m1.x<v0.x+v5.x));v15=v15.yx ^ v0.yx;v10=v10+v15+vec2(0u,u32(v10.x+v15.x<v10.x));v5=((v5 ^ v10).xy>>ROTATE_24)|((v5 ^ v10).yx<<ROTATE_8);v0=v0+v5+vec2(0u,u32(v0.x+v5.x<v0.x));v2=v2+v7+vec2(0u,u32(v2.x+v7.x<v2.x));v13=v13.yx ^ v2.yx;v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));v7=((v7 ^ v8).xy>>ROTATE_24)|((v7 ^ v8).yx<<ROTATE_8);v2=v2+v7+vec2(0u,u32(v2.x+v7.x<v2.x));v13=((v13 ^ v2).xy>>ROTATE_16)|((v13 ^ v2).yx<<ROTATE_16);v8=v8+v13+vec2(0u,u32(v8.x+v13.x<v8.x));if((BLAKE2B_IV32_1 ^ v0.y ^ v8.y)>threshold&&atomicLoad(&work.found)==0u){atomicStore(&work.found,1u);work.nonce=m0;}return;}";
+
+// src/shaders/gl-fragment.ts
+var NanoPowGlFragmentShader = `#version 300 es
+#pragma vscode_glsllint_stage: frag
+precision highp float;
+precision highp int;
+
+in vec2 uv_pos;
+out vec4 fragColor;
+
+// blockhash - array of precalculated block hash components
+// threshold - 0xfffffff8 for send/change blocks, 0xfffffe00 for all else
+// workload - Defines canvas size
+layout(std140) uniform UBO {
+ uint blockhash[8];
+ uint threshold;
+ float workload;
+};
+
+// Random work values
+// First 2 bytes will be overwritten by texture pixel position
+// Second 2 bytes will be modified if the canvas size is greater than 256x256
+// Last 4 bytes remain as generated externally
+layout(std140) uniform WORK {
+ uvec4 work[2];
+};
+
+// Defined separately from uint v[32] below as the original value is required
+// to calculate the second uint32 of the digest for threshold comparison
+const uint BLAKE2B_IV32_1 = 0x6A09E667u;
+
+// 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[24] ^= uint(INLEN);
+// It's always the "last" compression at this INLEN
+// v[28] = ~v[28];
+// v[29] = ~v[29];
+uint v[32] = uint[32](
+ 0xF2BDC900u, 0x6A09E667u, 0x84CAA73Bu, 0xBB67AE85u,
+ 0xFE94F82Bu, 0x3C6EF372u, 0x5F1D36F1u, 0xA54FF53Au,
+ 0xADE682D1u, 0x510E527Fu, 0x2B3E6C1Fu, 0x9B05688Cu,
+ 0xFB41BD6Bu, 0x1F83D9ABu, 0x137E2179u, 0x5BE0CD19u,
+ 0xF3BCC908u, 0x6A09E667u, 0x84CAA73Bu, 0xBB67AE85u,
+ 0xFE94F82Bu, 0x3C6EF372u, 0x5F1D36F1u, 0xA54FF53Au,
+ 0xADE682F9u, 0x510E527Fu, 0x2B3E6C1Fu, 0x9B05688Cu,
+ 0x04BE4294u, 0xE07C2654u, 0x137E2179u, 0x5BE0CD19u
+);
+// Input data buffer
+uint m[32];
+
+// These are offsets into the input data buffer for each mixing step.
+// They are multiplied by 2 from the original SIGMA values in
+// the C reference implementation, which refered to uint64s.
+const uint SIGMA82[192] = uint[192](
+ 0u,2u,4u,6u,8u,10u,12u,14u,16u,18u,20u,22u,24u,26u,28u,30u,
+ 28u,20u,8u,16u,18u,30u,26u,12u,2u,24u,0u,4u,22u,14u,10u,6u,
+ 22u,16u,24u,0u,10u,4u,30u,26u,20u,28u,6u,12u,14u,2u,18u,8u,
+ 14u,18u,6u,2u,26u,24u,22u,28u,4u,12u,10u,20u,8u,0u,30u,16u,
+ 18u,0u,10u,14u,4u,8u,20u,30u,28u,2u,22u,24u,12u,16u,6u,26u,
+ 4u,24u,12u,20u,0u,22u,16u,6u,8u,26u,14u,10u,30u,28u,2u,18u,
+ 24u,10u,2u,30u,28u,26u,8u,20u,0u,14u,12u,6u,18u,4u,16u,22u,
+ 26u,22u,14u,28u,24u,2u,6u,18u,10u,0u,30u,8u,16u,12u,4u,20u,
+ 12u,30u,28u,18u,22u,6u,0u,16u,24u,4u,26u,14u,2u,8u,20u,10u,
+ 20u,4u,16u,8u,14u,12u,2u,10u,30u,22u,18u,28u,6u,24u,26u,0u,
+ 0u,2u,4u,6u,8u,10u,12u,14u,16u,18u,20u,22u,24u,26u,28u,30u,
+ 28u,20u,8u,16u,18u,30u,26u,12u,2u,24u,0u,4u,22u,14u,10u,6u
+);
+
+// G mixing function
+void G (uint ix, uint iy, uint a, uint b, uint c, uint d) {
+ bool carry;
+ uint o0;
+ uint o1;
+ uint xor0;
+ uint xor1;
+
+ // a = a + b;
+ o0 = v[a] + v[b];
+ o1 = v[a+1u] + v[b+1u];
+ carry = o0 < v[a];
+ o1 = o1 + uint(mix(0.0, 1.0, float(carry)));
+ v[a] = o0;
+ v[a+1u] = o1;
+
+ // a = a + m[sigma[r][2*i+0]];
+ o0 = v[a] + m[ix];
+ o1 = v[a+1u] + m[ix+1u];
+ carry = o0 < v[a];
+ o1 = o1 + uint(mix(0.0, 1.0, float(carry)));
+ v[a] = o0;
+ v[a+1u] = o1;
+
+ // d = rotr64(d ^ a, 32);
+ xor0 = v[d] ^ v[a];
+ xor1 = v[d+1u] ^ v[a+1u];
+ v[d] = xor1;
+ v[d+1u] = xor0;
+
+ // c = c + d;
+ o0 = v[c] + v[d];
+ o1 = v[c+1u] + v[d+1u];
+ carry = o0 < v[c];
+ o1 = o1 + uint(mix(0.0, 1.0, float(carry)));
+ v[c] = o0;
+ v[c+1u] = o1;
+
+ // b = rotr64(b ^ c, 24);
+ xor0 = v[b] ^ v[c];
+ xor1 = v[b+1u] ^ v[c+1u];
+ v[b] = (xor0 >> 24u) ^ (xor1 << 8u);
+ v[b+1u] = (xor1 >> 24u) ^ (xor0 << 8u);
+
+ // a = a + b;
+ o0 = v[a] + v[b];
+ o1 = v[a+1u] + v[b+1u];
+ carry = o0 < v[a];
+ o1 = o1 + uint(mix(0.0, 1.0, float(carry)));
+ v[a] = o0;
+ v[a+1u] = o1;
+
+ // a = a + m[sigma[r][2*i+1]];
+ o0 = v[a] + m[iy];
+ o1 = v[a+1u] + m[iy+1u];
+ carry = o0 < v[a];
+ o1 = o1 + uint(mix(0.0, 1.0, float(carry)));
+ v[a] = o0;
+ v[a+1u] = o1;
+
+ // d = rotr64(d ^ a, 16)
+ xor0 = v[d] ^ v[a];
+ xor1 = v[d+1u] ^ v[a+1u];
+ v[d] = (xor0 >> 16u) ^ (xor1 << 16u);
+ v[d+1u] = (xor1 >> 16u) ^ (xor0 << 16u);
+
+ // c = c + d;
+ o0 = v[c] + v[d];
+ o1 = v[c+1u] + v[d+1u];
+ carry = o0 < v[c];
+ o1 = o1 + uint(mix(0.0, 1.0, float(carry)));
+ v[c] = o0;
+ v[c+1u] = o1;
+
+ // b = rotr64(b ^ c, 63)
+ xor0 = v[b] ^ v[c];
+ xor1 = v[b+1u] ^ v[c+1u];
+ v[b] = (xor1 >> 31u) ^ (xor0 << 1u);
+ v[b+1u] = (xor0 >> 31u) ^ (xor1 << 1u);
+}
+
+void main() {
+ int i;
+ uvec4 u_work0 = work[0u];
+ uvec4 u_work1 = work[1u];
+ uint uv_x = uint(uv_pos.x * workload);
+ uint uv_y = uint(uv_pos.y * workload);
+ uint x_pos = uv_x % 256u;
+ uint y_pos = uv_y % 256u;
+ uint x_index = (uv_x - x_pos) / 256u;
+ uint y_index = (uv_y - y_pos) / 256u;
+
+ // First 2 work bytes are the x,y pos within the 256x256 area, the next
+ // two bytes are modified from the random generated value, XOR'd with
+ // the x,y area index of where this pixel is located
+ m[0u] = (x_pos ^ (y_pos << 8u) ^ ((u_work0.b ^ x_index) << 16u) ^ ((u_work0.a ^ y_index) << 24u));
+
+ // Remaining bytes are un-modified from the random generated value
+ m[1u] = (u_work1.r ^ (u_work1.g << 8u) ^ (u_work1.b << 16u) ^ (u_work1.a << 24u));
+
+ // Block hash
+ for (uint i = 0u; i < 8u; i = i + 1u) {
+ m[i+2u] = blockhash[i];
+ }
+
+ // twelve rounds of mixing
+ for(uint i = 0u; i < 12u; i = i + 1u) {
+ G(SIGMA82[i * 16u + 0u], SIGMA82[i * 16u + 1u], 0u, 8u, 16u, 24u);
+ G(SIGMA82[i * 16u + 2u], SIGMA82[i * 16u + 3u], 2u, 10u, 18u, 26u);
+ G(SIGMA82[i * 16u + 4u], SIGMA82[i * 16u + 5u], 4u, 12u, 20u, 28u);
+ G(SIGMA82[i * 16u + 6u], SIGMA82[i * 16u + 7u], 6u, 14u, 22u, 30u);
+ G(SIGMA82[i * 16u + 8u], SIGMA82[i * 16u + 9u], 0u, 10u, 20u, 30u);
+ G(SIGMA82[i * 16u + 10u], SIGMA82[i * 16u + 11u], 2u, 12u, 22u, 24u);
+ G(SIGMA82[i * 16u + 12u], SIGMA82[i * 16u + 13u], 4u, 14u, 16u, 26u);
+ G(SIGMA82[i * 16u + 14u], SIGMA82[i * 16u + 15u], 6u, 8u, 18u, 28u);
+ }
+
+ // Pixel data is multipled by threshold test result (0 or 1)
+ // First 4 bytes insignificant, only calculate digest of second 4 bytes
+ if ((BLAKE2B_IV32_1 ^ v[1u] ^ v[17u]) > threshold) {
+ fragColor = vec4(
+ float(x_index + 1u)/255.0, // +1 to distinguish from 0 (unsuccessful) pixels
+ float(y_index + 1u)/255.0, // Same as previous
+ float(x_pos)/255.0, // Return the 2 custom bytes used in work value
+ float(y_pos)/255.0 // Second custom byte
+ );
+ } else {
+ discard;
+ }
+}
+`;
+
+// src/shaders/gl-vertex.ts
+var NanoPowGlVertexShader = `#version 300 es
+#pragma vscode_glsllint_stage: vert
+precision highp float;
+layout (location=0) in vec4 position;
+layout (location=1) in vec2 uv;
+
+out vec2 uv_pos;
+
+void main() {
+ uv_pos = uv;
+ gl_Position = position;
+}
+`;
+
+// src/classes/gl.ts
+var NanoPowGl = class _NanoPowGl {
+ static #busy = false;
+ /** Used to set canvas size. Must be a multiple of 256. */
+ static #WORKLOAD = 256 * Math.max(1, Math.floor(navigator.hardwareConcurrency));
+ static #hexify(arr) {
+ let out = "";
+ for (let i = arr.length - 1; i >= 0; i--) {
+ out += arr[i].toString(16).padStart(2, "0");
+ }
+ return out;
+ }
+ static #gl;
+ static #program;
+ static #vertexShader;
+ static #fragmentShader;
+ static #positionBuffer;
+ static #uvBuffer;
+ static #uboBuffer;
+ static #workBuffer;
+ static #query;
+ static #pixels;
+ /**Vertex Positions, 2 triangles */
+ static #positions = new Float32Array([
+ -1,
+ -1,
+ 0,
+ -1,
+ 1,
+ 0,
+ 1,
+ 1,
+ 0,
+ 1,
+ -1,
+ 0,
+ 1,
+ 1,
+ 0,
+ -1,
+ -1,
+ 0
+ ]);
+ /** Texture Positions */
+ static #uvPosArray = new Float32Array([
+ 1,
+ 1,
+ 1,
+ 0,
+ 0,
+ 0,
+ 0,
+ 1,
+ 0,
+ 0,
+ 1,
+ 1
+ ]);
+ /** Compile */
+ static async init() {
+ if (this.#busy) return;
+ this.#busy = true;
+ try {
+ this.#gl = new OffscreenCanvas(this.#WORKLOAD, this.#WORKLOAD).getContext("webgl2");
+ if (this.#gl == null) throw new Error("WebGL 2 is required");
+ this.#gl.clearColor(0, 0, 0, 1);
+ this.#program = this.#gl.createProgram();
+ if (this.#program == 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.#fragmentShader = this.#gl.createShader(this.#gl.FRAGMENT_SHADER);
+ if (this.#fragmentShader == null) throw new Error("Failed to create fragment shader");
+ this.#gl.shaderSource(this.#fragmentShader, NanoPowGlFragmentShader);
+ this.#gl.compileShader(this.#fragmentShader);
+ if (!this.#gl.getShaderParameter(this.#fragmentShader, this.#gl.COMPILE_STATUS))
+ throw new Error(this.#gl.getShaderInfoLog(this.#fragmentShader) ?? `Failed to compile fragment shader`);
+ this.#gl.attachShader(this.#program, this.#vertexShader);
+ this.#gl.attachShader(this.#program, this.#fragmentShader);
+ this.#gl.linkProgram(this.#program);
+ if (!this.#gl.getProgramParameter(this.#program, this.#gl.LINK_STATUS))
+ throw new Error(this.#gl.getProgramInfoLog(this.#program) ?? `Failed to link program`);
+ this.#gl.useProgram(this.#program);
+ 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, 3, this.#gl.FLOAT, false, 0, 0);
+ this.#gl.enableVertexAttribArray(0);
+ this.#uvBuffer = this.#gl.createBuffer();
+ this.#gl.bindBuffer(this.#gl.ARRAY_BUFFER, this.#uvBuffer);
+ this.#gl.bufferData(this.#gl.ARRAY_BUFFER, this.#uvPosArray, this.#gl.STATIC_DRAW);
+ this.#gl.vertexAttribPointer(1, 2, this.#gl.FLOAT, false, 0, 0);
+ this.#gl.enableVertexAttribArray(1);
+ 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.#program, this.#gl.getUniformBlockIndex(this.#program, "UBO"), 0);
+ this.#workBuffer = this.#gl.createBuffer();
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, this.#workBuffer);
+ this.#gl.bufferData(this.#gl.UNIFORM_BUFFER, 32, this.#gl.STREAM_DRAW);
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, null);
+ this.#gl.bindBufferBase(this.#gl.UNIFORM_BUFFER, 1, this.#workBuffer);
+ this.#gl.uniformBlockBinding(this.#program, this.#gl.getUniformBlockIndex(this.#program, "WORK"), 1);
+ this.#pixels = new Uint8Array(this.#gl.drawingBufferWidth * this.#gl.drawingBufferHeight * 4);
+ this.#query = this.#gl.createQuery();
+ } catch (err) {
+ throw new Error(`WebGL initialization failed. ${err}`);
+ } finally {
+ this.#busy = false;
+ }
+ }
+ static reset() {
+ _NanoPowGl.#query = null;
+ _NanoPowGl.#workBuffer = null;
+ _NanoPowGl.#uboBuffer = null;
+ _NanoPowGl.#uvBuffer = null;
+ _NanoPowGl.#positionBuffer = null;
+ _NanoPowGl.#fragmentShader = null;
+ _NanoPowGl.#vertexShader = null;
+ _NanoPowGl.#program = 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 > 2 && i > count * 0.1 && i < count * 0.9) truncated += times[i];
+ }
+ const averages = {
+ "Count (frames)": count,
+ "Total (ms)": sum,
+ "Rate (f/s)": 1e3 * count / (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.table(averages);
+ }
+ static #draw(work) {
+ if (this.#gl == null || this.#query == null) throw new Error("WebGL 2 is required to draw and query pixels");
+ if (this.#workBuffer == null) throw new Error("Work buffer is required to draw");
+ this.#gl.clear(this.#gl.COLOR_BUFFER_BIT);
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, this.#workBuffer);
+ this.#gl.bufferSubData(this.#gl.UNIFORM_BUFFER, 0, Uint32Array.from(work));
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, null);
+ this.#gl.beginQuery(this.#gl.ANY_SAMPLES_PASSED_CONSERVATIVE, this.#query);
+ this.#gl.drawArrays(this.#gl.TRIANGLES, 0, 6);
+ this.#gl.endQuery(this.#gl.ANY_SAMPLES_PASSED_CONSERVATIVE);
+ }
+ static async #checkQueryResult() {
+ return new Promise((resolve, reject) => {
+ try {
+ if (this.#gl == null || this.#query == null) throw new Error("WebGL 2 is required to check query results");
+ if (this.#gl.getQueryParameter(this.#query, this.#gl.QUERY_RESULT_AVAILABLE)) {
+ resolve(!!this.#gl.getQueryParameter(this.#query, this.#gl.QUERY_RESULT));
+ } else {
+ requestAnimationFrame(async () => {
+ const result = await _NanoPowGl.#checkQueryResult();
+ resolve(result);
+ });
+ }
+ } catch (err) {
+ reject(err);
+ }
+ });
+ }
+ /**
+ * 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 workBytes - Buffer with the original random nonce value
+ * @param workHex - Original nonce if provided for a validation call
+ * @returns Nonce as an 8-byte (16-char) hexadecimal string
+ */
+ static #readResult(workBytes, workHex) {
+ if (this.#gl == null) throw new Error("WebGL 2 is required to read pixels");
+ this.#gl.readPixels(0, 0, this.#gl.drawingBufferWidth, this.#gl.drawingBufferHeight, this.#gl.RGBA, this.#gl.UNSIGNED_BYTE, this.#pixels);
+ for (let i = 0; i < this.#pixels.length; i += 4) {
+ if (this.#pixels[i] !== 0) {
+ const hex = this.#hexify(workBytes.subarray(4, 8)) + this.#hexify([
+ this.#pixels[i + 2],
+ this.#pixels[i + 3],
+ workBytes[2] ^ this.#pixels[i] - 1,
+ workBytes[3] ^ this.#pixels[i + 1] - 1
+ ]);
+ 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) {
+ return new Promise((resolve) => {
+ setTimeout(async () => {
+ const result = this.search(hash, options);
+ resolve(result);
+ }, 100);
+ });
+ }
+ this.#busy = true;
+ if (_NanoPowGl.#gl == null) throw new Error("WebGL 2 is required");
+ if (this.#gl == null) throw new Error("WebGL 2 is required");
+ 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 > 16 ? 8 : options.effort;
+ const debug = !!options?.debug;
+ if (this.#WORKLOAD !== 256 * effort) {
+ this.#WORKLOAD = 256 * effort;
+ this.reset();
+ }
+ const uboView = new DataView(new ArrayBuffer(144));
+ for (let i = 0; i < 64; i += 8) {
+ const uint32 = hash.slice(i, i + 8);
+ uboView.setUint32(i * 2, parseInt(uint32, 16));
+ }
+ uboView.setUint32(128, threshold, true);
+ uboView.setFloat32(132, 256 * effort, true);
+ _NanoPowGl.#gl.bindBuffer(_NanoPowGl.#gl.UNIFORM_BUFFER, _NanoPowGl.#uboBuffer);
+ _NanoPowGl.#gl.bufferSubData(_NanoPowGl.#gl.UNIFORM_BUFFER, 0, uboView);
+ _NanoPowGl.#gl.bindBuffer(_NanoPowGl.#gl.UNIFORM_BUFFER, null);
+ let times = [];
+ let start = performance.now();
+ let nonce = null;
+ const seed = new Uint8Array(8);
+ while (nonce == null) {
+ start = performance.now();
+ crypto.getRandomValues(seed);
+ this.#draw(seed);
+ const found = await this.#checkQueryResult();
+ times.push(performance.now() - start);
+ if (found) {
+ nonce = this.#readResult(seed);
+ }
+ }
+ this.#busy = false;
+ if (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) {
+ return new Promise((resolve) => {
+ setTimeout(async () => {
+ const result = this.validate(work, hash, options);
+ resolve(result);
+ }, 100);
+ });
+ }
+ this.#busy = true;
+ if (_NanoPowGl.#gl == null) throw new Error("WebGL 2 is required");
+ if (this.#gl == null) throw new Error("WebGL 2 is required");
+ 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;
+ const debug = !!options?.debug;
+ const uboView = new DataView(new ArrayBuffer(144));
+ for (let i = 0; i < 64; i += 8) {
+ const uint32 = hash.slice(i, i + 8);
+ uboView.setUint32(i * 2, parseInt(uint32, 16));
+ }
+ uboView.setUint32(128, threshold, true);
+ uboView.setFloat32(132, _NanoPowGl.#WORKLOAD - 1, true);
+ _NanoPowGl.#gl.bindBuffer(_NanoPowGl.#gl.UNIFORM_BUFFER, _NanoPowGl.#uboBuffer);
+ _NanoPowGl.#gl.bufferSubData(_NanoPowGl.#gl.UNIFORM_BUFFER, 0, uboView);
+ _NanoPowGl.#gl.bindBuffer(_NanoPowGl.#gl.UNIFORM_BUFFER, null);
+ let nonce = null;
+ const data = new DataView(new ArrayBuffer(8));
+ data.setBigUint64(0, BigInt(`0x${work}`), true);
+ const seed = new Uint8Array(data.buffer);
+ this.#draw(seed);
+ let found = await this.#checkQueryResult();
+ if (found) {
+ try {
+ nonce = this.#readResult(seed, 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 #device = null;
+ static #uboBuffer;
+ static #gpuBuffer;
+ static #cpuBuffer;
+ 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. ${err}`);
+ } finally {
+ this.#busy = false;
+ }
+ }
+ static setup() {
+ if (this.#device == null) throw new Error(`WebGPU device failed to load.`);
+ this.#uboBuffer = this.#device.createBuffer({
+ size: 48,
+ usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST
+ });
+ 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.#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
+ }
+ });
+ }
+ 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, 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 > 2 && i > count * 0.1 && i < count * 0.9) truncated += times[i];
+ }
+ const averages = {
+ "Count (dispatches)": count,
+ "Total (ms)": sum,
+ "Rate (d/s)": 1e3 * count / (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.table(averages);
+ }
+ static async #dispatch(pipeline, seed, hash, threshold, passes) {
+ if (this.#device == null) throw new Error(`WebGPU device failed to load.`);
+ const uboView = new DataView(new ArrayBuffer(48));
+ for (let i = 0; i < 64; i += 8) {
+ const uint32 = hash.slice(i, i + 8);
+ uboView.setUint32(i / 2, parseInt(uint32, 16));
+ }
+ uboView.setBigUint64(32, seed, true);
+ uboView.setUint32(40, threshold, true);
+ this.#device.queue.writeBuffer(this.#uboBuffer, 0, uboView);
+ this.#device.queue.writeBuffer(this.#gpuBuffer, 0, new Uint32Array([0, 0, 0]));
+ 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 (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 (this.#busy) {
+ 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 TypeError(`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 ? 2048 : options.effort * 256;
+ const debug = !!options?.debug;
+ let loads = 0;
+ while (this.#device == null && loads < 20) {
+ await new Promise((resolve) => {
+ setTimeout(resolve, 500);
+ });
+ }
+ if (this.#device == null) throw new Error(`WebGPU device failed to load.`);
+ let times = [];
+ let start = performance.now();
+ let nonce = 0n;
+ do {
+ start = performance.now();
+ const random = Math.floor(Math.random() * 4294967295);
+ const seed = BigInt(random) << 32n | BigInt(random);
+ 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 (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 (this.#busy) {
+ 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 TypeError(`Invalid work ${work}`);
+ if (!/^[A-Fa-f0-9]{64}$/.test(hash)) throw new TypeError(`Invalid hash ${hash}`);
+ const 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) throw new Error(`WebGPU device failed to load.`);
+ const seed = BigInt(`0x${work}`);
+ const data = await this.#dispatch(this.#validatePipeline, seed, hash, threshold, 1);
+ const nonce = data.getBigUint64(0, true).toString(16).padStart(16, "0");
+ 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.`);
+ isGpuSupported = false;
+}
+try {
+ await NanoPowGl.init();
+ isGlSupported = true;
+} catch (err) {
+ console.warn(`WebGL is not supported in this environment.`);
+ isGlSupported = false;
+}
+var NanoPow = isGpuSupported ? NanoPowGpu : isGlSupported ? NanoPowGl : null;
+
+// src/main.ts
+var main_default = NanoPow;
+export {
+ NanoPow,
+ NanoPowGl,
+ NanoPowGpu,
+ main_default as default
+};