diff --git a/tests/shader/gen/message_passing.dxil b/tests/shader/gen/message_passing.dxil
index dc2eced..2be73da 100644
--- a/tests/shader/gen/message_passing.dxil
+++ b/tests/shader/gen/message_passing.dxil
Binary files differ
diff --git a/tests/shader/gen/message_passing.hlsl b/tests/shader/gen/message_passing.hlsl
index 3a48808..ba8ce5f 100644
--- a/tests/shader/gen/message_passing.hlsl
+++ b/tests/shader/gen/message_passing.hlsl
@@ -6,7 +6,7 @@
 
 static const uint3 gl_WorkGroupSize = uint3(256u, 1u, 1u);
 
-globallycoherent RWByteAddressBuffer data_buf : register(u0);
+RWByteAddressBuffer data_buf : register(u0);
 RWByteAddressBuffer control_buf : register(u1);
 
 static uint3 gl_GlobalInvocationID;
@@ -22,12 +22,13 @@
 
 void comp_main()
 {
-    data_buf.Store(gl_GlobalInvocationID.x * 8 + 0, 1u);
+    uint _76;
+    data_buf.InterlockedExchange(gl_GlobalInvocationID.x * 8 + 0, 1u, _76);
     DeviceMemoryBarrier();
     uint param = gl_GlobalInvocationID.x;
     uint write_flag_ix = permute_flag_ix(param);
-    uint _76;
-    data_buf.InterlockedExchange(write_flag_ix * 8 + 4, 1u, _76);
+    uint _77;
+    data_buf.InterlockedExchange(write_flag_ix * 8 + 4, 1u, _77);
     uint read_ix = (gl_GlobalInvocationID.x * 4099u) & 65535u;
     uint param_1 = read_ix;
     uint read_flag_ix = permute_flag_ix(param_1);
@@ -35,7 +36,9 @@
     data_buf.InterlockedAdd(read_flag_ix * 8 + 4, 0, _58);
     uint flag = _58;
     DeviceMemoryBarrier();
-    uint data = data_buf.Load(read_ix * 8 + 0);
+    uint _62;
+    data_buf.InterlockedAdd(read_ix * 8 + 0, 0, _62);
+    uint data = _62;
     if (flag > data)
     {
         uint _73;
diff --git a/tests/shader/gen/message_passing.msl b/tests/shader/gen/message_passing.msl
index 1bda181..e48f48a 100644
--- a/tests/shader/gen/message_passing.msl
+++ b/tests/shader/gen/message_passing.msl
@@ -31,20 +31,21 @@
     return (data_ix * 419u) & 65535u;
 }
 
-kernel void main0(volatile device DataBuf& data_buf [[buffer(0)]], device ControlBuf& control_buf [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+kernel void main0(device DataBuf& data_buf [[buffer(0)]], device ControlBuf& control_buf [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
-    data_buf.data[gl_GlobalInvocationID.x].data = 1u;
+    atomic_store_explicit((device atomic_uint*)&data_buf.data[gl_GlobalInvocationID.x].data, 1u, memory_order_relaxed);
     threadgroup_barrier(mem_flags::mem_device);
     uint param = gl_GlobalInvocationID.x;
     uint write_flag_ix = permute_flag_ix(param);
-    atomic_store_explicit((volatile device atomic_uint*)&data_buf.data[write_flag_ix].flag, 1u, memory_order_relaxed);
+    atomic_store_explicit((device atomic_uint*)&data_buf.data[write_flag_ix].flag, 1u, memory_order_relaxed);
     uint read_ix = (gl_GlobalInvocationID.x * 4099u) & 65535u;
     uint param_1 = read_ix;
     uint read_flag_ix = permute_flag_ix(param_1);
-    uint _58 = atomic_load_explicit((volatile device atomic_uint*)&data_buf.data[read_flag_ix].flag, memory_order_relaxed);
+    uint _58 = atomic_load_explicit((device atomic_uint*)&data_buf.data[read_flag_ix].flag, memory_order_relaxed);
     uint flag = _58;
     threadgroup_barrier(mem_flags::mem_device);
-    uint data = data_buf.data[read_ix].data;
+    uint _62 = atomic_load_explicit((device atomic_uint*)&data_buf.data[read_ix].data, memory_order_relaxed);
+    uint data = _62;
     if (flag > data)
     {
         uint _73 = atomic_fetch_add_explicit((device atomic_uint*)&control_buf.failures, 1u, memory_order_relaxed);
diff --git a/tests/shader/gen/message_passing.spv b/tests/shader/gen/message_passing.spv
index fa7fb3e..e5f56d6 100644
--- a/tests/shader/gen/message_passing.spv
+++ b/tests/shader/gen/message_passing.spv
Binary files differ
diff --git a/tests/shader/gen/message_passing_vkmm.spv b/tests/shader/gen/message_passing_vkmm.spv
index 58d3521..8527c2b 100644
--- a/tests/shader/gen/message_passing_vkmm.spv
+++ b/tests/shader/gen/message_passing_vkmm.spv
Binary files differ
diff --git a/tests/shader/gen/prefix.dxil b/tests/shader/gen/prefix.dxil
index 12d0d50..34f3d6a 100644
--- a/tests/shader/gen/prefix.dxil
+++ b/tests/shader/gen/prefix.dxil
Binary files differ
diff --git a/tests/shader/gen/prefix.hlsl b/tests/shader/gen/prefix.hlsl
index ba76110..3af5a96 100644
--- a/tests/shader/gen/prefix.hlsl
+++ b/tests/shader/gen/prefix.hlsl
@@ -12,11 +12,11 @@
 
 static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
 
-static const Monoid _183 = { 0u };
+static const Monoid _185 = { 0u };
 
 globallycoherent RWByteAddressBuffer _43 : register(u2);
 ByteAddressBuffer _67 : register(t0);
-RWByteAddressBuffer _367 : register(u1);
+RWByteAddressBuffer _372 : register(u1);
 
 static uint3 gl_LocalInvocationID;
 struct SPIRV_Cross_Input
@@ -76,12 +76,10 @@
     }
     if (gl_LocalInvocationID.x == 511u)
     {
-        uint _378;
-        _43.InterlockedExchange(part_ix * 12 + 8, agg.element, _378);
+        _43.Store(part_ix * 12 + 8, agg.element);
         if (part_ix == 0u)
         {
-            uint _379;
-            _43.InterlockedExchange(12, agg.element, _379);
+            _43.Store(12, agg.element);
         }
     }
     DeviceMemoryBarrier();
@@ -94,11 +92,12 @@
         }
         _43.Store(part_ix * 12 + 4, flag);
     }
-    Monoid exclusive = _183;
+    Monoid exclusive = _185;
     if (part_ix != 0u)
     {
         uint look_back_ix = part_ix - 1u;
         uint their_ix = 0u;
+        Monoid their_prefix;
         Monoid their_agg;
         Monoid m;
         while (true)
@@ -114,10 +113,9 @@
             {
                 if (gl_LocalInvocationID.x == 511u)
                 {
-                    uint _221;
-                    _43.InterlockedAdd(look_back_ix * 12 + 12, 0, _221);
-                    Monoid _222 = { _221 };
-                    Monoid their_prefix = _222;
+                    Monoid _223;
+                    _223.element = _43.Load(look_back_ix * 12 + 12);
+                    their_prefix.element = _223.element;
                     Monoid param_4 = their_prefix;
                     Monoid param_5 = exclusive;
                     exclusive = combine_monoid(param_4, param_5);
@@ -130,9 +128,9 @@
                 {
                     if (gl_LocalInvocationID.x == 511u)
                     {
-                        uint _242;
-                        _43.InterlockedAdd(look_back_ix * 12 + 8, 0, _242);
-                        their_agg.element = _242;
+                        Monoid _245;
+                        _245.element = _43.Load(look_back_ix * 12 + 8);
+                        their_agg.element = _245.element;
                         Monoid param_6 = their_agg;
                         Monoid param_7 = exclusive;
                         exclusive = combine_monoid(param_6, param_7);
@@ -144,9 +142,9 @@
             }
             if (gl_LocalInvocationID.x == 511u)
             {
-                Monoid _263;
-                _263.element = _67.Load(((look_back_ix * 8192u) + their_ix) * 4 + 0);
-                m.element = _263.element;
+                Monoid _267;
+                _267.element = _67.Load(((look_back_ix * 8192u) + their_ix) * 4 + 0);
+                m.element = _267.element;
                 if (their_ix == 0u)
                 {
                     their_agg = m;
@@ -187,8 +185,7 @@
             Monoid param_13 = agg;
             Monoid inclusive_prefix = combine_monoid(param_12, param_13);
             sh_prefix = exclusive;
-            uint _380;
-            _43.InterlockedExchange(part_ix * 12 + 12, inclusive_prefix.element, _380);
+            _43.Store(part_ix * 12 + 12, inclusive_prefix.element);
         }
         DeviceMemoryBarrier();
         if (gl_LocalInvocationID.x == 511u)
@@ -214,7 +211,7 @@
         Monoid param_16 = row;
         Monoid param_17 = local[i_2];
         Monoid m_1 = combine_monoid(param_16, param_17);
-        _367.Store((ix + i_2) * 4 + 0, m_1.element);
+        _372.Store((ix + i_2) * 4 + 0, m_1.element);
     }
 }
 
diff --git a/tests/shader/gen/prefix.msl b/tests/shader/gen/prefix.msl
index aaa7158..8e402a9 100644
--- a/tests/shader/gen/prefix.msl
+++ b/tests/shader/gen/prefix.msl
@@ -87,7 +87,7 @@
     return Monoid{ a.element + b.element };
 }
 
-kernel void main0(const device InBuf& _67 [[buffer(0)]], device OutBuf& _367 [[buffer(1)]], volatile device StateBuf& _43 [[buffer(2)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
+kernel void main0(const device InBuf& _67 [[buffer(0)]], device OutBuf& _372 [[buffer(1)]], volatile device StateBuf& _43 [[buffer(2)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
 {
     threadgroup uint sh_part_ix;
     threadgroup Monoid sh_scratch[512];
@@ -127,10 +127,10 @@
     }
     if (gl_LocalInvocationID.x == 511u)
     {
-        atomic_store_explicit((volatile device atomic_uint*)&_43.state[part_ix].aggregate.element, agg.element, memory_order_relaxed);
+        _43.state[part_ix].aggregate.element = agg.element;
         if (part_ix == 0u)
         {
-            atomic_store_explicit((volatile device atomic_uint*)&_43.state[0].prefix.element, agg.element, memory_order_relaxed);
+            _43.state[0].prefix.element = agg.element;
         }
     }
     threadgroup_barrier(mem_flags::mem_device);
@@ -148,6 +148,7 @@
     {
         uint look_back_ix = part_ix - 1u;
         uint their_ix = 0u;
+        Monoid their_prefix;
         Monoid their_agg;
         Monoid m;
         while (true)
@@ -163,8 +164,7 @@
             {
                 if (gl_LocalInvocationID.x == 511u)
                 {
-                    uint _221 = atomic_load_explicit((volatile device atomic_uint*)&_43.state[look_back_ix].prefix.element, memory_order_relaxed);
-                    Monoid their_prefix = Monoid{ _221 };
+                    their_prefix.element = _43.state[look_back_ix].prefix.element;
                     Monoid param_4 = their_prefix;
                     Monoid param_5 = exclusive;
                     exclusive = combine_monoid(param_4, param_5);
@@ -177,8 +177,7 @@
                 {
                     if (gl_LocalInvocationID.x == 511u)
                     {
-                        uint _242 = atomic_load_explicit((volatile device atomic_uint*)&_43.state[look_back_ix].aggregate.element, memory_order_relaxed);
-                        their_agg.element = _242;
+                        their_agg.element = _43.state[look_back_ix].aggregate.element;
                         Monoid param_6 = their_agg;
                         Monoid param_7 = exclusive;
                         exclusive = combine_monoid(param_6, param_7);
@@ -231,7 +230,7 @@
             Monoid param_13 = agg;
             Monoid inclusive_prefix = combine_monoid(param_12, param_13);
             sh_prefix = exclusive;
-            atomic_store_explicit((volatile device atomic_uint*)&_43.state[part_ix].prefix.element, inclusive_prefix.element, memory_order_relaxed);
+            _43.state[part_ix].prefix.element = inclusive_prefix.element;
         }
         threadgroup_barrier(mem_flags::mem_device);
         if (gl_LocalInvocationID.x == 511u)
@@ -257,7 +256,7 @@
         Monoid param_16 = row;
         Monoid param_17 = local[i_2];
         Monoid m_1 = combine_monoid(param_16, param_17);
-        _367.outbuf[ix + i_2].element = m_1.element;
+        _372.outbuf[ix + i_2].element = m_1.element;
     }
 }
 
diff --git a/tests/shader/gen/prefix.spv b/tests/shader/gen/prefix.spv
index 0da9ed8..d2c1aad 100644
--- a/tests/shader/gen/prefix.spv
+++ b/tests/shader/gen/prefix.spv
Binary files differ
diff --git a/tests/shader/gen/prefix_atomic.dxil b/tests/shader/gen/prefix_atomic.dxil
index 80c05a3..68f47e5 100644
--- a/tests/shader/gen/prefix_atomic.dxil
+++ b/tests/shader/gen/prefix_atomic.dxil
Binary files differ
diff --git a/tests/shader/gen/prefix_atomic.hlsl b/tests/shader/gen/prefix_atomic.hlsl
index 6105c01..10f7081 100644
--- a/tests/shader/gen/prefix_atomic.hlsl
+++ b/tests/shader/gen/prefix_atomic.hlsl
@@ -12,11 +12,11 @@
 
 static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
 
-static const Monoid _183 = { 0u };
+static const Monoid _185 = { 0u };
 
 globallycoherent RWByteAddressBuffer _43 : register(u2);
 ByteAddressBuffer _67 : register(t0);
-RWByteAddressBuffer _367 : register(u1);
+RWByteAddressBuffer _372 : register(u1);
 
 static uint3 gl_LocalInvocationID;
 struct SPIRV_Cross_Input
@@ -76,12 +76,10 @@
     }
     if (gl_LocalInvocationID.x == 511u)
     {
-        uint _378;
-        _43.InterlockedExchange(part_ix * 12 + 8, agg.element, _378);
+        _43.Store(part_ix * 12 + 8, agg.element);
         if (part_ix == 0u)
         {
-            uint _379;
-            _43.InterlockedExchange(12, agg.element, _379);
+            _43.Store(12, agg.element);
         }
     }
     DeviceMemoryBarrier();
@@ -92,23 +90,24 @@
         {
             flag = 2u;
         }
-        uint _380;
-        _43.InterlockedExchange(part_ix * 12 + 4, flag, _380);
+        uint _383;
+        _43.InterlockedExchange(part_ix * 12 + 4, flag, _383);
     }
-    Monoid exclusive = _183;
+    Monoid exclusive = _185;
     if (part_ix != 0u)
     {
         uint look_back_ix = part_ix - 1u;
         uint their_ix = 0u;
+        Monoid their_prefix;
         Monoid their_agg;
         Monoid m;
         while (true)
         {
             if (gl_LocalInvocationID.x == 511u)
             {
-                uint _206;
-                _43.InterlockedAdd(look_back_ix * 12 + 4, 0, _206);
-                sh_flag = _206;
+                uint _208;
+                _43.InterlockedAdd(look_back_ix * 12 + 4, 0, _208);
+                sh_flag = _208;
             }
             GroupMemoryBarrierWithGroupSync();
             DeviceMemoryBarrier();
@@ -117,10 +116,9 @@
             {
                 if (gl_LocalInvocationID.x == 511u)
                 {
-                    uint _221;
-                    _43.InterlockedAdd(look_back_ix * 12 + 12, 0, _221);
-                    Monoid _222 = { _221 };
-                    Monoid their_prefix = _222;
+                    Monoid _223;
+                    _223.element = _43.Load(look_back_ix * 12 + 12);
+                    their_prefix.element = _223.element;
                     Monoid param_4 = their_prefix;
                     Monoid param_5 = exclusive;
                     exclusive = combine_monoid(param_4, param_5);
@@ -133,9 +131,9 @@
                 {
                     if (gl_LocalInvocationID.x == 511u)
                     {
-                        uint _242;
-                        _43.InterlockedAdd(look_back_ix * 12 + 8, 0, _242);
-                        their_agg.element = _242;
+                        Monoid _245;
+                        _245.element = _43.Load(look_back_ix * 12 + 8);
+                        their_agg.element = _245.element;
                         Monoid param_6 = their_agg;
                         Monoid param_7 = exclusive;
                         exclusive = combine_monoid(param_6, param_7);
@@ -147,9 +145,9 @@
             }
             if (gl_LocalInvocationID.x == 511u)
             {
-                Monoid _263;
-                _263.element = _67.Load(((look_back_ix * 8192u) + their_ix) * 4 + 0);
-                m.element = _263.element;
+                Monoid _267;
+                _267.element = _67.Load(((look_back_ix * 8192u) + their_ix) * 4 + 0);
+                m.element = _267.element;
                 if (their_ix == 0u)
                 {
                     their_agg = m;
@@ -190,14 +188,13 @@
             Monoid param_13 = agg;
             Monoid inclusive_prefix = combine_monoid(param_12, param_13);
             sh_prefix = exclusive;
-            uint _381;
-            _43.InterlockedExchange(part_ix * 12 + 12, inclusive_prefix.element, _381);
+            _43.Store(part_ix * 12 + 12, inclusive_prefix.element);
         }
         DeviceMemoryBarrier();
         if (gl_LocalInvocationID.x == 511u)
         {
-            uint _382;
-            _43.InterlockedExchange(part_ix * 12 + 4, 2u, _382);
+            uint _384;
+            _43.InterlockedExchange(part_ix * 12 + 4, 2u, _384);
         }
     }
     GroupMemoryBarrierWithGroupSync();
@@ -218,7 +215,7 @@
         Monoid param_16 = row;
         Monoid param_17 = local[i_2];
         Monoid m_1 = combine_monoid(param_16, param_17);
-        _367.Store((ix + i_2) * 4 + 0, m_1.element);
+        _372.Store((ix + i_2) * 4 + 0, m_1.element);
     }
 }
 
diff --git a/tests/shader/gen/prefix_atomic.msl b/tests/shader/gen/prefix_atomic.msl
index d5de44a..6d7d155 100644
--- a/tests/shader/gen/prefix_atomic.msl
+++ b/tests/shader/gen/prefix_atomic.msl
@@ -87,7 +87,7 @@
     return Monoid{ a.element + b.element };
 }
 
-kernel void main0(const device InBuf& _67 [[buffer(0)]], device OutBuf& _367 [[buffer(1)]], volatile device StateBuf& _43 [[buffer(2)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
+kernel void main0(const device InBuf& _67 [[buffer(0)]], device OutBuf& _372 [[buffer(1)]], volatile device StateBuf& _43 [[buffer(2)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
 {
     threadgroup uint sh_part_ix;
     threadgroup Monoid sh_scratch[512];
@@ -127,10 +127,10 @@
     }
     if (gl_LocalInvocationID.x == 511u)
     {
-        atomic_store_explicit((volatile device atomic_uint*)&_43.state[part_ix].aggregate.element, agg.element, memory_order_relaxed);
+        _43.state[part_ix].aggregate.element = agg.element;
         if (part_ix == 0u)
         {
-            atomic_store_explicit((volatile device atomic_uint*)&_43.state[0].prefix.element, agg.element, memory_order_relaxed);
+            _43.state[0].prefix.element = agg.element;
         }
     }
     threadgroup_barrier(mem_flags::mem_device);
@@ -148,14 +148,15 @@
     {
         uint look_back_ix = part_ix - 1u;
         uint their_ix = 0u;
+        Monoid their_prefix;
         Monoid their_agg;
         Monoid m;
         while (true)
         {
             if (gl_LocalInvocationID.x == 511u)
             {
-                uint _206 = atomic_load_explicit((volatile device atomic_uint*)&_43.state[look_back_ix].flag, memory_order_relaxed);
-                sh_flag = _206;
+                uint _208 = atomic_load_explicit((volatile device atomic_uint*)&_43.state[look_back_ix].flag, memory_order_relaxed);
+                sh_flag = _208;
             }
             threadgroup_barrier(mem_flags::mem_threadgroup);
             threadgroup_barrier(mem_flags::mem_device);
@@ -164,8 +165,7 @@
             {
                 if (gl_LocalInvocationID.x == 511u)
                 {
-                    uint _221 = atomic_load_explicit((volatile device atomic_uint*)&_43.state[look_back_ix].prefix.element, memory_order_relaxed);
-                    Monoid their_prefix = Monoid{ _221 };
+                    their_prefix.element = _43.state[look_back_ix].prefix.element;
                     Monoid param_4 = their_prefix;
                     Monoid param_5 = exclusive;
                     exclusive = combine_monoid(param_4, param_5);
@@ -178,8 +178,7 @@
                 {
                     if (gl_LocalInvocationID.x == 511u)
                     {
-                        uint _242 = atomic_load_explicit((volatile device atomic_uint*)&_43.state[look_back_ix].aggregate.element, memory_order_relaxed);
-                        their_agg.element = _242;
+                        their_agg.element = _43.state[look_back_ix].aggregate.element;
                         Monoid param_6 = their_agg;
                         Monoid param_7 = exclusive;
                         exclusive = combine_monoid(param_6, param_7);
@@ -232,7 +231,7 @@
             Monoid param_13 = agg;
             Monoid inclusive_prefix = combine_monoid(param_12, param_13);
             sh_prefix = exclusive;
-            atomic_store_explicit((volatile device atomic_uint*)&_43.state[part_ix].prefix.element, inclusive_prefix.element, memory_order_relaxed);
+            _43.state[part_ix].prefix.element = inclusive_prefix.element;
         }
         threadgroup_barrier(mem_flags::mem_device);
         if (gl_LocalInvocationID.x == 511u)
@@ -258,7 +257,7 @@
         Monoid param_16 = row;
         Monoid param_17 = local[i_2];
         Monoid m_1 = combine_monoid(param_16, param_17);
-        _367.outbuf[ix + i_2].element = m_1.element;
+        _372.outbuf[ix + i_2].element = m_1.element;
     }
 }
 
diff --git a/tests/shader/gen/prefix_atomic.spv b/tests/shader/gen/prefix_atomic.spv
index 1b31f67..acca545 100644
--- a/tests/shader/gen/prefix_atomic.spv
+++ b/tests/shader/gen/prefix_atomic.spv
Binary files differ
diff --git a/tests/shader/gen/prefix_vkmm.spv b/tests/shader/gen/prefix_vkmm.spv
index c8b9ea2..0b8d475 100644
--- a/tests/shader/gen/prefix_vkmm.spv
+++ b/tests/shader/gen/prefix_vkmm.spv
Binary files differ
diff --git a/tests/shader/message_passing.comp b/tests/shader/message_passing.comp
index 521dedd..e5e53b6 100644
--- a/tests/shader/message_passing.comp
+++ b/tests/shader/message_passing.comp
@@ -23,7 +23,7 @@
     uint flag;
 };
 
-layout(binding = 0) coherent buffer DataBuf
+layout(binding = 0) buffer DataBuf
 {
     Element data[];
 } data_buf;
@@ -40,7 +40,7 @@
 
 void main()
 {
-    data_buf.data[gl_GlobalInvocationID.x].data = 1u;
+    atomicStore(data_buf.data[gl_GlobalInvocationID.x].data, 1u, gl_ScopeDevice, 0, 0);
 #ifndef VKMM
     memoryBarrierBuffer();
 #endif
@@ -52,7 +52,7 @@
 #ifndef VKMM
     memoryBarrierBuffer();
 #endif
-    uint data = data_buf.data[read_ix].data;
+    uint data = atomicLoad(data_buf.data[read_ix].data, gl_ScopeDevice, 0, 0);
     if (flag > data)
     {
         atomicAdd(control_buf.failures, 1u);
diff --git a/tests/shader/prefix.comp b/tests/shader/prefix.comp
index 0fc9909..a6a0d57 100644
--- a/tests/shader/prefix.comp
+++ b/tests/shader/prefix.comp
@@ -99,9 +99,9 @@
 
     // Publish aggregate for this partition
     if (gl_LocalInvocationID.x == WG_SIZE - 1) {
-        atomicStore(state[part_ix].aggregate.element, agg.element, gl_ScopeDevice, 0, 0);
+        state[part_ix].aggregate = agg;
         if (part_ix == 0) {
-            atomicStore(state[0].prefix.element, agg.element, gl_ScopeDevice, 0, 0);
+            state[0].prefix = agg;
         }
     }
     // Write flag with release semantics; this is done portably with a barrier.
@@ -147,13 +147,13 @@
 
             if (flag == FLAG_PREFIX_READY) {
                 if (gl_LocalInvocationID.x == WG_SIZE - 1) {
-                    Monoid their_prefix = Monoid(atomicLoad(state[look_back_ix].prefix.element, gl_ScopeDevice, 0, 0));
+                    Monoid their_prefix = state[look_back_ix].prefix;
                     exclusive = combine_monoid(their_prefix, exclusive);
                 }
                 break;
             } else if (flag == FLAG_AGGREGATE_READY) {
                 if (gl_LocalInvocationID.x == WG_SIZE - 1) {
-                    their_agg.element = atomicLoad(state[look_back_ix].aggregate.element, gl_ScopeDevice, 0, 0);
+                    their_agg = state[look_back_ix].aggregate;
                     exclusive = combine_monoid(their_agg, exclusive);
                 }
                 look_back_ix--;
@@ -193,7 +193,7 @@
         if (gl_LocalInvocationID.x == WG_SIZE - 1) {
             Monoid inclusive_prefix = combine_monoid(exclusive, agg);
             sh_prefix = exclusive;
-            atomicStore(state[part_ix].prefix.element, inclusive_prefix.element, gl_ScopeDevice, 0, 0);
+            state[part_ix].prefix = inclusive_prefix;
         }
 #ifndef VKMM
         memoryBarrierBuffer();