diff --git a/3rdparty/spirv-cross/CMakeLists.txt b/3rdparty/spirv-cross/CMakeLists.txt index 18d44842e..79566a86c 100644 --- a/3rdparty/spirv-cross/CMakeLists.txt +++ b/3rdparty/spirv-cross/CMakeLists.txt @@ -195,7 +195,7 @@ endif() if (SPIRV_CROSS_SHARED) set(spirv-cross-abi-major 0) - set(spirv-cross-abi-minor 1) + set(spirv-cross-abi-minor 2) set(spirv-cross-abi-patch 0) set(SPIRV_CROSS_VERSION ${spirv-cross-abi-major}.${spirv-cross-abi-minor}.${spirv-cross-abi-patch}) set(SPIRV_CROSS_INSTALL_LIB_DIR ${CMAKE_INSTALL_PREFIX}/lib) diff --git a/3rdparty/spirv-cross/main.cpp b/3rdparty/spirv-cross/main.cpp index 209ffd15a..9faa5a315 100644 --- a/3rdparty/spirv-cross/main.cpp +++ b/3rdparty/spirv-cross/main.cpp @@ -495,6 +495,8 @@ struct CLIArguments bool msl_ios = false; bool msl_pad_fragment_output = false; bool msl_domain_lower_left = false; + bool msl_argument_buffers = false; + vector msl_discrete_descriptor_sets; vector pls_in; vector pls_out; vector remaps; @@ -552,6 +554,8 @@ static void print_help() "\t[--msl-ios]\n" "\t[--msl-pad-fragment-output]\n" "\t[--msl-domain-lower-left]\n" + "\t[--msl-argument-buffers]\n" + "\t[--msl-discrete-descriptor-set ]\n" "\t[--hlsl]\n" "\t[--reflect]\n" "\t[--shader-model]\n" @@ -723,6 +727,9 @@ static int main_inner(int argc, char *argv[]) cbs.add("--msl-ios", [&args](CLIParser &) { args.msl_ios = true; }); cbs.add("--msl-pad-fragment-output", [&args](CLIParser &) { args.msl_pad_fragment_output = true; }); cbs.add("--msl-domain-lower-left", [&args](CLIParser &) { args.msl_domain_lower_left = true; }); + cbs.add("--msl-argument-buffers", [&args](CLIParser &) { args.msl_argument_buffers = true; }); + cbs.add("--msl-discrete-descriptor-set", + [&args](CLIParser &parser) { args.msl_discrete_descriptor_sets.push_back(parser.next_uint()); }); 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(); @@ -855,7 +862,10 @@ static int main_inner(int argc, char *argv[]) msl_opts.platform = CompilerMSL::Options::iOS; msl_opts.pad_fragment_output_components = args.msl_pad_fragment_output; msl_opts.tess_domain_origin_lower_left = args.msl_domain_lower_left; + msl_opts.argument_buffers = args.msl_argument_buffers; msl_comp->set_msl_options(msl_opts); + for (auto &v : args.msl_discrete_descriptor_sets) + msl_comp->add_discrete_descriptor_set(v); } else if (args.hlsl) compiler.reset(new CompilerHLSL(move(spirv_parser.get_parsed_ir()))); diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/argument-buffers-discrete.msl2.argument.discrete.comp b/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/argument-buffers-discrete.msl2.argument.discrete.comp new file mode 100644 index 000000000..98668f505 --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/argument-buffers-discrete.msl2.argument.discrete.comp @@ -0,0 +1,40 @@ +#include +#include + +using namespace metal; + +struct SSBO3 +{ + float4 v; +}; + +struct SSBO0 +{ + float4 v; +}; + +struct SSBO1 +{ + float4 v; +}; + +struct SSBO2 +{ + float4 v; +}; + +struct spvDescriptorSetBuffer0 +{ + const device SSBO0* ssbo0 [[id(0)]]; +}; + +struct spvDescriptorSetBuffer1 +{ + const device SSBO1* ssbo1 [[id(0)]]; +}; + +kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], const device SSBO2& ssbo2 [[buffer(5)]], device SSBO3& ssbo3 [[buffer(6)]]) +{ + ssbo3.v = ((*spvDescriptorSet0.ssbo0).v + (*spvDescriptorSet1.ssbo1).v) + ssbo2.v; +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/argument-buffers-image-load-store.msl2.argument.comp b/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/argument-buffers-image-load-store.msl2.argument.comp new file mode 100644 index 000000000..b1ebf27cf --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/argument-buffers-image-load-store.msl2.argument.comp @@ -0,0 +1,17 @@ +#include +#include + +using namespace metal; + +struct spvDescriptorSetBuffer0 +{ + texture2d uImage [[id(1)]]; + texture2d uImageRead [[id(2)]]; +}; + +kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + int2 _17 = int2(gl_GlobalInvocationID.xy); + spvDescriptorSet0.uImage.write(spvDescriptorSet0.uImageRead.read(uint2(_17)), uint2(_17)); +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/frag/argument-buffers.msl2.argument.frag b/3rdparty/spirv-cross/reference/opt/shaders-msl/frag/argument-buffers.msl2.argument.frag new file mode 100644 index 000000000..f00683898 --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/frag/argument-buffers.msl2.argument.frag @@ -0,0 +1,73 @@ +#include +#include + +using namespace metal; + +struct SSBO +{ + float4 ssbo; +}; + +struct SSBOs +{ + float4 ssbo; +}; + +struct Push +{ + float4 push; +}; + +struct UBO +{ + float4 ubo; +}; + +struct UBOs +{ + float4 ubo; +}; + +struct spvDescriptorSetBuffer0 +{ + texture2d uTexture [[id(2)]]; + sampler uTextureSmplr [[id(3)]]; + constant UBO* m_90 [[id(5)]]; + array, 2> uTextures [[id(6)]]; + array uTexturesSmplr [[id(8)]]; +}; + +struct spvDescriptorSetBuffer1 +{ + array, 4> uTexture2 [[id(3)]]; + device SSBO* m_60 [[id(7)]]; + const device SSBOs* ssbos [[id(8)]][2]; + array uSampler [[id(10)]]; +}; + +struct spvDescriptorSetBuffer2 +{ + constant UBOs* ubos [[id(4)]][4]; +}; + +struct main0_out +{ + float4 FragColor [[color(0)]]; +}; + +struct main0_in +{ + float2 vUV [[user(locn0)]]; +}; + +fragment main0_out main0(main0_in in [[stage_in]], constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], constant spvDescriptorSetBuffer2& spvDescriptorSet2 [[buffer(2)]], constant Push& registers [[buffer(3)]]) +{ + main0_out out = {}; + out.FragColor = ((((((spvDescriptorSet0.uTexture.sample(spvDescriptorSet0.uTextureSmplr, in.vUV) + spvDescriptorSet1.uTexture2[2].sample(spvDescriptorSet1.uSampler[1], in.vUV)) + spvDescriptorSet0.uTextures[1].sample(spvDescriptorSet0.uTexturesSmplr[1], in.vUV)) + (*spvDescriptorSet1.m_60).ssbo) + spvDescriptorSet1.ssbos[0]->ssbo) + registers.push) + (*spvDescriptorSet0.m_90).ubo) + spvDescriptorSet2.ubos[0]->ubo; + out.FragColor += (*spvDescriptorSet0.m_90).ubo; + out.FragColor += (*spvDescriptorSet1.m_60).ssbo; + out.FragColor += spvDescriptorSet2.ubos[1]->ubo; + out.FragColor += registers.push; + return out; +} + diff --git a/3rdparty/spirv-cross/reference/shaders-msl/comp/argument-buffers-discrete.msl2.argument.discrete.comp b/3rdparty/spirv-cross/reference/shaders-msl/comp/argument-buffers-discrete.msl2.argument.discrete.comp new file mode 100644 index 000000000..98668f505 --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-msl/comp/argument-buffers-discrete.msl2.argument.discrete.comp @@ -0,0 +1,40 @@ +#include +#include + +using namespace metal; + +struct SSBO3 +{ + float4 v; +}; + +struct SSBO0 +{ + float4 v; +}; + +struct SSBO1 +{ + float4 v; +}; + +struct SSBO2 +{ + float4 v; +}; + +struct spvDescriptorSetBuffer0 +{ + const device SSBO0* ssbo0 [[id(0)]]; +}; + +struct spvDescriptorSetBuffer1 +{ + const device SSBO1* ssbo1 [[id(0)]]; +}; + +kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], const device SSBO2& ssbo2 [[buffer(5)]], device SSBO3& ssbo3 [[buffer(6)]]) +{ + ssbo3.v = ((*spvDescriptorSet0.ssbo0).v + (*spvDescriptorSet1.ssbo1).v) + ssbo2.v; +} + diff --git a/3rdparty/spirv-cross/reference/shaders-msl/comp/argument-buffers-image-load-store.msl2.argument.comp b/3rdparty/spirv-cross/reference/shaders-msl/comp/argument-buffers-image-load-store.msl2.argument.comp new file mode 100644 index 000000000..58bb59eae --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-msl/comp/argument-buffers-image-load-store.msl2.argument.comp @@ -0,0 +1,17 @@ +#include +#include + +using namespace metal; + +struct spvDescriptorSetBuffer0 +{ + texture2d uImage [[id(1)]]; + texture2d uImageRead [[id(2)]]; +}; + +kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + int2 coord = int2(gl_GlobalInvocationID.xy); + spvDescriptorSet0.uImage.write(spvDescriptorSet0.uImageRead.read(uint2(coord)), uint2(coord)); +} + diff --git a/3rdparty/spirv-cross/reference/shaders-msl/frag/argument-buffers.msl2.argument.frag b/3rdparty/spirv-cross/reference/shaders-msl/frag/argument-buffers.msl2.argument.frag new file mode 100644 index 000000000..27b77a271 --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-msl/frag/argument-buffers.msl2.argument.frag @@ -0,0 +1,94 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct SSBO +{ + float4 ssbo; +}; + +struct SSBOs +{ + float4 ssbo; +}; + +struct Push +{ + float4 push; +}; + +struct UBO +{ + float4 ubo; +}; + +struct UBOs +{ + float4 ubo; +}; + +struct spvDescriptorSetBuffer0 +{ + texture2d uTexture [[id(2)]]; + sampler uTextureSmplr [[id(3)]]; + constant UBO* v_90 [[id(5)]]; + array, 2> uTextures [[id(6)]]; + array uTexturesSmplr [[id(8)]]; +}; + +struct spvDescriptorSetBuffer1 +{ + array, 4> uTexture2 [[id(3)]]; + device SSBO* v_60 [[id(7)]]; + const device SSBOs* ssbos [[id(8)]][2]; + array uSampler [[id(10)]]; +}; + +struct spvDescriptorSetBuffer2 +{ + constant UBOs* ubos [[id(4)]][4]; +}; + +struct main0_out +{ + float4 FragColor [[color(0)]]; +}; + +struct main0_in +{ + float2 vUV [[user(locn0)]]; +}; + +float4 sample_in_function2(thread texture2d uTexture, thread const sampler uTextureSmplr, thread float2& vUV, thread const array, 4> uTexture2, thread const array uSampler, thread const array, 2> uTextures, thread const array uTexturesSmplr, device SSBO& v_60, const device SSBOs* constant (&ssbos)[2], constant Push& registers) +{ + float4 ret = uTexture.sample(uTextureSmplr, vUV); + ret += uTexture2[2].sample(uSampler[1], vUV); + ret += uTextures[1].sample(uTexturesSmplr[1], vUV); + ret += v_60.ssbo; + ret += ssbos[0]->ssbo; + ret += registers.push; + return ret; +} + +float4 sample_in_function(thread texture2d uTexture, thread const sampler uTextureSmplr, thread float2& vUV, thread const array, 4> uTexture2, thread const array uSampler, thread const array, 2> uTextures, thread const array uTexturesSmplr, device SSBO& v_60, const device SSBOs* constant (&ssbos)[2], constant Push& registers, constant UBO& v_90, constant UBOs* constant (&ubos)[4]) +{ + float4 ret = sample_in_function2(uTexture, uTextureSmplr, vUV, uTexture2, uSampler, uTextures, uTexturesSmplr, v_60, ssbos, registers); + ret += v_90.ubo; + ret += ubos[0]->ubo; + return ret; +} + +fragment main0_out main0(main0_in in [[stage_in]], constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], constant spvDescriptorSetBuffer2& spvDescriptorSet2 [[buffer(2)]], constant Push& registers [[buffer(3)]]) +{ + main0_out out = {}; + out.FragColor = sample_in_function(spvDescriptorSet0.uTexture, spvDescriptorSet0.uTextureSmplr, in.vUV, spvDescriptorSet1.uTexture2, spvDescriptorSet1.uSampler, spvDescriptorSet0.uTextures, spvDescriptorSet0.uTexturesSmplr, (*spvDescriptorSet1.v_60), spvDescriptorSet1.ssbos, registers, (*spvDescriptorSet0.v_90), spvDescriptorSet2.ubos); + out.FragColor += (*spvDescriptorSet0.v_90).ubo; + out.FragColor += (*spvDescriptorSet1.v_60).ssbo; + out.FragColor += spvDescriptorSet2.ubos[1]->ubo; + out.FragColor += registers.push; + return out; +} + diff --git a/3rdparty/spirv-cross/shaders-msl/comp/argument-buffers-discrete.msl2.argument.discrete.comp b/3rdparty/spirv-cross/shaders-msl/comp/argument-buffers-discrete.msl2.argument.discrete.comp new file mode 100644 index 000000000..883f0019f --- /dev/null +++ b/3rdparty/spirv-cross/shaders-msl/comp/argument-buffers-discrete.msl2.argument.discrete.comp @@ -0,0 +1,27 @@ +#version 450 +layout(local_size_x = 1) in; + +layout(set = 0, binding = 0) readonly buffer SSBO0 +{ + vec4 v; +} ssbo0; + +layout(set = 1, binding = 0) readonly buffer SSBO1 +{ + vec4 v; +} ssbo1; + +layout(set = 2, binding = 5) readonly buffer SSBO2 +{ + vec4 v; +} ssbo2; + +layout(set = 3, binding = 6) writeonly buffer SSBO3 +{ + vec4 v; +} ssbo3; + +void main() +{ + ssbo3.v = ssbo0.v + ssbo1.v + ssbo2.v; +} diff --git a/3rdparty/spirv-cross/shaders-msl/comp/argument-buffers-image-load-store.msl2.argument.comp b/3rdparty/spirv-cross/shaders-msl/comp/argument-buffers-image-load-store.msl2.argument.comp new file mode 100644 index 000000000..72ca8899a --- /dev/null +++ b/3rdparty/spirv-cross/shaders-msl/comp/argument-buffers-image-load-store.msl2.argument.comp @@ -0,0 +1,10 @@ +#version 450 + +layout(set = 0, binding = 1, r32f) writeonly uniform image2D uImage; +layout(set = 0, binding = 2, r32f) readonly uniform image2D uImageRead; + +void main() +{ + ivec2 coord = ivec2(gl_GlobalInvocationID.xy); + imageStore(uImage, coord, imageLoad(uImageRead, coord)); +} diff --git a/3rdparty/spirv-cross/shaders-msl/frag/argument-buffers.msl2.argument.frag b/3rdparty/spirv-cross/shaders-msl/frag/argument-buffers.msl2.argument.frag new file mode 100644 index 000000000..0d9f6cdb4 --- /dev/null +++ b/3rdparty/spirv-cross/shaders-msl/frag/argument-buffers.msl2.argument.frag @@ -0,0 +1,61 @@ +#version 450 + +layout(std430, push_constant) uniform Push +{ + vec4 push; +} registers; + +layout(std140, set = 0, binding = 5) uniform UBO +{ + vec4 ubo; +}; + +layout(std430, set = 1, binding = 7) buffer SSBO +{ + vec4 ssbo; +}; + +layout(std430, set = 1, binding = 8) readonly buffer SSBOs +{ + vec4 ssbo; +} ssbos[2]; + +layout(std140, set = 2, binding = 4) uniform UBOs +{ + vec4 ubo; +} ubos[4]; + +layout(set = 0, binding = 2) uniform sampler2D uTexture; +layout(set = 0, binding = 6) uniform sampler2D uTextures[2]; +layout(set = 1, binding = 3) uniform texture2D uTexture2[4]; +layout(set = 1, binding = 10) uniform sampler uSampler[2]; +layout(location = 0) in vec2 vUV; +layout(location = 0) out vec4 FragColor; + +vec4 sample_in_function2() +{ + vec4 ret = texture(uTexture, vUV); + ret += texture(sampler2D(uTexture2[2], uSampler[1]), vUV); + ret += texture(uTextures[1], vUV); + ret += ssbo; + ret += ssbos[0].ssbo; + ret += registers.push; + return ret; +} + +vec4 sample_in_function() +{ + vec4 ret = sample_in_function2(); + ret += ubo; + ret += ubos[0].ubo; + return ret; +} + +void main() +{ + FragColor = sample_in_function(); + FragColor += ubo; + FragColor += ssbo; + FragColor += ubos[1].ubo; + FragColor += registers.push; +} diff --git a/3rdparty/spirv-cross/spirv_common.hpp b/3rdparty/spirv-cross/spirv_common.hpp index 11c65b522..d34843cf6 100644 --- a/3rdparty/spirv-cross/spirv_common.hpp +++ b/3rdparty/spirv-cross/spirv_common.hpp @@ -1404,8 +1404,9 @@ struct Meta { uint32_t packed_type = 0; bool packed = false; - uint32_t ib_member_index = static_cast(-1); + uint32_t ib_member_index = ~(0u); uint32_t ib_orig_id = 0; + uint32_t argument_buffer_id = ~(0u); } extended; }; diff --git a/3rdparty/spirv-cross/spirv_cross.cpp b/3rdparty/spirv-cross/spirv_cross.cpp index 004708d65..93f57f6ea 100644 --- a/3rdparty/spirv-cross/spirv_cross.cpp +++ b/3rdparty/spirv-cross/spirv_cross.cpp @@ -1091,6 +1091,11 @@ const std::string &Compiler::get_member_name(uint32_t id, uint32_t index) const return ir.get_member_name(id, index); } +void Compiler::set_qualified_name(uint32_t id, const string &name) +{ + ir.meta[id].decoration.qualified_alias = name; +} + void Compiler::set_member_qualified_name(uint32_t type_id, uint32_t index, const std::string &name) { ir.meta[type_id].members.resize(max(ir.meta[type_id].members.size(), size_t(index) + 1)); @@ -1156,6 +1161,10 @@ void Compiler::set_extended_decoration(uint32_t id, ExtendedDecorations decorati case SPIRVCrossDecorationInterfaceOrigID: dec.extended.ib_orig_id = value; break; + + case SPIRVCrossDecorationArgumentBufferID: + dec.extended.argument_buffer_id = value; + break; } } @@ -1182,6 +1191,10 @@ void Compiler::set_extended_member_decoration(uint32_t type, uint32_t index, Ext case SPIRVCrossDecorationInterfaceOrigID: dec.extended.ib_orig_id = value; break; + + case SPIRVCrossDecorationArgumentBufferID: + dec.extended.argument_buffer_id = value; + break; } } @@ -1205,6 +1218,9 @@ uint32_t Compiler::get_extended_decoration(uint32_t id, ExtendedDecorations deco case SPIRVCrossDecorationInterfaceOrigID: return dec.extended.ib_orig_id; + + case SPIRVCrossDecorationArgumentBufferID: + return dec.extended.argument_buffer_id; } return 0; @@ -1233,6 +1249,9 @@ uint32_t Compiler::get_extended_member_decoration(uint32_t type, uint32_t index, case SPIRVCrossDecorationInterfaceOrigID: return dec.extended.ib_orig_id; + + case SPIRVCrossDecorationArgumentBufferID: + return dec.extended.argument_buffer_id; } return 0; @@ -1258,6 +1277,9 @@ bool Compiler::has_extended_decoration(uint32_t id, ExtendedDecorations decorati case SPIRVCrossDecorationInterfaceOrigID: return dec.extended.ib_orig_id != 0; + + case SPIRVCrossDecorationArgumentBufferID: + return dec.extended.argument_buffer_id != 0; } return false; @@ -1286,6 +1308,9 @@ bool Compiler::has_extended_member_decoration(uint32_t type, uint32_t index, Ext case SPIRVCrossDecorationInterfaceOrigID: return dec.extended.ib_orig_id != 0; + + case SPIRVCrossDecorationArgumentBufferID: + return dec.extended.argument_buffer_id != uint32_t(-1); } return false; @@ -1305,12 +1330,16 @@ void Compiler::unset_extended_decoration(uint32_t id, ExtendedDecorations decora break; case SPIRVCrossDecorationInterfaceMemberIndex: - dec.extended.ib_member_index = -1; + dec.extended.ib_member_index = ~(0u); break; case SPIRVCrossDecorationInterfaceOrigID: dec.extended.ib_orig_id = 0; break; + + case SPIRVCrossDecorationArgumentBufferID: + dec.extended.argument_buffer_id = 0; + break; } } @@ -1330,12 +1359,16 @@ void Compiler::unset_extended_member_decoration(uint32_t type, uint32_t index, E break; case SPIRVCrossDecorationInterfaceMemberIndex: - dec.extended.ib_member_index = -1; + dec.extended.ib_member_index = ~(0u); break; case SPIRVCrossDecorationInterfaceOrigID: dec.extended.ib_orig_id = 0; break; + + case SPIRVCrossDecorationArgumentBufferID: + dec.extended.argument_buffer_id = 0; + break; } } diff --git a/3rdparty/spirv-cross/spirv_cross.hpp b/3rdparty/spirv-cross/spirv_cross.hpp index 16ca1d058..851a227bf 100644 --- a/3rdparty/spirv-cross/spirv_cross.hpp +++ b/3rdparty/spirv-cross/spirv_cross.hpp @@ -120,6 +120,7 @@ enum ExtendedDecorations SPIRVCrossDecorationPackedType, SPIRVCrossDecorationInterfaceMemberIndex, SPIRVCrossDecorationInterfaceOrigID, + SPIRVCrossDecorationArgumentBufferID }; class Compiler @@ -209,9 +210,6 @@ public: // or an empty string if no qualified alias exists const std::string &get_member_qualified_name(uint32_t type_id, uint32_t index) const; - // Sets the qualified member identifier for OpTypeStruct ID, member number "index". - void set_member_qualified_name(uint32_t type_id, uint32_t index, const std::string &name); - // Gets the decoration mask for a member of a struct, similar to get_decoration_mask. const Bitset &get_member_decoration_bitset(uint32_t id, uint32_t index) const; @@ -587,6 +585,10 @@ protected: // Gets the SPIR-V element type underlying an array variable. const SPIRType &get_variable_element_type(const SPIRVariable &var) const; + // Sets the qualified member identifier for OpTypeStruct ID, member number "index". + void set_member_qualified_name(uint32_t type_id, uint32_t index, const std::string &name); + void set_qualified_name(uint32_t id, const std::string &name); + // Returns if the given type refers to a sampled image. bool is_sampled_image_type(const SPIRType &type); diff --git a/3rdparty/spirv-cross/spirv_cross_c.cpp b/3rdparty/spirv-cross/spirv_cross_c.cpp index cf01ce589..b7943a2eb 100644 --- a/3rdparty/spirv-cross/spirv_cross_c.cpp +++ b/3rdparty/spirv-cross/spirv_cross_c.cpp @@ -470,6 +470,10 @@ spvc_result spvc_compiler_options_set_uint(spvc_compiler_options options, spvc_c options->msl.platform = static_cast(value); break; + case SPVC_COMPILER_OPTION_MSL_ARGUMENT_BUFFERS: + options->msl.argument_buffers = value != 0; + break; + default: options->context->report_error("Unknown option."); return SPVC_ERROR_INVALID_ARGUMENT; @@ -699,6 +703,19 @@ spvc_result spvc_compiler_msl_add_resource_binding(spvc_compiler compiler, return SPVC_SUCCESS; } +spvc_result spvc_compiler_msl_add_discrete_descriptor_set(spvc_compiler compiler, unsigned desc_set) +{ + if (compiler->backend != SPVC_BACKEND_MSL) + { + compiler->context->report_error("MSL function used on a non-MSL backend."); + return SPVC_ERROR_INVALID_ARGUMENT; + } + + auto &msl = *static_cast(compiler->compiler.get()); + msl.add_discrete_descriptor_set(desc_set); + return SPVC_SUCCESS; +} + spvc_bool spvc_compiler_msl_is_vertex_attribute_used(spvc_compiler compiler, unsigned location) { if (compiler->backend != SPVC_BACKEND_MSL) diff --git a/3rdparty/spirv-cross/spirv_cross_c.h b/3rdparty/spirv-cross/spirv_cross_c.h index 80ce0cb1a..8419aa51f 100644 --- a/3rdparty/spirv-cross/spirv_cross_c.h +++ b/3rdparty/spirv-cross/spirv_cross_c.h @@ -33,7 +33,7 @@ extern "C" { /* Bumped if ABI or API breaks backwards compatibility. */ #define SPVC_C_API_VERSION_MAJOR 0 /* Bumped if APIs or enumerations are added in a backwards compatible way. */ -#define SPVC_C_API_VERSION_MINOR 1 +#define SPVC_C_API_VERSION_MINOR 2 /* Bumped if internal implementation details change. */ #define SPVC_C_API_VERSION_PATCH 0 @@ -418,6 +418,7 @@ typedef enum spvc_compiler_option SPVC_COMPILER_OPTION_MSL_PAD_FRAGMENT_OUTPUT_COMPONENTS = 29 | SPVC_COMPILER_OPTION_MSL_BIT, SPVC_COMPILER_OPTION_MSL_TESS_DOMAIN_ORIGIN_LOWER_LEFT = 30 | SPVC_COMPILER_OPTION_MSL_BIT, SPVC_COMPILER_OPTION_MSL_PLATFORM = 31 | SPVC_COMPILER_OPTION_MSL_BIT, + SPVC_COMPILER_OPTION_MSL_ARGUMENT_BUFFERS = 32 | SPVC_COMPILER_OPTION_MSL_BIT, SPVC_COMPILER_OPTION_INT_MAX = 0x7fffffff } spvc_compiler_option; @@ -504,6 +505,7 @@ SPVC_PUBLIC_API spvc_result spvc_compiler_msl_add_vertex_attribute(spvc_compiler const spvc_msl_vertex_attribute *attrs); SPVC_PUBLIC_API spvc_result spvc_compiler_msl_add_resource_binding(spvc_compiler compiler, const spvc_msl_resource_binding *binding); +SPVC_PUBLIC_API spvc_result spvc_compiler_msl_add_discrete_descriptor_set(spvc_compiler compiler, unsigned desc_set); SPVC_PUBLIC_API spvc_bool spvc_compiler_msl_is_vertex_attribute_used(spvc_compiler compiler, unsigned location); SPVC_PUBLIC_API spvc_bool spvc_compiler_msl_is_resource_used(spvc_compiler compiler, SpvExecutionModel model, diff --git a/3rdparty/spirv-cross/spirv_msl.cpp b/3rdparty/spirv-cross/spirv_msl.cpp index 369fd9843..512196329 100644 --- a/3rdparty/spirv-cross/spirv_msl.cpp +++ b/3rdparty/spirv-cross/spirv_msl.cpp @@ -62,6 +62,12 @@ void CompilerMSL::add_msl_resource_binding(const MSLResourceBinding &binding) resource_bindings.push_back({ binding, false }); } +void CompilerMSL::add_discrete_descriptor_set(uint32_t desc_set) +{ + if (desc_set < kMaxArgumentBuffers) + argument_buffer_discrete_mask |= 1u << desc_set; +} + bool CompilerMSL::is_msl_vertex_attribute_used(uint32_t location) { return vtx_attrs_in_use.count(location) != 0; @@ -630,6 +636,15 @@ string CompilerMSL::compile() // the loop, so the hooks aren't added multiple times. fix_up_shader_inputs_outputs(); + // If we are using argument buffers, we create argument buffer structures for them here. + // These buffers will be used in the entry point, not the individual resources. + if (msl_options.argument_buffers) + { + if (!msl_options.supports_msl_version(2, 0)) + SPIRV_CROSS_THROW("Argument buffers can only be used with MSL 2.0 and up."); + analyze_argument_buffers(); + } + uint32_t pass_count = 0; do { @@ -4271,7 +4286,10 @@ void CompilerMSL::emit_function_prototype(SPIRFunction &func, const Bitset &) if (processing_entry_point) { - decl += entry_point_args(!func.arguments.empty()); + if (msl_options.argument_buffers) + decl += entry_point_args_argument_buffer(!func.arguments.empty()); + else + decl += entry_point_args_classic(!func.arguments.empty()); // If entry point function has variables that require early declaration, // ensure they each have an empty initializer, creating one if needed. @@ -4778,7 +4796,15 @@ string CompilerMSL::to_func_call_arg(uint32_t id) // Manufacture automatic sampler arg if the arg is a SampledImage texture. auto &type = expression_type(id); if (type.basetype == SPIRType::SampledImage && type.image.dim != DimBuffer) - arg_str += ", " + to_sampler_expression(id); + { + // Need to check the base variable in case we need to apply a qualified alias. + uint32_t var_id = 0; + auto *sampler_var = maybe_get(id); + if (sampler_var) + var_id = sampler_var->basevariable; + + arg_str += ", " + to_sampler_expression(var_id ? var_id : id); + } if (msl_options.swizzle_texture_samples && has_sampled_images && is_sampled_image_type(type)) arg_str += ", " + to_swizzle_expression(id); @@ -4963,6 +4989,10 @@ string CompilerMSL::to_struct_member(const SPIRType &type, uint32_t member_type_ const SPIRType *effective_membertype = &membertype; SPIRType override_type; + uint32_t orig_id = 0; + if (has_extended_member_decoration(type.self, index, SPIRVCrossDecorationInterfaceOrigID)) + orig_id = get_extended_member_decoration(type.self, index, SPIRVCrossDecorationInterfaceOrigID); + if (member_is_packed_type(type, index)) { // If we're packing a matrix, output an appropriate typedef @@ -4992,8 +5022,23 @@ string CompilerMSL::to_struct_member(const SPIRType &type, uint32_t member_type_ pack_pfx = "packed_"; } - return join(pack_pfx, type_to_glsl(*effective_membertype), " ", qualifier, to_member_name(type, index), - member_attribute_qualifier(type, index), type_to_array_glsl(membertype), ";"); + // Very specifically, image load-store in argument buffers are disallowed on MSL on iOS. + if (msl_options.is_ios() && membertype.basetype == SPIRType::Image && membertype.image.sampled == 2) + { + if (!has_decoration(orig_id, DecorationNonWritable)) + SPIRV_CROSS_THROW("Writable images are not allowed in argument buffers on iOS."); + } + + // Array information is baked into these types. + string array_type; + if (membertype.basetype != SPIRType::Image && membertype.basetype != SPIRType::Sampler && + membertype.basetype != SPIRType::SampledImage) + { + array_type = type_to_array_glsl(membertype); + } + + return join(pack_pfx, type_to_glsl(*effective_membertype, orig_id), " ", qualifier, to_member_name(type, index), + member_attribute_qualifier(type, index), array_type, ";"); } // Emit a structure member, padding and packing to maintain the correct memeber alignments. @@ -5014,6 +5059,10 @@ string CompilerMSL::member_attribute_qualifier(const SPIRType &type, uint32_t in BuiltIn builtin = BuiltInMax; bool is_builtin = is_member_builtin(type, index, &builtin); + if (has_extended_member_decoration(type.self, index, SPIRVCrossDecorationArgumentBufferID)) + return join(" [[id(", get_extended_member_decoration(type.self, index, SPIRVCrossDecorationArgumentBufferID), + ")]]"); + // Vertex function inputs if (execution.model == ExecutionModelVertex && type.storage == StorageClassInput) { @@ -5393,7 +5442,7 @@ string CompilerMSL::get_argument_address_space(const SPIRVariable &argument) return "thread"; } -string CompilerMSL::get_type_address_space(const SPIRType &type) +string CompilerMSL::get_type_address_space(const SPIRType &type, uint32_t id) { switch (type.storage) { @@ -5401,8 +5450,16 @@ string CompilerMSL::get_type_address_space(const SPIRType &type) return "threadgroup"; case StorageClassStorageBuffer: - // FIXME: Need to use 'const device' for pointers into non-writable SSBOs - return "device"; + { + // This can be called for variable pointer contexts as well, so be very careful about which method we choose. + Bitset flags; + if (ir.ids[id].get_type() == TypeVariable && has_decoration(type.self, DecorationBlock)) + flags = get_buffer_block_flags(id); + else + flags = get_decoration_bitset(id); + + return flags.get(DecorationNonWritable) ? "const device" : "device"; + } case StorageClassUniform: case StorageClassUniformConstant: @@ -5410,9 +5467,17 @@ string CompilerMSL::get_type_address_space(const SPIRType &type) if (type.basetype == SPIRType::Struct) { bool ssbo = has_decoration(type.self, DecorationBufferBlock); - // FIXME: Need to use 'const device' for pointers into non-writable SSBOs if (ssbo) - return "device"; + { + // This can be called for variable pointer contexts as well, so be very careful about which method we choose. + Bitset flags; + if (ir.ids[id].get_type() == TypeVariable && has_decoration(type.self, DecorationBlock)) + flags = get_buffer_block_flags(id); + else + flags = get_decoration_bitset(id); + + return flags.get(DecorationNonWritable) ? "const device" : "device"; + } else return "constant"; } @@ -5435,10 +5500,9 @@ string CompilerMSL::get_type_address_space(const SPIRType &type) return "thread"; } -// Returns a string containing a comma-delimited list of args for the entry point function -string CompilerMSL::entry_point_args(bool append_comma) +string CompilerMSL::entry_point_arg_stage_in() { - string ep_args; + string decl; // Stage-in structure uint32_t stage_in_id; @@ -5452,126 +5516,15 @@ string CompilerMSL::entry_point_args(bool append_comma) auto &var = get(stage_in_id); auto &type = get_variable_data_type(var); - if (!ep_args.empty()) - ep_args += ", "; - add_resource_name(var.self); - ep_args += join(type_to_glsl(type), " ", to_name(var.self), " [[stage_in]]"); + decl = join(type_to_glsl(type), " ", to_name(var.self), " [[stage_in]]"); } - // Output resources, sorted by resource index & type - // We need to sort to work around a bug on macOS 10.13 with NVidia drivers where switching between shaders - // with different order of buffers can result in issues with buffer assignments inside the driver. - struct Resource - { - Variant *id; - string name; - SPIRType::BaseType basetype; - uint32_t index; - }; - - vector resources; - - ir.for_each_typed_id([&](uint32_t self, SPIRVariable &var) { - auto &id = ir.ids[self]; - auto &type = get_variable_data_type(var); - - uint32_t var_id = var.self; - - if ((var.storage == StorageClassUniform || var.storage == StorageClassUniformConstant || - var.storage == StorageClassPushConstant || var.storage == StorageClassStorageBuffer) && - !is_hidden_variable(var)) - { - if (type.basetype == SPIRType::SampledImage) - { - add_resource_name(var_id); - resources.push_back( - { &id, to_name(var_id), SPIRType::Image, get_metal_resource_index(var, SPIRType::Image) }); - - if (type.image.dim != DimBuffer && constexpr_samplers.count(var_id) == 0) - { - resources.push_back({ &id, to_sampler_expression(var_id), SPIRType::Sampler, - get_metal_resource_index(var, SPIRType::Sampler) }); - } - } - else if (constexpr_samplers.count(var_id) == 0) - { - // constexpr samplers are not declared as resources. - add_resource_name(var_id); - resources.push_back( - { &id, to_name(var_id), type.basetype, get_metal_resource_index(var, type.basetype) }); - } - } - }); - - std::sort(resources.begin(), resources.end(), [](const Resource &lhs, const Resource &rhs) { - return tie(lhs.basetype, lhs.index) < tie(rhs.basetype, rhs.index); - }); - - for (auto &r : resources) - { - auto &var = r.id->get(); - auto &type = get_variable_data_type(var); - - uint32_t var_id = var.self; - - switch (r.basetype) - { - case SPIRType::Struct: - { - auto &m = ir.meta[type.self]; - if (m.members.size() == 0) - break; - if (!type.array.empty()) - { - if (type.array.size() > 1) - SPIRV_CROSS_THROW("Arrays of arrays of buffers are not supported."); - - // Metal doesn't directly support this, so we must expand the - // array. We'll declare a local array to hold these elements - // later. - uint32_t array_size = to_array_size_literal(type); - - if (array_size == 0) - SPIRV_CROSS_THROW("Unsized arrays of buffers are not supported in MSL."); - - buffer_arrays.push_back(var_id); - for (uint32_t i = 0; i < array_size; ++i) - { - if (!ep_args.empty()) - ep_args += ", "; - ep_args += get_argument_address_space(var) + " " + type_to_glsl(type) + "* " + r.name + "_" + - convert_to_string(i); - ep_args += " [[buffer(" + convert_to_string(r.index + i) + ")]]"; - } - } - else - { - if (!ep_args.empty()) - ep_args += ", "; - ep_args += get_argument_address_space(var) + " " + type_to_glsl(type) + "& " + r.name; - ep_args += " [[buffer(" + convert_to_string(r.index) + ")]]"; - } - break; - } - case SPIRType::Sampler: - if (!ep_args.empty()) - ep_args += ", "; - ep_args += sampler_type(type) + " " + r.name; - ep_args += " [[sampler(" + convert_to_string(r.index) + ")]]"; - break; - case SPIRType::Image: - if (!ep_args.empty()) - ep_args += ", "; - ep_args += image_type_glsl(type, var_id) + " " + r.name; - ep_args += " [[texture(" + convert_to_string(r.index) + ")]]"; - break; - default: - SPIRV_CROSS_THROW("Unexpected resource type"); - break; - } - } + return decl; +} +void CompilerMSL::entry_point_args_builtin(string &ep_args) +{ // Builtin variables ir.for_each_typed_id([&](uint32_t var_id, SPIRVariable &var) { BuiltIn bi_type = ir.meta[var_id].decoration.builtin_type; @@ -5650,6 +5603,170 @@ string CompilerMSL::entry_point_args(bool append_comma) } } } +} + +string CompilerMSL::entry_point_args_argument_buffer(bool append_comma) +{ + string ep_args = entry_point_arg_stage_in(); + + for (uint32_t i = 0; i < kMaxArgumentBuffers; i++) + { + uint32_t id = argument_buffer_ids[i]; + if (id == 0) + continue; + + add_resource_name(id); + auto &var = get(id); + auto &type = get_variable_data_type(var); + + if (!ep_args.empty()) + ep_args += ", "; + + ep_args += get_argument_address_space(var) + " " + type_to_glsl(type) + "& " + to_name(id); + ep_args += " [[buffer(" + convert_to_string(i) + ")]]"; + + // Makes it more practical for testing, since the push constant block can occupy the first available + // buffer slot if it's not bound explicitly. + next_metal_resource_index_buffer = i + 1; + } + + entry_point_args_discrete_descriptors(ep_args); + entry_point_args_builtin(ep_args); + + if (!ep_args.empty() && append_comma) + ep_args += ", "; + + return ep_args; +} + +void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args) +{ + // Output resources, sorted by resource index & type + // We need to sort to work around a bug on macOS 10.13 with NVidia drivers where switching between shaders + // with different order of buffers can result in issues with buffer assignments inside the driver. + struct Resource + { + SPIRVariable *var; + string name; + SPIRType::BaseType basetype; + uint32_t index; + }; + + vector resources; + + ir.for_each_typed_id([&](uint32_t, SPIRVariable &var) { + if ((var.storage == StorageClassUniform || var.storage == StorageClassUniformConstant || + var.storage == StorageClassPushConstant || var.storage == StorageClassStorageBuffer) && + !is_hidden_variable(var)) + { + auto &type = get_variable_data_type(var); + uint32_t var_id = var.self; + + if (var.storage != StorageClassPushConstant) + { + uint32_t desc_set = get_decoration(var_id, DecorationDescriptorSet); + if (descriptor_set_is_argument_buffer(desc_set)) + return; + } + + if (type.basetype == SPIRType::SampledImage) + { + add_resource_name(var_id); + resources.push_back( + { &var, to_name(var_id), SPIRType::Image, get_metal_resource_index(var, SPIRType::Image) }); + + if (type.image.dim != DimBuffer && constexpr_samplers.count(var_id) == 0) + { + resources.push_back({ &var, to_sampler_expression(var_id), SPIRType::Sampler, + get_metal_resource_index(var, SPIRType::Sampler) }); + } + } + else if (constexpr_samplers.count(var_id) == 0) + { + // constexpr samplers are not declared as resources. + add_resource_name(var_id); + resources.push_back( + { &var, to_name(var_id), type.basetype, get_metal_resource_index(var, type.basetype) }); + } + } + }); + + sort(resources.begin(), resources.end(), [](const Resource &lhs, const Resource &rhs) { + return tie(lhs.basetype, lhs.index) < tie(rhs.basetype, rhs.index); + }); + + for (auto &r : resources) + { + auto &var = *r.var; + auto &type = get_variable_data_type(var); + + uint32_t var_id = var.self; + + switch (r.basetype) + { + case SPIRType::Struct: + { + auto &m = ir.meta[type.self]; + if (m.members.size() == 0) + break; + if (!type.array.empty()) + { + if (type.array.size() > 1) + SPIRV_CROSS_THROW("Arrays of arrays of buffers are not supported."); + + // Metal doesn't directly support this, so we must expand the + // array. We'll declare a local array to hold these elements + // later. + uint32_t array_size = to_array_size_literal(type); + + if (array_size == 0) + SPIRV_CROSS_THROW("Unsized arrays of buffers are not supported in MSL."); + + buffer_arrays.push_back(var_id); + for (uint32_t i = 0; i < array_size; ++i) + { + if (!ep_args.empty()) + ep_args += ", "; + ep_args += get_argument_address_space(var) + " " + type_to_glsl(type) + "* " + r.name + "_" + + convert_to_string(i); + ep_args += " [[buffer(" + convert_to_string(r.index + i) + ")]]"; + } + } + else + { + if (!ep_args.empty()) + ep_args += ", "; + ep_args += get_argument_address_space(var) + " " + type_to_glsl(type) + "& " + r.name; + ep_args += " [[buffer(" + convert_to_string(r.index) + ")]]"; + } + break; + } + case SPIRType::Sampler: + if (!ep_args.empty()) + ep_args += ", "; + ep_args += sampler_type(type) + " " + r.name; + ep_args += " [[sampler(" + convert_to_string(r.index) + ")]]"; + break; + case SPIRType::Image: + if (!ep_args.empty()) + ep_args += ", "; + ep_args += image_type_glsl(type, var_id) + " " + r.name; + ep_args += " [[texture(" + convert_to_string(r.index) + ")]]"; + break; + default: + SPIRV_CROSS_THROW("Unexpected resource type"); + break; + } + } +} + +// Returns a string containing a comma-delimited list of args for the entry point function +// This is the "classic" method of MSL 1 when we don't have argument buffer support. +string CompilerMSL::entry_point_args_classic(bool append_comma) +{ + string ep_args = entry_point_arg_stage_in(); + entry_point_args_discrete_descriptors(ep_args); + entry_point_args_builtin(ep_args); if (!ep_args.empty() && append_comma) ep_args += ", "; @@ -5871,6 +5988,22 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg) // Arrays of images and samplers are special cased. if (!address_space.empty()) decl = join(address_space, " ", decl); + + if (msl_options.argument_buffers) + { + // An awkward case where we need to emit *more* address space declarations (yay!). + // An example is where we pass down an array of buffer pointers to leaf functions. + // It's a constant array containing pointers to constants. + // The pointer array is always constant however. E.g. + // device SSBO * constant (&array)[N]. + // const device SSBO * constant (&array)[N]. + // constant SSBO * constant (&array)[N]. + // However, this only matters for argument buffers, since for MSL 1.0 style codegen, + // we emit the buffer array on stack instead, and that seems to work just fine apparently. + if (storage == StorageClassUniform || storage == StorageClassStorageBuffer) + decl += " constant"; + } + decl += " (&"; decl += to_expression(name_id); decl += ")"; @@ -6230,9 +6363,16 @@ string CompilerMSL::to_member_reference(uint32_t base, const SPIRType &type, uin auto *var = maybe_get(base); // If this is a buffer array, we have to dereference the buffer pointers. // Otherwise, if this is a pointer expression, dereference it. - if ((var && ((var->storage == StorageClassUniform || var->storage == StorageClassStorageBuffer) && - is_array(get(var->basetype)))) || - (!ptr_chain && should_dereference(base))) + + bool declared_as_pointer = false; + + if (var) + { + bool is_buffer_variable = var->storage == StorageClassUniform || var->storage == StorageClassStorageBuffer; + declared_as_pointer = is_buffer_variable && is_array(get(var->basetype)); + } + + if (declared_as_pointer || (!ptr_chain && should_dereference(base))) return join("->", to_member_name(type, index)); else return join(".", to_member_name(type, index)); @@ -6259,7 +6399,7 @@ string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id) // Pointer? if (type.pointer) { - type_name = join(get_type_address_space(type), " ", type_to_glsl(get(type.parent_type), id)); + type_name = join(get_type_address_space(type, id), " ", type_to_glsl(get(type.parent_type), id)); switch (type.basetype) { case SPIRType::Image: @@ -7391,3 +7531,174 @@ std::string CompilerMSL::to_initializer_expression(const SPIRVariable &var) else return CompilerGLSL::to_initializer_expression(var); } + +bool CompilerMSL::descriptor_set_is_argument_buffer(uint32_t desc_set) const +{ + if (!msl_options.argument_buffers) + return false; + if (desc_set >= kMaxArgumentBuffers) + return false; + + return (argument_buffer_discrete_mask & (1u << desc_set)) == 0; +} + +void CompilerMSL::analyze_argument_buffers() +{ + // Gather all used resources and sort them out into argument buffers. + // Each argument buffer corresponds to a descriptor set in SPIR-V. + // The [[id(N)]] values used correspond to the resource mapping we have for MSL. + // Otherwise, the binding number is used, but this is generally not safe some types like + // combined image samplers and arrays of resources. Metal needs different indices here, + // while SPIR-V can have one descriptor set binding. To use argument buffers in practice, + // you will need to use the remapping from the API. + for (auto &id : argument_buffer_ids) + id = 0; + + // Output resources, sorted by resource index & type. + struct Resource + { + SPIRVariable *var; + string name; + SPIRType::BaseType basetype; + uint32_t index; + }; + vector resources_in_set[kMaxArgumentBuffers]; + + ir.for_each_typed_id([&](uint32_t self, SPIRVariable &var) { + if ((var.storage == StorageClassUniform || var.storage == StorageClassUniformConstant || + var.storage == StorageClassStorageBuffer) && + !is_hidden_variable(var)) + { + uint32_t desc_set = get_decoration(self, DecorationDescriptorSet); + // Ignore if it's part of a push descriptor set. + if (!descriptor_set_is_argument_buffer(desc_set)) + return; + + uint32_t var_id = var.self; + auto &type = get_variable_data_type(var); + + if (desc_set >= kMaxArgumentBuffers) + SPIRV_CROSS_THROW("Descriptor set index is out of range."); + + if (type.basetype == SPIRType::SampledImage) + { + add_resource_name(var_id); + + uint32_t image_resource_index = get_metal_resource_index(var, SPIRType::Image); + uint32_t sampler_resource_index = get_metal_resource_index(var, SPIRType::Sampler); + + // Avoid trivial conflicts where we didn't remap. + // This will let us at least compile test cases without having to instrument remaps. + if (sampler_resource_index == image_resource_index) + sampler_resource_index += type.array.empty() ? 1 : to_array_size_literal(type); + + resources_in_set[desc_set].push_back({ &var, to_name(var_id), SPIRType::Image, image_resource_index }); + + if (type.image.dim != DimBuffer && constexpr_samplers.count(var_id) == 0) + { + resources_in_set[desc_set].push_back( + { &var, to_sampler_expression(var_id), SPIRType::Sampler, sampler_resource_index }); + } + } + else if (constexpr_samplers.count(var_id) == 0) + { + // constexpr samplers are not declared as resources. + add_resource_name(var_id); + resources_in_set[desc_set].push_back( + { &var, to_name(var_id), type.basetype, get_metal_resource_index(var, type.basetype) }); + } + } + }); + + for (uint32_t desc_set = 0; desc_set < kMaxArgumentBuffers; desc_set++) + { + auto &resources = resources_in_set[desc_set]; + if (resources.empty()) + continue; + + assert(descriptor_set_is_argument_buffer(desc_set)); + + uint32_t next_id = ir.increase_bound_by(3); + uint32_t type_id = next_id + 1; + uint32_t ptr_type_id = next_id + 2; + argument_buffer_ids[desc_set] = next_id; + + auto &buffer_type = set(type_id); + buffer_type.storage = StorageClassUniform; + buffer_type.basetype = SPIRType::Struct; + set_name(type_id, join("spvDescriptorSetBuffer", desc_set)); + + auto &ptr_type = set(ptr_type_id); + ptr_type = buffer_type; + ptr_type.pointer = true; + ptr_type.pointer_depth = 1; + ptr_type.parent_type = type_id; + + uint32_t buffer_variable_id = next_id; + set(buffer_variable_id, ptr_type_id, StorageClassUniform); + set_name(buffer_variable_id, join("spvDescriptorSet", desc_set)); + + // Ids must be emitted in ID order. + sort(begin(resources), end(resources), [&](const Resource &lhs, const Resource &rhs) -> bool { + return tie(lhs.index, lhs.basetype) < tie(rhs.index, rhs.basetype); + }); + + uint32_t member_index = 0; + for (auto &resource : resources) + { + auto &var = *resource.var; + auto &type = get_variable_data_type(var); + string mbr_name = ensure_valid_name(resource.name, "m"); + set_member_name(buffer_type.self, member_index, mbr_name); + + if (resource.basetype == SPIRType::Sampler && type.basetype != SPIRType::Sampler) + { + // Have to synthesize a sampler type here. + + bool type_is_array = !type.array.empty(); + uint32_t sampler_type_id = ir.increase_bound_by(type_is_array ? 2 : 1); + auto &new_sampler_type = set(sampler_type_id); + new_sampler_type.basetype = SPIRType::Sampler; + new_sampler_type.storage = StorageClassUniformConstant; + + if (type_is_array) + { + uint32_t sampler_type_array_id = sampler_type_id + 1; + auto &sampler_type_array = set(sampler_type_array_id); + sampler_type_array = new_sampler_type; + sampler_type_array.array = type.array; + sampler_type_array.array_size_literal = type.array_size_literal; + sampler_type_array.parent_type = sampler_type_id; + buffer_type.member_types.push_back(sampler_type_array_id); + } + else + buffer_type.member_types.push_back(sampler_type_id); + } + else + { + if (resource.basetype == SPIRType::Image || resource.basetype == SPIRType::Sampler || + resource.basetype == SPIRType::SampledImage) + { + // Drop pointer information when we emit the resources into a struct. + buffer_type.member_types.push_back(get_variable_data_type_id(var)); + set_qualified_name(var.self, join(to_name(buffer_variable_id), ".", mbr_name)); + } + else + { + // Resources will be declared as pointers not references, so automatically dereference as appropriate. + buffer_type.member_types.push_back(var.basetype); + if (type.array.empty()) + set_qualified_name(var.self, join("(*", to_name(buffer_variable_id), ".", mbr_name, ")")); + else + set_qualified_name(var.self, join(to_name(buffer_variable_id), ".", mbr_name)); + } + } + + set_extended_member_decoration(buffer_type.self, member_index, SPIRVCrossDecorationArgumentBufferID, + resource.index); + set_extended_member_decoration(buffer_type.self, member_index, SPIRVCrossDecorationInterfaceOrigID, + var.self); + member_index++; + } + } +} diff --git a/3rdparty/spirv-cross/spirv_msl.hpp b/3rdparty/spirv-cross/spirv_msl.hpp index fda3fc1fd..38610e15c 100644 --- a/3rdparty/spirv-cross/spirv_msl.hpp +++ b/3rdparty/spirv-cross/spirv_msl.hpp @@ -54,6 +54,11 @@ struct MSLVertexAttr // Matches the binding index of a MSL resource for a binding within a descriptor set. // Taken together, the stage, desc_set and binding combine to form a reference to a resource // descriptor used in a particular shading stage. +// If using MSL 2.0 argument buffers, and the descriptor set is not marked as a discrete descriptor set, +// the binding reference we remap to will become an [[id(N)]] attribute within +// the "descriptor set" argument buffer structure. +// For resources which are bound in the "classic" MSL 1.0 way or discrete descriptors, the remap will become a +// [[buffer(N)]], [[texture(N)]] or [[sampler(N)]] depending on the resource types used. struct MSLResourceBinding { spv::ExecutionModel stage = spv::ExecutionModelMax; @@ -148,6 +153,8 @@ static const uint32_t kPushConstDescSet = ~(0u); // element to indicate the bindings for the push constants. static const uint32_t kPushConstBinding = 0; +static const uint32_t kMaxArgumentBuffers = 8; + // The current version of the aux buffer structure. It must be incremented any time a // new field is added to the aux buffer. #define SPIRV_CROSS_MSL_AUX_BUFFER_STRUCT_VERSION 1 @@ -180,6 +187,10 @@ public: bool swizzle_texture_samples = false; bool tess_domain_origin_lower_left = false; + // Enable use of MSL 2.0 indirect argument buffers. + // MSL 2.0 must also be enabled. + bool argument_buffers = false; + // Fragment output in MSL must have at least as many components as the render pass. // Add support to explicit pad out components. bool pad_fragment_output_components = false; @@ -275,6 +286,10 @@ public: // the set/binding combination was used by the MSL code. void add_msl_resource_binding(const MSLResourceBinding &resource); + // When using MSL argument buffers, we can force "classic" MSL 1.0 binding schemes for certain descriptor sets. + // This corresponds to VK_KHR_push_descriptor in Vulkan. + void add_discrete_descriptor_set(uint32_t desc_set); + // Query after compilation is done. This allows you to check if a location or set/binding combination was used by the shader. bool is_msl_vertex_attribute_used(uint32_t location); bool is_msl_resource_binding_used(spv::ExecutionModel model, uint32_t set, uint32_t binding); @@ -413,7 +428,11 @@ protected: void fix_up_shader_inputs_outputs(); std::string func_type_decl(SPIRType &type); - std::string entry_point_args(bool append_comma); + std::string entry_point_args_classic(bool append_comma); + std::string entry_point_args_argument_buffer(bool append_comma); + std::string entry_point_arg_stage_in(); + void entry_point_args_builtin(std::string &args); + void entry_point_args_discrete_descriptors(std::string &args); std::string to_qualified_member_name(const SPIRType &type, uint32_t index); std::string ensure_valid_name(std::string name, std::string pfx); std::string to_sampler_expression(uint32_t id); @@ -432,7 +451,7 @@ protected: bool is_member_packable(SPIRType &ib_type, uint32_t index); MSLStructMemberKey get_struct_member_key(uint32_t type_id, uint32_t index); std::string get_argument_address_space(const SPIRVariable &argument); - std::string get_type_address_space(const SPIRType &type); + std::string get_type_address_space(const SPIRType &type, uint32_t id); SPIRType &get_stage_in_struct_type(); SPIRType &get_stage_out_struct_type(); SPIRType &get_patch_stage_in_struct_type(); @@ -513,6 +532,11 @@ protected: std::unordered_map constexpr_samplers; std::vector buffer_arrays; + uint32_t argument_buffer_ids[kMaxArgumentBuffers]; + uint32_t argument_buffer_discrete_mask = 0; + void analyze_argument_buffers(); + bool descriptor_set_is_argument_buffer(uint32_t desc_set) const; + uint32_t get_target_components_for_fragment_location(uint32_t location) const; uint32_t build_extended_vector_type(uint32_t type_id, uint32_t components); diff --git a/3rdparty/spirv-cross/test_shaders.py b/3rdparty/spirv-cross/test_shaders.py index feefd2639..51095f0ad 100755 --- a/3rdparty/spirv-cross/test_shaders.py +++ b/3rdparty/spirv-cross/test_shaders.py @@ -169,6 +169,14 @@ def cross_compile_msl(shader, spirv, opt, paths): msl_args.append('--msl-capture-output') if '.domain.' in shader: msl_args.append('--msl-domain-lower-left') + if '.argument.' in shader: + msl_args.append('--msl-argument-buffers') + if '.discrete.' in shader: + # Arbitrary for testing purposes. + msl_args.append('--msl-discrete-descriptor-set') + msl_args.append('2') + msl_args.append('--msl-discrete-descriptor-set') + msl_args.append('3') subprocess.check_call(msl_args)