]> zoso.dev Git - nano-pow.git/commitdiff
Upload bundle and add URL to test page for mobile testing purposes.
authorChris Duncan <chris@zoso.dev>
Sat, 18 Jan 2025 23:06:27 +0000 (15:06 -0800)
committerChris Duncan <chris@zoso.dev>
Sat, 18 Jan 2025 23:06:27 +0000 (15:06 -0800)
main.min.js [new file with mode: 0644]
test.html

diff --git a/main.min.js b/main.min.js
new file mode 100644 (file)
index 0000000..9d6d6c9
--- /dev/null
@@ -0,0 +1,805 @@
+// 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
+};
index 64ef24add763b09029ca67647262d37d32939533..7f9b30d5c89555961027b7d23865649ea407c465 100644 (file)
--- a/test.html
+++ b/test.html
@@ -14,9 +14,14 @@ SPDX-License-Identifier: GPL-3.0-or-later
                } catch (err) {
                        console.warn(err)
                        try {
-                               ({ NanoPow, NanoPowGl, NanoPowGpu } = await import('https://cdn.jsdelivr.net/npm/nano-pow@latest/dist/main.min.js'))
+                               ({ NanoPow, NanoPowGl, NanoPowGpu } = await import('./main.min.js'))
                        } catch (err) {
-                               throw new Error(`Failed to load NanoPow ${err}`)
+                               console.warn(err)
+                               try {
+                                       ({ NanoPow, NanoPowGl, NanoPowGpu } = await import('https://cdn.jsdelivr.net/npm/nano-pow@latest/dist/main.min.js'))
+                               } catch (err) {
+                                       throw new Error(`Failed to load NanoPow ${err}`)
+                               }
                        }
                }