diff --git a/shader/coarse.wgsl b/shader/coarse.wgsl
index df09de9..0963c29 100644
--- a/shader/coarse.wgsl
+++ b/shader/coarse.wgsl
@@ -148,15 +148,14 @@ fn main(
     // Exit early if prior stages failed, as we can't run this stage.
     // We need to check only prior stages, as if this stage has failed in another workgroup, 
     // we still want to know this workgroup's memory requirement.   
+#ifdef have_uniform
     if local_id.x == 0u {
         // Reuse sh_part_count to hold failed flag, shmem is tight
         sh_part_count[0] = atomicLoad(&bump.failed);
     }
-#ifdef have_uniform
     let failed = workgroupUniformLoad(&sh_part_count[0]);
 #else
-    workgroupBarrier();
-    let failed = sh_part_count[0];
+    let failed = atomicLoad(&bump.failed);
 #endif
     if (failed & (STAGE_BINNING | STAGE_TILE_ALLOC | STAGE_PATH_COARSE)) != 0u {
         return;
diff --git a/shader/tile_alloc.wgsl b/shader/tile_alloc.wgsl
index 8d39e7c..8ec9217 100644
--- a/shader/tile_alloc.wgsl
+++ b/shader/tile_alloc.wgsl
@@ -29,7 +29,9 @@ let WG_SIZE = 256u;
 
 var<workgroup> sh_tile_count: array<u32, WG_SIZE>;
 var<workgroup> sh_tile_offset: u32;
+#ifdef have_uniform
 var<workgroup> sh_atomic_failed: u32;
+#endif
 
 @compute @workgroup_size(256)
 fn main(
@@ -39,14 +41,13 @@ fn main(
     // Exit early if prior stages failed, as we can't run this stage.
     // We need to check only prior stages, as if this stage has failed in another workgroup, 
     // we still want to know this workgroup's memory requirement.
+#ifdef have_uniform
     if local_id.x == 0u {
         sh_atomic_failed = atomicLoad(&bump.failed);
     }
-#ifdef have_uniform
     let failed = workgroupUniformLoad(&sh_atomic_failed);
 #else
-    workgroupBarrier();
-    let failed = sh_atomic_failed;
+    let failed = atomicLoad(&bump.failed);
 #endif
     if (failed & STAGE_BINNING) != 0u {
         return;