From 97b4fec1ff7a0ac76d0b71d19591270475c55dbf Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=D0=91=D1=80=D0=B0=D0=BD=D0=B8=D0=BC=D0=B8=D1=80=20=D0=9A?= =?UTF-8?q?=D0=B0=D1=80=D0=B0=D1=9F=D0=B8=D1=9B?= Date: Thu, 28 Nov 2019 20:20:59 -0800 Subject: [PATCH] Updated spirv-cross. --- .../comp/image-atomic-automatic-bindings.comp | 25 ++++++++++++++++++ .../asm/frag/texture-atomics.asm.frag | 2 +- ...re-atomics.asm.graphics-robust-access.frag | 2 +- .../comp/image-atomic-automatic-bindings.comp | 26 +++++++++++++++++++ ...c-conversion-signs.asm.nocompat.vk.comp.vk | 4 +++ .../asm/frag/texture-atomics.asm.frag | 2 +- ...re-atomics.asm.graphics-robust-access.frag | 2 +- .../comp/image-atomic-automatic-bindings.comp | 16 ++++++++++++ ...etic-conversion-signs.asm.nocompat.vk.comp | 12 +++++++++ 3rdparty/spirv-cross/spirv_cross.hpp | 8 +++--- 3rdparty/spirv-cross/spirv_glsl.cpp | 9 ++++--- 3rdparty/spirv-cross/spirv_msl.cpp | 11 +++++--- 3rdparty/spirv-cross/test_shaders.py | 14 ++++++++++ 13 files changed, 119 insertions(+), 14 deletions(-) create mode 100644 3rdparty/spirv-cross/reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.comp create mode 100644 3rdparty/spirv-cross/reference/shaders-msl/comp/image-atomic-automatic-bindings.comp create mode 100644 3rdparty/spirv-cross/shaders-msl/comp/image-atomic-automatic-bindings.comp diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.comp b/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.comp new file mode 100644 index 000000000..a8ade54be --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.comp @@ -0,0 +1,25 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" +#pragma clang diagnostic ignored "-Wunused-variable" + +#include +#include +#include + +using namespace metal; + +struct SSBO +{ + float4 outdata; +}; + +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + +// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics +#define spvImage2DAtomicCoord(tc, tex) (((tex).get_width() * (tc).x) + (tc).y) + +kernel void main0(device SSBO& _31 [[buffer(1)]], texture2d uImage [[texture(0)]], device atomic_uint* uImage_atomic [[buffer(0)]], texture2d uTexture [[texture(1)]], sampler uTextureSmplr [[sampler(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + uint _26 = atomic_fetch_add_explicit((device atomic_uint*)&uImage_atomic[spvImage2DAtomicCoord(int2(gl_GlobalInvocationID.xy), uImage)], 10u, memory_order_relaxed); + _31.outdata = uTexture.sample(uTextureSmplr, float2(gl_GlobalInvocationID.xy), level(0.0)) + float4(float(_26)); +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.frag b/3rdparty/spirv-cross/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.frag index a416259d8..98cdda44d 100644 --- a/3rdparty/spirv-cross/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.frag +++ b/3rdparty/spirv-cross/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.frag @@ -68,7 +68,7 @@ struct main0_in uint in_var_TEXCOORD0 [[user(locn0)]]; }; -fragment main0_out main0(main0_in in [[stage_in]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d RWShadowTileNumCulledObjects [[texture(2)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]]) +fragment main0_out main0(main0_in in [[stage_in]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d RWShadowTileNumCulledObjects [[texture(0)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]]) { main0_out out = {}; uint2 _77 = uint2(gl_FragCoord.xy); diff --git a/3rdparty/spirv-cross/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.graphics-robust-access.frag b/3rdparty/spirv-cross/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.graphics-robust-access.frag index 85e27bd1e..795581399 100644 --- a/3rdparty/spirv-cross/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.graphics-robust-access.frag +++ b/3rdparty/spirv-cross/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.graphics-robust-access.frag @@ -68,7 +68,7 @@ struct main0_in uint in_var_TEXCOORD0 [[user(locn0)]]; }; -fragment main0_out main0(main0_in in [[stage_in]], constant uint* spvBufferSizeConstants [[buffer(25)]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d RWShadowTileNumCulledObjects [[texture(2)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]]) +fragment main0_out main0(main0_in in [[stage_in]], constant uint* spvBufferSizeConstants [[buffer(25)]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d RWShadowTileNumCulledObjects [[texture(0)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]]) { main0_out out = {}; constant uint& CulledObjectBoxBoundsBufferSize = spvBufferSizeConstants[0]; diff --git a/3rdparty/spirv-cross/reference/shaders-msl/comp/image-atomic-automatic-bindings.comp b/3rdparty/spirv-cross/reference/shaders-msl/comp/image-atomic-automatic-bindings.comp new file mode 100644 index 000000000..40a8dfcd7 --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-msl/comp/image-atomic-automatic-bindings.comp @@ -0,0 +1,26 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" +#pragma clang diagnostic ignored "-Wunused-variable" + +#include +#include +#include + +using namespace metal; + +struct SSBO +{ + float4 outdata; +}; + +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + +// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics +#define spvImage2DAtomicCoord(tc, tex) (((tex).get_width() * (tc).x) + (tc).y) + +kernel void main0(device SSBO& _31 [[buffer(1)]], texture2d uImage [[texture(0)]], device atomic_uint* uImage_atomic [[buffer(0)]], texture2d uTexture [[texture(1)]], sampler uTextureSmplr [[sampler(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + uint _26 = atomic_fetch_add_explicit((device atomic_uint*)&uImage_atomic[spvImage2DAtomicCoord(int2(gl_GlobalInvocationID.xy), uImage)], 10u, memory_order_relaxed); + uint ret = _26; + _31.outdata = uTexture.sample(uTextureSmplr, float2(gl_GlobalInvocationID.xy), level(0.0)) + float4(float(ret)); +} + diff --git a/3rdparty/spirv-cross/reference/shaders-no-opt/asm/comp/arithmetic-conversion-signs.asm.nocompat.vk.comp.vk b/3rdparty/spirv-cross/reference/shaders-no-opt/asm/comp/arithmetic-conversion-signs.asm.nocompat.vk.comp.vk index c2fb39907..5f480728e 100644 --- a/3rdparty/spirv-cross/reference/shaders-no-opt/asm/comp/arithmetic-conversion-signs.asm.nocompat.vk.comp.vk +++ b/3rdparty/spirv-cross/reference/shaders-no-opt/asm/comp/arithmetic-conversion-signs.asm.nocompat.vk.comp.vk @@ -33,8 +33,12 @@ void main() _4.u16 = uint16_t(_30); _4.f32 = float(_31); _4.f32 = float(int16_t(_32)); + _4.f32 = float(_29); + _4.f32 = float(int(_30)); _4.f32 = float(uint16_t(_31)); _4.f32 = float(_32); + _4.f32 = float(uint(_29)); + _4.f32 = float(_30); _4.s16 = int16_t(_33); _4.u16 = uint16_t(int16_t(_33)); _4.u16 = uint16_t(_33); diff --git a/3rdparty/spirv-cross/reference/shaders-ue4/asm/frag/texture-atomics.asm.frag b/3rdparty/spirv-cross/reference/shaders-ue4/asm/frag/texture-atomics.asm.frag index a416259d8..98cdda44d 100644 --- a/3rdparty/spirv-cross/reference/shaders-ue4/asm/frag/texture-atomics.asm.frag +++ b/3rdparty/spirv-cross/reference/shaders-ue4/asm/frag/texture-atomics.asm.frag @@ -68,7 +68,7 @@ struct main0_in uint in_var_TEXCOORD0 [[user(locn0)]]; }; -fragment main0_out main0(main0_in in [[stage_in]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d RWShadowTileNumCulledObjects [[texture(2)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]]) +fragment main0_out main0(main0_in in [[stage_in]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d RWShadowTileNumCulledObjects [[texture(0)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]]) { main0_out out = {}; uint2 _77 = uint2(gl_FragCoord.xy); diff --git a/3rdparty/spirv-cross/reference/shaders-ue4/asm/frag/texture-atomics.asm.graphics-robust-access.frag b/3rdparty/spirv-cross/reference/shaders-ue4/asm/frag/texture-atomics.asm.graphics-robust-access.frag index a416259d8..98cdda44d 100644 --- a/3rdparty/spirv-cross/reference/shaders-ue4/asm/frag/texture-atomics.asm.graphics-robust-access.frag +++ b/3rdparty/spirv-cross/reference/shaders-ue4/asm/frag/texture-atomics.asm.graphics-robust-access.frag @@ -68,7 +68,7 @@ struct main0_in uint in_var_TEXCOORD0 [[user(locn0)]]; }; -fragment main0_out main0(main0_in in [[stage_in]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d RWShadowTileNumCulledObjects [[texture(2)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]]) +fragment main0_out main0(main0_in in [[stage_in]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d RWShadowTileNumCulledObjects [[texture(0)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]]) { main0_out out = {}; uint2 _77 = uint2(gl_FragCoord.xy); diff --git a/3rdparty/spirv-cross/shaders-msl/comp/image-atomic-automatic-bindings.comp b/3rdparty/spirv-cross/shaders-msl/comp/image-atomic-automatic-bindings.comp new file mode 100644 index 000000000..862cd2129 --- /dev/null +++ b/3rdparty/spirv-cross/shaders-msl/comp/image-atomic-automatic-bindings.comp @@ -0,0 +1,16 @@ +#version 450 +layout(local_size_x = 1) in; + +layout(set = 0, binding = 0, r32ui) uniform uimage2D uImage; +layout(set = 0, binding = 1) uniform sampler2D uTexture; + +layout(set = 0, binding = 2) buffer SSBO +{ + vec4 outdata; +}; + +void main() +{ + uint ret = imageAtomicAdd(uImage, ivec2(gl_GlobalInvocationID.xy), 10u); + outdata = textureLod(uTexture, vec2(gl_GlobalInvocationID.xy), 0.0) + float(ret); +} diff --git a/3rdparty/spirv-cross/shaders-no-opt/asm/comp/arithmetic-conversion-signs.asm.nocompat.vk.comp b/3rdparty/spirv-cross/shaders-no-opt/asm/comp/arithmetic-conversion-signs.asm.nocompat.vk.comp index 0e1ce235d..504a9546c 100644 --- a/3rdparty/spirv-cross/shaders-no-opt/asm/comp/arithmetic-conversion-signs.asm.nocompat.vk.comp +++ b/3rdparty/spirv-cross/shaders-no-opt/asm/comp/arithmetic-conversion-signs.asm.nocompat.vk.comp @@ -111,11 +111,23 @@ OpStore %ptr_f32 %s16_to_f32_signed %u16_to_f32_signed = OpConvertSToF %float %u16 OpStore %ptr_f32 %u16_to_f32_signed + + %s32_to_f32_signed = OpConvertSToF %float %s32 + OpStore %ptr_f32 %s32_to_f32_signed + %u32_to_f32_signed = OpConvertSToF %float %u32 + OpStore %ptr_f32 %u32_to_f32_signed + + ; UToF %s16_to_f32_unsigned = OpConvertUToF %float %s16 OpStore %ptr_f32 %s16_to_f32_unsigned %u16_to_f32_unsigned = OpConvertUToF %float %u16 OpStore %ptr_f32 %u16_to_f32_unsigned + %s32_to_f32_unsigned = OpConvertUToF %float %s32 + OpStore %ptr_f32 %s32_to_f32_unsigned + %u32_to_f32_unsigned = OpConvertUToF %float %u32 + OpStore %ptr_f32 %u32_to_f32_unsigned + ; FToS %f32_to_s16_signed = OpConvertFToS %short %f32 OpStore %ptr_s16 %f32_to_s16_signed diff --git a/3rdparty/spirv-cross/spirv_cross.hpp b/3rdparty/spirv-cross/spirv_cross.hpp index 73d9085c9..92379aef4 100644 --- a/3rdparty/spirv-cross/spirv_cross.hpp +++ b/3rdparty/spirv-cross/spirv_cross.hpp @@ -315,6 +315,10 @@ public: const std::string &get_cleansed_entry_point_name(const std::string &name, spv::ExecutionModel execution_model) const; + // Traverses all reachable opcodes and sets active_builtins to a bitmask of all builtin variables which are accessed in the shader. + void update_active_builtins(); + bool has_active_builtin(spv::BuiltIn builtin, spv::StorageClass storage); + // Query and modify OpExecutionMode. const Bitset &get_execution_mode_bitset() const; @@ -833,10 +837,6 @@ protected: uint32_t cull_distance_count = 0; bool position_invariant = false; - // Traverses all reachable opcodes and sets active_builtins to a bitmask of all builtin variables which are accessed in the shader. - void update_active_builtins(); - bool has_active_builtin(spv::BuiltIn builtin, spv::StorageClass storage); - void analyze_parameter_preservation( SPIRFunction &entry, const CFG &cfg, const std::unordered_map> &variable_to_blocks, diff --git a/3rdparty/spirv-cross/spirv_glsl.cpp b/3rdparty/spirv-cross/spirv_glsl.cpp index e9925165f..5bfec075f 100644 --- a/3rdparty/spirv-cross/spirv_glsl.cpp +++ b/3rdparty/spirv-cross/spirv_glsl.cpp @@ -9104,9 +9104,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) auto &arg_type = expression_type(ops[2]); auto func = type_to_glsl_constructor(type); - // If we're sign-extending or zero-extending, we need to make sure we cast from the correct type. - // For truncation, it does not matter, so don't emit useless casts. - if (arg_type.width < type.width) + if (arg_type.width < type.width || type_is_floating_point(type)) emit_unary_func_op_cast(result_type, id, ops[2], func.c_str(), input_type, type.basetype); else emit_unary_func_op(result_type, id, ops[2], func.c_str()); @@ -12119,7 +12117,12 @@ void CompilerGLSL::emit_block_chain(SPIRBlock &block) SPIRBlock::ContinueBlockType continue_type = SPIRBlock::ContinueNone; if (block.continue_block) + { continue_type = continue_block_type(get(block.continue_block)); + // If we know we cannot emit a loop, mark the block early as a complex loop so we don't force unnecessary recompiles. + if (continue_type == SPIRBlock::ComplexLoop) + block.complex_continue = true; + } // If we have loop variables, stop masking out access to the variable now. for (auto var_id : block.loop_variables) diff --git a/3rdparty/spirv-cross/spirv_msl.cpp b/3rdparty/spirv-cross/spirv_msl.cpp index b16e1e80e..c0db4f2f8 100644 --- a/3rdparty/spirv-cross/spirv_msl.cpp +++ b/3rdparty/spirv-cross/spirv_msl.cpp @@ -9627,9 +9627,14 @@ uint32_t CompilerMSL::get_metal_resource_index(SPIRVariable &var, SPIRType::Base // If a matching binding has been specified, find and use it. auto itr = resource_bindings.find({ execution.model, var_desc_set, var_binding }); - auto resource_decoration = var_type.basetype == SPIRType::SampledImage && basetype == SPIRType::Sampler ? - SPIRVCrossDecorationResourceIndexSecondary : - SPIRVCrossDecorationResourceIndexPrimary; + // Atomic helper buffers for image atomics need to use secondary bindings as well. + bool use_secondary_binding = (var_type.basetype == SPIRType::SampledImage && basetype == SPIRType::Sampler) || + basetype == SPIRType::AtomicCounter; + + auto resource_decoration = use_secondary_binding ? + SPIRVCrossDecorationResourceIndexSecondary : + SPIRVCrossDecorationResourceIndexPrimary; + if (plane == 1) resource_decoration = SPIRVCrossDecorationResourceIndexTertiary; if (plane == 2) diff --git a/3rdparty/spirv-cross/test_shaders.py b/3rdparty/spirv-cross/test_shaders.py index 4b2c4dc57..52a344f5e 100755 --- a/3rdparty/spirv-cross/test_shaders.py +++ b/3rdparty/spirv-cross/test_shaders.py @@ -1,5 +1,19 @@ #!/usr/bin/env python3 + # Copyright 2015-2019 Arm Limited + # + # Licensed under the Apache License, Version 2.0 (the "License"); + # you may not use this file except in compliance with the License. + # You may obtain a copy of the License at + # + # http://www.apache.org/licenses/LICENSE-2.0 + # + # Unless required by applicable law or agreed to in writing, software + # distributed under the License is distributed on an "AS IS" BASIS, + # WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + # See the License for the specific language governing permissions and + # limitations under the License. + import sys import os import os.path