From a90eab9386c90ebff856a2852eab759252a36072 Mon Sep 17 00:00:00 2001 From: Chris Duncan Date: Fri, 11 Apr 2025 22:13:42 -0700 Subject: [PATCH] Simplify GPU compute by reverting to vec2 and implementing G function since performance difference is almost nonexistent and the real benefit comes from native u64 types which do not yet exist in WGSL. Increase workgroup size to best value 96 found during testing on RTX 3070. Create --benchmark CLI argument. Store seed and blockhash in fast shared workgroup memory. Add benchmark convenience npm script. Fix documentation in inline help and manual page. Fix CLI default port collision with server default port. Change timestamp on server log files. Test concurrent curl requests to server. Fix action listed in server response error message. Update puppeteer and test page HTML with blank favicon to prevent unnecessary load. --- docs/nano-pow.1 | 33 +- package.json | 1 + src/bin/cli.ts | 48 +- src/bin/nano-pow.sh | 2 +- src/bin/server.ts | 11 +- src/lib/gpu/compute.wgsl | 1479 +++++--------------------------------- test/index.html | 2 +- test/script.sh | 9 +- 8 files changed, 261 insertions(+), 1324 deletions(-) diff --git a/docs/nano-pow.1 b/docs/nano-pow.1 index a246a1d..998dfef 100644 --- a/docs/nano-pow.1 +++ b/docs/nano-pow.1 @@ -28,19 +28,22 @@ Start work server (see SERVER below). Must be the first argument in order to be \fB\-h\fR, \fB\-\-help\fR Show this help dialog and exit. .TP -\fB\-d\fR, \fB\-\-debug\fR +\fB\-\-debug\fR Enable additional logging output. .TP -\fB\-j\fR, \fB\-\-json\fR -Format final output of all hashes as a JSON array of stringified values instead of incrementally returning Javascript objects for each result. +\fB\-\-benchmark\fR \fICOUNT\fR +Generate work for the specified number of random hashes. .TP -\fB\-e\fR, \fB\-\-effort\fR=\fIEFFORT\fR -Increase demand on GPU processing. Must be between 1 and 32 inclusive. +\fB\-b\fR, \fB\-\-batch\fR +Format final output of all hashes as a JSON array instead of incrementally returning an object for each result as soon as it is calculated. .TP -\fB\-t\fR, \fB\-\-difficulty\fR=\fIDIFFICULTY\fR +\fB\-t\fR, \fB\-\-difficulty\fR \fIDIFFICULTY\fR Override the minimum threshold value. Higher values increase difficulty. Must be a hexadecimal string between 1 and FFFFFFFFFFFFFFFF inclusive. .TP -\fB\-v\fR, \fB\-\-validate\fR=\fIWORK\fR +\fB\-e\fR, \fB\-\-effort\fR \fIEFFORT\fR +Increase demand on GPU processing. Must be between 1 and 32 inclusive. +.TP +\fB\-v\fR, \fB\-\-validate\fR \fIWORK\fR Check an existing work value instead of searching for one. If you pass multiple blockhashes, the work value will be validated against each one. .SH SERVER @@ -95,7 +98,7 @@ Validate an existing work value against a blockhash. Requires \fBwork\fR field c .SH EXAMPLES - CLI .PP -Search for a work nonce for a blockhash using the default difficulty 0xFFFFFFF800000000: +Search for a work nonce for a blockhash using the default difficulty FFFFFFF800000000: .EX $ nano-pow \fB0123456789abcdef0123456789abcdef0123456789abcdef0123456789abcdef\fR .EE @@ -103,19 +106,25 @@ $ nano-pow \fB0123456789abcdef0123456789abcdef0123456789abcdef0123456789abcdef\f .PP Search for a work nonce using a custom difficulty and increased effort: .EX -$ nano-pow \fB\-t fffffe0000000000 \-e 32 0123456789abcdef0123456789abcdef0123456789abcdef0123456789abcdef\fR +$ nano-pow \fB\-d fffffe0000000000 \-e 32 0123456789abcdef0123456789abcdef0123456789abcdef0123456789abcdef\fR .EE .PP -Read one or more lines of blockhashes from a file and output to another file: +Read one or more lines of blockhashes from a file and output results to another file: .EX -$ cat /path/to/file.txt | nano-pow --json > work.json +$ cat /path/to/file.txt | nano-pow \fB--batch\fR > work.json .EE .PP Validate an existing work nonce against a blockhash and show debugging output: .EX -$ nano-pow \fB\-d \-v fedcba9876543210 0123456789abcdef0123456789abcdef0123456789abcdef0123456789abcdef\fR +$ nano-pow \fB\--debug \-v fedcba9876543210 0123456789abcdef0123456789abcdef0123456789abcdef0123456789abcdef\fR +.EE + +.PP +Run a benchmark of generating work for 100 random sample blockhashes: +.EX +$ nano-pow \fB\--benchmark 100\fR .EE .SH EXAMPLES - SERVER diff --git a/package.json b/package.json index 90680aa..23620bb 100644 --- a/package.json +++ b/package.json @@ -45,6 +45,7 @@ "url": "git+https://zoso.dev/nano-pow.git" }, "scripts": { + "benchmark": "npm run build && ./dist/bin/nano-pow.sh --benchmark 1000", "build": "rm -rf {dist,types} && tsc && node esbuild.mjs && cp -p src/bin/nano-pow.sh dist/bin", "prepare": "npm run build", "start": "./dist/bin/nano-pow.sh --server", diff --git a/src/bin/cli.ts b/src/bin/cli.ts index cb1136e..a0b8699 100755 --- a/src/bin/cli.ts +++ b/src/bin/cli.ts @@ -2,14 +2,15 @@ //! SPDX-FileCopyrightText: 2025 Chris Duncan //! SPDX-License-Identifier: GPL-3.0-or-later -import * as readline from 'node:readline/promises' +import { getRandomValues } from 'node:crypto' +import { createInterface } from 'node:readline/promises' import type { WorkGenerateResponse, WorkValidateResponse } from '#types' process.title = 'NanoPow CLI' process.env.NANO_POW_DEBUG = '' process.env.NANO_POW_EFFORT = '' -process.env.NANO_POW_PORT = '5040' +process.env.NANO_POW_PORT = '5041' function log (...args: any[]): void { if (process.env.NANO_POW_DEBUG) console.log(new Date(Date.now()).toLocaleString(), 'NanoPow', args) @@ -19,7 +20,7 @@ const hashes: string[] = [] const stdinErrors: string[] = [] if (!process.stdin.isTTY) { - const stdin = readline.createInterface({ + const stdin = createInterface({ input: process.stdin }) let i = 0 @@ -44,11 +45,12 @@ If using --validate, results will also include validity properties. -h, --help show this dialog --debug enable additional logging output + --benchmark generate work for specified number of random hashes -b, --batch process all data before returning final results as array - -d, --difficulty= override the minimum difficulty value - -e, --effort= increase demand on GPU processing - -v, --validate= check an existing work value instead of searching for one + -d, --difficulty override the minimum difficulty value + -e, --effort increase demand on GPU processing + -v, --validate check an existing work value instead of searching for one If validating a nonce, it must be a 16-character hexadecimal value. Effort must be a decimal number between 1-32. @@ -67,12 +69,8 @@ while (/^[0-9A-Fa-f]{64}$/.test(args[args.length - 1] ?? '')) { } hashes.push(...inArgs) -if (hashes.length === 0) { - console.error('Invalid block hash input') - process.exit(1) -} - let isBatch = false +let isBenchmark = false const body: { [key: string]: any } = { action: 'work_generate' } @@ -101,6 +99,18 @@ for (let i = 0; i < args.length; i++) { process.env.NANO_POW_EFFORT = args[i + 1] break } + case ('--benchmark'): { + if (args[i + 1] == null) throw new Error('Missing argument for benchmark') + if (!(+args[i + 1] > 0)) throw new Error('Invalid benchmark count') + const random = new Uint8Array(32) + while (hashes.length < +args[i + 1]) { + getRandomValues(random) + const byteArray = [...random].map(byte => byte.toString(16).padStart(2, '0')) + hashes.push(byteArray.join('')) + } + isBenchmark = true + break + } case ('--debug'): { process.env.NANO_POW_DEBUG = 'true' break @@ -113,6 +123,11 @@ for (let i = 0; i < args.length; i++) { } } +if (hashes.length === 0) { + console.error('Invalid block hash input') + process.exit(1) +} + log('CLI args:', ...args) for (const stdinErr of stdinErrors) { log(stdinErr) @@ -124,8 +139,9 @@ await import('./server.js') // Execution must be sequential else GPU cannot map to CPU and will throw const results: (WorkGenerateResponse | WorkValidateResponse)[] = [] -const start = performance.now() const aborter = new AbortController() +if (isBenchmark) console.log('Starting benchmark...') +const start = performance.now() for (const hash of hashes) { try { body.hash = hash @@ -137,7 +153,7 @@ for (const hash of hashes) { }) clearTimeout(kill) const result = await response.json() - if (isBatch) { + if (isBatch || isBenchmark) { results.push(result) } else { console.log(result) @@ -147,6 +163,8 @@ for (const hash of hashes) { } } const end = performance.now() -if (isBatch) console.log(results) -log(end - start, 'ms total |', (end - start) / hashes.length, 'ms avg') +if (isBatch && !isBenchmark) console.log(results) +if (process.env.NANO_POW_DEBUG || isBenchmark) { + console.log(end - start, 'ms total |', (end - start) / hashes.length, 'ms avg') +} process.exit(0) diff --git a/src/bin/nano-pow.sh b/src/bin/nano-pow.sh index 4337ead..6229e98 100755 --- a/src/bin/nano-pow.sh +++ b/src/bin/nano-pow.sh @@ -10,7 +10,7 @@ NANO_POW_LOGS="$NANO_POW_HOME"/logs; mkdir -p "$NANO_POW_LOGS"; if [ "$1" = '--server' ]; then shift; - node "$SCRIPT_DIR"/server.js > "$NANO_POW_LOGS"/nano-pow-server-$(date +%s).log 2>&1 & echo "$!" > "$NANO_POW_HOME"/server.pid; + node "$SCRIPT_DIR"/server.js > "$NANO_POW_LOGS"/nano-pow-server-$(date -u -Iseconds).log 2>&1 & echo "$!" > "$NANO_POW_HOME"/server.pid; else node "$SCRIPT_DIR"/cli.js "$@"; fi; diff --git a/src/bin/server.ts b/src/bin/server.ts index 9124d7d..e4784be 100755 --- a/src/bin/server.ts +++ b/src/bin/server.ts @@ -48,6 +48,7 @@ async function respond (res: http.ServerResponse, data: Buffer[]): Promise if (action === 'work_validate' && !/^[0-9A-Fa-f]{16}$/.test(work ?? '')) { throw new Error('Invalid work. Must be a 16-character hex string.') } + response = `${action} failed` const options: NanoPowOptions = { debug: DEBUG, effort: EFFORT, @@ -156,7 +157,7 @@ page.on('console', msg => log(msg.text())) const path: string = new URL(import.meta.url).pathname const dir = path.slice(0, path.lastIndexOf('/')) const filename = join(dir, `${process.pid}.html`) -await writeFile(filename, '') +await writeFile(filename, '') await page.goto(import.meta.resolve(filename)) await page.waitForFunction(async (): Promise => { return await navigator['gpu'].requestAdapter() @@ -168,11 +169,9 @@ const enc = `sha256-${Buffer.from(hash).toString('base64')}` await page.setContent(` - - - - - + + + `) await unlink(filename) log('Puppeteer initialized') diff --git a/src/lib/gpu/compute.wgsl b/src/lib/gpu/compute.wgsl index 6a0ea13..9c46abe 100644 --- a/src/lib/gpu/compute.wgsl +++ b/src/lib/gpu/compute.wgsl @@ -81,41 +81,80 @@ const BLAKE2B_INIT = array, 16>( BLAKE2B_IV[7u] ); +fn G ( + a: ptr>, + b: ptr>, + c: ptr>, + d: ptr>, + m0: vec2, m1: vec2 +) { + *a += *b; + (*a).y += u32((*a).x < (*b).x); + *a += m0; + (*a).y += u32((*a).x < m0.x); + + *d = (*d ^ *a).yx; + + *c += *d; + (*c).y += u32((*c).x < (*d).x); + + *b ^= *c; + *b = (*b >> vec2(24u)) | (*b << vec2(8u)).yx; + + *a += *b; + (*a).y += u32((*a).x < (*b).x); + *a += m1; + (*a).y += u32((*a).x < m1.x); + + *d ^= *a; + *d = (*d >> vec2(16u)) | (*d << vec2(16u)).yx; + + *c += *d; + (*c).y += u32((*c).x < (*d).x); + + *b ^= *c; + *b = (*b >> vec2(31u)).yx | (*b << vec2(1u)); +} + /** * Used to fill partial `m` vec4 constructions. */ const Z = vec2(0u); -/** -* Used to apply boolean mask to swizzled result of carry bit comparison. -*/ -const CARRY = vec4(1u, 0u, 1u, 0u); - -/** -* Used to rotate bits by a fixed amount during G mixing. -*/ -const ROTATE_1 = vec4(1u); -const ROTATE_8 = vec4(8u); -const ROTATE_16 = vec4(16u); -const ROTATE_24 = vec4(24u); -const ROTATE_31 = vec4(31u); - /** * Shared flag to prevent execution for all workgroup threads based on the * atomicLoad() result of a single member thread. */ var found: bool; +/** +* Shared memory for seed and blockhash which do not change during execution and +* are eventually concatenated to form the hash input message. +*/ +var seed: vec2; +var m1: vec2; +var m2: vec2; +var m3: vec2; +var m4: vec2; + /** * Search compute function -* Calls main with a workgroup size of 64 which is generally considered optimal -* due to how warps and wavefronts are executed on modern GPUs. The entire -* workgroup exits immediately if a nonce was already found by a previous -* workgroup. +* Calls main with a workgroup size of 96 which in testing was the lowest value +* that would saturate the GPU active thread count and warp occupancy which +* provides a decent balance with the power-sensitive requirements of mobile +* devices. The entire workgroup exits immediately if a nonce was already found +* by a previous workgroup. */ -@compute @workgroup_size(32) +@compute @workgroup_size(96) fn search(@builtin(global_invocation_id) global_id: vec3, @builtin(local_invocation_id) local_id: vec3) { - found = (local_id.x == 0u && atomicLoad(&work.found) != 0u); + if (local_id.x == 0u) { + found = atomicLoad(&work.found) != 0u; + seed = ubo.seed; + m1 = ubo.blockhash[0u].xy; + m2 = ubo.blockhash[0u].zw; + m3 = ubo.blockhash[1u].xy; + m4 = ubo.blockhash[1u].zw; + } workgroupBarrier(); if (found) { return; } main(global_id, false); @@ -123,18 +162,23 @@ fn search(@builtin(global_invocation_id) global_id: vec3, @builtin(local_in /** * Validate compute function -* Calls main with a workgroup size of 1 so that only one value is tested +* Calls main with a workgroup size of 1 so that only one value is tested. */ @compute @workgroup_size(1) fn validate(@builtin(global_invocation_id) global_id: vec3) { + seed = ubo.seed; + m1 = ubo.blockhash[0u].xy; + m2 = ubo.blockhash[0u].zw; + m3 = ubo.blockhash[1u].xy; + m4 = ubo.blockhash[1u].zw; main(global_id, true); } /** * Main compute function -* A random u32 provided by the UBO is copied to form a pair. Each component of -* this 8-byte value is then XOR'd with a different dimensional index from -* the thread identifier. +* Each component of a random 8-byte value, provided by the UBO as a vec2, +* is XOR'd with a different dimensional index from the global thread identifier +* to create a unique nonce value for each thread. * * Where the reference implementation uses array lookups, the NanoPow * implementation assigns each array element to its own variable to enhance @@ -142,43 +186,40 @@ fn validate(@builtin(global_invocation_id) global_id: vec3) { */ fn main(id: vec3, validate: bool) { /** - * Initialize (nonce||blockhash) concatenation + * Initialize unique nonce */ - let m0: vec2 = ubo.seed ^ id.xy; - let m1: vec2 = ubo.blockhash[0u].xy; - let m2: vec2 = ubo.blockhash[0u].zw; - let m3: vec2 = ubo.blockhash[1u].xy; - let m4: vec2 = ubo.blockhash[1u].zw; + let m0: vec2 = seed ^ id.xy; /** * Compression buffer copied from the modified initialization vector. */ - var v01: vec4 = vec4(BLAKE2B_INIT[0u], BLAKE2B_INIT[1u]); - var v23: vec4 = vec4(BLAKE2B_INIT[2u], BLAKE2B_INIT[3u]); - var v45: vec4 = vec4(BLAKE2B_INIT[4u], BLAKE2B_INIT[5u]); - var v67: vec4 = vec4(BLAKE2B_INIT[6u], BLAKE2B_INIT[7u]); - var v89: vec4 = vec4(BLAKE2B_INIT[8u], BLAKE2B_INIT[9u]); - var vAB: vec4 = vec4(BLAKE2B_INIT[10u], BLAKE2B_INIT[11u]); - var vCD: vec4 = vec4(BLAKE2B_INIT[12u], BLAKE2B_INIT[13u]); - var vEF: vec4 = vec4(BLAKE2B_INIT[14u], BLAKE2B_INIT[15u]); - - /** - * Temporary variables used for subprocesses i=4 through i=7 - */ - var v56: vec4; - var vFC: vec4; - var v74: vec4; - var vDE: vec4; + var v0: vec2 = vec2(BLAKE2B_INIT[0u]); + var v1: vec2 = vec2(BLAKE2B_INIT[1u]); + var v2: vec2 = vec2(BLAKE2B_INIT[2u]); + var v3: vec2 = vec2(BLAKE2B_INIT[3u]); + var v4: vec2 = vec2(BLAKE2B_INIT[4u]); + var v5: vec2 = vec2(BLAKE2B_INIT[5u]); + var v6: vec2 = vec2(BLAKE2B_INIT[6u]); + var v7: vec2 = vec2(BLAKE2B_INIT[7u]); + var v8: vec2 = vec2(BLAKE2B_INIT[8u]); + var v9: vec2 = vec2(BLAKE2B_INIT[9u]); + var vA: vec2 = vec2(BLAKE2B_INIT[10u]); + var vB: vec2 = vec2(BLAKE2B_INIT[11u]); + var vC: vec2 = vec2(BLAKE2B_INIT[12u]); + var vD: vec2 = vec2(BLAKE2B_INIT[13u]); + var vE: vec2 = vec2(BLAKE2B_INIT[14u]); + var vF: vec2 = vec2(BLAKE2B_INIT[15u]); /** * Twelve rounds of G mixing as part of BLAKE2b compression step, each divided - * into eight subprocesses. Each statement of the first four subprocesses is - * executed in sequence so that the compiler can interleave independent - * instructions for improved scheduling. That is to say, the first statement - * `a = a + b` is executed for subprocesses 1-4, and then the next statement - * `a = a + m[sigma[r][2*i+0]]` is executed, and so on through all the steps of - * the mixing function. Once subprocesses 1-4 are done, computation on - * subprocesses 5-8 are executed in the same manner. + * into eight subprocesses, each of which is further paired to be processed in + * parallel by packing independent vec2 variables into vec4 variables. + * Each subprocess statement execution is alternated so that the compiler can + * interleave independent instructions for improved scheduling. That is to say, + * the first statement `a = a + b` is executed for each subprocess, and then + * the next statement `a = a + m[sigma[r][2*i+0]]` is executed, and so on + * through all the steps of the G mix function. Once subprocesses 1-4 are done, + * computation on subprocesses 5-8 are executed in the same manner. * * Each subprocess applies transformations to `m` and `v` variables based on a * defined set of index inputs. The algorithm for each subprocess is defined as @@ -202,10 +243,10 @@ fn main(id: vec3, validate: bool) { * b = rotr64(b ^ c, 63) * * Each sum step has an extra carry addition. Note that the m[sigma] sum is - * skipped if m[sigma] is zero since it effectively does nothing. + * skipped if m[sigma] is zero since it effectively does nothing. Also note + * that rotations must be applied differently from the reference implementation + * due to the lack of both a native rotate function and 64-bit support in WGSL. */ - var s0: vec4; - var s1: vec4; /**************************************************************************** * ROUND(0) * @@ -217,52 +258,10 @@ fn main(id: vec3, validate: bool) { * r=0, i=2, a=v[2], b=v[6], c=v[10], d=v[14], m[sigma]=4, m[sigma+1]=5 * r=0, i=3, a=v[3], b=v[7], c=v[11], d=v[15], m[sigma]=6, m[sigma+1]=7 */ - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - s0 = v01 + vec4(m0, m2); - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - v23 += vec4(m4, Z); - v23.y += u32(v23.x < m4.x); - - vCD = (vCD ^ v01).yxwz; - vEF = (vEF ^ v23).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_24) | (v45 << ROTATE_8).yxwz; - v67 ^= vAB; - v67 = (v67 >> ROTATE_24) | (v67 << ROTATE_8).yxwz; - - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - s0 = v01 + vec4(m1, m3); - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - // NOP - - vCD ^= v01; - vCD = (vCD >> ROTATE_16) | (vCD << ROTATE_16).yxwz; - vEF ^= v23; - vEF = (vEF >> ROTATE_16) | (vEF << ROTATE_16).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_31).yxwz | (v45 << ROTATE_1); - v67 ^= vAB; - v67 = (v67 >> ROTATE_31).yxwz | (v67 << ROTATE_1); + G(&v0, &v4, &v8, &vC, m0, m1); + G(&v1, &v5, &v9, &vD, m2, m3); + G(&v2, &v6, &vA, &vE, m4, Z); + G(&v3, &v7, &vB, &vF, Z, Z); /** * r=0, i=4, a=v[0], b=v[5], c=v[10], d=v[15], m[sigma]=8, m[sigma+1]=9 @@ -270,62 +269,10 @@ fn main(id: vec3, validate: bool) { * r=0, i=6, a=v[2], b=v[7], c=v[8], d=v[13], m[sigma]=12, m[sigma+1]=13 * r=0, i=7, a=v[3], b=v[4], c=v[9], d=v[14], m[sigma]=14, m[sigma+1]=15 */ - v56 = vec4(v45.zw, v67.xy); - v74 = vec4(v67.zw, v45.xy); - vFC = vec4(vEF.zw, vCD.xy); - vDE = vec4(vCD.zw, vEF.xy); - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - // NOP - - vFC = (vFC ^ v01).yxwz; - vDE = (vDE ^ v23).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_24) | (v56 << ROTATE_8).yxwz; - v74 ^= v89; - v74 = (v74 >> ROTATE_24) | (v74 << ROTATE_8).yxwz; - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - // NOP - - vFC ^= v01; - vFC = (vFC >> ROTATE_16) | (vFC << ROTATE_16).yxwz; - vDE ^= v23; - vDE = (vDE >> ROTATE_16) | (vDE << ROTATE_16).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_31).yxwz | (v56 << ROTATE_1); - v74 ^= v89; - v74 = (v74 >> ROTATE_31).yxwz | (v74 << ROTATE_1); - - v45 = vec4(v74.zw, v56.xy); - v67 = vec4(v56.zw, v74.xy); - vCD = vec4(vFC.zw, vDE.xy); - vEF = vec4(vDE.zw, vFC.xy); - - - + G(&v0, &v5, &vA, &vF, Z, Z); + G(&v1, &v6, &vB, &vC, Z, Z); + G(&v2, &v7, &v8, &vD, Z, Z); + G(&v3, &v4, &v9, &vE, Z, Z); /**************************************************************************** * ROUND(1) * @@ -337,50 +284,10 @@ fn main(id: vec3, validate: bool) { * r=1, i=2, a=v[2], b=v[6], c=v[10], d=v[14], m[sigma]=9, m[sigma+1]=15 * r=1, i=3, a=v[3], b=v[7], c=v[11], d=v[15], m[sigma]=13, m[sigma+1]=6 */ - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - v01 += vec4(Z, m4); - v01.w += u32(v01.z < m4.x); - // NOP - - vCD = (vCD ^ v01).yxwz; - vEF = (vEF ^ v23).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_24) | (v45 << ROTATE_8).yxwz; - v67 ^= vAB; - v67 = (v67 >> ROTATE_24) | (v67 << ROTATE_8).yxwz; - - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - // NOP - - vCD ^= v01; - vCD = (vCD >> ROTATE_16) | (vCD << ROTATE_16).yxwz; - vEF ^= v23; - vEF = (vEF >> ROTATE_16) | (vEF << ROTATE_16).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_31).yxwz | (v45 << ROTATE_1); - v67 ^= vAB; - v67 = (v67 >> ROTATE_31).yxwz | (v67 << ROTATE_1); + G(&v0, &v4, &v8, &vC, Z, Z); + G(&v1, &v5, &v9, &vD, m4, Z); + G(&v2, &v6, &vA, &vE, Z, Z); + G(&v3, &v7, &vB, &vF, Z, Z); /** * r=1, i=4, a=v[0], b=v[5], c=v[10], d=v[15], m[sigma]=1, m[sigma+1]=12 @@ -388,66 +295,10 @@ fn main(id: vec3, validate: bool) { * r=1, i=6, a=v[2], b=v[7], c=v[8], d=v[13], m[sigma]=11, m[sigma+1]=7 * r=1, i=7, a=v[3], b=v[4], c=v[9], d=v[14], m[sigma]=5, m[sigma+1]=3 */ - v56 = vec4(v45.zw, v67.xy); - v74 = vec4(v67.zw, v45.xy); - vFC = vec4(vEF.zw, vCD.xy); - vDE = vec4(vCD.zw, vEF.xy); - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - s0 = v01 + vec4(m1, m0); - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - // NOP - - vFC = (vFC ^ v01).yxwz; - vDE = (vDE ^ v23).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_24) | (v56 << ROTATE_8).yxwz; - v74 ^= v89; - v74 = (v74 >> ROTATE_24) | (v74 << ROTATE_8).yxwz; - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - v01 += vec4(Z, m2); - v01.w += u32(v01.z < m2.x); - v23 += vec4(Z, m3); - v23.w += u32(v23.z < m3.x); - - vFC ^= v01; - vFC = (vFC >> ROTATE_16) | (vFC << ROTATE_16).yxwz; - vDE ^= v23; - vDE = (vDE >> ROTATE_16) | (vDE << ROTATE_16).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_31).yxwz | (v56 << ROTATE_1); - v74 ^= v89; - v74 = (v74 >> ROTATE_31).yxwz | (v74 << ROTATE_1); - - v45 = vec4(v74.zw, v56.xy); - v67 = vec4(v56.zw, v74.xy); - vCD = vec4(vFC.zw, vDE.xy); - vEF = vec4(vDE.zw, vFC.xy); - - - - + G(&v0, &v5, &vA, &vF, m1, Z); + G(&v1, &v6, &vB, &vC, m0, m2); + G(&v2, &v7, &v8, &vD, Z, Z); + G(&v3, &v4, &v9, &vE, Z, m3); /**************************************************************************** * ROUND(2) * @@ -459,52 +310,10 @@ fn main(id: vec3, validate: bool) { * r=2, i=2, a=v[2], b=v[6], c=v[10], d=v[14], m[sigma]=5, m[sigma+1]=2 * r=2, i=3, a=v[3], b=v[7], c=v[11], d=v[15], m[sigma]=15, m[sigma+1]=13 */ - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - // NOP - - vCD = (vCD ^ v01).yxwz; - vEF = (vEF ^ v23).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_24) | (v45 << ROTATE_8).yxwz; - v67 ^= vAB; - v67 = (v67 >> ROTATE_24) | (v67 << ROTATE_8).yxwz; - - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - s0 = v01 + vec4(Z, m0); - v01= s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + vec4(m2, Z); - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - - vCD ^= v01; - vCD = (vCD >> ROTATE_16) | (vCD << ROTATE_16).yxwz; - vEF ^= v23; - vEF = (vEF >> ROTATE_16) | (vEF << ROTATE_16).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_31).yxwz | (v45 << ROTATE_1); - v67 ^= vAB; - v67 = (v67 >> ROTATE_31).yxwz | (v67 << ROTATE_1); + G(&v0, &v4, &v8, &vC, Z, Z); + G(&v1, &v5, &v9, &vD, Z, m0); + G(&v2, &v6, &vA, &vE, Z, m2); + G(&v3, &v7, &vB, &vF, Z, Z); /** * r=2, i=4, a=v[0], b=v[5], c=v[10], d=v[15], m[sigma]=10, m[sigma+1]=14 @@ -512,65 +321,10 @@ fn main(id: vec3, validate: bool) { * r=2, i=6, a=v[2], b=v[7], c=v[8], d=v[13], m[sigma]=7, m[sigma+1]=1 * r=2, i=7, a=v[3], b=v[4], c=v[9], d=v[14], m[sigma]=9, m[sigma+1]=4 */ - v56 = vec4(v45.zw, v67.xy); - v74 = vec4(v67.zw, v45.xy); - vFC = vec4(vEF.zw, vCD.xy); - vDE = vec4(vCD.zw, vEF.xy); - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // v01 += vec4(Z, m3) + vec4(Z, 0u, u32(v01.z + vec4(Z, m3).z < v01.z)); - v01.z += m3.x; - v01.w += m3.y + u32(v01.z < m3.x); - // NOP - - vFC = (vFC ^ v01).yxwz; - vDE = (vDE ^ v23).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_24) | (v56 << ROTATE_8).yxwz; - v74 ^= v89; - v74 = (v74 >> ROTATE_24) | (v74 << ROTATE_8).yxwz; - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - v23 += vec4(m1, m4) + vec4(0u, u32(v23.x + vec4(m1, m4).x < v23.x), 0u, u32(v23.z + vec4(m1, m4).z < v23.z)); - - vFC ^= v01; - vFC = (vFC >> ROTATE_16) | (vFC << ROTATE_16).yxwz; - vDE ^= v23; - vDE = (vDE >> ROTATE_16) | (vDE << ROTATE_16).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_31).yxwz | (v56 << ROTATE_1); - v74 ^= v89; - v74 = (v74 >> ROTATE_31).yxwz | (v74 << ROTATE_1); - - v45 = vec4(v74.zw, v56.xy); - v67 = vec4(v56.zw, v74.xy); - vCD = vec4(vFC.zw, vDE.xy); - vEF = vec4(vDE.zw, vFC.xy); - - - - + G(&v0, &v5, &vA, &vF, Z, Z); + G(&v1, &v6, &vB, &vC, m3, Z); + G(&v2, &v7, &v8, &vD, Z, m1); + G(&v3, &v4, &v9, &vE, Z, m4); /**************************************************************************** * ROUND(3) * @@ -582,49 +336,10 @@ fn main(id: vec3, validate: bool) { * r=3, i=2, a=v[2], b=v[6], c=v[10], d=v[14], m[sigma]=13, m[sigma+1]=12 * r=3, i=3, a=v[3], b=v[7], c=v[11], d=v[15], m[sigma]=11, m[sigma+1]=14 */ - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - v01 += vec4(Z, m3) + vec4(Z, 0u, u32(v01.z + vec4(Z, m3).z < v01.z)); - // NOP - - vCD = (vCD ^ v01).yxwz; - vEF = (vEF ^ v23).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_24) | (v45 << ROTATE_8).yxwz; - v67 ^= vAB; - v67 = (v67 >> ROTATE_24) | (v67 << ROTATE_8).yxwz; - - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - v01 += vec4(Z, m1) + vec4(Z, 0u, u32(v01.z + vec4(Z, m1).z < v01.z)); - // NOP - - vCD ^= v01; - vCD = (vCD >> ROTATE_16) | (vCD << ROTATE_16).yxwz; - vEF ^= v23; - vEF = (vEF >> ROTATE_16) | (vEF << ROTATE_16).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_31).yxwz | (v45 << ROTATE_1); - v67 ^= vAB; - v67 = (v67 >> ROTATE_31).yxwz | (v67 << ROTATE_1); + G(&v0, &v4, &v8, &vC, Z, Z); + G(&v1, &v5, &v9, &vD, m3, m1); + G(&v2, &v6, &vA, &vE, Z, Z); + G(&v3, &v7, &vB, &vF, Z, Z); /** * r=3, i=4, a=v[0], b=v[5], c=v[10], d=v[15], m[sigma]=2, m[sigma+1]=6 @@ -632,63 +347,10 @@ fn main(id: vec3, validate: bool) { * r=3, i=6, a=v[2], b=v[7], c=v[8], d=v[13], m[sigma]=4, m[sigma+1]=0 * r=3, i=7, a=v[3], b=v[4], c=v[9], d=v[14], m[sigma]=15, m[sigma+1]=8 */ - v56 = vec4(v45.zw, v67.xy); - v74 = vec4(v67.zw, v45.xy); - vFC = vec4(vEF.zw, vCD.xy); - vDE = vec4(vCD.zw, vEF.xy); - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - v01 += vec4(m2, Z) + vec4(0u, u32(v01.x + vec4(m2, Z).x < v01.x), Z); - v23 += vec4(m4, Z) + vec4(0u, u32(v23.x + vec4(m4, Z).x < v23.x), Z); - - vFC = (vFC ^ v01).yxwz; - vDE = (vDE ^ v23).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_24) | (v56 << ROTATE_8).yxwz; - v74 ^= v89; - v74 = (v74 >> ROTATE_24) | (v74 << ROTATE_8).yxwz; - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - v23 += vec4(m0, Z) + vec4(0u, u32(v23.x + vec4(m0, Z).x < v23.x), Z); - - vFC ^= v01; - vFC = (vFC >> ROTATE_16) | (vFC << ROTATE_16).yxwz; - vDE ^= v23; - vDE = (vDE >> ROTATE_16) | (vDE << ROTATE_16).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_31).yxwz | (v56 << ROTATE_1); - v74 ^= v89; - v74 = (v74 >> ROTATE_31).yxwz | (v74 << ROTATE_1); - - v45 = vec4(v74.zw, v56.xy); - v67 = vec4(v56.zw, v74.xy); - vCD = vec4(vFC.zw, vDE.xy); - vEF = vec4(vDE.zw, vFC.xy); - - - - + G(&v0, &v5, &vA, &vF, m2, Z); + G(&v1, &v6, &vB, &vC, Z, Z); + G(&v2, &v7, &v8, &vD, m4, m0); + G(&v3, &v4, &v9, &vE, Z, Z); /**************************************************************************** * ROUND(4) * @@ -700,50 +362,10 @@ fn main(id: vec3, validate: bool) { * r=4, i=2, a=v[2], b=v[6], c=v[10], d=v[14], m[sigma]=2, m[sigma+1]=4 * r=4, i=3, a=v[3], b=v[7], c=v[11], d=v[15], m[sigma]=10, m[sigma+1]=15 */ - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - s1 = v23 + vec4(m2, Z); - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - vCD = (vCD ^ v01).yxwz; - vEF = (vEF ^ v23).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_24) | (v45 << ROTATE_8).yxwz; - v67 ^= vAB; - v67 = (v67 >> ROTATE_24) | (v67 << ROTATE_8).yxwz; - - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - v01 += vec4(m0, Z) + vec4(0u, u32(v01.x + vec4(m0, Z).x < v01.x), Z); - v23 += vec4(m4, Z) + vec4(0u, u32(v23.x + vec4(m4, Z).x < v23.x), Z); - - vCD ^= v01; - vCD = (vCD >> ROTATE_16) | (vCD << ROTATE_16).yxwz; - vEF ^= v23; - vEF = (vEF >> ROTATE_16) | (vEF << ROTATE_16).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_31).yxwz | (v45 << ROTATE_1); - v67 ^= vAB; - v67 = (v67 >> ROTATE_31).yxwz | (v67 << ROTATE_1); + G(&v0, &v4, &v8, &vC, Z, m0); + G(&v1, &v5, &v9, &vD, Z, Z); + G(&v2, &v6, &vA, &vE, m2, m4); + G(&v3, &v7, &vB, &vF, Z, Z); /** */ @@ -754,64 +376,10 @@ fn main(id: vec3, validate: bool) { * r=4, i=6, a=v[2], b=v[7], c=v[8], d=v[13], m[sigma]=6, m[sigma+1]=8 * r=4, i=7, a=v[3], b=v[4], c=v[9], d=v[14], m[sigma]=3, m[sigma+1]=13 */ - v56 = vec4(v45.zw, v67.xy); - v74 = vec4(v67.zw, v45.xy); - vFC = vec4(vEF.zw, vCD.xy); - vDE = vec4(vCD.zw, vEF.xy); - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - s1 = v23 + vec4(Z, m3); - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - vFC = (vFC ^ v01).yxwz; - vDE = (vDE ^ v23).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_24) | (v56 << ROTATE_8).yxwz; - v74 ^= v89; - v74 = (v74 >> ROTATE_24) | (v74 << ROTATE_8).yxwz; - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - v01 += vec4(m1, Z) + vec4(0u, u32(v01.x + vec4(m1, Z).x < v01.x), Z); - // NOP - - vFC ^= v01; - vFC = (vFC >> ROTATE_16) | (vFC << ROTATE_16).yxwz; - vDE ^= v23; - vDE = (vDE >> ROTATE_16) | (vDE << ROTATE_16).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_31).yxwz | (v56 << ROTATE_1); - v74 ^= v89; - v74 = (v74 >> ROTATE_31).yxwz | (v74 << ROTATE_1); - - v45 = vec4(v74.zw, v56.xy); - v67 = vec4(v56.zw, v74.xy); - vCD = vec4(vFC.zw, vDE.xy); - vEF = vec4(vDE.zw, vFC.xy); - - - - + G(&v0, &v5, &vA, &vF, Z, m1); + G(&v1, &v6, &vB, &vC, Z, Z); + G(&v2, &v7, &v8, &vD, Z, Z); + G(&v3, &v4, &v9, &vE, m3, Z); /**************************************************************************** * ROUND(5) * @@ -823,50 +391,10 @@ fn main(id: vec3, validate: bool) { * r=5, i=2, a=v[2], b=v[6], c=v[10], d=v[14], m[sigma]=0, m[sigma+1]=11 * r=5, i=3, a=v[3], b=v[7], c=v[11], d=v[15], m[sigma]=8, m[sigma+1]=3 */ - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - v01 += vec4(m2, Z) + vec4(0u, u32(v01.x + vec4(m2, Z).x < v01.x), Z); - v23 += vec4(m0, Z) + vec4(0u, u32(v23.x + vec4(m0, Z).x < v23.x), Z); - - vCD = (vCD ^ v01).yxwz; - vEF = (vEF ^ v23).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_24) | (v45 << ROTATE_8).yxwz; - v67 ^= vAB; - v67 = (v67 >> ROTATE_24) | (v67 << ROTATE_8).yxwz; - - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - s1 = v23 + vec4(Z, m3); - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - vCD ^= v01; - vCD = (vCD >> ROTATE_16) | (vCD << ROTATE_16).yxwz; - vEF ^= v23; - vEF = (vEF >> ROTATE_16) | (vEF << ROTATE_16).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_31).yxwz | (v45 << ROTATE_1); - v67 ^= vAB; - v67 = (v67 >> ROTATE_31).yxwz | (v67 << ROTATE_1); + G(&v0, &v4, &v8, &vC, m2, Z); + G(&v1, &v5, &v9, &vD, Z, Z); + G(&v2, &v6, &vA, &vE, m0, Z); + G(&v3, &v7, &vB, &vF, Z, m3); /** * r=5, i=4, a=v[0], b=v[5], c=v[10], d=v[15], m[sigma]=4, m[sigma+1]=13 @@ -874,63 +402,10 @@ fn main(id: vec3, validate: bool) { * r=5, i=6, a=v[2], b=v[7], c=v[8], d=v[13], m[sigma]=15, m[sigma+1]=14 * r=5, i=7, a=v[3], b=v[4], c=v[9], d=v[14], m[sigma]=1, m[sigma+1]=9 */ - v56 = vec4(v45.zw, v67.xy); - v74 = vec4(v67.zw, v45.xy); - vFC = vec4(vEF.zw, vCD.xy); - vDE = vec4(vCD.zw, vEF.xy); - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - v01 += vec4(m4, Z) + vec4(0u, u32(v01.x + vec4(m4, Z).x < v01.x), Z); - v23 += vec4(Z, m1) + vec4(Z, 0u, u32(v23.z + vec4(Z, m1).z < v23.z)); - - vFC = (vFC ^ v01).yxwz; - vDE = (vDE ^ v23).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_24) | (v56 << ROTATE_8).yxwz; - v74 ^= v89; - v74 = (v74 >> ROTATE_24) | (v74 << ROTATE_8).yxwz; - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - // NOP - - vFC ^= v01; - vFC = (vFC >> ROTATE_16) | (vFC << ROTATE_16).yxwz; - vDE ^= v23; - vDE = (vDE >> ROTATE_16) | (vDE << ROTATE_16).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_31).yxwz | (v56 << ROTATE_1); - v74 ^= v89; - v74 = (v74 >> ROTATE_31).yxwz | (v74 << ROTATE_1); - - v45 = vec4(v74.zw, v56.xy); - v67 = vec4(v56.zw, v74.xy); - vCD = vec4(vFC.zw, vDE.xy); - vEF = vec4(vDE.zw, vFC.xy); - - - - + G(&v0, &v5, &vA, &vF, m4, Z); + G(&v1, &v6, &vB, &vC, Z, Z); + G(&v2, &v7, &v8, &vD, Z, Z); + G(&v3, &v4, &v9, &vE, m1, Z); /**************************************************************************** * ROUND(6) * @@ -942,49 +417,10 @@ fn main(id: vec3, validate: bool) { * r=6, i=2, a=v[2], b=v[6], c=v[10], d=v[14], m[sigma]=14, m[sigma+1]=13 * r=6, i=3, a=v[3], b=v[7], c=v[11], d=v[15], m[sigma]=4, m[sigma+1]=10 */ - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - v01 += vec4(Z, m1) + vec4(Z, 0u, u32(v01.z + vec4(Z, m1).z < v01.z)); - v23 += vec4(Z, m4) + vec4(Z, 0u, u32(v23.z + vec4(Z, m4).z < v23.z)); - - vCD = (vCD ^ v01).yxwz; - vEF = (vEF ^ v23).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_24) | (v45 << ROTATE_8).yxwz; - v67 ^= vAB; - v67 = (v67 >> ROTATE_24) | (v67 << ROTATE_8).yxwz; - - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - // NOP - - vCD ^= v01; - vCD = (vCD >> ROTATE_16) | (vCD << ROTATE_16).yxwz; - vEF ^= v23; - vEF = (vEF >> ROTATE_16) | (vEF << ROTATE_16).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_31).yxwz | (v45 << ROTATE_1); - v67 ^= vAB; - v67 = (v67 >> ROTATE_31).yxwz | (v67 << ROTATE_1); + G(&v0, &v4, &v8, &vC, Z, Z); + G(&v1, &v5, &v9, &vD, m1, Z); + G(&v2, &v6, &vA, &vE, Z, Z); + G(&v3, &v7, &vB, &vF, m4, Z); /** * r=6, i=4, a=v[0], b=v[5], c=v[10], d=v[15], m[sigma]=0, m[sigma+1]=7 @@ -992,64 +428,10 @@ fn main(id: vec3, validate: bool) { * r=6, i=6, a=v[2], b=v[7], c=v[8], d=v[13], m[sigma]=9, m[sigma+1]=2 * r=6, i=7, a=v[3], b=v[4], c=v[9], d=v[14], m[sigma]=8, m[sigma+1]=11 */ - v56 = vec4(v45.zw, v67.xy); - v74 = vec4(v67.zw, v45.xy); - vFC = vec4(vEF.zw, vCD.xy); - vDE = vec4(vCD.zw, vEF.xy); - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - v01 += vec4(m0, Z) + vec4(0u, u32(v01.x + vec4(m0, Z).x < v01.x), Z); - // NOP - - vFC = (vFC ^ v01).yxwz; - vDE = (vDE ^ v23).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_24) | (v56 << ROTATE_8).yxwz; - v74 ^= v89; - v74 = (v74 >> ROTATE_24) | (v74 << ROTATE_8).yxwz; - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - v01 += vec4(Z, m3) + vec4(Z, 0u, u32(v01.z + vec4(Z, m3).z < v01.z)); - s1 = v23 + vec4(m2, Z); - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - vFC ^= v01; - vFC = (vFC >> ROTATE_16) | (vFC << ROTATE_16).yxwz; - vDE ^= v23; - vDE = (vDE >> ROTATE_16) | (vDE << ROTATE_16).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_31).yxwz | (v56 << ROTATE_1); - v74 ^= v89; - v74 = (v74 >> ROTATE_31).yxwz | (v74 << ROTATE_1); - - v45 = vec4(v74.zw, v56.xy); - v67 = vec4(v56.zw, v74.xy); - vCD = vec4(vFC.zw, vDE.xy); - vEF = vec4(vDE.zw, vFC.xy); - - - - + G(&v0, &v5, &vA, &vF, m0, Z); + G(&v1, &v6, &vB, &vC, Z, m3); + G(&v2, &v7, &v8, &vD, Z, m2); + G(&v3, &v4, &v9, &vE, Z, Z); /**************************************************************************** * ROUND(7) * @@ -1061,50 +443,10 @@ fn main(id: vec3, validate: bool) { * r=7, i=2, a=v[2], b=v[6], c=v[10], d=v[14], m[sigma]=12, m[sigma+1]=1 * r=7, i=3, a=v[3], b=v[7], c=v[11], d=v[15], m[sigma]=3, m[sigma+1]=9 */ - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - s1 = v23 + vec4(Z, m3); - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - vCD = (vCD ^ v01).yxwz; - vEF = (vEF ^ v23).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_24) | (v45 << ROTATE_8).yxwz; - v67 ^= vAB; - v67 = (v67 >> ROTATE_24) | (v67 << ROTATE_8).yxwz; - - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - v23 += vec4(m1, Z) + vec4(0u, u32(v23.x + vec4(m1, Z).x < v23.x), Z); - - vCD ^= v01; - vCD = (vCD >> ROTATE_16) | (vCD << ROTATE_16).yxwz; - vEF ^= v23; - vEF = (vEF >> ROTATE_16) | (vEF << ROTATE_16).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_31).yxwz | (v45 << ROTATE_1); - v67 ^= vAB; - v67 = (v67 >> ROTATE_31).yxwz | (v67 << ROTATE_1); + G(&v0, &v4, &v8, &vC, Z, Z); + G(&v1, &v5, &v9, &vD, Z, Z); + G(&v2, &v6, &vA, &vE, Z, m1); + G(&v3, &v7, &vB, &vF, m3, Z); /** * r=7, i=4, a=v[0], b=v[5], c=v[10], d=v[15], m[sigma]=5, m[sigma+1]=0 @@ -1112,63 +454,10 @@ fn main(id: vec3, validate: bool) { * r=7, i=6, a=v[2], b=v[7], c=v[8], d=v[13], m[sigma]=8, m[sigma+1]=6 * r=7, i=7, a=v[3], b=v[4], c=v[9], d=v[14], m[sigma]=2, m[sigma+1]=10 */ - v56 = vec4(v45.zw, v67.xy); - v74 = vec4(v67.zw, v45.xy); - vFC = vec4(vEF.zw, vCD.xy); - vDE = vec4(vCD.zw, vEF.xy); - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - v23 += vec4(Z, m2) + vec4(Z, 0u, u32(v23.z + vec4(Z, m2).z < v23.z)); - - vFC = (vFC ^ v01).yxwz; - vDE = (vDE ^ v23).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_24) | (v56 << ROTATE_8).yxwz; - v74 ^= v89; - v74 = (v74 >> ROTATE_24) | (v74 << ROTATE_8).yxwz; - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - v01 += vec4(m0, m4) + vec4(0u, u32(v01.x + vec4(m0, m4).x < v01.x), 0u, u32(v01.z + vec4(m0, m4).z < v01.z)); - // NOP - - vFC ^= v01; - vFC = (vFC >> ROTATE_16) | (vFC << ROTATE_16).yxwz; - vDE ^= v23; - vDE = (vDE >> ROTATE_16) | (vDE << ROTATE_16).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_31).yxwz | (v56 << ROTATE_1); - v74 ^= v89; - v74 = (v74 >> ROTATE_31).yxwz | (v74 << ROTATE_1); - - v45 = vec4(v74.zw, v56.xy); - v67 = vec4(v56.zw, v74.xy); - vCD = vec4(vFC.zw, vDE.xy); - vEF = vec4(vDE.zw, vFC.xy); - - - - + G(&v0, &v5, &vA, &vF, Z, m0); + G(&v1, &v6, &vB, &vC, Z, m4); + G(&v2, &v7, &v8, &vD, Z, Z); + G(&v3, &v4, &v9, &vE, m2, Z); /**************************************************************************** * ROUND(8) * @@ -1180,50 +469,10 @@ fn main(id: vec3, validate: bool) { * r=8, i=2, a=v[2], b=v[6], c=v[10], d=v[14], m[sigma]=11, m[sigma+1]=3 * r=8, i=3, a=v[3], b=v[7], c=v[11], d=v[15], m[sigma]=0, m[sigma+1]=8 */ - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - v23 += vec4(Z, m0); - v23.w += u32(v23.z < m0.x); - - vCD = (vCD ^ v01).yxwz; - vEF = (vEF ^ v23).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_24) | (v45 << ROTATE_8).yxwz; - v67 ^= vAB; - v67 = (v67 >> ROTATE_24) | (v67 << ROTATE_8).yxwz; - - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - v23 += vec4(m3, Z) + vec4(0u, u32(v23.x + vec4(m3, Z).x < v23.x), Z); - - vCD ^= v01; - vCD = (vCD >> ROTATE_16) | (vCD << ROTATE_16).yxwz; - vEF ^= v23; - vEF = (vEF >> ROTATE_16) | (vEF << ROTATE_16).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_31).yxwz | (v45 << ROTATE_1); - v67 ^= vAB; - v67 = (v67 >> ROTATE_31).yxwz | (v67 << ROTATE_1); + G(&v0, &v4, &v8, &vC, Z, Z); + G(&v1, &v5, &v9, &vD, Z, Z); + G(&v2, &v6, &vA, &vE, Z, m3); + G(&v3, &v7, &vB, &vF, m0, Z); /** * r=8, i=4, a=v[0], b=v[5], c=v[10], d=v[15], m[sigma]=12, m[sigma+1]=2 @@ -1231,63 +480,10 @@ fn main(id: vec3, validate: bool) { * r=8, i=6, a=v[2], b=v[7], c=v[8], d=v[13], m[sigma]=1, m[sigma+1]=4 * r=8, i=7, a=v[3], b=v[4], c=v[9], d=v[14], m[sigma]=10, m[sigma+1]=5 */ - v56 = vec4(v45.zw, v67.xy); - v74 = vec4(v67.zw, v45.xy); - vFC = vec4(vEF.zw, vCD.xy); - vDE = vec4(vCD.zw, vEF.xy); - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - v23 += vec4(m1, Z) + vec4(0u, u32(v23.x + vec4(m1, Z).x < v23.x), Z); - - vFC = (vFC ^ v01).yxwz; - vDE = (vDE ^ v23).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_24) | (v56 << ROTATE_8).yxwz; - v74 ^= v89; - v74 = (v74 >> ROTATE_24) | (v74 << ROTATE_8).yxwz; - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - v01 += vec4(m2, Z) + vec4(0u, u32(v01.x + vec4(m2, Z).x < v01.x), Z); - v23 += vec4(m4, Z) + vec4(0u, u32(v23.x + vec4(m4, Z).x < v23.x), Z); - - vFC ^= v01; - vFC = (vFC >> ROTATE_16) | (vFC << ROTATE_16).yxwz; - vDE ^= v23; - vDE = (vDE >> ROTATE_16) | (vDE << ROTATE_16).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_31).yxwz | (v56 << ROTATE_1); - v74 ^= v89; - v74 = (v74 >> ROTATE_31).yxwz | (v74 << ROTATE_1); - - v45 = vec4(v74.zw, v56.xy); - v67 = vec4(v56.zw, v74.xy); - vCD = vec4(vFC.zw, vDE.xy); - vEF = vec4(vDE.zw, vFC.xy); - - - - + G(&v0, &v5, &vA, &vF, Z, m2); + G(&v1, &v6, &vB, &vC, Z, Z); + G(&v2, &v7, &v8, &vD, m1, m4); + G(&v3, &v4, &v9, &vE, Z, Z); /**************************************************************************** * ROUND(9) * @@ -1299,51 +495,10 @@ fn main(id: vec3, validate: bool) { * r=9, i=2, a=v[2], b=v[6], c=v[10], d=v[14], m[sigma]=7, m[sigma+1]=6 * r=9, i=3, a=v[3], b=v[7], c=v[11], d=v[15], m[sigma]=1, m[sigma+1]=5 */ - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - v23 += vec4(Z, m1); - v23.w += u32(v23.z < m1.x); - - vCD = (vCD ^ v01).yxwz; - vEF = (vEF ^ v23).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_24) | (v45 << ROTATE_8).yxwz; - v67 ^= vAB; - v67 = (v67 >> ROTATE_24) | (v67 << ROTATE_8).yxwz; - - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - s0 = v01 + vec4(m2, m4); - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - // NOP - - vCD ^= v01; - vCD = (vCD >> ROTATE_16) | (vCD << ROTATE_16).yxwz; - vEF ^= v23; - vEF = (vEF >> ROTATE_16) | (vEF << ROTATE_16).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_31).yxwz | (v45 << ROTATE_1); - v67 ^= vAB; - v67 = (v67 >> ROTATE_31).yxwz | (v67 << ROTATE_1); + G(&v0, &v4, &v8, &vC, Z, m2); + G(&v1, &v5, &v9, &vD, Z, m4); + G(&v2, &v6, &vA, &vE, Z, Z); + G(&v3, &v7, &vB, &vF, m1, Z); /** * r=9, i=4, a=v[0], b=v[5], c=v[10], d=v[15], m[sigma]=15, m[sigma+1]=11 @@ -1351,65 +506,10 @@ fn main(id: vec3, validate: bool) { * r=9, i=6, a=v[2], b=v[7], c=v[8], d=v[13], m[sigma]=3, m[sigma+1]=12 * r=9, i=7, a=v[3], b=v[4], c=v[9], d=v[14], m[sigma]=13, m[sigma+1]=0 */ - v56 = vec4(v45.zw, v67.xy); - v74 = vec4(v67.zw, v45.xy); - vFC = vec4(vEF.zw, vCD.xy); - vDE = vec4(vCD.zw, vEF.xy); - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - v23 += vec4(m3, Z); - v23.y += u32(v23.x < m3.x); - - vFC = (vFC ^ v01).yxwz; - vDE = (vDE ^ v23).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_24) | (v56 << ROTATE_8).yxwz; - v74 ^= v89; - v74 = (v74 >> ROTATE_24) | (v74 << ROTATE_8).yxwz; - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - v23 += vec4(Z, m0); - v23.w += u32(v23.z < m0.x); - - vFC ^= v01; - vFC = (vFC >> ROTATE_16) | (vFC << ROTATE_16).yxwz; - vDE ^= v23; - vDE = (vDE >> ROTATE_16) | (vDE << ROTATE_16).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_31).yxwz | (v56 << ROTATE_1); - v74 ^= v89; - v74 = (v74 >> ROTATE_31).yxwz | (v74 << ROTATE_1); - - v45 = vec4(v74.zw, v56.xy); - v67 = vec4(v56.zw, v74.xy); - vCD = vec4(vFC.zw, vDE.xy); - vEF = vec4(vDE.zw, vFC.xy); - - - - + G(&v0, &v5, &vA, &vF, Z, Z); + G(&v1, &v6, &vB, &vC, Z, Z); + G(&v2, &v7, &v8, &vD, m3, Z); + G(&v3, &v4, &v9, &vE, Z, m0); /**************************************************************************** * ROUND(10) * @@ -1421,52 +521,10 @@ fn main(id: vec3, validate: bool) { * r=10, i=2, a=v[2], b=v[6], c=v[10], d=v[14], m[sigma]=4, m[sigma+1]=5 * r=10, i=3, a=v[3], b=v[7], c=v[11], d=v[15], m[sigma]=6, m[sigma+1]=7 */ - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - s0 = v01 + vec4(m0, m2); - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - v23 += vec4(m4, Z); - v23.y += u32(v23.x < m4.x); - - vCD = (vCD ^ v01).yxwz; - vEF = (vEF ^ v23).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_24) | (v45 << ROTATE_8).yxwz; - v67 ^= vAB; - v67 = (v67 >> ROTATE_24) | (v67 << ROTATE_8).yxwz; - - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - s0 = v01 + vec4(m1, m3); - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - // NOP - - vCD ^= v01; - vCD = (vCD >> ROTATE_16) | (vCD << ROTATE_16).yxwz; - vEF ^= v23; - vEF = (vEF >> ROTATE_16) | (vEF << ROTATE_16).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_31).yxwz | (v45 << ROTATE_1); - v67 ^= vAB; - v67 = (v67 >> ROTATE_31).yxwz | (v67 << ROTATE_1); + G(&v0, &v4, &v8, &vC, m0, m1); + G(&v1, &v5, &v9, &vD, m2, m3); + G(&v2, &v6, &vA, &vE, m4, Z); + G(&v3, &v7, &vB, &vF, Z, Z); /** * r=10, i=4, a=v[0], b=v[5], c=v[10], d=v[15], m[sigma]=8, m[sigma+1]=9 @@ -1474,63 +532,10 @@ fn main(id: vec3, validate: bool) { * r=10, i=6, a=v[2], b=v[7], c=v[8], d=v[13], m[sigma]=12, m[sigma+1]=13 * r=10, i=7, a=v[3], b=v[4], c=v[9], d=v[14], m[sigma]=14, m[sigma+1]=15 */ - v56 = vec4(v45.zw, v67.xy); - v74 = vec4(v67.zw, v45.xy); - vFC = vec4(vEF.zw, vCD.xy); - vDE = vec4(vCD.zw, vEF.xy); - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - // NOP - - vFC = (vFC ^ v01).yxwz; - vDE = (vDE ^ v23).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_24) | (v56 << ROTATE_8).yxwz; - v74 ^= v89; - v74 = (v74 >> ROTATE_24) | (v74 << ROTATE_8).yxwz; - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - // NOP - - vFC ^= v01; - vFC = (vFC >> ROTATE_16) | (vFC << ROTATE_16).yxwz; - vDE ^= v23; - vDE = (vDE >> ROTATE_16) | (vDE << ROTATE_16).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_31).yxwz | (v56 << ROTATE_1); - v74 ^= v89; - v74 = (v74 >> ROTATE_31).yxwz | (v74 << ROTATE_1); - - v45 = vec4(v74.zw, v56.xy); - v67 = vec4(v56.zw, v74.xy); - vCD = vec4(vFC.zw, vDE.xy); - vEF = vec4(vDE.zw, vFC.xy); - - - - + G(&v0, &v5, &vA, &vF, Z, Z); + G(&v1, &v6, &vB, &vC, Z, Z); + G(&v2, &v7, &v8, &vD, Z, Z); + G(&v3, &v4, &v9, &vE, Z, Z); /**************************************************************************** * ROUND(11) * @@ -1542,50 +547,10 @@ fn main(id: vec3, validate: bool) { * r=11, i=2, a=v[2], b=v[6], c=v[10], d=v[14], m[sigma]=9, m[sigma+1]=15 * r=11, i=3, a=v[3], b=v[7], c=v[11], d=v[15], m[sigma]=13, m[sigma+1]=6 */ - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - v01 += vec4(Z, m4); - v01.w += u32(v01.z < m4.x); - // NOP - - vCD = (vCD ^ v01).yxwz; - vEF = (vEF ^ v23).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_24) | (v45 << ROTATE_8).yxwz; - v67 ^= vAB; - v67 = (v67 >> ROTATE_24) | (v67 << ROTATE_8).yxwz; - - s0 = v01 + v45; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v67; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // NOP - // NOP - - vCD ^= v01; - vCD = (vCD >> ROTATE_16) | (vCD << ROTATE_16).yxwz; - vEF ^= v23; - vEF = (vEF >> ROTATE_16) | (vEF << ROTATE_16).yxwz; - - s0 = v89 + vCD; - v89 = s0 + (vec4(s0 < v89) & CARRY).yxwz; - s1 = vAB + vEF; - vAB = s1 + (vec4(s1 < vAB) & CARRY).yxwz; - - v45 ^= v89; - v45 = (v45 >> ROTATE_31).yxwz | (v45 << ROTATE_1); - v67 ^= vAB; - v67 = (v67 >> ROTATE_31).yxwz | (v67 << ROTATE_1); + G(&v0, &v4, &v8, &vC, Z, Z); + G(&v1, &v5, &v9, &vD, m4, Z); + G(&v2, &v6, &vA, &vE, Z, Z); + G(&v3, &v7, &vB, &vF, Z, Z); /** * r=11, i=4, a=v[0], b=v[5], c=v[10], d=v[15], m[sigma]=1, m[sigma+1]=12 @@ -1593,66 +558,10 @@ fn main(id: vec3, validate: bool) { * r=11, i=6, a=v[2], b=v[7], c=v[8], d=v[13], m[sigma]=11, m[sigma+1]=7 * r=11, i=7, a=v[3], b=v[4], c=v[9], d=v[14], m[sigma]=5, m[sigma+1]=3 */ - v56 = vec4(v45.zw, v67.xy); - v74 = vec4(v67.zw, v45.xy); - vFC = vec4(vEF.zw, vCD.xy); - vDE = vec4(vCD.zw, vEF.xy); - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - s0 = v01 + vec4(m1, m0); - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - // NOP - - vFC = (vFC ^ v01).yxwz; - vDE = (vDE ^ v23).yxwz; - - s0 = vAB + vFC; - vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - v56 ^= vAB; - v56 = (v56 >> ROTATE_24) | (v56 << ROTATE_8).yxwz; - v74 ^= v89; - v74 = (v74 >> ROTATE_24) | (v74 << ROTATE_8).yxwz; - - s0 = v01 + v56; - v01 = s0 + (vec4(s0 < v01) & CARRY).yxwz; - s1 = v23 + v74; - v23 = s1 + (vec4(s1 < v23) & CARRY).yxwz; - - // v01 += vec4(Z, m2); - // v01.w += u32(v01.z < m2.x); - v23 += vec4(Z, m3); - v23.w += u32(v23.z < m3.x); - - // vFC ^= v01; - // vFC = (vFC >> ROTATE_16) | (vFC << ROTATE_16).yxwz; - vDE ^= v23; - vDE = (vDE >> ROTATE_16) | (vDE << ROTATE_16).yxwz; - - // s0 = vAB + vFC; - // vAB = s0 + (vec4(s0 < vAB) & CARRY).yxwz; - s1 = v89 + vDE; - v89 = s1 + (vec4(s1 < v89) & CARRY).yxwz; - - // v56 ^= vAB; - // v74 ^= v89; - // v56 = (v56 << ROTATE_1) | (v56 >> ROTATE_31).yxwz; - // v74 = (v74 << ROTATE_1) | (v74 >> ROTATE_31).yxwz; - - // v45 = vec4(v74.zw, v56.xy); - // v67 = vec4(v56.zw, v74.xy); - // vCD = vec4(vFC.zw, vDE.xy); - // vEF = vec4(vDE.zw, vFC.xy); - - - - + G(&v0, &v5, &vA, &vF, m1, Z); + // G(&v1, &v6, &vB, &vC, m0, m2); + G(&v2, &v7, &v8, &vD, Z, Z); + // G(&v3, &v4, &v9, &vE, Z, m3); /**************************************************************************** * NONCE CHECK * @@ -1661,7 +570,7 @@ fn main(id: vec3, validate: bool) { /** * Set nonce if it passes the difficulty threshold and no other thread has set it. */ - var result = BLAKE2B_INIT[0u] ^ v01.xy ^ v89.xy; + var result: vec2 = BLAKE2B_INIT[0u] ^ v0 ^ v8; if (validate || ((result.y > ubo.difficulty.y || (result.y == ubo.difficulty.y && result.y >= ubo.difficulty.y)) && atomicLoad(&work.found) == 0u)) { atomicStore(&work.found, 1u); work.nonce = m0; diff --git a/test/index.html b/test/index.html index 19d93a6..397ef8b 100644 --- a/test/index.html +++ b/test/index.html @@ -6,7 +6,7 @@ SPDX-License-Identifier: GPL-3.0-or-later - +