diff --git a/3rdparty/spirv-cross/.travis.yml b/3rdparty/spirv-cross/.travis.yml index fa14abfde..46b92c8fc 100644 --- a/3rdparty/spirv-cross/.travis.yml +++ b/3rdparty/spirv-cross/.travis.yml @@ -11,43 +11,62 @@ matrix: compiler: gcc env: - GENERATOR="Unix Makefiles" + - ARTIFACT=gcc-trusty-64bit - os: linux dist: trusty compiler: clang env: - GENERATOR="Unix Makefiles" + - ARTIFACT=clang-trusty-64bit - os: osx compiler: clang osx_image: xcode10 env: - GENERATOR="Unix Makefiles" + - ARTIFACT=clang-macos-64bit - os: windows before_install: - choco install python3 - - choco install python2 - - export PATH="/c/Python27:/c/Python27/Scripts:$PATH" - export PATH="/c/Python37:/c/Python37/Scripts:$PATH" env: - GENERATOR="Visual Studio 15 2017" + - ARTIFACT=vs2017-32bit - os: windows before_install: - choco install python3 - - choco install python2 - - export PATH="/c/Python27:/c/Python27/Scripts:$PATH" - export PATH="/c/Python37:/c/Python37/Scripts:$PATH" env: - GENERATOR="Visual Studio 15 2017 Win64" + - ARTIFACT=vs2017-64bit before_script: - - ./checkout_glslang_spirv_tools.sh + - "./checkout_glslang_spirv_tools.sh" script: - if [[ "$TRAVIS_OS_NAME" == "windows" ]]; then PYTHON3=$(which python); fi - if [[ "$TRAVIS_OS_NAME" != "windows" ]]; then PYTHON3=$(which python3); fi - - ./build_glslang_spirv_tools.sh Release + - "./build_glslang_spirv_tools.sh Release" - mkdir build - cd build - cmake .. -DSPIRV_CROSS_SHARED=ON -DCMAKE_INSTALL_PREFIX=output -DCMAKE_BUILD_TYPE=Release -G "${GENERATOR}" -DPYTHON_EXECUTABLE:FILEPATH="${PYTHON3}" -DSPIRV_CROSS_ENABLE_TESTS=ON - cmake --build . --config Release - cmake --build . --config Release --target install - ctest --verbose -C Release + - cd .. + +before_deploy: + - REV=${ARTIFACT}-$(git rev-parse --short=10 HEAD) + - cd build/output + - tar cf spirv-cross-${REV}.tar * + - gzip spirv-cross-${REV}.tar + - cd ../.. + - export FILE_TO_UPLOAD=build/output/spirv-cross-${REV}.tar.gz + +deploy: + provider: releases + api_key: + secure: c7YEOyzhE19TFo76UnbLWk/kikRQxsHsOxzkOqN6Q2aL8joNRw5kmcG84rGd+Rf6isX62cykCzA6qHkyJCv9QTIzcyXnLju17rLvgib7cXDcseaq8x4mFvet2yUxCglthDpFY2M2LB0Aqws71lPeYIrKXa6hCFEh8jO3AWxnaor7O3RYfNZylM9d33HgH6KLT3sDx/cukwBstmKeg7EG9OUnrSvairkPW0W2+jlq3SXPlq/WeVhf8hQs3Yg0BluExGbmLOwe9EaeUpeGuJMyHRxXypnToQv1/KwoScKpap5tYxdNWiwRGZ4lYcmKrjAYVvilTioh654oX5LQpn34mE/oe8Ko9AaATkSaoiisRFp6meWtnB39oFBoL5Yn15DqLQpRXPr1AJsnBXSGAac3aDBO1j4MIqTHmYlYlfRw3n2ZsBaFaTZnv++438SNQ54nkivyoDTIWjoOmYa9+K4mQc3415RDdQmjZTJM+lu+GAlMmNBTVbfNvrbU55Usu9Lo6BZJKKdUMvdBB78kJ5FHvcBlL+eMgmk1pABQY0IZROCt7NztHcv1UmAxoWNxveSFs5glydPNNjNS8bogc4dzBGYG0KMmILbBHihVbY2toA1M9CMdDHdp+LucfDMmzECmYSEmlx0h8win+Jjb74/qpOhaXuUZ0NnzVgCOyeUYuMQ= + file: "${FILE_TO_UPLOAD}" + skip_cleanup: true + on: + tags: true diff --git a/3rdparty/spirv-cross/CMakeLists.txt b/3rdparty/spirv-cross/CMakeLists.txt index 664ba4211..8bca53d3f 100644 --- a/3rdparty/spirv-cross/CMakeLists.txt +++ b/3rdparty/spirv-cross/CMakeLists.txt @@ -265,10 +265,11 @@ if (SPIRV_CROSS_STATIC) endif() endif() +set(spirv-cross-abi-major 0) +set(spirv-cross-abi-minor 9) +set(spirv-cross-abi-patch 0) + if (SPIRV_CROSS_SHARED) - set(spirv-cross-abi-major 0) - set(spirv-cross-abi-minor 6) - 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) set(SPIRV_CROSS_INSTALL_INC_DIR ${CMAKE_INSTALL_PREFIX}/include/spirv_cross) @@ -431,7 +432,10 @@ if (SPIRV_CROSS_CLI) target_compile_options(spirv-cross-c-api-test PRIVATE -std=c89 -Wall -Wextra) endif() add_test(NAME spirv-cross-c-api-test - COMMAND $ ${CMAKE_CURRENT_SOURCE_DIR}/tests-other/c_api_test.spv) + COMMAND $ ${CMAKE_CURRENT_SOURCE_DIR}/tests-other/c_api_test.spv + ${spirv-cross-abi-major} + ${spirv-cross-abi-minor} + ${spirv-cross-abi-patch}) add_test(NAME spirv-cross-small-vector-test COMMAND $) add_test(NAME spirv-cross-test diff --git a/3rdparty/spirv-cross/checkout_glslang_spirv_tools.sh b/3rdparty/spirv-cross/checkout_glslang_spirv_tools.sh index 491c40190..37b99a0e0 100755 --- a/3rdparty/spirv-cross/checkout_glslang_spirv_tools.sh +++ b/3rdparty/spirv-cross/checkout_glslang_spirv_tools.sh @@ -1,8 +1,8 @@ #!/bin/bash -GLSLANG_REV=ef807f4bc543e061f25dbbee6cb64dd5053b2adc -SPIRV_TOOLS_REV=12e4a7b649e6fe28683de9fc352200c82948a1f0 -SPIRV_HEADERS_REV=111a25e4ae45e2b4d7c18415e1d6884712b958c4 +GLSLANG_REV=e291f7a09f6733f6634fe077a228056fabee881e +SPIRV_TOOLS_REV=89fe836fe22c3e5c2a062ebeade012e2c2f0839b +SPIRV_HEADERS_REV=c4f8f65792d4bf2657ca751904c511bbcf2ac77b if [ -z $PROTOCOL ]; then PROTOCOL=git diff --git a/3rdparty/spirv-cross/main.cpp b/3rdparty/spirv-cross/main.cpp index 2cfb99dcf..cd57f176a 100644 --- a/3rdparty/spirv-cross/main.cpp +++ b/3rdparty/spirv-cross/main.cpp @@ -511,6 +511,7 @@ struct CLIArguments bool msl_argument_buffers = false; bool msl_texture_buffer_native = false; bool glsl_emit_push_constant_as_ubo = false; + bool glsl_emit_ubo_as_plain_uniforms = false; SmallVector msl_discrete_descriptor_sets; SmallVector pls_in; SmallVector pls_out; @@ -563,6 +564,7 @@ static void print_help() "\t[--cpp]\n" "\t[--cpp-interface-name ]\n" "\t[--glsl-emit-push-constant-as-ubo]\n" + "\t[--glsl-emit-ubo-as-plain-uniforms]\n" "\t[--msl]\n" "\t[--msl-version ]\n" "\t[--msl-capture-output]\n" @@ -854,6 +856,7 @@ static string compile_iteration(const CLIArguments &args, std::vector opts.vertex.flip_vert_y = args.yflip; opts.vertex.support_nonzero_base_instance = args.support_nonzero_baseinstance; opts.emit_push_constant_as_uniform_buffer = args.glsl_emit_push_constant_as_ubo; + opts.emit_uniform_buffer_as_plain_uniforms = args.glsl_emit_ubo_as_plain_uniforms; compiler->set_common_options(opts); // Set HLSL specific options. @@ -1025,6 +1028,7 @@ static int main_inner(int argc, char *argv[]) cbs.add("--cpp-interface-name", [&args](CLIParser &parser) { args.cpp_interface_name = parser.next_string(); }); 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("--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; }); diff --git a/3rdparty/spirv-cross/reference/opt/shaders-hlsl/comp/ssbo-array-length.comp b/3rdparty/spirv-cross/reference/opt/shaders-hlsl/comp/ssbo-array-length.comp new file mode 100644 index 000000000..2e3df626a --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-hlsl/comp/ssbo-array-length.comp @@ -0,0 +1,15 @@ +RWByteAddressBuffer _11 : register(u1); + +void comp_main() +{ + uint _14; + _11.GetDimensions(_14); + _14 = (_14 - 16) / 16; + _11.Store(0, uint(int(_14))); +} + +[numthreads(1, 1, 1)] +void main() +{ + comp_main(); +} diff --git a/3rdparty/spirv-cross/reference/opt/shaders-hlsl/frag/nonuniform-qualifier.nonuniformresource.sm51.frag b/3rdparty/spirv-cross/reference/opt/shaders-hlsl/frag/nonuniform-qualifier.nonuniformresource.sm51.frag new file mode 100644 index 000000000..544c57053 --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-hlsl/frag/nonuniform-qualifier.nonuniformresource.sm51.frag @@ -0,0 +1,46 @@ +struct UBO_1_1 +{ + float4 v[64]; +}; + +ConstantBuffer ubos[] : register(b0, space3); +ByteAddressBuffer ssbos[] : register(t0, space4); +Texture2D uSamplers[] : register(t0, space0); +SamplerState uSamps[] : register(s0, space2); +Texture2D uCombinedSamplers[] : register(t0, space1); +SamplerState _uCombinedSamplers_sampler[] : register(s0, space1); + +static int vIndex; +static float4 FragColor; +static float2 vUV; + +struct SPIRV_Cross_Input +{ + nointerpolation int vIndex : TEXCOORD0; + float2 vUV : TEXCOORD1; +}; + +struct SPIRV_Cross_Output +{ + float4 FragColor : SV_Target0; +}; + +void frag_main() +{ + int _22 = vIndex + 10; + int _32 = vIndex + 40; + FragColor = uSamplers[NonUniformResourceIndex(_22)].Sample(uSamps[NonUniformResourceIndex(_32)], vUV); + FragColor = uCombinedSamplers[NonUniformResourceIndex(_22)].Sample(_uCombinedSamplers_sampler[NonUniformResourceIndex(_22)], vUV); + FragColor += ubos[NonUniformResourceIndex(vIndex + 20)].v[_32]; + FragColor += asfloat(ssbos[NonUniformResourceIndex(vIndex + 50)].Load4((vIndex + 60) * 16 + 0)); +} + +SPIRV_Cross_Output main(SPIRV_Cross_Input stage_input) +{ + vIndex = stage_input.vIndex; + vUV = stage_input.vUV; + frag_main(); + SPIRV_Cross_Output stage_output; + stage_output.FragColor = FragColor; + return stage_output; +} diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/barriers.comp b/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/barriers.comp index 8822d0806..164cefd7b 100644 --- a/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/barriers.comp +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/barriers.comp @@ -8,13 +8,15 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 1u, 1u); kernel void main0() { threadgroup_barrier(mem_flags::mem_threadgroup); + threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture); + threadgroup_barrier(mem_flags::mem_texture); + threadgroup_barrier(mem_flags::mem_device); + threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture); threadgroup_barrier(mem_flags::mem_threadgroup); - threadgroup_barrier(mem_flags::mem_none); - threadgroup_barrier(mem_flags::mem_threadgroup); - threadgroup_barrier(mem_flags::mem_threadgroup); - threadgroup_barrier(mem_flags::mem_threadgroup); - threadgroup_barrier(mem_flags::mem_none); - threadgroup_barrier(mem_flags::mem_threadgroup); + threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture); + threadgroup_barrier(mem_flags::mem_texture); + threadgroup_barrier(mem_flags::mem_device); + threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture); threadgroup_barrier(mem_flags::mem_threadgroup); } diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/force-recompile-hooks.swizzle.comp b/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/force-recompile-hooks.swizzle.comp index c06d8ba1b..2d9d423e3 100644 --- a/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/force-recompile-hooks.swizzle.comp +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/force-recompile-hooks.swizzle.comp @@ -5,11 +5,6 @@ using namespace metal; -struct spvAux -{ - uint swizzleConst[1]; -}; - enum class spvSwizzle : uint { none = 0, @@ -130,9 +125,9 @@ inline vec spvGatherCompareSwizzle(sampler s, const thread Tex& t, Ts... p return t.gather_compare(s, spvForward(params)...); } -kernel void main0(constant spvAux& spvAuxBuffer [[buffer(30)]], texture2d foo [[texture(0)]], texture2d bar [[texture(1)]], sampler fooSmplr [[sampler(0)]]) +kernel void main0(constant uint* spvSwizzleConstants [[buffer(30)]], texture2d foo [[texture(0)]], texture2d bar [[texture(1)]], sampler fooSmplr [[sampler(0)]]) { - constant uint32_t& fooSwzl = spvAuxBuffer.swizzleConst[0]; + constant uint32_t& fooSwzl = spvSwizzleConstants[0]; bar.write(spvTextureSwizzle(foo.sample(fooSmplr, float2(1.0), level(0.0)), fooSwzl), uint2(int2(0))); } diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/frag/array-of-texture-swizzle.msl2.argument.discrete.swizzle.frag b/3rdparty/spirv-cross/reference/opt/shaders-msl/frag/array-of-texture-swizzle.msl2.argument.discrete.swizzle.frag new file mode 100644 index 000000000..b8e8f9d80 --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/frag/array-of-texture-swizzle.msl2.argument.discrete.swizzle.frag @@ -0,0 +1,156 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct spvDescriptorSetBuffer0 +{ + constant uint* spvSwizzleConstants [[id(0)]]; + array, 4> uSampler0 [[id(1)]]; + array uSampler0Smplr [[id(5)]]; +}; + +struct main0_out +{ + float4 FragColor [[color(0)]]; +}; + +struct main0_in +{ + float2 vUV [[user(locn0)]]; +}; + +enum class spvSwizzle : uint +{ + none = 0, + zero, + one, + red, + green, + blue, + alpha +}; + +template struct spvRemoveReference { typedef T type; }; +template struct spvRemoveReference { typedef T type; }; +template struct spvRemoveReference { typedef T type; }; +template inline constexpr thread T&& spvForward(thread typename spvRemoveReference::type& x) +{ + return static_cast(x); +} +template inline constexpr thread T&& spvForward(thread typename spvRemoveReference::type&& x) +{ + return static_cast(x); +} + +template +inline T spvGetSwizzle(vec x, T c, spvSwizzle s) +{ + switch (s) + { + case spvSwizzle::none: + return c; + case spvSwizzle::zero: + return 0; + case spvSwizzle::one: + return 1; + case spvSwizzle::red: + return x.r; + case spvSwizzle::green: + return x.g; + case spvSwizzle::blue: + return x.b; + case spvSwizzle::alpha: + return x.a; + } +} + +// Wrapper function that swizzles texture samples and fetches. +template +inline vec spvTextureSwizzle(vec x, uint s) +{ + if (!s) + return x; + return vec(spvGetSwizzle(x, x.r, spvSwizzle((s >> 0) & 0xFF)), spvGetSwizzle(x, x.g, spvSwizzle((s >> 8) & 0xFF)), spvGetSwizzle(x, x.b, spvSwizzle((s >> 16) & 0xFF)), spvGetSwizzle(x, x.a, spvSwizzle((s >> 24) & 0xFF))); +} + +template +inline T spvTextureSwizzle(T x, uint s) +{ + return spvTextureSwizzle(vec(x, 0, 0, 1), s).x; +} + +// Wrapper function that swizzles texture gathers. +template +inline vec spvGatherSwizzle(sampler s, const thread Tex& t, Ts... params, component c, uint sw) METAL_CONST_ARG(c) +{ + if (sw) + { + switch (spvSwizzle((sw >> (uint(c) * 8)) & 0xFF)) + { + case spvSwizzle::none: + break; + case spvSwizzle::zero: + return vec(0, 0, 0, 0); + case spvSwizzle::one: + return vec(1, 1, 1, 1); + case spvSwizzle::red: + return t.gather(s, spvForward(params)..., component::x); + case spvSwizzle::green: + return t.gather(s, spvForward(params)..., component::y); + case spvSwizzle::blue: + return t.gather(s, spvForward(params)..., component::z); + case spvSwizzle::alpha: + return t.gather(s, spvForward(params)..., component::w); + } + } + switch (c) + { + case component::x: + return t.gather(s, spvForward(params)..., component::x); + case component::y: + return t.gather(s, spvForward(params)..., component::y); + case component::z: + return t.gather(s, spvForward(params)..., component::z); + case component::w: + return t.gather(s, spvForward(params)..., component::w); + } +} + +// Wrapper function that swizzles depth texture gathers. +template +inline vec spvGatherCompareSwizzle(sampler s, const thread Tex& t, Ts... params, uint sw) +{ + if (sw) + { + switch (spvSwizzle(sw & 0xFF)) + { + case spvSwizzle::none: + case spvSwizzle::red: + break; + case spvSwizzle::zero: + case spvSwizzle::green: + case spvSwizzle::blue: + case spvSwizzle::alpha: + return vec(0, 0, 0, 0); + case spvSwizzle::one: + return vec(1, 1, 1, 1); + } + } + return t.gather_compare(s, spvForward(params)...); +} + +fragment main0_out main0(main0_in in [[stage_in]], constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant uint* spvSwizzleConstants [[buffer(30)]], texture2d uSampler1 [[texture(0)]], sampler uSampler1Smplr [[sampler(0)]]) +{ + main0_out out = {}; + constant uint32_t* spvDescriptorSet0_uSampler0Swzl = &spvDescriptorSet0.spvSwizzleConstants[1]; + constant uint32_t& uSampler1Swzl = spvSwizzleConstants[0]; + out.FragColor = spvTextureSwizzle(spvDescriptorSet0.uSampler0[2].sample(spvDescriptorSet0.uSampler0Smplr[2], in.vUV), spvDescriptorSet0_uSampler0Swzl[2]); + out.FragColor += spvTextureSwizzle(uSampler1.sample(uSampler1Smplr, in.vUV), uSampler1Swzl); + out.FragColor += spvTextureSwizzle(spvDescriptorSet0.uSampler0[1].sample(spvDescriptorSet0.uSampler0Smplr[1], in.vUV), spvDescriptorSet0_uSampler0Swzl[1]); + out.FragColor += spvTextureSwizzle(uSampler1.sample(uSampler1Smplr, in.vUV), uSampler1Swzl); + return out; +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/frag/array-of-texture-swizzle.msl2.swizzle.frag b/3rdparty/spirv-cross/reference/opt/shaders-msl/frag/array-of-texture-swizzle.msl2.swizzle.frag new file mode 100644 index 000000000..8e97cab98 --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/frag/array-of-texture-swizzle.msl2.swizzle.frag @@ -0,0 +1,146 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct main0_out +{ + float4 FragColor [[color(0)]]; +}; + +struct main0_in +{ + float2 vUV [[user(locn0)]]; +}; + +enum class spvSwizzle : uint +{ + none = 0, + zero, + one, + red, + green, + blue, + alpha +}; + +template struct spvRemoveReference { typedef T type; }; +template struct spvRemoveReference { typedef T type; }; +template struct spvRemoveReference { typedef T type; }; +template inline constexpr thread T&& spvForward(thread typename spvRemoveReference::type& x) +{ + return static_cast(x); +} +template inline constexpr thread T&& spvForward(thread typename spvRemoveReference::type&& x) +{ + return static_cast(x); +} + +template +inline T spvGetSwizzle(vec x, T c, spvSwizzle s) +{ + switch (s) + { + case spvSwizzle::none: + return c; + case spvSwizzle::zero: + return 0; + case spvSwizzle::one: + return 1; + case spvSwizzle::red: + return x.r; + case spvSwizzle::green: + return x.g; + case spvSwizzle::blue: + return x.b; + case spvSwizzle::alpha: + return x.a; + } +} + +// Wrapper function that swizzles texture samples and fetches. +template +inline vec spvTextureSwizzle(vec x, uint s) +{ + if (!s) + return x; + return vec(spvGetSwizzle(x, x.r, spvSwizzle((s >> 0) & 0xFF)), spvGetSwizzle(x, x.g, spvSwizzle((s >> 8) & 0xFF)), spvGetSwizzle(x, x.b, spvSwizzle((s >> 16) & 0xFF)), spvGetSwizzle(x, x.a, spvSwizzle((s >> 24) & 0xFF))); +} + +template +inline T spvTextureSwizzle(T x, uint s) +{ + return spvTextureSwizzle(vec(x, 0, 0, 1), s).x; +} + +// Wrapper function that swizzles texture gathers. +template +inline vec spvGatherSwizzle(sampler s, const thread Tex& t, Ts... params, component c, uint sw) METAL_CONST_ARG(c) +{ + if (sw) + { + switch (spvSwizzle((sw >> (uint(c) * 8)) & 0xFF)) + { + case spvSwizzle::none: + break; + case spvSwizzle::zero: + return vec(0, 0, 0, 0); + case spvSwizzle::one: + return vec(1, 1, 1, 1); + case spvSwizzle::red: + return t.gather(s, spvForward(params)..., component::x); + case spvSwizzle::green: + return t.gather(s, spvForward(params)..., component::y); + case spvSwizzle::blue: + return t.gather(s, spvForward(params)..., component::z); + case spvSwizzle::alpha: + return t.gather(s, spvForward(params)..., component::w); + } + } + switch (c) + { + case component::x: + return t.gather(s, spvForward(params)..., component::x); + case component::y: + return t.gather(s, spvForward(params)..., component::y); + case component::z: + return t.gather(s, spvForward(params)..., component::z); + case component::w: + return t.gather(s, spvForward(params)..., component::w); + } +} + +// Wrapper function that swizzles depth texture gathers. +template +inline vec spvGatherCompareSwizzle(sampler s, const thread Tex& t, Ts... params, uint sw) +{ + if (sw) + { + switch (spvSwizzle(sw & 0xFF)) + { + case spvSwizzle::none: + case spvSwizzle::red: + break; + case spvSwizzle::zero: + case spvSwizzle::green: + case spvSwizzle::blue: + case spvSwizzle::alpha: + return vec(0, 0, 0, 0); + case spvSwizzle::one: + return vec(1, 1, 1, 1); + } + } + return t.gather_compare(s, spvForward(params)...); +} + +fragment main0_out main0(main0_in in [[stage_in]], constant uint* spvSwizzleConstants [[buffer(30)]], array, 4> uSampler [[texture(0)]], array uSamplerSmplr [[sampler(0)]]) +{ + main0_out out = {}; + constant uint32_t* uSamplerSwzl = &spvSwizzleConstants[0]; + out.FragColor = spvTextureSwizzle(uSampler[2].sample(uSamplerSmplr[2], in.vUV), uSamplerSwzl[2]); + out.FragColor += spvTextureSwizzle(uSampler[1].sample(uSamplerSmplr[1], in.vUV), uSamplerSwzl[1]); + return out; +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/frag/nonuniform-qualifier.msl2.frag b/3rdparty/spirv-cross/reference/opt/shaders-msl/frag/nonuniform-qualifier.msl2.frag new file mode 100644 index 000000000..510d1ca58 --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/frag/nonuniform-qualifier.msl2.frag @@ -0,0 +1,50 @@ +#include +#include + +using namespace metal; + +struct UBO +{ + float4 v[64]; +}; + +struct SSBO +{ + float4 v[1]; +}; + +struct main0_out +{ + float4 FragColor [[color(0)]]; +}; + +struct main0_in +{ + int vIndex [[user(locn0)]]; + float2 vUV [[user(locn1)]]; +}; + +fragment main0_out main0(main0_in in [[stage_in]], constant UBO* ubos_0 [[buffer(0)]], constant UBO* ubos_1 [[buffer(1)]], const device SSBO* ssbos_0 [[buffer(2)]], const device SSBO* ssbos_1 [[buffer(3)]], array, 8> uSamplers [[texture(0)]], array, 8> uCombinedSamplers [[texture(8)]], array uSamps [[sampler(1)]], array uCombinedSamplersSmplr [[sampler(8)]]) +{ + constant UBO* ubos[] = + { + ubos_0, + ubos_1, + }; + + const device SSBO* ssbos[] = + { + ssbos_0, + ssbos_1, + }; + + main0_out out = {}; + int _24 = in.vIndex + 10; + int _35 = in.vIndex + 40; + out.FragColor = uSamplers[_24].sample(uSamps[_35], in.vUV); + out.FragColor = uCombinedSamplers[_24].sample(uCombinedSamplersSmplr[_24], in.vUV); + out.FragColor += ubos[(in.vIndex + 20)]->v[_35]; + out.FragColor += ssbos[(in.vIndex + 50)]->v[in.vIndex + 60]; + return out; +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.comp b/3rdparty/spirv-cross/reference/opt/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.comp new file mode 100644 index 000000000..908fde045 --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.comp @@ -0,0 +1,92 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct SSBO +{ + float FragColor; +}; + +inline uint4 spvSubgroupBallot(bool value) +{ + simd_vote vote = simd_ballot(value); + // simd_ballot() returns a 64-bit integer-like object, but + // SPIR-V callers expect a uint4. We must convert. + // FIXME: This won't include higher bits if Apple ever supports + // 128 lanes in an SIMD-group. + return uint4((uint)((simd_vote::vote_t)vote & 0xFFFFFFFF), (uint)(((simd_vote::vote_t)vote >> 32) & 0xFFFFFFFF), 0, 0); +} + +inline bool spvSubgroupBallotBitExtract(uint4 ballot, uint bit) +{ + return !!extract_bits(ballot[bit / 32], bit % 32, 1); +} + +inline uint spvSubgroupBallotFindLSB(uint4 ballot) +{ + return select(ctz(ballot.x), select(32 + ctz(ballot.y), select(64 + ctz(ballot.z), select(96 + ctz(ballot.w), uint(-1), ballot.w == 0), ballot.z == 0), ballot.y == 0), ballot.x == 0); +} + +inline uint spvSubgroupBallotFindMSB(uint4 ballot) +{ + return select(128 - (clz(ballot.w) + 1), select(96 - (clz(ballot.z) + 1), select(64 - (clz(ballot.y) + 1), select(32 - (clz(ballot.x) + 1), uint(-1), ballot.x == 0), ballot.y == 0), ballot.z == 0), ballot.w == 0); +} + +inline uint spvSubgroupBallotBitCount(uint4 ballot) +{ + return popcount(ballot.x) + popcount(ballot.y) + popcount(ballot.z) + popcount(ballot.w); +} + +inline uint spvSubgroupBallotInclusiveBitCount(uint4 ballot, uint gl_SubgroupInvocationID) +{ + uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID + 1, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID + 1 - 32, 0)), uint2(0)); + return spvSubgroupBallotBitCount(ballot & mask); +} + +inline uint spvSubgroupBallotExclusiveBitCount(uint4 ballot, uint gl_SubgroupInvocationID) +{ + uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID - 32, 0)), uint2(0)); + return spvSubgroupBallotBitCount(ballot & mask); +} + +template +inline bool spvSubgroupAllEqual(T value) +{ + return simd_all(value == simd_broadcast_first(value)); +} + +template<> +inline bool spvSubgroupAllEqual(bool value) +{ + return simd_all(value) || !simd_any(value); +} + +kernel void main0(device SSBO& _9 [[buffer(0)]], uint gl_NumSubgroups [[simdgroups_per_threadgroup]], uint gl_SubgroupID [[simdgroup_index_in_threadgroup]], uint gl_SubgroupSize [[thread_execution_width]], uint gl_SubgroupInvocationID [[thread_index_in_simdgroup]]) +{ + uint4 gl_SubgroupEqMask = 27 > 32 ? uint4(0, (1 << (gl_SubgroupInvocationID - 32)), uint2(0)) : uint4(1 << gl_SubgroupInvocationID, uint3(0)); + uint4 gl_SubgroupGeMask = uint4(extract_bits(0xFFFFFFFF, min(gl_SubgroupInvocationID, 32u), (uint)max(min((int)gl_SubgroupSize, 32) - (int)gl_SubgroupInvocationID, 0)), extract_bits(0xFFFFFFFF, (uint)max((int)gl_SubgroupInvocationID - 32, 0), (uint)max((int)gl_SubgroupSize - (int)max(gl_SubgroupInvocationID, 32u), 0)), uint2(0)); + uint4 gl_SubgroupGtMask = uint4(extract_bits(0xFFFFFFFF, min(gl_SubgroupInvocationID + 1, 32u), (uint)max(min((int)gl_SubgroupSize, 32) - (int)gl_SubgroupInvocationID - 1, 0)), extract_bits(0xFFFFFFFF, (uint)max((int)gl_SubgroupInvocationID + 1 - 32, 0), (uint)max((int)gl_SubgroupSize - (int)max(gl_SubgroupInvocationID + 1, 32u), 0)), uint2(0)); + uint4 gl_SubgroupLeMask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID + 1, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID + 1 - 32, 0)), uint2(0)); + uint4 gl_SubgroupLtMask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID - 32, 0)), uint2(0)); + _9.FragColor = float(gl_NumSubgroups); + _9.FragColor = float(gl_SubgroupID); + _9.FragColor = float(gl_SubgroupSize); + _9.FragColor = float(gl_SubgroupInvocationID); + simdgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture); + simdgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture); + simdgroup_barrier(mem_flags::mem_device); + simdgroup_barrier(mem_flags::mem_threadgroup); + simdgroup_barrier(mem_flags::mem_texture); + _9.FragColor = float4(gl_SubgroupEqMask).x; + _9.FragColor = float4(gl_SubgroupGeMask).x; + _9.FragColor = float4(gl_SubgroupGtMask).x; + _9.FragColor = float4(gl_SubgroupLeMask).x; + _9.FragColor = float4(gl_SubgroupLtMask).x; + uint4 _83 = spvSubgroupBallot(true); + float4 _165 = simd_prefix_inclusive_product(simd_product(float4(20.0))); + int4 _167 = simd_prefix_inclusive_product(simd_product(int4(20))); +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.ios.comp b/3rdparty/spirv-cross/reference/opt/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.ios.comp new file mode 100644 index 000000000..6d32de695 --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.ios.comp @@ -0,0 +1,23 @@ +#include +#include + +using namespace metal; + +struct SSBO +{ + float FragColor; +}; + +kernel void main0(device SSBO& _9 [[buffer(0)]], uint gl_NumSubgroups [[quadgroups_per_threadgroup]], uint gl_SubgroupID [[quadgroup_index_in_threadgroup]], uint gl_SubgroupSize [[thread_execution_width]], uint gl_SubgroupInvocationID [[thread_index_in_quadgroup]]) +{ + _9.FragColor = float(gl_NumSubgroups); + _9.FragColor = float(gl_SubgroupID); + _9.FragColor = float(gl_SubgroupSize); + _9.FragColor = float(gl_SubgroupInvocationID); + simdgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture); + simdgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture); + simdgroup_barrier(mem_flags::mem_device); + simdgroup_barrier(mem_flags::mem_threadgroup); + simdgroup_barrier(mem_flags::mem_texture); +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders/comp/ssbo-array-length.comp b/3rdparty/spirv-cross/reference/opt/shaders/comp/ssbo-array-length.comp new file mode 100644 index 000000000..ddc666e9b --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders/comp/ssbo-array-length.comp @@ -0,0 +1,14 @@ +#version 450 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(binding = 1, std140) buffer SSBO +{ + uint size; + float v[]; +} _11; + +void main() +{ + _11.size = uint(int(uint(_11.v.length()))); +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders/vulkan/comp/buffer-reference-bitcast.nocompat.vk.comp.vk b/3rdparty/spirv-cross/reference/opt/shaders/vulkan/comp/buffer-reference-bitcast.nocompat.vk.comp.vk new file mode 100644 index 000000000..5752f81b2 --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders/vulkan/comp/buffer-reference-bitcast.nocompat.vk.comp.vk @@ -0,0 +1,26 @@ +#version 450 +#extension GL_EXT_buffer_reference : require +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(buffer_reference) buffer PtrUint; +layout(buffer_reference) buffer PtrInt; +layout(buffer_reference, std430) buffer PtrUint +{ + uint value; +}; + +layout(buffer_reference, std430) buffer PtrInt +{ + int value; +}; + +layout(set = 0, binding = 0, std430) buffer Buf +{ + PtrUint ptr; +} _11; + +void main() +{ + PtrInt(_11.ptr).value = 10; +} + diff --git a/3rdparty/spirv-cross/reference/shaders-hlsl/comp/ssbo-array-length.comp b/3rdparty/spirv-cross/reference/shaders-hlsl/comp/ssbo-array-length.comp new file mode 100644 index 000000000..2e3df626a --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-hlsl/comp/ssbo-array-length.comp @@ -0,0 +1,15 @@ +RWByteAddressBuffer _11 : register(u1); + +void comp_main() +{ + uint _14; + _11.GetDimensions(_14); + _14 = (_14 - 16) / 16; + _11.Store(0, uint(int(_14))); +} + +[numthreads(1, 1, 1)] +void main() +{ + comp_main(); +} diff --git a/3rdparty/spirv-cross/reference/shaders-hlsl/frag/nonuniform-qualifier.nonuniformresource.sm51.frag b/3rdparty/spirv-cross/reference/shaders-hlsl/frag/nonuniform-qualifier.nonuniformresource.sm51.frag new file mode 100644 index 000000000..8f5e022eb --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-hlsl/frag/nonuniform-qualifier.nonuniformresource.sm51.frag @@ -0,0 +1,46 @@ +struct UBO_1_1 +{ + float4 v[64]; +}; + +ConstantBuffer ubos[] : register(b0, space3); +ByteAddressBuffer ssbos[] : register(t0, space4); +Texture2D uSamplers[] : register(t0, space0); +SamplerState uSamps[] : register(s0, space2); +Texture2D uCombinedSamplers[] : register(t0, space1); +SamplerState _uCombinedSamplers_sampler[] : register(s0, space1); + +static int vIndex; +static float4 FragColor; +static float2 vUV; + +struct SPIRV_Cross_Input +{ + nointerpolation int vIndex : TEXCOORD0; + float2 vUV : TEXCOORD1; +}; + +struct SPIRV_Cross_Output +{ + float4 FragColor : SV_Target0; +}; + +void frag_main() +{ + int i = vIndex; + FragColor = uSamplers[NonUniformResourceIndex(i + 10)].Sample(uSamps[NonUniformResourceIndex(i + 40)], vUV); + int _47 = i + 10; + FragColor = uCombinedSamplers[NonUniformResourceIndex(_47)].Sample(_uCombinedSamplers_sampler[NonUniformResourceIndex(_47)], vUV); + FragColor += ubos[NonUniformResourceIndex(i + 20)].v[i + 40]; + FragColor += asfloat(ssbos[NonUniformResourceIndex(i + 50)].Load4((i + 60) * 16 + 0)); +} + +SPIRV_Cross_Output main(SPIRV_Cross_Input stage_input) +{ + vIndex = stage_input.vIndex; + vUV = stage_input.vUV; + frag_main(); + SPIRV_Cross_Output stage_output; + stage_output.FragColor = FragColor; + return stage_output; +} diff --git a/3rdparty/spirv-cross/reference/shaders-msl-no-opt/asm/frag/texture-access.swizzle.asm.frag b/3rdparty/spirv-cross/reference/shaders-msl-no-opt/asm/frag/texture-access.swizzle.asm.frag index ea433e166..68445b792 100644 --- a/3rdparty/spirv-cross/reference/shaders-msl-no-opt/asm/frag/texture-access.swizzle.asm.frag +++ b/3rdparty/spirv-cross/reference/shaders-msl-no-opt/asm/frag/texture-access.swizzle.asm.frag @@ -5,11 +5,6 @@ using namespace metal; -struct spvAux -{ - uint swizzleConst[1]; -}; - // Returns 2D texture coords corresponding to 1D texel buffer coords uint2 spvTexelBufferCoord(uint tc) { @@ -136,18 +131,18 @@ inline vec spvGatherCompareSwizzle(sampler s, const thread Tex& t, Ts... p return t.gather_compare(s, spvForward(params)...); } -fragment void main0(constant spvAux& spvAuxBuffer [[buffer(30)]], texture1d tex1d [[texture(0)]], texture2d tex2d [[texture(1)]], texture3d tex3d [[texture(2)]], texturecube texCube [[texture(3)]], texture2d_array tex2dArray [[texture(4)]], texturecube_array texCubeArray [[texture(5)]], texture2d texBuffer [[texture(6)]], depth2d depth2d [[texture(7)]], depthcube depthCube [[texture(8)]], depth2d_array depth2dArray [[texture(9)]], depthcube_array depthCubeArray [[texture(10)]], sampler tex1dSamp [[sampler(0)]], sampler tex2dSamp [[sampler(1)]], sampler tex3dSamp [[sampler(2)]], sampler texCubeSamp [[sampler(3)]], sampler tex2dArraySamp [[sampler(4)]], sampler texCubeArraySamp [[sampler(5)]], sampler depth2dSamp [[sampler(7)]], sampler depthCubeSamp [[sampler(8)]], sampler depth2dArraySamp [[sampler(9)]], sampler depthCubeArraySamp [[sampler(10)]]) +fragment void main0(constant uint* spvSwizzleConstants [[buffer(30)]], texture1d tex1d [[texture(0)]], texture2d tex2d [[texture(1)]], texture3d tex3d [[texture(2)]], texturecube texCube [[texture(3)]], texture2d_array tex2dArray [[texture(4)]], texturecube_array texCubeArray [[texture(5)]], texture2d texBuffer [[texture(6)]], depth2d depth2d [[texture(7)]], depthcube depthCube [[texture(8)]], depth2d_array depth2dArray [[texture(9)]], depthcube_array depthCubeArray [[texture(10)]], sampler tex1dSamp [[sampler(0)]], sampler tex2dSamp [[sampler(1)]], sampler tex3dSamp [[sampler(2)]], sampler texCubeSamp [[sampler(3)]], sampler tex2dArraySamp [[sampler(4)]], sampler texCubeArraySamp [[sampler(5)]], sampler depth2dSamp [[sampler(7)]], sampler depthCubeSamp [[sampler(8)]], sampler depth2dArraySamp [[sampler(9)]], sampler depthCubeArraySamp [[sampler(10)]]) { - constant uint32_t& tex1dSwzl = spvAuxBuffer.swizzleConst[0]; - constant uint32_t& tex2dSwzl = spvAuxBuffer.swizzleConst[1]; - constant uint32_t& tex3dSwzl = spvAuxBuffer.swizzleConst[2]; - constant uint32_t& texCubeSwzl = spvAuxBuffer.swizzleConst[3]; - constant uint32_t& tex2dArraySwzl = spvAuxBuffer.swizzleConst[4]; - constant uint32_t& texCubeArraySwzl = spvAuxBuffer.swizzleConst[5]; - constant uint32_t& depth2dSwzl = spvAuxBuffer.swizzleConst[7]; - constant uint32_t& depthCubeSwzl = spvAuxBuffer.swizzleConst[8]; - constant uint32_t& depth2dArraySwzl = spvAuxBuffer.swizzleConst[9]; - constant uint32_t& depthCubeArraySwzl = spvAuxBuffer.swizzleConst[10]; + constant uint32_t& tex1dSwzl = spvSwizzleConstants[0]; + constant uint32_t& tex2dSwzl = spvSwizzleConstants[1]; + constant uint32_t& tex3dSwzl = spvSwizzleConstants[2]; + constant uint32_t& texCubeSwzl = spvSwizzleConstants[3]; + constant uint32_t& tex2dArraySwzl = spvSwizzleConstants[4]; + constant uint32_t& texCubeArraySwzl = spvSwizzleConstants[5]; + constant uint32_t& depth2dSwzl = spvSwizzleConstants[7]; + constant uint32_t& depthCubeSwzl = spvSwizzleConstants[8]; + constant uint32_t& depth2dArraySwzl = spvSwizzleConstants[9]; + constant uint32_t& depthCubeArraySwzl = spvSwizzleConstants[10]; float4 c = spvTextureSwizzle(tex1d.sample(tex1dSamp, 0.0), tex1dSwzl); c = spvTextureSwizzle(tex2d.sample(tex2dSamp, float2(0.0)), tex2dSwzl); c = spvTextureSwizzle(tex3d.sample(tex3dSamp, float3(0.0)), tex3dSwzl); diff --git a/3rdparty/spirv-cross/reference/shaders-msl-no-opt/frag/texture-access-int.swizzle.frag b/3rdparty/spirv-cross/reference/shaders-msl-no-opt/frag/texture-access-int.swizzle.frag index 997d5b55b..3f854e701 100644 --- a/3rdparty/spirv-cross/reference/shaders-msl-no-opt/frag/texture-access-int.swizzle.frag +++ b/3rdparty/spirv-cross/reference/shaders-msl-no-opt/frag/texture-access-int.swizzle.frag @@ -5,11 +5,6 @@ using namespace metal; -struct spvAux -{ - uint swizzleConst[1]; -}; - // Returns 2D texture coords corresponding to 1D texel buffer coords uint2 spvTexelBufferCoord(uint tc) { @@ -136,14 +131,14 @@ inline vec spvGatherCompareSwizzle(sampler s, const thread Tex& t, Ts... p return t.gather_compare(s, spvForward(params)...); } -fragment void main0(constant spvAux& spvAuxBuffer [[buffer(30)]], texture1d tex1d [[texture(0)]], texture2d tex2d [[texture(1)]], texture3d tex3d [[texture(2)]], texturecube texCube [[texture(3)]], texture2d_array tex2dArray [[texture(4)]], texturecube_array texCubeArray [[texture(5)]], texture2d texBuffer [[texture(6)]], sampler tex1dSmplr [[sampler(0)]], sampler tex2dSmplr [[sampler(1)]], sampler tex3dSmplr [[sampler(2)]], sampler texCubeSmplr [[sampler(3)]], sampler tex2dArraySmplr [[sampler(4)]], sampler texCubeArraySmplr [[sampler(5)]]) +fragment void main0(constant uint* spvSwizzleConstants [[buffer(30)]], texture1d tex1d [[texture(0)]], texture2d tex2d [[texture(1)]], texture3d tex3d [[texture(2)]], texturecube texCube [[texture(3)]], texture2d_array tex2dArray [[texture(4)]], texturecube_array texCubeArray [[texture(5)]], texture2d texBuffer [[texture(6)]], sampler tex1dSmplr [[sampler(0)]], sampler tex2dSmplr [[sampler(1)]], sampler tex3dSmplr [[sampler(2)]], sampler texCubeSmplr [[sampler(3)]], sampler tex2dArraySmplr [[sampler(4)]], sampler texCubeArraySmplr [[sampler(5)]]) { - constant uint32_t& tex1dSwzl = spvAuxBuffer.swizzleConst[0]; - constant uint32_t& tex2dSwzl = spvAuxBuffer.swizzleConst[1]; - constant uint32_t& tex3dSwzl = spvAuxBuffer.swizzleConst[2]; - constant uint32_t& texCubeSwzl = spvAuxBuffer.swizzleConst[3]; - constant uint32_t& tex2dArraySwzl = spvAuxBuffer.swizzleConst[4]; - constant uint32_t& texCubeArraySwzl = spvAuxBuffer.swizzleConst[5]; + constant uint32_t& tex1dSwzl = spvSwizzleConstants[0]; + constant uint32_t& tex2dSwzl = spvSwizzleConstants[1]; + constant uint32_t& tex3dSwzl = spvSwizzleConstants[2]; + constant uint32_t& texCubeSwzl = spvSwizzleConstants[3]; + constant uint32_t& tex2dArraySwzl = spvSwizzleConstants[4]; + constant uint32_t& texCubeArraySwzl = spvSwizzleConstants[5]; float4 c = float4(spvTextureSwizzle(tex1d.sample(tex1dSmplr, 0.0), tex1dSwzl)); c = float4(spvTextureSwizzle(tex2d.sample(tex2dSmplr, float2(0.0)), tex2dSwzl)); c = float4(spvTextureSwizzle(tex3d.sample(tex3dSmplr, float3(0.0)), tex3dSwzl)); diff --git a/3rdparty/spirv-cross/reference/shaders-msl-no-opt/frag/texture-access-leaf.swizzle.frag b/3rdparty/spirv-cross/reference/shaders-msl-no-opt/frag/texture-access-leaf.swizzle.frag index 076dcc177..01a5ae64d 100644 --- a/3rdparty/spirv-cross/reference/shaders-msl-no-opt/frag/texture-access-leaf.swizzle.frag +++ b/3rdparty/spirv-cross/reference/shaders-msl-no-opt/frag/texture-access-leaf.swizzle.frag @@ -5,11 +5,6 @@ using namespace metal; -struct spvAux -{ - uint swizzleConst[1]; -}; - // Returns 2D texture coords corresponding to 1D texel buffer coords uint2 spvTexelBufferCoord(uint tc) { @@ -183,18 +178,18 @@ float4 doSwizzle(thread texture1d tex1d, thread const sampler tex1dSmplr, return c; } -fragment void main0(constant spvAux& spvAuxBuffer [[buffer(30)]], texture1d tex1d [[texture(0)]], texture2d tex2d [[texture(1)]], texture3d tex3d [[texture(2)]], texturecube texCube [[texture(3)]], texture2d_array tex2dArray [[texture(4)]], texturecube_array texCubeArray [[texture(5)]], texture2d texBuffer [[texture(6)]], depth2d depth2d [[texture(7)]], depthcube depthCube [[texture(8)]], depth2d_array depth2dArray [[texture(9)]], depthcube_array depthCubeArray [[texture(10)]], sampler tex1dSmplr [[sampler(0)]], sampler tex2dSmplr [[sampler(1)]], sampler tex3dSmplr [[sampler(2)]], sampler texCubeSmplr [[sampler(3)]], sampler tex2dArraySmplr [[sampler(4)]], sampler texCubeArraySmplr [[sampler(5)]], sampler depth2dSmplr [[sampler(7)]], sampler depthCubeSmplr [[sampler(8)]], sampler depth2dArraySmplr [[sampler(9)]], sampler depthCubeArraySmplr [[sampler(10)]]) +fragment void main0(constant uint* spvSwizzleConstants [[buffer(30)]], texture1d tex1d [[texture(0)]], texture2d tex2d [[texture(1)]], texture3d tex3d [[texture(2)]], texturecube texCube [[texture(3)]], texture2d_array tex2dArray [[texture(4)]], texturecube_array texCubeArray [[texture(5)]], texture2d texBuffer [[texture(6)]], depth2d depth2d [[texture(7)]], depthcube depthCube [[texture(8)]], depth2d_array depth2dArray [[texture(9)]], depthcube_array depthCubeArray [[texture(10)]], sampler tex1dSmplr [[sampler(0)]], sampler tex2dSmplr [[sampler(1)]], sampler tex3dSmplr [[sampler(2)]], sampler texCubeSmplr [[sampler(3)]], sampler tex2dArraySmplr [[sampler(4)]], sampler texCubeArraySmplr [[sampler(5)]], sampler depth2dSmplr [[sampler(7)]], sampler depthCubeSmplr [[sampler(8)]], sampler depth2dArraySmplr [[sampler(9)]], sampler depthCubeArraySmplr [[sampler(10)]]) { - constant uint32_t& tex1dSwzl = spvAuxBuffer.swizzleConst[0]; - constant uint32_t& tex2dSwzl = spvAuxBuffer.swizzleConst[1]; - constant uint32_t& tex3dSwzl = spvAuxBuffer.swizzleConst[2]; - constant uint32_t& texCubeSwzl = spvAuxBuffer.swizzleConst[3]; - constant uint32_t& tex2dArraySwzl = spvAuxBuffer.swizzleConst[4]; - constant uint32_t& texCubeArraySwzl = spvAuxBuffer.swizzleConst[5]; - constant uint32_t& depth2dSwzl = spvAuxBuffer.swizzleConst[7]; - constant uint32_t& depthCubeSwzl = spvAuxBuffer.swizzleConst[8]; - constant uint32_t& depth2dArraySwzl = spvAuxBuffer.swizzleConst[9]; - constant uint32_t& depthCubeArraySwzl = spvAuxBuffer.swizzleConst[10]; + constant uint32_t& tex1dSwzl = spvSwizzleConstants[0]; + constant uint32_t& tex2dSwzl = spvSwizzleConstants[1]; + constant uint32_t& tex3dSwzl = spvSwizzleConstants[2]; + constant uint32_t& texCubeSwzl = spvSwizzleConstants[3]; + constant uint32_t& tex2dArraySwzl = spvSwizzleConstants[4]; + constant uint32_t& texCubeArraySwzl = spvSwizzleConstants[5]; + constant uint32_t& depth2dSwzl = spvSwizzleConstants[7]; + constant uint32_t& depthCubeSwzl = spvSwizzleConstants[8]; + constant uint32_t& depth2dArraySwzl = spvSwizzleConstants[9]; + constant uint32_t& depthCubeArraySwzl = spvSwizzleConstants[10]; float4 c = doSwizzle(tex1d, tex1dSmplr, tex1dSwzl, tex2d, tex2dSmplr, tex2dSwzl, tex3d, tex3dSmplr, tex3dSwzl, texCube, texCubeSmplr, texCubeSwzl, tex2dArray, tex2dArraySmplr, tex2dArraySwzl, texCubeArray, texCubeArraySmplr, texCubeArraySwzl, depth2d, depth2dSmplr, depth2dSwzl, depthCube, depthCubeSmplr, depthCubeSwzl, depth2dArray, depth2dArraySmplr, depth2dArraySwzl, depthCubeArray, depthCubeArraySmplr, depthCubeArraySwzl, texBuffer); } diff --git a/3rdparty/spirv-cross/reference/shaders-msl-no-opt/frag/texture-access-uint.swizzle.frag b/3rdparty/spirv-cross/reference/shaders-msl-no-opt/frag/texture-access-uint.swizzle.frag index 81ffec75b..5a3013ba8 100644 --- a/3rdparty/spirv-cross/reference/shaders-msl-no-opt/frag/texture-access-uint.swizzle.frag +++ b/3rdparty/spirv-cross/reference/shaders-msl-no-opt/frag/texture-access-uint.swizzle.frag @@ -5,11 +5,6 @@ using namespace metal; -struct spvAux -{ - uint swizzleConst[1]; -}; - // Returns 2D texture coords corresponding to 1D texel buffer coords uint2 spvTexelBufferCoord(uint tc) { @@ -136,14 +131,14 @@ inline vec spvGatherCompareSwizzle(sampler s, const thread Tex& t, Ts... p return t.gather_compare(s, spvForward(params)...); } -fragment void main0(constant spvAux& spvAuxBuffer [[buffer(30)]], texture1d tex1d [[texture(0)]], texture2d tex2d [[texture(1)]], texture3d tex3d [[texture(2)]], texturecube texCube [[texture(3)]], texture2d_array tex2dArray [[texture(4)]], texturecube_array texCubeArray [[texture(5)]], texture2d texBuffer [[texture(6)]], sampler tex1dSmplr [[sampler(0)]], sampler tex2dSmplr [[sampler(1)]], sampler tex3dSmplr [[sampler(2)]], sampler texCubeSmplr [[sampler(3)]], sampler tex2dArraySmplr [[sampler(4)]], sampler texCubeArraySmplr [[sampler(5)]]) +fragment void main0(constant uint* spvSwizzleConstants [[buffer(30)]], texture1d tex1d [[texture(0)]], texture2d tex2d [[texture(1)]], texture3d tex3d [[texture(2)]], texturecube texCube [[texture(3)]], texture2d_array tex2dArray [[texture(4)]], texturecube_array texCubeArray [[texture(5)]], texture2d texBuffer [[texture(6)]], sampler tex1dSmplr [[sampler(0)]], sampler tex2dSmplr [[sampler(1)]], sampler tex3dSmplr [[sampler(2)]], sampler texCubeSmplr [[sampler(3)]], sampler tex2dArraySmplr [[sampler(4)]], sampler texCubeArraySmplr [[sampler(5)]]) { - constant uint32_t& tex1dSwzl = spvAuxBuffer.swizzleConst[0]; - constant uint32_t& tex2dSwzl = spvAuxBuffer.swizzleConst[1]; - constant uint32_t& tex3dSwzl = spvAuxBuffer.swizzleConst[2]; - constant uint32_t& texCubeSwzl = spvAuxBuffer.swizzleConst[3]; - constant uint32_t& tex2dArraySwzl = spvAuxBuffer.swizzleConst[4]; - constant uint32_t& texCubeArraySwzl = spvAuxBuffer.swizzleConst[5]; + constant uint32_t& tex1dSwzl = spvSwizzleConstants[0]; + constant uint32_t& tex2dSwzl = spvSwizzleConstants[1]; + constant uint32_t& tex3dSwzl = spvSwizzleConstants[2]; + constant uint32_t& texCubeSwzl = spvSwizzleConstants[3]; + constant uint32_t& tex2dArraySwzl = spvSwizzleConstants[4]; + constant uint32_t& texCubeArraySwzl = spvSwizzleConstants[5]; float4 c = float4(spvTextureSwizzle(tex1d.sample(tex1dSmplr, 0.0), tex1dSwzl)); c = float4(spvTextureSwizzle(tex2d.sample(tex2dSmplr, float2(0.0)), tex2dSwzl)); c = float4(spvTextureSwizzle(tex3d.sample(tex3dSmplr, float3(0.0)), tex3dSwzl)); diff --git a/3rdparty/spirv-cross/reference/shaders-msl-no-opt/frag/texture-access.swizzle.frag b/3rdparty/spirv-cross/reference/shaders-msl-no-opt/frag/texture-access.swizzle.frag index 98ab50eab..befee5bcc 100644 --- a/3rdparty/spirv-cross/reference/shaders-msl-no-opt/frag/texture-access.swizzle.frag +++ b/3rdparty/spirv-cross/reference/shaders-msl-no-opt/frag/texture-access.swizzle.frag @@ -5,11 +5,6 @@ using namespace metal; -struct spvAux -{ - uint swizzleConst[1]; -}; - // Returns 2D texture coords corresponding to 1D texel buffer coords uint2 spvTexelBufferCoord(uint tc) { @@ -136,18 +131,18 @@ inline vec spvGatherCompareSwizzle(sampler s, const thread Tex& t, Ts... p return t.gather_compare(s, spvForward(params)...); } -fragment void main0(constant spvAux& spvAuxBuffer [[buffer(30)]], texture1d tex1d [[texture(0)]], texture2d tex2d [[texture(1)]], texture3d tex3d [[texture(2)]], texturecube texCube [[texture(3)]], texture2d_array tex2dArray [[texture(4)]], texturecube_array texCubeArray [[texture(5)]], texture2d texBuffer [[texture(6)]], depth2d depth2d [[texture(7)]], depthcube depthCube [[texture(8)]], depth2d_array depth2dArray [[texture(9)]], depthcube_array depthCubeArray [[texture(10)]], sampler tex1dSmplr [[sampler(0)]], sampler tex2dSmplr [[sampler(1)]], sampler tex3dSmplr [[sampler(2)]], sampler texCubeSmplr [[sampler(3)]], sampler tex2dArraySmplr [[sampler(4)]], sampler texCubeArraySmplr [[sampler(5)]], sampler depth2dSmplr [[sampler(7)]], sampler depthCubeSmplr [[sampler(8)]], sampler depth2dArraySmplr [[sampler(9)]], sampler depthCubeArraySmplr [[sampler(10)]]) +fragment void main0(constant uint* spvSwizzleConstants [[buffer(30)]], texture1d tex1d [[texture(0)]], texture2d tex2d [[texture(1)]], texture3d tex3d [[texture(2)]], texturecube texCube [[texture(3)]], texture2d_array tex2dArray [[texture(4)]], texturecube_array texCubeArray [[texture(5)]], texture2d texBuffer [[texture(6)]], depth2d depth2d [[texture(7)]], depthcube depthCube [[texture(8)]], depth2d_array depth2dArray [[texture(9)]], depthcube_array depthCubeArray [[texture(10)]], sampler tex1dSmplr [[sampler(0)]], sampler tex2dSmplr [[sampler(1)]], sampler tex3dSmplr [[sampler(2)]], sampler texCubeSmplr [[sampler(3)]], sampler tex2dArraySmplr [[sampler(4)]], sampler texCubeArraySmplr [[sampler(5)]], sampler depth2dSmplr [[sampler(7)]], sampler depthCubeSmplr [[sampler(8)]], sampler depth2dArraySmplr [[sampler(9)]], sampler depthCubeArraySmplr [[sampler(10)]]) { - constant uint32_t& tex1dSwzl = spvAuxBuffer.swizzleConst[0]; - constant uint32_t& tex2dSwzl = spvAuxBuffer.swizzleConst[1]; - constant uint32_t& tex3dSwzl = spvAuxBuffer.swizzleConst[2]; - constant uint32_t& texCubeSwzl = spvAuxBuffer.swizzleConst[3]; - constant uint32_t& tex2dArraySwzl = spvAuxBuffer.swizzleConst[4]; - constant uint32_t& texCubeArraySwzl = spvAuxBuffer.swizzleConst[5]; - constant uint32_t& depth2dSwzl = spvAuxBuffer.swizzleConst[7]; - constant uint32_t& depthCubeSwzl = spvAuxBuffer.swizzleConst[8]; - constant uint32_t& depth2dArraySwzl = spvAuxBuffer.swizzleConst[9]; - constant uint32_t& depthCubeArraySwzl = spvAuxBuffer.swizzleConst[10]; + constant uint32_t& tex1dSwzl = spvSwizzleConstants[0]; + constant uint32_t& tex2dSwzl = spvSwizzleConstants[1]; + constant uint32_t& tex3dSwzl = spvSwizzleConstants[2]; + constant uint32_t& texCubeSwzl = spvSwizzleConstants[3]; + constant uint32_t& tex2dArraySwzl = spvSwizzleConstants[4]; + constant uint32_t& texCubeArraySwzl = spvSwizzleConstants[5]; + constant uint32_t& depth2dSwzl = spvSwizzleConstants[7]; + constant uint32_t& depthCubeSwzl = spvSwizzleConstants[8]; + constant uint32_t& depth2dArraySwzl = spvSwizzleConstants[9]; + constant uint32_t& depthCubeArraySwzl = spvSwizzleConstants[10]; float4 c = spvTextureSwizzle(tex1d.sample(tex1dSmplr, 0.0), tex1dSwzl); c = spvTextureSwizzle(tex2d.sample(tex2dSmplr, float2(0.0)), tex2dSwzl); c = spvTextureSwizzle(tex3d.sample(tex3dSmplr, float3(0.0)), tex3dSwzl); diff --git a/3rdparty/spirv-cross/reference/shaders-msl-no-opt/vulkan/frag/texture-access-function.swizzle.vk.frag b/3rdparty/spirv-cross/reference/shaders-msl-no-opt/vulkan/frag/texture-access-function.swizzle.vk.frag index e99cd69a2..aa3aae238 100644 --- a/3rdparty/spirv-cross/reference/shaders-msl-no-opt/vulkan/frag/texture-access-function.swizzle.vk.frag +++ b/3rdparty/spirv-cross/reference/shaders-msl-no-opt/vulkan/frag/texture-access-function.swizzle.vk.frag @@ -5,11 +5,6 @@ using namespace metal; -struct spvAux -{ - uint swizzleConst[1]; -}; - struct main0_out { float4 fragColor [[color(0)]]; @@ -188,19 +183,19 @@ float4 do_samples(thread const texture1d t1, thread const sampler t1Smplr return c; } -fragment main0_out main0(constant spvAux& spvAuxBuffer [[buffer(30)]], texture1d tex1d [[texture(0)]], texture2d tex2d [[texture(1)]], texture3d tex3d [[texture(2)]], texturecube texCube [[texture(3)]], texture2d_array tex2dArray [[texture(4)]], texturecube_array texCubeArray [[texture(5)]], texture2d texBuffer [[texture(6)]], depth2d depth2d [[texture(7)]], depthcube depthCube [[texture(8)]], depth2d_array depth2dArray [[texture(9)]], depthcube_array depthCubeArray [[texture(10)]], sampler tex1dSmplr [[sampler(0)]], sampler tex3dSmplr [[sampler(2)]], sampler tex2dArraySmplr [[sampler(4)]], sampler texCubeArraySmplr [[sampler(5)]], sampler depth2dSmplr [[sampler(7)]], sampler depthCubeSmplr [[sampler(8)]], sampler depthCubeArraySmplr [[sampler(10)]], sampler defaultSampler [[sampler(11)]], sampler shadowSampler [[sampler(12)]]) +fragment main0_out main0(constant uint* spvSwizzleConstants [[buffer(30)]], texture1d tex1d [[texture(0)]], texture2d tex2d [[texture(1)]], texture3d tex3d [[texture(2)]], texturecube texCube [[texture(3)]], texture2d_array tex2dArray [[texture(4)]], texturecube_array texCubeArray [[texture(5)]], texture2d texBuffer [[texture(6)]], depth2d depth2d [[texture(7)]], depthcube depthCube [[texture(8)]], depth2d_array depth2dArray [[texture(9)]], depthcube_array depthCubeArray [[texture(10)]], sampler tex1dSmplr [[sampler(0)]], sampler tex3dSmplr [[sampler(2)]], sampler tex2dArraySmplr [[sampler(4)]], sampler texCubeArraySmplr [[sampler(5)]], sampler depth2dSmplr [[sampler(7)]], sampler depthCubeSmplr [[sampler(8)]], sampler depthCubeArraySmplr [[sampler(10)]], sampler defaultSampler [[sampler(11)]], sampler shadowSampler [[sampler(12)]]) { main0_out out = {}; - constant uint32_t& tex1dSwzl = spvAuxBuffer.swizzleConst[0]; - constant uint32_t& tex2dSwzl = spvAuxBuffer.swizzleConst[1]; - constant uint32_t& tex3dSwzl = spvAuxBuffer.swizzleConst[2]; - constant uint32_t& texCubeSwzl = spvAuxBuffer.swizzleConst[3]; - constant uint32_t& tex2dArraySwzl = spvAuxBuffer.swizzleConst[4]; - constant uint32_t& texCubeArraySwzl = spvAuxBuffer.swizzleConst[5]; - constant uint32_t& depth2dSwzl = spvAuxBuffer.swizzleConst[7]; - constant uint32_t& depthCubeSwzl = spvAuxBuffer.swizzleConst[8]; - constant uint32_t& depth2dArraySwzl = spvAuxBuffer.swizzleConst[9]; - constant uint32_t& depthCubeArraySwzl = spvAuxBuffer.swizzleConst[10]; + constant uint32_t& tex1dSwzl = spvSwizzleConstants[0]; + constant uint32_t& tex2dSwzl = spvSwizzleConstants[1]; + constant uint32_t& tex3dSwzl = spvSwizzleConstants[2]; + constant uint32_t& texCubeSwzl = spvSwizzleConstants[3]; + constant uint32_t& tex2dArraySwzl = spvSwizzleConstants[4]; + constant uint32_t& texCubeArraySwzl = spvSwizzleConstants[5]; + constant uint32_t& depth2dSwzl = spvSwizzleConstants[7]; + constant uint32_t& depthCubeSwzl = spvSwizzleConstants[8]; + constant uint32_t& depth2dArraySwzl = spvSwizzleConstants[9]; + constant uint32_t& depthCubeArraySwzl = spvSwizzleConstants[10]; out.fragColor = do_samples(tex1d, tex1dSmplr, tex1dSwzl, tex2d, tex2dSwzl, tex3d, tex3dSmplr, tex3dSwzl, texCube, texCubeSwzl, tex2dArray, tex2dArraySmplr, tex2dArraySwzl, texCubeArray, texCubeArraySmplr, texCubeArraySwzl, texBuffer, depth2d, depth2dSmplr, depth2dSwzl, depthCube, depthCubeSmplr, depthCubeSwzl, depth2dArray, depth2dArraySwzl, depthCubeArray, depthCubeArraySmplr, depthCubeArraySwzl, defaultSampler, shadowSampler); return out; } diff --git a/3rdparty/spirv-cross/reference/shaders-msl/comp/barriers.comp b/3rdparty/spirv-cross/reference/shaders-msl/comp/barriers.comp index fdaf87477..560fd8b53 100644 --- a/3rdparty/spirv-cross/reference/shaders-msl/comp/barriers.comp +++ b/3rdparty/spirv-cross/reference/shaders-msl/comp/barriers.comp @@ -14,17 +14,22 @@ void barrier_shared() void full_barrier() { - threadgroup_barrier(mem_flags::mem_threadgroup); + threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture); +} + +void image_barrier() +{ + threadgroup_barrier(mem_flags::mem_texture); } void buffer_barrier() { - threadgroup_barrier(mem_flags::mem_none); + threadgroup_barrier(mem_flags::mem_device); } void group_barrier() { - threadgroup_barrier(mem_flags::mem_threadgroup); + threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture); } void barrier_shared_exec() @@ -34,17 +39,22 @@ void barrier_shared_exec() void full_barrier_exec() { - threadgroup_barrier(mem_flags::mem_threadgroup); + threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture); +} + +void image_barrier_exec() +{ + threadgroup_barrier(mem_flags::mem_texture); } void buffer_barrier_exec() { - threadgroup_barrier(mem_flags::mem_none); + threadgroup_barrier(mem_flags::mem_device); } void group_barrier_exec() { - threadgroup_barrier(mem_flags::mem_threadgroup); + threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture); } void exec_barrier() @@ -56,10 +66,12 @@ kernel void main0() { barrier_shared(); full_barrier(); + image_barrier(); buffer_barrier(); group_barrier(); barrier_shared_exec(); full_barrier_exec(); + image_barrier_exec(); buffer_barrier_exec(); group_barrier_exec(); exec_barrier(); diff --git a/3rdparty/spirv-cross/reference/shaders-msl/comp/force-recompile-hooks.swizzle.comp b/3rdparty/spirv-cross/reference/shaders-msl/comp/force-recompile-hooks.swizzle.comp index d8b4af791..b784940e8 100644 --- a/3rdparty/spirv-cross/reference/shaders-msl/comp/force-recompile-hooks.swizzle.comp +++ b/3rdparty/spirv-cross/reference/shaders-msl/comp/force-recompile-hooks.swizzle.comp @@ -5,11 +5,6 @@ using namespace metal; -struct spvAux -{ - uint swizzleConst[1]; -}; - enum class spvSwizzle : uint { none = 0, @@ -130,9 +125,9 @@ inline vec spvGatherCompareSwizzle(sampler s, const thread Tex& t, Ts... p return t.gather_compare(s, spvForward(params)...); } -kernel void main0(constant spvAux& spvAuxBuffer [[buffer(30)]], texture2d foo [[texture(0)]], texture2d bar [[texture(1)]], sampler fooSmplr [[sampler(0)]]) +kernel void main0(constant uint* spvSwizzleConstants [[buffer(30)]], texture2d foo [[texture(0)]], texture2d bar [[texture(1)]], sampler fooSmplr [[sampler(0)]]) { - constant uint32_t& fooSwzl = spvAuxBuffer.swizzleConst[0]; + constant uint32_t& fooSwzl = spvSwizzleConstants[0]; float4 a = spvTextureSwizzle(foo.sample(fooSmplr, float2(1.0), level(0.0)), fooSwzl); bar.write(a, uint2(int2(0))); } diff --git a/3rdparty/spirv-cross/reference/shaders-msl/frag/array-of-texture-swizzle.msl2.argument.discrete.swizzle.frag b/3rdparty/spirv-cross/reference/shaders-msl/frag/array-of-texture-swizzle.msl2.argument.discrete.swizzle.frag new file mode 100644 index 000000000..5e749b3ab --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-msl/frag/array-of-texture-swizzle.msl2.argument.discrete.swizzle.frag @@ -0,0 +1,171 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct spvDescriptorSetBuffer0 +{ + constant uint* spvSwizzleConstants [[id(0)]]; + array, 4> uSampler0 [[id(1)]]; + array uSampler0Smplr [[id(5)]]; +}; + +struct main0_out +{ + float4 FragColor [[color(0)]]; +}; + +struct main0_in +{ + float2 vUV [[user(locn0)]]; +}; + +enum class spvSwizzle : uint +{ + none = 0, + zero, + one, + red, + green, + blue, + alpha +}; + +template struct spvRemoveReference { typedef T type; }; +template struct spvRemoveReference { typedef T type; }; +template struct spvRemoveReference { typedef T type; }; +template inline constexpr thread T&& spvForward(thread typename spvRemoveReference::type& x) +{ + return static_cast(x); +} +template inline constexpr thread T&& spvForward(thread typename spvRemoveReference::type&& x) +{ + return static_cast(x); +} + +template +inline T spvGetSwizzle(vec x, T c, spvSwizzle s) +{ + switch (s) + { + case spvSwizzle::none: + return c; + case spvSwizzle::zero: + return 0; + case spvSwizzle::one: + return 1; + case spvSwizzle::red: + return x.r; + case spvSwizzle::green: + return x.g; + case spvSwizzle::blue: + return x.b; + case spvSwizzle::alpha: + return x.a; + } +} + +// Wrapper function that swizzles texture samples and fetches. +template +inline vec spvTextureSwizzle(vec x, uint s) +{ + if (!s) + return x; + return vec(spvGetSwizzle(x, x.r, spvSwizzle((s >> 0) & 0xFF)), spvGetSwizzle(x, x.g, spvSwizzle((s >> 8) & 0xFF)), spvGetSwizzle(x, x.b, spvSwizzle((s >> 16) & 0xFF)), spvGetSwizzle(x, x.a, spvSwizzle((s >> 24) & 0xFF))); +} + +template +inline T spvTextureSwizzle(T x, uint s) +{ + return spvTextureSwizzle(vec(x, 0, 0, 1), s).x; +} + +// Wrapper function that swizzles texture gathers. +template +inline vec spvGatherSwizzle(sampler s, const thread Tex& t, Ts... params, component c, uint sw) METAL_CONST_ARG(c) +{ + if (sw) + { + switch (spvSwizzle((sw >> (uint(c) * 8)) & 0xFF)) + { + case spvSwizzle::none: + break; + case spvSwizzle::zero: + return vec(0, 0, 0, 0); + case spvSwizzle::one: + return vec(1, 1, 1, 1); + case spvSwizzle::red: + return t.gather(s, spvForward(params)..., component::x); + case spvSwizzle::green: + return t.gather(s, spvForward(params)..., component::y); + case spvSwizzle::blue: + return t.gather(s, spvForward(params)..., component::z); + case spvSwizzle::alpha: + return t.gather(s, spvForward(params)..., component::w); + } + } + switch (c) + { + case component::x: + return t.gather(s, spvForward(params)..., component::x); + case component::y: + return t.gather(s, spvForward(params)..., component::y); + case component::z: + return t.gather(s, spvForward(params)..., component::z); + case component::w: + return t.gather(s, spvForward(params)..., component::w); + } +} + +// Wrapper function that swizzles depth texture gathers. +template +inline vec spvGatherCompareSwizzle(sampler s, const thread Tex& t, Ts... params, uint sw) +{ + if (sw) + { + switch (spvSwizzle(sw & 0xFF)) + { + case spvSwizzle::none: + case spvSwizzle::red: + break; + case spvSwizzle::zero: + case spvSwizzle::green: + case spvSwizzle::blue: + case spvSwizzle::alpha: + return vec(0, 0, 0, 0); + case spvSwizzle::one: + return vec(1, 1, 1, 1); + } + } + return t.gather_compare(s, spvForward(params)...); +} + +float4 sample_in_func_1(thread const array, 4> uSampler0, thread const array uSampler0Smplr, constant uint32_t* uSampler0Swzl, thread float2& vUV) +{ + return spvTextureSwizzle(uSampler0[2].sample(uSampler0Smplr[2], vUV), uSampler0Swzl[2]); +} + +float4 sample_in_func_2(thread float2& vUV, thread texture2d uSampler1, thread const sampler uSampler1Smplr, constant uint32_t& uSampler1Swzl) +{ + return spvTextureSwizzle(uSampler1.sample(uSampler1Smplr, vUV), uSampler1Swzl); +} + +float4 sample_single_in_func(thread const texture2d s, thread const sampler sSmplr, constant uint32_t& sSwzl, thread float2& vUV) +{ + return spvTextureSwizzle(s.sample(sSmplr, vUV), sSwzl); +} + +fragment main0_out main0(main0_in in [[stage_in]], constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant uint* spvSwizzleConstants [[buffer(30)]], texture2d uSampler1 [[texture(0)]], sampler uSampler1Smplr [[sampler(0)]]) +{ + main0_out out = {}; + constant uint32_t* spvDescriptorSet0_uSampler0Swzl = &spvDescriptorSet0.spvSwizzleConstants[1]; + constant uint32_t& uSampler1Swzl = spvSwizzleConstants[0]; + out.FragColor = sample_in_func_1(spvDescriptorSet0.uSampler0, spvDescriptorSet0.uSampler0Smplr, spvDescriptorSet0_uSampler0Swzl, in.vUV); + out.FragColor += sample_in_func_2(in.vUV, uSampler1, uSampler1Smplr, uSampler1Swzl); + out.FragColor += sample_single_in_func(spvDescriptorSet0.uSampler0[1], spvDescriptorSet0.uSampler0Smplr[1], spvDescriptorSet0_uSampler0Swzl[1], in.vUV); + out.FragColor += sample_single_in_func(uSampler1, uSampler1Smplr, uSampler1Swzl, in.vUV); + return out; +} + diff --git a/3rdparty/spirv-cross/reference/shaders-msl/frag/array-of-texture-swizzle.msl2.swizzle.frag b/3rdparty/spirv-cross/reference/shaders-msl/frag/array-of-texture-swizzle.msl2.swizzle.frag new file mode 100644 index 000000000..0e4972c6d --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-msl/frag/array-of-texture-swizzle.msl2.swizzle.frag @@ -0,0 +1,156 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct main0_out +{ + float4 FragColor [[color(0)]]; +}; + +struct main0_in +{ + float2 vUV [[user(locn0)]]; +}; + +enum class spvSwizzle : uint +{ + none = 0, + zero, + one, + red, + green, + blue, + alpha +}; + +template struct spvRemoveReference { typedef T type; }; +template struct spvRemoveReference { typedef T type; }; +template struct spvRemoveReference { typedef T type; }; +template inline constexpr thread T&& spvForward(thread typename spvRemoveReference::type& x) +{ + return static_cast(x); +} +template inline constexpr thread T&& spvForward(thread typename spvRemoveReference::type&& x) +{ + return static_cast(x); +} + +template +inline T spvGetSwizzle(vec x, T c, spvSwizzle s) +{ + switch (s) + { + case spvSwizzle::none: + return c; + case spvSwizzle::zero: + return 0; + case spvSwizzle::one: + return 1; + case spvSwizzle::red: + return x.r; + case spvSwizzle::green: + return x.g; + case spvSwizzle::blue: + return x.b; + case spvSwizzle::alpha: + return x.a; + } +} + +// Wrapper function that swizzles texture samples and fetches. +template +inline vec spvTextureSwizzle(vec x, uint s) +{ + if (!s) + return x; + return vec(spvGetSwizzle(x, x.r, spvSwizzle((s >> 0) & 0xFF)), spvGetSwizzle(x, x.g, spvSwizzle((s >> 8) & 0xFF)), spvGetSwizzle(x, x.b, spvSwizzle((s >> 16) & 0xFF)), spvGetSwizzle(x, x.a, spvSwizzle((s >> 24) & 0xFF))); +} + +template +inline T spvTextureSwizzle(T x, uint s) +{ + return spvTextureSwizzle(vec(x, 0, 0, 1), s).x; +} + +// Wrapper function that swizzles texture gathers. +template +inline vec spvGatherSwizzle(sampler s, const thread Tex& t, Ts... params, component c, uint sw) METAL_CONST_ARG(c) +{ + if (sw) + { + switch (spvSwizzle((sw >> (uint(c) * 8)) & 0xFF)) + { + case spvSwizzle::none: + break; + case spvSwizzle::zero: + return vec(0, 0, 0, 0); + case spvSwizzle::one: + return vec(1, 1, 1, 1); + case spvSwizzle::red: + return t.gather(s, spvForward(params)..., component::x); + case spvSwizzle::green: + return t.gather(s, spvForward(params)..., component::y); + case spvSwizzle::blue: + return t.gather(s, spvForward(params)..., component::z); + case spvSwizzle::alpha: + return t.gather(s, spvForward(params)..., component::w); + } + } + switch (c) + { + case component::x: + return t.gather(s, spvForward(params)..., component::x); + case component::y: + return t.gather(s, spvForward(params)..., component::y); + case component::z: + return t.gather(s, spvForward(params)..., component::z); + case component::w: + return t.gather(s, spvForward(params)..., component::w); + } +} + +// Wrapper function that swizzles depth texture gathers. +template +inline vec spvGatherCompareSwizzle(sampler s, const thread Tex& t, Ts... params, uint sw) +{ + if (sw) + { + switch (spvSwizzle(sw & 0xFF)) + { + case spvSwizzle::none: + case spvSwizzle::red: + break; + case spvSwizzle::zero: + case spvSwizzle::green: + case spvSwizzle::blue: + case spvSwizzle::alpha: + return vec(0, 0, 0, 0); + case spvSwizzle::one: + return vec(1, 1, 1, 1); + } + } + return t.gather_compare(s, spvForward(params)...); +} + +float4 sample_in_func(thread const array, 4> uSampler, thread const array uSamplerSmplr, constant uint32_t* uSamplerSwzl, thread float2& vUV) +{ + return spvTextureSwizzle(uSampler[2].sample(uSamplerSmplr[2], vUV), uSamplerSwzl[2]); +} + +float4 sample_single_in_func(thread const texture2d s, thread const sampler sSmplr, constant uint32_t& sSwzl, thread float2& vUV) +{ + return spvTextureSwizzle(s.sample(sSmplr, vUV), sSwzl); +} + +fragment main0_out main0(main0_in in [[stage_in]], constant uint* spvSwizzleConstants [[buffer(30)]], array, 4> uSampler [[texture(0)]], array uSamplerSmplr [[sampler(0)]]) +{ + main0_out out = {}; + constant uint32_t* uSamplerSwzl = &spvSwizzleConstants[0]; + out.FragColor = sample_in_func(uSampler, uSamplerSmplr, uSamplerSwzl, in.vUV); + out.FragColor += sample_single_in_func(uSampler[1], uSamplerSmplr[1], uSamplerSwzl[1], in.vUV); + return out; +} + diff --git a/3rdparty/spirv-cross/reference/shaders-msl/frag/nonuniform-qualifier.msl2.frag b/3rdparty/spirv-cross/reference/shaders-msl/frag/nonuniform-qualifier.msl2.frag new file mode 100644 index 000000000..377a27d30 --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-msl/frag/nonuniform-qualifier.msl2.frag @@ -0,0 +1,51 @@ +#include +#include + +using namespace metal; + +struct UBO +{ + float4 v[64]; +}; + +struct SSBO +{ + float4 v[1]; +}; + +struct main0_out +{ + float4 FragColor [[color(0)]]; +}; + +struct main0_in +{ + int vIndex [[user(locn0)]]; + float2 vUV [[user(locn1)]]; +}; + +fragment main0_out main0(main0_in in [[stage_in]], constant UBO* ubos_0 [[buffer(0)]], constant UBO* ubos_1 [[buffer(1)]], const device SSBO* ssbos_0 [[buffer(2)]], const device SSBO* ssbos_1 [[buffer(3)]], array, 8> uSamplers [[texture(0)]], array, 8> uCombinedSamplers [[texture(8)]], array uSamps [[sampler(1)]], array uCombinedSamplersSmplr [[sampler(8)]]) +{ + constant UBO* ubos[] = + { + ubos_0, + ubos_1, + }; + + const device SSBO* ssbos[] = + { + ssbos_0, + ssbos_1, + }; + + main0_out out = {}; + int i = in.vIndex; + int _24 = i + 10; + out.FragColor = uSamplers[_24].sample(uSamps[i + 40], in.vUV); + int _50 = i + 10; + out.FragColor = uCombinedSamplers[_50].sample(uCombinedSamplersSmplr[_50], in.vUV); + out.FragColor += ubos[(i + 20)]->v[i + 40]; + out.FragColor += ssbos[(i + 50)]->v[i + 60]; + return out; +} + diff --git a/3rdparty/spirv-cross/reference/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.comp b/3rdparty/spirv-cross/reference/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.comp new file mode 100644 index 000000000..e52fb209f --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.comp @@ -0,0 +1,146 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct SSBO +{ + float FragColor; +}; + +inline uint4 spvSubgroupBallot(bool value) +{ + simd_vote vote = simd_ballot(value); + // simd_ballot() returns a 64-bit integer-like object, but + // SPIR-V callers expect a uint4. We must convert. + // FIXME: This won't include higher bits if Apple ever supports + // 128 lanes in an SIMD-group. + return uint4((uint)((simd_vote::vote_t)vote & 0xFFFFFFFF), (uint)(((simd_vote::vote_t)vote >> 32) & 0xFFFFFFFF), 0, 0); +} + +inline bool spvSubgroupBallotBitExtract(uint4 ballot, uint bit) +{ + return !!extract_bits(ballot[bit / 32], bit % 32, 1); +} + +inline uint spvSubgroupBallotFindLSB(uint4 ballot) +{ + return select(ctz(ballot.x), select(32 + ctz(ballot.y), select(64 + ctz(ballot.z), select(96 + ctz(ballot.w), uint(-1), ballot.w == 0), ballot.z == 0), ballot.y == 0), ballot.x == 0); +} + +inline uint spvSubgroupBallotFindMSB(uint4 ballot) +{ + return select(128 - (clz(ballot.w) + 1), select(96 - (clz(ballot.z) + 1), select(64 - (clz(ballot.y) + 1), select(32 - (clz(ballot.x) + 1), uint(-1), ballot.x == 0), ballot.y == 0), ballot.z == 0), ballot.w == 0); +} + +inline uint spvSubgroupBallotBitCount(uint4 ballot) +{ + return popcount(ballot.x) + popcount(ballot.y) + popcount(ballot.z) + popcount(ballot.w); +} + +inline uint spvSubgroupBallotInclusiveBitCount(uint4 ballot, uint gl_SubgroupInvocationID) +{ + uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID + 1, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID + 1 - 32, 0)), uint2(0)); + return spvSubgroupBallotBitCount(ballot & mask); +} + +inline uint spvSubgroupBallotExclusiveBitCount(uint4 ballot, uint gl_SubgroupInvocationID) +{ + uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID - 32, 0)), uint2(0)); + return spvSubgroupBallotBitCount(ballot & mask); +} + +template +inline bool spvSubgroupAllEqual(T value) +{ + return simd_all(value == simd_broadcast_first(value)); +} + +template<> +inline bool spvSubgroupAllEqual(bool value) +{ + return simd_all(value) || !simd_any(value); +} + +kernel void main0(device SSBO& _9 [[buffer(0)]], uint gl_NumSubgroups [[simdgroups_per_threadgroup]], uint gl_SubgroupID [[simdgroup_index_in_threadgroup]], uint gl_SubgroupSize [[thread_execution_width]], uint gl_SubgroupInvocationID [[thread_index_in_simdgroup]]) +{ + uint4 gl_SubgroupEqMask = 27 > 32 ? uint4(0, (1 << (gl_SubgroupInvocationID - 32)), uint2(0)) : uint4(1 << gl_SubgroupInvocationID, uint3(0)); + uint4 gl_SubgroupGeMask = uint4(extract_bits(0xFFFFFFFF, min(gl_SubgroupInvocationID, 32u), (uint)max(min((int)gl_SubgroupSize, 32) - (int)gl_SubgroupInvocationID, 0)), extract_bits(0xFFFFFFFF, (uint)max((int)gl_SubgroupInvocationID - 32, 0), (uint)max((int)gl_SubgroupSize - (int)max(gl_SubgroupInvocationID, 32u), 0)), uint2(0)); + uint4 gl_SubgroupGtMask = uint4(extract_bits(0xFFFFFFFF, min(gl_SubgroupInvocationID + 1, 32u), (uint)max(min((int)gl_SubgroupSize, 32) - (int)gl_SubgroupInvocationID - 1, 0)), extract_bits(0xFFFFFFFF, (uint)max((int)gl_SubgroupInvocationID + 1 - 32, 0), (uint)max((int)gl_SubgroupSize - (int)max(gl_SubgroupInvocationID + 1, 32u), 0)), uint2(0)); + uint4 gl_SubgroupLeMask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID + 1, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID + 1 - 32, 0)), uint2(0)); + uint4 gl_SubgroupLtMask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID - 32, 0)), uint2(0)); + _9.FragColor = float(gl_NumSubgroups); + _9.FragColor = float(gl_SubgroupID); + _9.FragColor = float(gl_SubgroupSize); + _9.FragColor = float(gl_SubgroupInvocationID); + simdgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture); + simdgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture); + simdgroup_barrier(mem_flags::mem_device); + simdgroup_barrier(mem_flags::mem_threadgroup); + simdgroup_barrier(mem_flags::mem_texture); + bool elected = simd_is_first(); + _9.FragColor = float4(gl_SubgroupEqMask).x; + _9.FragColor = float4(gl_SubgroupGeMask).x; + _9.FragColor = float4(gl_SubgroupGtMask).x; + _9.FragColor = float4(gl_SubgroupLeMask).x; + _9.FragColor = float4(gl_SubgroupLtMask).x; + float4 broadcasted = simd_broadcast(float4(10.0), 8u); + float3 first = simd_broadcast_first(float3(20.0)); + uint4 ballot_value = spvSubgroupBallot(true); + bool inverse_ballot_value = spvSubgroupBallotBitExtract(ballot_value, gl_SubgroupInvocationID); + bool bit_extracted = spvSubgroupBallotBitExtract(uint4(10u), 8u); + uint bit_count = spvSubgroupBallotBitCount(ballot_value); + uint inclusive_bit_count = spvSubgroupBallotInclusiveBitCount(ballot_value, gl_SubgroupInvocationID); + uint exclusive_bit_count = spvSubgroupBallotExclusiveBitCount(ballot_value, gl_SubgroupInvocationID); + uint lsb = spvSubgroupBallotFindLSB(ballot_value); + uint msb = spvSubgroupBallotFindMSB(ballot_value); + uint shuffled = simd_shuffle(10u, 8u); + uint shuffled_xor = simd_shuffle_xor(30u, 8u); + uint shuffled_up = simd_shuffle_up(20u, 4u); + uint shuffled_down = simd_shuffle_down(20u, 4u); + bool has_all = simd_all(true); + bool has_any = simd_any(true); + bool has_equal = spvSubgroupAllEqual(0); + has_equal = spvSubgroupAllEqual(true); + float4 added = simd_sum(float4(20.0)); + int4 iadded = simd_sum(int4(20)); + float4 multiplied = simd_product(float4(20.0)); + int4 imultiplied = simd_product(int4(20)); + float4 lo = simd_min(float4(20.0)); + float4 hi = simd_max(float4(20.0)); + int4 slo = simd_min(int4(20)); + int4 shi = simd_max(int4(20)); + uint4 ulo = simd_min(uint4(20u)); + uint4 uhi = simd_max(uint4(20u)); + uint4 anded = simd_and(ballot_value); + uint4 ored = simd_or(ballot_value); + uint4 xored = simd_xor(ballot_value); + added = simd_prefix_inclusive_sum(added); + iadded = simd_prefix_inclusive_sum(iadded); + multiplied = simd_prefix_inclusive_product(multiplied); + imultiplied = simd_prefix_inclusive_product(imultiplied); + added = simd_prefix_exclusive_sum(multiplied); + multiplied = simd_prefix_exclusive_product(multiplied); + iadded = simd_prefix_exclusive_sum(imultiplied); + imultiplied = simd_prefix_exclusive_product(imultiplied); + added = quad_sum(added); + multiplied = quad_product(multiplied); + iadded = quad_sum(iadded); + imultiplied = quad_product(imultiplied); + lo = quad_min(lo); + hi = quad_max(hi); + ulo = quad_min(ulo); + uhi = quad_max(uhi); + slo = quad_min(slo); + shi = quad_max(shi); + anded = quad_and(anded); + ored = quad_or(ored); + xored = quad_xor(xored); + float4 swap_horiz = quad_shuffle_xor(float4(20.0), 1u); + float4 swap_vertical = quad_shuffle_xor(float4(20.0), 2u); + float4 swap_diagonal = quad_shuffle_xor(float4(20.0), 3u); + float4 quad_broadcast0 = quad_broadcast(float4(20.0), 3u); +} + diff --git a/3rdparty/spirv-cross/reference/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.ios.comp b/3rdparty/spirv-cross/reference/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.ios.comp new file mode 100644 index 000000000..84fcb9c3a --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.ios.comp @@ -0,0 +1,31 @@ +#include +#include + +using namespace metal; + +struct SSBO +{ + float FragColor; +}; + +kernel void main0(device SSBO& _9 [[buffer(0)]], uint gl_NumSubgroups [[quadgroups_per_threadgroup]], uint gl_SubgroupID [[quadgroup_index_in_threadgroup]], uint gl_SubgroupSize [[thread_execution_width]], uint gl_SubgroupInvocationID [[thread_index_in_quadgroup]]) +{ + _9.FragColor = float(gl_NumSubgroups); + _9.FragColor = float(gl_SubgroupID); + _9.FragColor = float(gl_SubgroupSize); + _9.FragColor = float(gl_SubgroupInvocationID); + simdgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture); + simdgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture); + simdgroup_barrier(mem_flags::mem_device); + simdgroup_barrier(mem_flags::mem_threadgroup); + simdgroup_barrier(mem_flags::mem_texture); + uint shuffled = quad_shuffle(10u, 8u); + uint shuffled_xor = quad_shuffle_xor(30u, 8u); + uint shuffled_up = quad_shuffle_up(20u, 4u); + uint shuffled_down = quad_shuffle_down(20u, 4u); + float4 swap_horiz = quad_shuffle_xor(float4(20.0), 1u); + float4 swap_vertical = quad_shuffle_xor(float4(20.0), 2u); + float4 swap_diagonal = quad_shuffle_xor(float4(20.0), 3u); + float4 quad_broadcast0 = quad_broadcast(float4(20.0), 3u); +} + diff --git a/3rdparty/spirv-cross/reference/shaders/comp/ssbo-array-length.comp b/3rdparty/spirv-cross/reference/shaders/comp/ssbo-array-length.comp new file mode 100644 index 000000000..ddc666e9b --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders/comp/ssbo-array-length.comp @@ -0,0 +1,14 @@ +#version 450 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(binding = 1, std140) buffer SSBO +{ + uint size; + float v[]; +} _11; + +void main() +{ + _11.size = uint(int(uint(_11.v.length()))); +} + diff --git a/3rdparty/spirv-cross/reference/shaders/vulkan/comp/buffer-reference-bitcast.nocompat.vk.comp.vk b/3rdparty/spirv-cross/reference/shaders/vulkan/comp/buffer-reference-bitcast.nocompat.vk.comp.vk new file mode 100644 index 000000000..5752f81b2 --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders/vulkan/comp/buffer-reference-bitcast.nocompat.vk.comp.vk @@ -0,0 +1,26 @@ +#version 450 +#extension GL_EXT_buffer_reference : require +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(buffer_reference) buffer PtrUint; +layout(buffer_reference) buffer PtrInt; +layout(buffer_reference, std430) buffer PtrUint +{ + uint value; +}; + +layout(buffer_reference, std430) buffer PtrInt +{ + int value; +}; + +layout(set = 0, binding = 0, std430) buffer Buf +{ + PtrUint ptr; +} _11; + +void main() +{ + PtrInt(_11.ptr).value = 10; +} + diff --git a/3rdparty/spirv-cross/reference/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp.vk b/3rdparty/spirv-cross/reference/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp.vk index 610d60cb4..e22974114 100644 --- a/3rdparty/spirv-cross/reference/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp.vk +++ b/3rdparty/spirv-cross/reference/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp.vk @@ -22,7 +22,7 @@ void copy_node(restrict Node dst, restrict Node a, restrict Node b) dst.value = a.value + b.value; } -void overwrite_node(out Node dst, Node src) +void overwrite_node(out restrict Node dst, restrict Node src) { dst = src; } diff --git a/3rdparty/spirv-cross/shaders-hlsl/comp/ssbo-array-length.comp b/3rdparty/spirv-cross/shaders-hlsl/comp/ssbo-array-length.comp new file mode 100644 index 000000000..3ad4b9515 --- /dev/null +++ b/3rdparty/spirv-cross/shaders-hlsl/comp/ssbo-array-length.comp @@ -0,0 +1,12 @@ +#version 450 +layout(local_size_x = 1) in; +layout(set = 0, binding = 1, std140) buffer SSBO +{ + uint size; + float v[]; +}; + +void main() +{ + size = v.length(); +} diff --git a/3rdparty/spirv-cross/shaders-hlsl/frag/nonuniform-qualifier.nonuniformresource.sm51.frag b/3rdparty/spirv-cross/shaders-hlsl/frag/nonuniform-qualifier.nonuniformresource.sm51.frag new file mode 100644 index 000000000..0aadd1488 --- /dev/null +++ b/3rdparty/spirv-cross/shaders-hlsl/frag/nonuniform-qualifier.nonuniformresource.sm51.frag @@ -0,0 +1,28 @@ +#version 450 +#extension GL_EXT_nonuniform_qualifier : require + +layout(set = 0, binding = 0) uniform texture2D uSamplers[]; +layout(set = 1, binding = 0) uniform sampler2D uCombinedSamplers[]; +layout(set = 2, binding = 0) uniform sampler uSamps[]; +layout(location = 0) flat in int vIndex; +layout(location = 1) in vec2 vUV; +layout(location = 0) out vec4 FragColor; + +layout(set = 3, binding = 0) uniform UBO +{ + vec4 v[64]; +} ubos[]; + +layout(set = 4, binding = 0) readonly buffer SSBO +{ + vec4 v[]; +} ssbos[]; + +void main() +{ + int i = vIndex; + FragColor = texture(sampler2D(uSamplers[nonuniformEXT(i + 10)], uSamps[nonuniformEXT(i + 40)]), vUV); + FragColor = texture(uCombinedSamplers[nonuniformEXT(i + 10)], vUV); + FragColor += ubos[nonuniformEXT(i + 20)].v[nonuniformEXT(i + 40)]; + FragColor += ssbos[nonuniformEXT(i + 50)].v[nonuniformEXT(i + 60)]; +} diff --git a/3rdparty/spirv-cross/shaders-msl/comp/barriers.comp b/3rdparty/spirv-cross/shaders-msl/comp/barriers.comp index c49b62636..7e0ea42d4 100644 --- a/3rdparty/spirv-cross/shaders-msl/comp/barriers.comp +++ b/3rdparty/spirv-cross/shaders-msl/comp/barriers.comp @@ -11,12 +11,10 @@ void full_barrier() memoryBarrier(); } -#if 0 void image_barrier() { memoryBarrierImage(); } -#endif void buffer_barrier() { @@ -40,13 +38,11 @@ void full_barrier_exec() barrier(); } -#if 0 void image_barrier_exec() { memoryBarrierImage(); barrier(); } -#endif void buffer_barrier_exec() { @@ -69,13 +65,13 @@ void main() { barrier_shared(); full_barrier(); - //image_barrier(); + image_barrier(); buffer_barrier(); group_barrier(); barrier_shared_exec(); full_barrier_exec(); - //image_barrier_exec(); + image_barrier_exec(); buffer_barrier_exec(); group_barrier_exec(); diff --git a/3rdparty/spirv-cross/shaders-msl/frag/array-of-texture-swizzle.msl2.argument.discrete.swizzle.frag b/3rdparty/spirv-cross/shaders-msl/frag/array-of-texture-swizzle.msl2.argument.discrete.swizzle.frag new file mode 100644 index 000000000..556cc9ca3 --- /dev/null +++ b/3rdparty/spirv-cross/shaders-msl/frag/array-of-texture-swizzle.msl2.argument.discrete.swizzle.frag @@ -0,0 +1,31 @@ +#version 450 + +layout(set = 0, binding = 1) uniform sampler2D uSampler0[4]; +layout(set = 2, binding = 0) uniform sampler2D uSampler1; +layout(set = 1, binding = 4) uniform sampler2D uSamp; +layout(location = 0) in vec2 vUV; + +layout(location = 0) out vec4 FragColor; + +vec4 sample_in_func_1() +{ + return texture(uSampler0[2], vUV); +} + +vec4 sample_in_func_2() +{ + return texture(uSampler1, vUV); +} + +vec4 sample_single_in_func(sampler2D s) +{ + return texture(s, vUV); +} + +void main() +{ + FragColor = sample_in_func_1(); + FragColor += sample_in_func_2(); + FragColor += sample_single_in_func(uSampler0[1]); + FragColor += sample_single_in_func(uSampler1); +} diff --git a/3rdparty/spirv-cross/shaders-msl/frag/array-of-texture-swizzle.msl2.swizzle.frag b/3rdparty/spirv-cross/shaders-msl/frag/array-of-texture-swizzle.msl2.swizzle.frag new file mode 100644 index 000000000..4694aa3ea --- /dev/null +++ b/3rdparty/spirv-cross/shaders-msl/frag/array-of-texture-swizzle.msl2.swizzle.frag @@ -0,0 +1,23 @@ +#version 450 + +layout(set = 0, binding = 0) uniform sampler2D uSampler[4]; +layout(set = 0, binding = 1) uniform sampler2D uSamp; +layout(location = 0) in vec2 vUV; + +layout(location = 0) out vec4 FragColor; + +vec4 sample_in_func() +{ + return texture(uSampler[2], vUV); +} + +vec4 sample_single_in_func(sampler2D s) +{ + return texture(s, vUV); +} + +void main() +{ + FragColor = sample_in_func(); + FragColor += sample_single_in_func(uSampler[1]); +} diff --git a/3rdparty/spirv-cross/shaders-msl/frag/nonuniform-qualifier.msl2.frag b/3rdparty/spirv-cross/shaders-msl/frag/nonuniform-qualifier.msl2.frag new file mode 100644 index 000000000..ba9dd7fb2 --- /dev/null +++ b/3rdparty/spirv-cross/shaders-msl/frag/nonuniform-qualifier.msl2.frag @@ -0,0 +1,28 @@ +#version 450 +#extension GL_EXT_nonuniform_qualifier : require + +layout(binding = 0) uniform texture2D uSamplers[8]; +layout(binding = 8) uniform sampler2D uCombinedSamplers[8]; +layout(binding = 1) uniform sampler uSamps[7]; +layout(location = 0) flat in int vIndex; +layout(location = 1) in vec2 vUV; +layout(location = 0) out vec4 FragColor; + +layout(set = 0, binding = 0) uniform UBO +{ + vec4 v[64]; +} ubos[2]; + +layout(set = 0, binding = 2) readonly buffer SSBO +{ + vec4 v[]; +} ssbos[2]; + +void main() +{ + int i = vIndex; + FragColor = texture(sampler2D(uSamplers[nonuniformEXT(i + 10)], uSamps[nonuniformEXT(i + 40)]), vUV); + FragColor = texture(uCombinedSamplers[nonuniformEXT(i + 10)], vUV); + FragColor += ubos[nonuniformEXT(i + 20)].v[nonuniformEXT(i + 40)]; + FragColor += ssbos[nonuniformEXT(i + 50)].v[nonuniformEXT(i + 60)]; +} diff --git a/3rdparty/spirv-cross/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.comp b/3rdparty/spirv-cross/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.comp new file mode 100644 index 000000000..f8f5133f8 --- /dev/null +++ b/3rdparty/spirv-cross/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.comp @@ -0,0 +1,126 @@ +#version 450 +#extension GL_KHR_shader_subgroup_basic : require +#extension GL_KHR_shader_subgroup_ballot : require +#extension GL_KHR_shader_subgroup_vote : require +#extension GL_KHR_shader_subgroup_shuffle : require +#extension GL_KHR_shader_subgroup_shuffle_relative : require +#extension GL_KHR_shader_subgroup_arithmetic : require +#extension GL_KHR_shader_subgroup_clustered : require +#extension GL_KHR_shader_subgroup_quad : require +layout(local_size_x = 1) in; + +layout(std430, binding = 0) buffer SSBO +{ + float FragColor; +}; + +void main() +{ + // basic + FragColor = float(gl_NumSubgroups); + FragColor = float(gl_SubgroupID); + FragColor = float(gl_SubgroupSize); + FragColor = float(gl_SubgroupInvocationID); + subgroupBarrier(); + subgroupMemoryBarrier(); + subgroupMemoryBarrierBuffer(); + subgroupMemoryBarrierShared(); + subgroupMemoryBarrierImage(); + bool elected = subgroupElect(); + + // ballot + FragColor = float(gl_SubgroupEqMask); + FragColor = float(gl_SubgroupGeMask); + FragColor = float(gl_SubgroupGtMask); + FragColor = float(gl_SubgroupLeMask); + FragColor = float(gl_SubgroupLtMask); + vec4 broadcasted = subgroupBroadcast(vec4(10.0), 8u); + vec3 first = subgroupBroadcastFirst(vec3(20.0)); + uvec4 ballot_value = subgroupBallot(true); + bool inverse_ballot_value = subgroupInverseBallot(ballot_value); + bool bit_extracted = subgroupBallotBitExtract(uvec4(10u), 8u); + uint bit_count = subgroupBallotBitCount(ballot_value); + uint inclusive_bit_count = subgroupBallotInclusiveBitCount(ballot_value); + uint exclusive_bit_count = subgroupBallotExclusiveBitCount(ballot_value); + uint lsb = subgroupBallotFindLSB(ballot_value); + uint msb = subgroupBallotFindMSB(ballot_value); + + // shuffle + uint shuffled = subgroupShuffle(10u, 8u); + uint shuffled_xor = subgroupShuffleXor(30u, 8u); + + // shuffle relative + uint shuffled_up = subgroupShuffleUp(20u, 4u); + uint shuffled_down = subgroupShuffleDown(20u, 4u); + + // vote + bool has_all = subgroupAll(true); + bool has_any = subgroupAny(true); + bool has_equal = subgroupAllEqual(0); + has_equal = subgroupAllEqual(true); + + // arithmetic + vec4 added = subgroupAdd(vec4(20.0)); + ivec4 iadded = subgroupAdd(ivec4(20)); + vec4 multiplied = subgroupMul(vec4(20.0)); + ivec4 imultiplied = subgroupMul(ivec4(20)); + vec4 lo = subgroupMin(vec4(20.0)); + vec4 hi = subgroupMax(vec4(20.0)); + ivec4 slo = subgroupMin(ivec4(20)); + ivec4 shi = subgroupMax(ivec4(20)); + uvec4 ulo = subgroupMin(uvec4(20)); + uvec4 uhi = subgroupMax(uvec4(20)); + uvec4 anded = subgroupAnd(ballot_value); + uvec4 ored = subgroupOr(ballot_value); + uvec4 xored = subgroupXor(ballot_value); + + added = subgroupInclusiveAdd(added); + iadded = subgroupInclusiveAdd(iadded); + multiplied = subgroupInclusiveMul(multiplied); + imultiplied = subgroupInclusiveMul(imultiplied); + //lo = subgroupInclusiveMin(lo); // FIXME: Unsupported by Metal + //hi = subgroupInclusiveMax(hi); + //slo = subgroupInclusiveMin(slo); + //shi = subgroupInclusiveMax(shi); + //ulo = subgroupInclusiveMin(ulo); + //uhi = subgroupInclusiveMax(uhi); + //anded = subgroupInclusiveAnd(anded); + //ored = subgroupInclusiveOr(ored); + //xored = subgroupInclusiveXor(ored); + //added = subgroupExclusiveAdd(lo); + + added = subgroupExclusiveAdd(multiplied); + multiplied = subgroupExclusiveMul(multiplied); + iadded = subgroupExclusiveAdd(imultiplied); + imultiplied = subgroupExclusiveMul(imultiplied); + //lo = subgroupExclusiveMin(lo); // FIXME: Unsupported by Metal + //hi = subgroupExclusiveMax(hi); + //ulo = subgroupExclusiveMin(ulo); + //uhi = subgroupExclusiveMax(uhi); + //slo = subgroupExclusiveMin(slo); + //shi = subgroupExclusiveMax(shi); + //anded = subgroupExclusiveAnd(anded); + //ored = subgroupExclusiveOr(ored); + //xored = subgroupExclusiveXor(ored); + + // clustered + added = subgroupClusteredAdd(added, 4u); + multiplied = subgroupClusteredMul(multiplied, 4u); + iadded = subgroupClusteredAdd(iadded, 4u); + imultiplied = subgroupClusteredMul(imultiplied, 4u); + lo = subgroupClusteredMin(lo, 4u); + hi = subgroupClusteredMax(hi, 4u); + ulo = subgroupClusteredMin(ulo, 4u); + uhi = subgroupClusteredMax(uhi, 4u); + slo = subgroupClusteredMin(slo, 4u); + shi = subgroupClusteredMax(shi, 4u); + anded = subgroupClusteredAnd(anded, 4u); + ored = subgroupClusteredOr(ored, 4u); + xored = subgroupClusteredXor(xored, 4u); + + // quad + vec4 swap_horiz = subgroupQuadSwapHorizontal(vec4(20.0)); + vec4 swap_vertical = subgroupQuadSwapVertical(vec4(20.0)); + vec4 swap_diagonal = subgroupQuadSwapDiagonal(vec4(20.0)); + vec4 quad_broadcast = subgroupQuadBroadcast(vec4(20.0), 3u); +} diff --git a/3rdparty/spirv-cross/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.ios.comp b/3rdparty/spirv-cross/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.ios.comp new file mode 100644 index 000000000..66eb4a207 --- /dev/null +++ b/3rdparty/spirv-cross/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.ios.comp @@ -0,0 +1,41 @@ +#version 450 +#extension GL_KHR_shader_subgroup_basic : require +#extension GL_KHR_shader_subgroup_shuffle : require +#extension GL_KHR_shader_subgroup_shuffle_relative : require +#extension GL_KHR_shader_subgroup_quad : require +layout(local_size_x = 1) in; + +layout(std430, binding = 0) buffer SSBO +{ + float FragColor; +}; + +// Reduced test for functionality exposed on iOS. + +void main() +{ + // basic + FragColor = float(gl_NumSubgroups); + FragColor = float(gl_SubgroupID); + FragColor = float(gl_SubgroupSize); + FragColor = float(gl_SubgroupInvocationID); + subgroupBarrier(); + subgroupMemoryBarrier(); + subgroupMemoryBarrierBuffer(); + subgroupMemoryBarrierShared(); + subgroupMemoryBarrierImage(); + + // shuffle + uint shuffled = subgroupShuffle(10u, 8u); + uint shuffled_xor = subgroupShuffleXor(30u, 8u); + + // shuffle relative + uint shuffled_up = subgroupShuffleUp(20u, 4u); + uint shuffled_down = subgroupShuffleDown(20u, 4u); + + // quad + vec4 swap_horiz = subgroupQuadSwapHorizontal(vec4(20.0)); + vec4 swap_vertical = subgroupQuadSwapVertical(vec4(20.0)); + vec4 swap_diagonal = subgroupQuadSwapDiagonal(vec4(20.0)); + vec4 quad_broadcast = subgroupQuadBroadcast(vec4(20.0), 3u); +} diff --git a/3rdparty/spirv-cross/shaders/comp/ssbo-array-length.comp b/3rdparty/spirv-cross/shaders/comp/ssbo-array-length.comp new file mode 100644 index 000000000..3ad4b9515 --- /dev/null +++ b/3rdparty/spirv-cross/shaders/comp/ssbo-array-length.comp @@ -0,0 +1,12 @@ +#version 450 +layout(local_size_x = 1) in; +layout(set = 0, binding = 1, std140) buffer SSBO +{ + uint size; + float v[]; +}; + +void main() +{ + size = v.length(); +} diff --git a/3rdparty/spirv-cross/shaders/vulkan/comp/buffer-reference-bitcast.nocompat.vk.comp b/3rdparty/spirv-cross/shaders/vulkan/comp/buffer-reference-bitcast.nocompat.vk.comp new file mode 100644 index 000000000..eda904ee3 --- /dev/null +++ b/3rdparty/spirv-cross/shaders/vulkan/comp/buffer-reference-bitcast.nocompat.vk.comp @@ -0,0 +1,22 @@ +#version 450 +#extension GL_EXT_buffer_reference: require + +layout(buffer_reference) buffer PtrUint +{ + uint value; +}; + +layout(buffer_reference) buffer PtrInt +{ + int value; +}; + +layout(set = 0, binding = 0) buffer Buf +{ + PtrUint ptr; +}; + +void main() +{ + PtrInt(ptr).value = 10; +} diff --git a/3rdparty/spirv-cross/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp b/3rdparty/spirv-cross/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp index 624b8c0a0..f08e11139 100644 --- a/3rdparty/spirv-cross/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp +++ b/3rdparty/spirv-cross/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp @@ -21,7 +21,7 @@ void copy_node(restrict Node dst, restrict Node a, restrict Node b) dst.value = a.value + b.value; } -void overwrite_node(out Node dst, Node src) +void overwrite_node(out restrict Node dst, restrict Node src) { dst = src; } diff --git a/3rdparty/spirv-cross/spirv_cpp.cpp b/3rdparty/spirv-cross/spirv_cpp.cpp index 19d93b7b4..90566c1cf 100644 --- a/3rdparty/spirv-cross/spirv_cpp.cpp +++ b/3rdparty/spirv-cross/spirv_cpp.cpp @@ -317,7 +317,7 @@ string CompilerCPP::compile() backend.basic_uint_type = "uint32_t"; backend.swizzle_is_function = true; backend.shared_is_implied = true; - backend.flexible_member_array_supported = false; + backend.unsized_array_supported = false; backend.explicit_struct_type = true; backend.use_initializer_list = true; diff --git a/3rdparty/spirv-cross/spirv_cross.cpp b/3rdparty/spirv-cross/spirv_cross.cpp index 97090af4f..6b66b7460 100644 --- a/3rdparty/spirv-cross/spirv_cross.cpp +++ b/3rdparty/spirv-cross/spirv_cross.cpp @@ -4217,7 +4217,7 @@ Compiler::PhysicalStorageBufferPointerHandler::PhysicalStorageBufferPointerHandl bool Compiler::PhysicalStorageBufferPointerHandler::handle(Op op, const uint32_t *args, uint32_t) { - if (op == OpConvertUToPtr) + if (op == OpConvertUToPtr || op == OpBitcast) { auto &type = compiler.get(args[0]); if (type.storage == StorageClassPhysicalStorageBufferEXT && type.pointer && type.pointer_depth == 1) diff --git a/3rdparty/spirv-cross/spirv_cross_c.cpp b/3rdparty/spirv-cross/spirv_cross_c.cpp index b1f825f39..6811d0274 100644 --- a/3rdparty/spirv-cross/spirv_cross_c.cpp +++ b/3rdparty/spirv-cross/spirv_cross_c.cpp @@ -442,6 +442,9 @@ spvc_result spvc_compiler_options_set_uint(spvc_compiler_options options, spvc_c case SPVC_COMPILER_OPTION_GLSL_EMIT_PUSH_CONSTANT_AS_UNIFORM_BUFFER: options->glsl.emit_push_constant_as_uniform_buffer = value != 0; break; + case SPVC_COMPILER_OPTION_GLSL_EMIT_UNIFORM_BUFFER_AS_PLAIN_UNIFORMS: + options->glsl.emit_uniform_buffer_as_plain_uniforms = value != 0; + break; #endif #if SPIRV_CROSS_C_API_HLSL @@ -471,8 +474,8 @@ spvc_result spvc_compiler_options_set_uint(spvc_compiler_options options, spvc_c options->msl.texel_buffer_texture_width = value; break; - case SPVC_COMPILER_OPTION_MSL_AUX_BUFFER_INDEX: - options->msl.aux_buffer_index = value; + case SPVC_COMPILER_OPTION_MSL_SWIZZLE_BUFFER_INDEX: + options->msl.swizzle_buffer_index = value; break; case SPVC_COMPILER_OPTION_MSL_INDIRECT_PARAMS_BUFFER_INDEX: @@ -723,7 +726,7 @@ spvc_bool spvc_compiler_msl_is_rasterization_disabled(spvc_compiler compiler) #endif } -spvc_bool spvc_compiler_msl_needs_aux_buffer(spvc_compiler compiler) +spvc_bool spvc_compiler_msl_needs_swizzle_buffer(spvc_compiler compiler) { #if SPIRV_CROSS_C_API_MSL if (compiler->backend != SPVC_BACKEND_MSL) @@ -733,13 +736,18 @@ spvc_bool spvc_compiler_msl_needs_aux_buffer(spvc_compiler compiler) } auto &msl = *static_cast(compiler->compiler.get()); - return msl.needs_aux_buffer() ? SPVC_TRUE : SPVC_FALSE; + return msl.needs_swizzle_buffer() ? SPVC_TRUE : SPVC_FALSE; #else compiler->context->report_error("MSL function used on a non-MSL backend."); return SPVC_FALSE; #endif } +spvc_bool spvc_compiler_msl_needs_aux_buffer(spvc_compiler compiler) +{ + return spvc_compiler_msl_needs_swizzle_buffer(compiler); +} + spvc_bool spvc_compiler_msl_needs_output_buffer(spvc_compiler compiler) { #if SPIRV_CROSS_C_API_MSL @@ -1273,6 +1281,11 @@ const char *spvc_compiler_get_member_decoration_string(spvc_compiler compiler, s .c_str(); } +const char *spvc_compiler_get_member_name(spvc_compiler compiler, spvc_type_id id, unsigned member_index) +{ + return compiler->compiler->get_member_name(id, member_index).c_str(); +} + spvc_result spvc_compiler_get_entry_points(spvc_compiler compiler, const spvc_entry_point **entry_points, size_t *num_entry_points) { @@ -1410,7 +1423,7 @@ unsigned spvc_type_get_bit_width(spvc_type type) return type->width; } -unsigned spvc_type_get_SmallVector_size(spvc_type type) +unsigned spvc_type_get_vector_size(spvc_type type) { return type->vecsize; } @@ -1640,6 +1653,32 @@ spvc_constant_id spvc_compiler_get_work_group_size_specialization_constants(spvc return ret; } +spvc_result spvc_compiler_get_active_buffer_ranges(spvc_compiler compiler, + spvc_variable_id id, + const spvc_buffer_range **ranges, + size_t *num_ranges) +{ + SPVC_BEGIN_SAFE_SCOPE + { + auto active_ranges = compiler->compiler->get_active_buffer_ranges(id); + SmallVector translated; + translated.reserve(active_ranges.size()); + for (auto &r : active_ranges) + { + spvc_buffer_range trans = { r.index, r.offset, r.range }; + translated.push_back(trans); + } + + auto ptr = spvc_allocate>(); + ptr->buffer = std::move(translated); + *ranges = ptr->buffer.data(); + *num_ranges = ptr->buffer.size(); + compiler->context->allocations.push_back(std::move(ptr)); + } + SPVC_END_SAFE_SCOPE(compiler->context, SPVC_ERROR_OUT_OF_MEMORY) + return SPVC_SUCCESS; +} + float spvc_constant_get_scalar_fp16(spvc_constant constant, unsigned column, unsigned row) { return constant->scalar_f16(column, row); diff --git a/3rdparty/spirv-cross/spirv_cross_c.h b/3rdparty/spirv-cross/spirv_cross_c.h index ee8d15d9e..43eb924fb 100644 --- a/3rdparty/spirv-cross/spirv_cross_c.h +++ b/3rdparty/spirv-cross/spirv_cross_c.h @@ -33,7 +33,7 @@ extern "C" { /* Bumped if ABI or API breaks backwards compatibility. */ #define SPVC_C_API_VERSION_MAJOR 0 /* Bumped if APIs or enumerations are added in a backwards compatible way. */ -#define SPVC_C_API_VERSION_MINOR 6 +#define SPVC_C_API_VERSION_MINOR 9 /* Bumped if internal implementation details change. */ #define SPVC_C_API_VERSION_PATCH 0 @@ -111,6 +111,14 @@ typedef struct spvc_specialization_constant unsigned constant_id; } spvc_specialization_constant; +/* See C++ API. */ +typedef struct spvc_buffer_range +{ + unsigned index; + size_t offset; + size_t range; +} spvc_buffer_range; + /* See C++ API. */ typedef struct spvc_hlsl_root_constants { @@ -290,9 +298,12 @@ SPVC_PUBLIC_API void spvc_msl_resource_binding_init(spvc_msl_resource_binding *b #define SPVC_MSL_PUSH_CONSTANT_DESC_SET (~(0u)) #define SPVC_MSL_PUSH_CONSTANT_BINDING (0) +#define SPVC_MSL_SWIZZLE_BUFFER_BINDING (~(1u)) + +/* Obsolete. Sticks around for backwards compatibility. */ #define SPVC_MSL_AUX_BUFFER_STRUCT_VERSION 1 -/* Runtime check for incompatibility. */ +/* Runtime check for incompatibility. Obsolete. */ SPVC_PUBLIC_API unsigned spvc_msl_get_aux_buffer_struct_version(void); /* Maps to C++ API. */ @@ -407,7 +418,11 @@ typedef enum spvc_compiler_option SPVC_COMPILER_OPTION_MSL_VERSION = 17 | SPVC_COMPILER_OPTION_MSL_BIT, SPVC_COMPILER_OPTION_MSL_TEXEL_BUFFER_TEXTURE_WIDTH = 18 | SPVC_COMPILER_OPTION_MSL_BIT, + + /* Obsolete, use SWIZZLE_BUFFER_INDEX instead. */ SPVC_COMPILER_OPTION_MSL_AUX_BUFFER_INDEX = 19 | SPVC_COMPILER_OPTION_MSL_BIT, + SPVC_COMPILER_OPTION_MSL_SWIZZLE_BUFFER_INDEX = 19 | SPVC_COMPILER_OPTION_MSL_BIT, + SPVC_COMPILER_OPTION_MSL_INDIRECT_PARAMS_BUFFER_INDEX = 20 | SPVC_COMPILER_OPTION_MSL_BIT, SPVC_COMPILER_OPTION_MSL_SHADER_OUTPUT_BUFFER_INDEX = 21 | SPVC_COMPILER_OPTION_MSL_BIT, SPVC_COMPILER_OPTION_MSL_SHADER_PATCH_OUTPUT_BUFFER_INDEX = 22 | SPVC_COMPILER_OPTION_MSL_BIT, @@ -426,6 +441,8 @@ typedef enum spvc_compiler_option SPVC_COMPILER_OPTION_MSL_TEXTURE_BUFFER_NATIVE = 34 | SPVC_COMPILER_OPTION_MSL_BIT, + SPVC_COMPILER_OPTION_GLSL_EMIT_UNIFORM_BUFFER_AS_PLAIN_UNIFORMS = 35 | SPVC_COMPILER_OPTION_GLSL_BIT, + SPVC_COMPILER_OPTION_INT_MAX = 0x7fffffff } spvc_compiler_option; @@ -503,7 +520,11 @@ SPVC_PUBLIC_API spvc_variable_id spvc_compiler_hlsl_remap_num_workgroups_builtin * Maps to C++ API. */ SPVC_PUBLIC_API spvc_bool spvc_compiler_msl_is_rasterization_disabled(spvc_compiler compiler); + +/* Obsolete. Renamed to needs_swizzle_buffer. */ SPVC_PUBLIC_API spvc_bool spvc_compiler_msl_needs_aux_buffer(spvc_compiler compiler); +SPVC_PUBLIC_API spvc_bool spvc_compiler_msl_needs_swizzle_buffer(spvc_compiler compiler); + SPVC_PUBLIC_API spvc_bool spvc_compiler_msl_needs_output_buffer(spvc_compiler compiler); SPVC_PUBLIC_API spvc_bool spvc_compiler_msl_needs_patch_output_buffer(spvc_compiler compiler); SPVC_PUBLIC_API spvc_bool spvc_compiler_msl_needs_input_threadgroup_mem(spvc_compiler compiler); @@ -565,6 +586,7 @@ SPVC_PUBLIC_API unsigned spvc_compiler_get_member_decoration(spvc_compiler compi unsigned member_index, SpvDecoration decoration); SPVC_PUBLIC_API const char *spvc_compiler_get_member_decoration_string(spvc_compiler compiler, spvc_type_id id, unsigned member_index, SpvDecoration decoration); +SPVC_PUBLIC_API const char *spvc_compiler_get_member_name(spvc_compiler compiler, spvc_type_id id, unsigned member_index); /* * Entry points. @@ -657,6 +679,15 @@ SPVC_PUBLIC_API spvc_constant_id spvc_compiler_get_work_group_size_specializatio spvc_specialization_constant *y, spvc_specialization_constant *z); +/* + * Buffer ranges + * Maps to C++ API. + */ +SPVC_PUBLIC_API spvc_result spvc_compiler_get_active_buffer_ranges(spvc_compiler compiler, + spvc_variable_id id, + const spvc_buffer_range **ranges, + size_t *num_ranges); + /* * No stdint.h until C99, sigh :( * For smaller types, the result is sign or zero-extended as appropriate. diff --git a/3rdparty/spirv-cross/spirv_glsl.cpp b/3rdparty/spirv-cross/spirv_glsl.cpp index 630b7e040..32582fb62 100644 --- a/3rdparty/spirv-cross/spirv_glsl.cpp +++ b/3rdparty/spirv-cross/spirv_glsl.cpp @@ -1459,9 +1459,19 @@ string CompilerGLSL::layout_for_variable(const SPIRVariable &var) attr.push_back(join("set = ", dec.set)); } + bool push_constant_block = options.vulkan_semantics && var.storage == StorageClassPushConstant; + bool ssbo_block = var.storage == StorageClassStorageBuffer || + (var.storage == StorageClassUniform && typeflags.get(DecorationBufferBlock)); + bool emulated_ubo = var.storage == StorageClassPushConstant && options.emit_push_constant_as_uniform_buffer; + bool ubo_block = var.storage == StorageClassUniform && typeflags.get(DecorationBlock); + // GL 3.0/GLSL 1.30 is not considered legacy, but it doesn't have UBOs ... bool can_use_buffer_blocks = (options.es && options.version >= 300) || (!options.es && options.version >= 140); + // pretend no UBOs when options say so + if (ubo_block && options.emit_uniform_buffer_as_plain_uniforms) + can_use_buffer_blocks = false; + bool can_use_binding; if (options.es) can_use_binding = options.version >= 310; @@ -1478,12 +1488,6 @@ string CompilerGLSL::layout_for_variable(const SPIRVariable &var) if (flags.get(DecorationOffset)) attr.push_back(join("offset = ", dec.offset)); - bool push_constant_block = options.vulkan_semantics && var.storage == StorageClassPushConstant; - bool ssbo_block = var.storage == StorageClassStorageBuffer || - (var.storage == StorageClassUniform && typeflags.get(DecorationBufferBlock)); - bool emulated_ubo = var.storage == StorageClassPushConstant && options.emit_push_constant_as_uniform_buffer; - bool ubo_block = var.storage == StorageClassUniform && typeflags.get(DecorationBlock); - // Instead of adding explicit offsets for every element here, just assume we're using std140 or std430. // If SPIR-V does not comply with either layout, we cannot really work around it. if (can_use_buffer_blocks && (ubo_block || emulated_ubo)) @@ -1611,9 +1615,13 @@ void CompilerGLSL::emit_push_constant_block_glsl(const SPIRVariable &var) void CompilerGLSL::emit_buffer_block(const SPIRVariable &var) { + auto &type = get(var.basetype); + bool ubo_block = var.storage == StorageClassUniform && has_decoration(type.self, DecorationBlock); + if (flattened_buffer_blocks.count(var.self)) emit_buffer_block_flattened(var); - else if (is_legacy() || (!options.es && options.version == 130)) + else if (is_legacy() || (!options.es && options.version == 130) || + (ubo_block && options.emit_uniform_buffer_as_plain_uniforms)) emit_buffer_block_legacy(var); else emit_buffer_block_native(var); @@ -5772,6 +5780,10 @@ case OpGroupNonUniform##op: \ string CompilerGLSL::bitcast_glsl_op(const SPIRType &out_type, const SPIRType &in_type) { + // OpBitcast can deal with pointers. + if (out_type.pointer || in_type.pointer) + return type_to_glsl(out_type); + if (out_type.basetype == in_type.basetype) return ""; @@ -7566,7 +7578,8 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) uint32_t result_type = ops[0]; uint32_t id = ops[1]; auto e = access_chain_internal(ops[2], &ops[3], length - 3, ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, nullptr); - set(id, e + ".length()", result_type, true); + set(id, join(type_to_glsl(get(result_type)), "(", e, ".length())"), result_type, + true); break; } @@ -9913,7 +9926,7 @@ string CompilerGLSL::to_array_size(const SPIRType &type, uint32_t index) return to_expression(size); else if (size) return convert_to_string(size); - else if (!backend.flexible_member_array_supported) + else if (!backend.unsized_array_supported) { // For runtime-sized arrays, we can work around // lack of standard support for this by simply having @@ -11832,6 +11845,9 @@ void CompilerGLSL::bitcast_to_builtin_store(uint32_t target_id, std::string &exp void CompilerGLSL::convert_non_uniform_expression(const SPIRType &type, std::string &expr) { + if (*backend.nonuniform_qualifier == '\0') + return; + // Handle SPV_EXT_descriptor_indexing. if (type.basetype == SPIRType::Sampler || type.basetype == SPIRType::SampledImage || type.basetype == SPIRType::Image) diff --git a/3rdparty/spirv-cross/spirv_glsl.hpp b/3rdparty/spirv-cross/spirv_glsl.hpp index 79fd0444f..184bbbd81 100644 --- a/3rdparty/spirv-cross/spirv_glsl.hpp +++ b/3rdparty/spirv-cross/spirv_glsl.hpp @@ -99,6 +99,10 @@ public: // In non-Vulkan GLSL, emit push constant blocks as UBOs rather than plain uniforms. bool emit_push_constant_as_uniform_buffer = false; + // Always emit uniform blocks as plain uniforms, regardless of the GLSL version, even when UBOs are supported. + // Does not apply to shader storage or push constant blocks. + bool emit_uniform_buffer_as_plain_uniforms = false; + enum Precision { DontCare, @@ -375,7 +379,7 @@ protected: const char *nonuniform_qualifier = "nonuniformEXT"; bool swizzle_is_function = false; bool shared_is_implied = false; - bool flexible_member_array_supported = true; + bool unsized_array_supported = true; bool explicit_struct_type = false; bool use_initializer_list = false; bool use_typed_initializer_list = false; diff --git a/3rdparty/spirv-cross/spirv_hlsl.cpp b/3rdparty/spirv-cross/spirv_hlsl.cpp index 0cd2fb97a..46613c50f 100644 --- a/3rdparty/spirv-cross/spirv_hlsl.cpp +++ b/3rdparty/spirv-cross/spirv_hlsl.cpp @@ -4508,6 +4508,25 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction) HLSL_UFOP(reversebits); break; + case OpArrayLength: + { + auto *var = maybe_get(ops[2]); + if (!var) + SPIRV_CROSS_THROW("Array length must point directly to an SSBO block."); + + auto &type = get(var->basetype); + if (!has_decoration(type.self, DecorationBlock) && !has_decoration(type.self, DecorationBufferBlock)) + SPIRV_CROSS_THROW("Array length expression must point to a block type."); + + // This must be 32-bit uint, so we're good to go. + emit_uninitialized_temporary_expression(ops[0], ops[1]); + statement(to_expression(ops[2]), ".GetDimensions(", to_expression(ops[1]), ");"); + uint32_t offset = type_struct_member_offset(type, ops[3]); + uint32_t stride = type_struct_member_array_stride(type, ops[3]); + statement(to_expression(ops[1]), " = (", to_expression(ops[1]), " - ", offset, ") / ", stride, ";"); + break; + } + default: CompilerGLSL::emit_instruction(instruction); break; @@ -4628,6 +4647,28 @@ uint32_t CompilerHLSL::remap_num_workgroups_builtin() return variable_id; } +void CompilerHLSL::validate_shader_model() +{ + // Check for nonuniform qualifier. + // Instead of looping over all decorations to find this, just look at capabilities. + for (auto &cap : ir.declared_capabilities) + { + switch (cap) + { + case CapabilityShaderNonUniformEXT: + case CapabilityRuntimeDescriptorArrayEXT: + if (hlsl_options.shader_model < 51) + SPIRV_CROSS_THROW( + "Shader model 5.1 or higher is required to use bindless resources or NonUniformResourceIndex."); + default: + break; + } + } + + if (ir.addressing_model != AddressingModelLogical) + SPIRV_CROSS_THROW("Only Logical addressing model can be used with HLSL."); +} + string CompilerHLSL::compile() { // Do not deal with ES-isms like precision, older extensions and such. @@ -4644,7 +4685,7 @@ string CompilerHLSL::compile() backend.basic_uint_type = "uint"; backend.swizzle_is_function = false; backend.shared_is_implied = true; - backend.flexible_member_array_supported = false; + backend.unsized_array_supported = true; backend.explicit_struct_type = false; backend.use_initializer_list = true; backend.use_constructor_splatting = false; @@ -4653,8 +4694,10 @@ string CompilerHLSL::compile() backend.can_declare_struct_inline = false; backend.can_declare_arrays_inline = false; backend.can_return_array = false; + backend.nonuniform_qualifier = "NonUniformResourceIndex"; build_function_control_flow_graphs_and_analyze(); + validate_shader_model(); update_active_builtins(); analyze_image_and_sampler_usage(); diff --git a/3rdparty/spirv-cross/spirv_hlsl.hpp b/3rdparty/spirv-cross/spirv_hlsl.hpp index 50bade21d..d96c911f8 100644 --- a/3rdparty/spirv-cross/spirv_hlsl.hpp +++ b/3rdparty/spirv-cross/spirv_hlsl.hpp @@ -220,6 +220,8 @@ private: // Custom root constant layout, which should be emitted // when translating push constant ranges. std::vector root_constants_layout; + + void validate_shader_model(); }; } // namespace SPIRV_CROSS_NAMESPACE diff --git a/3rdparty/spirv-cross/spirv_msl.cpp b/3rdparty/spirv-cross/spirv_msl.cpp index aaa65d4b5..1d57cceca 100644 --- a/3rdparty/spirv-cross/spirv_msl.cpp +++ b/3rdparty/spirv-cross/spirv_msl.cpp @@ -28,8 +28,6 @@ using namespace std; static const uint32_t k_unknown_location = ~0u; static const uint32_t k_unknown_component = ~0u; -static const uint32_t k_aux_mbr_idx_swizzle_const = 0u; - CompilerMSL::CompilerMSL(std::vector spirv_) : CompilerGLSL(move(spirv_)) { @@ -93,7 +91,14 @@ void CompilerMSL::build_implicit_builtins() bool need_sample_pos = active_input_builtins.get(BuiltInSamplePosition); bool need_vertex_params = capture_output_to_buffer && get_execution_model() == ExecutionModelVertex; bool need_tesc_params = get_execution_model() == ExecutionModelTessellationControl; - if (need_subpass_input || need_sample_pos || need_vertex_params || need_tesc_params) + bool need_subgroup_mask = + active_input_builtins.get(BuiltInSubgroupEqMask) || active_input_builtins.get(BuiltInSubgroupGeMask) || + active_input_builtins.get(BuiltInSubgroupGtMask) || active_input_builtins.get(BuiltInSubgroupLeMask) || + active_input_builtins.get(BuiltInSubgroupLtMask); + bool need_subgroup_ge_mask = !msl_options.is_ios() && (active_input_builtins.get(BuiltInSubgroupGeMask) || + active_input_builtins.get(BuiltInSubgroupGtMask)); + if (need_subpass_input || need_sample_pos || need_subgroup_mask || need_vertex_params || need_tesc_params || + needs_subgroup_invocation_id) { bool has_frag_coord = false; bool has_sample_id = false; @@ -103,18 +108,21 @@ void CompilerMSL::build_implicit_builtins() bool has_base_instance = false; bool has_invocation_id = false; bool has_primitive_id = false; + bool has_subgroup_invocation_id = false; + bool has_subgroup_size = false; ir.for_each_typed_id([&](uint32_t, SPIRVariable &var) { if (var.storage != StorageClassInput || !ir.meta[var.self].decoration.builtin) return; - if (need_subpass_input && ir.meta[var.self].decoration.builtin_type == BuiltInFragCoord) + BuiltIn builtin = ir.meta[var.self].decoration.builtin_type; + if (need_subpass_input && builtin == BuiltInFragCoord) { builtin_frag_coord_id = var.self; has_frag_coord = true; } - if (need_sample_pos && ir.meta[var.self].decoration.builtin_type == BuiltInSampleId) + if (need_sample_pos && builtin == BuiltInSampleId) { builtin_sample_id_id = var.self; has_sample_id = true; @@ -122,7 +130,7 @@ void CompilerMSL::build_implicit_builtins() if (need_vertex_params) { - switch (ir.meta[var.self].decoration.builtin_type) + switch (builtin) { case BuiltInVertexIndex: builtin_vertex_idx_id = var.self; @@ -147,7 +155,7 @@ void CompilerMSL::build_implicit_builtins() if (need_tesc_params) { - switch (ir.meta[var.self].decoration.builtin_type) + switch (builtin) { case BuiltInInvocationId: builtin_invocation_id_id = var.self; @@ -161,6 +169,18 @@ void CompilerMSL::build_implicit_builtins() break; } } + + if ((need_subgroup_mask || needs_subgroup_invocation_id) && builtin == BuiltInSubgroupLocalInvocationId) + { + builtin_subgroup_invocation_id_id = var.self; + has_subgroup_invocation_id = true; + } + + if (need_subgroup_ge_mask && builtin == BuiltInSubgroupSize) + { + builtin_subgroup_size_id = var.self; + has_subgroup_size = true; + } }); if (!has_frag_coord && need_subpass_input) @@ -311,16 +331,67 @@ void CompilerMSL::build_implicit_builtins() builtin_primitive_id_id = var_id; } } + + if (!has_subgroup_invocation_id && (need_subgroup_mask || needs_subgroup_invocation_id)) + { + uint32_t offset = ir.increase_bound_by(3); + uint32_t type_id = offset; + uint32_t type_ptr_id = offset + 1; + uint32_t var_id = offset + 2; + + // Create gl_SubgroupInvocationID. + SPIRType uint_type; + uint_type.basetype = SPIRType::UInt; + uint_type.width = 32; + set(type_id, uint_type); + + SPIRType uint_type_ptr; + uint_type_ptr = uint_type; + uint_type_ptr.pointer = true; + uint_type_ptr.parent_type = type_id; + uint_type_ptr.storage = StorageClassInput; + auto &ptr_type = set(type_ptr_id, uint_type_ptr); + ptr_type.self = type_id; + + set(var_id, type_ptr_id, StorageClassInput); + set_decoration(var_id, DecorationBuiltIn, BuiltInSubgroupLocalInvocationId); + builtin_subgroup_invocation_id_id = var_id; + } + + if (!has_subgroup_size && need_subgroup_ge_mask) + { + uint32_t offset = ir.increase_bound_by(3); + uint32_t type_id = offset; + uint32_t type_ptr_id = offset + 1; + uint32_t var_id = offset + 2; + + // Create gl_SubgroupSize. + SPIRType uint_type; + uint_type.basetype = SPIRType::UInt; + uint_type.width = 32; + set(type_id, uint_type); + + SPIRType uint_type_ptr; + uint_type_ptr = uint_type; + uint_type_ptr.pointer = true; + uint_type_ptr.parent_type = type_id; + uint_type_ptr.storage = StorageClassInput; + auto &ptr_type = set(type_ptr_id, uint_type_ptr); + ptr_type.self = type_id; + + set(var_id, type_ptr_id, StorageClassInput); + set_decoration(var_id, DecorationBuiltIn, BuiltInSubgroupSize); + builtin_subgroup_size_id = var_id; + } } - if (needs_aux_buffer_def) + if (needs_swizzle_buffer_def) { - uint32_t offset = ir.increase_bound_by(5); + uint32_t offset = ir.increase_bound_by(4); uint32_t type_id = offset; - uint32_t type_arr_id = offset + 1; - uint32_t struct_id = offset + 2; - uint32_t struct_ptr_id = offset + 3; - uint32_t var_id = offset + 4; + uint32_t type_ptr_id = offset + 1; + uint32_t type_ptr_ptr_id = offset + 2; + uint32_t var_id = offset + 3; // Create a buffer to hold extra data, including the swizzle constants. SPIRType uint_type; @@ -328,36 +399,25 @@ void CompilerMSL::build_implicit_builtins() uint_type.width = 32; set(type_id, uint_type); - SPIRType uint_type_arr = uint_type; - uint_type_arr.array.push_back(0); - uint_type_arr.array_size_literal.push_back(true); - uint_type_arr.parent_type = type_id; - set(type_arr_id, uint_type_arr); - set_decoration(type_arr_id, DecorationArrayStride, 4); + SPIRType uint_type_pointer = uint_type; + uint_type_pointer.pointer = true; + uint_type_pointer.pointer_depth = 1; + uint_type_pointer.parent_type = type_id; + uint_type_pointer.storage = StorageClassUniform; + set(type_ptr_id, uint_type_pointer); + set_decoration(type_ptr_id, DecorationArrayStride, 4); - SPIRType struct_type; - struct_type.basetype = SPIRType::Struct; - struct_type.member_types.push_back(type_arr_id); - auto &type = set(struct_id, struct_type); - type.self = struct_id; - set_decoration(struct_id, DecorationBlock); - set_name(struct_id, "spvAux"); - set_member_name(struct_id, k_aux_mbr_idx_swizzle_const, "swizzleConst"); - set_member_decoration(struct_id, k_aux_mbr_idx_swizzle_const, DecorationOffset, 0); + SPIRType uint_type_pointer2 = uint_type_pointer; + uint_type_pointer2.pointer_depth++; + uint_type_pointer2.parent_type = type_ptr_id; + set(type_ptr_ptr_id, uint_type_pointer2); - SPIRType struct_type_ptr = struct_type; - struct_type_ptr.pointer = true; - struct_type_ptr.parent_type = struct_id; - struct_type_ptr.storage = StorageClassUniform; - auto &ptr_type = set(struct_ptr_id, struct_type_ptr); - ptr_type.self = struct_id; - - set(var_id, struct_ptr_id, StorageClassUniform); - set_name(var_id, "spvAuxBuffer"); + set(var_id, type_ptr_ptr_id, StorageClassUniformConstant); + set_name(var_id, "spvSwizzleConstants"); // This should never match anything. set_decoration(var_id, DecorationDescriptorSet, 0xFFFFFFFE); - set_decoration(var_id, DecorationBinding, msl_options.aux_buffer_index); - aux_buffer_id = var_id; + set_decoration(var_id, DecorationBinding, msl_options.swizzle_buffer_index); + swizzle_buffer_id = var_id; } } @@ -577,7 +637,7 @@ string CompilerMSL::compile() backend.use_initializer_list = true; backend.use_typed_initializer_list = true; backend.native_row_major_matrix = false; - backend.flexible_member_array_supported = false; + backend.unsized_array_supported = false; backend.can_declare_arrays_inline = false; backend.can_return_array = false; backend.boolean_mix_support = false; @@ -585,6 +645,7 @@ string CompilerMSL::compile() backend.array_is_value_type = false; backend.comparison_image_samples_scalar = true; backend.native_pointers = true; + backend.nonuniform_qualifier = ""; capture_output_to_buffer = msl_options.capture_output_to_buffer; is_rasterization_disabled = msl_options.disable_rasterization || capture_output_to_buffer; @@ -597,16 +658,14 @@ string CompilerMSL::compile() update_active_builtins(); analyze_image_and_sampler_usage(); analyze_sampled_image_usage(); + preprocess_op_codes(); build_implicit_builtins(); fixup_image_load_store_access(); set_enabled_interface_variables(get_active_interface_variables()); - if (aux_buffer_id) - active_interface_variables.insert(aux_buffer_id); - - // Preprocess OpCodes to extract the need to output additional header content - preprocess_op_codes(); + if (swizzle_buffer_id) + active_interface_variables.insert(swizzle_buffer_id); // Create structs to hold input, output and uniform variables. // Do output first to ensure out. is declared at top of entry function. @@ -699,6 +758,9 @@ void CompilerMSL::preprocess_op_codes() is_rasterization_disabled = true; capture_output_to_buffer = true; } + + if (preproc.needs_subgroup_invocation_id) + needs_subgroup_invocation_id = true; } // Move the Private and Workgroup global variables to the entry function. @@ -2876,6 +2938,90 @@ void CompilerMSL::emit_custom_functions() statement("return t.gather_compare(s, spvForward(params)...);"); end_scope(); statement(""); + break; + + case SPVFuncImplSubgroupBallot: + statement("inline uint4 spvSubgroupBallot(bool value)"); + begin_scope(); + statement("simd_vote vote = simd_ballot(value);"); + statement("// simd_ballot() returns a 64-bit integer-like object, but"); + statement("// SPIR-V callers expect a uint4. We must convert."); + statement("// FIXME: This won't include higher bits if Apple ever supports"); + statement("// 128 lanes in an SIMD-group."); + statement("return uint4((uint)((simd_vote::vote_t)vote & 0xFFFFFFFF), (uint)(((simd_vote::vote_t)vote >> " + "32) & 0xFFFFFFFF), 0, 0);"); + end_scope(); + statement(""); + break; + + case SPVFuncImplSubgroupBallotBitExtract: + statement("inline bool spvSubgroupBallotBitExtract(uint4 ballot, uint bit)"); + begin_scope(); + statement("return !!extract_bits(ballot[bit / 32], bit % 32, 1);"); + end_scope(); + statement(""); + break; + + case SPVFuncImplSubgroupBallotFindLSB: + statement("inline uint spvSubgroupBallotFindLSB(uint4 ballot)"); + begin_scope(); + statement("return select(ctz(ballot.x), select(32 + ctz(ballot.y), select(64 + ctz(ballot.z), select(96 + " + "ctz(ballot.w), uint(-1), ballot.w == 0), ballot.z == 0), ballot.y == 0), ballot.x == 0);"); + end_scope(); + statement(""); + break; + + case SPVFuncImplSubgroupBallotFindMSB: + statement("inline uint spvSubgroupBallotFindMSB(uint4 ballot)"); + begin_scope(); + statement("return select(128 - (clz(ballot.w) + 1), select(96 - (clz(ballot.z) + 1), select(64 - " + "(clz(ballot.y) + 1), select(32 - (clz(ballot.x) + 1), uint(-1), ballot.x == 0), ballot.y == 0), " + "ballot.z == 0), ballot.w == 0);"); + end_scope(); + statement(""); + break; + + case SPVFuncImplSubgroupBallotBitCount: + statement("inline uint spvSubgroupBallotBitCount(uint4 ballot)"); + begin_scope(); + statement("return popcount(ballot.x) + popcount(ballot.y) + popcount(ballot.z) + popcount(ballot.w);"); + end_scope(); + statement(""); + statement("inline uint spvSubgroupBallotInclusiveBitCount(uint4 ballot, uint gl_SubgroupInvocationID)"); + begin_scope(); + statement("uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID + 1, 32u)), " + "extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID + 1 - 32, 0)), " + "uint2(0));"); + statement("return spvSubgroupBallotBitCount(ballot & mask);"); + end_scope(); + statement(""); + statement("inline uint spvSubgroupBallotExclusiveBitCount(uint4 ballot, uint gl_SubgroupInvocationID)"); + begin_scope(); + statement("uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID, 32u)), " + "extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID - 32, 0)), uint2(0));"); + statement("return spvSubgroupBallotBitCount(ballot & mask);"); + end_scope(); + statement(""); + break; + + case SPVFuncImplSubgroupAllEqual: + // Metal doesn't provide a function to evaluate this directly. But, we can + // implement this by comparing every thread's value to one thread's value + // (in this case, the value of the first active thread). Then, by the transitive + // property of equality, if all comparisons return true, then they are all equal. + statement("template"); + statement("inline bool spvSubgroupAllEqual(T value)"); + begin_scope(); + statement("return simd_all(value == simd_broadcast_first(value));"); + end_scope(); + statement(""); + statement("template<>"); + statement("inline bool spvSubgroupAllEqual(bool value)"); + begin_scope(); + statement("return simd_all(value) || !simd_any(value);"); + end_scope(); + statement(""); + break; default: break; @@ -3894,33 +4040,70 @@ void CompilerMSL::emit_barrier(uint32_t id_exe_scope, uint32_t id_mem_scope, uin if (get_execution_model() != ExecutionModelGLCompute && get_execution_model() != ExecutionModelTessellationControl) return; - string bar_stmt = "threadgroup_barrier(mem_flags::"; + uint32_t exe_scope = id_exe_scope ? get(id_exe_scope).scalar() : uint32_t(ScopeInvocation); + uint32_t mem_scope = id_mem_scope ? get(id_mem_scope).scalar() : uint32_t(ScopeInvocation); + // Use the wider of the two scopes (smaller value) + exe_scope = min(exe_scope, mem_scope); + + string bar_stmt; + if ((msl_options.is_ios() && msl_options.supports_msl_version(1, 2)) || msl_options.supports_msl_version(2)) + bar_stmt = exe_scope < ScopeSubgroup ? "threadgroup_barrier" : "simdgroup_barrier"; + else + bar_stmt = "threadgroup_barrier"; + bar_stmt += "("; uint32_t mem_sem = id_mem_sem ? get(id_mem_sem).scalar() : uint32_t(MemorySemanticsMaskNone); - if (get_execution_model() == ExecutionModelTessellationControl) + // Use the | operator to combine flags if we can. + if (msl_options.supports_msl_version(1, 2)) + { + string mem_flags = ""; // For tesc shaders, this also affects objects in the Output storage class. // Since in Metal, these are placed in a device buffer, we have to sync device memory here. - bar_stmt += "mem_device"; - else if (mem_sem & MemorySemanticsCrossWorkgroupMemoryMask) - bar_stmt += "mem_device"; - else if (mem_sem & (MemorySemanticsSubgroupMemoryMask | MemorySemanticsWorkgroupMemoryMask | - MemorySemanticsAtomicCounterMemoryMask)) - bar_stmt += "mem_threadgroup"; - else if (mem_sem & MemorySemanticsImageMemoryMask) - bar_stmt += "mem_texture"; + if (get_execution_model() == ExecutionModelTessellationControl || + (mem_sem & (MemorySemanticsUniformMemoryMask | MemorySemanticsCrossWorkgroupMemoryMask))) + mem_flags += "mem_flags::mem_device"; + if (mem_sem & (MemorySemanticsSubgroupMemoryMask | MemorySemanticsWorkgroupMemoryMask | + MemorySemanticsAtomicCounterMemoryMask)) + { + if (!mem_flags.empty()) + mem_flags += " | "; + mem_flags += "mem_flags::mem_threadgroup"; + } + if (mem_sem & MemorySemanticsImageMemoryMask) + { + if (!mem_flags.empty()) + mem_flags += " | "; + mem_flags += "mem_flags::mem_texture"; + } + + if (mem_flags.empty()) + mem_flags = "mem_flags::mem_none"; + + bar_stmt += mem_flags; + } else - bar_stmt += "mem_none"; + { + if ((mem_sem & (MemorySemanticsUniformMemoryMask | MemorySemanticsCrossWorkgroupMemoryMask)) && + (mem_sem & (MemorySemanticsSubgroupMemoryMask | MemorySemanticsWorkgroupMemoryMask | + MemorySemanticsAtomicCounterMemoryMask))) + bar_stmt += "mem_flags::mem_device_and_threadgroup"; + else if (mem_sem & (MemorySemanticsUniformMemoryMask | MemorySemanticsCrossWorkgroupMemoryMask)) + bar_stmt += "mem_flags::mem_device"; + else if (mem_sem & (MemorySemanticsSubgroupMemoryMask | MemorySemanticsWorkgroupMemoryMask | + MemorySemanticsAtomicCounterMemoryMask)) + bar_stmt += "mem_flags::mem_threadgroup"; + else if (mem_sem & MemorySemanticsImageMemoryMask) + bar_stmt += "mem_flags::mem_texture"; + else + bar_stmt += "mem_flags::mem_none"; + } if (msl_options.is_ios() && (msl_options.supports_msl_version(2) && !msl_options.supports_msl_version(2, 1))) { bar_stmt += ", "; - // Use the wider of the two scopes (smaller value) - uint32_t exe_scope = id_exe_scope ? get(id_exe_scope).scalar() : uint32_t(ScopeInvocation); - uint32_t mem_scope = id_mem_scope ? get(id_mem_scope).scalar() : uint32_t(ScopeInvocation); - uint32_t scope = min(exe_scope, mem_scope); - switch (scope) + switch (mem_scope) { case ScopeCrossDevice: case ScopeDevice: @@ -4373,7 +4556,10 @@ void CompilerMSL::emit_function_prototype(SPIRFunction &func, const Bitset &) // Manufacture automatic swizzle arg. if (msl_options.swizzle_texture_samples && has_sampled_images && is_sampled_image_type(arg_type)) - decl += join(", constant uint32_t& ", to_swizzle_expression(arg.id)); + { + bool arg_is_array = !arg_type.array.empty(); + decl += join(", constant uint32_t", arg_is_array ? "* " : "& ", to_swizzle_expression(arg.id)); + } if (&arg != &func.arguments.back()) decl += ", "; @@ -4760,7 +4946,7 @@ string CompilerMSL::to_function_args(uint32_t img, const SPIRType &imgtype, bool if (!is_gather) farg_str += ")"; farg_str += ", " + to_swizzle_expression(img); - used_aux_buffer = true; + used_swizzle_buffer = true; } *p_forward = forward; @@ -4851,8 +5037,17 @@ string CompilerMSL::to_func_call_arg(uint32_t id) 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); + { + // Need to check the base variable in case we need to apply a qualified alias. + uint32_t var_id = 0; + auto *sampler_var = maybe_get(id); + if (sampler_var) + var_id = sampler_var->basevariable; + + arg_str += ", " + to_swizzle_expression(var_id ? var_id : id); + } return arg_str; } @@ -4883,9 +5078,15 @@ string CompilerMSL::to_sampler_expression(uint32_t id) string CompilerMSL::to_swizzle_expression(uint32_t id) { auto *combined = maybe_get(id); + auto expr = to_expression(combined ? combined->image : id); auto index = expr.find_first_of('['); + // If an image is part of an argument buffer translate this to a legal identifier. + for (auto &c : expr) + if (c == '.') + c = '_'; + if (index == string::npos) return expr + swizzle_name_suffix; else @@ -5184,6 +5385,8 @@ string CompilerMSL::member_attribute_qualifier(const SPIRType &type, uint32_t in { case BuiltInInvocationId: case BuiltInPrimitiveId: + case BuiltInSubgroupLocalInvocationId: // FIXME: Should work in any stage + case BuiltInSubgroupSize: // FIXME: Should work in any stage return string(" [[") + builtin_qualifier(builtin) + "]]" + (mbr_type.array.empty() ? "" : " "); case BuiltInPatchVertices: return ""; @@ -5343,6 +5546,10 @@ string CompilerMSL::member_attribute_qualifier(const SPIRType &type, uint32_t in case BuiltInNumWorkgroups: case BuiltInLocalInvocationId: case BuiltInLocalInvocationIndex: + case BuiltInNumSubgroups: + case BuiltInSubgroupId: + case BuiltInSubgroupLocalInvocationId: // FIXME: Should work in any stage + case BuiltInSubgroupSize: // FIXME: Should work in any stage return string(" [[") + builtin_qualifier(builtin) + "]]"; default: @@ -5532,7 +5739,8 @@ string CompilerMSL::get_type_address_space(const SPIRType &type, uint32_t id) else return "constant"; } - break; + else + return "constant"; case StorageClassFunction: case StorageClassGeneric: @@ -5589,7 +5797,9 @@ void CompilerMSL::entry_point_args_builtin(string &ep_args) if (bi_type != BuiltInSamplePosition && bi_type != BuiltInHelperInvocation && bi_type != BuiltInPatchVertices && bi_type != BuiltInTessLevelInner && bi_type != BuiltInTessLevelOuter && bi_type != BuiltInPosition && bi_type != BuiltInPointSize && - bi_type != BuiltInClipDistance && bi_type != BuiltInCullDistance) + bi_type != BuiltInClipDistance && bi_type != BuiltInCullDistance && bi_type != BuiltInSubgroupEqMask && + bi_type != BuiltInSubgroupGeMask && bi_type != BuiltInSubgroupGtMask && + bi_type != BuiltInSubgroupLeMask && bi_type != BuiltInSubgroupLtMask) { if (!ep_args.empty()) ep_args += ", "; @@ -5812,7 +6022,10 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args) ep_args += " [[texture(" + convert_to_string(r.index) + ")]]"; break; default: - SPIRV_CROSS_THROW("Unexpected resource type"); + if (!ep_args.empty()) + ep_args += ", "; + ep_args += type_to_glsl(type, var_id) + " " + r.name; + ep_args += " [[buffer(" + convert_to_string(r.index) + ")]]"; break; } } @@ -5847,11 +6060,24 @@ void CompilerMSL::fix_up_shader_inputs_outputs() if (msl_options.swizzle_texture_samples && has_sampled_images && is_sampled_image_type(type)) { auto &entry_func = this->get(ir.default_entry_point); - entry_func.fixup_hooks_in.push_back([this, &var, var_id]() { - auto &aux_type = expression_type(aux_buffer_id); - statement("constant uint32_t& ", to_swizzle_expression(var_id), " = ", to_name(aux_buffer_id), ".", - to_member_name(aux_type, k_aux_mbr_idx_swizzle_const), "[", - convert_to_string(get_metal_resource_index(var, SPIRType::Image)), "];"); + entry_func.fixup_hooks_in.push_back([this, &type, &var, var_id]() { + bool is_array_type = !type.array.empty(); + + uint32_t desc_set = get_decoration(var_id, DecorationDescriptorSet); + if (descriptor_set_is_argument_buffer(desc_set)) + { + statement("constant uint32_t", is_array_type ? "* " : "& ", to_swizzle_expression(var_id), + is_array_type ? " = &" : " = ", to_name(argument_buffer_ids[desc_set]), + ".spvSwizzleConstants", "[", + convert_to_string(get_metal_resource_index(var, SPIRType::Image)), "];"); + } + else + { + // If we have an array of images, we need to be able to index into it, so take a pointer instead. + statement("constant uint32_t", is_array_type ? "* " : "& ", to_swizzle_expression(var_id), + is_array_type ? " = &" : " = ", to_name(swizzle_buffer_id), "[", + convert_to_string(get_metal_resource_index(var, SPIRType::Image)), "];"); + } }); } } @@ -5903,6 +6129,94 @@ void CompilerMSL::fix_up_shader_inputs_outputs() entry_func.fixup_hooks_in.push_back([=]() { statement(tc, ".y = 1.0 - ", tc, ".y;"); }); } break; + case BuiltInSubgroupEqMask: + if (msl_options.is_ios()) + SPIRV_CROSS_THROW("Subgroup ballot functionality is unavailable on iOS."); + if (!msl_options.supports_msl_version(2, 1)) + SPIRV_CROSS_THROW("Subgroup ballot functionality requires Metal 2.1."); + entry_func.fixup_hooks_in.push_back([=]() { + statement(builtin_type_decl(bi_type), " ", to_expression(var_id), " = ", + builtin_subgroup_invocation_id_id, " > 32 ? uint4(0, (1 << (", + to_expression(builtin_subgroup_invocation_id_id), " - 32)), uint2(0)) : uint4(1 << ", + to_expression(builtin_subgroup_invocation_id_id), ", uint3(0));"); + }); + break; + case BuiltInSubgroupGeMask: + if (msl_options.is_ios()) + SPIRV_CROSS_THROW("Subgroup ballot functionality is unavailable on iOS."); + if (!msl_options.supports_msl_version(2, 1)) + SPIRV_CROSS_THROW("Subgroup ballot functionality requires Metal 2.1."); + entry_func.fixup_hooks_in.push_back([=]() { + // Case where index < 32, size < 32: + // mask0 = bfe(0xFFFFFFFF, index, size - index); + // mask1 = bfe(0xFFFFFFFF, 0, 0); // Gives 0 + // Case where index < 32 but size >= 32: + // mask0 = bfe(0xFFFFFFFF, index, 32 - index); + // mask1 = bfe(0xFFFFFFFF, 0, size - 32); + // Case where index >= 32: + // mask0 = bfe(0xFFFFFFFF, 32, 0); // Gives 0 + // mask1 = bfe(0xFFFFFFFF, index - 32, size - index); + // This is expressed without branches to avoid divergent + // control flow--hence the complicated min/max expressions. + // This is further complicated by the fact that if you attempt + // to bfe out-of-bounds on Metal, undefined behavior is the + // result. + statement(builtin_type_decl(bi_type), " ", to_expression(var_id), + " = uint4(extract_bits(0xFFFFFFFF, min(", + to_expression(builtin_subgroup_invocation_id_id), ", 32u), (uint)max(min((int)", + to_expression(builtin_subgroup_size_id), ", 32) - (int)", + to_expression(builtin_subgroup_invocation_id_id), + ", 0)), extract_bits(0xFFFFFFFF, (uint)max((int)", + to_expression(builtin_subgroup_invocation_id_id), " - 32, 0), (uint)max((int)", + to_expression(builtin_subgroup_size_id), " - (int)max(", + to_expression(builtin_subgroup_invocation_id_id), ", 32u), 0)), uint2(0));"); + }); + break; + case BuiltInSubgroupGtMask: + if (msl_options.is_ios()) + SPIRV_CROSS_THROW("Subgroup ballot functionality is unavailable on iOS."); + if (!msl_options.supports_msl_version(2, 1)) + SPIRV_CROSS_THROW("Subgroup ballot functionality requires Metal 2.1."); + entry_func.fixup_hooks_in.push_back([=]() { + // The same logic applies here, except now the index is one + // more than the subgroup invocation ID. + statement(builtin_type_decl(bi_type), " ", to_expression(var_id), + " = uint4(extract_bits(0xFFFFFFFF, min(", + to_expression(builtin_subgroup_invocation_id_id), " + 1, 32u), (uint)max(min((int)", + to_expression(builtin_subgroup_size_id), ", 32) - (int)", + to_expression(builtin_subgroup_invocation_id_id), + " - 1, 0)), extract_bits(0xFFFFFFFF, (uint)max((int)", + to_expression(builtin_subgroup_invocation_id_id), " + 1 - 32, 0), (uint)max((int)", + to_expression(builtin_subgroup_size_id), " - (int)max(", + to_expression(builtin_subgroup_invocation_id_id), " + 1, 32u), 0)), uint2(0));"); + }); + break; + case BuiltInSubgroupLeMask: + if (msl_options.is_ios()) + SPIRV_CROSS_THROW("Subgroup ballot functionality is unavailable on iOS."); + if (!msl_options.supports_msl_version(2, 1)) + SPIRV_CROSS_THROW("Subgroup ballot functionality requires Metal 2.1."); + entry_func.fixup_hooks_in.push_back([=]() { + statement(builtin_type_decl(bi_type), " ", to_expression(var_id), + " = uint4(extract_bits(0xFFFFFFFF, 0, min(", + to_expression(builtin_subgroup_invocation_id_id), + " + 1, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)", + to_expression(builtin_subgroup_invocation_id_id), " + 1 - 32, 0)), uint2(0));"); + }); + break; + case BuiltInSubgroupLtMask: + if (msl_options.is_ios()) + SPIRV_CROSS_THROW("Subgroup ballot functionality is unavailable on iOS."); + if (!msl_options.supports_msl_version(2, 1)) + SPIRV_CROSS_THROW("Subgroup ballot functionality requires Metal 2.1."); + entry_func.fixup_hooks_in.push_back([=]() { + statement(builtin_type_decl(bi_type), " ", to_expression(var_id), + " = uint4(extract_bits(0xFFFFFFFF, 0, min(", + to_expression(builtin_subgroup_invocation_id_id), + ", 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)", + to_expression(builtin_subgroup_invocation_id_id), " - 32, 0)), uint2(0));"); + }); + break; default: break; } @@ -5930,20 +6244,23 @@ uint32_t CompilerMSL::get_metal_resource_index(SPIRVariable &var, SPIRType::Base itr->second = true; switch (basetype) { - case SPIRType::Struct: - return itr->first.msl_buffer; case SPIRType::Image: return itr->first.msl_texture; case SPIRType::Sampler: return itr->first.msl_sampler; default: - return 0; + return itr->first.msl_buffer; } } // If there is no explicit mapping of bindings to MSL, use the declared binding. if (has_decoration(var.self, DecorationBinding)) - return get_decoration(var.self, DecorationBinding); + { + var_binding = get_decoration(var.self, DecorationBinding); + // Avoid emitting sentinel bindings. + if (var_binding < 0x80000000u) + return var_binding; + } uint32_t binding_stride = 1; auto &type = get(var.basetype); @@ -5954,10 +6271,6 @@ uint32_t CompilerMSL::get_metal_resource_index(SPIRVariable &var, SPIRType::Base uint32_t resource_index; switch (basetype) { - case SPIRType::Struct: - resource_index = next_metal_resource_index_buffer; - next_metal_resource_index_buffer += binding_stride; - break; case SPIRType::Image: resource_index = next_metal_resource_index_texture; next_metal_resource_index_texture += binding_stride; @@ -5967,7 +6280,8 @@ uint32_t CompilerMSL::get_metal_resource_index(SPIRVariable &var, SPIRType::Base next_metal_resource_index_sampler += binding_stride; break; default: - resource_index = 0; + resource_index = next_metal_resource_index_buffer; + next_metal_resource_index_buffer += binding_stride; break; } return resource_index; @@ -6258,6 +6572,7 @@ void CompilerMSL::replace_illegal_names() "M_2_SQRTPI", "M_SQRT2", "M_SQRT1_2", + "quad_broadcast", }; static const unordered_set illegal_func_names = { @@ -6740,6 +7055,234 @@ string CompilerMSL::image_type_glsl(const SPIRType &type, uint32_t id) return img_type_name; } +void CompilerMSL::emit_subgroup_op(const Instruction &i) +{ + const uint32_t *ops = stream(i); + auto op = static_cast(i.op); + + // Metal 2.0 is required. iOS only supports quad ops. macOS only supports + // broadcast and shuffle on 10.13 (2.0), with full support in 10.14 (2.1). + // Note that iOS makes no distinction between a quad-group and a subgroup; + // all subgroups are quad-groups there. + if (!msl_options.supports_msl_version(2)) + SPIRV_CROSS_THROW("Subgroups are only supported in Metal 2.0 and up."); + + if (msl_options.is_ios()) + { + switch (op) + { + default: + SPIRV_CROSS_THROW("iOS only supports quad-group operations."); + case OpGroupNonUniformBroadcast: + case OpGroupNonUniformShuffle: + case OpGroupNonUniformShuffleXor: + case OpGroupNonUniformShuffleUp: + case OpGroupNonUniformShuffleDown: + case OpGroupNonUniformQuadSwap: + case OpGroupNonUniformQuadBroadcast: + break; + } + } + + if (msl_options.is_macos() && !msl_options.supports_msl_version(2, 1)) + { + switch (op) + { + default: + SPIRV_CROSS_THROW("Subgroup ops beyond broadcast and shuffle on macOS require Metal 2.0 and up."); + case OpGroupNonUniformBroadcast: + case OpGroupNonUniformShuffle: + case OpGroupNonUniformShuffleXor: + case OpGroupNonUniformShuffleUp: + case OpGroupNonUniformShuffleDown: + break; + } + } + + uint32_t result_type = ops[0]; + uint32_t id = ops[1]; + + auto scope = static_cast(get(ops[2]).scalar()); + if (scope != ScopeSubgroup) + SPIRV_CROSS_THROW("Only subgroup scope is supported."); + + switch (op) + { + case OpGroupNonUniformElect: + emit_op(result_type, id, "simd_is_first()", true); + break; + + case OpGroupNonUniformBroadcast: + emit_binary_func_op(result_type, id, ops[3], ops[4], + msl_options.is_ios() ? "quad_broadcast" : "simd_broadcast"); + break; + + case OpGroupNonUniformBroadcastFirst: + emit_unary_func_op(result_type, id, ops[3], "simd_broadcast_first"); + break; + + case OpGroupNonUniformBallot: + emit_unary_func_op(result_type, id, ops[3], "spvSubgroupBallot"); + break; + + case OpGroupNonUniformInverseBallot: + emit_binary_func_op(result_type, id, ops[3], builtin_subgroup_invocation_id_id, "spvSubgroupBallotBitExtract"); + break; + + case OpGroupNonUniformBallotBitExtract: + emit_binary_func_op(result_type, id, ops[3], ops[4], "spvSubgroupBallotBitExtract"); + break; + + case OpGroupNonUniformBallotFindLSB: + emit_unary_func_op(result_type, id, ops[3], "spvSubgroupBallotFindLSB"); + break; + + case OpGroupNonUniformBallotFindMSB: + emit_unary_func_op(result_type, id, ops[3], "spvSubgroupBallotFindMSB"); + break; + + case OpGroupNonUniformBallotBitCount: + { + auto operation = static_cast(ops[3]); + if (operation == GroupOperationReduce) + emit_unary_func_op(result_type, id, ops[4], "spvSubgroupBallotBitCount"); + else if (operation == GroupOperationInclusiveScan) + emit_binary_func_op(result_type, id, ops[4], builtin_subgroup_invocation_id_id, + "spvSubgroupBallotInclusiveBitCount"); + else if (operation == GroupOperationExclusiveScan) + emit_binary_func_op(result_type, id, ops[4], builtin_subgroup_invocation_id_id, + "spvSubgroupBallotExclusiveBitCount"); + else + SPIRV_CROSS_THROW("Invalid BitCount operation."); + break; + } + + case OpGroupNonUniformShuffle: + emit_binary_func_op(result_type, id, ops[3], ops[4], msl_options.is_ios() ? "quad_shuffle" : "simd_shuffle"); + break; + + case OpGroupNonUniformShuffleXor: + emit_binary_func_op(result_type, id, ops[3], ops[4], + msl_options.is_ios() ? "quad_shuffle_xor" : "simd_shuffle_xor"); + break; + + case OpGroupNonUniformShuffleUp: + emit_binary_func_op(result_type, id, ops[3], ops[4], + msl_options.is_ios() ? "quad_shuffle_up" : "simd_shuffle_up"); + break; + + case OpGroupNonUniformShuffleDown: + emit_binary_func_op(result_type, id, ops[3], ops[4], + msl_options.is_ios() ? "quad_shuffle_down" : "simd_shuffle_down"); + break; + + case OpGroupNonUniformAll: + emit_unary_func_op(result_type, id, ops[3], "simd_all"); + break; + + case OpGroupNonUniformAny: + emit_unary_func_op(result_type, id, ops[3], "simd_any"); + break; + + case OpGroupNonUniformAllEqual: + emit_unary_func_op(result_type, id, ops[3], "spvSubgroupAllEqual"); + break; + + // clang-format off +#define MSL_GROUP_OP(op, msl_op) \ +case OpGroupNonUniform##op: \ + { \ + auto operation = static_cast(ops[3]); \ + if (operation == GroupOperationReduce) \ + emit_unary_func_op(result_type, id, ops[4], "simd_" #msl_op); \ + else if (operation == GroupOperationInclusiveScan) \ + emit_unary_func_op(result_type, id, ops[4], "simd_prefix_inclusive_" #msl_op); \ + else if (operation == GroupOperationExclusiveScan) \ + emit_unary_func_op(result_type, id, ops[4], "simd_prefix_exclusive_" #msl_op); \ + else if (operation == GroupOperationClusteredReduce) \ + { \ + /* Only cluster sizes of 4 are supported. */ \ + uint32_t cluster_size = get(ops[5]).scalar(); \ + if (cluster_size != 4) \ + SPIRV_CROSS_THROW("Metal only supports quad ClusteredReduce."); \ + emit_unary_func_op(result_type, id, ops[4], "quad_" #msl_op); \ + } \ + else \ + SPIRV_CROSS_THROW("Invalid group operation."); \ + break; \ + } + MSL_GROUP_OP(FAdd, sum) + MSL_GROUP_OP(FMul, product) + MSL_GROUP_OP(IAdd, sum) + MSL_GROUP_OP(IMul, product) +#undef MSL_GROUP_OP + // The others, unfortunately, don't support InclusiveScan or ExclusiveScan. +#define MSL_GROUP_OP(op, msl_op) \ +case OpGroupNonUniform##op: \ + { \ + auto operation = static_cast(ops[3]); \ + if (operation == GroupOperationReduce) \ + emit_unary_func_op(result_type, id, ops[4], "simd_" #msl_op); \ + else if (operation == GroupOperationInclusiveScan) \ + SPIRV_CROSS_THROW("Metal doesn't support InclusiveScan for OpGroupNonUniform" #op "."); \ + else if (operation == GroupOperationExclusiveScan) \ + SPIRV_CROSS_THROW("Metal doesn't support ExclusiveScan for OpGroupNonUniform" #op "."); \ + else if (operation == GroupOperationClusteredReduce) \ + { \ + /* Only cluster sizes of 4 are supported. */ \ + uint32_t cluster_size = get(ops[5]).scalar(); \ + if (cluster_size != 4) \ + SPIRV_CROSS_THROW("Metal only supports quad ClusteredReduce."); \ + emit_unary_func_op(result_type, id, ops[4], "quad_" #msl_op); \ + } \ + else \ + SPIRV_CROSS_THROW("Invalid group operation."); \ + break; \ + } + MSL_GROUP_OP(FMin, min) + MSL_GROUP_OP(FMax, max) + MSL_GROUP_OP(SMin, min) + MSL_GROUP_OP(SMax, max) + MSL_GROUP_OP(UMin, min) + MSL_GROUP_OP(UMax, max) + MSL_GROUP_OP(BitwiseAnd, and) + MSL_GROUP_OP(BitwiseOr, or) + MSL_GROUP_OP(BitwiseXor, xor) + MSL_GROUP_OP(LogicalAnd, and) + MSL_GROUP_OP(LogicalOr, or) + MSL_GROUP_OP(LogicalXor, xor) + // clang-format on + + case OpGroupNonUniformQuadSwap: + { + // We can implement this easily based on the following table giving + // the target lane ID from the direction and current lane ID: + // Direction + // | 0 | 1 | 2 | + // ---+---+---+---+ + // L 0 | 1 2 3 + // a 1 | 0 3 2 + // n 2 | 3 0 1 + // e 3 | 2 1 0 + // Notice that target = source ^ (direction + 1). + uint32_t mask = get(ops[4]).scalar() + 1; + uint32_t mask_id = ir.increase_bound_by(1); + set(mask_id, expression_type_id(ops[4]), mask, false); + emit_binary_func_op(result_type, id, ops[3], mask_id, "quad_shuffle_xor"); + break; + } + + case OpGroupNonUniformQuadBroadcast: + emit_binary_func_op(result_type, id, ops[3], ops[4], "quad_broadcast"); + break; + + default: + SPIRV_CROSS_THROW("Invalid opcode for subgroup."); + } + + register_control_dependent_expression(id); +} + string CompilerMSL::bitcast_glsl_op(const SPIRType &out_type, const SPIRType &in_type) { if (out_type.basetype == in_type.basetype) @@ -6946,6 +7489,32 @@ string CompilerMSL::builtin_qualifier(BuiltIn builtin) case BuiltInLocalInvocationIndex: return "thread_index_in_threadgroup"; + case BuiltInSubgroupSize: + return "thread_execution_width"; + + case BuiltInNumSubgroups: + if (!msl_options.supports_msl_version(2)) + SPIRV_CROSS_THROW("Subgroup builtins require Metal 2.0."); + return msl_options.is_ios() ? "quadgroups_per_threadgroup" : "simdgroups_per_threadgroup"; + + case BuiltInSubgroupId: + if (!msl_options.supports_msl_version(2)) + SPIRV_CROSS_THROW("Subgroup builtins require Metal 2.0."); + return msl_options.is_ios() ? "quadgroup_index_in_threadgroup" : "simdgroup_index_in_threadgroup"; + + case BuiltInSubgroupLocalInvocationId: + if (!msl_options.supports_msl_version(2)) + SPIRV_CROSS_THROW("Subgroup builtins require Metal 2.0."); + return msl_options.is_ios() ? "thread_index_in_quadgroup" : "thread_index_in_simdgroup"; + + case BuiltInSubgroupEqMask: + case BuiltInSubgroupGeMask: + case BuiltInSubgroupGtMask: + case BuiltInSubgroupLeMask: + case BuiltInSubgroupLtMask: + // Shouldn't be reached. + SPIRV_CROSS_THROW("Subgroup ballot masks are handled specially in MSL."); + default: return "unsupported-built-in"; } @@ -7034,7 +7603,17 @@ string CompilerMSL::builtin_type_decl(BuiltIn builtin) case BuiltInWorkgroupId: return "uint3"; case BuiltInLocalInvocationIndex: + case BuiltInNumSubgroups: + case BuiltInSubgroupId: + case BuiltInSubgroupSize: + case BuiltInSubgroupLocalInvocationId: return "uint"; + case BuiltInSubgroupEqMask: + case BuiltInSubgroupGeMask: + case BuiltInSubgroupGtMask: + case BuiltInSubgroupLeMask: + case BuiltInSubgroupLtMask: + return "uint4"; case BuiltInHelperInvocation: return "bool"; @@ -7199,7 +7778,7 @@ bool CompilerMSL::SampledImageScanner::handle(spv::Op opcode, const uint32_t *ar case OpImageDrefGather: compiler.has_sampled_images = compiler.has_sampled_images || compiler.is_sampled_image_type(compiler.expression_type(args[2])); - compiler.needs_aux_buffer_def = compiler.needs_aux_buffer_def || compiler.has_sampled_images; + compiler.needs_swizzle_buffer_def = compiler.needs_swizzle_buffer_def || compiler.has_sampled_images; break; default: break; @@ -7259,6 +7838,15 @@ bool CompilerMSL::OpCodePreprocessor::handle(Op opcode, const uint32_t *args, ui uses_atomics = true; break; + case OpGroupNonUniformInverseBallot: + needs_subgroup_invocation_id = true; + break; + + case OpGroupNonUniformBallotBitCount: + if (args[3] != GroupOperationReduce) + needs_subgroup_invocation_id = true; + break; + default: break; } @@ -7425,6 +8013,25 @@ CompilerMSL::SPVFuncImpl CompilerMSL::OpCodePreprocessor::get_spv_func_impl(Op o break; } + case OpGroupNonUniformBallot: + return SPVFuncImplSubgroupBallot; + + case OpGroupNonUniformInverseBallot: + case OpGroupNonUniformBallotBitExtract: + return SPVFuncImplSubgroupBallotBitExtract; + + case OpGroupNonUniformBallotFindLSB: + return SPVFuncImplSubgroupBallotFindLSB; + + case OpGroupNonUniformBallotFindMSB: + return SPVFuncImplSubgroupBallotFindMSB; + + case OpGroupNonUniformBallotBitCount: + return SPVFuncImplSubgroupBallotBitCount; + + case OpGroupNonUniformAllEqual: + return SPVFuncImplSubgroupAllEqual; + default: break; } @@ -7634,6 +8241,8 @@ void CompilerMSL::analyze_argument_buffers() }; SmallVector resources_in_set[kMaxArgumentBuffers]; + bool set_needs_swizzle_buffer[kMaxArgumentBuffers] = {}; + ir.for_each_typed_id([&](uint32_t self, SPIRVariable &var) { if ((var.storage == StorageClassUniform || var.storage == StorageClassUniformConstant || var.storage == StorageClassStorageBuffer) && @@ -7677,9 +8286,54 @@ void CompilerMSL::analyze_argument_buffers() resources_in_set[desc_set].push_back( { &var, to_name(var_id), type.basetype, get_metal_resource_index(var, type.basetype) }); } + + // Check if this descriptor set needs a swizzle buffer. + if (needs_swizzle_buffer_def && is_sampled_image_type(type)) + set_needs_swizzle_buffer[desc_set] = true; } }); + if (needs_swizzle_buffer_def) + { + uint32_t swizzle_buffer_type_id = 0; + + // We might have to add a swizzle buffer resource to the set. + for (uint32_t desc_set = 0; desc_set < kMaxArgumentBuffers; desc_set++) + { + if (!set_needs_swizzle_buffer[desc_set]) + continue; + + if (swizzle_buffer_type_id == 0) + { + uint32_t offset = ir.increase_bound_by(2); + uint32_t type_id = offset; + swizzle_buffer_type_id = offset + 1; + + // Create a buffer to hold extra data, including the swizzle constants. + SPIRType uint_type; + uint_type.basetype = SPIRType::UInt; + uint_type.width = 32; + set(type_id, uint_type); + + SPIRType uint_type_pointer = uint_type; + uint_type_pointer.pointer = true; + uint_type_pointer.pointer_depth = 1; + uint_type_pointer.parent_type = type_id; + uint_type_pointer.storage = StorageClassUniform; + set(swizzle_buffer_type_id, uint_type_pointer); + set_decoration(swizzle_buffer_type_id, DecorationArrayStride, 4); + } + + uint32_t var_id = ir.increase_bound_by(1); + auto &var = set(var_id, swizzle_buffer_type_id, StorageClassUniformConstant); + set_name(var_id, "spvSwizzleConstants"); + set_decoration(var_id, DecorationDescriptorSet, desc_set); + set_decoration(var_id, DecorationBinding, kSwizzleBufferBinding); + resources_in_set[desc_set].push_back( + { &var, to_name(var_id), SPIRType::UInt, get_metal_resource_index(var, SPIRType::UInt) }); + } + } + for (uint32_t desc_set = 0; desc_set < kMaxArgumentBuffers; desc_set++) { auto &resources = resources_in_set[desc_set]; diff --git a/3rdparty/spirv-cross/spirv_msl.hpp b/3rdparty/spirv-cross/spirv_msl.hpp index 8d3a8ad8d..6e7c83c3b 100644 --- a/3rdparty/spirv-cross/spirv_msl.hpp +++ b/3rdparty/spirv-cross/spirv_msl.hpp @@ -152,11 +152,11 @@ 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; +// Special constant used in a MSLResourceBinding binding +// element to indicate the buffer binding for swizzle buffers. +static const uint32_t kSwizzleBufferBinding = ~(1u); -// 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 +static const uint32_t kMaxArgumentBuffers = 8; // Decompiles SPIR-V to Metal Shading Language class CompilerMSL : public CompilerGLSL @@ -174,7 +174,7 @@ public: Platform platform = macOS; uint32_t msl_version = make_msl_version(1, 2); uint32_t texel_buffer_texture_width = 4096; // Width of 2D Metal textures used as 1D texel buffers - uint32_t aux_buffer_index = 30; + uint32_t swizzle_buffer_index = 30; uint32_t indirect_params_buffer_index = 29; uint32_t shader_output_buffer_index = 28; uint32_t shader_patch_output_buffer_index = 27; @@ -243,10 +243,10 @@ public: } // Provide feedback to calling API to allow it to pass an auxiliary - // buffer if the shader needs it. - bool needs_aux_buffer() const + // swizzle buffer if the shader needs it. + bool needs_swizzle_buffer() const { - return used_aux_buffer; + return used_swizzle_buffer; } // Provide feedback to calling API to allow it to pass an output @@ -344,6 +344,12 @@ protected: SPVFuncImplRowMajor4x2, SPVFuncImplRowMajor4x3, SPVFuncImplTextureSwizzle, + SPVFuncImplSubgroupBallot, + SPVFuncImplSubgroupBallotBitExtract, + SPVFuncImplSubgroupBallotFindLSB, + SPVFuncImplSubgroupBallotFindMSB, + SPVFuncImplSubgroupBallotBitCount, + SPVFuncImplSubgroupAllEqual, SPVFuncImplArrayCopyMultidimMax = 6 }; @@ -354,6 +360,7 @@ protected: void emit_header() override; void emit_function_prototype(SPIRFunction &func, const Bitset &return_flags) override; void emit_sampled_image_op(uint32_t result_type, uint32_t result_id, uint32_t image_id, uint32_t samp_id) override; + void emit_subgroup_op(const Instruction &i) override; void emit_fixup() override; std::string to_struct_member(const SPIRType &type, uint32_t member_type_id, uint32_t index, const std::string &qualifier = ""); @@ -477,7 +484,9 @@ protected: uint32_t builtin_base_instance_id = 0; uint32_t builtin_invocation_id_id = 0; uint32_t builtin_primitive_id_id = 0; - uint32_t aux_buffer_id = 0; + uint32_t builtin_subgroup_invocation_id_id = 0; + uint32_t builtin_subgroup_size_id = 0; + uint32_t swizzle_buffer_id = 0; void bitcast_to_builtin_store(uint32_t target_id, std::string &expr, const SPIRType &expr_type) override; void bitcast_from_builtin_load(uint32_t source_id, std::string &expr, const SPIRType &expr_type) override; @@ -515,9 +524,10 @@ protected: bool needs_instance_idx_arg = false; bool is_rasterization_disabled = false; bool capture_output_to_buffer = false; - bool needs_aux_buffer_def = false; - bool used_aux_buffer = false; + bool needs_swizzle_buffer_def = false; + bool used_swizzle_buffer = false; bool added_builtin_tess_level = false; + bool needs_subgroup_invocation_id = false; std::string qual_pos_var_name; std::string stage_in_var_name = "in"; std::string stage_out_var_name = "out"; @@ -561,6 +571,7 @@ protected: bool suppress_missing_prototypes = false; bool uses_atomics = false; bool uses_resource_write = false; + bool needs_subgroup_invocation_id = false; }; // OpcodeHandler that scans for uses of sampled images diff --git a/3rdparty/spirv-cross/spirv_parser.cpp b/3rdparty/spirv-cross/spirv_parser.cpp index ea356a6ac..1c0a830f2 100644 --- a/3rdparty/spirv-cross/spirv_parser.cpp +++ b/3rdparty/spirv-cross/spirv_parser.cpp @@ -59,6 +59,7 @@ static bool is_valid_spirv_version(uint32_t version) case 0x10100: // SPIR-V 1.1 case 0x10200: // SPIR-V 1.2 case 0x10300: // SPIR-V 1.3 + case 0x10400: // SPIR-V 1.4 return true; default: diff --git a/3rdparty/spirv-cross/test_shaders.py b/3rdparty/spirv-cross/test_shaders.py index d6037a8ea..5c9e31e66 100755 --- a/3rdparty/spirv-cross/test_shaders.py +++ b/3rdparty/spirv-cross/test_shaders.py @@ -143,7 +143,7 @@ def cross_compile_msl(shader, spirv, opt, iterations, paths): spirv_path = create_temporary() msl_path = create_temporary(os.path.basename(shader)) - spirv_cmd = [paths.spirv_as, '-o', spirv_path, shader] + spirv_cmd = [paths.spirv_as, '--target-env', 'vulkan1.1', '-o', spirv_path, shader] if '.preserve.' in shader: spirv_cmd.append('--preserve-numeric-ids') @@ -220,13 +220,18 @@ def shader_to_win_path(shader): ignore_fxc = False def validate_shader_hlsl(shader, force_no_external_validation, paths): - subprocess.check_call([paths.glslang, '-e', 'main', '-D', '--target-env', 'vulkan1.1', '-V', shader]) + if not '.nonuniformresource' in shader: + # glslang HLSL does not support this, so rely on fxc to test it. + subprocess.check_call([paths.glslang, '-e', 'main', '-D', '--target-env', 'vulkan1.1', '-V', shader]) is_no_fxc = '.nofxc.' in shader global ignore_fxc if (not ignore_fxc) and (not force_no_external_validation) and (not is_no_fxc): try: win_path = shader_to_win_path(shader) - subprocess.check_call(['fxc', '-nologo', shader_model_hlsl(shader), win_path]) + args = ['fxc', '-nologo', shader_model_hlsl(shader), win_path] + if '.nonuniformresource.' in shader: + args.append('/enable_unbounded_descriptor_tables') + subprocess.check_call(args) except OSError as oe: if (oe.errno != errno.ENOENT): # Ignore not found errors print('Failed to run FXC.') @@ -253,7 +258,7 @@ def cross_compile_hlsl(shader, spirv, opt, force_no_external_validation, iterati spirv_path = create_temporary() hlsl_path = create_temporary(os.path.basename(shader)) - spirv_cmd = [paths.spirv_as, '-o', spirv_path, shader] + spirv_cmd = [paths.spirv_as, '--target-env', 'vulkan1.1', '-o', spirv_path, shader] if '.preserve.' in shader: spirv_cmd.append('--preserve-numeric-ids') @@ -281,7 +286,7 @@ def cross_compile_reflect(shader, spirv, opt, iterations, paths): spirv_path = create_temporary() reflect_path = create_temporary(os.path.basename(shader)) - spirv_cmd = [paths.spirv_as, '-o', spirv_path, shader] + spirv_cmd = [paths.spirv_as, '--target-env', 'vulkan1.1', '-o', spirv_path, shader] if '.preserve.' in shader: spirv_cmd.append('--preserve-numeric-ids') @@ -312,7 +317,7 @@ def cross_compile(shader, vulkan, spirv, invalid_spirv, eliminate, is_legacy, fl if vulkan or spirv: vulkan_glsl_path = create_temporary('vk' + os.path.basename(shader)) - spirv_cmd = [paths.spirv_as, '-o', spirv_path, shader] + spirv_cmd = [paths.spirv_as, '--target-env', 'vulkan1.1', '-o', spirv_path, shader] if '.preserve.' in shader: spirv_cmd.append('--preserve-numeric-ids') diff --git a/3rdparty/spirv-cross/tests-other/c_api_test.c b/3rdparty/spirv-cross/tests-other/c_api_test.c index 2d643c7eb..db396256d 100644 --- a/3rdparty/spirv-cross/tests-other/c_api_test.c +++ b/3rdparty/spirv-cross/tests-other/c_api_test.c @@ -116,12 +116,32 @@ int main(int argc, char **argv) SpvId *buffer = NULL; size_t word_count = 0; - if (argc != 2) + if (argc != 5) return 1; if (read_file(argv[1], &buffer, &word_count) < 0) return 1; + unsigned abi_major, abi_minor, abi_patch; + spvc_get_version(&abi_major, &abi_minor, &abi_patch); + if (abi_major != strtoul(argv[2], NULL, 0)) + { + fprintf(stderr, "VERSION_MAJOR mismatch!\n"); + return 1; + } + + if (abi_minor != strtoul(argv[3], NULL, 0)) + { + fprintf(stderr, "VERSION_MINOR mismatch!\n"); + return 1; + } + + if (abi_patch != strtoul(argv[4], NULL, 0)) + { + fprintf(stderr, "VERSION_PATCH mismatch!\n"); + return 1; + } + SPVC_CHECKED_CALL(spvc_context_create(&context)); spvc_context_set_error_callback(context, error_callback, NULL); SPVC_CHECKED_CALL(spvc_context_parse_spirv(context, buffer, word_count, &ir));