MSL: Fix struct declaration order with complex type aliases.

MSL generally emits the aliases, which means we cannot always place the
master type first, unlike GLSL and HLSL. The logic fix is just to
reorder after we have tagged types with packing information, rather than
doing it in the parser fixup.

reference/opt/shaders-msl/comp/composite-array-initialization.comp

@@ -11,23 +11,23 @@ struct Data
     float b;
 };
 
+constant float X_tmp [[function_constant(0)]];
+constant float X = is_function_constant_defined(X_tmp) ? X_tmp : 4.0;
+
 struct Data_1
 {
     float a;
     float b;
 };
 
-constant float X_tmp [[function_constant(0)]];
-constant float X = is_function_constant_defined(X_tmp) ? X_tmp : 4.0;
-
 struct SSBO
 {
-    Data outdata[1];
+    Data_1 outdata[1];
 };
 
 constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(2u, 1u, 1u);
 
-constant Data_1 _25[2] = { Data_1{ 1.0, 2.0 }, Data_1{ 3.0, 4.0 } };
+constant Data _25[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } };
 
 // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
 template<typename T, uint N>
@@ -44,8 +44,8 @@ void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])
 
 kernel void main0(device SSBO& _53 [[buffer(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
 {
-    Data_1 _31[2] = { Data_1{ X, 2.0 }, Data_1{ 3.0, 5.0 } };
-    Data_1 data2[2];
+    Data _31[2] = { Data{ X, 2.0 }, Data{ 3.0, 5.0 } };
+    Data data2[2];
     spvArrayCopyFromStack1(data2, _31);
     _53.outdata[gl_WorkGroupID.x].a = _25[gl_LocalInvocationID.x].a + data2[gl_LocalInvocationID.x].a;
     _53.outdata[gl_WorkGroupID.x].b = _25[gl_LocalInvocationID.x].b + data2[gl_LocalInvocationID.x].b;

reference/shaders-msl/asm/frag/extract-packed-from-composite.asm.frag

@@ -7,19 +7,19 @@ using namespace metal;
 
 struct Foo
 {
-    packed_float3 a;
+    float3 a;
     float b;
 };
 
 struct Foo_1
 {
-    float3 a;
+    packed_float3 a;
     float b;
 };
 
 struct buf
 {
-    Foo results[16];
+    Foo_1 results[16];
     float4 bar;
 };
 
@@ -31,7 +31,7 @@ struct main0_out
 float4 _main(thread const float4& pos, constant buf& v_11)
 {
     int _46 = int(pos.x) % 16;
-    Foo_1 foo;
+    Foo foo;
     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);

reference/shaders-msl/asm/vert/extract-transposed-matrix-from-struct.asm.vert

@@ -25,7 +25,7 @@ struct InstanceData_1
 
 struct gInstanceData
 {
-    InstanceData _data[1];
+    InstanceData_1 _data[1];
 };
 
 struct main0_out
@@ -41,7 +41,7 @@ struct main0_in
 
 V2F _VS(thread const float3& PosL, thread const uint& instanceID, const device gInstanceData& gInstanceData_1)
 {
-    InstanceData_1 instData;
+    InstanceData instData;
     instData.MATRIX_MVP = transpose(gInstanceData_1._data[instanceID].MATRIX_MVP);
     instData.Color = gInstanceData_1._data[instanceID].Color;
     V2F v2f;

reference/shaders-msl/comp/composite-array-initialization.comp

@@ -11,23 +11,23 @@ struct Data
     float b;
 };
 
+constant float X_tmp [[function_constant(0)]];
+constant float X = is_function_constant_defined(X_tmp) ? X_tmp : 4.0;
+
 struct Data_1
 {
     float a;
     float b;
 };
 
-constant float X_tmp [[function_constant(0)]];
-constant float X = is_function_constant_defined(X_tmp) ? X_tmp : 4.0;
-
 struct SSBO
 {
-    Data outdata[1];
+    Data_1 outdata[1];
 };
 
 constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(2u, 1u, 1u);
 
-constant Data_1 _25[2] = { Data_1{ 1.0, 2.0 }, Data_1{ 3.0, 4.0 } };
+constant Data _25[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } };
 
 // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
 template<typename T, uint N>
@@ -42,20 +42,20 @@ void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
 
-Data_1 combine(thread const Data_1& a, thread const Data_1& b)
+Data combine(thread const Data& a, thread const Data& b)
 {
-    return Data_1{ a.a + b.a, a.b + b.b };
+    return Data{ a.a + b.a, a.b + b.b };
 }
 
 kernel void main0(device SSBO& _53 [[buffer(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
 {
-    Data_1 data[2] = { Data_1{ 1.0, 2.0 }, Data_1{ 3.0, 4.0 } };
-    Data_1 _31[2] = { Data_1{ X, 2.0 }, Data_1{ 3.0, 5.0 } };
-    Data_1 data2[2];
+    Data data[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } };
+    Data _31[2] = { Data{ X, 2.0 }, Data{ 3.0, 5.0 } };
+    Data data2[2];
     spvArrayCopyFromStack1(data2, _31);
-    Data_1 param = data[gl_LocalInvocationID.x];
-    Data_1 param_1 = data2[gl_LocalInvocationID.x];
-    Data_1 _73 = combine(param, param_1);
+    Data param = data[gl_LocalInvocationID.x];
+    Data param_1 = data2[gl_LocalInvocationID.x];
+    Data _73 = combine(param, param_1);
     _53.outdata[gl_WorkGroupID.x].a = _73.a;
     _53.outdata[gl_WorkGroupID.x].b = _73.b;
 }

reference/shaders-msl/comp/packing-test-1.comp

@@ -5,19 +5,19 @@ using namespace metal;
 
 struct T1
 {
-    packed_float3 a;
+    float3 a;
     float b;
 };
 
 struct T1_1
 {
-    float3 a;
+    packed_float3 a;
     float b;
 };
 
 struct Buffer0
 {
-    T1 buf0[1];
+    T1_1 buf0[1];
 };
 
 struct Buffer1
@@ -29,7 +29,7 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = 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;
+    T1 v;
     v.a = float3(_15.buf0[0].a);
     v.b = _15.buf0[0].b;
     float x = v.b;

reference/shaders-msl/comp/storage-buffer-std140-vector-array.comp

@@ -4,14 +4,6 @@
 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];
@@ -19,14 +11,22 @@ struct Sub_1
     float4 f4[2];
 };
 
+struct Sub_1
+{
+    float4 f[2];
+    float4 f2[2];
+    float3 f3[2];
+    float4 f4[2];
+};
+
 struct SSBO
 {
-    Sub sub[2];
+    Sub_1 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;
+    Sub 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;

reference/shaders-msl/comp/struct-nested.comp

@@ -13,24 +13,24 @@ struct s2
     s1 b;
 };
 
-struct s2_1
-{
-    s1 b;
-};
-
 struct s1_1
 {
     int a;
 };
 
+struct s2_1
+{
+    s1_1 b;
+};
+
 struct dstbuffer
 {
-    s2 test[1];
+    s2_1 test[1];
 };
 
 kernel void main0(device dstbuffer& _19 [[buffer(1)]])
 {
-    s2_1 testVal;
+    s2 testVal;
     testVal.b.a = 0;
     _19.test[0].b.a = testVal.b.a;
 }

reference/shaders-msl/comp/type-alias.comp

@@ -10,14 +10,19 @@ struct S0
     float4 a;
 };
 
+struct S1
+{
+    float4 a;
+};
+
 struct S0_1
 {
     float4 a;
 };
 
-struct S1
+struct SSBO0
 {
-    float4 a;
+    S0_1 s0s[1];
 };
 
 struct S1_1
@@ -25,14 +30,9 @@ struct S1_1
     float4 a;
 };
 
-struct SSBO0
-{
-    S0 s0s[1];
-};
-
 struct SSBO1
 {
-    S1 s1s[1];
+    S1_1 s1s[1];
 };
 
 struct SSBO2
@@ -40,24 +40,24 @@ struct SSBO2
     float4 outputs[1];
 };
 
-float4 overload(thread const S0_1& s0)
+float4 overload(thread const S0& s0)
 {
     return s0.a;
 }
 
-float4 overload(thread const S1_1& s1)
+float4 overload(thread const S1& s1)
 {
     return s1.a;
 }
 
 kernel void main0(device SSBO0& _36 [[buffer(0)]], device SSBO1& _55 [[buffer(1)]], device SSBO2& _66 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
-    S0_1 s0;
+    S0 s0;
     s0.a = _36.s0s[gl_GlobalInvocationID.x].a;
-    S1_1 s1;
+    S1 s1;
     s1.a = _55.s1s[gl_GlobalInvocationID.x].a;
-    S0_1 param = s0;
-    S1_1 param_1 = s1;
+    S0 param = s0;
+    S1 param_1 = s1;
     _66.outputs[gl_GlobalInvocationID.x] = overload(param) + overload(param_1);
 }
 

reference/shaders-msl/frag/packing-test-3.frag

@@ -12,19 +12,19 @@ struct VertexOutput
 
 struct TestStruct
 {
-    packed_float3 position;
+    float3 position;
     float radius;
 };
 
 struct TestStruct_1
 {
-    float3 position;
+    packed_float3 position;
     float radius;
 };
 
 struct CB0
 {
-    TestStruct CB0[16];
+    TestStruct_1 CB0[16];
 };
 
 struct main0_out
@@ -34,7 +34,7 @@ struct main0_out
 
 float4 _main(thread const VertexOutput& IN, constant CB0& v_26)
 {
-    TestStruct_1 st;
+    TestStruct st;
     st.position = float3(v_26.CB0[1].position);
     st.radius = v_26.CB0[1].radius;
     float4 col = float4(st.position, st.radius);

spirv_cpp.cpp

@@ -321,6 +321,8 @@ string CompilerCPP::compile()
 	backend.explicit_struct_type = true;
 	backend.use_initializer_list = true;
 
+	fixup_type_alias();
+	reorder_type_alias();
 	build_function_control_flow_graphs_and_analyze();
 	update_active_builtins();
 

spirv_cross.cpp

@@ -872,65 +872,6 @@ bool Compiler::type_is_block_like(const SPIRType &type) const
 	return false;
 }
 
-void Compiler::fixup_type_alias()
-{
-	// Due to how some backends work, the "master" type of type_alias must be a block-like type if it exists.
-	// FIXME: Multiple alias types which are both block-like will be awkward, for now, it's best to just drop the type
-	// alias if the slave type is a block type.
-	ir.for_each_typed_id<SPIRType>([&](uint32_t self, SPIRType &type) {
-		if (type.type_alias && type_is_block_like(type))
-		{
-			// Become the master.
-			ir.for_each_typed_id<SPIRType>([&](uint32_t other_id, SPIRType &other_type) {
-				if (other_id == type.self)
-					return;
-
-				if (other_type.type_alias == type.type_alias)
-					other_type.type_alias = type.self;
-			});
-
-			this->get<SPIRType>(type.type_alias).type_alias = self;
-			type.type_alias = 0;
-		}
-	});
-
-	ir.for_each_typed_id<SPIRType>([&](uint32_t, SPIRType &type) {
-		if (type.type_alias && type_is_block_like(type))
-		{
-			// This is not allowed, drop the type_alias.
-			type.type_alias = 0;
-		}
-	});
-
-	// Reorder declaration of types so that the master of the type alias is always emitted first.
-	// We need this in case a type B depends on type A (A must come before in the vector), but A is an alias of a type Abuffer, which
-	// means declaration of A doesn't happen (yet), and order would be B, ABuffer and not ABuffer, B. Fix this up here.
-	auto &type_ids = ir.ids_for_type[TypeType];
-	for (auto alias_itr = begin(type_ids); alias_itr != end(type_ids); ++alias_itr)
-	{
-		auto &type = get<SPIRType>(*alias_itr);
-		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);
-			assert(master_itr != end(type_ids));
-
-			if (alias_itr < master_itr)
-			{
-				// Must also swap the type order for the constant-type joined array.
-				auto &joined_types = ir.ids_for_constant_or_type;
-				auto alt_alias_itr = find(begin(joined_types), end(joined_types), *alias_itr);
-				auto alt_master_itr = find(begin(joined_types), end(joined_types), *master_itr);
-				assert(alt_alias_itr != end(joined_types));
-				assert(alt_master_itr != end(joined_types));
-
-				swap(*alias_itr, *master_itr);
-				swap(*alt_alias_itr, *alt_master_itr);
-			}
-		}
-	}
-}
-
 void Compiler::parse_fixup()
 {
 	// Figure out specialization constants for work group sizes.
@@ -964,8 +905,6 @@ void Compiler::parse_fixup()
 				aliased_variables.push_back(var.self);
 		}
 	}
-
-	fixup_type_alias();
 }
 
 void Compiler::update_name_cache(unordered_set<string> &cache_primary, const unordered_set<string> &cache_secondary,

spirv_cross.hpp

@@ -972,15 +972,13 @@ protected:
 	void unset_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration);
 
 	bool type_is_array_of_pointers(const SPIRType &type) const;
+	bool type_is_block_like(const SPIRType &type) const;
+	bool type_is_opaque_value(const SPIRType &type) const;
 
 private:
 	// Used only to implement the old deprecated get_entry_point() interface.
 	const SPIREntryPoint &get_first_entry_point(const std::string &name) const;
 	SPIREntryPoint &get_first_entry_point(const std::string &name);
-
-	void fixup_type_alias();
-	bool type_is_block_like(const SPIRType &type) const;
-	bool type_is_opaque_value(const SPIRType &type) const;
 };
 } // namespace SPIRV_CROSS_NAMESPACE
 

spirv_glsl.cpp

@@ -494,6 +494,8 @@ string CompilerGLSL::compile()
 	backend.supports_extensions = true;
 
 	// Scan the SPIR-V to find trivial uses of extensions.
+	fixup_type_alias();
+	reorder_type_alias();
 	build_function_control_flow_graphs_and_analyze();
 	find_static_extensions();
 	fixup_image_load_store_access();
@@ -11898,3 +11900,66 @@ void CompilerGLSL::reset_name_caches()
 	block_names.clear();
 	function_overloads.clear();
 }
+
+void CompilerGLSL::fixup_type_alias()
+{
+	// Due to how some backends work, the "master" type of type_alias must be a block-like type if it exists.
+	// FIXME: Multiple alias types which are both block-like will be awkward, for now, it's best to just drop the type
+	// alias if the slave type is a block type.
+	ir.for_each_typed_id<SPIRType>([&](uint32_t self, SPIRType &type) {
+		if (type.type_alias && type_is_block_like(type))
+		{
+			// Become the master.
+			ir.for_each_typed_id<SPIRType>([&](uint32_t other_id, SPIRType &other_type) {
+				if (other_id == type.self)
+					return;
+
+				if (other_type.type_alias == type.type_alias)
+					other_type.type_alias = type.self;
+			});
+
+			this->get<SPIRType>(type.type_alias).type_alias = self;
+			type.type_alias = 0;
+		}
+	});
+
+	ir.for_each_typed_id<SPIRType>([&](uint32_t, SPIRType &type) {
+		if (type.type_alias && type_is_block_like(type))
+		{
+			// This is not allowed, drop the type_alias.
+			type.type_alias = 0;
+		}
+	});
+}
+
+void CompilerGLSL::reorder_type_alias()
+{
+	// Reorder declaration of types so that the master of the type alias is always emitted first.
+	// We need this in case a type B depends on type A (A must come before in the vector), but A is an alias of a type Abuffer, which
+	// means declaration of A doesn't happen (yet), and order would be B, ABuffer and not ABuffer, B. Fix this up here.
+	auto &type_ids = ir.ids_for_type[TypeType];
+	for (auto alias_itr = begin(type_ids); alias_itr != end(type_ids); ++alias_itr)
+	{
+		auto &type = get<SPIRType>(*alias_itr);
+		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);
+			assert(master_itr != end(type_ids));
+
+			if (alias_itr < master_itr)
+			{
+				// Must also swap the type order for the constant-type joined array.
+				auto &joined_types = ir.ids_for_constant_or_type;
+				auto alt_alias_itr = find(begin(joined_types), end(joined_types), *alias_itr);
+				auto alt_master_itr = find(begin(joined_types), end(joined_types), *master_itr);
+				assert(alt_alias_itr != end(joined_types));
+				assert(alt_master_itr != end(joined_types));
+
+				swap(*alias_itr, *master_itr);
+				swap(*alt_alias_itr, *alt_master_itr);
+			}
+		}
+	}
+}
+

spirv_glsl.hpp

@@ -661,6 +661,9 @@ protected:
 
 	char current_locale_radix_character = '.';
 
+	void fixup_type_alias();
+	void reorder_type_alias();
+
 private:
 	void init();
 };

spirv_hlsl.cpp

@@ -4696,6 +4696,8 @@ string CompilerHLSL::compile()
 	backend.can_return_array = false;
 	backend.nonuniform_qualifier = "NonUniformResourceIndex";
 
+	fixup_type_alias();
+	reorder_type_alias();
 	build_function_control_flow_graphs_and_analyze();
 	validate_shader_model();
 	update_active_builtins();

spirv_msl.cpp

@@ -650,6 +650,7 @@ string CompilerMSL::compile()
 	capture_output_to_buffer = msl_options.capture_output_to_buffer;
 	is_rasterization_disabled = msl_options.disable_rasterization || capture_output_to_buffer;
 
+	fixup_type_alias();
 	replace_illegal_names();
 
 	struct_member_padding.clear();
@@ -691,6 +692,7 @@ string CompilerMSL::compile()
 
 	// Mark any non-stage-in structs to be tightly packed.
 	mark_packable_structs();
+	reorder_type_alias();
 
 	// Add fixup hooks required by shader inputs and outputs. This needs to happen before
 	// the loop, so the hooks aren't added multiple times.

spirv_reflect.cpp

@@ -256,6 +256,8 @@ string CompilerReflection::compile()
 	json_stream = std::make_shared<simple_json::Stream>();
 	json_stream->set_current_locale_radix_character(current_locale_radix_character);
 	json_stream->begin_json_object();
+	fixup_type_alias();
+	reorder_type_alias();
 	emit_entry_points();
 	emit_types();
 	emit_resources();