From: Chris Duncan <chris@zoso.dev>
Date: Sat, 25 Jan 2025 07:30:23 +0000 (-0800)
Subject: Check found flag only once per workgroup to avoid numerous competing atomic loads.
X-Git-Tag: v3.0.0~69
X-Git-Url: https://zoso.dev/?a=commitdiff_plain;h=7ee0cbb57e9f30304efb2f0c050762b6011fe70e;p=nano-pow.git

Check found flag only once per workgroup to avoid numerous competing atomic loads.
---

diff --git a/src/shaders/compute.wgsl b/src/shaders/compute.wgsl
index 5c2e222..23db45f 100644
--- a/src/shaders/compute.wgsl
+++ b/src/shaders/compute.wgsl
@@ -14,6 +14,12 @@ struct WORK {
 };
 @group(0) @binding(1) var<storage, read_write> work: WORK;
 
+/**
+* Shared flag to prevent execution for all workgroup threads based on the
+* atomicLoad() result of a single member thread.
+*/
+var<workgroup> found: bool;
+
 /**
 * Defined separately from `v0` because the original value is required to
 * calculate the digest and compare it to the threshold.
@@ -34,8 +40,10 @@ const ROTATE_31 = vec2(31u, 31u);
 * Calls main with a workgroup size of 64 which has been tested as optimal
 */
 @compute @workgroup_size(64)
-fn search(@builtin(global_invocation_id) global_id: vec3<u32>) {
-	if (atomicLoad(&work.found) != 0u) { return; }
+fn search(@builtin(global_invocation_id) global_id: vec3<u32>, @builtin(local_invocation_id) local_id: vec3<u32>) {
+	found = (local_id.x == 0u && atomicLoad(&work.found) != 0u);
+	workgroupBarrier();
+	if (found) { return; }
 	main(global_id);
 }
 
@@ -1434,7 +1442,7 @@ fn main(id: vec3<u32>) {
 	if ((BLAKE2B_IV32_0.y ^ v0.y ^ v8.y) > threshold) {
 		let wasFound: u32 = atomicExchange(&work.found, 1u);
 		if (wasFound == 0u) {
-		work.nonce = m0;
+			work.nonce = m0;
 		}
 	}
 	return;