From fc633353b0d403a375517cf4f14225afa6377960 Mon Sep 17 00:00:00 2001
From: Chris Duncan <chris@zoso.dev>
Date: Wed, 1 Jan 2025 02:36:53 -0800
Subject: [PATCH] Prepare to write out entire work value. Revert to global
 invocation id so larger workloads can be supported later.

---
 src/lib/workers/powgpu.ts | 25 +++++++++++--------------
 1 file changed, 11 insertions(+), 14 deletions(-)

diff --git a/src/lib/workers/powgpu.ts b/src/lib/workers/powgpu.ts
index 7897f26..d334032 100644
--- a/src/lib/workers/powgpu.ts
+++ b/src/lib/workers/powgpu.ts
@@ -49,7 +49,7 @@ export class PowGpu extends WorkerInterface {
 			threshold: u32
 		};
 		@group(0) @binding(0) var<uniform> ubo: UBO;
-		@group(0) @binding(1) var<storage, read_write> work: u32;
+		@group(0) @binding(1) var<storage, read_write> work: vec2<u32>;
 
 		/**
 		* Defined separately from uint v[32] below as the original value is required
@@ -138,17 +138,13 @@ export class PowGpu extends WorkerInterface {
 		* Main compute function
 		*/
 		@compute @workgroup_size(256)
-		fn main(
-			@builtin(global_invocation_id) global_id: vec3<u32>,
-			@builtin(workgroup_id) workgroup_id: vec3<u32>,
-			@builtin(local_invocation_id) local_id: vec3<u32>
-		) {
+		fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
 			var m: array<u32, 16>;
 
 			// 8-byte work is split into two 4-byte u32
-			// First 3 bytes provided, last byte defined by this compute index
-			m[0u] = (ubo.rand.x << 8u) ^ workgroup_id.x;
-			m[1u] = (ubo.rand.y << 8u) ^ local_id.x;
+			// First 6 bytes provided, last 2 bytes defined by this compute index
+			m[0u] = ubo.rand.x;
+			m[1u] = (ubo.rand.y << 16u) ^ global_id.x;
 
 			// Block hash
 			m[2u] = ubo.blockhash[0u].x;
@@ -195,8 +191,9 @@ export class PowGpu extends WorkerInterface {
 			}
 
 			// Store the result directly into work array
-			if ((BLAKE2B_IV32_1 ^ v[1u] ^ v[17u]) > 0u) {
-				work = global_id.x;
+			if ((BLAKE2B_IV32_1 ^ v[1u] ^ v[17u]) > 0xffffffffu) {
+				work.x = ubo.rand.x;
+				work.y = (ubo.rand.y << 16u) ^ global_id.x;
 			}
 		}
 	`;
@@ -295,6 +292,7 @@ export class PowGpu extends WorkerInterface {
 			uboView.setUint32(i / 2, parseInt(uint32, 16))
 		}
 		const rand = crypto.getRandomValues(new Uint32Array(2))
+		console.log(rand)
 		uboView.setUint32(32, rand[0], true)
 		uboView.setUint32(36, rand[1], true)
 		uboView.setUint32(40, threshold, true)
@@ -352,10 +350,9 @@ export class PowGpu extends WorkerInterface {
 		const result = new Uint32Array(PowGpu.#cpuBuffer.getMappedRange()).slice()
 		PowGpu.#cpuBuffer.unmap()
 
-		console.log(`result`)
-		console.dir(result)
+		console.log(`result: ${[...result]}`)
 		if (result[0] !== 0 || result[1] !== 0) {
-			const hex = PowGpu.#hexify([...result])
+			const hex = PowGpu.#hexify([result[0], result[1]])
 			typeof callback === 'function' && callback(hex)
 			return
 		}
-- 
2.34.1