From: Chris Duncan Date: Sat, 11 Jan 2025 06:02:26 +0000 (-0800) Subject: Remove build artifacts. X-Git-Tag: v0.0.1~2 X-Git-Url: https://zoso.dev/?a=commitdiff_plain;h=5c2d82852311cf29fb4e74c4c6c7d3a31c02043e;p=nano-pow.git Remove build artifacts. --- diff --git a/.gitignore b/.gitignore index f0cedc5..cb6bf43 100644 --- a/.gitignore +++ b/.gitignore @@ -17,4 +17,3 @@ node_modules/ # Build artifacts build/ -dist/ diff --git a/dist/global.min.js b/dist/global.min.js deleted file mode 100644 index 5c1386d..0000000 --- a/dist/global.min.js +++ /dev/null @@ -1,598 +0,0 @@ -var __defProp = Object.defineProperty; -var __export = (target, all) => { - for (var name in all) - __defProp(target, name, { get: all[name], enumerable: true }); -}; - -// build/main.js -var main_exports = {}; -__export(main_exports, { - NanoPowGl: () => NanoPowGl, - NanoPowGpu: () => NanoPowGpu -}); - -// src/shaders/compute.wgsl -var compute_default = "struct UBO{blockhash:array,2>,random:u32,threshold:u32};@group(0)@binding(0)var ubo:UBO;struct WORK{nonce:vec2,found:atomic};@group(0)@binding(1)varwork:WORK;const BLAKE2B_IV32_1:u32=0x6A09E667u;@compute @workgroup_size(256)fn main(@builtin(global_invocation_id)id:vec3){if(atomicLoad(&work.found)!=0u){return;}let threshold:u32=ubo.threshold;var m0:u32=ubo.random ^ id.x;var m1:u32=ubo.random ^ id.y;var m2:u32=ubo.blockhash[0u].x;var m3:u32=ubo.blockhash[0u].y;var m4:u32=ubo.blockhash[0u].z;var m5:u32=ubo.blockhash[0u].w;var m6:u32=ubo.blockhash[1u].x;var m7:u32=ubo.blockhash[1u].y;var m8:u32=ubo.blockhash[1u].z;var m9:u32=ubo.blockhash[1u].w;var v0:u32=0xF2BDC900u;var v1:u32=0x6A09E667u;var v2:u32=0x84CAA73Bu;var v3:u32=0xBB67AE85u;var v4:u32=0xFE94F82Bu;var v5:u32=0x3C6EF372u;var v6:u32=0x5F1D36F1u;var v7:u32=0xA54FF53Au;var v8:u32=0xADE682D1u;var v9:u32=0x510E527Fu;var v10:u32=0x2B3E6C1Fu;var v11:u32=0x9B05688Cu;var v12:u32=0xFB41BD6Bu;var v13:u32=0x1F83D9ABu;var v14:u32=0x137E2179u;var v15:u32=0x5BE0CD19u;var v16:u32=0xF3BCC908u;var v17:u32=0x6A09E667u;var v18:u32=0x84CAA73Bu;var v19:u32=0xBB67AE85u;var v20:u32=0xFE94F82Bu;var v21:u32=0x3C6EF372u;var v22:u32=0x5F1D36F1u;var v23:u32=0xA54FF53Au;var v24:u32=0xADE682F9u;var v25:u32=0x510E527Fu;var v26:u32=0x2B3E6C1Fu;var v27:u32=0x9B05688Cu;var v28:u32=0x04BE4294u;var v29:u32=0xE07C2654u;var v30:u32=0x137E2179u;var v31:u32=0x5BE0CD19u;var o0:u32;var o1:u32;var xor0:u32;var xor1:u32;o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);if((BLAKE2B_IV32_1 ^ v1 ^ v17)>threshold&&atomicLoad(&work.found)==0u){atomicStore(&work.found,1u);work.nonce.x=m0;work.nonce.y=m1;}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 { - /** 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 - ]); - static { - 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(); - } - /** - * 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, threshold = 4294967288) { - if (_NanoPowGl.#gl == null) throw new Error("WebGL 2 is required"); - if (!/^[A-F-a-f0-9]{64}$/.test(hash)) throw new Error(`invalid_hash ${hash}`); - if (typeof threshold !== "number") throw new TypeError(`Invalid threshold ${threshold}`); - if (this.#gl == null) throw new Error("WebGL 2 is required"); - 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 work = new Uint8Array(8); - while (nonce == null) { - this.draw(work); - const found = await this.checkQueryResult(); - if (found) { - nonce = this.readResult(work); - } - } - return nonce; - } - 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); - crypto.getRandomValues(work); - 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() { - 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)) { - return this.#gl.getQueryParameter(this.#query, this.#gl.QUERY_RESULT); - } - return new Promise((resolve, reject) => { - try { - 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 work - Buffer with the original random nonce value - * @returns Nonce as an 8-byte (16-char) hexadecimal string - */ - static readResult(work) { - 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(work.subarray(4, 8)) + this.#hexify([ - this.#pixels[i + 2], - this.#pixels[i + 3], - work[2] ^ this.#pixels[i] - 1, - work[3] ^ this.#pixels[i + 1] - 1 - ]); - return hex; - } - } - throw new Error("Query reported result but nonce value not 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 #pipeline; - static { - this.init(); - } - // Initialize WebGPU - static async init() { - if (this.#busy) return; - this.#busy = true; - if (navigator.gpu == null) { - throw new Error("WebGPU is not supported in this browser."); - } - try { - 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" } - } - ] - }); - this.#pipeline = this.#device.createComputePipeline({ - layout: this.#device.createPipelineLayout({ - bindGroupLayouts: [this.#bindGroupLayout] - }), - compute: { - entryPoint: "main", - module: this.#device.createShaderModule({ - code: compute_default - }) - } - }); - } - static reset() { - console.warn(`GPU device lost. Reinitializing...`); - _NanoPowGpu.#cpuBuffer?.destroy(); - _NanoPowGpu.#gpuBuffer?.destroy(); - _NanoPowGpu.#uboBuffer?.destroy(); - _NanoPowGpu.#busy = false; - _NanoPowGpu.init(); - } - /** - * 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, threshold = 4294967288) { - if (this.#busy) { - return new Promise((resolve) => { - setTimeout(async () => { - const result = this.search(hash, threshold); - resolve(result); - }, 100); - }); - } - this.#busy = true; - if (!/^[A-Fa-f0-9]{64}$/.test(hash)) throw new TypeError(`Invalid hash ${hash}`); - if (typeof threshold !== "number") throw new TypeError(`Invalid threshold ${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.`); - let nonce = 0n; - do { - 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)); - } - const random = Math.floor(Math.random() * 4294967295); - uboView.setUint32(32, random, true); - uboView.setUint32(36, 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(this.#pipeline); - passEncoder.setBindGroup(0, bindGroup); - passEncoder.dispatchWorkgroups(255, 255); - 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.search(hash, threshold); - } - if (data == null) throw new Error(`Failed to get data from buffer.`); - nonce = data.getBigUint64(0, true); - this.#busy = !data.getUint32(8); - } while (this.#busy); - return nonce.toString(16).padStart(16, "0"); - } -}; - -// build/global.js -globalThis.NanoPow ??= main_exports; diff --git a/dist/main.min.js b/dist/main.min.js deleted file mode 100644 index 850d6e7..0000000 --- a/dist/main.min.js +++ /dev/null @@ -1,586 +0,0 @@ -// src/shaders/compute.wgsl -var compute_default = "struct UBO{blockhash:array,2>,random:u32,threshold:u32};@group(0)@binding(0)var ubo:UBO;struct WORK{nonce:vec2,found:atomic};@group(0)@binding(1)varwork:WORK;const BLAKE2B_IV32_1:u32=0x6A09E667u;@compute @workgroup_size(256)fn main(@builtin(global_invocation_id)id:vec3){if(atomicLoad(&work.found)!=0u){return;}let threshold:u32=ubo.threshold;var m0:u32=ubo.random ^ id.x;var m1:u32=ubo.random ^ id.y;var m2:u32=ubo.blockhash[0u].x;var m3:u32=ubo.blockhash[0u].y;var m4:u32=ubo.blockhash[0u].z;var m5:u32=ubo.blockhash[0u].w;var m6:u32=ubo.blockhash[1u].x;var m7:u32=ubo.blockhash[1u].y;var m8:u32=ubo.blockhash[1u].z;var m9:u32=ubo.blockhash[1u].w;var v0:u32=0xF2BDC900u;var v1:u32=0x6A09E667u;var v2:u32=0x84CAA73Bu;var v3:u32=0xBB67AE85u;var v4:u32=0xFE94F82Bu;var v5:u32=0x3C6EF372u;var v6:u32=0x5F1D36F1u;var v7:u32=0xA54FF53Au;var v8:u32=0xADE682D1u;var v9:u32=0x510E527Fu;var v10:u32=0x2B3E6C1Fu;var v11:u32=0x9B05688Cu;var v12:u32=0xFB41BD6Bu;var v13:u32=0x1F83D9ABu;var v14:u32=0x137E2179u;var v15:u32=0x5BE0CD19u;var v16:u32=0xF3BCC908u;var v17:u32=0x6A09E667u;var v18:u32=0x84CAA73Bu;var v19:u32=0xBB67AE85u;var v20:u32=0xFE94F82Bu;var v21:u32=0x3C6EF372u;var v22:u32=0x5F1D36F1u;var v23:u32=0xA54FF53Au;var v24:u32=0xADE682F9u;var v25:u32=0x510E527Fu;var v26:u32=0x2B3E6C1Fu;var v27:u32=0x9B05688Cu;var v28:u32=0x04BE4294u;var v29:u32=0xE07C2654u;var v30:u32=0x137E2179u;var v31:u32=0x5BE0CD19u;var o0:u32;var o1:u32;var xor0:u32;var xor1:u32;o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v0+v8;o1=v1+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v16+v24;o1=v17+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v2+v10;o1=v3+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v18+v26;o1=v19+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v4+v12;o1=v5+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v20+v28;o1=v21+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v6+v14;o1=v7+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v22+v30;o1=v23+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v11=(xor1>>24u)^(xor0<<8u);o0=v0+v10;o1=v1+v11;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v31=(xor1>>16u)^(xor0<<16u);o0=v20+v30;o1=v21+v31;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v11=(xor0>>31u)^(xor1<<1u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v13=(xor1>>24u)^(xor0<<8u);o0=v2+v12;o1=v3+v13;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v25=(xor1>>16u)^(xor0<<16u);o0=v22+v24;o1=v23+v25;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v13=(xor0>>31u)^(xor1<<1u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v15=(xor1>>24u)^(xor0<<8u);o0=v4+v14;o1=v5+v15;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v27=(xor1>>16u)^(xor0<<16u);o0=v16+v26;o1=v17+v27;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v15=(xor0>>31u)^(xor1<<1u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>24u)^(xor1<<8u);v9=(xor1>>24u)^(xor0<<8u);o0=v6+v8;o1=v7+v9;o1=o1+select(0u,1u,o0>16u)^(xor1<<16u);v29=(xor1>>16u)^(xor0<<16u);o0=v18+v28;o1=v19+v29;o1=o1+select(0u,1u,o0>31u)^(xor0<<1u);v9=(xor0>>31u)^(xor1<<1u);if((BLAKE2B_IV32_1 ^ v1 ^ v17)>threshold&&atomicLoad(&work.found)==0u){atomicStore(&work.found,1u);work.nonce.x=m0;work.nonce.y=m1;}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 { - /** 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 - ]); - static { - 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(); - } - /** - * 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, threshold = 4294967288) { - if (_NanoPowGl.#gl == null) throw new Error("WebGL 2 is required"); - if (!/^[A-F-a-f0-9]{64}$/.test(hash)) throw new Error(`invalid_hash ${hash}`); - if (typeof threshold !== "number") throw new TypeError(`Invalid threshold ${threshold}`); - if (this.#gl == null) throw new Error("WebGL 2 is required"); - 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 work = new Uint8Array(8); - while (nonce == null) { - this.draw(work); - const found = await this.checkQueryResult(); - if (found) { - nonce = this.readResult(work); - } - } - return nonce; - } - 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); - crypto.getRandomValues(work); - 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() { - 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)) { - return this.#gl.getQueryParameter(this.#query, this.#gl.QUERY_RESULT); - } - return new Promise((resolve, reject) => { - try { - 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 work - Buffer with the original random nonce value - * @returns Nonce as an 8-byte (16-char) hexadecimal string - */ - static readResult(work) { - 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(work.subarray(4, 8)) + this.#hexify([ - this.#pixels[i + 2], - this.#pixels[i + 3], - work[2] ^ this.#pixels[i] - 1, - work[3] ^ this.#pixels[i + 1] - 1 - ]); - return hex; - } - } - throw new Error("Query reported result but nonce value not 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 #pipeline; - static { - this.init(); - } - // Initialize WebGPU - static async init() { - if (this.#busy) return; - this.#busy = true; - if (navigator.gpu == null) { - throw new Error("WebGPU is not supported in this browser."); - } - try { - 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" } - } - ] - }); - this.#pipeline = this.#device.createComputePipeline({ - layout: this.#device.createPipelineLayout({ - bindGroupLayouts: [this.#bindGroupLayout] - }), - compute: { - entryPoint: "main", - module: this.#device.createShaderModule({ - code: compute_default - }) - } - }); - } - static reset() { - console.warn(`GPU device lost. Reinitializing...`); - _NanoPowGpu.#cpuBuffer?.destroy(); - _NanoPowGpu.#gpuBuffer?.destroy(); - _NanoPowGpu.#uboBuffer?.destroy(); - _NanoPowGpu.#busy = false; - _NanoPowGpu.init(); - } - /** - * 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, threshold = 4294967288) { - if (this.#busy) { - return new Promise((resolve) => { - setTimeout(async () => { - const result = this.search(hash, threshold); - resolve(result); - }, 100); - }); - } - this.#busy = true; - if (!/^[A-Fa-f0-9]{64}$/.test(hash)) throw new TypeError(`Invalid hash ${hash}`); - if (typeof threshold !== "number") throw new TypeError(`Invalid threshold ${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.`); - let nonce = 0n; - do { - 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)); - } - const random = Math.floor(Math.random() * 4294967295); - uboView.setUint32(32, random, true); - uboView.setUint32(36, 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(this.#pipeline); - passEncoder.setBindGroup(0, bindGroup); - passEncoder.dispatchWorkgroups(255, 255); - 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.search(hash, threshold); - } - if (data == null) throw new Error(`Failed to get data from buffer.`); - nonce = data.getBigUint64(0, true); - this.#busy = !data.getUint32(8); - } while (this.#busy); - return nonce.toString(16).padStart(16, "0"); - } -}; -export { - NanoPowGl, - NanoPowGpu -};