From 7ee0cbb57e9f30304efb2f0c050762b6011fe70e Mon Sep 17 00:00:00 2001 From: Chris Duncan Date: Fri, 24 Jan 2025 23:30:23 -0800 Subject: [PATCH] Check found flag only once per workgroup to avoid numerous competing atomic loads. --- src/shaders/compute.wgsl | 14 +++++++++++--- 1 file changed, 11 insertions(+), 3 deletions(-) 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 work: WORK; +/** +* Shared flag to prevent execution for all workgroup threads based on the +* atomicLoad() result of a single member thread. +*/ +var 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) { - if (atomicLoad(&work.found) != 0u) { return; } +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); + workgroupBarrier(); + if (found) { return; } main(global_id); } @@ -1434,7 +1442,7 @@ fn main(id: vec3) { 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; -- 2.34.1