diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/asm/comp/buffer-write.asm.comp b/3rdparty/spirv-cross/reference/opt/shaders-msl/asm/comp/buffer-write.asm.comp index ddf958289..ab375a397 100644 --- a/3rdparty/spirv-cross/reference/opt/shaders-msl/asm/comp/buffer-write.asm.comp +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/asm/comp/buffer-write.asm.comp @@ -18,6 +18,6 @@ uint2 spvTexelBufferCoord(uint tc) kernel void main0(constant cb& _6 [[buffer(7)]], texture2d _buffer [[texture(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]]) { - _buffer.write(_6.value, spvTexelBufferCoord(((32u * gl_WorkGroupID.x) + gl_LocalInvocationIndex))); + _buffer.write(float4(_6.value), spvTexelBufferCoord(((32u * gl_WorkGroupID.x) + gl_LocalInvocationIndex))); } diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/asm/comp/image-load-store-short-vector.asm.comp b/3rdparty/spirv-cross/reference/opt/shaders-msl/asm/comp/image-load-store-short-vector.asm.comp new file mode 100644 index 000000000..fb97d0da9 --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/asm/comp/image-load-store-short-vector.asm.comp @@ -0,0 +1,10 @@ +#include +#include + +using namespace metal; + +kernel void main0(texture2d TargetTexture [[texture(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) +{ + TargetTexture.write((TargetTexture.read(uint2(gl_WorkGroupID.xy)).xy + float2(1.0)).xyyy, uint2((gl_WorkGroupID.xy + uint2(1u)))); +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/storage-buffer-std140-vector-array.comp b/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/storage-buffer-std140-vector-array.comp new file mode 100644 index 000000000..905222d39 --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/storage-buffer-std140-vector-array.comp @@ -0,0 +1,48 @@ +#include +#include + +using namespace metal; + +struct Sub +{ + float4 f[2]; + float4 f2[2]; + float3 f3[2]; + float4 f4[2]; +}; + +struct SSBO +{ + Sub sub[2]; +}; + +kernel void main0(device SSBO& _27 [[buffer(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + float _153[2]; + _153[0] = _27.sub[gl_WorkGroupID.x].f[0].x; + _153[1] = _27.sub[gl_WorkGroupID.x].f[1].x; + float2 _154[2]; + _154[0] = _27.sub[gl_WorkGroupID.x].f2[0].xy; + _154[1] = _27.sub[gl_WorkGroupID.x].f2[1].xy; + float3 _155[2]; + _155[0] = _27.sub[gl_WorkGroupID.x].f3[0]; + _155[1] = _27.sub[gl_WorkGroupID.x].f3[1]; + float4 _156[2]; + _156[0] = _27.sub[gl_WorkGroupID.x].f4[0]; + _156[1] = _27.sub[gl_WorkGroupID.x].f4[1]; + _153[gl_GlobalInvocationID.x] += 1.0; + _154[gl_GlobalInvocationID.x] += float2(2.0); + _155[gl_GlobalInvocationID.x] += float3(3.0); + _156[gl_GlobalInvocationID.x] += float4(4.0); + _27.sub[gl_WorkGroupID.x].f[0].x = _153[0]; + _27.sub[gl_WorkGroupID.x].f[1].x = _153[1]; + _27.sub[gl_WorkGroupID.x].f2[0].xy = _154[0]; + _27.sub[gl_WorkGroupID.x].f2[1].xy = _154[1]; + _27.sub[gl_WorkGroupID.x].f3[0] = _155[0]; + _27.sub[gl_WorkGroupID.x].f3[1] = _155[1]; + _27.sub[gl_WorkGroupID.x].f4[0] = _156[0]; + _27.sub[gl_WorkGroupID.x].f4[1] = _156[1]; + _27.sub[0].f[0].x += 5.0; + _27.sub[0].f2[1].xy += float2(5.0); +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/struct-packing.comp b/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/struct-packing.comp index 2b37844fe..468eb7e6e 100644 --- a/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/struct-packing.comp +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/comp/struct-packing.comp @@ -69,7 +69,7 @@ struct SSBO1 struct S0_1 { - float2 a[1]; + float4 a[1]; float b; }; @@ -115,21 +115,21 @@ struct SSBO0 Content_1 content; Content_1 content1[2]; Content_1 content2; - float array[1]; + float4 array[1]; }; kernel void main0(device SSBO0& ssbo_140 [[buffer(0)]], device SSBO1& ssbo_430 [[buffer(1)]]) { Content_1 _60 = ssbo_140.content; - ssbo_430.content.m0s[0].a[0] = _60.m0s[0].a[0]; + ssbo_430.content.m0s[0].a[0] = _60.m0s[0].a[0].xy; ssbo_430.content.m0s[0].b = _60.m0s[0].b; - ssbo_430.content.m1s[0].a = _60.m1s[0].a; + ssbo_430.content.m1s[0].a = float3(_60.m1s[0].a); ssbo_430.content.m1s[0].b = _60.m1s[0].b; ssbo_430.content.m2s[0].a[0] = _60.m2s[0].a[0]; ssbo_430.content.m2s[0].b = _60.m2s[0].b; - ssbo_430.content.m0.a[0] = _60.m0.a[0]; + ssbo_430.content.m0.a[0] = _60.m0.a[0].xy; ssbo_430.content.m0.b = _60.m0.b; - ssbo_430.content.m1.a = _60.m1.a; + ssbo_430.content.m1.a = float3(_60.m1.a); ssbo_430.content.m1.b = _60.m1.b; ssbo_430.content.m2.a[0] = _60.m2.a[0]; ssbo_430.content.m2.b = _60.m2.b; diff --git a/3rdparty/spirv-cross/reference/shaders-msl/asm/comp/buffer-write.asm.comp b/3rdparty/spirv-cross/reference/shaders-msl/asm/comp/buffer-write.asm.comp index ddf958289..ab375a397 100644 --- a/3rdparty/spirv-cross/reference/shaders-msl/asm/comp/buffer-write.asm.comp +++ b/3rdparty/spirv-cross/reference/shaders-msl/asm/comp/buffer-write.asm.comp @@ -18,6 +18,6 @@ uint2 spvTexelBufferCoord(uint tc) kernel void main0(constant cb& _6 [[buffer(7)]], texture2d _buffer [[texture(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]]) { - _buffer.write(_6.value, spvTexelBufferCoord(((32u * gl_WorkGroupID.x) + gl_LocalInvocationIndex))); + _buffer.write(float4(_6.value), spvTexelBufferCoord(((32u * gl_WorkGroupID.x) + gl_LocalInvocationIndex))); } diff --git a/3rdparty/spirv-cross/reference/shaders-msl/asm/comp/image-load-store-short-vector.asm.comp b/3rdparty/spirv-cross/reference/shaders-msl/asm/comp/image-load-store-short-vector.asm.comp new file mode 100644 index 000000000..c90faf9ef --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-msl/asm/comp/image-load-store-short-vector.asm.comp @@ -0,0 +1,21 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +void _main(thread const uint3& id, thread texture2d TargetTexture) +{ + float2 loaded = TargetTexture.read(uint2(id.xy)).xy; + float2 storeTemp = loaded + float2(1.0); + TargetTexture.write(storeTemp.xyyy, uint2((id.xy + uint2(1u)))); +} + +kernel void main0(texture2d TargetTexture [[texture(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) +{ + uint3 id = gl_WorkGroupID; + uint3 param = id; + _main(param, TargetTexture); +} + diff --git a/3rdparty/spirv-cross/reference/shaders-msl/asm/frag/extract-packed-from-composite.asm.frag b/3rdparty/spirv-cross/reference/shaders-msl/asm/frag/extract-packed-from-composite.asm.frag index 7d51b1538..957b77f2e 100644 --- a/3rdparty/spirv-cross/reference/shaders-msl/asm/frag/extract-packed-from-composite.asm.frag +++ b/3rdparty/spirv-cross/reference/shaders-msl/asm/frag/extract-packed-from-composite.asm.frag @@ -32,7 +32,7 @@ float4 _main(thread const float4& pos, constant buf& v_11) { int _46 = int(pos.x) % 16; Foo_1 foo; - foo.a = v_11.results[_46].a; + foo.a = float3(v_11.results[_46].a); foo.b = v_11.results[_46].b; return float4(dot(foo.a, v_11.bar.xyz), foo.b, 0.0, 0.0); } diff --git a/3rdparty/spirv-cross/reference/shaders-msl/comp/packing-test-1.comp b/3rdparty/spirv-cross/reference/shaders-msl/comp/packing-test-1.comp index 92351c67a..06160ebc3 100644 --- a/3rdparty/spirv-cross/reference/shaders-msl/comp/packing-test-1.comp +++ b/3rdparty/spirv-cross/reference/shaders-msl/comp/packing-test-1.comp @@ -30,7 +30,7 @@ constant uint3 gl_WorkGroupSize = uint3(32u, 1u, 1u); kernel void main0(device Buffer0& _15 [[buffer(1)]], device Buffer1& _34 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { T1_1 v; - v.a = _15.buf0[0].a; + v.a = float3(_15.buf0[0].a); v.b = _15.buf0[0].b; float x = v.b; _34.buf1[gl_GlobalInvocationID.x] = x; diff --git a/3rdparty/spirv-cross/reference/shaders-msl/comp/storage-buffer-std140-vector-array.comp b/3rdparty/spirv-cross/reference/shaders-msl/comp/storage-buffer-std140-vector-array.comp new file mode 100644 index 000000000..a5a1cf9ec --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-msl/comp/storage-buffer-std140-vector-array.comp @@ -0,0 +1,53 @@ +#include +#include + +using namespace metal; + +struct Sub +{ + float4 f[2]; + float4 f2[2]; + float3 f3[2]; + float4 f4[2]; +}; + +struct Sub_1 +{ + float f[2]; + float2 f2[2]; + float3 f3[2]; + float4 f4[2]; +}; + +struct SSBO +{ + Sub sub[2]; +}; + +kernel void main0(device SSBO& _27 [[buffer(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + Sub_1 foo; + foo.f[0] = _27.sub[gl_WorkGroupID.x].f[0].x; + foo.f[1] = _27.sub[gl_WorkGroupID.x].f[1].x; + foo.f2[0] = _27.sub[gl_WorkGroupID.x].f2[0].xy; + foo.f2[1] = _27.sub[gl_WorkGroupID.x].f2[1].xy; + foo.f3[0] = _27.sub[gl_WorkGroupID.x].f3[0]; + foo.f3[1] = _27.sub[gl_WorkGroupID.x].f3[1]; + foo.f4[0] = _27.sub[gl_WorkGroupID.x].f4[0]; + foo.f4[1] = _27.sub[gl_WorkGroupID.x].f4[1]; + foo.f[gl_GlobalInvocationID.x] += 1.0; + foo.f2[gl_GlobalInvocationID.x] += float2(2.0); + foo.f3[gl_GlobalInvocationID.x] += float3(3.0); + foo.f4[gl_GlobalInvocationID.x] += float4(4.0); + _27.sub[gl_WorkGroupID.x].f[0].x = foo.f[0]; + _27.sub[gl_WorkGroupID.x].f[1].x = foo.f[1]; + _27.sub[gl_WorkGroupID.x].f2[0].xy = foo.f2[0]; + _27.sub[gl_WorkGroupID.x].f2[1].xy = foo.f2[1]; + _27.sub[gl_WorkGroupID.x].f3[0] = foo.f3[0]; + _27.sub[gl_WorkGroupID.x].f3[1] = foo.f3[1]; + _27.sub[gl_WorkGroupID.x].f4[0] = foo.f4[0]; + _27.sub[gl_WorkGroupID.x].f4[1] = foo.f4[1]; + _27.sub[0].f[0].x += 5.0; + _27.sub[0].f2[1].xy += float2(5.0); +} + diff --git a/3rdparty/spirv-cross/reference/shaders-msl/comp/struct-packing.comp b/3rdparty/spirv-cross/reference/shaders-msl/comp/struct-packing.comp index 2b37844fe..468eb7e6e 100644 --- a/3rdparty/spirv-cross/reference/shaders-msl/comp/struct-packing.comp +++ b/3rdparty/spirv-cross/reference/shaders-msl/comp/struct-packing.comp @@ -69,7 +69,7 @@ struct SSBO1 struct S0_1 { - float2 a[1]; + float4 a[1]; float b; }; @@ -115,21 +115,21 @@ struct SSBO0 Content_1 content; Content_1 content1[2]; Content_1 content2; - float array[1]; + float4 array[1]; }; kernel void main0(device SSBO0& ssbo_140 [[buffer(0)]], device SSBO1& ssbo_430 [[buffer(1)]]) { Content_1 _60 = ssbo_140.content; - ssbo_430.content.m0s[0].a[0] = _60.m0s[0].a[0]; + ssbo_430.content.m0s[0].a[0] = _60.m0s[0].a[0].xy; ssbo_430.content.m0s[0].b = _60.m0s[0].b; - ssbo_430.content.m1s[0].a = _60.m1s[0].a; + ssbo_430.content.m1s[0].a = float3(_60.m1s[0].a); ssbo_430.content.m1s[0].b = _60.m1s[0].b; ssbo_430.content.m2s[0].a[0] = _60.m2s[0].a[0]; ssbo_430.content.m2s[0].b = _60.m2s[0].b; - ssbo_430.content.m0.a[0] = _60.m0.a[0]; + ssbo_430.content.m0.a[0] = _60.m0.a[0].xy; ssbo_430.content.m0.b = _60.m0.b; - ssbo_430.content.m1.a = _60.m1.a; + ssbo_430.content.m1.a = float3(_60.m1.a); ssbo_430.content.m1.b = _60.m1.b; ssbo_430.content.m2.a[0] = _60.m2.a[0]; ssbo_430.content.m2.b = _60.m2.b; diff --git a/3rdparty/spirv-cross/reference/shaders-msl/frag/packing-test-3.frag b/3rdparty/spirv-cross/reference/shaders-msl/frag/packing-test-3.frag index a02884e85..916ab8757 100644 --- a/3rdparty/spirv-cross/reference/shaders-msl/frag/packing-test-3.frag +++ b/3rdparty/spirv-cross/reference/shaders-msl/frag/packing-test-3.frag @@ -35,7 +35,7 @@ struct main0_out float4 _main(thread const VertexOutput& IN, constant CB0& v_26) { TestStruct_1 st; - st.position = v_26.CB0[1].position; + st.position = float3(v_26.CB0[1].position); st.radius = v_26.CB0[1].radius; float4 col = float4(st.position, st.radius); return col; diff --git a/3rdparty/spirv-cross/reference/shaders-msl/vert/copy.flatten.vert b/3rdparty/spirv-cross/reference/shaders-msl/vert/copy.flatten.vert index a87b44782..3230e4fbd 100644 --- a/3rdparty/spirv-cross/reference/shaders-msl/vert/copy.flatten.vert +++ b/3rdparty/spirv-cross/reference/shaders-msl/vert/copy.flatten.vert @@ -43,7 +43,7 @@ vertex main0_out main0(main0_in in [[stage_in]], constant UBO& _21 [[buffer(0)]] for (int i = 0; i < 4; i++) { Light_1 light; - light.Position = _21.lights[i].Position; + light.Position = float3(_21.lights[i].Position); light.Radius = _21.lights[i].Radius; light.Color = _21.lights[i].Color; float3 L = in.aVertex.xyz - light.Position; diff --git a/3rdparty/spirv-cross/shaders-msl/asm/comp/image-load-store-short-vector.asm.comp b/3rdparty/spirv-cross/shaders-msl/asm/comp/image-load-store-short-vector.asm.comp new file mode 100644 index 000000000..8f759293e --- /dev/null +++ b/3rdparty/spirv-cross/shaders-msl/asm/comp/image-load-store-short-vector.asm.comp @@ -0,0 +1,75 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 7 +; Bound: 44 +; Schema: 0 + OpCapability Shader + OpCapability StorageImageExtendedFormats + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %id_1 + OpExecutionMode %main LocalSize 1 1 1 + OpSource HLSL 500 + OpName %main "main" + OpName %_main_vu3_ "@main(vu3;" + OpName %id "id" + OpName %loaded "loaded" + OpName %TargetTexture "TargetTexture" + OpName %storeTemp "storeTemp" + OpName %id_0 "id" + OpName %id_1 "id" + OpName %param "param" + OpDecorate %TargetTexture DescriptorSet 0 + OpDecorate %TargetTexture Binding 0 + OpDecorate %id_1 BuiltIn WorkgroupId + %void = OpTypeVoid + %3 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 +%_ptr_Function_v3uint = OpTypePointer Function %v3uint + %9 = OpTypeFunction %void %_ptr_Function_v3uint + %float = OpTypeFloat 32 + %v2float = OpTypeVector %float 2 +%_ptr_Function_v2float = OpTypePointer Function %v2float + %17 = OpTypeImage %float 2D 0 0 0 2 Rg32f +%_ptr_UniformConstant_17 = OpTypePointer UniformConstant %17 +%TargetTexture = OpVariable %_ptr_UniformConstant_17 UniformConstant + %v2uint = OpTypeVector %uint 2 + %float_1 = OpConstant %float 1 + %uint_1 = OpConstant %uint 1 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %id_1 = OpVariable %_ptr_Input_v3uint Input + %main = OpFunction %void None %3 + %5 = OpLabel + %id_0 = OpVariable %_ptr_Function_v3uint Function + %param = OpVariable %_ptr_Function_v3uint Function + %40 = OpLoad %v3uint %id_1 + OpStore %id_0 %40 + %42 = OpLoad %v3uint %id_0 + OpStore %param %42 + %43 = OpFunctionCall %void %_main_vu3_ %param + OpReturn + OpFunctionEnd + %_main_vu3_ = OpFunction %void None %9 + %id = OpFunctionParameter %_ptr_Function_v3uint + %12 = OpLabel + %loaded = OpVariable %_ptr_Function_v2float Function + %storeTemp = OpVariable %_ptr_Function_v2float Function + %20 = OpLoad %17 %TargetTexture + %22 = OpLoad %v3uint %id + %23 = OpVectorShuffle %v2uint %22 %22 0 1 + %24 = OpImageRead %v2float %20 %23 + OpStore %loaded %24 + %26 = OpLoad %v2float %loaded + %28 = OpCompositeConstruct %v2float %float_1 %float_1 + %29 = OpFAdd %v2float %26 %28 + OpStore %storeTemp %29 + %30 = OpLoad %17 %TargetTexture + %31 = OpLoad %v3uint %id + %32 = OpVectorShuffle %v2uint %31 %31 0 1 + %34 = OpCompositeConstruct %v2uint %uint_1 %uint_1 + %35 = OpIAdd %v2uint %32 %34 + %36 = OpLoad %v2float %storeTemp + OpImageWrite %30 %35 %36 + OpReturn + OpFunctionEnd diff --git a/3rdparty/spirv-cross/shaders-msl/comp/storage-buffer-std140-vector-array.comp b/3rdparty/spirv-cross/shaders-msl/comp/storage-buffer-std140-vector-array.comp new file mode 100644 index 000000000..7e786ec6c --- /dev/null +++ b/3rdparty/spirv-cross/shaders-msl/comp/storage-buffer-std140-vector-array.comp @@ -0,0 +1,30 @@ +#version 450 +layout(local_size_x = 1) in; + +struct Sub +{ + float f[2]; + vec2 f2[2]; + vec3 f3[2]; + vec4 f4[2]; +}; + +layout(std140, binding = 0) buffer SSBO +{ + Sub sub[2]; +}; + +void main() +{ + Sub foo = sub[gl_WorkGroupID.x]; + + foo.f[gl_GlobalInvocationID.x] += 1.0; + foo.f2[gl_GlobalInvocationID.x] += 2.0; + foo.f3[gl_GlobalInvocationID.x] += 3.0; + foo.f4[gl_GlobalInvocationID.x] += 4.0; + sub[gl_WorkGroupID.x] = foo; + + sub[0].f[0] += 5.0; + sub[0].f2[1] += 5.0; +} + diff --git a/3rdparty/spirv-cross/spirv_common.hpp b/3rdparty/spirv-cross/spirv_common.hpp index 57820d052..6852542ee 100644 --- a/3rdparty/spirv-cross/spirv_common.hpp +++ b/3rdparty/spirv-cross/spirv_common.hpp @@ -1339,6 +1339,7 @@ T &variant_set(Variant &var, P &&... args) struct AccessChainMeta { + uint32_t storage_packed_type = 0; bool need_transpose = false; bool storage_is_packed = false; bool storage_is_invariant = false; @@ -1365,6 +1366,12 @@ struct Meta uint32_t index = 0; spv::FPRoundingMode fp_rounding_mode = spv::FPRoundingModeMax; bool builtin = false; + + struct + { + uint32_t packed_type = 0; + bool packed = false; + } extended; }; Decoration decoration; diff --git a/3rdparty/spirv-cross/spirv_cross.cpp b/3rdparty/spirv-cross/spirv_cross.cpp index ecddc056e..ef3ec249c 100644 --- a/3rdparty/spirv-cross/spirv_cross.cpp +++ b/3rdparty/spirv-cross/spirv_cross.cpp @@ -169,7 +169,7 @@ string Compiler::to_name(uint32_t id, bool allow_alias) const { // If the alias master has been specially packed, we will have emitted a clean variant as well, // so skip the name aliasing here. - if (!has_decoration(type.type_alias, DecorationCPacked)) + if (!has_extended_decoration(type.type_alias, SPIRVCrossDecorationPacked)) return to_name(type.type_alias); } } @@ -515,7 +515,7 @@ bool Compiler::is_member_builtin(const SPIRType &type, uint32_t index, BuiltIn * bool Compiler::is_scalar(const SPIRType &type) const { - return type.vecsize == 1 && type.columns == 1; + return type.basetype != SPIRType::Struct && type.vecsize == 1 && type.columns == 1; } bool Compiler::is_vector(const SPIRType &type) const @@ -868,7 +868,7 @@ void Compiler::fixup_type_alias() for (auto alias_itr = begin(type_ids); alias_itr != end(type_ids); ++alias_itr) { auto &type = get(*alias_itr); - if (type.type_alias != 0 && !has_decoration(type.type_alias, DecorationCPacked)) + if (type.type_alias != 0 && !has_extended_decoration(type.type_alias, SPIRVCrossDecorationPacked)) { // We will skip declaring this type, so make sure the type_alias type comes before. auto master_itr = find(begin(type_ids), end(type_ids), type.type_alias); @@ -1165,6 +1165,153 @@ void Compiler::set_decoration(uint32_t id, Decoration decoration, uint32_t argum ir.set_decoration(id, decoration, argument); } +void Compiler::set_extended_decoration(uint32_t id, ExtendedDecorations decoration, uint32_t value) +{ + auto &dec = ir.meta[id].decoration; + switch (decoration) + { + case SPIRVCrossDecorationPacked: + dec.extended.packed = true; + break; + + case SPIRVCrossDecorationPackedType: + dec.extended.packed_type = value; + break; + } +} + +void Compiler::set_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration, + uint32_t value) +{ + ir.meta[type].members.resize(max(ir.meta[type].members.size(), size_t(index) + 1)); + auto &dec = ir.meta[type].members[index]; + + switch (decoration) + { + case SPIRVCrossDecorationPacked: + dec.extended.packed = true; + break; + + case SPIRVCrossDecorationPackedType: + dec.extended.packed_type = value; + break; + } +} + +uint32_t Compiler::get_extended_decoration(uint32_t id, ExtendedDecorations decoration) const +{ + auto *m = ir.find_meta(id); + if (!m) + return 0; + + auto &dec = m->decoration; + switch (decoration) + { + case SPIRVCrossDecorationPacked: + return uint32_t(dec.extended.packed); + + case SPIRVCrossDecorationPackedType: + return dec.extended.packed_type; + } + + return 0; +} + +uint32_t Compiler::get_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration) const +{ + auto *m = ir.find_meta(type); + if (!m) + return 0; + + if (index >= m->members.size()) + return 0; + + auto &dec = m->members[index]; + switch (decoration) + { + case SPIRVCrossDecorationPacked: + return uint32_t(dec.extended.packed); + + case SPIRVCrossDecorationPackedType: + return dec.extended.packed_type; + } + + return 0; +} + +bool Compiler::has_extended_decoration(uint32_t id, ExtendedDecorations decoration) const +{ + auto *m = ir.find_meta(id); + if (!m) + return false; + + auto &dec = m->decoration; + switch (decoration) + { + case SPIRVCrossDecorationPacked: + return dec.extended.packed; + + case SPIRVCrossDecorationPackedType: + return dec.extended.packed_type != 0; + } + + return false; +} + +bool Compiler::has_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration) const +{ + auto *m = ir.find_meta(type); + if (!m) + return false; + + if (index >= m->members.size()) + return false; + + auto &dec = m->members[index]; + switch (decoration) + { + case SPIRVCrossDecorationPacked: + return dec.extended.packed; + + case SPIRVCrossDecorationPackedType: + return dec.extended.packed_type != 0; + } + + return false; +} + +void Compiler::unset_extended_decoration(uint32_t id, ExtendedDecorations decoration) +{ + auto &dec = ir.meta[id].decoration; + switch (decoration) + { + case SPIRVCrossDecorationPacked: + dec.extended.packed = false; + break; + + case SPIRVCrossDecorationPackedType: + dec.extended.packed_type = 0; + break; + } +} + +void Compiler::unset_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration) +{ + ir.meta[type].members.resize(max(ir.meta[type].members.size(), size_t(index) + 1)); + auto &dec = ir.meta[type].members[index]; + + switch (decoration) + { + case SPIRVCrossDecorationPacked: + dec.extended.packed = false; + break; + + case SPIRVCrossDecorationPackedType: + dec.extended.packed_type = 0; + break; + } +} + StorageClass Compiler::get_storage_class(uint32_t id) const { return get(id).storage; diff --git a/3rdparty/spirv-cross/spirv_cross.hpp b/3rdparty/spirv-cross/spirv_cross.hpp index 82e2e3331..d587e5095 100644 --- a/3rdparty/spirv-cross/spirv_cross.hpp +++ b/3rdparty/spirv-cross/spirv_cross.hpp @@ -114,6 +114,12 @@ struct EntryPoint spv::ExecutionModel execution_model; }; +enum ExtendedDecorations +{ + SPIRVCrossDecorationPacked, + SPIRVCrossDecorationPackedType +}; + class Compiler { public: @@ -946,6 +952,17 @@ protected: bool image_is_comparison(const SPIRType &type, uint32_t id) const; + void set_extended_decoration(uint32_t id, ExtendedDecorations decoration, uint32_t value = 0); + uint32_t get_extended_decoration(uint32_t id, ExtendedDecorations decoration) const; + bool has_extended_decoration(uint32_t id, ExtendedDecorations decoration) const; + void unset_extended_decoration(uint32_t id, ExtendedDecorations decoration); + + void set_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration, + uint32_t value = 0); + uint32_t get_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration) const; + bool has_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration) const; + void unset_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration); + private: // Used only to implement the old deprecated get_entry_point() interface. const SPIREntryPoint &get_first_entry_point(const std::string &name) const; diff --git a/3rdparty/spirv-cross/spirv_glsl.cpp b/3rdparty/spirv-cross/spirv_glsl.cpp index 1a3f08cd4..131bc803d 100644 --- a/3rdparty/spirv-cross/spirv_glsl.cpp +++ b/3rdparty/spirv-cross/spirv_glsl.cpp @@ -696,7 +696,7 @@ void CompilerGLSL::emit_struct(SPIRType &type) // Type-punning with these types is legal, which complicates things // when we are storing struct and array types in an SSBO for example. // If the type master is packed however, we can no longer assume that the struct declaration will be redundant. - if (type.type_alias != 0 && !has_decoration(type.type_alias, DecorationCPacked)) + if (type.type_alias != 0 && !has_extended_decoration(type.type_alias, SPIRVCrossDecorationPacked)) return; add_resource_name(type.self); @@ -810,9 +810,9 @@ string CompilerGLSL::layout_for_member(const SPIRType &type, uint32_t index) SPIRV_CROSS_THROW("Component decoration is not supported in ES targets."); } - // DecorationCPacked is set by layout_for_variable earlier to mark that we need to emit offset qualifiers. + // SPIRVCrossDecorationPacked is set by layout_for_variable earlier to mark that we need to emit offset qualifiers. // This is only done selectively in GLSL as needed. - if (has_decoration(type.self, DecorationCPacked) && dec.decoration_flags.get(DecorationOffset)) + if (has_extended_decoration(type.self, SPIRVCrossDecorationPacked) && dec.decoration_flags.get(DecorationOffset)) attr.push_back(join("offset = ", dec.offset)); if (attr.empty()) @@ -1370,7 +1370,7 @@ string CompilerGLSL::layout_for_variable(const SPIRVariable &var) // layout_for_variable() will be called before the actual buffer emit. // The alternative is a full pass before codegen where we deduce this decoration, // but then we are just doing the exact same work twice, and more complexity. - set_decoration(type.self, DecorationCPacked); + set_extended_decoration(type.self, SPIRVCrossDecorationPacked); } else { @@ -1398,7 +1398,7 @@ string CompilerGLSL::layout_for_variable(const SPIRVariable &var) if (!options.es && !options.vulkan_semantics && options.version < 440) require_extension_internal("GL_ARB_enhanced_layouts"); - set_decoration(type.self, DecorationCPacked); + set_extended_decoration(type.self, SPIRVCrossDecorationPacked); } else if (buffer_is_packing_standard(type, BufferPackingStd430EnhancedLayout)) { @@ -1409,7 +1409,7 @@ string CompilerGLSL::layout_for_variable(const SPIRVariable &var) if (!options.es && !options.vulkan_semantics && options.version < 440) require_extension_internal("GL_ARB_enhanced_layouts"); - set_decoration(type.self, DecorationCPacked); + set_extended_decoration(type.self, SPIRVCrossDecorationPacked); } else { @@ -2384,7 +2384,7 @@ void CompilerGLSL::handle_invalid_expression(uint32_t id) // by wrapping the expression in a constructor of the appropriate type. // GLSL does not support packed formats, so simply return the expression. // Subclasses that do will override -string CompilerGLSL::unpack_expression_type(string expr_str, const SPIRType &) +string CompilerGLSL::unpack_expression_type(string expr_str, const SPIRType &, uint32_t) { return expr_str; } @@ -2486,26 +2486,28 @@ string CompilerGLSL::to_enclosed_expression(uint32_t id, bool register_expressio return enclose_expression(to_expression(id, register_expression_read)); } -string CompilerGLSL::to_unpacked_expression(uint32_t id) +string CompilerGLSL::to_unpacked_expression(uint32_t id, bool register_expression_read) { // If we need to transpose, it will also take care of unpacking rules. auto *e = maybe_get(id); bool need_transpose = e && e->need_transpose; - if (!need_transpose && has_decoration(id, DecorationCPacked)) - return unpack_expression_type(to_expression(id), expression_type(id)); + if (!need_transpose && has_extended_decoration(id, SPIRVCrossDecorationPacked)) + return unpack_expression_type(to_expression(id, register_expression_read), expression_type(id), + get_extended_decoration(id, SPIRVCrossDecorationPackedType)); else - return to_expression(id); + return to_expression(id, register_expression_read); } -string CompilerGLSL::to_enclosed_unpacked_expression(uint32_t id) +string CompilerGLSL::to_enclosed_unpacked_expression(uint32_t id, bool register_expression_read) { // If we need to transpose, it will also take care of unpacking rules. auto *e = maybe_get(id); bool need_transpose = e && e->need_transpose; - if (!need_transpose && has_decoration(id, DecorationCPacked)) - return unpack_expression_type(to_expression(id), expression_type(id)); + if (!need_transpose && has_extended_decoration(id, SPIRVCrossDecorationPacked)) + return unpack_expression_type(to_expression(id, register_expression_read), expression_type(id), + get_extended_decoration(id, SPIRVCrossDecorationPackedType)); else - return to_enclosed_expression(id); + return to_enclosed_expression(id, register_expression_read); } string CompilerGLSL::to_dereferenced_expression(uint32_t id, bool register_expression_read) @@ -2517,28 +2519,28 @@ string CompilerGLSL::to_dereferenced_expression(uint32_t id, bool register_expre return to_expression(id, register_expression_read); } -string CompilerGLSL::to_pointer_expression(uint32_t id) +string CompilerGLSL::to_pointer_expression(uint32_t id, bool register_expression_read) { auto &type = expression_type(id); if (type.pointer && expression_is_lvalue(id) && !should_dereference(id)) - return address_of_expression(to_enclosed_expression(id)); + return address_of_expression(to_enclosed_expression(id, register_expression_read)); else - return to_expression(id); + return to_unpacked_expression(id, register_expression_read); } -string CompilerGLSL::to_enclosed_pointer_expression(uint32_t id) +string CompilerGLSL::to_enclosed_pointer_expression(uint32_t id, bool register_expression_read) { auto &type = expression_type(id); if (type.pointer && expression_is_lvalue(id) && !should_dereference(id)) - return address_of_expression(to_enclosed_expression(id)); + return address_of_expression(to_enclosed_expression(id, register_expression_read)); else - return to_enclosed_expression(id); + return to_enclosed_unpacked_expression(id, register_expression_read); } string CompilerGLSL::to_extract_component_expression(uint32_t id, uint32_t index) { auto expr = to_enclosed_expression(id); - if (has_decoration(id, DecorationCPacked)) + if (has_extended_decoration(id, SPIRVCrossDecorationPacked)) return join(expr, "[", index, "]"); else return join(expr, ".", index_to_swizzle(index)); @@ -2581,7 +2583,7 @@ string CompilerGLSL::to_expression(uint32_t id, bool register_expression_read) return to_enclosed_expression(e.base_expression) + e.expression; else if (e.need_transpose) { - bool is_packed = has_decoration(id, DecorationCPacked); + bool is_packed = has_extended_decoration(id, SPIRVCrossDecorationPacked); return convert_row_major_matrix(e.expression, get(e.expression_type), is_packed); } else @@ -2776,8 +2778,8 @@ string CompilerGLSL::constant_op_expression(const SPIRConstantOp &cop) case OpCompositeExtract: { - auto expr = - access_chain_internal(cop.arguments[0], &cop.arguments[1], uint32_t(cop.arguments.size() - 1), true, false); + auto expr = access_chain_internal(cop.arguments[0], &cop.arguments[1], uint32_t(cop.arguments.size() - 1), + ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, nullptr); return expr; } @@ -4210,6 +4212,10 @@ void CompilerGLSL::emit_texture_op(const Instruction &i) if (is_legacy() && image_is_comparison(imgtype, img)) expr += ".r"; + // Deals with reads from MSL. We might need to downconvert to fewer components. + if (op == OpImageRead) + expr = remap_swizzle(get(result_type), 4, expr); + emit_op(result_type, id, expr, forward); for (auto &inherit : inherited_expressions) inherit_expression_dependencies(id, inherit); @@ -5482,10 +5488,15 @@ const char *CompilerGLSL::index_to_swizzle(uint32_t index) } string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indices, uint32_t count, - bool index_is_literal, bool chain_only, bool ptr_chain, - AccessChainMeta *meta, bool register_expression_read) + AccessChainFlags flags, AccessChainMeta *meta) { string expr; + + bool index_is_literal = (flags & ACCESS_CHAIN_INDEX_IS_LITERAL_BIT) != 0; + bool chain_only = (flags & ACCESS_CHAIN_CHAIN_ONLY_BIT) != 0; + bool ptr_chain = (flags & ACCESS_CHAIN_PTR_CHAIN_BIT) != 0; + bool register_expression_read = (flags & ACCESS_CHAIN_SKIP_REGISTER_EXPRESSION_READ_BIT) == 0; + if (!chain_only) expr = to_enclosed_expression(base, register_expression_read); @@ -5496,7 +5507,8 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice bool access_chain_is_arrayed = expr.find_first_of('[') != string::npos; bool row_major_matrix_needs_conversion = is_non_native_row_major_matrix(base); - bool is_packed = has_decoration(base, DecorationCPacked); + bool is_packed = has_extended_decoration(base, SPIRVCrossDecorationPacked); + uint32_t packed_type = get_extended_decoration(base, SPIRVCrossDecorationPackedType); bool is_invariant = has_decoration(base, DecorationInvariant); bool pending_array_enclose = false; bool dimension_flatten = false; @@ -5672,6 +5684,11 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice is_invariant = true; is_packed = member_is_packed_type(*type, index); + if (is_packed) + packed_type = get_extended_member_decoration(type->self, index, SPIRVCrossDecorationPackedType); + else + packed_type = 0; + row_major_matrix_needs_conversion = member_is_non_native_row_major_matrix(*type, index); type = &get(type->member_types[index]); } @@ -5683,6 +5700,7 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice expr = convert_row_major_matrix(expr, *type, is_packed); row_major_matrix_needs_conversion = false; is_packed = false; + packed_type = 0; } expr += "["; @@ -5722,6 +5740,7 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice } is_packed = false; + packed_type = 0; type_id = type->parent_type; type = &get(type_id); } @@ -5741,6 +5760,7 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice meta->need_transpose = row_major_matrix_needs_conversion; meta->storage_is_packed = is_packed; meta->storage_is_invariant = is_invariant; + meta->storage_packed_type = packed_type; } return expr; @@ -5772,7 +5792,11 @@ string CompilerGLSL::access_chain(uint32_t base, const uint32_t *indices, uint32 } else if (flattened_structs.count(base) && count > 0) { - auto chain = access_chain_internal(base, indices, count, false, true, ptr_chain, nullptr, false).substr(1); + AccessChainFlags flags = ACCESS_CHAIN_CHAIN_ONLY_BIT | ACCESS_CHAIN_SKIP_REGISTER_EXPRESSION_READ_BIT; + if (ptr_chain) + flags |= ACCESS_CHAIN_PTR_CHAIN_BIT; + + auto chain = access_chain_internal(base, indices, count, flags, nullptr).substr(1); if (meta) { meta->need_transpose = false; @@ -5782,7 +5806,10 @@ string CompilerGLSL::access_chain(uint32_t base, const uint32_t *indices, uint32 } else { - return access_chain_internal(base, indices, count, false, false, ptr_chain, meta, false); + AccessChainFlags flags = ACCESS_CHAIN_SKIP_REGISTER_EXPRESSION_READ_BIT; + if (ptr_chain) + flags |= ACCESS_CHAIN_PTR_CHAIN_BIT; + return access_chain_internal(base, indices, count, flags, meta); } } @@ -6603,6 +6630,30 @@ void CompilerGLSL::handle_store_to_invariant_variable(uint32_t store_id, uint32_ disallow_forwarding_in_expression_chain(*expr); } +void CompilerGLSL::emit_store_statement(uint32_t lhs_expression, uint32_t rhs_expression) +{ + auto rhs = to_pointer_expression(rhs_expression); + + // Statements to OpStore may be empty if it is a struct with zero members. Just forward the store to /dev/null. + if (!rhs.empty()) + { + handle_store_to_invariant_variable(lhs_expression, rhs_expression); + + auto lhs = to_dereferenced_expression(lhs_expression); + + // We might need to bitcast in order to store to a builtin. + bitcast_to_builtin_store(lhs_expression, rhs, expression_type(rhs_expression)); + + // Tries to optimize assignments like " = op expr". + // While this is purely cosmetic, this is important for legacy ESSL where loop + // variable increments must be in either i++ or i += const-expr. + // Without this, we end up with i = i + 1, which is correct GLSL, but not correct GLES 2.0. + if (!optimize_read_modify_write(expression_type(rhs_expression), lhs, rhs)) + statement(lhs, " = ", rhs, ";"); + register_write(lhs_expression); + } +} + void CompilerGLSL::emit_instruction(const Instruction &instruction) { auto ops = stream(instruction); @@ -6674,8 +6725,12 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) register_read(id, ptr, forward); // Pass through whether the result is of a packed type. - if (has_decoration(ptr, DecorationCPacked)) - set_decoration(id, DecorationCPacked); + if (has_extended_decoration(ptr, SPIRVCrossDecorationPacked)) + { + set_extended_decoration(id, SPIRVCrossDecorationPacked); + set_extended_decoration(id, SPIRVCrossDecorationPackedType, + get_extended_decoration(ptr, SPIRVCrossDecorationPackedType)); + } inherit_expression_dependencies(id, ptr); if (forward) @@ -6706,14 +6761,11 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) // Mark the result as being packed. Some platforms handled packed vectors differently than non-packed. if (meta.storage_is_packed) - set_decoration(ops[1], DecorationCPacked); - else - unset_decoration(ops[1], DecorationCPacked); - + set_extended_decoration(ops[1], SPIRVCrossDecorationPacked); + if (meta.storage_packed_type != 0) + set_extended_decoration(ops[1], SPIRVCrossDecorationPackedType, meta.storage_packed_type); if (meta.storage_is_invariant) set_decoration(ops[1], DecorationInvariant); - else - unset_decoration(ops[1], DecorationInvariant); for (uint32_t i = 2; i < length; i++) { @@ -6743,27 +6795,9 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) } else { - auto rhs = to_pointer_expression(ops[1]); - - // Statements to OpStore may be empty if it is a struct with zero members. Just forward the store to /dev/null. - if (!rhs.empty()) - { - handle_store_to_invariant_variable(ops[0], ops[1]); - - auto lhs = to_dereferenced_expression(ops[0]); - - // We might need to bitcast in order to store to a builtin. - bitcast_to_builtin_store(ops[0], rhs, expression_type(ops[1])); - - // Tries to optimize assignments like " = op expr". - // While this is purely cosmetic, this is important for legacy ESSL where loop - // variable increments must be in either i++ or i += const-expr. - // Without this, we end up with i = i + 1, which is correct GLSL, but not correct GLES 2.0. - if (!optimize_read_modify_write(expression_type(ops[1]), lhs, rhs)) - statement(lhs, " = ", rhs, ";"); - register_write(ops[0]); - } + emit_store_statement(ops[0], ops[1]); } + // Storing a pointer results in a variable pointer, so we must conservatively assume // we can write through it. if (expression_type(ops[1]).pointer) @@ -6775,7 +6809,7 @@ 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, true); + 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); break; } @@ -7007,7 +7041,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) // Make a copy, then use access chain to store the variable. statement(declare_temporary(result_type, id), to_expression(vec), ";"); set(id, to_name(id), result_type, true); - auto chain = access_chain_internal(id, &index, 1, false); + auto chain = access_chain_internal(id, &index, 1, 0, nullptr); statement(chain, " = ", to_expression(comp), ";"); break; } @@ -7017,7 +7051,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) uint32_t result_type = ops[0]; uint32_t id = ops[1]; - auto expr = access_chain_internal(ops[2], &ops[3], 1, false); + auto expr = access_chain_internal(ops[2], &ops[3], 1, 0, nullptr); emit_op(result_type, id, expr, should_forward(ops[2])); inherit_expression_dependencies(id, ops[2]); inherit_expression_dependencies(id, ops[3]); @@ -7041,7 +7075,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) allow_base_expression = false; // Packed expressions cannot be split up. - if (has_decoration(ops[2], DecorationCPacked)) + if (has_extended_decoration(ops[2], SPIRVCrossDecorationPacked)) allow_base_expression = false; AccessChainMeta meta; @@ -7062,14 +7096,15 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) // // Including the base will prevent this and would trigger multiple reads // from expression causing it to be forced to an actual temporary in GLSL. - auto expr = access_chain_internal(ops[2], &ops[3], length, true, true, false, &meta); + auto expr = access_chain_internal(ops[2], &ops[3], length, + ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_CHAIN_ONLY_BIT, &meta); e = &emit_op(result_type, id, expr, true, !expression_is_forwarded(ops[2])); inherit_expression_dependencies(id, ops[2]); e->base_expression = ops[2]; } else { - auto expr = access_chain_internal(ops[2], &ops[3], length, true, false, false, &meta); + auto expr = access_chain_internal(ops[2], &ops[3], length, ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, &meta); e = &emit_op(result_type, id, expr, should_forward(ops[2]), !expression_is_forwarded(ops[2])); inherit_expression_dependencies(id, ops[2]); } @@ -7079,7 +7114,9 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) // instead of loading everything through an access chain. e->need_transpose = meta.need_transpose; if (meta.storage_is_packed) - set_decoration(id, DecorationCPacked); + set_extended_decoration(id, SPIRVCrossDecorationPacked); + if (meta.storage_packed_type != 0) + set_extended_decoration(id, SPIRVCrossDecorationPackedType, meta.storage_packed_type); if (meta.storage_is_invariant) set_decoration(id, DecorationInvariant); @@ -7100,7 +7137,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) // Make a copy, then use access chain to store the variable. statement(declare_temporary(result_type, id), to_expression(composite), ";"); set(id, to_name(id), result_type, true); - auto chain = access_chain_internal(id, elems, length, true); + auto chain = access_chain_internal(id, elems, length, ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, nullptr); statement(chain, " = ", to_expression(obj), ";"); break; @@ -7167,7 +7204,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) shuffle = true; // Cannot use swizzles with packed expressions, force shuffle path. - if (!shuffle && has_decoration(vec0, DecorationCPacked)) + if (!shuffle && has_extended_decoration(vec0, SPIRVCrossDecorationPacked)) shuffle = true; string expr; @@ -8796,7 +8833,7 @@ bool CompilerGLSL::member_is_non_native_row_major_matrix(const SPIRType &type, u // GLSL does not define packed data types, but certain subclasses do. bool CompilerGLSL::member_is_packed_type(const SPIRType &type, uint32_t index) const { - return has_member_decoration(type.self, index, DecorationCPacked); + return has_extended_member_decoration(type.self, index, SPIRVCrossDecorationPacked); } // Wraps the expression string in a function call that converts the diff --git a/3rdparty/spirv-cross/spirv_glsl.hpp b/3rdparty/spirv-cross/spirv_glsl.hpp index 6192e7a49..e19228c24 100644 --- a/3rdparty/spirv-cross/spirv_glsl.hpp +++ b/3rdparty/spirv-cross/spirv_glsl.hpp @@ -51,6 +51,15 @@ struct PlsRemap PlsFormat format; }; +enum AccessChainFlagBits +{ + ACCESS_CHAIN_INDEX_IS_LITERAL_BIT = 1 << 0, + ACCESS_CHAIN_CHAIN_ONLY_BIT = 1 << 1, + ACCESS_CHAIN_PTR_CHAIN_BIT = 1 << 2, + ACCESS_CHAIN_SKIP_REGISTER_EXPRESSION_READ_BIT = 1 << 3 +}; +typedef uint32_t AccessChainFlags; + class CompilerGLSL : public Compiler { public: @@ -261,7 +270,7 @@ protected: virtual void emit_buffer_block(const SPIRVariable &type); virtual void emit_push_constant_block(const SPIRVariable &var); virtual void emit_uniform(const SPIRVariable &var); - virtual std::string unpack_expression_type(std::string expr_str, const SPIRType &type); + virtual std::string unpack_expression_type(std::string expr_str, const SPIRType &type, uint32_t packed_type_id); std::unique_ptr buffer; @@ -447,9 +456,10 @@ protected: bool expression_is_forwarded(uint32_t id); SPIRExpression &emit_op(uint32_t result_type, uint32_t result_id, const std::string &rhs, bool forward_rhs, bool suppress_usage_tracking = false); - std::string access_chain_internal(uint32_t base, const uint32_t *indices, uint32_t count, bool index_is_literal, - bool chain_only = false, bool ptr_chain = false, AccessChainMeta *meta = nullptr, - bool register_expression_read = true); + + std::string access_chain_internal(uint32_t base, const uint32_t *indices, uint32_t count, AccessChainFlags flags, + AccessChainMeta *meta); + std::string access_chain(uint32_t base, const uint32_t *indices, uint32_t count, const SPIRType &target_type, AccessChainMeta *meta = nullptr, bool ptr_chain = false); @@ -476,11 +486,11 @@ protected: void append_global_func_args(const SPIRFunction &func, uint32_t index, std::vector &arglist); std::string to_expression(uint32_t id, bool register_expression_read = true); std::string to_enclosed_expression(uint32_t id, bool register_expression_read = true); - std::string to_unpacked_expression(uint32_t id); - std::string to_enclosed_unpacked_expression(uint32_t id); + std::string to_unpacked_expression(uint32_t id, bool register_expression_read = true); + std::string to_enclosed_unpacked_expression(uint32_t id, bool register_expression_read = true); std::string to_dereferenced_expression(uint32_t id, bool register_expression_read = true); - std::string to_pointer_expression(uint32_t id); - std::string to_enclosed_pointer_expression(uint32_t id); + std::string to_pointer_expression(uint32_t id, bool register_expression_read = true); + std::string to_enclosed_pointer_expression(uint32_t id, bool register_expression_read = true); std::string to_extract_component_expression(uint32_t id, uint32_t index); std::string enclose_expression(const std::string &expr); std::string dereference_expression(const std::string &expr); @@ -624,6 +634,7 @@ protected: void disallow_forwarding_in_expression_chain(const SPIRExpression &expr); bool expression_is_constant_null(uint32_t id) const; + virtual void emit_store_statement(uint32_t lhs_expression, uint32_t rhs_expression); private: void init() diff --git a/3rdparty/spirv-cross/spirv_hlsl.cpp b/3rdparty/spirv-cross/spirv_hlsl.cpp index dc2e03087..4b470a1be 100644 --- a/3rdparty/spirv-cross/spirv_hlsl.cpp +++ b/3rdparty/spirv-cross/spirv_hlsl.cpp @@ -1835,7 +1835,7 @@ void CompilerHLSL::emit_struct_member(const SPIRType &type, uint32_t member_type string packing_offset; bool is_push_constant = type.storage == StorageClassPushConstant; - if ((has_decoration(type.self, DecorationCPacked) || is_push_constant) && + if ((has_extended_decoration(type.self, SPIRVCrossDecorationPacked) || is_push_constant) && has_member_decoration(type.self, index, DecorationOffset)) { uint32_t offset = memb[index].offset - base_offset; @@ -1870,7 +1870,7 @@ void CompilerHLSL::emit_buffer_block(const SPIRVariable &var) if (type.array.empty()) { if (buffer_is_packing_standard(type, BufferPackingHLSLCbufferPackOffset)) - set_decoration(type.self, DecorationCPacked); + set_extended_decoration(type.self, SPIRVCrossDecorationPacked); else SPIRV_CROSS_THROW("cbuffer cannot be expressed with either HLSL packing layout or packoffset."); @@ -1952,7 +1952,7 @@ void CompilerHLSL::emit_push_constant_block(const SPIRVariable &var) auto &type = get(var.basetype); if (buffer_is_packing_standard(type, BufferPackingHLSLCbufferPackOffset, layout.start, layout.end)) - set_decoration(type.self, DecorationCPacked); + set_extended_decoration(type.self, SPIRVCrossDecorationPacked); else SPIRV_CROSS_THROW( "root constant cbuffer cannot be expressed with either HLSL packing layout or packoffset."); diff --git a/3rdparty/spirv-cross/spirv_msl.cpp b/3rdparty/spirv-cross/spirv_msl.cpp index 4f81ae388..87b00ea1f 100644 --- a/3rdparty/spirv-cross/spirv_msl.cpp +++ b/3rdparty/spirv-cross/spirv_msl.cpp @@ -745,7 +745,7 @@ void CompilerMSL::mark_packable_structs() } // If the specified type is a struct, it and any nested structs -// are marked as packable with the DecorationCPacked decoration, +// are marked as packable with the SPIRVCrossDecorationPacked decoration, void CompilerMSL::mark_as_packable(SPIRType &type) { // If this is not the base type (eg. it's a pointer or array), tunnel down @@ -757,7 +757,7 @@ void CompilerMSL::mark_as_packable(SPIRType &type) if (type.basetype == SPIRType::Struct) { - set_decoration(type.self, DecorationCPacked); + set_extended_decoration(type.self, SPIRVCrossDecorationPacked); // Recurse size_t mbr_cnt = type.member_types.size(); @@ -1479,6 +1479,7 @@ uint32_t CompilerMSL::ensure_correct_attribute_type(uint32_t type_id, uint32_t l auto &base_type = set(base_type_id); base_type = type; base_type.basetype = type.basetype == SPIRType::Short ? SPIRType::UShort : SPIRType::UInt; + base_type.pointer = false; if (!type.pointer) return base_type_id; @@ -1508,6 +1509,7 @@ uint32_t CompilerMSL::ensure_correct_attribute_type(uint32_t type_id, uint32_t l auto &base_type = set(base_type_id); base_type = type; base_type.basetype = SPIRType::UInt; + base_type.pointer = false; if (!type.pointer) return base_type_id; @@ -1551,7 +1553,11 @@ void CompilerMSL::align_struct(SPIRType &ib_type) for (uint32_t mbr_idx = 0; mbr_idx < mbr_cnt; mbr_idx++) { if (is_member_packable(ib_type, mbr_idx)) - set_member_decoration(ib_type_id, mbr_idx, DecorationCPacked); + { + set_extended_member_decoration(ib_type_id, mbr_idx, SPIRVCrossDecorationPacked); + set_extended_member_decoration(ib_type_id, mbr_idx, SPIRVCrossDecorationPackedType, + ib_type.member_types[mbr_idx]); + } // Align current offset to the current member's default alignment. size_t align_mask = get_declared_struct_member_alignment(ib_type, mbr_idx) - 1; @@ -1580,11 +1586,30 @@ void CompilerMSL::align_struct(SPIRType &ib_type) bool CompilerMSL::is_member_packable(SPIRType &ib_type, uint32_t index) { // We've already marked it as packable - if (has_member_decoration(ib_type.self, index, DecorationCPacked)) + if (has_extended_member_decoration(ib_type.self, index, SPIRVCrossDecorationPacked)) return true; auto &mbr_type = get(ib_type.member_types[index]); + uint32_t component_size = mbr_type.width / 8; + uint32_t unpacked_mbr_size; + if (mbr_type.vecsize == 3) + unpacked_mbr_size = component_size * (mbr_type.vecsize + 1) * mbr_type.columns; + else + unpacked_mbr_size = component_size * mbr_type.vecsize * mbr_type.columns; + + // Special case for packing. Check for float[] or vec2[] in std140 layout. Here we actually need to pad out instead, + // but we will use the same mechanism. + if (is_array(mbr_type) && (is_scalar(mbr_type) || is_vector(mbr_type)) && mbr_type.vecsize <= 2 && + type_struct_member_array_stride(ib_type, index) == 4 * component_size) + { + return true; + } + + // TODO: Another sanity check for matrices. We currently do not support std140 matrices which need to be padded out per column. + //if (is_matrix(mbr_type) && mbr_type.vecsize <= 2 && type_struct_member_matrix_stride(ib_type, index) == 16) + // SPIRV_CROSS_THROW("Currently cannot support matrices with small vector size in std140 layout."); + // Only vectors or 3-row matrices need to be packed. if (mbr_type.vecsize == 1 || (is_matrix(mbr_type) && mbr_type.vecsize != 3)) return false; @@ -1593,12 +1618,6 @@ bool CompilerMSL::is_member_packable(SPIRType &ib_type, uint32_t index) if (is_matrix(mbr_type) && !has_member_decoration(ib_type.self, index, DecorationRowMajor)) return false; - uint32_t component_size = mbr_type.width / 8; - uint32_t unpacked_mbr_size; - if (mbr_type.vecsize == 3) - unpacked_mbr_size = component_size * (mbr_type.vecsize + 1) * mbr_type.columns; - else - unpacked_mbr_size = component_size * mbr_type.vecsize * mbr_type.columns; if (is_array(mbr_type)) { // If member is an array, and the array stride is larger than the type needs, don't pack it. @@ -1640,11 +1659,49 @@ MSLStructMemberKey CompilerMSL::get_struct_member_key(uint32_t type_id, uint32_t return k; } +void CompilerMSL::emit_store_statement(uint32_t lhs_expression, uint32_t rhs_expression) +{ + if (!has_extended_decoration(lhs_expression, SPIRVCrossDecorationPacked) || + get_extended_decoration(lhs_expression, SPIRVCrossDecorationPackedType) == 0) + { + CompilerGLSL::emit_store_statement(lhs_expression, rhs_expression); + } + else + { + // Special handling when storing to a float[] or float2[] in std140 layout. + + auto &type = get(get_extended_decoration(lhs_expression, SPIRVCrossDecorationPackedType)); + string lhs = to_dereferenced_expression(lhs_expression); + string rhs = to_pointer_expression(rhs_expression); + + // Unpack the expression so we can store to it with a float or float2. + // It's still an l-value, so it's fine. Most other unpacking of expressions turn them into r-values instead. + if (is_scalar(type) && is_array(type)) + lhs = enclose_expression(lhs) + ".x"; + else if (is_vector(type) && type.vecsize == 2 && is_array(type)) + lhs = enclose_expression(lhs) + ".xy"; + + if (!optimize_read_modify_write(expression_type(rhs_expression), lhs, rhs)) + statement(lhs, " = ", rhs, ";"); + register_write(lhs_expression); + } +} + // Converts the format of the current expression from packed to unpacked, // by wrapping the expression in a constructor of the appropriate type. -string CompilerMSL::unpack_expression_type(string expr_str, const SPIRType &type) +string CompilerMSL::unpack_expression_type(string expr_str, const SPIRType &type, uint32_t packed_type_id) { - return join(type_to_glsl(type), "(", expr_str, ")"); + const SPIRType *packed_type = nullptr; + if (packed_type_id) + packed_type = &get(packed_type_id); + + // float[] and float2[] cases are really just padding, so directly swizzle from the backing float4 instead. + if (packed_type && is_array(*packed_type) && is_scalar(*packed_type)) + return enclose_expression(expr_str) + ".x"; + else if (packed_type && is_array(*packed_type) && is_vector(*packed_type) && packed_type->vecsize == 2) + return enclose_expression(expr_str) + ".xy"; + else + return join(type_to_glsl(type), "(", expr_str, ")"); } // Emits the file header info @@ -2109,8 +2166,9 @@ void CompilerMSL::emit_custom_functions() statement(""); statement("// Wrapper function that swizzles texture gathers."); statement("template"); - statement("inline vec spvGatherSwizzle(sampler s, const thread Tex& t, Ts... params, component c, uint sw) " - "METAL_CONST_ARG(c)"); + statement( + "inline vec spvGatherSwizzle(sampler s, const thread Tex& t, Ts... params, component c, uint sw) " + "METAL_CONST_ARG(c)"); begin_scope(); statement("if (sw)"); begin_scope(); @@ -2149,7 +2207,8 @@ void CompilerMSL::emit_custom_functions() statement(""); statement("// Wrapper function that swizzles depth texture gathers."); statement("template"); - statement("inline vec spvGatherCompareSwizzle(sampler s, const thread Tex& t, Ts... params, uint sw) "); + statement( + "inline vec spvGatherCompareSwizzle(sampler s, const thread Tex& t, Ts... params, uint sw) "); begin_scope(); statement("if (sw)"); begin_scope(); @@ -2330,7 +2389,7 @@ void CompilerMSL::emit_specialization_constants_and_structs() declared_structs.insert(type_id); - if (has_decoration(type_id, DecorationCPacked)) + if (has_extended_decoration(type_id, SPIRVCrossDecorationPacked)) align_struct(type); // Make sure we declare the underlying struct type, and not the "decorated" type with pointers, etc. @@ -2652,8 +2711,13 @@ void CompilerMSL::emit_instruction(const Instruction &instruction) test(bias, ImageOperandsBiasMask); test(lod, ImageOperandsLodMask); + auto &texel_type = expression_type(texel_id); + auto store_type = texel_type; + store_type.vecsize = 4; + statement(join( - to_expression(img_id), ".write(", to_expression(texel_id), ", ", + to_expression(img_id), ".write(", remap_swizzle(store_type, texel_type.vecsize, to_expression(texel_id)), + ", ", to_function_args(img_id, img_type, true, false, false, coord_id, 0, 0, 0, 0, lod, 0, 0, 0, 0, 0, &forward), ");")); @@ -2813,7 +2877,7 @@ void CompilerMSL::emit_instruction(const Instruction &instruction) uint32_t mtx_id = ops[opcode == OpMatrixTimesVector ? 2 : 3]; auto *e = maybe_get(mtx_id); auto &t = expression_type(mtx_id); - bool is_packed = has_decoration(mtx_id, DecorationCPacked); + bool is_packed = has_extended_decoration(mtx_id, SPIRVCrossDecorationPacked); if (e && e->need_transpose && (t.columns == t.vecsize || is_packed)) { e->need_transpose = false; @@ -2822,11 +2886,11 @@ void CompilerMSL::emit_instruction(const Instruction &instruction) // are generally transposed, so unpacking using a constructor argument // will result in an error. // The simplest solution for now is to just avoid unpacking the matrix in this operation. - unset_decoration(mtx_id, DecorationCPacked); + unset_extended_decoration(mtx_id, SPIRVCrossDecorationPacked); emit_binary_op(ops[0], ops[1], ops[3], ops[2], "*"); if (is_packed) - set_decoration(mtx_id, DecorationCPacked); + set_extended_decoration(mtx_id, SPIRVCrossDecorationPacked); e->need_transpose = true; } else @@ -3859,7 +3923,7 @@ bool CompilerMSL::is_non_native_row_major_matrix(uint32_t id) // Generate a function that will swap matrix elements from row-major to column-major. // Packed row-matrix should just use transpose() function. - if (!has_decoration(id, DecorationCPacked)) + if (!has_extended_decoration(id, SPIRVCrossDecorationPacked)) { const auto type = expression_type(id); add_convert_row_major_matrix_function(type.columns, type.vecsize); @@ -3881,7 +3945,7 @@ bool CompilerMSL::member_is_non_native_row_major_matrix(const SPIRType &type, ui // Generate a function that will swap matrix elements from row-major to column-major. // Packed row-matrix should just use transpose() function. - if (!has_member_decoration(type.self, index, DecorationCPacked)) + if (!has_extended_member_decoration(type.self, index, SPIRVCrossDecorationPacked)) { const auto mbr_type = get(type.member_types[index]); add_convert_row_major_matrix_function(mbr_type.columns, mbr_type.vecsize); @@ -3964,13 +4028,16 @@ void CompilerMSL::emit_struct_member(const SPIRType &type, uint32_t member_type_ // If this member is packed, mark it as so. string pack_pfx = ""; + + const SPIRType *effective_membertype = &membertype; + SPIRType override_type; + if (member_is_packed_type(type, index)) { - pack_pfx = "packed_"; - // If we're packing a matrix, output an appropriate typedef if (membertype.vecsize > 1 && membertype.columns > 1) { + pack_pfx = "packed_"; string base_type = membertype.width == 16 ? "half" : "float"; string td_line = "typedef "; td_line += base_type + to_string(membertype.vecsize) + "x" + to_string(membertype.columns); @@ -3979,10 +4046,19 @@ void CompilerMSL::emit_struct_member(const SPIRType &type, uint32_t member_type_ td_line += ";"; add_typedef_line(td_line); } + else if (is_array(membertype) && membertype.vecsize <= 2 && membertype.basetype != SPIRType::Struct) + { + // A "packed" float array, but we pad here instead to 4-vector. + override_type = membertype; + override_type.vecsize = 4; + effective_membertype = &override_type; + } + else + pack_pfx = "packed_"; } - statement(pack_pfx, type_to_glsl(membertype), " ", qualifier, to_member_name(type, index), - member_attribute_qualifier(type, index), type_to_array_glsl(membertype), ";"); + statement(pack_pfx, type_to_glsl(*effective_membertype), " ", qualifier, to_member_name(type, index), + member_attribute_qualifier(type, index), type_to_array_glsl(*effective_membertype), ";"); } // Return a MSL qualifier for the specified function attribute member @@ -5576,7 +5652,7 @@ size_t CompilerMSL::get_declared_struct_member_size(const SPIRType &struct_type, uint32_t columns = type.columns; // An unpacked 3-element vector or matrix column is the same memory size as a 4-element. - if (vecsize == 3 && !has_member_decoration(struct_type.self, index, DecorationCPacked)) + if (vecsize == 3 && !has_extended_member_decoration(struct_type.self, index, SPIRVCrossDecorationPacked)) vecsize = 4; return component_size * vecsize * columns; diff --git a/3rdparty/spirv-cross/spirv_msl.hpp b/3rdparty/spirv-cross/spirv_msl.hpp index f4520ddc0..8fcb224b0 100644 --- a/3rdparty/spirv-cross/spirv_msl.hpp +++ b/3rdparty/spirv-cross/spirv_msl.hpp @@ -346,7 +346,7 @@ protected: uint32_t grad_y, uint32_t lod, uint32_t coffset, uint32_t offset, uint32_t bias, uint32_t comp, uint32_t sample, bool *p_forward) override; std::string to_initializer_expression(const SPIRVariable &var) override; - std::string unpack_expression_type(std::string expr_str, const SPIRType &type) override; + std::string unpack_expression_type(std::string expr_str, const SPIRType &type, uint32_t packed_type_id) override; std::string bitcast_glsl_op(const SPIRType &result_type, const SPIRType &argument_type) override; bool skip_argument(uint32_t id) const override; std::string to_member_reference(uint32_t base, const SPIRType &type, uint32_t index, bool ptr_chain) override; @@ -430,6 +430,7 @@ protected: 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; + void emit_store_statement(uint32_t lhs_expression, uint32_t rhs_expression) override; void analyze_sampled_image_usage();