diff --git a/tests/shader/gen/stack_leaf.dxil b/tests/shader/gen/stack_leaf.dxil
index 75423cd..14658c3 100644
--- a/tests/shader/gen/stack_leaf.dxil
+++ b/tests/shader/gen/stack_leaf.dxil
Binary files differ
diff --git a/tests/shader/gen/stack_leaf.hlsl b/tests/shader/gen/stack_leaf.hlsl
index fb56399..4969add 100644
--- a/tests/shader/gen/stack_leaf.hlsl
+++ b/tests/shader/gen/stack_leaf.hlsl
@@ -4,14 +4,14 @@
     uint b;
 };
 
-static const uint3 gl_WorkGroupSize = uint3(64u, 1u, 1u);
+static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
 
-static const Bic _237 = { 0u, 0u };
+static const Bic _157 = { 0u, 0u };
 
-ByteAddressBuffer _250 : register(t1);
-ByteAddressBuffer _512 : register(t2);
-ByteAddressBuffer _593 : register(t0);
-RWByteAddressBuffer _751 : register(u3);
+ByteAddressBuffer _170 : register(t1);
+ByteAddressBuffer _298 : register(t2);
+ByteAddressBuffer _314 : register(t0);
+RWByteAddressBuffer _399 : register(u3);
 
 static uint3 gl_WorkGroupID;
 static uint3 gl_LocalInvocationID;
@@ -23,40 +23,23 @@
     uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
 };
 
-groupshared Bic sh_bic[126];
-groupshared uint sh_bitmaps[64];
+groupshared Bic sh_bic[1022];
 groupshared uint sh_stack[512];
-groupshared uint sh_link[64];
-groupshared uint sh_next[64];
 
 Bic bic_combine(Bic x, Bic y)
 {
     uint m = min(x.b, y.a);
-    Bic _47 = { (x.a + y.a) - m, (x.b + y.b) - m };
-    return _47;
-}
-
-uint search_bit_set(uint bitmask, uint ix)
-{
-    uint result = 0u;
-    for (uint j = 0u; j < 5u; j++)
-    {
-        uint _step = 1u << (4u - j);
-        if (uint(int(countbits(bitmask & ((1u << (result + _step)) - 1u)))) <= ix)
-        {
-            result += _step;
-        }
-    }
-    return result;
+    Bic _42 = { (x.a + y.a) - m, (x.b + y.b) - m };
+    return _42;
 }
 
 uint search_link(inout Bic bic)
 {
     uint ix = gl_LocalInvocationID.x;
     uint j = 0u;
-    while (j < 6u)
+    while (j < 9u)
     {
-        uint base = 128u - (2u << (6u - j));
+        uint base = 1024u - (2u << (9u - j));
         if (((ix >> j) & 1u) != 0u)
         {
             Bic param = sh_bic[(base + (ix >> j)) - 1u];
@@ -76,7 +59,7 @@
         while (j > 0u)
         {
             j--;
-            uint base_1 = 128u - (2u << (6u - j));
+            uint base_1 = 1024u - (2u << (9u - j));
             Bic param_2 = sh_bic[(base_1 + (ix >> j)) - 1u];
             Bic param_3 = bic;
             Bic test_1 = bic_combine(param_2, param_3);
@@ -87,55 +70,41 @@
             }
         }
     }
-    if (ix > 0u)
-    {
-        ix--;
-        Bic param_4 = sh_bic[ix];
-        Bic param_5 = bic;
-        Bic test_2 = bic_combine(param_4, param_5);
-        uint param_6 = sh_bitmaps[ix];
-        uint param_7 = test_2.b - 1u;
-        uint ix_in_chunk = search_bit_set(param_6, param_7);
-        return (ix * 8u) + ix_in_chunk;
-    }
-    else
-    {
-        return 4294967295u - bic.a;
-    }
+    return ix;
 }
 
 void comp_main()
 {
     uint th = gl_LocalInvocationID.x;
-    Bic bic = _237;
-    if ((th * 8u) < gl_WorkGroupID.x)
+    Bic bic = _157;
+    if ((th * 1u) < gl_WorkGroupID.x)
     {
-        Bic _255;
-        _255.a = _250.Load((th * 8u) * 8 + 0);
-        _255.b = _250.Load((th * 8u) * 8 + 4);
-        bic.a = _255.a;
-        bic.b = _255.b;
+        Bic _175;
+        _175.a = _170.Load((th * 1u) * 8 + 0);
+        _175.b = _170.Load((th * 1u) * 8 + 4);
+        bic.a = _175.a;
+        bic.b = _175.b;
     }
     Bic other;
-    for (uint i = 1u; i < 8u; i++)
+    for (uint i = 1u; i < 1u; i++)
     {
-        if (((th * 8u) + i) < gl_WorkGroupID.x)
+        if (((th * 1u) + i) < gl_WorkGroupID.x)
         {
-            Bic _283;
-            _283.a = _250.Load(((th * 8u) + i) * 8 + 0);
-            _283.b = _250.Load(((th * 8u) + i) * 8 + 4);
-            other.a = _283.a;
-            other.b = _283.b;
+            Bic _203;
+            _203.a = _170.Load(((th * 1u) + i) * 8 + 0);
+            _203.b = _170.Load(((th * 1u) + i) * 8 + 4);
+            other.a = _203.a;
+            other.b = _203.b;
             Bic param = bic;
             Bic param_1 = other;
             bic = bic_combine(param, param_1);
         }
     }
     sh_bic[th] = bic;
-    for (uint i_1 = 0u; i_1 < 6u; i_1++)
+    for (uint i_1 = 0u; i_1 < 9u; i_1++)
     {
         GroupMemoryBarrierWithGroupSync();
-        if ((th + (1u << i_1)) < 64u)
+        if ((th + (1u << i_1)) < 512u)
         {
             Bic other_1 = sh_bic[th + (1u << i_1)];
             Bic param_2 = bic;
@@ -146,212 +115,58 @@
         sh_bic[th] = bic;
     }
     GroupMemoryBarrierWithGroupSync();
-    if (th == 63u)
-    {
-        bic = _237;
-    }
-    else
-    {
-        bic = sh_bic[th + 1u];
-    }
-    uint last_b = bic.b;
-    uint bitmap = 0u;
-    Bic param_4;
-    for (uint i_2 = 0u; i_2 < 8u; i_2++)
-    {
-        uint this_ix = (((th * 8u) + 8u) - 1u) - i_2;
-        if (this_ix < gl_WorkGroupID.x)
-        {
-            Bic _369;
-            _369.a = _250.Load(this_ix * 8 + 0);
-            _369.b = _250.Load(this_ix * 8 + 4);
-            param_4.a = _369.a;
-            param_4.b = _369.b;
-            Bic param_5 = bic;
-            bic = bic_combine(param_4, param_5);
-        }
-        sh_stack[this_ix] = bic.b;
-        if (bic.b > last_b)
-        {
-            bitmap |= (1u << (7u - i_2));
-        }
-        last_b = bic.b;
-    }
-    sh_bitmaps[th] = bitmap;
-    uint link = 0u;
-    if (bitmap != 0u)
-    {
-        link = (th * 8u) + uint(int(firstbithigh(bitmap)));
-    }
-    sh_link[th] = link;
-    for (uint i_3 = 0u; i_3 < 6u; i_3++)
-    {
-        GroupMemoryBarrierWithGroupSync();
-        if (th >= (1u << i_3))
-        {
-            link = max(link, sh_link[th - (1u << i_3)]);
-        }
-        GroupMemoryBarrierWithGroupSync();
-        sh_link[th] = link;
-    }
-    GroupMemoryBarrierWithGroupSync();
-    uint sp = 504u - (th * 8u);
+    uint sp = 511u - th;
     uint ix = 0u;
-    for (uint i_4 = 0u; i_4 < 9u; i_4++)
+    for (uint i_2 = 0u; i_2 < 9u; i_2++)
     {
-        uint probe = ix + (256u >> i_4);
-        if (sp < sh_stack[probe])
+        uint probe = ix + (256u >> i_2);
+        if (sp < sh_bic[probe].b)
         {
             ix = probe;
         }
     }
-    uint b = sh_stack[ix];
-    uint local_stack[8];
-    for (uint i_5 = 0u; i_5 < 8u; i_5++)
+    uint b = sh_bic[ix].b;
+    if (sp < b)
     {
-        local_stack[i_5] = 0u;
-    }
-    uint i_6 = 0u;
-    while ((sp + i_6) < b)
-    {
-        local_stack[7u - i_6] = _512.Load(((((ix * 512u) + b) - (sp + i_6)) - 1u) * 4 + 0);
-        i_6++;
-        if (i_6 == 8u)
-        {
-            break;
-        }
-        if ((sp + i_6) == b)
-        {
-            uint bits = sh_bitmaps[ix / 8u] & ((1u << (ix % 8u)) - 1u);
-            if (bits == 0u)
-            {
-                ix = sh_link[max((ix / 8u), 1u) - 1u];
-            }
-            else
-            {
-                ix = (ix & 4294967288u) + uint(int(firstbithigh(bits)));
-            }
-            b = sh_stack[ix];
-        }
+        sh_stack[th] = _298.Load(((((ix * 512u) + b) - sp) - 1u) * 4 + 0);
     }
     GroupMemoryBarrierWithGroupSync();
-    for (uint i_7 = 0u; i_7 < 8u; i_7++)
-    {
-        sh_stack[(th * 8u) + i_7] = local_stack[i_7];
-    }
-    uint inp = _593.Load((((gl_GlobalInvocationID.x * 8u) + 8u) - 1u) * 4 + 0);
-    Bic _605 = { 1u - inp, inp };
-    bic = _605;
-    bitmap = inp << uint(7);
-    for (uint i_8 = 7u; i_8 > 0u; i_8--)
-    {
-        inp = _593.Load((((gl_GlobalInvocationID.x * 8u) + i_8) - 1u) * 4 + 0);
-        bool _626 = inp == 1u;
-        bool _632;
-        if (_626)
-        {
-            _632 = bic.a == 0u;
-        }
-        else
-        {
-            _632 = _626;
-        }
-        if (_632)
-        {
-            bitmap |= (1u << (i_8 - 1u));
-        }
-        Bic _644 = { 1u - inp, inp };
-        Bic other_2 = _644;
-        Bic param_6 = other_2;
-        Bic param_7 = bic;
-        bic = bic_combine(param_6, param_7);
-    }
-    sh_bitmaps[th] = bitmap;
+    uint inp = _314.Load((((gl_GlobalInvocationID.x * 1u) + 1u) - 1u) * 4 + 0);
+    Bic _326 = { 1u - inp, inp };
+    bic = _326;
     sh_bic[th] = bic;
     uint inbase = 0u;
-    for (uint i_9 = 0u; i_9 < 5u; i_9++)
+    for (uint i_3 = 0u; i_3 < 8u; i_3++)
     {
-        uint outbase = 128u - (1u << (6u - i_9));
+        uint outbase = 1024u - (1u << (9u - i_3));
         GroupMemoryBarrierWithGroupSync();
-        if (th < (1u << (5u - i_9)))
+        if (th < (1u << (8u - i_3)))
         {
-            Bic param_8 = sh_bic[inbase + (th * 2u)];
-            Bic param_9 = sh_bic[(inbase + (th * 2u)) + 1u];
-            sh_bic[outbase + th] = bic_combine(param_8, param_9);
+            Bic param_4 = sh_bic[inbase + (th * 2u)];
+            Bic param_5 = sh_bic[(inbase + (th * 2u)) + 1u];
+            sh_bic[outbase + th] = bic_combine(param_4, param_5);
         }
         inbase = outbase;
     }
     GroupMemoryBarrierWithGroupSync();
-    bic.b = 0u;
-    Bic param_10 = bic;
-    uint _706 = search_link(param_10);
-    bic = param_10;
-    sh_link[th] = _706;
-    bic = _237;
-    Bic param_11 = bic;
-    uint _711 = search_link(param_11);
-    bic = param_11;
-    ix = _711;
-    uint loc_sp = 0u;
+    bic = _157;
+    Bic param_6 = bic;
+    uint _377 = search_link(param_6);
+    bic = param_6;
+    ix = _377;
     uint outp;
-    uint loc_stack[8];
-    for (uint i_10 = 0u; i_10 < 8u; i_10++)
+    if (ix > 0u)
     {
-        if (loc_sp > 0u)
-        {
-            outp = loc_stack[loc_sp - 1u];
-        }
-        else
-        {
-            if (int(ix) >= 0)
-            {
-                outp = (gl_WorkGroupID.x * 512u) + ix;
-            }
-            else
-            {
-                outp = sh_stack[512u + ix];
-            }
-        }
-        _751.Store(((gl_GlobalInvocationID.x * 8u) + i_10) * 4 + 0, outp);
-        inp = _593.Load(((gl_GlobalInvocationID.x * 8u) + i_10) * 4 + 0);
-        if (inp == 1u)
-        {
-            loc_stack[loc_sp] = (gl_GlobalInvocationID.x * 8u) + i_10;
-            loc_sp++;
-        }
-        else
-        {
-            if (inp == 0u)
-            {
-                if (loc_sp > 0u)
-                {
-                    loc_sp--;
-                }
-                else
-                {
-                    if (int(ix) >= 0)
-                    {
-                        uint bits_1 = sh_bitmaps[ix / 8u] & ((1u << (ix % 8u)) - 1u);
-                        if (bits_1 == 0u)
-                        {
-                            ix = sh_link[ix / 8u];
-                        }
-                        else
-                        {
-                            ix = (ix & 4294967288u) + uint(int(firstbithigh(bits_1)));
-                        }
-                    }
-                    else
-                    {
-                        ix--;
-                    }
-                }
-            }
-        }
+        outp = ((gl_WorkGroupID.x * 512u) + ix) - 1u;
     }
+    else
+    {
+        outp = sh_stack[511u - bic.a];
+    }
+    _399.Store(gl_GlobalInvocationID.x * 4 + 0, outp);
 }
 
-[numthreads(64, 1, 1)]
+[numthreads(512, 1, 1)]
 void main(SPIRV_Cross_Input stage_input)
 {
     gl_WorkGroupID = stage_input.gl_WorkGroupID;
diff --git a/tests/shader/gen/stack_leaf.msl b/tests/shader/gen/stack_leaf.msl
index cbf72a3..13afb0a 100644
--- a/tests/shader/gen/stack_leaf.msl
+++ b/tests/shader/gen/stack_leaf.msl
@@ -1,56 +1,10 @@
 #pragma clang diagnostic ignored "-Wmissing-prototypes"
-#pragma clang diagnostic ignored "-Wmissing-braces"
 
 #include <metal_stdlib>
 #include <simd/simd.h>
 
 using namespace metal;
 
-template<typename T, size_t Num>
-struct spvUnsafeArray
-{
-    T elements[Num ? Num : 1];
-    
-    thread T& operator [] (size_t pos) thread
-    {
-        return elements[pos];
-    }
-    constexpr const thread T& operator [] (size_t pos) const thread
-    {
-        return elements[pos];
-    }
-    
-    device T& operator [] (size_t pos) device
-    {
-        return elements[pos];
-    }
-    constexpr const device T& operator [] (size_t pos) const device
-    {
-        return elements[pos];
-    }
-    
-    constexpr const constant T& operator [] (size_t pos) const constant
-    {
-        return elements[pos];
-    }
-    
-    threadgroup T& operator [] (size_t pos) threadgroup
-    {
-        return elements[pos];
-    }
-    constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
-    {
-        return elements[pos];
-    }
-};
-
-// Implementation of the unsigned GLSL findMSB() function
-template<typename T>
-inline T spvFindUMSB(T x)
-{
-    return select(clz(T(0)) - (clz(x) + T(1)), T(-1), x == T(0));
-}
-
 struct Bic
 {
     uint a;
@@ -83,7 +37,7 @@
     uint outbuf[1];
 };
 
-constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(512u, 1u, 1u);
 
 static inline __attribute__((always_inline))
 Bic bic_combine(thread const Bic& x, thread const Bic& y)
@@ -93,28 +47,13 @@
 }
 
 static inline __attribute__((always_inline))
-uint search_bit_set(thread const uint& bitmask, thread const uint& ix)
-{
-    uint result = 0u;
-    for (uint j = 0u; j < 5u; j++)
-    {
-        uint _step = 1u << (4u - j);
-        if (uint(int(popcount(bitmask & ((1u << (result + _step)) - 1u)))) <= ix)
-        {
-            result += _step;
-        }
-    }
-    return result;
-}
-
-static inline __attribute__((always_inline))
-uint search_link(thread Bic& bic, thread uint3& gl_LocalInvocationID, threadgroup Bic (&sh_bic)[126], threadgroup uint (&sh_bitmaps)[64])
+uint search_link(thread Bic& bic, thread uint3& gl_LocalInvocationID, threadgroup Bic (&sh_bic)[1022])
 {
     uint ix = gl_LocalInvocationID.x;
     uint j = 0u;
-    while (j < 6u)
+    while (j < 9u)
     {
-        uint base = 128u - (2u << (6u - j));
+        uint base = 1024u - (2u << (9u - j));
         if (((ix >> j) & 1u) != 0u)
         {
             Bic param = sh_bic[(base + (ix >> j)) - 1u];
@@ -134,7 +73,7 @@
         while (j > 0u)
         {
             j--;
-            uint base_1 = 128u - (2u << (6u - j));
+            uint base_1 = 1024u - (2u << (9u - j));
             Bic param_2 = sh_bic[(base_1 + (ix >> j)) - 1u];
             Bic param_3 = bic;
             Bic test_1 = bic_combine(param_2, param_3);
@@ -145,56 +84,39 @@
             }
         }
     }
-    if (ix > 0u)
-    {
-        ix--;
-        Bic param_4 = sh_bic[ix];
-        Bic param_5 = bic;
-        Bic test_2 = bic_combine(param_4, param_5);
-        uint param_6 = sh_bitmaps[ix];
-        uint param_7 = test_2.b - 1u;
-        uint ix_in_chunk = search_bit_set(param_6, param_7);
-        return (ix * 8u) + ix_in_chunk;
-    }
-    else
-    {
-        return 4294967295u - bic.a;
-    }
+    return ix;
 }
 
-kernel void main0(const device InBuf& _593 [[buffer(0)]], const device BicBuf& _250 [[buffer(1)]], const device StackBuf& _512 [[buffer(2)]], device OutBuf& _751 [[buffer(3)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+kernel void main0(const device InBuf& _314 [[buffer(0)]], const device BicBuf& _170 [[buffer(1)]], const device StackBuf& _298 [[buffer(2)]], device OutBuf& _399 [[buffer(3)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
-    threadgroup Bic sh_bic[126];
-    threadgroup uint sh_bitmaps[64];
+    threadgroup Bic sh_bic[1022];
     threadgroup uint sh_stack[512];
-    threadgroup uint sh_link[64];
-    threadgroup uint sh_next[64];
     uint th = gl_LocalInvocationID.x;
     Bic bic = Bic{ 0u, 0u };
-    if ((th * 8u) < gl_WorkGroupID.x)
+    if ((th * 1u) < gl_WorkGroupID.x)
     {
-        uint _252 = th * 8u;
-        bic.a = _250.bicbuf[_252].a;
-        bic.b = _250.bicbuf[_252].b;
+        uint _172 = th * 1u;
+        bic.a = _170.bicbuf[_172].a;
+        bic.b = _170.bicbuf[_172].b;
     }
     Bic other;
-    for (uint i = 1u; i < 8u; i++)
+    for (uint i = 1u; i < 1u; i++)
     {
-        if (((th * 8u) + i) < gl_WorkGroupID.x)
+        if (((th * 1u) + i) < gl_WorkGroupID.x)
         {
-            uint _281 = (th * 8u) + i;
-            other.a = _250.bicbuf[_281].a;
-            other.b = _250.bicbuf[_281].b;
+            uint _201 = (th * 1u) + i;
+            other.a = _170.bicbuf[_201].a;
+            other.b = _170.bicbuf[_201].b;
             Bic param = bic;
             Bic param_1 = other;
             bic = bic_combine(param, param_1);
         }
     }
     sh_bic[th] = bic;
-    for (uint i_1 = 0u; i_1 < 6u; i_1++)
+    for (uint i_1 = 0u; i_1 < 9u; i_1++)
     {
         threadgroup_barrier(mem_flags::mem_threadgroup);
-        if ((th + (1u << i_1)) < 64u)
+        if ((th + (1u << i_1)) < 512u)
         {
             Bic other_1 = sh_bic[th + (1u << i_1)];
             Bic param_2 = bic;
@@ -205,203 +127,53 @@
         sh_bic[th] = bic;
     }
     threadgroup_barrier(mem_flags::mem_threadgroup);
-    if (th == 63u)
-    {
-        bic = Bic{ 0u, 0u };
-    }
-    else
-    {
-        bic = sh_bic[th + 1u];
-    }
-    uint last_b = bic.b;
-    uint bitmap = 0u;
-    Bic param_4;
-    for (uint i_2 = 0u; i_2 < 8u; i_2++)
-    {
-        uint this_ix = (((th * 8u) + 8u) - 1u) - i_2;
-        if (this_ix < gl_WorkGroupID.x)
-        {
-            param_4.a = _250.bicbuf[this_ix].a;
-            param_4.b = _250.bicbuf[this_ix].b;
-            Bic param_5 = bic;
-            bic = bic_combine(param_4, param_5);
-        }
-        sh_stack[this_ix] = bic.b;
-        if (bic.b > last_b)
-        {
-            bitmap |= (1u << (7u - i_2));
-        }
-        last_b = bic.b;
-    }
-    sh_bitmaps[th] = bitmap;
-    uint link = 0u;
-    if (bitmap != 0u)
-    {
-        link = (th * 8u) + uint(int(spvFindUMSB(bitmap)));
-    }
-    sh_link[th] = link;
-    for (uint i_3 = 0u; i_3 < 6u; i_3++)
-    {
-        threadgroup_barrier(mem_flags::mem_threadgroup);
-        if (th >= (1u << i_3))
-        {
-            link = max(link, sh_link[th - (1u << i_3)]);
-        }
-        threadgroup_barrier(mem_flags::mem_threadgroup);
-        sh_link[th] = link;
-    }
-    threadgroup_barrier(mem_flags::mem_threadgroup);
-    uint sp = 504u - (th * 8u);
+    uint sp = 511u - th;
     uint ix = 0u;
-    for (uint i_4 = 0u; i_4 < 9u; i_4++)
+    for (uint i_2 = 0u; i_2 < 9u; i_2++)
     {
-        uint probe = ix + (256u >> i_4);
-        if (sp < sh_stack[probe])
+        uint probe = ix + (256u >> i_2);
+        if (sp < sh_bic[probe].b)
         {
             ix = probe;
         }
     }
-    uint b = sh_stack[ix];
-    spvUnsafeArray<uint, 8> local_stack;
-    for (uint i_5 = 0u; i_5 < 8u; i_5++)
+    uint b = sh_bic[ix].b;
+    if (sp < b)
     {
-        local_stack[i_5] = 0u;
-    }
-    uint i_6 = 0u;
-    while ((sp + i_6) < b)
-    {
-        local_stack[7u - i_6] = _512.stack[(((ix * 512u) + b) - (sp + i_6)) - 1u];
-        i_6++;
-        if (i_6 == 8u)
-        {
-            break;
-        }
-        if ((sp + i_6) == b)
-        {
-            uint bits = sh_bitmaps[ix / 8u] & ((1u << (ix % 8u)) - 1u);
-            if (bits == 0u)
-            {
-                ix = sh_link[max((ix / 8u), 1u) - 1u];
-            }
-            else
-            {
-                ix = (ix & 4294967288u) + uint(int(spvFindUMSB(bits)));
-            }
-            b = sh_stack[ix];
-        }
+        sh_stack[th] = _298.stack[(((ix * 512u) + b) - sp) - 1u];
     }
     threadgroup_barrier(mem_flags::mem_threadgroup);
-    for (uint i_7 = 0u; i_7 < 8u; i_7++)
-    {
-        sh_stack[(th * 8u) + i_7] = local_stack[i_7];
-    }
-    uint inp = _593.inbuf[((gl_GlobalInvocationID.x * 8u) + 8u) - 1u];
+    uint inp = _314.inbuf[((gl_GlobalInvocationID.x * 1u) + 1u) - 1u];
     bic = Bic{ 1u - inp, inp };
-    bitmap = inp << uint(7);
-    for (uint i_8 = 7u; i_8 > 0u; i_8--)
-    {
-        inp = _593.inbuf[((gl_GlobalInvocationID.x * 8u) + i_8) - 1u];
-        bool _626 = inp == 1u;
-        bool _632;
-        if (_626)
-        {
-            _632 = bic.a == 0u;
-        }
-        else
-        {
-            _632 = _626;
-        }
-        if (_632)
-        {
-            bitmap |= (1u << (i_8 - 1u));
-        }
-        Bic other_2 = Bic{ 1u - inp, inp };
-        Bic param_6 = other_2;
-        Bic param_7 = bic;
-        bic = bic_combine(param_6, param_7);
-    }
-    sh_bitmaps[th] = bitmap;
     sh_bic[th] = bic;
     uint inbase = 0u;
-    for (uint i_9 = 0u; i_9 < 5u; i_9++)
+    for (uint i_3 = 0u; i_3 < 8u; i_3++)
     {
-        uint outbase = 128u - (1u << (6u - i_9));
+        uint outbase = 1024u - (1u << (9u - i_3));
         threadgroup_barrier(mem_flags::mem_threadgroup);
-        if (th < (1u << (5u - i_9)))
+        if (th < (1u << (8u - i_3)))
         {
-            Bic param_8 = sh_bic[inbase + (th * 2u)];
-            Bic param_9 = sh_bic[(inbase + (th * 2u)) + 1u];
-            sh_bic[outbase + th] = bic_combine(param_8, param_9);
+            Bic param_4 = sh_bic[inbase + (th * 2u)];
+            Bic param_5 = sh_bic[(inbase + (th * 2u)) + 1u];
+            sh_bic[outbase + th] = bic_combine(param_4, param_5);
         }
         inbase = outbase;
     }
     threadgroup_barrier(mem_flags::mem_threadgroup);
-    bic.b = 0u;
-    Bic param_10 = bic;
-    uint _706 = search_link(param_10, gl_LocalInvocationID, sh_bic, sh_bitmaps);
-    bic = param_10;
-    sh_link[th] = _706;
     bic = Bic{ 0u, 0u };
-    Bic param_11 = bic;
-    uint _711 = search_link(param_11, gl_LocalInvocationID, sh_bic, sh_bitmaps);
-    bic = param_11;
-    ix = _711;
-    uint loc_sp = 0u;
+    Bic param_6 = bic;
+    uint _377 = search_link(param_6, gl_LocalInvocationID, sh_bic);
+    bic = param_6;
+    ix = _377;
     uint outp;
-    spvUnsafeArray<uint, 8> loc_stack;
-    for (uint i_10 = 0u; i_10 < 8u; i_10++)
+    if (ix > 0u)
     {
-        if (loc_sp > 0u)
-        {
-            outp = loc_stack[loc_sp - 1u];
-        }
-        else
-        {
-            if (int(ix) >= 0)
-            {
-                outp = (gl_WorkGroupID.x * 512u) + ix;
-            }
-            else
-            {
-                outp = sh_stack[512u + ix];
-            }
-        }
-        _751.outbuf[(gl_GlobalInvocationID.x * 8u) + i_10] = outp;
-        inp = _593.inbuf[(gl_GlobalInvocationID.x * 8u) + i_10];
-        if (inp == 1u)
-        {
-            loc_stack[loc_sp] = (gl_GlobalInvocationID.x * 8u) + i_10;
-            loc_sp++;
-        }
-        else
-        {
-            if (inp == 0u)
-            {
-                if (loc_sp > 0u)
-                {
-                    loc_sp--;
-                }
-                else
-                {
-                    if (int(ix) >= 0)
-                    {
-                        uint bits_1 = sh_bitmaps[ix / 8u] & ((1u << (ix % 8u)) - 1u);
-                        if (bits_1 == 0u)
-                        {
-                            ix = sh_link[ix / 8u];
-                        }
-                        else
-                        {
-                            ix = (ix & 4294967288u) + uint(int(spvFindUMSB(bits_1)));
-                        }
-                    }
-                    else
-                    {
-                        ix--;
-                    }
-                }
-            }
-        }
+        outp = ((gl_WorkGroupID.x * 512u) + ix) - 1u;
     }
+    else
+    {
+        outp = sh_stack[511u - bic.a];
+    }
+    _399.outbuf[gl_GlobalInvocationID.x] = outp;
 }
 
diff --git a/tests/shader/gen/stack_leaf.spv b/tests/shader/gen/stack_leaf.spv
index 827ef14..527d20c 100644
--- a/tests/shader/gen/stack_leaf.spv
+++ b/tests/shader/gen/stack_leaf.spv
Binary files differ
diff --git a/tests/shader/gen/stack_reduce.dxil b/tests/shader/gen/stack_reduce.dxil
index fd97e87..50ceb82 100644
--- a/tests/shader/gen/stack_reduce.dxil
+++ b/tests/shader/gen/stack_reduce.dxil
Binary files differ
diff --git a/tests/shader/gen/stack_reduce.hlsl b/tests/shader/gen/stack_reduce.hlsl
index 0eab9cb..5df49cd 100644
--- a/tests/shader/gen/stack_reduce.hlsl
+++ b/tests/shader/gen/stack_reduce.hlsl
@@ -4,13 +4,13 @@
     uint b;
 };
 
-static const uint3 gl_WorkGroupSize = uint3(64u, 1u, 1u);
+static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
 
-static const Bic _175 = { 0u, 0u };
+static const Bic _174 = { 0u, 0u };
 
 ByteAddressBuffer _48 : register(t0);
-RWByteAddressBuffer _160 : register(u1);
-RWByteAddressBuffer _223 : register(u2);
+RWByteAddressBuffer _159 : register(u1);
+RWByteAddressBuffer _221 : register(u2);
 
 static uint3 gl_WorkGroupID;
 static uint3 gl_LocalInvocationID;
@@ -22,7 +22,7 @@
     uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
 };
 
-groupshared Bic sh_bic[64];
+groupshared Bic sh_bic[512];
 
 Bic bic_combine(Bic x, Bic y)
 {
@@ -33,24 +33,24 @@
 
 void comp_main()
 {
-    uint inp[8];
-    inp[0] = _48.Load((gl_GlobalInvocationID.x * 8u) * 4 + 0);
-    Bic _68 = { 1u - inp[0], inp[0] };
-    Bic bic = _68;
-    for (uint i = 1u; i < 8u; i++)
+    uint inp[1];
+    inp[0] = _48.Load((gl_GlobalInvocationID.x * 1u) * 4 + 0);
+    Bic _67 = { 1u - inp[0], inp[0] };
+    Bic bic = _67;
+    for (uint i = 1u; i < 1u; i++)
     {
-        inp[i] = _48.Load(((gl_GlobalInvocationID.x * 8u) + i) * 4 + 0);
-        Bic _95 = { 1u - inp[i], inp[i] };
-        Bic other = _95;
+        inp[i] = _48.Load(((gl_GlobalInvocationID.x * 1u) + i) * 4 + 0);
+        Bic _94 = { 1u - inp[i], inp[i] };
+        Bic other = _94;
         Bic param = bic;
         Bic param_1 = other;
         bic = bic_combine(param, param_1);
     }
     sh_bic[gl_LocalInvocationID.x] = bic;
-    for (uint i_1 = 0u; i_1 < 6u; i_1++)
+    for (uint i_1 = 0u; i_1 < 9u; i_1++)
     {
         GroupMemoryBarrierWithGroupSync();
-        if ((gl_LocalInvocationID.x + (1u << i_1)) < 64u)
+        if ((gl_LocalInvocationID.x + (1u << i_1)) < 512u)
         {
             Bic other_1 = sh_bic[gl_LocalInvocationID.x + (1u << i_1)];
             Bic param_2 = bic;
@@ -62,43 +62,43 @@
     }
     if (gl_LocalInvocationID.x == 0u)
     {
-        _160.Store(gl_WorkGroupID.x * 8 + 0, bic.a);
-        _160.Store(gl_WorkGroupID.x * 8 + 4, bic.b);
+        _159.Store(gl_WorkGroupID.x * 8 + 0, bic.a);
+        _159.Store(gl_WorkGroupID.x * 8 + 4, bic.b);
     }
     GroupMemoryBarrierWithGroupSync();
     uint size = sh_bic[0].b;
-    bic = _175;
-    if ((gl_LocalInvocationID.x + 1u) < 64u)
+    bic = _174;
+    if ((gl_LocalInvocationID.x + 1u) < 512u)
     {
         bic = sh_bic[gl_LocalInvocationID.x + 1u];
     }
     uint out_ix = ((gl_WorkGroupID.x * 512u) + size) - bic.b;
-    for (uint i_2 = 8u; i_2 > 0u; i_2--)
+    for (uint i_2 = 1u; i_2 > 0u; i_2--)
     {
-        bool _209 = inp[i_2 - 1u] == 1u;
-        bool _215;
-        if (_209)
+        bool _207 = inp[i_2 - 1u] == 1u;
+        bool _213;
+        if (_207)
         {
-            _215 = bic.a == 0u;
+            _213 = bic.a == 0u;
         }
         else
         {
-            _215 = _209;
+            _213 = _207;
         }
-        if (_215)
+        if (_213)
         {
             out_ix--;
-            _223.Store(out_ix * 4 + 0, ((gl_GlobalInvocationID.x * 8u) + i_2) - 1u);
+            _221.Store(out_ix * 4 + 0, ((gl_GlobalInvocationID.x * 1u) + i_2) - 1u);
         }
-        Bic _242 = { 1u - inp[i_2 - 1u], inp[i_2 - 1u] };
-        Bic other_2 = _242;
+        Bic _240 = { 1u - inp[i_2 - 1u], inp[i_2 - 1u] };
+        Bic other_2 = _240;
         Bic param_4 = other_2;
         Bic param_5 = bic;
         bic = bic_combine(param_4, param_5);
     }
 }
 
-[numthreads(64, 1, 1)]
+[numthreads(512, 1, 1)]
 void main(SPIRV_Cross_Input stage_input)
 {
     gl_WorkGroupID = stage_input.gl_WorkGroupID;
diff --git a/tests/shader/gen/stack_reduce.msl b/tests/shader/gen/stack_reduce.msl
index 2723dee..13ce9fa 100644
--- a/tests/shader/gen/stack_reduce.msl
+++ b/tests/shader/gen/stack_reduce.msl
@@ -71,7 +71,7 @@
     uint stack[1];
 };
 
-constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(512u, 1u, 1u);
 
 static inline __attribute__((always_inline))
 Bic bic_combine(thread const Bic& x, thread const Bic& y)
@@ -80,25 +80,25 @@
     return Bic{ (x.a + y.a) - m, (x.b + y.b) - m };
 }
 
-kernel void main0(const device InBuf& _48 [[buffer(0)]], device OutBuf& _160 [[buffer(1)]], device StackBuf& _223 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+kernel void main0(const device InBuf& _48 [[buffer(0)]], device OutBuf& _159 [[buffer(1)]], device StackBuf& _221 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
 {
-    threadgroup Bic sh_bic[64];
-    spvUnsafeArray<uint, 8> inp;
-    inp[0] = _48.inbuf[gl_GlobalInvocationID.x * 8u];
+    threadgroup Bic sh_bic[512];
+    spvUnsafeArray<uint, 1> inp;
+    inp[0] = _48.inbuf[gl_GlobalInvocationID.x * 1u];
     Bic bic = Bic{ 1u - inp[0], inp[0] };
-    for (uint i = 1u; i < 8u; i++)
+    for (uint i = 1u; i < 1u; i++)
     {
-        inp[i] = _48.inbuf[(gl_GlobalInvocationID.x * 8u) + i];
+        inp[i] = _48.inbuf[(gl_GlobalInvocationID.x * 1u) + i];
         Bic other = Bic{ 1u - inp[i], inp[i] };
         Bic param = bic;
         Bic param_1 = other;
         bic = bic_combine(param, param_1);
     }
     sh_bic[gl_LocalInvocationID.x] = bic;
-    for (uint i_1 = 0u; i_1 < 6u; i_1++)
+    for (uint i_1 = 0u; i_1 < 9u; i_1++)
     {
         threadgroup_barrier(mem_flags::mem_threadgroup);
-        if ((gl_LocalInvocationID.x + (1u << i_1)) < 64u)
+        if ((gl_LocalInvocationID.x + (1u << i_1)) < 512u)
         {
             Bic other_1 = sh_bic[gl_LocalInvocationID.x + (1u << i_1)];
             Bic param_2 = bic;
@@ -110,33 +110,33 @@
     }
     if (gl_LocalInvocationID.x == 0u)
     {
-        _160.outbuf[gl_WorkGroupID.x].a = bic.a;
-        _160.outbuf[gl_WorkGroupID.x].b = bic.b;
+        _159.outbuf[gl_WorkGroupID.x].a = bic.a;
+        _159.outbuf[gl_WorkGroupID.x].b = bic.b;
     }
     threadgroup_barrier(mem_flags::mem_threadgroup);
     uint size = sh_bic[0].b;
     bic = Bic{ 0u, 0u };
-    if ((gl_LocalInvocationID.x + 1u) < 64u)
+    if ((gl_LocalInvocationID.x + 1u) < 512u)
     {
         bic = sh_bic[gl_LocalInvocationID.x + 1u];
     }
     uint out_ix = ((gl_WorkGroupID.x * 512u) + size) - bic.b;
-    for (uint i_2 = 8u; i_2 > 0u; i_2--)
+    for (uint i_2 = 1u; i_2 > 0u; i_2--)
     {
-        bool _209 = inp[i_2 - 1u] == 1u;
-        bool _215;
-        if (_209)
+        bool _207 = inp[i_2 - 1u] == 1u;
+        bool _213;
+        if (_207)
         {
-            _215 = bic.a == 0u;
+            _213 = bic.a == 0u;
         }
         else
         {
-            _215 = _209;
+            _213 = _207;
         }
-        if (_215)
+        if (_213)
         {
             out_ix--;
-            _223.stack[out_ix] = ((gl_GlobalInvocationID.x * 8u) + i_2) - 1u;
+            _221.stack[out_ix] = ((gl_GlobalInvocationID.x * 1u) + i_2) - 1u;
         }
         Bic other_2 = Bic{ 1u - inp[i_2 - 1u], inp[i_2 - 1u] };
         Bic param_4 = other_2;
diff --git a/tests/shader/gen/stack_reduce.spv b/tests/shader/gen/stack_reduce.spv
index 90be43a..733c3a0 100644
--- a/tests/shader/gen/stack_reduce.spv
+++ b/tests/shader/gen/stack_reduce.spv
Binary files differ
diff --git a/tests/shader/intersection_leaf.comp b/tests/shader/intersection_leaf.comp
index c3e594b..04918f1 100644
--- a/tests/shader/intersection_leaf.comp
+++ b/tests/shader/intersection_leaf.comp
@@ -6,7 +6,7 @@
 
 #define LG_N_SEQ 0
 #define N_SEQ (1 << LG_N_SEQ)
-#define LG_WG_SIZE 8
+#define LG_WG_SIZE 9
 #define WG_SIZE (1 << LG_WG_SIZE)
 #define PART_SIZE (WG_SIZE * N_SEQ)
 
diff --git a/tests/shader/intersection_reduce.comp b/tests/shader/intersection_reduce.comp
index cc1e600..19e568b 100644
--- a/tests/shader/intersection_reduce.comp
+++ b/tests/shader/intersection_reduce.comp
@@ -5,7 +5,7 @@
 #version 450
 
 // At least for now, N_SEQ is hardwired at 1
-#define LG_WG_SIZE 8
+#define LG_WG_SIZE 9
 #define WG_SIZE (1 << LG_WG_SIZE)
 #define PART_SIZE WG_SIZE
 
diff --git a/tests/shader/stack_leaf.comp b/tests/shader/stack_leaf.comp
index 1cdb099..23bf47a 100644
--- a/tests/shader/stack_leaf.comp
+++ b/tests/shader/stack_leaf.comp
@@ -4,9 +4,9 @@
 
 #version 450
 
-#define LG_N_SEQ 3
+#define LG_N_SEQ 0
 #define N_SEQ (1 << LG_N_SEQ)
-#define LG_WG_SIZE 6
+#define LG_WG_SIZE 9
 #define WG_SIZE (1 << LG_WG_SIZE)
 #define PART_SIZE (WG_SIZE * N_SEQ)
 
@@ -232,6 +232,7 @@
         sh_stack[th] = stack[ix * PART_SIZE + b - sp - 1];
     }
 #endif
+    barrier();
 
     // Do tree reduction of bicyclic semigroups (up-sweep)
     uint inp = inbuf[gl_GlobalInvocationID.x * N_SEQ + N_SEQ - 1];
diff --git a/tests/shader/stack_reduce.comp b/tests/shader/stack_reduce.comp
index b066135..71da8b6 100644
--- a/tests/shader/stack_reduce.comp
+++ b/tests/shader/stack_reduce.comp
@@ -4,8 +4,8 @@
 
 #version 450
 
-#define N_SEQ 8
-#define LG_WG_SIZE 6
+#define N_SEQ 1
+#define LG_WG_SIZE 9
 #define WG_SIZE (1 << LG_WG_SIZE)
 #define PART_SIZE (WG_SIZE * N_SEQ)
 
diff --git a/tests/shader/union_leaf.comp b/tests/shader/union_leaf.comp
index e40bc6d..49775eb 100644
--- a/tests/shader/union_leaf.comp
+++ b/tests/shader/union_leaf.comp
@@ -6,7 +6,7 @@
 
 #define LG_N_SEQ 0
 #define N_SEQ (1 << LG_N_SEQ)
-#define LG_WG_SIZE 8
+#define LG_WG_SIZE 9
 #define WG_SIZE (1 << LG_WG_SIZE)
 #define PART_SIZE (WG_SIZE * N_SEQ)
 
diff --git a/tests/shader/union_reduce.comp b/tests/shader/union_reduce.comp
index f584ae3..fec376c 100644
--- a/tests/shader/union_reduce.comp
+++ b/tests/shader/union_reduce.comp
@@ -5,7 +5,7 @@
 #version 450
 
 // At least for now, N_SEQ is hardwired at 1
-#define LG_WG_SIZE 8
+#define LG_WG_SIZE 9
 #define WG_SIZE (1 << LG_WG_SIZE)
 #define PART_SIZE WG_SIZE
 
diff --git a/tests/src/bbox_intersection.rs b/tests/src/bbox_intersection.rs
index fac9704..2c006a8 100644
--- a/tests/src/bbox_intersection.rs
+++ b/tests/src/bbox_intersection.rs
@@ -24,7 +24,7 @@
 use crate::runner::Runner;
 use crate::test_result::TestResult;
 
-const WG_SIZE: u64 = 256;
+const WG_SIZE: u64 = 512;
 const N_ROWS: u64 = 1;
 const ELEMENTS_PER_WG: u64 = WG_SIZE * N_ROWS;
 
diff --git a/tests/src/bbox_union.rs b/tests/src/bbox_union.rs
index ca2224c..bf77f59 100644
--- a/tests/src/bbox_union.rs
+++ b/tests/src/bbox_union.rs
@@ -24,7 +24,7 @@
 use crate::runner::Runner;
 use crate::test_result::TestResult;
 
-const WG_SIZE: u64 = 256;
+const WG_SIZE: u64 = 512;
 const N_ROWS: u64 = 1;
 const ELEMENTS_PER_WG: u64 = WG_SIZE * N_ROWS;
 
diff --git a/tests/src/stack.rs b/tests/src/stack.rs
index aba875b..923010d 100644
--- a/tests/src/stack.rs
+++ b/tests/src/stack.rs
@@ -21,8 +21,8 @@
 use crate::runner::Runner;
 use crate::test_result::TestResult;
 
-const WG_SIZE: u64 = 64;
-const N_ROWS: u64 = 8;
+const WG_SIZE: u64 = 512;
+const N_ROWS: u64 = 1;
 const ELEMENTS_PER_WG: u64 = WG_SIZE * N_ROWS;
 
 struct StackCode {