Support XFB in MoltenVK

main.cpp

@@ -681,6 +681,7 @@ struct CLIArguments
 	bool msl_readwrite_texture_fences = true;
 	bool msl_agx_manual_cube_grad_fixup = false;
 	const char *msl_combined_sampler_suffix = nullptr;
+	CompilerMSL::Options::PrimitiveType msl_xfb_primitive_type = CompilerMSL::Options::PrimitiveType::Dynamic;
 	bool glsl_emit_push_constant_as_ubo = false;
 	bool glsl_emit_ubo_as_plain_uniforms = false;
 	bool glsl_force_flattened_io_blocks = false;
@@ -968,7 +969,11 @@ static void print_help_msl()
 	                "\t\tAll released Apple Silicon GPUs to date ignore one of the three partial derivatives\n"
 	                "\t\tbased on the selected major axis, and expect the remaining derivatives to be\n"
 	                "\t\tpartially transformed. This fixup gives correct results on Apple Silicon.\n"
-	                "\t[--msl-combined-sampler-suffix <suffix>]:\n\t\tUses a custom suffix for combined samplers.\n");
+	                "\t[--msl-combined-sampler-suffix <suffix>]:\n\t\tUses a custom suffix for combined samplers.\n"
+	                "\t[--msl-xfb-primitive-type <type>]:\n\t\tGenerates code in a vertex shader to capture primitives of the\n\t\t"
+	                "specified type for transform feedback. <type> may be one of dynamic,\n\t\t"
+	                "point-list, line-list, line-strip, triangle-list, triangle-strip, or\n\t\t"
+	                "triangle-fan. The default is \"dynamic\".\n");
 	// clang-format on
 }
 
@@ -1243,6 +1248,7 @@ static string compile_iteration(const CLIArguments &args, std::vector<uint32_t>
 		msl_opts.manual_helper_invocation_updates = args.msl_manual_helper_invocation_updates;
 		msl_opts.check_discarded_frag_stores = args.msl_check_discarded_frag_stores;
 		msl_opts.sample_dref_lod_array_as_grad = args.msl_sample_dref_lod_array_as_grad;
+		msl_opts.xfb_primitive_type = args.msl_xfb_primitive_type;
 		msl_opts.ios_support_base_vertex_instance = true;
 		msl_opts.runtime_array_rich_descriptor = args.msl_runtime_array_rich_descriptor;
 		msl_opts.replace_recursive_inputs = args.msl_replace_recursive_inputs;
@@ -1811,6 +1817,25 @@ static int main_inner(int argc, char *argv[])
 	        [&args](CLIParser &) { args.msl_runtime_array_rich_descriptor = true; });
 	cbs.add("--msl-replace-recursive-inputs",
 	        [&args](CLIParser &) { args.msl_replace_recursive_inputs = true; });
+	cbs.add("--msl-xfb-primitive-type",
+	        [&args](CLIParser &parser)
+	        {
+		        const char *type = parser.next_value_string("dynamic");
+		        if (strcmp(type, "dynamic") == 0)
+			        args.msl_xfb_primitive_type = CompilerMSL::Options::PrimitiveType::Dynamic;
+		        else if (strcmp(type, "point-list") == 0)
+			        args.msl_xfb_primitive_type = CompilerMSL::Options::PrimitiveType::PointList;
+		        else if (strcmp(type, "line-list") == 0)
+			        args.msl_xfb_primitive_type = CompilerMSL::Options::PrimitiveType::LineList;
+		        else if (strcmp(type, "line-strip") == 0)
+			        args.msl_xfb_primitive_type = CompilerMSL::Options::PrimitiveType::LineStrip;
+		        else if (strcmp(type, "triangle-list") == 0)
+			        args.msl_xfb_primitive_type = CompilerMSL::Options::PrimitiveType::TriangleList;
+		        else if (strcmp(type, "triangle-strip") == 0)
+			        args.msl_xfb_primitive_type = CompilerMSL::Options::PrimitiveType::TriangleStrip;
+		        else if (strcmp(type, "triangle-fan") == 0)
+			        args.msl_xfb_primitive_type = CompilerMSL::Options::PrimitiveType::TriangleFan;
+	        });
 	cbs.add("--extension", [&args](CLIParser &parser) { args.extensions.push_back(parser.next_string()); });
 	cbs.add("--rename-entry-point", [&args](CLIParser &parser) {
 		auto old_name = parser.next_string();

reference/opt/shaders-msl/vert/transform-feedback-decorations.xfb-line-list.vert

@@ -0,0 +1,62 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct VertOut
+{
+    float4 vBar;
+};
+
+struct spvXfbBuffer1
+{
+    char _m0_pad[4];
+    packed_float4 gl_Position;
+};
+
+struct spvXfbBuffer2
+{
+    char _m0_pad[16];
+    float4 vFoo;
+};
+
+struct spvXfbBuffer3
+{
+    float4 vBar;
+};
+
+kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 spvStageInputSize [[grid_size]], device atomic_uint* spvXfbCounter1 [[buffer(17)]], device spvXfbBuffer1* spvXfb1 [[buffer(13)]], device atomic_uint* spvXfbCounter2 [[buffer(18)]], device spvXfbBuffer2* spvXfb2 [[buffer(14)]], device atomic_uint* spvXfbCounter3 [[buffer(19)]], device spvXfbBuffer3* spvXfb3 [[buffer(15)]])
+{
+    spvXfbBuffer1 spvXfbOutput1 = {};
+    spvXfbBuffer2 spvXfbOutput2 = {};
+    spvXfbBuffer3 spvXfbOutput3 = {};
+    VertOut _20 = {};
+    if (any(gl_GlobalInvocationID >= spvStageInputSize))
+        return;
+    spvXfbOutput1.gl_Position = float4(1.0);
+    spvXfbOutput2.vFoo = float4(3.0);
+    _20.vBar = float4(5.0);
+    spvXfbOutput3.vBar = _20.vBar;
+    uint spvXfbIndex = gl_GlobalInvocationID.y * (spvStageInputSize.x & ~1u) + gl_GlobalInvocationID.x;
+    uint spvInitOffset1 = atomic_load_explicit(spvXfbCounter1, memory_order_relaxed);
+    spvXfb1 = reinterpret_cast<device spvXfbBuffer1*>(reinterpret_cast<device char*>(spvXfb1) + spvInitOffset1);
+    if ((gl_GlobalInvocationID.x & 1) || gl_GlobalInvocationID.x < spvStageInputSize.x - 1u)
+        spvXfb1[spvXfbIndex] = spvXfbOutput1;
+    uint spvInitOffset2 = atomic_load_explicit(spvXfbCounter2, memory_order_relaxed);
+    spvXfb2 = reinterpret_cast<device spvXfbBuffer2*>(reinterpret_cast<device char*>(spvXfb2) + spvInitOffset2);
+    if ((gl_GlobalInvocationID.x & 1) || gl_GlobalInvocationID.x < spvStageInputSize.x - 1u)
+        spvXfb2[spvXfbIndex] = spvXfbOutput2;
+    uint spvInitOffset3 = atomic_load_explicit(spvXfbCounter3, memory_order_relaxed);
+    spvXfb3 = reinterpret_cast<device spvXfbBuffer3*>(reinterpret_cast<device char*>(spvXfb3) + spvInitOffset3);
+    if ((gl_GlobalInvocationID.x & 1) || gl_GlobalInvocationID.x < spvStageInputSize.x - 1u)
+        spvXfb3[spvXfbIndex] = spvXfbOutput3;
+    threadgroup_barrier(mem_flags::mem_device);
+    if (all(gl_GlobalInvocationID.xy == 0))
+    {
+        uint spvWritten = (spvStageInputSize.x & ~1u) * spvStageInputSize.y;
+        atomic_store_explicit(spvXfbCounter1, spvInitOffset1 + sizeof(*spvXfb1) * spvWritten, memory_order_relaxed);
+        atomic_store_explicit(spvXfbCounter2, spvInitOffset2 + sizeof(*spvXfb2) * spvWritten, memory_order_relaxed);
+        atomic_store_explicit(spvXfbCounter3, spvInitOffset3 + sizeof(*spvXfb3) * spvWritten, memory_order_relaxed);
+    }
+}
+

reference/opt/shaders-msl/vert/transform-feedback-decorations.xfb-line-strip.vert

@@ -0,0 +1,68 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct VertOut
+{
+    float4 vBar;
+};
+
+struct spvXfbBuffer1
+{
+    char _m0_pad[4];
+    packed_float4 gl_Position;
+};
+
+struct spvXfbBuffer2
+{
+    char _m0_pad[16];
+    float4 vFoo;
+};
+
+struct spvXfbBuffer3
+{
+    float4 vBar;
+};
+
+kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 spvStageInputSize [[grid_size]], device atomic_uint* spvXfbCounter1 [[buffer(17)]], device spvXfbBuffer1* spvXfb1 [[buffer(13)]], device atomic_uint* spvXfbCounter2 [[buffer(18)]], device spvXfbBuffer2* spvXfb2 [[buffer(14)]], device atomic_uint* spvXfbCounter3 [[buffer(19)]], device spvXfbBuffer3* spvXfb3 [[buffer(15)]])
+{
+    spvXfbBuffer1 spvXfbOutput1 = {};
+    spvXfbBuffer2 spvXfbOutput2 = {};
+    spvXfbBuffer3 spvXfbOutput3 = {};
+    VertOut _25 = {};
+    if (any(gl_GlobalInvocationID >= spvStageInputSize))
+        return;
+    spvXfbOutput1.gl_Position = float4(1.0);
+    spvXfbOutput2.vFoo = float4(3.0);
+    _25.vBar = float4(5.0);
+    spvXfbOutput3.vBar = _25.vBar;
+    uint spvXfbIndex = 2 * gl_GlobalInvocationID.y * (spvStageInputSize.x - 1u) + 2 * gl_GlobalInvocationID.x;
+    uint spvInitOffset1 = atomic_load_explicit(spvXfbCounter1, memory_order_relaxed);
+    spvXfb1 = reinterpret_cast<device spvXfbBuffer1*>(reinterpret_cast<device char*>(spvXfb1) + spvInitOffset1);
+    if (gl_GlobalInvocationID.x != spvStageInputSize.x - 1u)
+        spvXfb1[spvXfbIndex] = spvXfbOutput1;
+    if (gl_GlobalInvocationID.x != 0)
+        spvXfb1[spvXfbIndex - 1u] = spvXfbOutput1;
+    uint spvInitOffset2 = atomic_load_explicit(spvXfbCounter2, memory_order_relaxed);
+    spvXfb2 = reinterpret_cast<device spvXfbBuffer2*>(reinterpret_cast<device char*>(spvXfb2) + spvInitOffset2);
+    if (gl_GlobalInvocationID.x != spvStageInputSize.x - 1u)
+        spvXfb2[spvXfbIndex] = spvXfbOutput2;
+    if (gl_GlobalInvocationID.x != 0)
+        spvXfb2[spvXfbIndex - 1u] = spvXfbOutput2;
+    uint spvInitOffset3 = atomic_load_explicit(spvXfbCounter3, memory_order_relaxed);
+    spvXfb3 = reinterpret_cast<device spvXfbBuffer3*>(reinterpret_cast<device char*>(spvXfb3) + spvInitOffset3);
+    if (gl_GlobalInvocationID.x != spvStageInputSize.x - 1u)
+        spvXfb3[spvXfbIndex] = spvXfbOutput3;
+    if (gl_GlobalInvocationID.x != 0)
+        spvXfb3[spvXfbIndex - 1u] = spvXfbOutput3;
+    threadgroup_barrier(mem_flags::mem_device);
+    if (all(gl_GlobalInvocationID.xy == 0))
+    {
+        uint spvWritten = 2 * (spvStageInputSize.x - 1u) * spvStageInputSize.y;
+        atomic_store_explicit(spvXfbCounter1, spvInitOffset1 + sizeof(*spvXfb1) * spvWritten, memory_order_relaxed);
+        atomic_store_explicit(spvXfbCounter2, spvInitOffset2 + sizeof(*spvXfb2) * spvWritten, memory_order_relaxed);
+        atomic_store_explicit(spvXfbCounter3, spvInitOffset3 + sizeof(*spvXfb3) * spvWritten, memory_order_relaxed);
+    }
+}
+

reference/opt/shaders-msl/vert/transform-feedback-decorations.xfb-point-list.vert

@@ -0,0 +1,59 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct VertOut
+{
+    float4 vBar;
+};
+
+struct spvXfbBuffer1
+{
+    char _m0_pad[4];
+    packed_float4 gl_Position;
+};
+
+struct spvXfbBuffer2
+{
+    char _m0_pad[16];
+    float4 vFoo;
+};
+
+struct spvXfbBuffer3
+{
+    float4 vBar;
+};
+
+kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 spvStageInputSize [[grid_size]], device atomic_uint* spvXfbCounter1 [[buffer(17)]], device spvXfbBuffer1* spvXfb1 [[buffer(13)]], device atomic_uint* spvXfbCounter2 [[buffer(18)]], device spvXfbBuffer2* spvXfb2 [[buffer(14)]], device atomic_uint* spvXfbCounter3 [[buffer(19)]], device spvXfbBuffer3* spvXfb3 [[buffer(15)]])
+{
+    spvXfbBuffer1 spvXfbOutput1 = {};
+    spvXfbBuffer2 spvXfbOutput2 = {};
+    spvXfbBuffer3 spvXfbOutput3 = {};
+    VertOut _25 = {};
+    if (any(gl_GlobalInvocationID >= spvStageInputSize))
+        return;
+    spvXfbOutput1.gl_Position = float4(1.0);
+    spvXfbOutput2.vFoo = float4(3.0);
+    _25.vBar = float4(5.0);
+    spvXfbOutput3.vBar = _25.vBar;
+    uint spvXfbIndex = gl_GlobalInvocationID.y * spvStageInputSize.x + gl_GlobalInvocationID.x;
+    uint spvInitOffset1 = atomic_load_explicit(spvXfbCounter1, memory_order_relaxed);
+    spvXfb1 = reinterpret_cast<device spvXfbBuffer1*>(reinterpret_cast<device char*>(spvXfb1) + spvInitOffset1);
+    spvXfb1[spvXfbIndex] = spvXfbOutput1;
+    uint spvInitOffset2 = atomic_load_explicit(spvXfbCounter2, memory_order_relaxed);
+    spvXfb2 = reinterpret_cast<device spvXfbBuffer2*>(reinterpret_cast<device char*>(spvXfb2) + spvInitOffset2);
+    spvXfb2[spvXfbIndex] = spvXfbOutput2;
+    uint spvInitOffset3 = atomic_load_explicit(spvXfbCounter3, memory_order_relaxed);
+    spvXfb3 = reinterpret_cast<device spvXfbBuffer3*>(reinterpret_cast<device char*>(spvXfb3) + spvInitOffset3);
+    spvXfb3[spvXfbIndex] = spvXfbOutput3;
+    threadgroup_barrier(mem_flags::mem_device);
+    if (all(gl_GlobalInvocationID.xy == 0))
+    {
+        uint spvWritten = spvStageInputSize.x * spvStageInputSize.y;
+        atomic_store_explicit(spvXfbCounter1, spvInitOffset1 + sizeof(*spvXfb1) * spvWritten, memory_order_relaxed);
+        atomic_store_explicit(spvXfbCounter2, spvInitOffset2 + sizeof(*spvXfb2) * spvWritten, memory_order_relaxed);
+        atomic_store_explicit(spvXfbCounter3, spvInitOffset3 + sizeof(*spvXfb3) * spvWritten, memory_order_relaxed);
+    }
+}
+

reference/opt/shaders-msl/vert/transform-feedback-decorations.xfb-triangle-fan.vert

@@ -0,0 +1,93 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct VertOut
+{
+    float4 vBar;
+};
+
+struct spvXfbBuffer1
+{
+    char _m0_pad[4];
+    packed_float4 gl_Position;
+};
+
+struct spvXfbBuffer2
+{
+    char _m0_pad[16];
+    float4 vFoo;
+};
+
+struct spvXfbBuffer3
+{
+    float4 vBar;
+};
+
+kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 spvStageInputSize [[grid_size]], device atomic_uint* spvXfbCounter1 [[buffer(17)]], device spvXfbBuffer1* spvXfb1 [[buffer(13)]], device atomic_uint* spvXfbCounter2 [[buffer(18)]], device spvXfbBuffer2* spvXfb2 [[buffer(14)]], device atomic_uint* spvXfbCounter3 [[buffer(19)]], device spvXfbBuffer3* spvXfb3 [[buffer(15)]])
+{
+    spvXfbBuffer1 spvXfbOutput1 = {};
+    spvXfbBuffer2 spvXfbOutput2 = {};
+    spvXfbBuffer3 spvXfbOutput3 = {};
+    VertOut _25 = {};
+    if (any(gl_GlobalInvocationID >= spvStageInputSize))
+        return;
+    spvXfbOutput1.gl_Position = float4(1.0);
+    spvXfbOutput2.vFoo = float4(3.0);
+    _25.vBar = float4(5.0);
+    spvXfbOutput3.vBar = _25.vBar;
+    uint spvXfbBaseIndex = 3 * gl_GlobalInvocationID.y * subsat(spvStageInputSize.x, 2u);
+    uint spvXfbIndex = spvXfbBaseIndex + 3 * gl_GlobalInvocationID.x - 2u;
+    uint spvInitOffset1 = atomic_load_explicit(spvXfbCounter1, memory_order_relaxed);
+    spvXfb1 = reinterpret_cast<device spvXfbBuffer1*>(reinterpret_cast<device char*>(spvXfb1) + spvInitOffset1);
+    if (gl_GlobalInvocationID.x == 0)
+    {
+        for (uint i = 0; i < subsat(spvStageInputSize.x, 2u); ++i)
+            spvXfb1[spvXfbBaseIndex + 3 * i] = spvXfbOutput1;
+    }
+    else
+    {
+        if (gl_GlobalInvocationID.x != spvStageInputSize.x - 1u)
+            spvXfb1[spvXfbIndex] = spvXfbOutput1;
+        if (gl_GlobalInvocationID.x != 1)
+            spvXfb1[spvXfbIndex - 2u] = spvXfbOutput1;
+    }
+    uint spvInitOffset2 = atomic_load_explicit(spvXfbCounter2, memory_order_relaxed);
+    spvXfb2 = reinterpret_cast<device spvXfbBuffer2*>(reinterpret_cast<device char*>(spvXfb2) + spvInitOffset2);
+    if (gl_GlobalInvocationID.x == 0)
+    {
+        for (uint i = 0; i < subsat(spvStageInputSize.x, 2u); ++i)
+            spvXfb2[spvXfbBaseIndex + 3 * i] = spvXfbOutput2;
+    }
+    else
+    {
+        if (gl_GlobalInvocationID.x != spvStageInputSize.x - 1u)
+            spvXfb2[spvXfbIndex] = spvXfbOutput2;
+        if (gl_GlobalInvocationID.x != 1)
+            spvXfb2[spvXfbIndex - 2u] = spvXfbOutput2;
+    }
+    uint spvInitOffset3 = atomic_load_explicit(spvXfbCounter3, memory_order_relaxed);
+    spvXfb3 = reinterpret_cast<device spvXfbBuffer3*>(reinterpret_cast<device char*>(spvXfb3) + spvInitOffset3);
+    if (gl_GlobalInvocationID.x == 0)
+    {
+        for (uint i = 0; i < subsat(spvStageInputSize.x, 2u); ++i)
+            spvXfb3[spvXfbBaseIndex + 3 * i] = spvXfbOutput3;
+    }
+    else
+    {
+        if (gl_GlobalInvocationID.x != spvStageInputSize.x - 1u)
+            spvXfb3[spvXfbIndex] = spvXfbOutput3;
+        if (gl_GlobalInvocationID.x != 1)
+            spvXfb3[spvXfbIndex - 2u] = spvXfbOutput3;
+    }
+    threadgroup_barrier(mem_flags::mem_device);
+    if (all(gl_GlobalInvocationID.xy == 0))
+    {
+        uint spvWritten = 3 * subsat(spvStageInputSize.x, 2u) * spvStageInputSize.y;
+        atomic_store_explicit(spvXfbCounter1, spvInitOffset1 + sizeof(*spvXfb1) * spvWritten, memory_order_relaxed);
+        atomic_store_explicit(spvXfbCounter2, spvInitOffset2 + sizeof(*spvXfb2) * spvWritten, memory_order_relaxed);
+        atomic_store_explicit(spvXfbCounter3, spvInitOffset3 + sizeof(*spvXfb3) * spvWritten, memory_order_relaxed);
+    }
+}
+