From 4df759d8e1545e0f984b4eca9b2eb66c4626988e Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=D0=91=D1=80=D0=B0=D0=BD=D0=B8=D0=BC=D0=B8=D1=80=20=D0=9A?= =?UTF-8?q?=D0=B0=D1=80=D0=B0=D1=9F=D0=B8=D1=9B?= Date: Fri, 9 Aug 2019 20:33:38 -0700 Subject: [PATCH] Updated spirv-cross. --- 3rdparty/spirv-cross/main.cpp | 10 +- .../asm/frag/line-directive.line.asm.frag | 12 +- .../asm/frag/line-directive.line.asm.frag | 12 +- .../shaders-msl/comp/basic.dispatchbase.comp | 38 +++ .../comp/basic.dispatchbase.msl11.comp | 34 +++ .../phi-temporary-copy-loop-variable.asm.comp | 4 - ...etch-no-sampler.no-samplerless.asm.vk.frag | 13 + ...h-no-sampler.no-samplerless.asm.vk.frag.vk | 14 ++ ...uery-no-sampler.no-samplerless.vk.asm.frag | 6 + ...y-no-sampler.no-samplerless.vk.asm.frag.vk | 6 + .../asm/frag/line-directive.line.asm.frag | 12 +- ...op-body-dominator-continue-access.asm.frag | 9 +- .../inliner-dominator-inside-loop.asm.frag | 12 +- ...-frexp-scalar-access-chain-output.asm.frag | 17 ++ .../shaders-msl/comp/basic.dispatchbase.comp | 41 +++ .../comp/basic.dispatchbase.msl11.comp | 37 +++ .../asm/frag/do-while-continue-phi.asm.frag | 37 +++ .../inliner-dominator-inside-loop.asm.frag | 12 +- .../phi-temporary-copy-loop-variable.asm.comp | 4 - ...etch-no-sampler.no-samplerless.asm.vk.frag | 38 +++ ...h-no-sampler.no-samplerless.asm.vk.frag.vk | 37 +++ ...uery-no-sampler.no-samplerless.vk.asm.frag | 13 + ...y-no-sampler.no-samplerless.vk.asm.frag.vk | 14 ++ ...-frexp-scalar-access-chain-output.asm.frag | 36 +++ .../shaders-msl/comp/basic.dispatchbase.comp | 29 +++ .../comp/basic.dispatchbase.msl11.comp | 29 +++ .../asm/frag/do-while-continue-phi.asm.frag | 64 +++++ ...etch-no-sampler.no-samplerless.asm.vk.frag | 163 ++++++++++++ ...uery-no-sampler.no-samplerless.vk.asm.frag | 57 +++++ 3rdparty/spirv-cross/spirv_common.hpp | 4 + 3rdparty/spirv-cross/spirv_cross.cpp | 33 ++- 3rdparty/spirv-cross/spirv_cross.hpp | 3 + 3rdparty/spirv-cross/spirv_glsl.cpp | 64 +++-- 3rdparty/spirv-cross/spirv_glsl.hpp | 1 - 3rdparty/spirv-cross/spirv_msl.cpp | 234 +++++++++++------- 3rdparty/spirv-cross/spirv_msl.hpp | 13 +- 3rdparty/spirv-cross/test_shaders.py | 4 + 37 files changed, 993 insertions(+), 173 deletions(-) create mode 100644 3rdparty/spirv-cross/reference/opt/shaders-msl/comp/basic.dispatchbase.comp create mode 100644 3rdparty/spirv-cross/reference/opt/shaders-msl/comp/basic.dispatchbase.msl11.comp create mode 100644 3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag create mode 100644 3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag.vk create mode 100644 3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag create mode 100644 3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag.vk create mode 100644 3rdparty/spirv-cross/reference/shaders-msl-no-opt/asm/frag/modf-frexp-scalar-access-chain-output.asm.frag create mode 100644 3rdparty/spirv-cross/reference/shaders-msl/comp/basic.dispatchbase.comp create mode 100644 3rdparty/spirv-cross/reference/shaders-msl/comp/basic.dispatchbase.msl11.comp create mode 100644 3rdparty/spirv-cross/reference/shaders-no-opt/asm/frag/do-while-continue-phi.asm.frag create mode 100644 3rdparty/spirv-cross/reference/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag create mode 100644 3rdparty/spirv-cross/reference/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag.vk create mode 100644 3rdparty/spirv-cross/reference/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag create mode 100644 3rdparty/spirv-cross/reference/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag.vk create mode 100644 3rdparty/spirv-cross/shaders-msl-no-opt/asm/frag/modf-frexp-scalar-access-chain-output.asm.frag create mode 100644 3rdparty/spirv-cross/shaders-msl/comp/basic.dispatchbase.comp create mode 100644 3rdparty/spirv-cross/shaders-msl/comp/basic.dispatchbase.msl11.comp create mode 100644 3rdparty/spirv-cross/shaders-no-opt/asm/frag/do-while-continue-phi.asm.frag create mode 100644 3rdparty/spirv-cross/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag create mode 100644 3rdparty/spirv-cross/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag diff --git a/3rdparty/spirv-cross/main.cpp b/3rdparty/spirv-cross/main.cpp index 00db1be1f..3d97247e1 100644 --- a/3rdparty/spirv-cross/main.cpp +++ b/3rdparty/spirv-cross/main.cpp @@ -516,8 +516,10 @@ struct CLIArguments bool msl_texture_buffer_native = false; bool msl_multiview = false; bool msl_view_index_from_device_index = false; + bool msl_dispatch_base = false; bool glsl_emit_push_constant_as_ubo = false; bool glsl_emit_ubo_as_plain_uniforms = false; + bool vulkan_glsl_disable_ext_samplerless_texture_functions = false; bool emit_line_directives = false; SmallVector msl_discrete_descriptor_sets; SmallVector pls_in; @@ -584,6 +586,7 @@ static void print_help() "\t[--cpp-interface-name ]\n" "\t[--glsl-emit-push-constant-as-ubo]\n" "\t[--glsl-emit-ubo-as-plain-uniforms]\n" + "\t[--vulkan-glsl-disable-ext-samplerless-texture-functions]\n" "\t[--msl]\n" "\t[--msl-version ]\n" "\t[--msl-capture-output]\n" @@ -596,6 +599,7 @@ static void print_help() "\t[--msl-discrete-descriptor-set ]\n" "\t[--msl-multiview]\n" "\t[--msl-view-index-from-device-index]\n" + "\t[--msl-dispatch-base]\n" "\t[--hlsl]\n" "\t[--reflect]\n" "\t[--shader-model]\n" @@ -756,6 +760,7 @@ static string compile_iteration(const CLIArguments &args, std::vector msl_opts.texture_buffer_native = args.msl_texture_buffer_native; msl_opts.multiview = args.msl_multiview; msl_opts.view_index_from_device_index = args.msl_view_index_from_device_index; + msl_opts.dispatch_base = args.msl_dispatch_base; msl_comp->set_msl_options(msl_opts); for (auto &v : args.msl_discrete_descriptor_sets) msl_comp->add_discrete_descriptor_set(v); @@ -765,7 +770,7 @@ static string compile_iteration(const CLIArguments &args, std::vector else { combined_image_samplers = !args.vulkan_semantics; - if (!args.vulkan_semantics) + if (!args.vulkan_semantics || args.vulkan_glsl_disable_ext_samplerless_texture_functions) build_dummy_sampler = true; compiler.reset(new CompilerGLSL(move(spirv_parser.get_parsed_ir()))); } @@ -1058,6 +1063,8 @@ static int main_inner(int argc, char *argv[]) cbs.add("--metal", [&args](CLIParser &) { args.msl = true; }); // Legacy compatibility cbs.add("--glsl-emit-push-constant-as-ubo", [&args](CLIParser &) { args.glsl_emit_push_constant_as_ubo = true; }); cbs.add("--glsl-emit-ubo-as-plain-uniforms", [&args](CLIParser &) { args.glsl_emit_ubo_as_plain_uniforms = true; }); + cbs.add("--vulkan-glsl-disable-ext-samplerless-texture-functions", + [&args](CLIParser &) { args.vulkan_glsl_disable_ext_samplerless_texture_functions = true; }); cbs.add("--msl", [&args](CLIParser &) { args.msl = true; }); cbs.add("--hlsl", [&args](CLIParser &) { args.hlsl = true; }); cbs.add("--hlsl-enable-compat", [&args](CLIParser &) { args.hlsl_compat = true; }); @@ -1078,6 +1085,7 @@ static int main_inner(int argc, char *argv[]) cbs.add("--msl-multiview", [&args](CLIParser &) { args.msl_multiview = true; }); cbs.add("--msl-view-index-from-device-index", [&args](CLIParser &) { args.msl_view_index_from_device_index = true; }); + cbs.add("--msl-dispatch-base", [&args](CLIParser &) { args.msl_dispatch_base = true; }); 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(); diff --git a/3rdparty/spirv-cross/reference/opt/shaders-hlsl/asm/frag/line-directive.line.asm.frag b/3rdparty/spirv-cross/reference/opt/shaders-hlsl/asm/frag/line-directive.line.asm.frag index b596a8446..93bb69eb3 100644 --- a/3rdparty/spirv-cross/reference/opt/shaders-hlsl/asm/frag/line-directive.line.asm.frag +++ b/3rdparty/spirv-cross/reference/opt/shaders-hlsl/asm/frag/line-directive.line.asm.frag @@ -14,14 +14,12 @@ struct SPIRV_Cross_Output #line 8 "test.frag" void frag_main() { - float _80; #line 8 "test.frag" FragColor = 1.0f; #line 9 "test.frag" FragColor = 2.0f; #line 10 "test.frag" - _80 = vColor; - if (_80 < 0.0f) + if (vColor < 0.0f) { #line 12 "test.frag" FragColor = 3.0f; @@ -31,16 +29,16 @@ void frag_main() #line 16 "test.frag" FragColor = 4.0f; } - for (int _126 = 0; float(_126) < (40.0f + _80); ) + for (int _126 = 0; float(_126) < (40.0f + vColor); ) { #line 21 "test.frag" FragColor += 0.20000000298023223876953125f; #line 22 "test.frag" FragColor += 0.300000011920928955078125f; - _126 += (int(_80) + 5); + _126 += (int(vColor) + 5); continue; } - switch (int(_80)) + switch (int(vColor)) { case 0: { @@ -66,7 +64,7 @@ void frag_main() } for (;;) { - FragColor += (10.0f + _80); + FragColor += (10.0f + vColor); #line 43 "test.frag" if (FragColor < 100.0f) { diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/asm/frag/line-directive.line.asm.frag b/3rdparty/spirv-cross/reference/opt/shaders-msl/asm/frag/line-directive.line.asm.frag index 30018aad4..0d9666d66 100644 --- a/3rdparty/spirv-cross/reference/opt/shaders-msl/asm/frag/line-directive.line.asm.frag +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/asm/frag/line-directive.line.asm.frag @@ -17,14 +17,12 @@ struct main0_in fragment main0_out main0(main0_in in [[stage_in]]) { main0_out out = {}; - float _80; #line 8 "test.frag" out.FragColor = 1.0; #line 9 "test.frag" out.FragColor = 2.0; #line 10 "test.frag" - _80 = in.vColor; - if (_80 < 0.0) + if (in.vColor < 0.0) { #line 12 "test.frag" out.FragColor = 3.0; @@ -34,16 +32,16 @@ fragment main0_out main0(main0_in in [[stage_in]]) #line 16 "test.frag" out.FragColor = 4.0; } - for (int _126 = 0; float(_126) < (40.0 + _80); ) + for (int _126 = 0; float(_126) < (40.0 + in.vColor); ) { #line 21 "test.frag" out.FragColor += 0.20000000298023223876953125; #line 22 "test.frag" out.FragColor += 0.300000011920928955078125; - _126 += (int(_80) + 5); + _126 += (int(in.vColor) + 5); continue; } - switch (int(_80)) + switch (int(in.vColor)) { case 0: { @@ -69,7 +67,7 @@ fragment main0_out main0(main0_in in [[stage_in]]) } for (;;) { - out.FragColor += (10.0 + _80); + out.FragColor += (10.0 + in.vColor); #line 43 "test.frag" if (out.FragColor < 100.0) { diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/basic.dispatchbase.comp b/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/basic.dispatchbase.comp new file mode 100644 index 000000000..ebbc144c7 --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/basic.dispatchbase.comp @@ -0,0 +1,38 @@ +#pragma clang diagnostic ignored "-Wunused-variable" + +#include +#include +#include + +using namespace metal; + +struct SSBO +{ + float4 in_data[1]; +}; + +struct SSBO2 +{ + float4 out_data[1]; +}; + +struct SSBO3 +{ + uint counter; +}; + +constant uint _59_tmp [[function_constant(10)]]; +constant uint _59 = is_function_constant_defined(_59_tmp) ? _59_tmp : 1u; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(_59, 1u, 1u); + +kernel void main0(const device SSBO& _27 [[buffer(0)]], device SSBO2& _49 [[buffer(1)]], device SSBO3& _52 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 spvDispatchBase [[grid_origin]]) +{ + gl_GlobalInvocationID += spvDispatchBase * gl_WorkGroupSize; + float4 _33 = _27.in_data[gl_GlobalInvocationID.x]; + if (dot(_33, float4(1.0, 5.0, 6.0, 2.0)) > 8.19999980926513671875) + { + uint _56 = atomic_fetch_add_explicit((device atomic_uint*)&_52.counter, 1u, memory_order_relaxed); + _49.out_data[_56] = _33; + } +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/basic.dispatchbase.msl11.comp b/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/basic.dispatchbase.msl11.comp new file mode 100644 index 000000000..8c3d25761 --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/basic.dispatchbase.msl11.comp @@ -0,0 +1,34 @@ +#pragma clang diagnostic ignored "-Wunused-variable" + +#include +#include +#include + +using namespace metal; + +struct SSBO +{ + float4 in_data[1]; +}; + +struct SSBO2 +{ + float4 out_data[1]; +}; + +struct SSBO3 +{ + uint counter; +}; + +kernel void main0(constant uint3& spvDispatchBase [[buffer(29)]], const device SSBO& _27 [[buffer(0)]], device SSBO2& _49 [[buffer(1)]], device SSBO3& _52 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + gl_GlobalInvocationID += spvDispatchBase * uint3(1, 1, 1); + float4 _33 = _27.in_data[gl_GlobalInvocationID.x]; + if (dot(_33, float4(1.0, 5.0, 6.0, 2.0)) > 8.19999980926513671875) + { + uint _56 = atomic_fetch_add_explicit((device atomic_uint*)&_52.counter, 1u, memory_order_relaxed); + _49.out_data[_56] = _33; + } +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders/asm/comp/phi-temporary-copy-loop-variable.asm.comp b/3rdparty/spirv-cross/reference/opt/shaders/asm/comp/phi-temporary-copy-loop-variable.asm.comp index 9ae8d6fd7..9266982b7 100644 --- a/3rdparty/spirv-cross/reference/opt/shaders/asm/comp/phi-temporary-copy-loop-variable.asm.comp +++ b/3rdparty/spirv-cross/reference/opt/shaders/asm/comp/phi-temporary-copy-loop-variable.asm.comp @@ -14,10 +14,6 @@ void main() { break; } - else - { - continue; - } continue; } imageStore(outImageTexture, ivec2(gl_GlobalInvocationID.xy), vec4(float(_30 - 1), float(_30), 1.0, 1.0)); diff --git a/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag b/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag new file mode 100644 index 000000000..452fd6fb9 --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag @@ -0,0 +1,13 @@ +#version 450 + +uniform sampler2D SPIRV_Cross_CombinedSampledImageSPIRV_Cross_DummySampler; +uniform sampler2D SPIRV_Cross_CombinedSampledImageSampler; + +layout(location = 0) out vec4 _entryPointOutput; + +void main() +{ + ivec2 _152 = ivec3(int(gl_FragCoord.x * 1280.0), int(gl_FragCoord.y * 720.0), 0).xy; + _entryPointOutput = ((texelFetch(SPIRV_Cross_CombinedSampledImageSPIRV_Cross_DummySampler, _152, 0) + texelFetch(SPIRV_Cross_CombinedSampledImageSPIRV_Cross_DummySampler, _152, 0)) + texture(SPIRV_Cross_CombinedSampledImageSampler, gl_FragCoord.xy)) + texture(SPIRV_Cross_CombinedSampledImageSampler, gl_FragCoord.xy); +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag.vk b/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag.vk new file mode 100644 index 000000000..23acab0b1 --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag.vk @@ -0,0 +1,14 @@ +#version 450 + +layout(set = 0, binding = 0) uniform sampler Sampler; +layout(set = 0, binding = 0) uniform texture2D SampledImage; +layout(set = 0, binding = 0) uniform sampler SPIRV_Cross_DummySampler; + +layout(location = 0) out vec4 _entryPointOutput; + +void main() +{ + ivec2 _152 = ivec3(int(gl_FragCoord.x * 1280.0), int(gl_FragCoord.y * 720.0), 0).xy; + _entryPointOutput = ((texelFetch(sampler2D(SampledImage, SPIRV_Cross_DummySampler), _152, 0) + texelFetch(sampler2D(SampledImage, SPIRV_Cross_DummySampler), _152, 0)) + texture(sampler2D(SampledImage, Sampler), gl_FragCoord.xy)) + texture(sampler2D(SampledImage, Sampler), gl_FragCoord.xy); +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag b/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag new file mode 100644 index 000000000..05ce10adf --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag @@ -0,0 +1,6 @@ +#version 450 + +void main() +{ +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag.vk b/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag.vk new file mode 100644 index 000000000..05ce10adf --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag.vk @@ -0,0 +1,6 @@ +#version 450 + +void main() +{ +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/line-directive.line.asm.frag b/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/line-directive.line.asm.frag index 30be934fc..74eb62bbe 100644 --- a/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/line-directive.line.asm.frag +++ b/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/line-directive.line.asm.frag @@ -7,14 +7,12 @@ layout(location = 0) in float vColor; #line 8 "test.frag" void main() { - float _80; #line 8 "test.frag" FragColor = 1.0; #line 9 "test.frag" FragColor = 2.0; #line 10 "test.frag" - _80 = vColor; - if (_80 < 0.0) + if (vColor < 0.0) { #line 12 "test.frag" FragColor = 3.0; @@ -24,16 +22,16 @@ void main() #line 16 "test.frag" FragColor = 4.0; } - for (int _126 = 0; float(_126) < (40.0 + _80); ) + for (int _126 = 0; float(_126) < (40.0 + vColor); ) { #line 21 "test.frag" FragColor += 0.20000000298023223876953125; #line 22 "test.frag" FragColor += 0.300000011920928955078125; - _126 += (int(_80) + 5); + _126 += (int(vColor) + 5); continue; } - switch (int(_80)) + switch (int(vColor)) { case 0: { @@ -59,7 +57,7 @@ void main() } for (;;) { - FragColor += (10.0 + _80); + FragColor += (10.0 + vColor); #line 43 "test.frag" if (FragColor < 100.0) { diff --git a/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/loop-body-dominator-continue-access.asm.frag b/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/loop-body-dominator-continue-access.asm.frag index ad7913294..10a5e4780 100644 --- a/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/loop-body-dominator-continue-access.asm.frag +++ b/3rdparty/spirv-cross/reference/opt/shaders/asm/frag/loop-body-dominator-continue-access.asm.frag @@ -46,14 +46,7 @@ void main() _231 = true; break; } - else - { - uint _204 = _227 + uint(1); - _227 = _204; - continue; - } - uint _204 = _227 + uint(1); - _227 = _204; + _227++; continue; } else diff --git a/3rdparty/spirv-cross/reference/shaders-msl-no-opt/asm/frag/inliner-dominator-inside-loop.asm.frag b/3rdparty/spirv-cross/reference/shaders-msl-no-opt/asm/frag/inliner-dominator-inside-loop.asm.frag index f63c9ab5c..168129ae8 100644 --- a/3rdparty/spirv-cross/reference/shaders-msl-no-opt/asm/frag/inliner-dominator-inside-loop.asm.frag +++ b/3rdparty/spirv-cross/reference/shaders-msl-no-opt/asm/frag/inliner-dominator-inside-loop.asm.frag @@ -137,7 +137,7 @@ fragment main0_out main0(main0_in in [[stage_in]], constant CB0& _19 [[buffer(0) float2 _166 = in.IN_Uv_EdgeDistance1.xy * 1.0; bool _173; float4 _193; - do + for (;;) { _173 = 0.0 == 0.0; if (_173) @@ -153,9 +153,9 @@ fragment main0_out main0(main0_in in [[stage_in]], constant CB0& _19 [[buffer(0) } _193 = _192; break; - } while (false); + } float4 _220; - do + for (;;) { if (_173) { @@ -170,7 +170,7 @@ fragment main0_out main0(main0_in in [[stage_in]], constant CB0& _19 [[buffer(0) } _220 = _219; break; - } while (false); + } float2 _223 = float2(1.0); float2 _224 = (_220.wy * 2.0) - _223; float3 _232 = float3(_224, sqrt(fast::clamp(1.0 + dot(-_224, _224), 0.0, 1.0))); @@ -181,7 +181,7 @@ fragment main0_out main0(main0_in in [[stage_in]], constant CB0& _19 [[buffer(0) float3 _256 = float3(_255.x, _255.y, _253.z); float3 _271 = ((in.IN_Color.xyz * (_193 * 1.0).xyz) * (1.0 + (_256.x * 0.300000011920928955078125))) * (StudsMapTexture.sample(StudsMapSampler, _156.UvStuds).x * 2.0); float4 _298; - do + for (;;) { if (0.75 == 0.0) { @@ -196,7 +196,7 @@ fragment main0_out main0(main0_in in [[stage_in]], constant CB0& _19 [[buffer(0) } _298 = _297; break; - } while (false); + } float2 _303 = mix(float2(0.800000011920928955078125, 120.0), (_298.xy * float2(2.0, 256.0)) + float2(0.0, 0.00999999977648258209228515625), float2(_165)); Surface _304 = _125; _304.albedo = _271; diff --git a/3rdparty/spirv-cross/reference/shaders-msl-no-opt/asm/frag/modf-frexp-scalar-access-chain-output.asm.frag b/3rdparty/spirv-cross/reference/shaders-msl-no-opt/asm/frag/modf-frexp-scalar-access-chain-output.asm.frag new file mode 100644 index 000000000..910c8fa73 --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-msl-no-opt/asm/frag/modf-frexp-scalar-access-chain-output.asm.frag @@ -0,0 +1,17 @@ +#include +#include + +using namespace metal; + +fragment void main0() +{ + float3 col; + int2 _18; + float _23; + float _21 = modf(0.1500000059604644775390625, _23); + col.x = _23; + int _24; + float _22 = frexp(0.1500000059604644775390625, _24); + _18.y = _24; +} + diff --git a/3rdparty/spirv-cross/reference/shaders-msl/comp/basic.dispatchbase.comp b/3rdparty/spirv-cross/reference/shaders-msl/comp/basic.dispatchbase.comp new file mode 100644 index 000000000..92d517cff --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-msl/comp/basic.dispatchbase.comp @@ -0,0 +1,41 @@ +#pragma clang diagnostic ignored "-Wunused-variable" + +#include +#include +#include + +using namespace metal; + +struct SSBO +{ + float4 in_data[1]; +}; + +struct SSBO2 +{ + float4 out_data[1]; +}; + +struct SSBO3 +{ + uint counter; +}; + +constant uint _59_tmp [[function_constant(10)]]; +constant uint _59 = is_function_constant_defined(_59_tmp) ? _59_tmp : 1u; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(_59, 1u, 1u); + +kernel void main0(const device SSBO& _27 [[buffer(0)]], device SSBO2& _49 [[buffer(1)]], device SSBO3& _52 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 spvDispatchBase [[grid_origin]]) +{ + gl_GlobalInvocationID += spvDispatchBase * gl_WorkGroupSize; + gl_WorkGroupID += spvDispatchBase; + uint ident = gl_GlobalInvocationID.x; + uint workgroup = gl_WorkGroupID.x; + float4 idata = _27.in_data[ident]; + if (dot(idata, float4(1.0, 5.0, 6.0, 2.0)) > 8.19999980926513671875) + { + uint _56 = atomic_fetch_add_explicit((device atomic_uint*)&_52.counter, 1u, memory_order_relaxed); + _49.out_data[_56] = idata; + } +} + diff --git a/3rdparty/spirv-cross/reference/shaders-msl/comp/basic.dispatchbase.msl11.comp b/3rdparty/spirv-cross/reference/shaders-msl/comp/basic.dispatchbase.msl11.comp new file mode 100644 index 000000000..084518a52 --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-msl/comp/basic.dispatchbase.msl11.comp @@ -0,0 +1,37 @@ +#pragma clang diagnostic ignored "-Wunused-variable" + +#include +#include +#include + +using namespace metal; + +struct SSBO +{ + float4 in_data[1]; +}; + +struct SSBO2 +{ + float4 out_data[1]; +}; + +struct SSBO3 +{ + uint counter; +}; + +kernel void main0(constant uint3& spvDispatchBase [[buffer(29)]], const device SSBO& _27 [[buffer(0)]], device SSBO2& _49 [[buffer(1)]], device SSBO3& _52 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) +{ + gl_GlobalInvocationID += spvDispatchBase * uint3(1, 1, 1); + gl_WorkGroupID += spvDispatchBase; + uint ident = gl_GlobalInvocationID.x; + uint workgroup = gl_WorkGroupID.x; + float4 idata = _27.in_data[ident]; + if (dot(idata, float4(1.0, 5.0, 6.0, 2.0)) > 8.19999980926513671875) + { + uint _56 = atomic_fetch_add_explicit((device atomic_uint*)&_52.counter, 1u, memory_order_relaxed); + _49.out_data[_56] = idata; + } +} + diff --git a/3rdparty/spirv-cross/reference/shaders-no-opt/asm/frag/do-while-continue-phi.asm.frag b/3rdparty/spirv-cross/reference/shaders-no-opt/asm/frag/do-while-continue-phi.asm.frag new file mode 100644 index 000000000..2024c302e --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-no-opt/asm/frag/do-while-continue-phi.asm.frag @@ -0,0 +1,37 @@ +#version 310 es +precision mediump float; +precision highp int; + +layout(location = 0) out highp vec4 _GLF_color; + +void main() +{ + for (;;) + { + bool _32; + for (;;) + { + if (gl_FragCoord.x != gl_FragCoord.x) + { + _32 = true; + break; + } + if (false) + { + continue; + } + else + { + _32 = false; + break; + } + } + if (_32) + { + break; + } + _GLF_color = vec4(1.0, 0.0, 0.0, 1.0); + break; + } +} + diff --git a/3rdparty/spirv-cross/reference/shaders-no-opt/asm/frag/inliner-dominator-inside-loop.asm.frag b/3rdparty/spirv-cross/reference/shaders-no-opt/asm/frag/inliner-dominator-inside-loop.asm.frag index 98116cfdc..29653cbb4 100644 --- a/3rdparty/spirv-cross/reference/shaders-no-opt/asm/frag/inliner-dominator-inside-loop.asm.frag +++ b/3rdparty/spirv-cross/reference/shaders-no-opt/asm/frag/inliner-dominator-inside-loop.asm.frag @@ -130,7 +130,7 @@ void main() vec2 _166 = IN_Uv_EdgeDistance1.xy * 1.0; bool _173; vec4 _193; - do + for (;;) { _173 = 0.0 == 0.0; if (_173) @@ -146,9 +146,9 @@ void main() } _193 = _192; break; - } while (false); + } vec4 _220; - do + for (;;) { if (_173) { @@ -163,7 +163,7 @@ void main() } _220 = _219; break; - } while (false); + } vec2 _223 = vec2(1.0); vec2 _224 = (_220.wy * 2.0) - _223; vec3 _232 = vec3(_224, sqrt(clamp(1.0 + dot(-_224, _224), 0.0, 1.0))); @@ -174,7 +174,7 @@ void main() vec3 _256 = vec3(_255.x, _255.y, _253.z); vec3 _271 = ((IN_Color.xyz * (_193 * 1.0).xyz) * (1.0 + (_256.x * 0.300000011920928955078125))) * (texture(SPIRV_Cross_CombinedStudsMapTextureStudsMapSampler, _156.UvStuds).x * 2.0); vec4 _298; - do + for (;;) { if (0.75 == 0.0) { @@ -189,7 +189,7 @@ void main() } _298 = _297; break; - } while (false); + } vec2 _303 = mix(vec2(0.800000011920928955078125, 120.0), (_298.xy * vec2(2.0, 256.0)) + vec2(0.0, 0.00999999977648258209228515625), vec2(_165)); Surface _304 = _125; _304.albedo = _271; diff --git a/3rdparty/spirv-cross/reference/shaders/asm/comp/phi-temporary-copy-loop-variable.asm.comp b/3rdparty/spirv-cross/reference/shaders/asm/comp/phi-temporary-copy-loop-variable.asm.comp index 9ae8d6fd7..9266982b7 100644 --- a/3rdparty/spirv-cross/reference/shaders/asm/comp/phi-temporary-copy-loop-variable.asm.comp +++ b/3rdparty/spirv-cross/reference/shaders/asm/comp/phi-temporary-copy-loop-variable.asm.comp @@ -14,10 +14,6 @@ void main() { break; } - else - { - continue; - } continue; } imageStore(outImageTexture, ivec2(gl_GlobalInvocationID.xy), vec4(float(_30 - 1), float(_30), 1.0, 1.0)); diff --git a/3rdparty/spirv-cross/reference/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag b/3rdparty/spirv-cross/reference/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag new file mode 100644 index 000000000..60bb78aa5 --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag @@ -0,0 +1,38 @@ +#version 450 + +uniform sampler2D SPIRV_Cross_CombinedparamSPIRV_Cross_DummySampler; +uniform sampler2D SPIRV_Cross_CombinedSampledImageSPIRV_Cross_DummySampler; +uniform sampler2D SPIRV_Cross_CombinedparamSampler; +uniform sampler2D SPIRV_Cross_CombinedSampledImageSampler; + +layout(location = 0) out vec4 _entryPointOutput; + +vec4 sample_fetch(ivec3 UV, sampler2D SPIRV_Cross_CombinedtexSPIRV_Cross_DummySampler) +{ + return texelFetch(SPIRV_Cross_CombinedtexSPIRV_Cross_DummySampler, UV.xy, UV.z); +} + +vec4 sample_sampler(vec2 UV, sampler2D SPIRV_Cross_CombinedtexSampler) +{ + return texture(SPIRV_Cross_CombinedtexSampler, UV); +} + +vec4 _main(vec4 xIn) +{ + ivec3 coord = ivec3(int(xIn.x * 1280.0), int(xIn.y * 720.0), 0); + ivec3 param = coord; + vec4 value = sample_fetch(param, SPIRV_Cross_CombinedparamSPIRV_Cross_DummySampler); + value += texelFetch(SPIRV_Cross_CombinedSampledImageSPIRV_Cross_DummySampler, coord.xy, coord.z); + vec2 param_1 = xIn.xy; + value += sample_sampler(param_1, SPIRV_Cross_CombinedparamSampler); + value += texture(SPIRV_Cross_CombinedSampledImageSampler, xIn.xy); + return value; +} + +void main() +{ + vec4 xIn = gl_FragCoord; + vec4 param = xIn; + _entryPointOutput = _main(param); +} + diff --git a/3rdparty/spirv-cross/reference/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag.vk b/3rdparty/spirv-cross/reference/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag.vk new file mode 100644 index 000000000..e4d9fc454 --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag.vk @@ -0,0 +1,37 @@ +#version 450 + +layout(set = 0, binding = 0) uniform sampler Sampler; +layout(set = 0, binding = 0) uniform texture2D SampledImage; +layout(set = 0, binding = 0) uniform sampler SPIRV_Cross_DummySampler; + +layout(location = 0) out vec4 _entryPointOutput; + +vec4 sample_fetch(texture2D tex, ivec3 UV) +{ + return texelFetch(sampler2D(tex, SPIRV_Cross_DummySampler), UV.xy, UV.z); +} + +vec4 sample_sampler(texture2D tex, vec2 UV) +{ + return texture(sampler2D(tex, Sampler), UV); +} + +vec4 _main(vec4 xIn) +{ + ivec3 coord = ivec3(int(xIn.x * 1280.0), int(xIn.y * 720.0), 0); + ivec3 param = coord; + vec4 value = sample_fetch(SampledImage, param); + value += texelFetch(sampler2D(SampledImage, SPIRV_Cross_DummySampler), coord.xy, coord.z); + vec2 param_1 = xIn.xy; + value += sample_sampler(SampledImage, param_1); + value += texture(sampler2D(SampledImage, Sampler), xIn.xy); + return value; +} + +void main() +{ + vec4 xIn = gl_FragCoord; + vec4 param = xIn; + _entryPointOutput = _main(param); +} + diff --git a/3rdparty/spirv-cross/reference/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag b/3rdparty/spirv-cross/reference/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag new file mode 100644 index 000000000..2040dd1af --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag @@ -0,0 +1,13 @@ +#version 450 + +uniform sampler2D SPIRV_Cross_CombineduSampler2DSPIRV_Cross_DummySampler; +uniform sampler2DMS SPIRV_Cross_CombineduSampler2DMSSPIRV_Cross_DummySampler; + +void main() +{ + ivec2 b = textureSize(SPIRV_Cross_CombineduSampler2DSPIRV_Cross_DummySampler, 0); + ivec2 c = textureSize(SPIRV_Cross_CombineduSampler2DMSSPIRV_Cross_DummySampler); + int l1 = textureQueryLevels(SPIRV_Cross_CombineduSampler2DSPIRV_Cross_DummySampler); + int s0 = textureSamples(SPIRV_Cross_CombineduSampler2DMSSPIRV_Cross_DummySampler); +} + diff --git a/3rdparty/spirv-cross/reference/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag.vk b/3rdparty/spirv-cross/reference/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag.vk new file mode 100644 index 000000000..828d2a872 --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag.vk @@ -0,0 +1,14 @@ +#version 450 + +layout(set = 0, binding = 0) uniform texture2D uSampler2D; +layout(set = 0, binding = 0) uniform texture2DMS uSampler2DMS; +layout(set = 0, binding = 0) uniform sampler SPIRV_Cross_DummySampler; + +void main() +{ + ivec2 b = textureSize(sampler2D(uSampler2D, SPIRV_Cross_DummySampler), 0); + ivec2 c = textureSize(sampler2DMS(uSampler2DMS, SPIRV_Cross_DummySampler)); + int l1 = textureQueryLevels(sampler2D(uSampler2D, SPIRV_Cross_DummySampler)); + int s0 = textureSamples(sampler2DMS(uSampler2DMS, SPIRV_Cross_DummySampler)); +} + diff --git a/3rdparty/spirv-cross/shaders-msl-no-opt/asm/frag/modf-frexp-scalar-access-chain-output.asm.frag b/3rdparty/spirv-cross/shaders-msl-no-opt/asm/frag/modf-frexp-scalar-access-chain-output.asm.frag new file mode 100644 index 000000000..707fa550b --- /dev/null +++ b/3rdparty/spirv-cross/shaders-msl-no-opt/asm/frag/modf-frexp-scalar-access-chain-output.asm.frag @@ -0,0 +1,36 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 7 +; Bound: 17 +; Schema: 0 + OpCapability Shader + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint Fragment %main "main" + OpExecutionMode %main OriginUpperLeft + OpSource ESSL 310 + OpName %main "main" + OpName %col "col" + %void = OpTypeVoid + %3 = OpTypeFunction %void + %float = OpTypeFloat 32 +%_ptr_Function_float = OpTypePointer Function %float +%float_0_150000006 = OpConstant %float 0.150000006 + %v3float = OpTypeVector %float 3 +%_ptr_Function_v3float = OpTypePointer Function %v3float + %int = OpTypeInt 32 1 + %int_0 = OpConstant %int 0 + %int_1 = OpConstant %int 1 + %v2int = OpTypeVector %int 2 +%_ptr_Function_v2int = OpTypePointer Function %v2int +%_ptr_Function_int = OpTypePointer Function %int + %main = OpFunction %void None %3 + %5 = OpLabel + %col = OpVariable %_ptr_Function_v3float Function + %icol = OpVariable %_ptr_Function_v2int Function + %ptr_x = OpAccessChain %_ptr_Function_float %col %int_0 + %ptr_y = OpAccessChain %_ptr_Function_int %icol %int_1 + %16 = OpExtInst %float %1 Modf %float_0_150000006 %ptr_x + %17 = OpExtInst %float %1 Frexp %float_0_150000006 %ptr_y + OpReturn + OpFunctionEnd diff --git a/3rdparty/spirv-cross/shaders-msl/comp/basic.dispatchbase.comp b/3rdparty/spirv-cross/shaders-msl/comp/basic.dispatchbase.comp new file mode 100644 index 000000000..2c873468c --- /dev/null +++ b/3rdparty/spirv-cross/shaders-msl/comp/basic.dispatchbase.comp @@ -0,0 +1,29 @@ +#version 310 es +layout(local_size_x_id = 10) in; + +layout(std430, binding = 0) readonly buffer SSBO +{ + vec4 in_data[]; +}; + +layout(std430, binding = 1) writeonly buffer SSBO2 +{ + vec4 out_data[]; +}; + +layout(std430, binding = 2) buffer SSBO3 +{ + uint counter; +}; + +void main() +{ + uint ident = gl_GlobalInvocationID.x; + uint workgroup = gl_WorkGroupID.x; + vec4 idata = in_data[ident]; + if (dot(idata, vec4(1.0, 5.0, 6.0, 2.0)) > 8.2) + { + out_data[atomicAdd(counter, 1u)] = idata; + } +} + diff --git a/3rdparty/spirv-cross/shaders-msl/comp/basic.dispatchbase.msl11.comp b/3rdparty/spirv-cross/shaders-msl/comp/basic.dispatchbase.msl11.comp new file mode 100644 index 000000000..91453332a --- /dev/null +++ b/3rdparty/spirv-cross/shaders-msl/comp/basic.dispatchbase.msl11.comp @@ -0,0 +1,29 @@ +#version 310 es +layout(local_size_x = 1) in; + +layout(std430, binding = 0) readonly buffer SSBO +{ + vec4 in_data[]; +}; + +layout(std430, binding = 1) writeonly buffer SSBO2 +{ + vec4 out_data[]; +}; + +layout(std430, binding = 2) buffer SSBO3 +{ + uint counter; +}; + +void main() +{ + uint ident = gl_GlobalInvocationID.x; + uint workgroup = gl_WorkGroupID.x; + vec4 idata = in_data[ident]; + if (dot(idata, vec4(1.0, 5.0, 6.0, 2.0)) > 8.2) + { + out_data[atomicAdd(counter, 1u)] = idata; + } +} + diff --git a/3rdparty/spirv-cross/shaders-no-opt/asm/frag/do-while-continue-phi.asm.frag b/3rdparty/spirv-cross/shaders-no-opt/asm/frag/do-while-continue-phi.asm.frag new file mode 100644 index 000000000..97400dfb1 --- /dev/null +++ b/3rdparty/spirv-cross/shaders-no-opt/asm/frag/do-while-continue-phi.asm.frag @@ -0,0 +1,64 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 7 +; Bound: 42 +; Schema: 0 + OpCapability Shader + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint Fragment %main "main" %gl_FragCoord %_GLF_color + OpExecutionMode %main OriginUpperLeft + OpSource ESSL 310 + OpName %main "main" + OpName %gl_FragCoord "gl_FragCoord" + OpName %_GLF_color "_GLF_color" + OpDecorate %gl_FragCoord BuiltIn FragCoord + OpDecorate %_GLF_color Location 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 +%_ptr_Input_v4float = OpTypePointer Input %v4float +%gl_FragCoord = OpVariable %_ptr_Input_v4float Input + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 +%_ptr_Input_float = OpTypePointer Input %float + %bool = OpTypeBool + %false = OpConstantFalse %bool +%_ptr_Output_v4float = OpTypePointer Output %v4float + %_GLF_color = OpVariable %_ptr_Output_v4float Output + %float_1 = OpConstant %float 1 + %float_0 = OpConstant %float 0 + %31 = OpConstantComposite %v4float %float_1 %float_0 %float_0 %float_1 + %true = OpConstantTrue %bool + %main = OpFunction %void None %3 + %5 = OpLabel + OpBranch %33 + %33 = OpLabel + OpLoopMerge %32 %35 None + OpBranch %6 + %6 = OpLabel + OpLoopMerge %8 %24 None + OpBranch %7 + %7 = OpLabel + %17 = OpAccessChain %_ptr_Input_float %gl_FragCoord %uint_0 + %18 = OpLoad %float %17 + %22 = OpFOrdNotEqual %bool %18 %18 + OpSelectionMerge %24 None + OpBranchConditional %22 %23 %24 + %23 = OpLabel + OpBranch %8 + %24 = OpLabel + OpBranchConditional %false %6 %8 + %8 = OpLabel + %41 = OpPhi %bool %true %23 %false %24 + OpSelectionMerge %39 None + OpBranchConditional %41 %32 %39 + %39 = OpLabel + OpStore %_GLF_color %31 + OpBranch %32 + %35 = OpLabel + OpBranch %33 + %32 = OpLabel + OpReturn + OpFunctionEnd diff --git a/3rdparty/spirv-cross/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag b/3rdparty/spirv-cross/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag new file mode 100644 index 000000000..a3d64c09d --- /dev/null +++ b/3rdparty/spirv-cross/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag @@ -0,0 +1,163 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 2 +; Bound: 113 +; Schema: 0 + OpCapability Shader + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint Fragment %main "main" %xIn_1 %_entryPointOutput + OpExecutionMode %main OriginUpperLeft + OpSource HLSL 500 + OpName %main "main" + OpName %sample_fetch_t21_vi3_ "sample_fetch(t21;vi3;" + OpName %tex "tex" + OpName %UV "UV" + OpName %sample_sampler_t21_vf2_ "sample_sampler(t21;vf2;" + OpName %tex_0 "tex" + OpName %UV_0 "UV" + OpName %_main_vf4_ "@main(vf4;" + OpName %xIn "xIn" + OpName %Sampler "Sampler" + OpName %coord "coord" + OpName %value "value" + OpName %SampledImage "SampledImage" + OpName %param "param" + OpName %param_0 "param" + OpName %param_1 "param" + OpName %param_2 "param" + OpName %xIn_0 "xIn" + OpName %xIn_1 "xIn" + OpName %_entryPointOutput "@entryPointOutput" + OpName %param_3 "param" + OpDecorate %Sampler DescriptorSet 0 + OpDecorate %Sampler Binding 0 + OpDecorate %SampledImage DescriptorSet 0 + OpDecorate %SampledImage Binding 0 + OpDecorate %xIn_1 BuiltIn FragCoord + OpDecorate %_entryPointOutput Location 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %float = OpTypeFloat 32 + %7 = OpTypeImage %float 2D 0 0 0 1 Unknown +%_ptr_Function_7 = OpTypePointer Function %7 + %int = OpTypeInt 32 1 + %v3int = OpTypeVector %int 3 +%_ptr_Function_v3int = OpTypePointer Function %v3int + %v4float = OpTypeVector %float 4 + %13 = OpTypeFunction %v4float %_ptr_Function_7 %_ptr_Function_v3int + %v2float = OpTypeVector %float 2 +%_ptr_Function_v2float = OpTypePointer Function %v2float + %20 = OpTypeFunction %v4float %_ptr_Function_7 %_ptr_Function_v2float +%_ptr_Function_v4float = OpTypePointer Function %v4float + %26 = OpTypeFunction %v4float %_ptr_Function_v4float + %v2int = OpTypeVector %int 2 + %uint = OpTypeInt 32 0 + %uint_2 = OpConstant %uint 2 +%_ptr_Function_int = OpTypePointer Function %int + %43 = OpTypeSampler +%_ptr_UniformConstant_43 = OpTypePointer UniformConstant %43 + %Sampler = OpVariable %_ptr_UniformConstant_43 UniformConstant + %47 = OpTypeSampledImage %7 + %uint_0 = OpConstant %uint 0 +%_ptr_Function_float = OpTypePointer Function %float + %float_1280 = OpConstant %float 1280 + %uint_1 = OpConstant %uint 1 + %float_720 = OpConstant %float 720 + %int_0 = OpConstant %int 0 +%_ptr_UniformConstant_7 = OpTypePointer UniformConstant %7 +%SampledImage = OpVariable %_ptr_UniformConstant_7 UniformConstant +%_ptr_Input_v4float = OpTypePointer Input %v4float + %xIn_1 = OpVariable %_ptr_Input_v4float Input +%_ptr_Output_v4float = OpTypePointer Output %v4float +%_entryPointOutput = OpVariable %_ptr_Output_v4float Output + %main = OpFunction %void None %3 + %5 = OpLabel + %xIn_0 = OpVariable %_ptr_Function_v4float Function + %param_3 = OpVariable %_ptr_Function_v4float Function + %107 = OpLoad %v4float %xIn_1 + OpStore %xIn_0 %107 + %111 = OpLoad %v4float %xIn_0 + OpStore %param_3 %111 + %112 = OpFunctionCall %v4float %_main_vf4_ %param_3 + OpStore %_entryPointOutput %112 + OpReturn + OpFunctionEnd +%sample_fetch_t21_vi3_ = OpFunction %v4float None %13 + %tex = OpFunctionParameter %_ptr_Function_7 + %UV = OpFunctionParameter %_ptr_Function_v3int + %17 = OpLabel + %30 = OpLoad %7 %tex + %32 = OpLoad %v3int %UV + %33 = OpVectorShuffle %v2int %32 %32 0 1 + %37 = OpAccessChain %_ptr_Function_int %UV %uint_2 + %38 = OpLoad %int %37 + %39 = OpImageFetch %v4float %30 %33 Lod %38 + OpReturnValue %39 + OpFunctionEnd +%sample_sampler_t21_vf2_ = OpFunction %v4float None %20 + %tex_0 = OpFunctionParameter %_ptr_Function_7 + %UV_0 = OpFunctionParameter %_ptr_Function_v2float + %24 = OpLabel + %42 = OpLoad %7 %tex_0 + %46 = OpLoad %43 %Sampler + %48 = OpSampledImage %47 %42 %46 + %49 = OpLoad %v2float %UV_0 + %50 = OpImageSampleImplicitLod %v4float %48 %49 + OpReturnValue %50 + OpFunctionEnd + %_main_vf4_ = OpFunction %v4float None %26 + %xIn = OpFunctionParameter %_ptr_Function_v4float + %29 = OpLabel + %coord = OpVariable %_ptr_Function_v3int Function + %value = OpVariable %_ptr_Function_v4float Function + %param = OpVariable %_ptr_Function_7 Function + %param_0 = OpVariable %_ptr_Function_v3int Function + %param_1 = OpVariable %_ptr_Function_7 Function + %param_2 = OpVariable %_ptr_Function_v2float Function + %56 = OpAccessChain %_ptr_Function_float %xIn %uint_0 + %57 = OpLoad %float %56 + %59 = OpFMul %float %57 %float_1280 + %60 = OpConvertFToS %int %59 + %62 = OpAccessChain %_ptr_Function_float %xIn %uint_1 + %63 = OpLoad %float %62 + %65 = OpFMul %float %63 %float_720 + %66 = OpConvertFToS %int %65 + %68 = OpCompositeConstruct %v3int %60 %66 %int_0 + OpStore %coord %68 + %73 = OpLoad %7 %SampledImage + OpStore %param %73 + %75 = OpLoad %v3int %coord + OpStore %param_0 %75 + %76 = OpFunctionCall %v4float %sample_fetch_t21_vi3_ %param %param_0 + OpStore %value %76 + %77 = OpLoad %7 %SampledImage + %78 = OpLoad %v3int %coord + %79 = OpVectorShuffle %v2int %78 %78 0 1 + %80 = OpAccessChain %_ptr_Function_int %coord %uint_2 + %81 = OpLoad %int %80 + %82 = OpImageFetch %v4float %77 %79 Lod %81 + %83 = OpLoad %v4float %value + %84 = OpFAdd %v4float %83 %82 + OpStore %value %84 + %86 = OpLoad %7 %SampledImage + OpStore %param_1 %86 + %88 = OpLoad %v4float %xIn + %89 = OpVectorShuffle %v2float %88 %88 0 1 + OpStore %param_2 %89 + %90 = OpFunctionCall %v4float %sample_sampler_t21_vf2_ %param_1 %param_2 + %91 = OpLoad %v4float %value + %92 = OpFAdd %v4float %91 %90 + OpStore %value %92 + %93 = OpLoad %7 %SampledImage + %94 = OpLoad %43 %Sampler + %95 = OpSampledImage %47 %93 %94 + %96 = OpLoad %v4float %xIn + %97 = OpVectorShuffle %v2float %96 %96 0 1 + %98 = OpImageSampleImplicitLod %v4float %95 %97 + %99 = OpLoad %v4float %value + %100 = OpFAdd %v4float %99 %98 + OpStore %value %100 + %101 = OpLoad %v4float %value + OpReturnValue %101 + OpFunctionEnd diff --git a/3rdparty/spirv-cross/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag b/3rdparty/spirv-cross/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag new file mode 100644 index 000000000..a232bd489 --- /dev/null +++ b/3rdparty/spirv-cross/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag @@ -0,0 +1,57 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 6 +; Bound: 36 +; Schema: 0 + OpCapability Shader + OpCapability ImageQuery + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint Fragment %main "main" + OpExecutionMode %main OriginUpperLeft + OpSource GLSL 450 + OpName %main "main" + OpName %b "b" + OpName %uSampler2D "uSampler2D" + OpName %c "c" + OpName %uSampler2DMS "uSampler2DMS" + OpName %l1 "l1" + OpName %s0 "s0" + OpDecorate %uSampler2D DescriptorSet 0 + OpDecorate %uSampler2D Binding 0 + OpDecorate %uSampler2DMS DescriptorSet 0 + OpDecorate %uSampler2DMS Binding 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %v2int = OpTypeVector %int 2 +%_ptr_Function_v2int = OpTypePointer Function %v2int + %float = OpTypeFloat 32 + %11 = OpTypeImage %float 2D 0 0 0 1 Unknown +%_ptr_UniformConstant_12 = OpTypePointer UniformConstant %11 + %uSampler2D = OpVariable %_ptr_UniformConstant_12 UniformConstant + %int_0 = OpConstant %int 0 + %20 = OpTypeImage %float 2D 0 0 1 1 Unknown +%_ptr_UniformConstant_21 = OpTypePointer UniformConstant %20 +%uSampler2DMS = OpVariable %_ptr_UniformConstant_21 UniformConstant +%_ptr_Function_int = OpTypePointer Function %int + %main = OpFunction %void None %3 + %5 = OpLabel + %b = OpVariable %_ptr_Function_v2int Function + %c = OpVariable %_ptr_Function_v2int Function + %l1 = OpVariable %_ptr_Function_int Function + %s0 = OpVariable %_ptr_Function_int Function + %15 = OpLoad %11 %uSampler2D + %18 = OpImageQuerySizeLod %v2int %15 %int_0 + OpStore %b %18 + %24 = OpLoad %20 %uSampler2DMS + %26 = OpImageQuerySize %v2int %24 + OpStore %c %26 + %29 = OpLoad %11 %uSampler2D + %31 = OpImageQueryLevels %int %29 + OpStore %l1 %31 + %33 = OpLoad %20 %uSampler2DMS + %35 = OpImageQuerySamples %int %33 + OpStore %s0 %35 + OpReturn + OpFunctionEnd diff --git a/3rdparty/spirv-cross/spirv_common.hpp b/3rdparty/spirv-cross/spirv_common.hpp index bc626436a..3db55cf5c 100644 --- a/3rdparty/spirv-cross/spirv_common.hpp +++ b/3rdparty/spirv-cross/spirv_common.hpp @@ -1433,6 +1433,10 @@ enum ExtendedDecorations // Marks a buffer block for using explicit offsets (GLSL/HLSL). SPIRVCrossDecorationExplicitOffset, + // Apply to a variable in the Input storage class; marks it as holding the base group passed to vkCmdDispatchBase(). + // In MSL, this is used to adjust the WorkgroupId and GlobalInvocationId variables. + SPIRVCrossDecorationBuiltInDispatchBase, + SPIRVCrossDecorationCount }; diff --git a/3rdparty/spirv-cross/spirv_cross.cpp b/3rdparty/spirv-cross/spirv_cross.cpp index 715e514d5..3a9670ed4 100644 --- a/3rdparty/spirv-cross/spirv_cross.cpp +++ b/3rdparty/spirv-cross/spirv_cross.cpp @@ -1501,6 +1501,12 @@ SPIRBlock::ContinueBlockType Compiler::continue_block_type(const SPIRBlock &bloc const auto *true_block = maybe_get(block.true_block); const auto *merge_block = maybe_get(dominator.merge_block); + // If we need to flush Phi in this block, we cannot have a DoWhile loop. + bool flush_phi_to_false = false_block && flush_phi_required(block.self, block.false_block); + bool flush_phi_to_true = true_block && flush_phi_required(block.self, block.true_block); + if (flush_phi_to_false || flush_phi_to_true) + return SPIRBlock::ComplexLoop; + bool positive_do_while = block.true_block == dominator.self && (block.false_block == dominator.merge_block || (false_block && merge_block && execution_is_noop(*false_block, *merge_block))); @@ -3288,10 +3294,11 @@ void Compiler::analyze_variable_scope(SPIRFunction &entry, AnalyzeVariableScopeA { builder.add_block(block); - // If a temporary is used in more than one block, we might have to lift continue block - // access up to loop header like we did for variables. if (blocks.size() != 1 && is_continue(block)) { + // The risk here is that inner loop can dominate the continue block. + // Any temporary we access in the continue block must be declared before the loop. + // This is moot for complex loops however. auto &loop_header_block = get(ir.continue_block_to_loop_header[block]); assert(loop_header_block.merge == SPIRBlock::MergeLoop); @@ -3299,14 +3306,17 @@ void Compiler::analyze_variable_scope(SPIRFunction &entry, AnalyzeVariableScopeA if (!loop_header_block.complex_continue) builder.add_block(loop_header_block.self); } - else if (blocks.size() != 1 && is_single_block_loop(block)) - { - // Awkward case, because the loop header is also the continue block. - force_temporary = true; - } } uint32_t dominating_block = builder.get_dominator(); + + if (blocks.size() != 1 && is_single_block_loop(dominating_block)) + { + // Awkward case, because the loop header is also the continue block, + // so hoisting to loop header does not help. + force_temporary = true; + } + if (dominating_block) { // If we touch a variable in the dominating block, this is the expected setup. @@ -4246,3 +4256,12 @@ bool Compiler::type_is_array_of_pointers(const SPIRType &type) const // If parent type has same pointer depth, we must have an array of pointers. return type.pointer_depth == get(type.parent_type).pointer_depth; } + +bool Compiler::flush_phi_required(uint32_t from, uint32_t to) const +{ + auto &child = get(to); + for (auto &phi : child.phi_variables) + if (phi.parent == from) + return true; + return false; +} diff --git a/3rdparty/spirv-cross/spirv_cross.hpp b/3rdparty/spirv-cross/spirv_cross.hpp index deb757afe..ca75dc66d 100644 --- a/3rdparty/spirv-cross/spirv_cross.hpp +++ b/3rdparty/spirv-cross/spirv_cross.hpp @@ -816,6 +816,7 @@ protected: std::unordered_set forwarded_temporaries; std::unordered_set suppressed_usage_tracking; std::unordered_set hoisted_temporaries; + std::unordered_set forced_invariant_temporaries; Bitset active_input_builtins; Bitset active_output_builtins; @@ -972,6 +973,8 @@ protected: bool reflection_ssbo_instance_name_is_significant() const; std::string get_remapped_declared_block_name(uint32_t id, bool fallback_prefer_instance_name) const; + bool flush_phi_required(uint32_t from, uint32_t to) const; + private: // Used only to implement the old deprecated get_entry_point() interface. const SPIREntryPoint &get_first_entry_point(const std::string &name) const; diff --git a/3rdparty/spirv-cross/spirv_glsl.cpp b/3rdparty/spirv-cross/spirv_glsl.cpp index 4fe7d5889..bf33d83b0 100644 --- a/3rdparty/spirv-cross/spirv_glsl.cpp +++ b/3rdparty/spirv-cross/spirv_glsl.cpp @@ -4654,16 +4654,16 @@ void CompilerGLSL::emit_sampled_image_op(uint32_t result_type, uint32_t result_i { emit_binary_func_op(result_type, result_id, image_id, samp_id, type_to_glsl(get(result_type), result_id).c_str()); - - // Make sure to suppress usage tracking and any expression invalidation. - // It is illegal to create temporaries of opaque types. - forwarded_temporaries.erase(result_id); } else { // Make sure to suppress usage tracking. It is illegal to create temporaries of opaque types. emit_op(result_type, result_id, to_combined_image_sampler(image_id, samp_id), true, true); } + + // Make sure to suppress usage tracking and any expression invalidation. + // It is illegal to create temporaries of opaque types. + forwarded_temporaries.erase(result_id); } static inline bool image_opcode_is_sample_no_dref(Op op) @@ -4976,10 +4976,19 @@ std::string CompilerGLSL::convert_separate_image_to_expression(uint32_t id) { if (options.vulkan_semantics) { - // Newer glslang supports this extension to deal with texture2D as argument to texture functions. if (dummy_sampler_id) - SPIRV_CROSS_THROW("Vulkan GLSL should not have a dummy sampler for combining."); - require_extension_internal("GL_EXT_samplerless_texture_functions"); + { + // Don't need to consider Shadow state since the dummy sampler is always non-shadow. + auto sampled_type = type; + sampled_type.basetype = SPIRType::SampledImage; + return join(type_to_glsl(sampled_type), "(", to_expression(id), ", ", + to_expression(dummy_sampler_id), ")"); + } + else + { + // Newer glslang supports this extension to deal with texture2D as argument to texture functions. + require_extension_internal("GL_EXT_samplerless_texture_functions"); + } } else { @@ -5277,7 +5286,6 @@ void CompilerGLSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop, case GLSLstd450ModfStruct: { - forced_temporaries.insert(id); auto &type = get(result_type); emit_uninitialized_temporary_expression(result_type, id); statement(to_expression(id), ".", to_member_name(type, 0), " = ", "modf(", to_expression(args[0]), ", ", @@ -5417,7 +5425,6 @@ void CompilerGLSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop, case GLSLstd450FrexpStruct: { - forced_temporaries.insert(id); auto &type = get(result_type); emit_uninitialized_temporary_expression(result_type, id); statement(to_expression(id), ".", to_member_name(type, 0), " = ", "frexp(", to_expression(args[0]), ", ", @@ -7544,14 +7551,16 @@ void CompilerGLSL::disallow_forwarding_in_expression_chain(const SPIRExpression // Allow trivially forwarded expressions like OpLoad or trivial shuffles, // these will be marked as having suppressed usage tracking. // Our only concern is to make sure arithmetic operations are done in similar ways. - if (expression_is_forwarded(expr.self) && !expression_suppresses_usage_tracking(expr.self)) + if (expression_is_forwarded(expr.self) && !expression_suppresses_usage_tracking(expr.self) && + forced_invariant_temporaries.count(expr.self) == 0) { forced_temporaries.insert(expr.self); + forced_invariant_temporaries.insert(expr.self); force_recompile(); - } - for (auto &dependent : expr.expression_dependencies) - disallow_forwarding_in_expression_chain(get(dependent)); + for (auto &dependent : expr.expression_dependencies) + disallow_forwarding_in_expression_chain(get(dependent)); + } } void CompilerGLSL::handle_store_to_invariant_variable(uint32_t store_id, uint32_t value_id) @@ -8525,7 +8534,6 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) uint32_t result_id = ops[1]; uint32_t op0 = ops[2]; uint32_t op1 = ops[3]; - forced_temporaries.insert(result_id); auto &type = get(result_type); emit_uninitialized_temporary_expression(result_type, result_id); const char *op = opcode == OpUMulExtended ? "umulExtended" : "imulExtended"; @@ -11060,15 +11068,6 @@ void CompilerGLSL::emit_fixup() } } -bool CompilerGLSL::flush_phi_required(uint32_t from, uint32_t to) -{ - auto &child = get(to); - for (auto &phi : child.phi_variables) - if (phi.parent == from) - return true; - return false; -} - void CompilerGLSL::flush_phi(uint32_t from, uint32_t to) { auto &child = get(to); @@ -11238,10 +11237,17 @@ void CompilerGLSL::branch(uint32_t from, uint32_t to) void CompilerGLSL::branch(uint32_t from, uint32_t cond, uint32_t true_block, uint32_t false_block) { - // If we branch directly to a selection merge target, we don't really need a code path. + auto &from_block = get(from); + uint32_t merge_block = from_block.merge == SPIRBlock::MergeSelection ? from_block.next_block : 0; + + // If we branch directly to a selection merge target, we don't need a code path. + // This covers both merge out of if () / else () as well as a break for switch blocks. bool true_sub = !is_conditional(true_block); bool false_sub = !is_conditional(false_block); + bool true_block_is_selection_merge = true_block == merge_block; + bool false_block_is_selection_merge = false_block == merge_block; + if (true_sub) { emit_block_hints(get(from)); @@ -11250,7 +11256,11 @@ void CompilerGLSL::branch(uint32_t from, uint32_t cond, uint32_t true_block, uin branch(from, true_block); end_scope(); - if (false_sub || is_continue(false_block) || is_break(false_block)) + // If we merge to continue, we handle that explicitly in emit_block_chain(), + // so there is no need to branch to it directly here. + // break; is required to handle ladder fallthrough cases, so keep that in for now, even + // if we could potentially handle it in emit_block_chain(). + if (false_sub || (!false_block_is_selection_merge && is_continue(false_block)) || is_break(false_block)) { statement("else"); begin_scope(); @@ -11265,7 +11275,7 @@ void CompilerGLSL::branch(uint32_t from, uint32_t cond, uint32_t true_block, uin end_scope(); } } - else if (false_sub && !true_sub) + else if (false_sub) { // Only need false path, use negative conditional. emit_block_hints(get(from)); @@ -11274,7 +11284,7 @@ void CompilerGLSL::branch(uint32_t from, uint32_t cond, uint32_t true_block, uin branch(from, false_block); end_scope(); - if (is_continue(true_block) || is_break(true_block)) + if ((!true_block_is_selection_merge && is_continue(true_block)) || is_break(true_block)) { statement("else"); begin_scope(); diff --git a/3rdparty/spirv-cross/spirv_glsl.hpp b/3rdparty/spirv-cross/spirv_glsl.hpp index 3a4f48593..f5582f630 100644 --- a/3rdparty/spirv-cross/spirv_glsl.hpp +++ b/3rdparty/spirv-cross/spirv_glsl.hpp @@ -436,7 +436,6 @@ protected: void branch_to_continue(uint32_t from, uint32_t to); void branch(uint32_t from, uint32_t cond, uint32_t true_block, uint32_t false_block); void flush_phi(uint32_t from, uint32_t to); - bool flush_phi_required(uint32_t from, uint32_t to); void flush_variable_declaration(uint32_t id); void flush_undeclared_variables(SPIRBlock &block); void emit_variable_temporary_copies(const SPIRVariable &var); diff --git a/3rdparty/spirv-cross/spirv_msl.cpp b/3rdparty/spirv-cross/spirv_msl.cpp index 9f33034ff..10dc85497 100644 --- a/3rdparty/spirv-cross/spirv_msl.cpp +++ b/3rdparty/spirv-cross/spirv_msl.cpp @@ -107,8 +107,11 @@ void CompilerMSL::build_implicit_builtins() active_input_builtins.get(BuiltInSubgroupGtMask)); bool need_multiview = get_execution_model() == ExecutionModelVertex && !msl_options.view_index_from_device_index && (msl_options.multiview || active_input_builtins.get(BuiltInViewIndex)); + bool need_dispatch_base = + msl_options.dispatch_base && get_execution_model() == ExecutionModelGLCompute && + (active_input_builtins.get(BuiltInWorkgroupId) || active_input_builtins.get(BuiltInGlobalInvocationId)); if (need_subpass_input || need_sample_pos || need_subgroup_mask || need_vertex_params || need_tesc_params || - need_multiview || needs_subgroup_invocation_id) + need_multiview || need_dispatch_base || needs_subgroup_invocation_id) { bool has_frag_coord = false; bool has_sample_id = false; @@ -121,6 +124,7 @@ void CompilerMSL::build_implicit_builtins() bool has_subgroup_invocation_id = false; bool has_subgroup_size = false; bool has_view_idx = false; + uint32_t workgroup_id_type = 0; ir.for_each_typed_id([&](uint32_t, SPIRVariable &var) { if (var.storage != StorageClassInput || !ir.meta[var.self].decoration.builtin) @@ -208,6 +212,13 @@ void CompilerMSL::build_implicit_builtins() has_view_idx = true; } } + + // The base workgroup needs to have the same type and vector size + // as the workgroup or invocation ID, so keep track of the type that + // was used. + if (need_dispatch_base && workgroup_id_type == 0 && + (builtin == BuiltInWorkgroupId || builtin == BuiltInGlobalInvocationId)) + workgroup_id_type = var.basetype; }); if (!has_frag_coord && need_subpass_input) @@ -457,6 +468,42 @@ void CompilerMSL::build_implicit_builtins() builtin_subgroup_size_id = var_id; mark_implicit_builtin(StorageClassInput, BuiltInSubgroupSize, var_id); } + + if (need_dispatch_base) + { + uint32_t var_id; + if (msl_options.supports_msl_version(1, 2)) + { + // If we have MSL 1.2, we can (ab)use the [[grid_origin]] builtin + // to convey this information and save a buffer slot. + uint32_t offset = ir.increase_bound_by(1); + var_id = offset; + + set(var_id, workgroup_id_type, StorageClassInput); + set_extended_decoration(var_id, SPIRVCrossDecorationBuiltInDispatchBase); + get_entry_point().interface_variables.push_back(var_id); + } + else + { + // Otherwise, we need to fall back to a good ol' fashioned buffer. + uint32_t offset = ir.increase_bound_by(2); + var_id = offset; + uint32_t type_id = offset + 1; + + SPIRType var_type = get(workgroup_id_type); + var_type.storage = StorageClassUniform; + set(type_id, var_type); + + set(var_id, type_id, StorageClassUniform); + // This should never match anything. + set_decoration(var_id, DecorationDescriptorSet, ~(5u)); + set_decoration(var_id, DecorationBinding, msl_options.indirect_params_buffer_index); + set_extended_decoration(var_id, SPIRVCrossDecorationResourceIndexPrimary, + msl_options.indirect_params_buffer_index); + } + set_name(var_id, "spvDispatchBase"); + builtin_dispatch_base_id = var_id; + } } if (needs_swizzle_buffer_def) @@ -802,6 +849,8 @@ string CompilerMSL::compile() active_interface_variables.insert(view_mask_buffer_id); if (builtin_layer_id) active_interface_variables.insert(builtin_layer_id); + if (builtin_dispatch_base_id && !msl_options.supports_msl_version(1, 2)) + active_interface_variables.insert(builtin_dispatch_base_id); // Create structs to hold input, output and uniform variables. // Do output first to ensure out. is declared at top of entry function. @@ -4600,10 +4649,8 @@ void CompilerMSL::emit_instruction(const Instruction &instruction) uint32_t result_id = ops[1]; uint32_t op0 = ops[2]; uint32_t op1 = ops[3]; - forced_temporaries.insert(result_id); auto &type = get(result_type); - statement(variable_decl(type, to_name(result_id)), ";"); - set(result_id, to_name(result_id), result_type, true); + emit_uninitialized_temporary_expression(result_type, result_id); auto &res_type = get(type.member_types[1]); if (opcode == OpIAddCarry) @@ -4632,10 +4679,8 @@ void CompilerMSL::emit_instruction(const Instruction &instruction) uint32_t result_id = ops[1]; uint32_t op0 = ops[2]; uint32_t op1 = ops[3]; - forced_temporaries.insert(result_id); auto &type = get(result_type); - statement(variable_decl(type, to_name(result_id)), ";"); - set(result_id, to_name(result_id), result_type, true); + emit_uninitialized_temporary_expression(result_type, result_id); statement(to_expression(result_id), ".", to_member_name(type, 0), " = ", to_enclosed_expression(op0), " * ", to_enclosed_expression(op1), ";"); @@ -4917,8 +4962,6 @@ void CompilerMSL::emit_atomic_func_op(uint32_t result_type, uint32_t result_id, uint32_t mem_order_2, bool has_mem_order_2, uint32_t obj, uint32_t op1, bool op1_is_pointer, bool op1_is_literal, uint32_t op2) { - forced_temporaries.insert(result_id); - string exp = string(op) + "("; auto &type = get_pointee_type(expression_type(obj)); @@ -4957,12 +5000,11 @@ void CompilerMSL::emit_atomic_func_op(uint32_t result_type, uint32_t result_id, // the CAS loop, otherwise it will loop infinitely, with the comparison test always failing. // The function updates the comparitor value from the memory value, so the additional // comparison test evaluates the memory value against the expected value. - statement(variable_decl(type, to_name(result_id)), ";"); + emit_uninitialized_temporary_expression(result_type, result_id); statement("do"); begin_scope(); statement(to_name(result_id), " = ", to_expression(op1), ";"); end_scope_decl(join("while (!", exp, " && ", to_name(result_id), " == ", to_enclosed_expression(op1), ")")); - set(result_id, to_name(result_id), result_type, true); } else { @@ -5211,6 +5253,32 @@ void CompilerMSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop, CompilerGLSL::emit_glsl_op(result_type, id, eop, args, count); break; + case GLSLstd450Modf: + case GLSLstd450Frexp: + { + // Special case. If the variable is a scalar access chain, we cannot use it directly. We have to emit a temporary. + auto *ptr = maybe_get(args[1]); + if (ptr && ptr->access_chain && is_scalar(expression_type(args[1]))) + { + register_call_out_argument(args[1]); + forced_temporaries.insert(id); + + // Need to create temporaries and copy over to access chain after. + // We cannot directly take the reference of a vector swizzle in MSL, even if it's scalar ... + uint32_t &tmp_id = extra_sub_expressions[id]; + if (!tmp_id) + tmp_id = ir.increase_bound_by(1); + + uint32_t tmp_type_id = get_pointee_type_id(ptr->expression_type); + emit_uninitialized_temporary_expression(tmp_type_id, tmp_id); + emit_binary_func_op(result_type, id, args[0], tmp_id, eop == GLSLstd450Modf ? "modf" : "frexp"); + statement(to_expression(args[1]), " = ", to_expression(tmp_id), ";"); + } + else + CompilerGLSL::emit_glsl_op(result_type, id, eop, args, count); + break; + } + default: CompilerGLSL::emit_glsl_op(result_type, id, eop, args, count); break; @@ -6527,80 +6595,15 @@ string CompilerMSL::func_type_decl(SPIRType &type) string CompilerMSL::get_argument_address_space(const SPIRVariable &argument) { const auto &type = get(argument.basetype); - Bitset flags; - if (type.basetype == SPIRType::Struct && - (has_decoration(type.self, DecorationBlock) || has_decoration(type.self, DecorationBufferBlock))) - flags = ir.get_buffer_block_flags(argument); - else - flags = get_decoration_bitset(argument.self); - const char *addr_space = nullptr; - - switch (type.storage) - { - case StorageClassWorkgroup: - addr_space = "threadgroup"; - break; - - case StorageClassStorageBuffer: - { - // For arguments from variable pointers, we use the write count deduction, so - // we should not assume any constness here. Only for global SSBOs. - bool readonly = false; - if (has_decoration(type.self, DecorationBlock)) - readonly = flags.get(DecorationNonWritable); - - addr_space = readonly ? "const device" : "device"; - break; - } - - case StorageClassUniform: - case StorageClassUniformConstant: - case StorageClassPushConstant: - if (type.basetype == SPIRType::Struct) - { - bool ssbo = has_decoration(type.self, DecorationBufferBlock); - if (ssbo) - { - bool readonly = flags.get(DecorationNonWritable); - addr_space = readonly ? "const device" : "device"; - } - else - addr_space = "constant"; - break; - } - break; - - case StorageClassFunction: - case StorageClassGeneric: - // No address space for plain values. - addr_space = type.pointer ? "thread" : ""; - break; - - case StorageClassInput: - if (get_execution_model() == ExecutionModelTessellationControl && argument.basevariable == stage_in_ptr_var_id) - addr_space = "threadgroup"; - break; - - case StorageClassOutput: - if (capture_output_to_buffer) - addr_space = "device"; - break; - - default: - break; - } - - if (!addr_space) - addr_space = "thread"; - - return join(flags.get(DecorationVolatile) || flags.get(DecorationCoherent) ? "volatile " : "", addr_space); + return get_type_address_space(type, argument.self, true); } -string CompilerMSL::get_type_address_space(const SPIRType &type, uint32_t id) +string CompilerMSL::get_type_address_space(const SPIRType &type, uint32_t id, bool argument) { // 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 && type.basetype == SPIRType::Struct && + auto *var = maybe_get(id); + if (var && type.basetype == SPIRType::Struct && (has_decoration(type.self, DecorationBlock) || has_decoration(type.self, DecorationBufferBlock))) flags = get_buffer_block_flags(id); else @@ -6614,8 +6617,16 @@ string CompilerMSL::get_type_address_space(const SPIRType &type, uint32_t id) break; case StorageClassStorageBuffer: - addr_space = flags.get(DecorationNonWritable) ? "const device" : "device"; + { + // For arguments from variable pointers, we use the write count deduction, so + // we should not assume any constness here. Only for global SSBOs. + bool readonly = false; + if (!var || has_decoration(type.self, DecorationBlock)) + readonly = flags.get(DecorationNonWritable); + + addr_space = readonly ? "const device" : "device"; break; + } case StorageClassUniform: case StorageClassUniformConstant: @@ -6628,14 +6639,18 @@ string CompilerMSL::get_type_address_space(const SPIRType &type, uint32_t id) else addr_space = "constant"; } - else + else if (!argument) addr_space = "constant"; break; case StorageClassFunction: case StorageClassGeneric: - // No address space for plain values. - addr_space = type.pointer ? "thread" : ""; + break; + + case StorageClassInput: + if (get_execution_model() == ExecutionModelTessellationControl && var && + var->basevariable == stage_in_ptr_var_id) + addr_space = "threadgroup"; break; case StorageClassOutput: @@ -6648,7 +6663,8 @@ string CompilerMSL::get_type_address_space(const SPIRType &type, uint32_t id) } if (!addr_space) - addr_space = "thread"; + // No address space for plain values. + addr_space = type.pointer || (argument && type.basetype == SPIRType::ControlPointArray) ? "thread" : ""; return join(flags.get(DecorationVolatile) || flags.get(DecorationCoherent) ? "volatile " : "", addr_space); } @@ -6748,6 +6764,19 @@ void CompilerMSL::entry_point_args_builtin(string &ep_args) ep_args += "]]"; } } + + if (var.storage == StorageClassInput && + has_extended_decoration(var_id, SPIRVCrossDecorationBuiltInDispatchBase)) + { + // This is a special implicit builtin, not corresponding to any SPIR-V builtin, + // which holds the base that was passed to vkCmdDispatchBase(). If it's present, + // assume we emitted it for a good reason. + assert(msl_options.supports_msl_version(1, 2)); + if (!ep_args.empty()) + ep_args += ", "; + + ep_args += type_to_glsl(get_variable_data_type(var)) + " " + to_expression(var_id) + " [[grid_origin]]"; + } }); // Correct the types of all encountered active builtins. We couldn't do this before @@ -7023,7 +7052,11 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args) default: if (!ep_args.empty()) ep_args += ", "; - ep_args += type_to_glsl(type, var_id) + " " + r.name; + if (!type.pointer) + ep_args += get_type_address_space(get(var.basetype), var_id) + " " + + type_to_glsl(type, var_id) + "& " + r.name; + else + ep_args += type_to_glsl(type, var_id) + " " + r.name; ep_args += " [[buffer(" + convert_to_string(r.index) + ")]]"; break; } @@ -7343,6 +7376,35 @@ void CompilerMSL::fix_up_shader_inputs_outputs() msl_options.device_index, ";"); }); break; + case BuiltInWorkgroupId: + if (!msl_options.dispatch_base || !active_input_builtins.get(BuiltInWorkgroupId)) + break; + + // The vkCmdDispatchBase() command lets the client set the base value + // of WorkgroupId. Metal has no direct equivalent; we must make this + // adjustment ourselves. + entry_func.fixup_hooks_in.push_back([=]() { + statement(to_expression(var_id), " += ", to_dereferenced_expression(builtin_dispatch_base_id), ";"); + }); + break; + case BuiltInGlobalInvocationId: + if (!msl_options.dispatch_base || !active_input_builtins.get(BuiltInGlobalInvocationId)) + break; + + // GlobalInvocationId is defined as LocalInvocationId + WorkgroupId * WorkgroupSize. + // This needs to be adjusted too. + entry_func.fixup_hooks_in.push_back([=]() { + auto &execution = this->get_entry_point(); + uint32_t workgroup_size_id = execution.workgroup_size.constant; + if (workgroup_size_id) + statement(to_expression(var_id), " += ", to_dereferenced_expression(builtin_dispatch_base_id), + " * ", to_expression(workgroup_size_id), ";"); + else + statement(to_expression(var_id), " += ", to_dereferenced_expression(builtin_dispatch_base_id), + " * uint3(", execution.workgroup_size.x, ", ", execution.workgroup_size.y, ", ", + execution.workgroup_size.z, ");"); + }); + break; default: break; } diff --git a/3rdparty/spirv-cross/spirv_msl.hpp b/3rdparty/spirv-cross/spirv_msl.hpp index abd481b33..f0858c9da 100644 --- a/3rdparty/spirv-cross/spirv_msl.hpp +++ b/3rdparty/spirv-cross/spirv_msl.hpp @@ -198,6 +198,7 @@ public: bool tess_domain_origin_lower_left = false; bool multiview = false; bool view_index_from_device_index = false; + bool dispatch_base = false; // Enable use of MSL 2.0 indirect argument buffers. // MSL 2.0 must also be enabled. @@ -225,7 +226,7 @@ public: msl_version = make_msl_version(major, minor, patch); } - bool supports_msl_version(uint32_t major, uint32_t minor = 0, uint32_t patch = 0) + bool supports_msl_version(uint32_t major, uint32_t minor = 0, uint32_t patch = 0) const { return msl_version >= make_msl_version(major, minor, patch); } @@ -276,6 +277,13 @@ public: return msl_options.multiview && !msl_options.view_index_from_device_index; } + // Provide feedback to calling API to allow it to pass a buffer + // containing the dispatch base workgroup ID. + bool needs_dispatch_base_buffer() const + { + return msl_options.dispatch_base && !msl_options.supports_msl_version(1, 2); + } + // Provide feedback to calling API to allow it to pass an output // buffer if the shader needs it. bool needs_output_buffer() const @@ -533,7 +541,7 @@ protected: void ensure_member_packing_rules_msl(SPIRType &ib_type, uint32_t index); bool validate_member_packing_rules_msl(const SPIRType &type, uint32_t index) const; std::string get_argument_address_space(const SPIRVariable &argument); - std::string get_type_address_space(const SPIRType &type, uint32_t id); + std::string get_type_address_space(const SPIRType &type, uint32_t id, bool argument = false); const char *to_restrict(uint32_t id, bool space = true); SPIRType &get_stage_in_struct_type(); SPIRType &get_stage_out_struct_type(); @@ -563,6 +571,7 @@ protected: uint32_t builtin_primitive_id_id = 0; uint32_t builtin_subgroup_invocation_id_id = 0; uint32_t builtin_subgroup_size_id = 0; + uint32_t builtin_dispatch_base_id = 0; uint32_t swizzle_buffer_id = 0; uint32_t buffer_size_buffer_id = 0; uint32_t view_mask_buffer_id = 0; diff --git a/3rdparty/spirv-cross/test_shaders.py b/3rdparty/spirv-cross/test_shaders.py index b3f692533..27466b398 100755 --- a/3rdparty/spirv-cross/test_shaders.py +++ b/3rdparty/spirv-cross/test_shaders.py @@ -207,6 +207,8 @@ def cross_compile_msl(shader, spirv, opt, iterations, paths): msl_args.append('--msl-multiview') if '.viewfromdev.' in shader: msl_args.append('--msl-view-index-from-device-index') + if '.dispatchbase.' in shader: + msl_args.append('--msl-dispatch-base') subprocess.check_call(msl_args) @@ -383,6 +385,8 @@ def cross_compile(shader, vulkan, spirv, invalid_spirv, eliminate, is_legacy, fl extra_args += ['--glsl-emit-push-constant-as-ubo'] if '.line.' in shader: extra_args += ['--emit-line-directives'] + if '.no-samplerless.' in shader: + extra_args += ['--vulkan-glsl-disable-ext-samplerless-texture-functions'] spirv_cross_path = paths.spirv_cross