--- /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_IV_0=vec2(0xF2BDC900u,0x6A09E667u);const ROTATE_1=vec2(1u);const ROTATE_8=vec2(8u);const ROTATE_16=vec2(16u);const ROTATE_24=vec2(24u);const ROTATE_31=vec2(31u);var<workgroup> found:bool;@compute @workgroup_size(32)fn search(@builtin(global_invocation_id)global_id:vec3<u32>,@builtin(local_invocation_id)local_id:vec3<u32>){main(global_id);}@compute @workgroup_size(1)fn validate(@builtin(global_invocation_id)global_id:vec3<u32>){main(global_id);}fn main(id:vec3<u32>){let m0:vec2<u32>=ubo.random ^ id.xy;let m1:vec2<u32>=ubo.blockhash[0u].xy;let m2:vec2<u32>=ubo.blockhash[0u].zw;let m3:vec2<u32>=ubo.blockhash[1u].xy;let m4:vec2<u32>=ubo.blockhash[1u].zw;var v0:vec2<u32>=BLAKE2B_IV_0;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+vec2(0u,u32(v0.x+v4.x<v0.x));v0=v0+m0+vec2(0u,u32(v0.x+m0.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));v0=v0+m1+vec2(0u,u32(v0.x+m1.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));v1=v1+m2+vec2(0u,u32(v1.x+m2.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));v1=v1+m3+vec2(0u,u32(v1.x+m3.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));v2=v2+m4+vec2(0u,u32(v2.x+m4.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+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+vec2(0u,u32(v1.x+v5.x<v1.x));v1=v1+m4+vec2(0u,u32(v1.x+m4.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+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));v0=v0+m1+vec2(0u,u32(v0.x+m1.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));v1=v1+m0+vec2(0u,u32(v1.x+m0.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));v1=v1+m2+vec2(0u,u32(v1.x+m2.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));v3=v3+m3+vec2(0u,u32(v3.x+m3.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));v1=v1+m0+vec2(0u,u32(v1.x+m0.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));v2=v2+m2+vec2(0u,u32(v2.x+m2.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));v1=v1+m3+vec2(0u,u32(v1.x+m3.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));v2=v2+m1+vec2(0u,u32(v2.x+m1.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));v3=v3+m4+vec2(0u,u32(v3.x+m4.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));v1=v1+m3+vec2(0u,u32(v1.x+m3.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));v1=v1+m1+vec2(0u,u32(v1.x+m1.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+vec2(0u,u32(v0.x+v5.x<v0.x));v0=v0+m2+vec2(0u,u32(v0.x+m2.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));v2=v2+m4+vec2(0u,u32(v2.x+m4.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));v2=v2+m0+vec2(0u,u32(v2.x+m0.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));v0=v0+m0+vec2(0u,u32(v0.x+m0.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));v2=v2+m2+vec2(0u,u32(v2.x+m2.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));v2=v2+m4+vec2(0u,u32(v2.x+m4.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));v0=v0+m1+vec2(0u,u32(v0.x+m1.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));v3=v3+m3+vec2(0u,u32(v3.x+m3.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));v0=v0+m2+vec2(0u,u32(v0.x+m2.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));v2=v2+m0+vec2(0u,u32(v2.x+m0.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));v3=v3+m3+vec2(0u,u32(v3.x+m3.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));v0=v0+m4+vec2(0u,u32(v0.x+m4.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));v3=v3+m1+vec2(0u,u32(v3.x+m1.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));v1=v1+m1+vec2(0u,u32(v1.x+m1.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+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));v3=v3+m4+vec2(0u,u32(v3.x+m4.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));v0=v0+m0+vec2(0u,u32(v0.x+m0.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));v1=v1+m3+vec2(0u,u32(v1.x+m3.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));v2=v2+m2+vec2(0u,u32(v2.x+m2.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+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+vec2(0u,u32(v2.x+v6.x<v2.x));v2=v2+m1+vec2(0u,u32(v2.x+m1.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));v3=v3+m3+vec2(0u,u32(v3.x+m3.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));v0=v0+m0+vec2(0u,u32(v0.x+m0.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));v1=v1+m4+vec2(0u,u32(v1.x+m4.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));v3=v3+m2+vec2(0u,u32(v3.x+m2.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+vec2(0u,u32(v2.x+v6.x<v2.x));v2=v2+m3+vec2(0u,u32(v2.x+m3.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));v3=v3+m0+vec2(0u,u32(v3.x+m0.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));v0=v0+m2+vec2(0u,u32(v0.x+m2.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));v2=v2+m1+vec2(0u,u32(v2.x+m1.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));v2=v2+m4+vec2(0u,u32(v2.x+m4.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));v0=v0+m2+vec2(0u,u32(v0.x+m2.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));v1=v1+m4+vec2(0u,u32(v1.x+m4.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));v3=v3+m1+vec2(0u,u32(v3.x+m1.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));v2=v2+m3+vec2(0u,u32(v2.x+m3.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));v3=v3+m0+vec2(0u,u32(v3.x+m0.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));v0=v0+m0+vec2(0u,u32(v0.x+m0.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));v0=v0+m1+vec2(0u,u32(v0.x+m1.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));v1=v1+m2+vec2(0u,u32(v1.x+m2.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));v1=v1+m3+vec2(0u,u32(v1.x+m3.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));v2=v2+m4+vec2(0u,u32(v2.x+m4.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+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+vec2(0u,u32(v1.x+v5.x<v1.x));v1=v1+m4+vec2(0u,u32(v1.x+m4.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+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+vec2(0u,u32(v0.x+v5.x<v0.x));v0=v0+m1+vec2(0u,u32(v0.x+m1.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));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_IV_0.y ^ v0.y ^ v8.y)>ubo.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;
+
+in vec2 uv_pos;
+out uvec4 nonce;
+
+// 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
+layout(std140) uniform WORK {
+ uvec2 work;
+};
+
+// 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;
+
+// Used during G for vector bit rotations
+const uvec2 ROTATE_1 = uvec2(1u, 1u);
+const uvec2 ROTATE_8 = uvec2(8u, 8u);
+const uvec2 ROTATE_16 = uvec2(16u, 16u);
+const uvec2 ROTATE_24 = uvec2(24u, 24u);
+const uvec2 ROTATE_31 = uvec2(31u, 31u);
+
+// Both buffers represent 16 uint64s as 32 uint32s
+// because that's what GLSL offers, just like Javascript
+
+// Compression buffer, intialized to 2 instances of the initialization vector
+// The following values have been modified from the BLAKE2B_IV:
+// OUTLEN is constant 8 bytes
+// v[0] ^= 0x01010000u ^ uint(OUTLEN);
+// INLEN is constant 40 bytes: work value (8) + block hash (32)
+// v[24] ^= uint(INLEN);
+// It's always the "last" compression at this INLEN
+// v[28] = ~v[28];
+// v[29] = ~v[29];
+uvec2 v[16] = uvec2[16](
+ uvec2(0xF2BDC900u, 0x6A09E667u),
+ uvec2(0x84CAA73Bu, 0xBB67AE85u),
+ uvec2(0xFE94F82Bu, 0x3C6EF372u),
+ uvec2(0x5F1D36F1u, 0xA54FF53Au),
+ uvec2(0xADE682D1u, 0x510E527Fu),
+ uvec2(0x2B3E6C1Fu, 0x9B05688Cu),
+ uvec2(0xFB41BD6Bu, 0x1F83D9ABu),
+ uvec2(0x137E2179u, 0x5BE0CD19u),
+ uvec2(0xF3BCC908u, 0x6A09E667u),
+ uvec2(0x84CAA73Bu, 0xBB67AE85u),
+ uvec2(0xFE94F82Bu, 0x3C6EF372u),
+ uvec2(0x5F1D36F1u, 0xA54FF53Au),
+ uvec2(0xADE682F9u, 0x510E527Fu),
+ uvec2(0x2B3E6C1Fu, 0x9B05688Cu),
+ uvec2(0x04BE4294u, 0xE07C2654u),
+ uvec2(0x137E2179u, 0x5BE0CD19u)
+);
+
+// Input data buffer
+uvec2 m[16];
+
+// Offsets into the input data buffer for each mixing step
+const uint SIGMA[192] = uint[192](
+ 0u,1u,2u,3u,4u,5u,6u,7u,8u,9u,10u,11u,12u,13u,14u,15u,
+ 14u,10u,4u,8u,9u,15u,13u,6u,1u,12u,0u,2u,11u,7u,5u,3u,
+ 11u,8u,12u,0u,5u,2u,15u,13u,10u,14u,3u,6u,7u,1u,9u,4u,
+ 7u,9u,3u,1u,13u,12u,11u,14u,2u,6u,5u,10u,4u,0u,15u,8u,
+ 9u,0u,5u,7u,2u,4u,10u,15u,14u,1u,11u,12u,6u,8u,3u,13u,
+ 2u,12u,6u,10u,0u,11u,8u,3u,4u,13u,7u,5u,15u,14u,1u,9u,
+ 12u,5u,1u,15u,14u,13u,4u,10u,0u,7u,6u,3u,9u,2u,8u,11u,
+ 13u,11u,7u,14u,12u,1u,3u,9u,5u,0u,15u,4u,8u,6u,2u,10u,
+ 6u,15u,14u,9u,11u,3u,0u,8u,12u,2u,13u,7u,1u,4u,10u,5u,
+ 10u,2u,8u,4u,7u,6u,1u,5u,15u,11u,9u,14u,3u,12u,13u,0u,
+ 0u,1u,2u,3u,4u,5u,6u,7u,8u,9u,10u,11u,12u,13u,14u,15u,
+ 14u,10u,4u,8u,9u,15u,13u,6u,1u,12u,0u,2u,11u,7u,5u,3u
+);
+
+// G mixing function
+void G (uint a, uint b, uint c, uint d, uint x, uint y) {
+ v[a] = v[a] + v[b] + uvec2(0u, uint(v[a].x + v[b].x < v[b].x));
+ v[a] = v[a] + m[x] + uvec2(0u, uint(v[a].x + m[x].x < m[x].x));
+ v[d] = (v[d] ^ v[a]).yx;
+ v[c] = v[c] + v[d] + uvec2(0u, uint(v[c].x + v[d].x < v[d].x));
+ v[b] = ((v[b] ^ v[c]) >> ROTATE_24) | ((v[b] ^ v[c]).yx << ROTATE_8);
+ v[a] = v[a] + v[b] + uvec2(0u, uint(v[a].x + v[b].x < v[b].x));
+ v[a] = v[a] + m[y] + uvec2(0u, uint(v[a].x + m[y].x < m[y].x));
+ v[d] = ((v[d] ^ v[a]) >> ROTATE_16) | ((v[d] ^ v[a]).yx << ROTATE_16);
+ v[c] = v[c] + v[d] + uvec2(0u, uint(v[c].x + v[d].x < v[d].x));
+ v[b] = ((v[b] ^ v[c]).yx >> ROTATE_31) | ((v[b] ^ v[c]) << ROTATE_1);
+}
+
+void main() {
+ // Nonce uniquely differentiated by pixel location
+ m[0u].x = work.x ^ uint(uv_pos.x * workload);
+ m[0u].y = work.y ^ uint(uv_pos.y * workload);
+
+ // Block hash
+ m[1u] = uvec2(blockhash[0u], blockhash[1u]);
+ m[2u] = uvec2(blockhash[2u], blockhash[3u]);
+ m[3u] = uvec2(blockhash[4u], blockhash[5u]);
+ m[4u] = uvec2(blockhash[6u], blockhash[7u]);
+
+ // twelve rounds of mixing
+ for(uint i = 0u; i < 12u; i = i + 1u) {
+ G(0u, 4u, 8u, 12u, SIGMA[i * 16u + 0u], SIGMA[i * 16u + 1u]);
+ G(1u, 5u, 9u, 13u, SIGMA[i * 16u + 2u], SIGMA[i * 16u + 3u]);
+ G(2u, 6u, 10u, 14u, SIGMA[i * 16u + 4u], SIGMA[i * 16u + 5u]);
+ G(3u, 7u, 11u, 15u, SIGMA[i * 16u + 6u], SIGMA[i * 16u + 7u]);
+ G(0u, 5u, 10u, 15u, SIGMA[i * 16u + 8u], SIGMA[i * 16u + 9u]);
+ G(1u, 6u, 11u, 12u, SIGMA[i * 16u + 10u], SIGMA[i * 16u + 11u]);
+ G(2u, 7u, 8u, 13u, SIGMA[i * 16u + 12u], SIGMA[i * 16u + 13u]);
+ G(3u, 4u, 9u, 14u, SIGMA[i * 16u + 14u], SIGMA[i * 16u + 15u]);
+ }
+
+ // Pixel data set from work values
+ // Finalize digest from high bits, low bits can be safely ignored
+ if ((BLAKE2B_IV32_1 ^ v[0u].y ^ v[8u].y) > threshold) {
+ nonce = uvec4(1u, m[0].y, m[0].x, 1u);
+ } 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 #gl;
+ static #program;
+ static #vertexShader;
+ static #fragmentShader;
+ static #texture;
+ static #framebuffer;
+ 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.#texture = this.#gl.createTexture();
+ this.#gl.bindTexture(this.#gl.TEXTURE_2D, this.#texture);
+ this.#gl.texImage2D(this.#gl.TEXTURE_2D, 0, this.#gl.RGBA32UI, this.#gl.drawingBufferWidth, this.#gl.drawingBufferHeight, 0, this.#gl.RGBA_INTEGER, this.#gl.UNSIGNED_INT, null);
+ this.#gl.texParameteri(this.#gl.TEXTURE_2D, this.#gl.TEXTURE_MIN_FILTER, this.#gl.NEAREST);
+ this.#gl.texParameteri(this.#gl.TEXTURE_2D, this.#gl.TEXTURE_MAG_FILTER, this.#gl.NEAREST);
+ this.#gl.bindTexture(this.#gl.TEXTURE_2D, null);
+ this.#framebuffer = this.#gl.createFramebuffer();
+ this.#gl.bindFramebuffer(this.#gl.FRAMEBUFFER, this.#framebuffer);
+ this.#gl.framebufferTexture2D(this.#gl.FRAMEBUFFER, this.#gl.COLOR_ATTACHMENT0, this.#gl.TEXTURE_2D, this.#texture, 0);
+ if (this.#gl.checkFramebufferStatus(this.#gl.FRAMEBUFFER) !== this.#gl.FRAMEBUFFER_COMPLETE)
+ throw new Error(`Failed to create framebuffer`);
+ this.#gl.bindFramebuffer(this.#gl.FRAMEBUFFER, null);
+ 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.#gl.bindBuffer(this.#gl.ARRAY_BUFFER, null);
+ 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.#gl.bindBuffer(this.#gl.ARRAY_BUFFER, null);
+ this.#uboBuffer = this.#gl.createBuffer();
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, this.#uboBuffer);
+ this.#gl.bufferData(this.#gl.UNIFORM_BUFFER, 144, this.#gl.DYNAMIC_DRAW);
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, null);
+ this.#gl.bindBufferBase(this.#gl.UNIFORM_BUFFER, 0, this.#uboBuffer);
+ this.#gl.uniformBlockBinding(this.#program, this.#gl.getUniformBlockIndex(this.#program, "UBO"), 0);
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, null);
+ 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.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, null);
+ this.#pixels = new Uint32Array(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.#framebuffer = null;
+ _NanoPowGl.#texture = 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 < 3 || i > count * 0.1 && i < count * 0.9) truncated += times[i];
+ }
+ const averages = {
+ "Count (frames)": count,
+ "Total (ms)": sum,
+ "Rate (f/s)": 1e3 * count * 0.8 / (truncated || sum),
+ "Minimum (ms)": min,
+ "Maximum (ms)": max,
+ "Arithmetic Mean (ms)": sum / count,
+ "Truncated Mean (ms)": truncated / count,
+ "Harmonic Mean (ms)": count / reciprocals,
+ "Geometric Mean (ms)": Math.exp(logarithms / count)
+ };
+ console.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, work);
+ this.#gl.bindBuffer(this.#gl.UNIFORM_BUFFER, null);
+ this.#gl.beginQuery(this.#gl.ANY_SAMPLES_PASSED_CONSERVATIVE, this.#query);
+ this.#gl.bindFramebuffer(this.#gl.FRAMEBUFFER, this.#framebuffer);
+ this.#gl.drawArrays(this.#gl.TRIANGLES, 0, 6);
+ this.#gl.bindFramebuffer(this.#gl.FRAMEBUFFER, null);
+ 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 workHex - Original nonce if provided for a validation call
+ * @returns Nonce as an 8-byte (16-char) hexadecimal string
+ */
+ static #readResult(workHex) {
+ if (this.#gl == null) throw new Error("WebGL 2 is required to read pixels");
+ this.#gl.bindFramebuffer(this.#gl.FRAMEBUFFER, this.#framebuffer);
+ this.#gl.readPixels(0, 0, this.#gl.drawingBufferWidth, this.#gl.drawingBufferHeight, this.#gl.RGBA_INTEGER, this.#gl.UNSIGNED_INT, this.#pixels);
+ this.#gl.bindFramebuffer(this.#gl.FRAMEBUFFER, null);
+ for (let i = 0; i < this.#pixels.length; i += 4) {
+ if (this.#pixels[i] !== 0) {
+ const hex = `${this.#pixels[i + 1].toString(16).padStart(8, "0")}${this.#pixels[i + 2].toString(16).padStart(8, "0")}`;
+ if (workHex == null || workHex == hex) return hex;
+ }
+ }
+ throw new Error("Query reported result but nonce value not found");
+ }
+ /**
+ * Finds a nonce that satisfies the Nano proof-of-work requirements.
+ *
+ * @param {string} hash - Hexadecimal hash of previous block, or public key for new accounts
+ * @param {number} [threshold=0xfffffff8] - Difficulty of proof-of-work calculation
+ */
+ static async search(hash, options) {
+ if (_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}`);
+ if (this.#busy) {
+ return new Promise((resolve) => {
+ setTimeout(async () => {
+ const result = this.search(hash, options);
+ resolve(result);
+ }, 100);
+ });
+ }
+ this.#busy = true;
+ const threshold = typeof options?.threshold !== "number" || options.threshold < 0 || options.threshold > 4294967295 ? 4294967288 : options.threshold;
+ const effort = typeof options?.effort !== "number" || options.effort < 1 || options.effort > 32 ? 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();
+ }
+ }
+ 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 (_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}`);
+ if (this.#busy) {
+ return new Promise((resolve) => {
+ setTimeout(async () => {
+ const result = this.validate(work, hash, options);
+ resolve(result);
+ }, 100);
+ });
+ }
+ this.#busy = true;
+ const threshold = typeof options?.threshold !== "number" || options.threshold < 0 || options.threshold > 4294967295 ? 4294967288 : options.threshold;
+ const debug = !!options?.debug;
+ if (this.#WORKLOAD !== 1) {
+ this.#WORKLOAD = 1;
+ 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, _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(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, truncatedCount = 0, truncated = 0, sum = 0, reciprocals = 0, logarithms = 0, min = Number.MAX_SAFE_INTEGER, max = 0, median = 0, rate = 0;
+ times.sort();
+ for (let i = 0; i < count; i++) {
+ sum += times[i];
+ reciprocals += 1 / times[i];
+ logarithms += Math.log(times[i]);
+ min = Math.min(min, times[i]);
+ max = Math.max(max, times[i]);
+ if (i === Math.ceil(count / 2)) {
+ median = times[i];
+ }
+ if (count < 3 || i > 0.1 * count && i < 0.9 * (count - 1)) {
+ truncated += times[i];
+ truncatedCount++;
+ }
+ }
+ const averages = {
+ "Count (dispatches)": count,
+ "Total (ms)": sum,
+ "Rate (d/s)": 1e3 * truncatedCount / (truncated || sum),
+ "Minimum (ms)": min,
+ "Maximum (ms)": max,
+ "Median (ms)": median,
+ "Arithmetic Mean (ms)": sum / count,
+ "Truncated Mean (ms)": truncated / truncatedCount,
+ "Harmonic Mean (ms)": count / reciprocals,
+ "Geometric Mean (ms)": Math.exp(logarithms / count)
+ };
+ console.table(averages);
+ }
+ static async #dispatch(pipeline, seed, hash, threshold, passes) {
+ if (this.#device == null) throw new Error(`WebGPU device failed to load.`);
+ const uboView = new DataView(new ArrayBuffer(48));
+ for (let i = 0; i < 64; i += 16) {
+ const u64 = hash.slice(i, i + 16);
+ uboView.setBigUint64(i / 2, BigInt(`0x${u64}`));
+ }
+ 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 (!/^[A-Fa-f0-9]{64}$/.test(hash)) throw new TypeError(`Invalid hash ${hash}`);
+ if (this.#busy) {
+ return new Promise((resolve) => {
+ setTimeout(async () => {
+ const result = this.search(hash, options);
+ resolve(result);
+ }, 100);
+ });
+ }
+ this.#busy = true;
+ const threshold = typeof options?.threshold !== "number" || options.threshold < 0 || options.threshold > 4294967295 ? 4294967288 : options.threshold;
+ const effort = typeof options?.effort !== "number" || options.effort < 1 || options.effort > 32 ? 2048 : options.effort * 256;
+ const debug = !!options?.debug;
+ let loads = 0;
+ while (this.#device == null && loads < 20) {
+ await new Promise((resolve) => {
+ setTimeout(resolve, 500);
+ });
+ }
+ if (this.#device == null) {
+ this.#busy = false;
+ throw new Error(`WebGPU device failed to load.`);
+ }
+ let times = [];
+ let start = performance.now();
+ let nonce = 0n;
+ do {
+ start = performance.now();
+ const 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 (!/^[A-Fa-f0-9]{16}$/.test(work)) throw new TypeError(`Invalid work ${work}`);
+ if (!/^[A-Fa-f0-9]{64}$/.test(hash)) throw new TypeError(`Invalid hash ${hash}`);
+ if (this.#busy) {
+ return new Promise((resolve) => {
+ setTimeout(async () => {
+ const result = this.validate(work, hash, options);
+ resolve(result);
+ }, 100);
+ });
+ }
+ this.#busy = true;
+ 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) {
+ this.#busy = false;
+ 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.error(err);
+ console.warn(`WebGPU is not supported in this environment.`);
+ isGpuSupported = false;
+}
+try {
+ await NanoPowGl.init();
+ isGlSupported = true;
+} catch (err) {
+ console.error(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