From 066cb0bf6f2275c34c70accfaecd992a8c2a7c9a Mon Sep 17 00:00:00 2001 From: Chris Duncan Date: Fri, 10 Jan 2025 21:23:02 -0800 Subject: [PATCH] Finalize testing page. Track dist folder so that mobile testing can see its files. --- .gitignore | 3 +- dist/global.min.js | 598 +++++++++++++++++++++++++++++++++++++++++++++ dist/main.min.js | 586 ++++++++++++++++++++++++++++++++++++++++++++ test.html | 108 ++++---- 4 files changed, 1241 insertions(+), 54 deletions(-) create mode 100644 dist/global.min.js create mode 100644 dist/main.min.js diff --git a/.gitignore b/.gitignore index d42d158..03b8885 100644 --- a/.gitignore +++ b/.gitignore @@ -15,6 +15,5 @@ coverage.info # Dependency directories node_modules/ -# esbuild +# tsc build/ -dist/ diff --git a/dist/global.min.js b/dist/global.min.js new file mode 100644 index 0000000..5c1386d --- /dev/null +++ b/dist/global.min.js @@ -0,0 +1,598 @@ +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 new file mode 100644 index 0000000..850d6e7 --- /dev/null +++ b/dist/main.min.js @@ -0,0 +1,586 @@ +// 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 +}; diff --git a/test.html b/test.html index dbe1a14..bbf90f1 100644 --- a/test.html +++ b/test.html @@ -13,68 +13,72 @@ SPDX-License-Identifier: GPL-3.0-or-later

nano-pow

-
- +
+

+	
+

+
 
-	
+
-- 
2.34.1