Merge pull request #890 from KhronosGroup/msl-argument-buffers
MSL: Add support for Metal 2 indirect argument buffers.
This commit is contained in:
Hans-Kristian Arntzen
GitHub
commit
62dce70c2e
@ -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)
|
||||
|
||||
10
main.cpp
10
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<uint32_t> msl_discrete_descriptor_sets;
|
||||
vector<PLSArg> pls_in;
|
||||
vector<PLSArg> pls_out;
|
||||
vector<Remap> 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 <index>]\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())));
|
||||
|
||||
@ -0,0 +1,40 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
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;
|
||||
}
|
||||
|
||||
@ -0,0 +1,17 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct spvDescriptorSetBuffer0
|
||||
{
|
||||
texture2d<float, access::write> uImage [[id(1)]];
|
||||
texture2d<float> 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));
|
||||
}
|
||||
|
||||
@ -0,0 +1,73 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
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<float> uTexture [[id(2)]];
|
||||
sampler uTextureSmplr [[id(3)]];
|
||||
constant UBO* m_90 [[id(5)]];
|
||||
array<texture2d<float>, 2> uTextures [[id(6)]];
|
||||
array<sampler, 2> uTexturesSmplr [[id(8)]];
|
||||
};
|
||||
|
||||
struct spvDescriptorSetBuffer1
|
||||
{
|
||||
array<texture2d<float>, 4> uTexture2 [[id(3)]];
|
||||
device SSBO* m_60 [[id(7)]];
|
||||
const device SSBOs* ssbos [[id(8)]][2];
|
||||
array<sampler, 2> 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;
|
||||
}
|
||||
|
||||
@ -0,0 +1,40 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
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;
|
||||
}
|
||||
|
||||
@ -0,0 +1,17 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct spvDescriptorSetBuffer0
|
||||
{
|
||||
texture2d<float, access::write> uImage [[id(1)]];
|
||||
texture2d<float> 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));
|
||||
}
|
||||
|
||||
@ -0,0 +1,94 @@
|
||||
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
||||
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
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<float> uTexture [[id(2)]];
|
||||
sampler uTextureSmplr [[id(3)]];
|
||||
constant UBO* v_90 [[id(5)]];
|
||||
array<texture2d<float>, 2> uTextures [[id(6)]];
|
||||
array<sampler, 2> uTexturesSmplr [[id(8)]];
|
||||
};
|
||||
|
||||
struct spvDescriptorSetBuffer1
|
||||
{
|
||||
array<texture2d<float>, 4> uTexture2 [[id(3)]];
|
||||
device SSBO* v_60 [[id(7)]];
|
||||
const device SSBOs* ssbos [[id(8)]][2];
|
||||
array<sampler, 2> 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<float> uTexture, thread const sampler uTextureSmplr, thread float2& vUV, thread const array<texture2d<float>, 4> uTexture2, thread const array<sampler, 2> uSampler, thread const array<texture2d<float>, 2> uTextures, thread const array<sampler, 2> 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<float> uTexture, thread const sampler uTextureSmplr, thread float2& vUV, thread const array<texture2d<float>, 4> uTexture2, thread const array<sampler, 2> uSampler, thread const array<texture2d<float>, 2> uTextures, thread const array<sampler, 2> 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;
|
||||
}
|
||||
|
||||
@ -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;
|
||||
}
|
||||
@ -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));
|
||||
}
|
||||
61
shaders-msl/frag/argument-buffers.msl2.argument.frag
Normal file
61
shaders-msl/frag/argument-buffers.msl2.argument.frag
Normal file
@ -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;
|
||||
}
|
||||
@ -1404,8 +1404,9 @@ struct Meta
|
||||
{
|
||||
uint32_t packed_type = 0;
|
||||
bool packed = false;
|
||||
uint32_t ib_member_index = static_cast<uint32_t>(-1);
|
||||
uint32_t ib_member_index = ~(0u);
|
||||
uint32_t ib_orig_id = 0;
|
||||
uint32_t argument_buffer_id = ~(0u);
|
||||
} extended;
|
||||
};
|
||||
|
||||
|
||||
@ -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;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@ -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);
|
||||
|
||||
|
||||
@ -470,6 +470,10 @@ spvc_result spvc_compiler_options_set_uint(spvc_compiler_options options, spvc_c
|
||||
options->msl.platform = static_cast<CompilerMSL::Options::Platform>(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<CompilerMSL *>(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)
|
||||
|
||||
@ -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,
|
||||
|
||||
575
spirv_msl.cpp
575
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<SPIRVariable>(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<SPIRVariable>(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<Resource> resources;
|
||||
|
||||
ir.for_each_typed_id<SPIRVariable>([&](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<SPIRVariable>();
|
||||
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<SPIRVariable>([&](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<SPIRVariable>(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<Resource> resources;
|
||||
|
||||
ir.for_each_typed_id<SPIRVariable>([&](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<SPIRVariable>(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<SPIRType>(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<SPIRType>(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<SPIRType>(type.parent_type), id));
|
||||
type_name = join(get_type_address_space(type, id), " ", type_to_glsl(get<SPIRType>(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<Resource> resources_in_set[kMaxArgumentBuffers];
|
||||
|
||||
ir.for_each_typed_id<SPIRVariable>([&](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<SPIRType>(type_id);
|
||||
buffer_type.storage = StorageClassUniform;
|
||||
buffer_type.basetype = SPIRType::Struct;
|
||||
set_name(type_id, join("spvDescriptorSetBuffer", desc_set));
|
||||
|
||||
auto &ptr_type = set<SPIRType>(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<SPIRVariable>(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<SPIRType>(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<SPIRType>(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++;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@ -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<uint32_t, MSLConstexprSampler> constexpr_samplers;
|
||||
std::vector<uint32_t> 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);
|
||||
|
||||
|
||||
@ -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)
|
||||
|
||||
|
||||
Loading…
Reference in New Issue
Block a user