Merge pull request #842 from KhronosGroup/fix-838

MSL: Use correct alignment for structs which are members of other structs.

reference/opt/shaders-msl/asm/comp/block-name-alias-global.asm.comp

@@ -22,7 +22,7 @@ struct A_2
 
 struct A_3
 {
-    A_2 Data[1024];
+    /* FIXME: A padded struct is needed here. If you see this message, file a bug! */ A_2 Data[1024];
 };
 
 struct B
@@ -32,7 +32,7 @@ struct B
 
 struct B_1
 {
-    A_2 Data[1024];
+    /* FIXME: A padded struct is needed here. If you see this message, file a bug! */ A_2 Data[1024];
 };
 
 kernel void main0(device B& C3 [[buffer(0)]], device A_1& C1 [[buffer(1)]], constant A_3& C2 [[buffer(2)]], constant B_1& C4 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])

reference/opt/shaders-msl/asm/frag/vector-shuffle-oom.asm.frag

@@ -73,7 +73,7 @@ struct _18
     float2 _m24;
     float2 _m25;
     float2 _m26;
-    char pad27[8];
+    char _m27_pad[8];
     packed_float3 _m27;
     float _m28;
     float _m29;

reference/opt/shaders-msl/comp/struct-packing.comp

@@ -43,7 +43,6 @@ struct Content
     S1 m1;
     S2 m2;
     S3 m3;
-    char pad7[4];
     float m4;
     S4 m3s[8];
 };
@@ -53,7 +52,6 @@ struct SSBO1
     Content content;
     Content content1[2];
     Content content2;
-    char pad3[8];
     float2x2 m0;
     float2x2 m1;
     float2x3 m2[4];
@@ -61,9 +59,9 @@ struct SSBO1
     float2x2 m4;
     float2x2 m5[9];
     packed_float2x3 m6[4][2];
-    char pad10[8];
+    char _m10_pad[8];
     float3x2 m7;
-    char pad11[8];
+    char _m11_pad[8];
     float array[1];
 };
 
@@ -105,9 +103,9 @@ struct Content_1
     S1_1 m1;
     S2_1 m2;
     S3_1 m3;
-    char pad7[4];
     float m4;
-    S4_1 m3s[8];
+    char _m8_pad[12];
+    /* FIXME: A padded struct is needed here. If you see this message, file a bug! */ S4_1 m3s[8];
 };
 
 struct SSBO0

reference/opt/shaders-msl/vert/packed_matrix.vert

@@ -16,7 +16,7 @@ struct _42
     float4x4 _m0;
     float4x4 _m1;
     float _m2;
-    char pad3[12];
+    char _m3_pad[12];
     packed_float3 _m3;
     float _m4;
     packed_float3 _m5;

reference/opt/shaders-msl/vert/ubo.alignment.vert

@@ -7,7 +7,7 @@ struct UBO
 {
     float4x4 mvp;
     float2 targSize;
-    char pad2[8];
+    char _m2_pad[8];
     packed_float3 color;
     float opacity;
 };

reference/shaders-msl/asm/comp/block-name-alias-global.asm.comp

@@ -22,7 +22,7 @@ struct A_2
 
 struct A_3
 {
-    A_2 Data[1024];
+    /* FIXME: A padded struct is needed here. If you see this message, file a bug! */ A_2 Data[1024];
 };
 
 struct B
@@ -32,7 +32,7 @@ struct B
 
 struct B_1
 {
-    A_2 Data[1024];
+    /* FIXME: A padded struct is needed here. If you see this message, file a bug! */ A_2 Data[1024];
 };
 
 kernel void main0(device B& C3 [[buffer(0)]], device A_1& C1 [[buffer(1)]], constant A_3& C2 [[buffer(2)]], constant B_1& C4 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])

reference/shaders-msl/asm/frag/vector-shuffle-oom.asm.frag

@@ -73,7 +73,7 @@ struct _18
     float2 _m24;
     float2 _m25;
     float2 _m26;
-    char pad27[8];
+    char _m27_pad[8];
     packed_float3 _m27;
     float _m28;
     float _m29;

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

@@ -43,7 +43,6 @@ struct Content
     S1 m1;
     S2 m2;
     S3 m3;
-    char pad7[4];
     float m4;
     S4 m3s[8];
 };
@@ -53,7 +52,6 @@ struct SSBO1
     Content content;
     Content content1[2];
     Content content2;
-    char pad3[8];
     float2x2 m0;
     float2x2 m1;
     float2x3 m2[4];
@@ -61,9 +59,9 @@ struct SSBO1
     float2x2 m4;
     float2x2 m5[9];
     packed_float2x3 m6[4][2];
-    char pad10[8];
+    char _m10_pad[8];
     float3x2 m7;
-    char pad11[8];
+    char _m11_pad[8];
     float array[1];
 };
 
@@ -105,9 +103,9 @@ struct Content_1
     S1_1 m1;
     S2_1 m2;
     S3_1 m3;
-    char pad7[4];
     float m4;
-    S4_1 m3s[8];
+    char _m8_pad[12];
+    /* FIXME: A padded struct is needed here. If you see this message, file a bug! */ S4_1 m3s[8];
 };
 
 struct SSBO0

reference/shaders-msl/vert/packed_matrix.vert

@@ -16,7 +16,7 @@ struct _42
     float4x4 _m0;
     float4x4 _m1;
     float _m2;
-    char pad3[12];
+    char _m3_pad[12];
     packed_float3 _m3;
     float _m4;
     packed_float3 _m5;

reference/shaders-msl/vert/ubo.alignment.vert

@@ -7,7 +7,7 @@ struct UBO
 {
     float4x4 mvp;
     float2 targSize;
-    char pad2[8];
+    char _m2_pad[8];
     packed_float3 color;
     float opacity;
 };

spirv_glsl.cpp

@@ -960,7 +960,7 @@ uint32_t CompilerGLSL::type_to_packed_alignment(const SPIRType &type, const Bits
 	if (type.basetype == SPIRType::Struct)
 	{
 		// Rule 9. Structs alignments are maximum alignment of its members.
-		uint32_t alignment = 0;
+		uint32_t alignment = 1;
 		for (uint32_t i = 0; i < type.member_types.size(); i++)
 		{
 			auto member_flags = ir.meta[type.self].members[i].decoration_flags;

spirv_msl.cpp

@@ -1565,11 +1565,11 @@ void CompilerMSL::align_struct(SPIRType &ib_type)
 
 		// Align current offset to the current member's default alignment.
 		size_t align_mask = get_declared_struct_member_alignment(ib_type, mbr_idx) - 1;
-		curr_offset = uint32_t((curr_offset + align_mask) & ~align_mask);
+		uint32_t aligned_curr_offset = uint32_t((curr_offset + align_mask) & ~align_mask);
 
 		// Fetch the member offset as declared in the SPIRV.
 		uint32_t mbr_offset = get_member_decoration(ib_type_id, mbr_idx, DecorationOffset);
-		if (mbr_offset > curr_offset)
+		if (mbr_offset > aligned_curr_offset)
 		{
 			// Since MSL and SPIR-V have slightly different struct member alignment and
 			// size rules, we'll pad to standard C-packing rules. If the member is farther
@@ -1610,6 +1610,18 @@ bool CompilerMSL::is_member_packable(SPIRType &ib_type, uint32_t index)
 		return true;
 	}
 
+	// Check for array of struct, where the SPIR-V declares an array stride which is larger than the struct itself.
+	// This can happen for struct A { float a }; A a[]; in std140 layout.
+	// TODO: Emit a padded struct which can be used for this purpose.
+	if (is_array(mbr_type) && mbr_type.basetype == SPIRType::Struct)
+	{
+		size_t declared_struct_size = get_declared_struct_size(mbr_type);
+		size_t alignment = get_declared_struct_member_alignment(ib_type, index);
+		declared_struct_size = (declared_struct_size + alignment - 1) & ~(alignment - 1);
+		if (type_struct_member_array_stride(ib_type, index) > declared_struct_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.");
@@ -4030,7 +4042,7 @@ void CompilerMSL::emit_struct_member(const SPIRType &type, uint32_t member_type_
 	MSLStructMemberKey key = get_struct_member_key(type.self, index);
 	uint32_t pad_len = struct_member_padding[key];
 	if (pad_len > 0)
-		statement("char pad", to_string(index), "[", to_string(pad_len), "];");
+		statement("char _m", index, "_pad", "[", to_string(pad_len), "];");
 
 	// If this member is packed, mark it as so.
 	string pack_pfx = "";
@@ -4041,7 +4053,11 @@ void CompilerMSL::emit_struct_member(const SPIRType &type, uint32_t member_type_
 	if (member_is_packed_type(type, index))
 	{
 		// If we're packing a matrix, output an appropriate typedef
-		if (membertype.vecsize > 1 && membertype.columns > 1)
+		if (membertype.basetype == SPIRType::Struct)
+		{
+			pack_pfx = "/* FIXME: A padded struct is needed here. If you see this message, file a bug! */ ";
+		}
+		else if (membertype.vecsize > 1 && membertype.columns > 1)
 		{
 			pack_pfx = "packed_";
 			string base_type = membertype.width == 16 ? "half" : "float";
@@ -5681,7 +5697,12 @@ size_t CompilerMSL::get_declared_struct_member_size(const SPIRType &struct_type,
 		}
 
 		if (type.basetype == SPIRType::Struct)
-			return get_declared_struct_size(type);
+		{
+			// The size of a struct in Metal is aligned up to its natural alignment.
+			auto size = get_declared_struct_size(type);
+			auto alignment = get_declared_struct_member_alignment(struct_type, index);
+			return (size + alignment - 1) & ~(alignment - 1);
+		}
 
 		uint32_t component_size = type.width / 8;
 		uint32_t vecsize = type.vecsize;
@@ -5712,7 +5733,13 @@ size_t CompilerMSL::get_declared_struct_member_alignment(const SPIRType &struct_
 		SPIRV_CROSS_THROW("Querying alignment of opaque object.");
 
 	case SPIRType::Struct:
-		return 16; // Per Vulkan spec section 14.5.4
+	{
+		// In MSL, a struct's alignment is equal to the maximum alignment of any of its members.
+		uint32_t alignment = 1;
+		for (uint32_t i = 0; i < type.member_types.size(); i++)
+			alignment = max(alignment, uint32_t(get_declared_struct_member_alignment(type, i)));
+		return alignment;
+	}
 
 	default:
 	{
@@ -5720,7 +5747,18 @@ size_t CompilerMSL::get_declared_struct_member_alignment(const SPIRType &struct_
 		// Alignment of unpacked type is the same as the vector size.
 		// Alignment of 3-elements vector is the same as 4-elements (including packed using column).
 		if (member_is_packed_type(struct_type, index))
-			return (type.width / 8) * (type.columns == 3 ? 4 : type.columns);
+		{
+			// This is getting pretty complicated.
+			// The special case of array of float/float2 needs to be handled here.
+			uint32_t packed_type_id =
+			    get_extended_member_decoration(struct_type.self, index, SPIRVCrossDecorationPackedType);
+			const SPIRType *packed_type = packed_type_id != 0 ? &get<SPIRType>(packed_type_id) : nullptr;
+			if (packed_type && is_array(*packed_type) && !is_matrix(*packed_type) &&
+			    packed_type->basetype != SPIRType::Struct)
+				return (packed_type->width / 8) * 4;
+			else
+				return (type.width / 8) * (type.columns == 3 ? 4 : type.columns);
+		}
 		else
 			return (type.width / 8) * (type.vecsize == 3 ? 4 : type.vecsize);
 	}