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 5c2e2229b1d36bee358afcff38858fd59b5322ce..23db45f42fc8034ac874258115122276b6ed6183 100644 (file)
--- 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;