From 38f97c28cf26024cda2c4919ca1492b323431080 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=D0=91=D1=80=D0=B0=D0=BD=D0=B8=D0=BC=D0=B8=D1=80=20=D0=9A?= =?UTF-8?q?=D0=B0=D1=80=D0=B0=D1=9F=D0=B8=D1=9B?= Date: Fri, 18 Jan 2019 17:21:31 -0800 Subject: [PATCH] Updated spirv-cross. --- .../struct-resource-name-aliasing.asm.comp | 16 ++ .../comp/force-recompile-hooks.swizzle.comp | 138 +++++++++++++++++ .../struct-resource-name-aliasing.asm.comp | 23 +++ .../comp/force-recompile-hooks.swizzle.comp | 139 ++++++++++++++++++ .../struct-resource-name-aliasing.asm.comp | 47 ++++++ .../comp/force-recompile-hooks.swizzle.comp | 9 ++ 3rdparty/spirv-cross/spirv_msl.cpp | 98 ++++++++---- 3rdparty/spirv-cross/spirv_msl.hpp | 1 + 8 files changed, 439 insertions(+), 32 deletions(-) create mode 100644 3rdparty/spirv-cross/reference/opt/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp create mode 100644 3rdparty/spirv-cross/reference/opt/shaders-msl/comp/force-recompile-hooks.swizzle.comp create mode 100644 3rdparty/spirv-cross/reference/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp create mode 100644 3rdparty/spirv-cross/reference/shaders-msl/comp/force-recompile-hooks.swizzle.comp create mode 100644 3rdparty/spirv-cross/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp create mode 100644 3rdparty/spirv-cross/shaders-msl/comp/force-recompile-hooks.swizzle.comp diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp b/3rdparty/spirv-cross/reference/opt/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp new file mode 100644 index 000000000..5b1f0a036 --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp @@ -0,0 +1,16 @@ +#include +#include + +using namespace metal; + +struct bufA +{ + uint _data[1]; +}; + +kernel void main0(device bufA& bufA_1 [[buffer(0)]], device bufA& bufB [[buffer(1)]]) +{ + bufA_1._data[0] = 0u; + bufB._data[0] = 0u; +} + 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 new file mode 100644 index 000000000..267cc518b --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/force-recompile-hooks.swizzle.comp @@ -0,0 +1,138 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct spvAux +{ + uint swizzleConst[1]; +}; + +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)...); +} + +kernel void main0(constant spvAux& spvAuxBuffer [[buffer(0)]], texture2d foo [[texture(0)]], texture2d bar [[texture(1)]], sampler fooSmplr [[sampler(0)]]) +{ + constant uint32_t& fooSwzl = spvAuxBuffer.swizzleConst[0]; + bar.write(spvTextureSwizzle(foo.sample(fooSmplr, float2(1.0), level(0.0)), fooSwzl), uint2(int2(0))); +} + diff --git a/3rdparty/spirv-cross/reference/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp b/3rdparty/spirv-cross/reference/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp new file mode 100644 index 000000000..a3323bf25 --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp @@ -0,0 +1,23 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct bufA +{ + uint _data[1]; +}; + +void _main(device bufA& bufA_1, device bufA& bufB) +{ + bufA_1._data[0] = 0u; + bufB._data[0] = 0u; +} + +kernel void main0(device bufA& bufA_1 [[buffer(0)]], device bufA& bufB [[buffer(1)]]) +{ + _main(bufA_1, bufB); +} + 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 new file mode 100644 index 000000000..667819dc2 --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-msl/comp/force-recompile-hooks.swizzle.comp @@ -0,0 +1,139 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct spvAux +{ + uint swizzleConst[1]; +}; + +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)...); +} + +kernel void main0(constant spvAux& spvAuxBuffer [[buffer(0)]], texture2d foo [[texture(0)]], texture2d bar [[texture(1)]], sampler fooSmplr [[sampler(0)]]) +{ + constant uint32_t& fooSwzl = spvAuxBuffer.swizzleConst[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/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp b/3rdparty/spirv-cross/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp new file mode 100644 index 000000000..98d31537e --- /dev/null +++ b/3rdparty/spirv-cross/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp @@ -0,0 +1,47 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 7 +; Bound: 21 +; Schema: 0 + OpCapability Shader + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 8 8 1 + OpSource HLSL 500 + OpName %main "main" + OpName %_main_ "@main(" + OpName %bufA "bufA" + OpMemberName %bufA 0 "@data" + OpName %bufA_0 "bufA" + OpName %bufB "bufB" + OpDecorate %_runtimearr_uint ArrayStride 4 + OpMemberDecorate %bufA 0 Offset 0 + OpDecorate %bufA BufferBlock + OpDecorate %bufA_0 DescriptorSet 0 + OpDecorate %bufB DescriptorSet 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %uint = OpTypeInt 32 0 +%_runtimearr_uint = OpTypeRuntimeArray %uint + %bufA = OpTypeStruct %_runtimearr_uint +%_ptr_Uniform_bufA = OpTypePointer Uniform %bufA + %bufA_0 = OpVariable %_ptr_Uniform_bufA Uniform + %int = OpTypeInt 32 1 + %int_0 = OpConstant %int 0 + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %bufB = OpVariable %_ptr_Uniform_bufA Uniform + %main = OpFunction %void None %3 + %5 = OpLabel + %20 = OpFunctionCall %void %_main_ + OpReturn + OpFunctionEnd + %_main_ = OpFunction %void None %3 + %7 = OpLabel + %17 = OpAccessChain %_ptr_Uniform_uint %bufA_0 %int_0 %int_0 + OpStore %17 %uint_0 + %19 = OpAccessChain %_ptr_Uniform_uint %bufB %int_0 %int_0 + OpStore %19 %uint_0 + OpReturn + OpFunctionEnd diff --git a/3rdparty/spirv-cross/shaders-msl/comp/force-recompile-hooks.swizzle.comp b/3rdparty/spirv-cross/shaders-msl/comp/force-recompile-hooks.swizzle.comp new file mode 100644 index 000000000..2752d3051 --- /dev/null +++ b/3rdparty/spirv-cross/shaders-msl/comp/force-recompile-hooks.swizzle.comp @@ -0,0 +1,9 @@ +#version 450 + +layout(binding = 0) uniform sampler2D foo; +layout(binding = 1, rgba8) uniform image2D bar; + +void main() { + vec4 a = texture(foo, vec2(1, 1)); + imageStore(bar, ivec2(0, 0), a); +} diff --git a/3rdparty/spirv-cross/spirv_msl.cpp b/3rdparty/spirv-cross/spirv_msl.cpp index 87b00ea1f..98fe3ae05 100644 --- a/3rdparty/spirv-cross/spirv_msl.cpp +++ b/3rdparty/spirv-cross/spirv_msl.cpp @@ -447,6 +447,10 @@ string CompilerMSL::compile() // Mark any non-stage-in structs to be tightly packed. mark_packable_structs(); + // Add fixup hooks required by shader inputs and outputs. This needs to happen before + // the loop, so the hooks aren't added multiple times. + fix_up_shader_inputs_outputs(); + uint32_t pass_count = 0; do { @@ -4420,6 +4424,7 @@ string CompilerMSL::entry_point_args(bool append_comma) if (!ep_args.empty()) ep_args += ", "; + add_resource_name(var.self); ep_args += type_to_glsl(type) + " " + to_name(var.self) + " [[stage_in]]"; } @@ -4448,6 +4453,7 @@ string CompilerMSL::entry_point_args(bool append_comma) { if (type.basetype == SPIRType::SampledImage) { + add_resource_name(var_id); resources.push_back( { &id, to_name(var_id), SPIRType::Image, get_metal_resource_index(var, SPIRType::Image) }); @@ -4460,20 +4466,10 @@ string CompilerMSL::entry_point_args(bool append_comma) else if (constexpr_samplers.count(var_id) == 0) { // constexpr samplers are not declared as resources. + add_resource_name(var_id); resources.push_back( { &id, to_name(var_id), type.basetype, get_metal_resource_index(var, type.basetype) }); } - - 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)), "];"); - }); - } } }); @@ -4554,27 +4550,7 @@ string CompilerMSL::entry_point_args(bool append_comma) // point, we get that by calling get_sample_position() on the sample ID. if (var.storage == StorageClassInput && is_builtin_variable(var)) { - if (bi_type == BuiltInSamplePosition) - { - auto &entry_func = this->get(ir.default_entry_point); - entry_func.fixup_hooks_in.push_back([=]() { - statement(builtin_type_decl(bi_type), " ", to_expression(var_id), " = get_sample_position(", - to_expression(builtin_sample_id_id), ");"); - }); - } - else if (bi_type == BuiltInHelperInvocation) - { - if (msl_options.is_ios()) - SPIRV_CROSS_THROW("simd_is_helper_thread() is only supported on macOS."); - else if (msl_options.is_macos() && !msl_options.supports_msl_version(2, 1)) - SPIRV_CROSS_THROW("simd_is_helper_thread() requires version 2.1 on macOS."); - - auto &entry_func = this->get(ir.default_entry_point); - entry_func.fixup_hooks_in.push_back([=]() { - statement(builtin_type_decl(bi_type), " ", to_expression(var_id), " = simd_is_helper_thread();"); - }); - } - else + if (bi_type != BuiltInSamplePosition && bi_type != BuiltInHelperInvocation) { if (!ep_args.empty()) ep_args += ", "; @@ -4598,6 +4574,64 @@ string CompilerMSL::entry_point_args(bool append_comma) return ep_args; } +void CompilerMSL::fix_up_shader_inputs_outputs() +{ + // Look for sampled images. Add hooks to set up the swizzle constants. + ir.for_each_typed_id([&](uint32_t, SPIRVariable &var) { + auto &type = get_variable_data_type(var); + + uint32_t var_id = var.self; + + if ((var.storage == StorageClassUniform || var.storage == StorageClassUniformConstant || + var.storage == StorageClassPushConstant || var.storage == StorageClassStorageBuffer) && + !is_hidden_variable(var)) + { + if (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)), "];"); + }); + } + } + }); + + // Builtin variables + ir.for_each_typed_id([&](uint32_t, SPIRVariable &var) { + uint32_t var_id = var.self; + BuiltIn bi_type = ir.meta[var_id].decoration.builtin_type; + + if (var.storage == StorageClassInput && is_builtin_variable(var)) + { + auto &entry_func = this->get(ir.default_entry_point); + switch (bi_type) + { + case BuiltInSamplePosition: + entry_func.fixup_hooks_in.push_back([=]() { + statement(builtin_type_decl(bi_type), " ", to_expression(var_id), " = get_sample_position(", + to_expression(builtin_sample_id_id), ");"); + }); + break; + case BuiltInHelperInvocation: + if (msl_options.is_ios()) + SPIRV_CROSS_THROW("simd_is_helper_thread() is only supported on macOS."); + else if (msl_options.is_macos() && !msl_options.supports_msl_version(2, 1)) + SPIRV_CROSS_THROW("simd_is_helper_thread() requires version 2.1 on macOS."); + + entry_func.fixup_hooks_in.push_back([=]() { + statement(builtin_type_decl(bi_type), " ", to_expression(var_id), " = simd_is_helper_thread();"); + }); + break; + default: + break; + } + } + }); +} + // Returns the Metal index of the resource of the specified type as used by the specified variable. uint32_t CompilerMSL::get_metal_resource_index(SPIRVariable &var, SPIRType::BaseType basetype) { diff --git a/3rdparty/spirv-cross/spirv_msl.hpp b/3rdparty/spirv-cross/spirv_msl.hpp index 8fcb224b0..3ea5b4d09 100644 --- a/3rdparty/spirv-cross/spirv_msl.hpp +++ b/3rdparty/spirv-cross/spirv_msl.hpp @@ -392,6 +392,7 @@ protected: void emit_interface_block(uint32_t ib_var_id); bool maybe_emit_array_assignment(uint32_t id_lhs, uint32_t id_rhs); void add_convert_row_major_matrix_function(uint32_t cols, uint32_t rows); + void fix_up_shader_inputs_outputs(); std::string func_type_decl(SPIRType &type); std::string entry_point_args(bool append_comma);