Compile shaders on Windows computer
diff --git a/tests/shader/gen/intersection_leaf.dxil b/tests/shader/gen/intersection_leaf.dxil
new file mode 100644
index 0000000..ca4e6ec
--- /dev/null
+++ b/tests/shader/gen/intersection_leaf.dxil
Binary files differ
diff --git a/tests/shader/gen/intersection_leaf.hlsl b/tests/shader/gen/intersection_leaf.hlsl
new file mode 100644
index 0000000..2b48f28
--- /dev/null
+++ b/tests/shader/gen/intersection_leaf.hlsl
@@ -0,0 +1,234 @@
+struct Bic
+{
+ uint a;
+ uint b;
+};
+
+struct Node
+{
+ uint node_type;
+ uint pad1;
+ uint pad2;
+ uint pad3;
+ float4 bbox;
+};
+
+static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
+
+static const Bic _76 = { 0u, 0u };
+
+ByteAddressBuffer _89 : register(t1);
+ByteAddressBuffer _167 : register(t2);
+ByteAddressBuffer _285 : register(t0);
+RWByteAddressBuffer _520 : register(u3);
+
+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 Bic sh_bic[1022];
+groupshared float4 sh_bbox[512];
+groupshared float4 sh_stack[512];
+groupshared uint sh_link[512];
+
+Bic bic_combine(Bic x, Bic y)
+{
+ uint m = min(x.b, y.a);
+ Bic _46 = { (x.a + y.a) - m, (x.b + y.b) - m };
+ return _46;
+}
+
+float4 bbox_intersect(float4 a, float4 b)
+{
+ return float4(max(a.xy, b.xy), min(a.zw, b.zw));
+}
+
+void comp_main()
+{
+ uint th = gl_LocalInvocationID.x;
+ Bic bic = _76;
+ if (th < gl_WorkGroupID.x)
+ {
+ Bic _93;
+ _93.a = _89.Load(th * 8 + 0);
+ _93.b = _89.Load(th * 8 + 4);
+ bic.a = _93.a;
+ bic.b = _93.b;
+ }
+ sh_bic[th] = bic;
+ for (uint i = 0u; i < 9u; i++)
+ {
+ GroupMemoryBarrierWithGroupSync();
+ uint other_ix = th + (1u << i);
+ if (other_ix < 512u)
+ {
+ Bic param = bic;
+ Bic param_1 = sh_bic[other_ix];
+ bic = bic_combine(param, param_1);
+ }
+ GroupMemoryBarrierWithGroupSync();
+ sh_bic[th] = bic;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ uint size = sh_bic[0].b;
+ uint bic_next_b = 0u;
+ if ((th + 1u) < 512u)
+ {
+ bic_next_b = sh_bic[th + 1u].b;
+ }
+ float4 bbox = float4(-1000000000.0f, -1000000000.0f, 1000000000.0f, 1000000000.0f);
+ if (bic.b > bic_next_b)
+ {
+ bbox = asfloat(_167.Load4(((((th * 512u) + bic.b) - bic_next_b) - 1u) * 16 + 0));
+ }
+ for (uint i_1 = 0u; i_1 < 9u; i_1++)
+ {
+ sh_bbox[th] = bbox;
+ GroupMemoryBarrierWithGroupSync();
+ if (th >= (1u << i_1))
+ {
+ float4 param_2 = sh_bbox[th - (1u << i_1)];
+ float4 param_3 = bbox;
+ bbox = bbox_intersect(param_2, param_3);
+ }
+ GroupMemoryBarrierWithGroupSync();
+ }
+ sh_bbox[th] = bbox;
+ GroupMemoryBarrierWithGroupSync();
+ uint sp = 511u - th;
+ uint ix = 0u;
+ for (uint i_2 = 0u; i_2 < 9u; i_2++)
+ {
+ uint probe = ix + (256u >> i_2);
+ if (sp < sh_bic[probe].b)
+ {
+ ix = probe;
+ }
+ }
+ uint b = sh_bic[ix].b;
+ if (sp < b)
+ {
+ bbox = asfloat(_167.Load4(((((ix * 512u) + b) - sp) - 1u) * 16 + 0));
+ if (ix > 0u)
+ {
+ float4 param_4 = sh_bbox[ix - 1u];
+ float4 param_5 = bbox;
+ bbox = bbox_intersect(param_4, param_5);
+ }
+ sh_stack[th] = bbox;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ Node _291;
+ _291.node_type = _285.Load(gl_GlobalInvocationID.x * 32 + 0);
+ _291.pad1 = _285.Load(gl_GlobalInvocationID.x * 32 + 4);
+ _291.pad2 = _285.Load(gl_GlobalInvocationID.x * 32 + 8);
+ _291.pad3 = _285.Load(gl_GlobalInvocationID.x * 32 + 12);
+ _291.bbox = asfloat(_285.Load4(gl_GlobalInvocationID.x * 32 + 16));
+ Node inp;
+ inp.node_type = _291.node_type;
+ inp.pad1 = _291.pad1;
+ inp.pad2 = _291.pad2;
+ inp.pad3 = _291.pad3;
+ inp.bbox = _291.bbox;
+ uint node_type = inp.node_type;
+ Bic _314 = { uint(node_type == 1u), uint(node_type == 0u) };
+ bic = _314;
+ sh_bic[th] = bic;
+ uint inbase = 0u;
+ for (uint i_3 = 0u; i_3 < 8u; i_3++)
+ {
+ uint outbase = 1024u - (1u << (9u - i_3));
+ GroupMemoryBarrierWithGroupSync();
+ if (th < (1u << (8u - i_3)))
+ {
+ Bic param_6 = sh_bic[inbase + (th * 2u)];
+ Bic param_7 = sh_bic[(inbase + (th * 2u)) + 1u];
+ sh_bic[outbase + th] = bic_combine(param_6, param_7);
+ }
+ inbase = outbase;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ ix = th;
+ bic = _76;
+ uint j = 0u;
+ while (j < 9u)
+ {
+ uint base = 1024u - (2u << (9u - j));
+ if (((ix >> j) & 1u) != 0u)
+ {
+ Bic param_8 = sh_bic[(base + (ix >> j)) - 1u];
+ Bic param_9 = bic;
+ Bic test = bic_combine(param_8, param_9);
+ if (test.b > 0u)
+ {
+ break;
+ }
+ bic = test;
+ ix -= (1u << j);
+ }
+ j++;
+ }
+ if (ix > 0u)
+ {
+ while (j > 0u)
+ {
+ j--;
+ uint base_1 = 1024u - (2u << (9u - j));
+ Bic param_10 = sh_bic[(base_1 + (ix >> j)) - 1u];
+ Bic param_11 = bic;
+ Bic test_1 = bic_combine(param_10, param_11);
+ if (test_1.b == 0u)
+ {
+ bic = test_1;
+ ix -= (1u << j);
+ }
+ }
+ }
+ uint _455;
+ if (ix > 0u)
+ {
+ _455 = ix - 1u;
+ }
+ else
+ {
+ _455 = 4294967295u - bic.a;
+ }
+ uint link = _455;
+ bbox = inp.bbox;
+ for (uint i_4 = 0u; i_4 < 9u; i_4++)
+ {
+ sh_link[th] = link;
+ sh_bbox[th] = bbox;
+ GroupMemoryBarrierWithGroupSync();
+ if (int(link) >= 0)
+ {
+ float4 param_12 = sh_bbox[link];
+ float4 param_13 = bbox;
+ bbox = bbox_intersect(param_12, param_13);
+ link = sh_link[link];
+ }
+ GroupMemoryBarrierWithGroupSync();
+ }
+ if (int(link + size) >= 0)
+ {
+ float4 param_14 = sh_stack[512u + link];
+ float4 param_15 = bbox;
+ bbox = bbox_intersect(param_14, param_15);
+ }
+ _520.Store4(gl_GlobalInvocationID.x * 16 + 0, asuint(bbox));
+}
+
+[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/intersection_leaf.msl b/tests/shader/gen/intersection_leaf.msl
new file mode 100644
index 0000000..0bf83f1
--- /dev/null
+++ b/tests/shader/gen/intersection_leaf.msl
@@ -0,0 +1,241 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct Bic
+{
+ uint a;
+ uint b;
+};
+
+struct Bic_1
+{
+ uint a;
+ uint b;
+};
+
+struct BicBuf
+{
+ Bic_1 bicbuf[1];
+};
+
+struct StackBuf
+{
+ float4 stack[1];
+};
+
+struct Node
+{
+ uint node_type;
+ uint pad1;
+ uint pad2;
+ uint pad3;
+ float4 bbox;
+};
+
+struct Node_1
+{
+ uint node_type;
+ uint pad1;
+ uint pad2;
+ uint pad3;
+ float4 bbox;
+};
+
+struct InBuf
+{
+ Node_1 inbuf[1];
+};
+
+struct OutBuf
+{
+ float4 outbuf[1];
+};
+
+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)
+{
+ uint m = min(x.b, y.a);
+ return Bic{ (x.a + y.a) - m, (x.b + y.b) - m };
+}
+
+static inline __attribute__((always_inline))
+float4 bbox_intersect(thread const float4& a, thread const float4& b)
+{
+ return float4(fast::max(a.xy, b.xy), fast::min(a.zw, b.zw));
+}
+
+kernel void main0(const device InBuf& _285 [[buffer(0)]], const device BicBuf& _89 [[buffer(1)]], const device StackBuf& _167 [[buffer(2)]], device OutBuf& _520 [[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[1022];
+ threadgroup float4 sh_bbox[512];
+ threadgroup float4 sh_stack[512];
+ threadgroup uint sh_link[512];
+ uint th = gl_LocalInvocationID.x;
+ Bic bic = Bic{ 0u, 0u };
+ if (th < gl_WorkGroupID.x)
+ {
+ bic.a = _89.bicbuf[th].a;
+ bic.b = _89.bicbuf[th].b;
+ }
+ sh_bic[th] = bic;
+ for (uint i = 0u; i < 9u; i++)
+ {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ uint other_ix = th + (1u << i);
+ if (other_ix < 512u)
+ {
+ Bic param = bic;
+ Bic param_1 = sh_bic[other_ix];
+ bic = bic_combine(param, param_1);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ sh_bic[th] = bic;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ uint size = sh_bic[0].b;
+ uint bic_next_b = 0u;
+ if ((th + 1u) < 512u)
+ {
+ bic_next_b = sh_bic[th + 1u].b;
+ }
+ float4 bbox = float4(-1000000000.0, -1000000000.0, 1000000000.0, 1000000000.0);
+ if (bic.b > bic_next_b)
+ {
+ bbox = _167.stack[(((th * 512u) + bic.b) - bic_next_b) - 1u];
+ }
+ for (uint i_1 = 0u; i_1 < 9u; i_1++)
+ {
+ sh_bbox[th] = bbox;
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ if (th >= (1u << i_1))
+ {
+ float4 param_2 = sh_bbox[th - (1u << i_1)];
+ float4 param_3 = bbox;
+ bbox = bbox_intersect(param_2, param_3);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ }
+ sh_bbox[th] = bbox;
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ uint sp = 511u - th;
+ uint ix = 0u;
+ for (uint i_2 = 0u; i_2 < 9u; i_2++)
+ {
+ uint probe = ix + (256u >> i_2);
+ if (sp < sh_bic[probe].b)
+ {
+ ix = probe;
+ }
+ }
+ uint b = sh_bic[ix].b;
+ if (sp < b)
+ {
+ bbox = _167.stack[(((ix * 512u) + b) - sp) - 1u];
+ if (ix > 0u)
+ {
+ float4 param_4 = sh_bbox[ix - 1u];
+ float4 param_5 = bbox;
+ bbox = bbox_intersect(param_4, param_5);
+ }
+ sh_stack[th] = bbox;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ Node inp;
+ inp.node_type = _285.inbuf[gl_GlobalInvocationID.x].node_type;
+ inp.pad1 = _285.inbuf[gl_GlobalInvocationID.x].pad1;
+ inp.pad2 = _285.inbuf[gl_GlobalInvocationID.x].pad2;
+ inp.pad3 = _285.inbuf[gl_GlobalInvocationID.x].pad3;
+ inp.bbox = _285.inbuf[gl_GlobalInvocationID.x].bbox;
+ uint node_type = inp.node_type;
+ bic = Bic{ uint(node_type == 1u), uint(node_type == 0u) };
+ sh_bic[th] = bic;
+ uint inbase = 0u;
+ for (uint i_3 = 0u; i_3 < 8u; i_3++)
+ {
+ uint outbase = 1024u - (1u << (9u - i_3));
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ if (th < (1u << (8u - i_3)))
+ {
+ Bic param_6 = sh_bic[inbase + (th * 2u)];
+ Bic param_7 = sh_bic[(inbase + (th * 2u)) + 1u];
+ sh_bic[outbase + th] = bic_combine(param_6, param_7);
+ }
+ inbase = outbase;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ ix = th;
+ bic = Bic{ 0u, 0u };
+ uint j = 0u;
+ while (j < 9u)
+ {
+ uint base = 1024u - (2u << (9u - j));
+ if (((ix >> j) & 1u) != 0u)
+ {
+ Bic param_8 = sh_bic[(base + (ix >> j)) - 1u];
+ Bic param_9 = bic;
+ Bic test = bic_combine(param_8, param_9);
+ if (test.b > 0u)
+ {
+ break;
+ }
+ bic = test;
+ ix -= (1u << j);
+ }
+ j++;
+ }
+ if (ix > 0u)
+ {
+ while (j > 0u)
+ {
+ j--;
+ uint base_1 = 1024u - (2u << (9u - j));
+ Bic param_10 = sh_bic[(base_1 + (ix >> j)) - 1u];
+ Bic param_11 = bic;
+ Bic test_1 = bic_combine(param_10, param_11);
+ if (test_1.b == 0u)
+ {
+ bic = test_1;
+ ix -= (1u << j);
+ }
+ }
+ }
+ uint _455;
+ if (ix > 0u)
+ {
+ _455 = ix - 1u;
+ }
+ else
+ {
+ _455 = 4294967295u - bic.a;
+ }
+ uint link = _455;
+ bbox = inp.bbox;
+ for (uint i_4 = 0u; i_4 < 9u; i_4++)
+ {
+ sh_link[th] = link;
+ sh_bbox[th] = bbox;
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ if (int(link) >= 0)
+ {
+ float4 param_12 = sh_bbox[link];
+ float4 param_13 = bbox;
+ bbox = bbox_intersect(param_12, param_13);
+ link = sh_link[link];
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ }
+ if (int(link + size) >= 0)
+ {
+ float4 param_14 = sh_stack[512u + link];
+ float4 param_15 = bbox;
+ bbox = bbox_intersect(param_14, param_15);
+ }
+ _520.outbuf[gl_GlobalInvocationID.x] = bbox;
+}
+
diff --git a/tests/shader/gen/intersection_leaf.spv b/tests/shader/gen/intersection_leaf.spv
new file mode 100644
index 0000000..001a7d7
--- /dev/null
+++ b/tests/shader/gen/intersection_leaf.spv
Binary files differ
diff --git a/tests/shader/gen/intersection_reduce.dxil b/tests/shader/gen/intersection_reduce.dxil
new file mode 100644
index 0000000..e8737e8
--- /dev/null
+++ b/tests/shader/gen/intersection_reduce.dxil
Binary files differ
diff --git a/tests/shader/gen/intersection_reduce.hlsl b/tests/shader/gen/intersection_reduce.hlsl
new file mode 100644
index 0000000..90f9540
--- /dev/null
+++ b/tests/shader/gen/intersection_reduce.hlsl
@@ -0,0 +1,152 @@
+struct Bic
+{
+ uint a;
+ uint b;
+};
+
+struct Node
+{
+ uint node_type;
+ uint pad1;
+ uint pad2;
+ uint pad3;
+ float4 bbox;
+};
+
+static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
+
+static const Bic _181 = { 0u, 0u };
+
+ByteAddressBuffer _82 : register(t0);
+RWByteAddressBuffer _165 : register(u1);
+RWByteAddressBuffer _273 : register(u2);
+
+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 Bic sh_bic[512];
+groupshared float4 sh_bbox[512];
+
+Bic bic_combine(Bic x, Bic y)
+{
+ uint m = min(x.b, y.a);
+ Bic _46 = { (x.a + y.a) - m, (x.b + y.b) - m };
+ return _46;
+}
+
+float4 bbox_intersect(float4 a, float4 b)
+{
+ return float4(max(a.xy, b.xy), min(a.zw, b.zw));
+}
+
+void comp_main()
+{
+ uint th = gl_LocalInvocationID.x;
+ Node _88;
+ _88.node_type = _82.Load(gl_GlobalInvocationID.x * 32 + 0);
+ _88.pad1 = _82.Load(gl_GlobalInvocationID.x * 32 + 4);
+ _88.pad2 = _82.Load(gl_GlobalInvocationID.x * 32 + 8);
+ _88.pad3 = _82.Load(gl_GlobalInvocationID.x * 32 + 12);
+ _88.bbox = asfloat(_82.Load4(gl_GlobalInvocationID.x * 32 + 16));
+ Node inp;
+ inp.node_type = _88.node_type;
+ inp.pad1 = _88.pad1;
+ inp.pad2 = _88.pad2;
+ inp.pad3 = _88.pad3;
+ inp.bbox = _88.bbox;
+ uint node_type = inp.node_type;
+ Bic _114 = { uint(node_type == 1u), uint(node_type == 0u) };
+ Bic bic = _114;
+ sh_bic[gl_LocalInvocationID.x] = bic;
+ for (uint i = 0u; i < 9u; i++)
+ {
+ GroupMemoryBarrierWithGroupSync();
+ uint other_ix = gl_LocalInvocationID.x + (1u << i);
+ if (other_ix < 512u)
+ {
+ Bic param = bic;
+ Bic param_1 = sh_bic[other_ix];
+ bic = bic_combine(param, param_1);
+ }
+ GroupMemoryBarrierWithGroupSync();
+ sh_bic[th] = bic;
+ }
+ if (th == 0u)
+ {
+ _165.Store(gl_WorkGroupID.x * 8 + 0, bic.a);
+ _165.Store(gl_WorkGroupID.x * 8 + 4, bic.b);
+ }
+ GroupMemoryBarrierWithGroupSync();
+ uint size = sh_bic[0].b;
+ bic = _181;
+ if ((th + 1u) < 512u)
+ {
+ bic = sh_bic[th + 1u];
+ }
+ bool _193 = inp.node_type == 0u;
+ bool _199;
+ if (_193)
+ {
+ _199 = bic.a == 0u;
+ }
+ else
+ {
+ _199 = _193;
+ }
+ if (_199)
+ {
+ uint out_ix = (size - bic.b) - 1u;
+ sh_bbox[out_ix] = inp.bbox;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ float4 bbox;
+ if (th < size)
+ {
+ bbox = sh_bbox[th];
+ }
+ for (uint i_1 = 0u; i_1 < 9u; i_1++)
+ {
+ bool _235 = th < size;
+ bool _242;
+ if (_235)
+ {
+ _242 = th >= (1u << i_1);
+ }
+ else
+ {
+ _242 = _235;
+ }
+ if (_242)
+ {
+ float4 param_2 = sh_bbox[th - (1u << i_1)];
+ float4 param_3 = bbox;
+ bbox = bbox_intersect(param_2, param_3);
+ }
+ GroupMemoryBarrierWithGroupSync();
+ if (th < size)
+ {
+ sh_bbox[th] = bbox;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ }
+ if (th < size)
+ {
+ _273.Store4(gl_GlobalInvocationID.x * 16 + 0, asuint(bbox));
+ }
+}
+
+[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/intersection_reduce.msl b/tests/shader/gen/intersection_reduce.msl
new file mode 100644
index 0000000..357a525
--- /dev/null
+++ b/tests/shader/gen/intersection_reduce.msl
@@ -0,0 +1,158 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct Bic
+{
+ uint a;
+ uint b;
+};
+
+struct Node
+{
+ uint node_type;
+ uint pad1;
+ uint pad2;
+ uint pad3;
+ float4 bbox;
+};
+
+struct Node_1
+{
+ uint node_type;
+ uint pad1;
+ uint pad2;
+ uint pad3;
+ float4 bbox;
+};
+
+struct InBuf
+{
+ Node_1 inbuf[1];
+};
+
+struct Bic_1
+{
+ uint a;
+ uint b;
+};
+
+struct BicBuf
+{
+ Bic_1 bicbuf[1];
+};
+
+struct StackBuf
+{
+ float4 stack[1];
+};
+
+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)
+{
+ uint m = min(x.b, y.a);
+ return Bic{ (x.a + y.a) - m, (x.b + y.b) - m };
+}
+
+static inline __attribute__((always_inline))
+float4 bbox_intersect(thread const float4& a, thread const float4& b)
+{
+ return float4(fast::max(a.xy, b.xy), fast::min(a.zw, b.zw));
+}
+
+kernel void main0(const device InBuf& _82 [[buffer(0)]], device BicBuf& _165 [[buffer(1)]], device StackBuf& _273 [[buffer(2)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+{
+ threadgroup Bic sh_bic[512];
+ threadgroup float4 sh_bbox[512];
+ uint th = gl_LocalInvocationID.x;
+ Node inp;
+ inp.node_type = _82.inbuf[gl_GlobalInvocationID.x].node_type;
+ inp.pad1 = _82.inbuf[gl_GlobalInvocationID.x].pad1;
+ inp.pad2 = _82.inbuf[gl_GlobalInvocationID.x].pad2;
+ inp.pad3 = _82.inbuf[gl_GlobalInvocationID.x].pad3;
+ inp.bbox = _82.inbuf[gl_GlobalInvocationID.x].bbox;
+ uint node_type = inp.node_type;
+ Bic bic = Bic{ uint(node_type == 1u), uint(node_type == 0u) };
+ sh_bic[gl_LocalInvocationID.x] = bic;
+ for (uint i = 0u; i < 9u; i++)
+ {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ uint other_ix = gl_LocalInvocationID.x + (1u << i);
+ if (other_ix < 512u)
+ {
+ Bic param = bic;
+ Bic param_1 = sh_bic[other_ix];
+ bic = bic_combine(param, param_1);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ sh_bic[th] = bic;
+ }
+ if (th == 0u)
+ {
+ _165.bicbuf[gl_WorkGroupID.x].a = bic.a;
+ _165.bicbuf[gl_WorkGroupID.x].b = bic.b;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ uint size = sh_bic[0].b;
+ bic = Bic{ 0u, 0u };
+ if ((th + 1u) < 512u)
+ {
+ bic = sh_bic[th + 1u];
+ }
+ bool _193 = inp.node_type == 0u;
+ bool _199;
+ if (_193)
+ {
+ _199 = bic.a == 0u;
+ }
+ else
+ {
+ _199 = _193;
+ }
+ if (_199)
+ {
+ uint out_ix = (size - bic.b) - 1u;
+ sh_bbox[out_ix] = inp.bbox;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ float4 bbox;
+ if (th < size)
+ {
+ bbox = sh_bbox[th];
+ }
+ for (uint i_1 = 0u; i_1 < 9u; i_1++)
+ {
+ bool _235 = th < size;
+ bool _242;
+ if (_235)
+ {
+ _242 = th >= (1u << i_1);
+ }
+ else
+ {
+ _242 = _235;
+ }
+ if (_242)
+ {
+ float4 param_2 = sh_bbox[th - (1u << i_1)];
+ float4 param_3 = bbox;
+ bbox = bbox_intersect(param_2, param_3);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ if (th < size)
+ {
+ sh_bbox[th] = bbox;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ }
+ if (th < size)
+ {
+ _273.stack[gl_GlobalInvocationID.x] = bbox;
+ }
+}
+
diff --git a/tests/shader/gen/intersection_reduce.spv b/tests/shader/gen/intersection_reduce.spv
new file mode 100644
index 0000000..587476d
--- /dev/null
+++ b/tests/shader/gen/intersection_reduce.spv
Binary files differ
diff --git a/tests/shader/gen/stack_leaf.dxil b/tests/shader/gen/stack_leaf.dxil
index 14658c3..16867ad 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_reduce.dxil b/tests/shader/gen/stack_reduce.dxil
index 50ceb82..440b017 100644
--- a/tests/shader/gen/stack_reduce.dxil
+++ b/tests/shader/gen/stack_reduce.dxil
Binary files differ
diff --git a/tests/shader/gen/union_leaf.dxil b/tests/shader/gen/union_leaf.dxil
new file mode 100644
index 0000000..03d8011
--- /dev/null
+++ b/tests/shader/gen/union_leaf.dxil
Binary files differ
diff --git a/tests/shader/gen/union_leaf.hlsl b/tests/shader/gen/union_leaf.hlsl
new file mode 100644
index 0000000..16ee5b1
--- /dev/null
+++ b/tests/shader/gen/union_leaf.hlsl
@@ -0,0 +1,223 @@
+struct Bic
+{
+ uint a;
+ uint b;
+};
+
+struct BicBbox
+{
+ Bic bic;
+ uint pad2;
+ uint pad3;
+ float4 bbox;
+};
+
+struct Node
+{
+ uint node_type;
+ uint pad1;
+ uint pad2;
+ uint pad3;
+ float4 bbox;
+};
+
+static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
+
+static const Bic _76 = { 0u, 0u };
+
+ByteAddressBuffer _94 : register(t1);
+ByteAddressBuffer _213 : register(t2);
+ByteAddressBuffer _249 : register(t0);
+RWByteAddressBuffer _492 : register(u3);
+
+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 Bic sh_bic[1022];
+groupshared float4 sh_bbox[1022];
+groupshared float4 sh_stack[512];
+
+Bic bic_combine(Bic x, Bic y)
+{
+ uint m = min(x.b, y.a);
+ Bic _46 = { (x.a + y.a) - m, (x.b + y.b) - m };
+ return _46;
+}
+
+float4 bbox_union(float4 a, float4 b)
+{
+ return float4(min(a.xy, b.xy), max(a.zw, b.zw));
+}
+
+void comp_main()
+{
+ uint th = gl_LocalInvocationID.x;
+ Bic bic = _76;
+ float4 bbox = float4(1000000000.0f, 1000000000.0f, -1000000000.0f, -1000000000.0f);
+ if (th < gl_WorkGroupID.x)
+ {
+ Bic _98;
+ _98.a = _94.Load(th * 32 + 0);
+ _98.b = _94.Load(th * 32 + 4);
+ bic.a = _98.a;
+ bic.b = _98.b;
+ bbox = asfloat(_94.Load4(th * 32 + 16));
+ }
+ sh_bic[th] = bic;
+ sh_bbox[th] = bbox;
+ for (uint i = 0u; i < 9u; i++)
+ {
+ GroupMemoryBarrierWithGroupSync();
+ uint other_ix = th + (1u << i);
+ if (other_ix < 512u)
+ {
+ Bic param = bic;
+ Bic param_1 = sh_bic[other_ix];
+ bic = bic_combine(param, param_1);
+ float4 param_2 = bbox;
+ float4 param_3 = sh_bbox[other_ix];
+ bbox = bbox_union(param_2, param_3);
+ }
+ GroupMemoryBarrierWithGroupSync();
+ sh_bic[th] = bic;
+ sh_bbox[th] = bbox;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ uint size = sh_bic[0].b;
+ uint sp = 511u - th;
+ uint ix = 0u;
+ for (uint i_1 = 0u; i_1 < 9u; i_1++)
+ {
+ uint probe = ix + (256u >> i_1);
+ if (sp < sh_bic[probe].b)
+ {
+ ix = probe;
+ }
+ }
+ uint b = sh_bic[ix].b;
+ if (sp < b)
+ {
+ float4 bbox_1 = asfloat(_213.Load4(((((ix * 512u) + b) - sp) - 1u) * 16 + 0));
+ if ((ix + 1u) < 512u)
+ {
+ float4 param_4 = bbox_1;
+ float4 param_5 = sh_bbox[ix + 1u];
+ bbox_1 = bbox_union(param_4, param_5);
+ }
+ sh_stack[th] = bbox_1;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ Node _255;
+ _255.node_type = _249.Load(gl_GlobalInvocationID.x * 32 + 0);
+ _255.pad1 = _249.Load(gl_GlobalInvocationID.x * 32 + 4);
+ _255.pad2 = _249.Load(gl_GlobalInvocationID.x * 32 + 8);
+ _255.pad3 = _249.Load(gl_GlobalInvocationID.x * 32 + 12);
+ _255.bbox = asfloat(_249.Load4(gl_GlobalInvocationID.x * 32 + 16));
+ Node inp;
+ inp.node_type = _255.node_type;
+ inp.pad1 = _255.pad1;
+ inp.pad2 = _255.pad2;
+ inp.pad3 = _255.pad3;
+ inp.bbox = _255.bbox;
+ uint node_type = inp.node_type;
+ Bic _277 = { uint(node_type == 1u), uint(node_type == 0u) };
+ bic = _277;
+ sh_bic[th] = bic;
+ sh_bbox[th] = inp.bbox;
+ uint inbase = 0u;
+ for (uint i_2 = 0u; i_2 < 8u; i_2++)
+ {
+ uint outbase = 1024u - (1u << (9u - i_2));
+ GroupMemoryBarrierWithGroupSync();
+ if (th < (1u << (8u - i_2)))
+ {
+ Bic param_6 = sh_bic[inbase + (th * 2u)];
+ Bic param_7 = sh_bic[(inbase + (th * 2u)) + 1u];
+ sh_bic[outbase + th] = bic_combine(param_6, param_7);
+ float4 param_8 = sh_bbox[inbase + (th * 2u)];
+ float4 param_9 = sh_bbox[(inbase + (th * 2u)) + 1u];
+ sh_bbox[outbase + th] = bbox_union(param_8, param_9);
+ }
+ inbase = outbase;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ ix = th;
+ bbox = inp.bbox;
+ bic = _76;
+ if (node_type == 1u)
+ {
+ uint j = 0u;
+ while (j < 9u)
+ {
+ uint base = 1024u - (2u << (9u - j));
+ if (((ix >> j) & 1u) != 0u)
+ {
+ Bic param_10 = sh_bic[(base + (ix >> j)) - 1u];
+ Bic param_11 = bic;
+ Bic test = bic_combine(param_10, param_11);
+ if (test.b > 0u)
+ {
+ break;
+ }
+ bic = test;
+ float4 param_12 = sh_bbox[(base + (ix >> j)) - 1u];
+ float4 param_13 = bbox;
+ bbox = bbox_union(param_12, param_13);
+ ix -= (1u << j);
+ }
+ j++;
+ }
+ if (ix > 0u)
+ {
+ while (j > 0u)
+ {
+ j--;
+ uint base_1 = 1024u - (2u << (9u - j));
+ Bic param_14 = sh_bic[(base_1 + (ix >> j)) - 1u];
+ Bic param_15 = bic;
+ Bic test_1 = bic_combine(param_14, param_15);
+ if (test_1.b == 0u)
+ {
+ bic = test_1;
+ float4 param_16 = sh_bbox[(base_1 + (ix >> j)) - 1u];
+ float4 param_17 = bbox;
+ bbox = bbox_union(param_16, param_17);
+ ix -= (1u << j);
+ }
+ }
+ }
+ bool _470 = ix == 0u;
+ bool _477;
+ if (_470)
+ {
+ _477 = bic.a < size;
+ }
+ else
+ {
+ _477 = _470;
+ }
+ if (_477)
+ {
+ float4 param_18 = sh_stack[511u - bic.a];
+ float4 param_19 = bbox;
+ bbox = bbox_union(param_18, param_19);
+ }
+ }
+ _492.Store4(gl_GlobalInvocationID.x * 16 + 0, asuint(bbox));
+}
+
+[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/union_leaf.msl b/tests/shader/gen/union_leaf.msl
new file mode 100644
index 0000000..d42b1eb
--- /dev/null
+++ b/tests/shader/gen/union_leaf.msl
@@ -0,0 +1,230 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct Bic
+{
+ uint a;
+ uint b;
+};
+
+struct Bic_1
+{
+ uint a;
+ uint b;
+};
+
+struct BicBbox
+{
+ Bic_1 bic;
+ uint pad2;
+ uint pad3;
+ float4 bbox;
+};
+
+struct BicBuf
+{
+ BicBbox bicbuf[1];
+};
+
+struct StackBuf
+{
+ float4 stack[1];
+};
+
+struct Node
+{
+ uint node_type;
+ uint pad1;
+ uint pad2;
+ uint pad3;
+ float4 bbox;
+};
+
+struct Node_1
+{
+ uint node_type;
+ uint pad1;
+ uint pad2;
+ uint pad3;
+ float4 bbox;
+};
+
+struct InBuf
+{
+ Node_1 inbuf[1];
+};
+
+struct OutBuf
+{
+ float4 outbuf[1];
+};
+
+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)
+{
+ uint m = min(x.b, y.a);
+ return Bic{ (x.a + y.a) - m, (x.b + y.b) - m };
+}
+
+static inline __attribute__((always_inline))
+float4 bbox_union(thread const float4& a, thread const float4& b)
+{
+ return float4(fast::min(a.xy, b.xy), fast::max(a.zw, b.zw));
+}
+
+kernel void main0(const device InBuf& _249 [[buffer(0)]], const device BicBuf& _94 [[buffer(1)]], const device StackBuf& _213 [[buffer(2)]], device OutBuf& _492 [[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[1022];
+ threadgroup float4 sh_bbox[1022];
+ threadgroup float4 sh_stack[512];
+ uint th = gl_LocalInvocationID.x;
+ Bic bic = Bic{ 0u, 0u };
+ float4 bbox = float4(1000000000.0, 1000000000.0, -1000000000.0, -1000000000.0);
+ if (th < gl_WorkGroupID.x)
+ {
+ bic.a = _94.bicbuf[th].bic.a;
+ bic.b = _94.bicbuf[th].bic.b;
+ bbox = _94.bicbuf[th].bbox;
+ }
+ sh_bic[th] = bic;
+ sh_bbox[th] = bbox;
+ for (uint i = 0u; i < 9u; i++)
+ {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ uint other_ix = th + (1u << i);
+ if (other_ix < 512u)
+ {
+ Bic param = bic;
+ Bic param_1 = sh_bic[other_ix];
+ bic = bic_combine(param, param_1);
+ float4 param_2 = bbox;
+ float4 param_3 = sh_bbox[other_ix];
+ bbox = bbox_union(param_2, param_3);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ sh_bic[th] = bic;
+ sh_bbox[th] = bbox;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ uint size = sh_bic[0].b;
+ uint sp = 511u - th;
+ uint ix = 0u;
+ for (uint i_1 = 0u; i_1 < 9u; i_1++)
+ {
+ uint probe = ix + (256u >> i_1);
+ if (sp < sh_bic[probe].b)
+ {
+ ix = probe;
+ }
+ }
+ uint b = sh_bic[ix].b;
+ if (sp < b)
+ {
+ float4 bbox_1 = _213.stack[(((ix * 512u) + b) - sp) - 1u];
+ if ((ix + 1u) < 512u)
+ {
+ float4 param_4 = bbox_1;
+ float4 param_5 = sh_bbox[ix + 1u];
+ bbox_1 = bbox_union(param_4, param_5);
+ }
+ sh_stack[th] = bbox_1;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ Node inp;
+ inp.node_type = _249.inbuf[gl_GlobalInvocationID.x].node_type;
+ inp.pad1 = _249.inbuf[gl_GlobalInvocationID.x].pad1;
+ inp.pad2 = _249.inbuf[gl_GlobalInvocationID.x].pad2;
+ inp.pad3 = _249.inbuf[gl_GlobalInvocationID.x].pad3;
+ inp.bbox = _249.inbuf[gl_GlobalInvocationID.x].bbox;
+ uint node_type = inp.node_type;
+ bic = Bic{ uint(node_type == 1u), uint(node_type == 0u) };
+ sh_bic[th] = bic;
+ sh_bbox[th] = inp.bbox;
+ uint inbase = 0u;
+ for (uint i_2 = 0u; i_2 < 8u; i_2++)
+ {
+ uint outbase = 1024u - (1u << (9u - i_2));
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ if (th < (1u << (8u - i_2)))
+ {
+ Bic param_6 = sh_bic[inbase + (th * 2u)];
+ Bic param_7 = sh_bic[(inbase + (th * 2u)) + 1u];
+ sh_bic[outbase + th] = bic_combine(param_6, param_7);
+ float4 param_8 = sh_bbox[inbase + (th * 2u)];
+ float4 param_9 = sh_bbox[(inbase + (th * 2u)) + 1u];
+ sh_bbox[outbase + th] = bbox_union(param_8, param_9);
+ }
+ inbase = outbase;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ ix = th;
+ bbox = inp.bbox;
+ bic = Bic{ 0u, 0u };
+ if (node_type == 1u)
+ {
+ uint j = 0u;
+ while (j < 9u)
+ {
+ uint base = 1024u - (2u << (9u - j));
+ if (((ix >> j) & 1u) != 0u)
+ {
+ Bic param_10 = sh_bic[(base + (ix >> j)) - 1u];
+ Bic param_11 = bic;
+ Bic test = bic_combine(param_10, param_11);
+ if (test.b > 0u)
+ {
+ break;
+ }
+ bic = test;
+ float4 param_12 = sh_bbox[(base + (ix >> j)) - 1u];
+ float4 param_13 = bbox;
+ bbox = bbox_union(param_12, param_13);
+ ix -= (1u << j);
+ }
+ j++;
+ }
+ if (ix > 0u)
+ {
+ while (j > 0u)
+ {
+ j--;
+ uint base_1 = 1024u - (2u << (9u - j));
+ Bic param_14 = sh_bic[(base_1 + (ix >> j)) - 1u];
+ Bic param_15 = bic;
+ Bic test_1 = bic_combine(param_14, param_15);
+ if (test_1.b == 0u)
+ {
+ bic = test_1;
+ float4 param_16 = sh_bbox[(base_1 + (ix >> j)) - 1u];
+ float4 param_17 = bbox;
+ bbox = bbox_union(param_16, param_17);
+ ix -= (1u << j);
+ }
+ }
+ }
+ bool _470 = ix == 0u;
+ bool _477;
+ if (_470)
+ {
+ _477 = bic.a < size;
+ }
+ else
+ {
+ _477 = _470;
+ }
+ if (_477)
+ {
+ float4 param_18 = sh_stack[511u - bic.a];
+ float4 param_19 = bbox;
+ bbox = bbox_union(param_18, param_19);
+ }
+ }
+ _492.outbuf[gl_GlobalInvocationID.x] = bbox;
+}
+
diff --git a/tests/shader/gen/union_leaf.spv b/tests/shader/gen/union_leaf.spv
new file mode 100644
index 0000000..e81edb2
--- /dev/null
+++ b/tests/shader/gen/union_leaf.spv
Binary files differ
diff --git a/tests/shader/gen/union_reduce.dxil b/tests/shader/gen/union_reduce.dxil
new file mode 100644
index 0000000..22fbb6a
--- /dev/null
+++ b/tests/shader/gen/union_reduce.dxil
Binary files differ
diff --git a/tests/shader/gen/union_reduce.hlsl b/tests/shader/gen/union_reduce.hlsl
new file mode 100644
index 0000000..696d217
--- /dev/null
+++ b/tests/shader/gen/union_reduce.hlsl
@@ -0,0 +1,135 @@
+struct Bic
+{
+ uint a;
+ uint b;
+};
+
+struct Node
+{
+ uint node_type;
+ uint pad1;
+ uint pad2;
+ uint pad3;
+ float4 bbox;
+};
+
+struct BicBbox
+{
+ Bic bic;
+ uint pad2;
+ uint pad3;
+ float4 bbox;
+};
+
+static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
+
+static const Bic _219 = { 0u, 0u };
+
+ByteAddressBuffer _74 : register(t0);
+RWByteAddressBuffer _193 : register(u1);
+RWByteAddressBuffer _255 : register(u2);
+
+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 Bic sh_bic[512];
+groupshared float4 sh_bbox[512];
+
+Bic bic_combine(Bic x, Bic y)
+{
+ uint m = min(x.b, y.a);
+ Bic _46 = { (x.a + y.a) - m, (x.b + y.b) - m };
+ return _46;
+}
+
+float4 bbox_union(float4 a, float4 b)
+{
+ return float4(min(a.xy, b.xy), max(a.zw, b.zw));
+}
+
+void comp_main()
+{
+ Node _84;
+ _84.node_type = _74.Load(gl_GlobalInvocationID.x * 32 + 0);
+ _84.pad1 = _74.Load(gl_GlobalInvocationID.x * 32 + 4);
+ _84.pad2 = _74.Load(gl_GlobalInvocationID.x * 32 + 8);
+ _84.pad3 = _74.Load(gl_GlobalInvocationID.x * 32 + 12);
+ _84.bbox = asfloat(_74.Load4(gl_GlobalInvocationID.x * 32 + 16));
+ Node inp;
+ inp.node_type = _84.node_type;
+ inp.pad1 = _84.pad1;
+ inp.pad2 = _84.pad2;
+ inp.pad3 = _84.pad3;
+ inp.bbox = _84.bbox;
+ uint node_type = inp.node_type;
+ float4 bbox = inp.bbox;
+ Bic _113 = { uint(node_type == 1u), uint(node_type == 0u) };
+ Bic bic = _113;
+ sh_bic[gl_LocalInvocationID.x] = bic;
+ sh_bbox[gl_LocalInvocationID.x] = bbox;
+ for (uint i = 0u; i < 9u; i++)
+ {
+ GroupMemoryBarrierWithGroupSync();
+ uint other_ix = gl_LocalInvocationID.x + (1u << i);
+ if (other_ix < 512u)
+ {
+ Bic param = bic;
+ Bic param_1 = sh_bic[other_ix];
+ bic = bic_combine(param, param_1);
+ float4 param_2 = bbox;
+ float4 param_3 = sh_bbox[other_ix];
+ bbox = bbox_union(param_2, param_3);
+ }
+ GroupMemoryBarrierWithGroupSync();
+ sh_bic[gl_LocalInvocationID.x] = bic;
+ sh_bbox[gl_LocalInvocationID.x] = bbox;
+ }
+ if (gl_LocalInvocationID.x == 0u)
+ {
+ BicBbox _187 = { bic, 0u, 0u, bbox };
+ BicBbox bic_bbox = _187;
+ _193.Store(gl_WorkGroupID.x * 32 + 0, bic_bbox.bic.a);
+ _193.Store(gl_WorkGroupID.x * 32 + 4, bic_bbox.bic.b);
+ _193.Store(gl_WorkGroupID.x * 32 + 8, bic_bbox.pad2);
+ _193.Store(gl_WorkGroupID.x * 32 + 12, bic_bbox.pad3);
+ _193.Store4(gl_WorkGroupID.x * 32 + 16, asuint(bic_bbox.bbox));
+ }
+ GroupMemoryBarrierWithGroupSync();
+ uint size = sh_bic[0].b;
+ bic = _219;
+ if ((gl_LocalInvocationID.x + 1u) < 512u)
+ {
+ bic = sh_bic[gl_LocalInvocationID.x + 1u];
+ }
+ bool _233 = inp.node_type == 0u;
+ bool _239;
+ if (_233)
+ {
+ _239 = bic.a == 0u;
+ }
+ else
+ {
+ _239 = _233;
+ }
+ if (_239)
+ {
+ uint out_ix = (((gl_WorkGroupID.x * 512u) + size) - bic.b) - 1u;
+ _255.Store4(out_ix * 16 + 0, asuint(sh_bbox[gl_LocalInvocationID.x]));
+ }
+}
+
+[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/union_reduce.msl b/tests/shader/gen/union_reduce.msl
new file mode 100644
index 0000000..bc87386
--- /dev/null
+++ b/tests/shader/gen/union_reduce.msl
@@ -0,0 +1,148 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct Bic
+{
+ uint a;
+ uint b;
+};
+
+struct Node
+{
+ uint node_type;
+ uint pad1;
+ uint pad2;
+ uint pad3;
+ float4 bbox;
+};
+
+struct Node_1
+{
+ uint node_type;
+ uint pad1;
+ uint pad2;
+ uint pad3;
+ float4 bbox;
+};
+
+struct InBuf
+{
+ Node_1 inbuf[1];
+};
+
+struct BicBbox
+{
+ Bic bic;
+ uint pad2;
+ uint pad3;
+ float4 bbox;
+};
+
+struct Bic_1
+{
+ uint a;
+ uint b;
+};
+
+struct BicBbox_1
+{
+ Bic_1 bic;
+ uint pad2;
+ uint pad3;
+ float4 bbox;
+};
+
+struct BicBuf
+{
+ BicBbox_1 bicbuf[1];
+};
+
+struct StackBuf
+{
+ float4 stack[1];
+};
+
+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)
+{
+ uint m = min(x.b, y.a);
+ return Bic{ (x.a + y.a) - m, (x.b + y.b) - m };
+}
+
+static inline __attribute__((always_inline))
+float4 bbox_union(thread const float4& a, thread const float4& b)
+{
+ return float4(fast::min(a.xy, b.xy), fast::max(a.zw, b.zw));
+}
+
+kernel void main0(const device InBuf& _74 [[buffer(0)]], device BicBuf& _193 [[buffer(1)]], device StackBuf& _255 [[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[512];
+ threadgroup float4 sh_bbox[512];
+ Node inp;
+ inp.node_type = _74.inbuf[gl_GlobalInvocationID.x].node_type;
+ inp.pad1 = _74.inbuf[gl_GlobalInvocationID.x].pad1;
+ inp.pad2 = _74.inbuf[gl_GlobalInvocationID.x].pad2;
+ inp.pad3 = _74.inbuf[gl_GlobalInvocationID.x].pad3;
+ inp.bbox = _74.inbuf[gl_GlobalInvocationID.x].bbox;
+ uint node_type = inp.node_type;
+ float4 bbox = inp.bbox;
+ Bic bic = Bic{ uint(node_type == 1u), uint(node_type == 0u) };
+ sh_bic[gl_LocalInvocationID.x] = bic;
+ sh_bbox[gl_LocalInvocationID.x] = bbox;
+ for (uint i = 0u; i < 9u; i++)
+ {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ uint other_ix = gl_LocalInvocationID.x + (1u << i);
+ if (other_ix < 512u)
+ {
+ Bic param = bic;
+ Bic param_1 = sh_bic[other_ix];
+ bic = bic_combine(param, param_1);
+ float4 param_2 = bbox;
+ float4 param_3 = sh_bbox[other_ix];
+ bbox = bbox_union(param_2, param_3);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ sh_bic[gl_LocalInvocationID.x] = bic;
+ sh_bbox[gl_LocalInvocationID.x] = bbox;
+ }
+ if (gl_LocalInvocationID.x == 0u)
+ {
+ BicBbox bic_bbox = BicBbox{ bic, 0u, 0u, bbox };
+ _193.bicbuf[gl_WorkGroupID.x].bic.a = bic_bbox.bic.a;
+ _193.bicbuf[gl_WorkGroupID.x].bic.b = bic_bbox.bic.b;
+ _193.bicbuf[gl_WorkGroupID.x].pad2 = bic_bbox.pad2;
+ _193.bicbuf[gl_WorkGroupID.x].pad3 = bic_bbox.pad3;
+ _193.bicbuf[gl_WorkGroupID.x].bbox = bic_bbox.bbox;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ uint size = sh_bic[0].b;
+ bic = Bic{ 0u, 0u };
+ if ((gl_LocalInvocationID.x + 1u) < 512u)
+ {
+ bic = sh_bic[gl_LocalInvocationID.x + 1u];
+ }
+ bool _233 = inp.node_type == 0u;
+ bool _239;
+ if (_233)
+ {
+ _239 = bic.a == 0u;
+ }
+ else
+ {
+ _239 = _233;
+ }
+ if (_239)
+ {
+ uint out_ix = (((gl_WorkGroupID.x * 512u) + size) - bic.b) - 1u;
+ _255.stack[out_ix] = sh_bbox[gl_LocalInvocationID.x];
+ }
+}
+
diff --git a/tests/shader/gen/union_reduce.spv b/tests/shader/gen/union_reduce.spv
new file mode 100644
index 0000000..a049391
--- /dev/null
+++ b/tests/shader/gen/union_reduce.spv
Binary files differ
