+++ /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
-};
projects / nano-pow.git / commitdiff