From b36ca7fc2e75f25abd752477e7766a47a3280968 Mon Sep 17 00:00:00 2001
From: Raph Levien <raph.levien@gmail.com>
Date: Sat, 6 Nov 2021 16:25:56 -0700
Subject: [PATCH] Add generated shaders

---
 tests/shader/gen/prefix_reduce.hlsl |  72 ++++++++++++++++
 tests/shader/gen/prefix_reduce.msl  |  68 +++++++++++++++
 tests/shader/gen/prefix_reduce.spv  | Bin 0 -> 3504 bytes
 tests/shader/gen/prefix_root.hlsl   |  80 ++++++++++++++++++
 tests/shader/gen/prefix_root.msl    | 112 +++++++++++++++++++++++++
 tests/shader/gen/prefix_root.spv    | Bin 0 -> 4104 bytes
 tests/shader/gen/prefix_scan.hlsl   |  92 +++++++++++++++++++++
 tests/shader/gen/prefix_scan.msl    | 123 ++++++++++++++++++++++++++++
 tests/shader/gen/prefix_scan.spv    | Bin 0 -> 4736 bytes
 9 files changed, 547 insertions(+)
 create mode 100644 tests/shader/gen/prefix_reduce.hlsl
 create mode 100644 tests/shader/gen/prefix_reduce.msl
 create mode 100644 tests/shader/gen/prefix_reduce.spv
 create mode 100644 tests/shader/gen/prefix_root.hlsl
 create mode 100644 tests/shader/gen/prefix_root.msl
 create mode 100644 tests/shader/gen/prefix_root.spv
 create mode 100644 tests/shader/gen/prefix_scan.hlsl
 create mode 100644 tests/shader/gen/prefix_scan.msl
 create mode 100644 tests/shader/gen/prefix_scan.spv

diff --git a/tests/shader/gen/prefix_reduce.hlsl b/tests/shader/gen/prefix_reduce.hlsl
new file mode 100644
index 0000000..837a75a
--- /dev/null
+++ b/tests/shader/gen/prefix_reduce.hlsl
@@ -0,0 +1,72 @@
+struct Monoid
+{
+    uint element;
+};
+
+static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
+
+ByteAddressBuffer _40 : register(t0);
+RWByteAddressBuffer _129 : register(u1);
+
+static uint3 gl_WorkGroupID;
+static uint3 gl_LocalInvocationID;
+static uint3 gl_GlobalInvocationID;
+struct SPIRV_Cross_Input
+{
+    uint3 gl_WorkGroupID : SV_GroupID;
+    uint3 gl_LocalInvocationID : SV_GroupThreadID;
+    uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
+};
+
+groupshared Monoid sh_scratch[512];
+
+Monoid combine_monoid(Monoid a, Monoid b)
+{
+    Monoid _22 = { a.element + b.element };
+    return _22;
+}
+
+void comp_main()
+{
+    uint ix = gl_GlobalInvocationID.x * 8u;
+    Monoid _44;
+    _44.element = _40.Load(ix * 4 + 0);
+    Monoid agg;
+    agg.element = _44.element;
+    Monoid param_1;
+    for (uint i = 1u; i < 8u; i++)
+    {
+        Monoid param = agg;
+        Monoid _64;
+        _64.element = _40.Load((ix + i) * 4 + 0);
+        param_1.element = _64.element;
+        agg = combine_monoid(param, param_1);
+    }
+    sh_scratch[gl_LocalInvocationID.x] = agg;
+    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    {
+        GroupMemoryBarrierWithGroupSync();
+        if ((gl_LocalInvocationID.x + uint(1 << int(i_1))) < 512u)
+        {
+            Monoid other = sh_scratch[gl_LocalInvocationID.x + uint(1 << int(i_1))];
+            Monoid param_2 = agg;
+            Monoid param_3 = other;
+            agg = combine_monoid(param_2, param_3);
+        }
+        GroupMemoryBarrierWithGroupSync();
+        sh_scratch[gl_LocalInvocationID.x] = agg;
+    }
+    if (gl_LocalInvocationID.x == 0u)
+    {
+        _129.Store(gl_WorkGroupID.x * 4 + 0, agg.element);
+    }
+}
+
+[numthreads(512, 1, 1)]
+void main(SPIRV_Cross_Input stage_input)
+{
+    gl_WorkGroupID = stage_input.gl_WorkGroupID;
+    gl_LocalInvocationID = stage_input.gl_LocalInvocationID;
+    gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID;
+    comp_main();
+}
diff --git a/tests/shader/gen/prefix_reduce.msl b/tests/shader/gen/prefix_reduce.msl
new file mode 100644
index 0000000..e1ed0ce
--- /dev/null
+++ b/tests/shader/gen/prefix_reduce.msl
@@ -0,0 +1,68 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct Monoid
+{
+    uint element;
+};
+
+struct Monoid_1
+{
+    uint element;
+};
+
+struct InBuf
+{
+    Monoid_1 inbuf[1];
+};
+
+struct OutBuf
+{
+    Monoid_1 outbuf[1];
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(512u, 1u, 1u);
+
+static inline __attribute__((always_inline))
+Monoid combine_monoid(thread const Monoid& a, thread const Monoid& b)
+{
+    return Monoid{ a.element + b.element };
+}
+
+kernel void main0(const device InBuf& _40 [[buffer(0)]], device OutBuf& _129 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+{
+    threadgroup Monoid sh_scratch[512];
+    uint ix = gl_GlobalInvocationID.x * 8u;
+    Monoid agg;
+    agg.element = _40.inbuf[ix].element;
+    Monoid param_1;
+    for (uint i = 1u; i < 8u; i++)
+    {
+        Monoid param = agg;
+        param_1.element = _40.inbuf[ix + i].element;
+        agg = combine_monoid(param, param_1);
+    }
+    sh_scratch[gl_LocalInvocationID.x] = agg;
+    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    {
+        threadgroup_barrier(mem_flags::mem_threadgroup);
+        if ((gl_LocalInvocationID.x + uint(1 << int(i_1))) < 512u)
+        {
+            Monoid other = sh_scratch[gl_LocalInvocationID.x + uint(1 << int(i_1))];
+            Monoid param_2 = agg;
+            Monoid param_3 = other;
+            agg = combine_monoid(param_2, param_3);
+        }
+        threadgroup_barrier(mem_flags::mem_threadgroup);
+        sh_scratch[gl_LocalInvocationID.x] = agg;
+    }
+    if (gl_LocalInvocationID.x == 0u)
+    {
+        _129.outbuf[gl_WorkGroupID.x].element = agg.element;
+    }
+}
+
diff --git a/tests/shader/gen/prefix_reduce.spv b/tests/shader/gen/prefix_reduce.spv
new file mode 100644
index 0000000000000000000000000000000000000000..d1db3aab8ab019deb8a9f24a8d497a576e57c22e
GIT binary patch
literal 3504
zcmZ{ldsmco5XT?6iDCk&UDZI!D21fFMpBps6ozJHcU)kFP1vPn*Q_jstgPNifAtzV
zy^vl-w{z<B`8@kfJZ-0+XPlYuTz)h2n|)|)8Q7X6TaxWbXYyUrIK4>=NRo6V9Vwrf
zzBoO+R$CZ8dhCcEJCoMbQJ-$&+Tm@`N}*KNuorq6+JJh<p{*ApZJN)xR{o(TKcckb
zk&joWsN`IwTq!LiWN1%yprp85Tq%}oAl=Ch^!dulT&Y}~T}gch*J{=Jd~K)^IaD7R
z8T~IpY%BdHNdZ#sg3o30{qWLt@y@i5#pT(F<;q-Pd8&M^GGC~bD&?uOnePC+u(*hC
zTk<e+zFu|DX1#66LG-EenfevDJ*j(2<+&8Imca&;*7bxo$o$X1R}0m`N~VwfOK(qJ
zLSI{&U7N2KYV%8t`h5fHpQi5S^3(hmut)wPa;3IZtirv|hs@u%+N9q^pQ+c<X6!}X
zQ>oX|X6$9X0b%V(cY#_iRjOAfs+IbxL!96qcg6Y3C-IEAIa`lDkH9;iHfWOnV{Z4Q
z`sN%rV_H)mx72uNdI0i--rrDD-8osqgPHwv_~w0^z2)=UZ{)Yne13KFZ-KmT18)8Y
zGX2|3ch`3DD`<iDlg)g}dx-J=KFG5li@%2k*2ZQgsV>&f%(i8^xMciJcpqd9hcdkz
z+0%#K1CN~dqFbB!twFo9Fjjmt(|x1Jr|&+<qyK)mHgWHh_e0iW{l04>rZs&}b@796
z-`5<E!tGzoJJkm>O#9=RY;4@&2g!SP6L})T?fr!e55Ea?lkX+ZB%-~GnRg1_*saFB
z>i(X@qIY-1d-R{lZ03mkZzB5+p}&o6FQLDK+)09c^z42`|L&2xd2hju5%XU4bx17g
z*gz&t>VGG*IoHEX*SB)F?8~`G$Jr0#?ZxB!{sz(xH>Z2x9@>k#?=_Iqwfh~)+x-^h
z^|73;zu%R-z1XDtZOPkLGku(@+@psecg_7eO-=UK2Z`UtNBscQ2lYc^bY%}>;r|%2
z*f`$q%^}EL@;9JOKi}Xmr0qhMcPPVP#I~pN8Af&{#(I~!e#*|~2=qHM9f3Pr`7h*u
z3VsywXg>zmE{{5oBb(y~?s9)mLSpvy2RYQAgv@2Ea>SlOw#U$)MmE>2td~(_eZ=)~
z=BFWhv3B24J!bSQa?HqBW1P`*khYl7c!t6Lnb8?!_aJ7ZuHPx<@4e4M{=N22{k-2g
z#GQ?EvOjtBa~3(@&x`0{{<dTGV)@xSr`V%uWPi8$JvxUj9`kqw**c7wfuD!`?d#_q
z>iQ_V*XHuRSK-zx_pM)pzYck{zX8`SkA5y8n<L)(Eo3pj9lo{tMaW#nD*p`bw?y(A
z61}~P+}zuHVB&G#`^et+6TR((&qB^l{04db9bJaRV`mD;=D7?-opZ?I6VR{BVFB*3
zW^MD3m^CX$jaQJHYg_~qj~bVd^EH;xtx-H``~X=zYP^bU9`Sg$Wn|w?TfFTGvbb;O
zH%L8f736`;R)(vG?L%b!jq}^29(!Cx-T@h-O&{}!`|dFZ`TxD$8n))S)WF1}hB~r&
zqrYp&_NSkBs_WyO?wzxj8}GY*1Uc83+4T%_w#I7{%kQ##7jO45a&wKJfQf&feXBl2
z7V|gZUVH}kxEI=PKw{RW9J#+h-hhmYdT%3(NB>_UH~MdUzrR9Xhq@ts<6HDKvc4Yk
z-+_Mv>94Q7s_SR3QIGsj{*s=A{2RFy-FNkDgH90F3E!TRyO8bWF1UB~z<VKmyeoRS
kJHx|g7rL=|pWW!<@jiQyjT1l5{H*65Nc=DERJPW?0lO#@AOHXW

literal 0
HcmV?d00001

diff --git a/tests/shader/gen/prefix_root.hlsl b/tests/shader/gen/prefix_root.hlsl
new file mode 100644
index 0000000..2ad617c
--- /dev/null
+++ b/tests/shader/gen/prefix_root.hlsl
@@ -0,0 +1,80 @@
+struct Monoid
+{
+    uint element;
+};
+
+static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
+
+static const Monoid _133 = { 0u };
+
+RWByteAddressBuffer _42 : register(u0);
+
+static uint3 gl_LocalInvocationID;
+static uint3 gl_GlobalInvocationID;
+struct SPIRV_Cross_Input
+{
+    uint3 gl_LocalInvocationID : SV_GroupThreadID;
+    uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
+};
+
+groupshared Monoid sh_scratch[512];
+
+Monoid combine_monoid(Monoid a, Monoid b)
+{
+    Monoid _22 = { a.element + b.element };
+    return _22;
+}
+
+void comp_main()
+{
+    uint ix = gl_GlobalInvocationID.x * 8u;
+    Monoid _46;
+    _46.element = _42.Load(ix * 4 + 0);
+    Monoid local[8];
+    local[0].element = _46.element;
+    Monoid param_1;
+    for (uint i = 1u; i < 8u; i++)
+    {
+        Monoid param = local[i - 1u];
+        Monoid _71;
+        _71.element = _42.Load((ix + i) * 4 + 0);
+        param_1.element = _71.element;
+        local[i] = combine_monoid(param, param_1);
+    }
+    Monoid agg = local[7];
+    sh_scratch[gl_LocalInvocationID.x] = agg;
+    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    {
+        GroupMemoryBarrierWithGroupSync();
+        if (gl_LocalInvocationID.x >= uint(1 << int(i_1)))
+        {
+            Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))];
+            Monoid param_2 = other;
+            Monoid param_3 = agg;
+            agg = combine_monoid(param_2, param_3);
+        }
+        GroupMemoryBarrierWithGroupSync();
+        sh_scratch[gl_LocalInvocationID.x] = agg;
+    }
+    GroupMemoryBarrierWithGroupSync();
+    Monoid row = _133;
+    if (gl_LocalInvocationID.x > 0u)
+    {
+        row = sh_scratch[gl_LocalInvocationID.x - 1u];
+    }
+    for (uint i_2 = 0u; i_2 < 8u; i_2++)
+    {
+        Monoid param_4 = row;
+        Monoid param_5 = local[i_2];
+        Monoid m = combine_monoid(param_4, param_5);
+        _42.Store((ix + i_2) * 4 + 0, m.element);
+    }
+}
+
+[numthreads(512, 1, 1)]
+void main(SPIRV_Cross_Input stage_input)
+{
+    gl_LocalInvocationID = stage_input.gl_LocalInvocationID;
+    gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID;
+    comp_main();
+}
diff --git a/tests/shader/gen/prefix_root.msl b/tests/shader/gen/prefix_root.msl
new file mode 100644
index 0000000..ff02287
--- /dev/null
+++ b/tests/shader/gen/prefix_root.msl
@@ -0,0 +1,112 @@
+#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];
+    }
+};
+
+struct Monoid
+{
+    uint element;
+};
+
+struct Monoid_1
+{
+    uint element;
+};
+
+struct DataBuf
+{
+    Monoid_1 data[1];
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(512u, 1u, 1u);
+
+static inline __attribute__((always_inline))
+Monoid combine_monoid(thread const Monoid& a, thread const Monoid& b)
+{
+    return Monoid{ a.element + b.element };
+}
+
+kernel void main0(device DataBuf& _42 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
+{
+    threadgroup Monoid sh_scratch[512];
+    uint ix = gl_GlobalInvocationID.x * 8u;
+    spvUnsafeArray<Monoid, 8> local;
+    local[0].element = _42.data[ix].element;
+    Monoid param_1;
+    for (uint i = 1u; i < 8u; i++)
+    {
+        Monoid param = local[i - 1u];
+        param_1.element = _42.data[ix + i].element;
+        local[i] = combine_monoid(param, param_1);
+    }
+    Monoid agg = local[7];
+    sh_scratch[gl_LocalInvocationID.x] = agg;
+    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    {
+        threadgroup_barrier(mem_flags::mem_threadgroup);
+        if (gl_LocalInvocationID.x >= uint(1 << int(i_1)))
+        {
+            Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))];
+            Monoid param_2 = other;
+            Monoid param_3 = agg;
+            agg = combine_monoid(param_2, param_3);
+        }
+        threadgroup_barrier(mem_flags::mem_threadgroup);
+        sh_scratch[gl_LocalInvocationID.x] = agg;
+    }
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    Monoid row = Monoid{ 0u };
+    if (gl_LocalInvocationID.x > 0u)
+    {
+        row = sh_scratch[gl_LocalInvocationID.x - 1u];
+    }
+    for (uint i_2 = 0u; i_2 < 8u; i_2++)
+    {
+        Monoid param_4 = row;
+        Monoid param_5 = local[i_2];
+        Monoid m = combine_monoid(param_4, param_5);
+        _42.data[ix + i_2].element = m.element;
+    }
+}
+
diff --git a/tests/shader/gen/prefix_root.spv b/tests/shader/gen/prefix_root.spv
new file mode 100644
index 0000000000000000000000000000000000000000..70ba31c92ccab887aaf9be7a7fc71ee83cc25c2a
GIT binary patch
literal 4104
zcmZ{miE>m`5QcA@B!oo*0$M0w0tm<^D2s?7AR5phAnse3EF`0oOiU&$ii#NA7u*q1
zal^GNpTY7$d<wTJEx+&1ZMlr)r8?dJ-+lV@>C^WPg^LzUNRozRQqq!qm(-r=$we?p
zGC65Zbw}sE&XuFp-j!?Ct}<e3Qb+@hX-y_1P3T5suvi+>F$cL9nZ|ng=?E#)T&^kb
z519M}Xd)uNR!l+3uF6oQ)SHl@DaBApU%79vZ>S2>noI%rR0g|CLw#L?Y3!2G>d076
zbwzFEim}zJH~yDlZGp2UNfFVWjPB0t+33<?`Ihty{pGHXa;3Xi-Z^xz(o?LKDnmQB
zr?HL6Rp>HgIm3(4`I^N(=WA|EmVmbxtHo_&2hwwieR|>4*0WTDNBdfIDYYAuZRp|R
zNO3U3w_Jb=$!>75zn^tY$vxPk16`v%BgJaZK<zx<2j}Uu<htXedp(Rl_8P{nR0sM-
z(B8)pbbjwAFTjOl96VAvl;$~!UCUEQPGLvAr!QF3n4HPs4)$<nybJXXGDZIufPIWv
zjBZAp7h|G-=VJ;f=X4rxFXtmZf!k{^#rDKKugLu0X4pBH!Gt=a?c9U;v}cji{o4?q
zd3c=F8rR6wBq^3_r@o09mRDRq72Sq3Aj>k`itW<|o`#N`mw_(_o8KDrTdQm3H)hzL
zMLuKaB0k2?L+g`wKXp4|J#)bJu2yd?eqn~?J>Md9EoUL^zgT_=dNJZ2SEKEZoHdE<
zy`28#neE!B;W}*Vnn~0R8E<cPWqkBz57_!*-9Bs+Tl-y^PtJYrMmrnw)*<#?kc(b<
z=k9C#gPAX%zZcGXi~SB@dyc^U*exVj&U)9KK<u~q&Y_KMMwau}J94J1&+{tlm<wa9
z$;W&B9ma2{HG2QfpSi^g(SB<J>-W15*n8Yohy6C>{eBw)yWV#_u<M8Gu<v``@B5y|
z<2h`;lR0cY-}`+0={oGYuYW72c0R8}oMHR8i8I^3Hbj1$81a0h4QWR<(=~e}7xBxm
z<+c*-ysbc-LEkNX#w|jxMD)40wO)hRU*q;>d#%DY&c`^<xEe9mIQJJDr|o@sw|{Vt
zYtY`O`d9W`hhC5P=)WGVUmdk?z&6K^?B%<#5s~w(f09GI7BQD=wPWp#*scw{DZ|IJ
zvusX%wfZ9FCOCT)cnh{Uj%91MVjHssi9X$oZJ%oSYO}u`EN>5-L2>l|R_y4%`P|F?
zcOd$r*E=!}^>43tVmrgpYq4<~xLfDuHpJh5_cShQydB#b?TIysjkWLIfxT15xjV7*
z=e`3h=Qk+sRW5(8-a+iQ7u)^zAm+5TI}v%;@6NFC&iQ?a-=F-P-=FcZ{sFM-T;uoY
zA;fQ(apn{oqwRZaF6(;)?fKNs-F|ct@zLLn)~}B9_h6eN=B^K0t`~9c#9hdK#I@SL
zz#l*>{Yac`0K5KdB{=!m?@?^``<b&%ML&kvYx$$(ZAO<7`5%b!T^K~nBmWy`uAqIa
zPu~zCXMNgH&*Rwj^^CyDM?It1`Fg5g>yeLo#<1n%eL9G39{ISVL)hMtzPQK3*z(@f
zFj^eGC$Q!89YKr3cNE+E*XMUv9CQB^_7uc5`iwDF-aC#ysQ>Stp2k<-muKMQqlRPH
z=8f|`i|zT0b5HRx#66ucd#`rAcX%AJukk&3F5~R2>-EXy=XO81zGh>`;ZGoOug_!G
z*Z%^X{P)?+yofC~7jb4@M*BE3`d&ihtXn&Bzk*$#`&Bskxc}F%<^0}9e_qG7Kl-AF
zZ(z&E8N{A9_EqaUiMCFCZ=$V9?K}K7`W?ha|GQ}Y>ge@**yf0L_<d}-cEt0F-$Kmg
zTJ2c-0k&(S<`1#$w{@LHi#=lt`Xj_|oogcg<BT`{47NCC^%HDo)jax)Ifa&w{m)|0
z&f!n7ttaN@Gwid7eC+u-wt3=r<_qklNGoD|{H}b7ZLE)HIEVfUG2U3~5gTW3?X&$+
zM}1#o=j-_fEEjXEPcDCV_W7@*f#3Ir^#6Jj!1lpsBC?)+?dPPNJsI1*FGgGIH1u@D
o7;BAsF3I?anE`ffK4vCZK6*F{+ja7Lm<Q{*6p{a%J+-a%ABikjC;$Ke

literal 0
HcmV?d00001

diff --git a/tests/shader/gen/prefix_scan.hlsl b/tests/shader/gen/prefix_scan.hlsl
new file mode 100644
index 0000000..feeff2e
--- /dev/null
+++ b/tests/shader/gen/prefix_scan.hlsl
@@ -0,0 +1,92 @@
+struct Monoid
+{
+    uint element;
+};
+
+static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
+
+static const Monoid _133 = { 0u };
+
+RWByteAddressBuffer _42 : register(u0);
+RWByteAddressBuffer _143 : register(u1);
+
+static uint3 gl_WorkGroupID;
+static uint3 gl_LocalInvocationID;
+static uint3 gl_GlobalInvocationID;
+struct SPIRV_Cross_Input
+{
+    uint3 gl_WorkGroupID : SV_GroupID;
+    uint3 gl_LocalInvocationID : SV_GroupThreadID;
+    uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
+};
+
+groupshared Monoid sh_scratch[512];
+
+Monoid combine_monoid(Monoid a, Monoid b)
+{
+    Monoid _22 = { a.element + b.element };
+    return _22;
+}
+
+void comp_main()
+{
+    uint ix = gl_GlobalInvocationID.x * 8u;
+    Monoid _46;
+    _46.element = _42.Load(ix * 4 + 0);
+    Monoid local[8];
+    local[0].element = _46.element;
+    Monoid param_1;
+    for (uint i = 1u; i < 8u; i++)
+    {
+        Monoid param = local[i - 1u];
+        Monoid _71;
+        _71.element = _42.Load((ix + i) * 4 + 0);
+        param_1.element = _71.element;
+        local[i] = combine_monoid(param, param_1);
+    }
+    Monoid agg = local[7];
+    sh_scratch[gl_LocalInvocationID.x] = agg;
+    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    {
+        GroupMemoryBarrierWithGroupSync();
+        if (gl_LocalInvocationID.x >= uint(1 << int(i_1)))
+        {
+            Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))];
+            Monoid param_2 = other;
+            Monoid param_3 = agg;
+            agg = combine_monoid(param_2, param_3);
+        }
+        GroupMemoryBarrierWithGroupSync();
+        sh_scratch[gl_LocalInvocationID.x] = agg;
+    }
+    GroupMemoryBarrierWithGroupSync();
+    Monoid row = _133;
+    if (gl_WorkGroupID.x > 0u)
+    {
+        Monoid _148;
+        _148.element = _143.Load((gl_WorkGroupID.x - 1u) * 4 + 0);
+        row.element = _148.element;
+    }
+    if (gl_LocalInvocationID.x > 0u)
+    {
+        Monoid param_4 = row;
+        Monoid param_5 = sh_scratch[gl_LocalInvocationID.x - 1u];
+        row = combine_monoid(param_4, param_5);
+    }
+    for (uint i_2 = 0u; i_2 < 8u; i_2++)
+    {
+        Monoid param_6 = row;
+        Monoid param_7 = local[i_2];
+        Monoid m = combine_monoid(param_6, param_7);
+        _42.Store((ix + i_2) * 4 + 0, m.element);
+    }
+}
+
+[numthreads(512, 1, 1)]
+void main(SPIRV_Cross_Input stage_input)
+{
+    gl_WorkGroupID = stage_input.gl_WorkGroupID;
+    gl_LocalInvocationID = stage_input.gl_LocalInvocationID;
+    gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID;
+    comp_main();
+}
diff --git a/tests/shader/gen/prefix_scan.msl b/tests/shader/gen/prefix_scan.msl
new file mode 100644
index 0000000..c1efb22
--- /dev/null
+++ b/tests/shader/gen/prefix_scan.msl
@@ -0,0 +1,123 @@
+#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];
+    }
+};
+
+struct Monoid
+{
+    uint element;
+};
+
+struct Monoid_1
+{
+    uint element;
+};
+
+struct DataBuf
+{
+    Monoid_1 data[1];
+};
+
+struct ParentBuf
+{
+    Monoid_1 parent[1];
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(512u, 1u, 1u);
+
+static inline __attribute__((always_inline))
+Monoid combine_monoid(thread const Monoid& a, thread const Monoid& b)
+{
+    return Monoid{ a.element + b.element };
+}
+
+kernel void main0(device DataBuf& _42 [[buffer(0)]], device ParentBuf& _143 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+{
+    threadgroup Monoid sh_scratch[512];
+    uint ix = gl_GlobalInvocationID.x * 8u;
+    spvUnsafeArray<Monoid, 8> local;
+    local[0].element = _42.data[ix].element;
+    Monoid param_1;
+    for (uint i = 1u; i < 8u; i++)
+    {
+        Monoid param = local[i - 1u];
+        param_1.element = _42.data[ix + i].element;
+        local[i] = combine_monoid(param, param_1);
+    }
+    Monoid agg = local[7];
+    sh_scratch[gl_LocalInvocationID.x] = agg;
+    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    {
+        threadgroup_barrier(mem_flags::mem_threadgroup);
+        if (gl_LocalInvocationID.x >= uint(1 << int(i_1)))
+        {
+            Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))];
+            Monoid param_2 = other;
+            Monoid param_3 = agg;
+            agg = combine_monoid(param_2, param_3);
+        }
+        threadgroup_barrier(mem_flags::mem_threadgroup);
+        sh_scratch[gl_LocalInvocationID.x] = agg;
+    }
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    Monoid row = Monoid{ 0u };
+    if (gl_WorkGroupID.x > 0u)
+    {
+        row.element = _143.parent[gl_WorkGroupID.x - 1u].element;
+    }
+    if (gl_LocalInvocationID.x > 0u)
+    {
+        Monoid param_4 = row;
+        Monoid param_5 = sh_scratch[gl_LocalInvocationID.x - 1u];
+        row = combine_monoid(param_4, param_5);
+    }
+    for (uint i_2 = 0u; i_2 < 8u; i_2++)
+    {
+        Monoid param_6 = row;
+        Monoid param_7 = local[i_2];
+        Monoid m = combine_monoid(param_6, param_7);
+        _42.data[ix + i_2].element = m.element;
+    }
+}
+
diff --git a/tests/shader/gen/prefix_scan.spv b/tests/shader/gen/prefix_scan.spv
new file mode 100644
index 0000000000000000000000000000000000000000..d4216e95f91038753486e70cb2a8db3ac3beca00
GIT binary patch
literal 4736
zcmZ{mhjUa_6vm%y5`!QC5giMfC?L`l6)ac)(TE0#Vn+>|LK01~akH@(#29<;6;ZL%
z?22t1{}BHQ+l-FC-`jWeWgTzk%=y0WoO|!N=iaxI!lYyOPm=wT1Cqw%hom}YB$H5*
zWJ=PI`qs7$ZA(Tf-AhhccA^>6l0uqjP80L$@paf>v46<Gk=T{kI5wR%jLpDUHeJto
z1^%GPkC1vY^7G{sm8>fbmHN9AR;W)oQqnWfGuSg!L1{{+!n;a?o&7^S9fN7^{E<p|
zw5zhXI&<;p(xof@t1!2~d6T4wiKpN@Gkh4n|5o+J^c=kd9jybU&f>t@p_@xx#Y%r^
zXziLbw=S8BA3z+)_&j{RX1U|vn(LDJ@HNFsarNlt^jvaBH=5ph76=&Plkoj1u1i+q
zhl}OnV8+jzzzfNGc(J#adG*N^;E}$Lk*;#F($!Zzk9Wd(+AO)|IO$&35s$rw!KF%H
zPZ{rh+=kEZea8e|NXFph(k<jRr1#D~SC+~*ww6ny!#2#C9wVo{SmydXef3_|9IWlN
zl2u5as7m0vWLq`%&zZMR;DuxdyjpW1*$Ix`?4FQUm+Z~)R@U{6dhhxR*eUMnSh&NS
z<M0hw9cE74)xSA~RP$>v-+Ma-a|Cax(v;hO_dFttk7xYwj6aIc?-#wv@23{~*{A$|
za`)R0bN}!0?&q7H$=5as-@<I?i)WGJe$ANUC?bA$RjrOMAW6Ae3+JDladpl4)9}rh
zdoIj)6WGxVpN@|;4~2VH*S7}a_QJX9D>Ck##Cqn=#vJA!g*T?|e)<;7dX9wKt7^T~
z{5cs{_k8p4)in$0{mRwn<B!AK<9NLNQL`qwcdKT6QHGrxHJl8#u7k-sBh&58x=fGW
zw8O10=4}AGu(e;7#njyAa=dR(-8$rc*J{x#-<12BzcGvD*Y8I2-eSMaV9yb}7u?7K
zi<s}c`!V|+>yLw7e>OexeMY~0hjRCQ9PgTH&XfBVqhI?n+<;$1_q}W1rsuS8+H&5P
z`JT<;`}hr|9`BmIBj2$#%IDzc=D2ZxD<kfESY5-<tKsW&-1+|g<>wFAaDM~y@jGhx
zSPl2LE1$nT$K7v7j$7}}8os-R`y1wZtLTc~!O@uCiuZLEJ+Z&dnEDtw@?)@OtOYxp
zuG&Mj$X^InTSd0-dNJl(_Zv25o_D?kGv?mbdIDzu&ATMq>qM}54)Z+YQp{ZQ++S{<
z*!%MCf9IZ0!F%8OU)Xaw{#49i{4~6AebjzB*fsXBm*3n9OwF_Y!5Z>qm}@y#9CObE
zJ2&`Q8Q+?n<?J+8tuJ!UL9<uESAtz*OEzZ}*qoJE^yyr%eX6ck{l?b7)$M_AR3805
z9~}L6J@>Nzt(dXs^#z%R`={4y!M^e6wcNbZxm(}Oh1gWgJ<W?6F9KVmJ+UUax%S;V
zuy^`6cN;i=?u+4S>#(?2wfwz$2eIEJVE1dsT+`Yv#nhd@KI7*5&acA!oy))Tt1~_3
zUjuiZb2j3y#r&-_&o$-di2cr8%lfXzdp^DIZWF$UIgEGWjqBt5U0~OUch>_}>&AR{
z@(yeh=3Md5=$r9cFBWI(1J|CdA5A^>y8-NeKXJBc_!}{Mt$ru#HsA*^^<CunEev9=
zqy8&rF5w;4XKV;lvp#Xua}&6>o-&$x)H4Fk*HeL8k9yQI3RaKr>1MF&sK*`M0``uK
z#Xa(`ccSh+4ddkzyB(}%>^8hSVt0bQe`EeG%j4bO1)hpI$Cx?hs(Z)L2mSxu)7`{s
z`*IJOdepE5?7DHjd%>R1Jol7u!Q9g~X7BaR_YSvW_BDP_?#ncL>wIHs`FFbsUR$%d
zW9au|ajy@6$FbR1?Drs8{iE#7JOoxV$2ao`-r<`u_AsVq-QrmLNpS7jPob&D{XY#>
z^Y=b}>z)BWfi+>qt!p0sS<G1c);*VLxOC6y_~$Xt=v>b$H_zTgjryqZ1#s>Cy@;kB
z=YI*T7UzEj?C|`?UdGfszc|)@6<oXaYiR0m{@1~3o<H8p8(`01{oC+z*W8GI6Z7pj
zFY@2Ybn~}^<?+Vf2K&Zc$5^}pb$b$Z>y7Qedro8T;60Px-@y0q?_&<*AK;Da<L!J1
zc8&NA{0OYpg4r+myO?V^R~&P9ft?%Y+5?V0>;}v2LnHoUY%b=U$p0kMJ;z?KJo@k{
z*gm+9F>`j})not9z=!4d=V0rJ@9qn*IqI?Jmtfb4xBC@%0p|NNKi=xsU~?UwVITe*
z%=~z(a`Wu1eYQXPsP9{FzMk*kYVkW@Of7$R_W4i#Q$G#!f9d_<_Q5e3JC#1z&jWIJ
z3fR35#9OQ1)eOuWYmIsi%Jj&Y33qNj=U}*c^l%o~dFt)F2kSWmQ~!%S#n$>avB;sb

literal 0
HcmV?d00001