MSL: Avoid packed arrays in more cases.

HansKristian-Work · HansKristian-Work · commit 107ab7c2b734 · 2020-05-06T10:27:12.000+02:00
Extend the array stride relaxation to non-packed arrays as well, as
long as the array in question contains a single array element.
diff --git a/reference/opt/shaders-msl/asm/vert/packed-bool-to-uint.asm.vert b/reference/opt/shaders-msl/asm/vert/packed-bool-to-uint.asm.vert
@@ -29,7 +29,7 @@ vertex main0_out main0(main0_in in [[stage_in]], constant defaultUniformsVS& _9
 {
     main0_out out = {};
     out.gl_Position = _9.umatrix * float4(_9.uquad[int(gl_VertexIndex)].x, _9.uquad[int(gl_VertexIndex)].y, in.a_position.z, in.a_position.w);
-    if (uint(_9.flags.flags[0]) != 0u)
+    if (_9.flags.flags[0] != 0u)
     {
         out.gl_Position.z = 0.0;
     }
diff --git a/reference/opt/shaders-msl/asm/vert/packed-bool2-to-packed_uint2.asm.vert b/reference/opt/shaders-msl/asm/vert/packed-bool2-to-packed_uint2.asm.vert
@@ -5,7 +5,7 @@ using namespace metal;
 
 struct Struct
 {
-    packed_uint2 flags[1];
+    uint2 flags[1];
 };
 
 struct defaultUniformsVS
@@ -29,7 +29,7 @@ vertex main0_out main0(main0_in in [[stage_in]], constant defaultUniformsVS& _9
 {
     main0_out out = {};
     out.gl_Position = _9.umatrix * float4(_9.uquad[int(gl_VertexIndex)].x, _9.uquad[int(gl_VertexIndex)].y, in.a_position.z, in.a_position.w);
-    if (_9.flags.flags[0][0u] != 0u)
+    if (_9.flags.flags[0].x != 0u)
     {
         out.gl_Position.z = 0.0;
     }
diff --git a/reference/opt/shaders-msl/comp/struct-packing.comp b/reference/opt/shaders-msl/comp/struct-packing.comp
@@ -66,7 +66,7 @@ struct SSBO1
 
 struct S0_1
 {
-    packed_float2 a[1];
+    float2 a[1];
     char _m1_pad[8];
     float b;
     char _m0_final_padding[12];
@@ -124,13 +124,13 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
 kernel void main0(device SSBO1& ssbo_430 [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]])
 {
     Content_1 _60 = ssbo_140.content;
-    ssbo_430.content.m0s[0].a[0] = float2(_60.m0s[0].a[0]);
+    ssbo_430.content.m0s[0].a[0] = _60.m0s[0].a[0];
     ssbo_430.content.m0s[0].b = _60.m0s[0].b;
     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] = float2(_60.m0.a[0]);
+    ssbo_430.content.m0.a[0] = _60.m0.a[0];
     ssbo_430.content.m0.b = _60.m0.b;
     ssbo_430.content.m1.a = float3(_60.m1.a);
     ssbo_430.content.m1.b = _60.m1.b;
diff --git a/reference/opt/shaders-msl/vert/packed-bool-to-uint.vert b/reference/opt/shaders-msl/vert/packed-bool-to-uint.vert
@@ -29,7 +29,7 @@ vertex main0_out main0(main0_in in [[stage_in]], constant defaultUniformsVS& _24
 {
     main0_out out = {};
     out.gl_Position = _24.umatrix * float4(_24.uquad[int(gl_VertexIndex)].x, _24.uquad[int(gl_VertexIndex)].y, in.a_position.z, in.a_position.w);
-    if (uint(_24.flags.flags[0]) != 0u)
+    if (_24.flags.flags[0] != 0u)
     {
         out.gl_Position.z = 0.0;
     }
diff --git a/reference/opt/shaders-msl/vert/packed-bool2-to-packed_uint2.vert b/reference/opt/shaders-msl/vert/packed-bool2-to-packed_uint2.vert
@@ -5,7 +5,7 @@ using namespace metal;
 
 struct Struct
 {
-    packed_uint2 flags[1];
+    uint2 flags[1];
 };
 
 struct defaultUniformsVS
@@ -29,7 +29,7 @@ vertex main0_out main0(main0_in in [[stage_in]], constant defaultUniformsVS& _25
 {
     main0_out out = {};
     out.gl_Position = _25.umatrix * float4(_25.uquad[int(gl_VertexIndex)].x, _25.uquad[int(gl_VertexIndex)].y, in.a_position.z, in.a_position.w);
-    if (_25.flags.flags[0][0u] != 0u)
+    if (_25.flags.flags[0].x != 0u)
     {
         out.gl_Position.z = 0.0;
     }
diff --git a/reference/shaders-msl-no-opt/comp/struct-packing-scalar.nocompat.invalid.vk.comp b/reference/shaders-msl-no-opt/comp/struct-packing-scalar.nocompat.invalid.vk.comp
@@ -61,7 +61,7 @@ struct SSBO1
 
 struct S0_1
 {
-    packed_float2 a[1];
+    float2 a[1];
     char _m1_pad[8];
     float b;
     char _m0_final_padding[12];
@@ -126,13 +126,13 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
 
 kernel void main0(device SSBO1& ssbo_scalar [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]], device SSBO2& ssbo_scalar2 [[buffer(2)]])
 {
-    ssbo_scalar.content.m0s[0].a[0] = float2(ssbo_140.content.m0s[0].a[0]);
+    ssbo_scalar.content.m0s[0].a[0] = ssbo_140.content.m0s[0].a[0];
     ssbo_scalar.content.m0s[0].b = ssbo_140.content.m0s[0].b;
     ssbo_scalar.content.m1s[0].a = float3(ssbo_140.content.m1s[0].a);
     ssbo_scalar.content.m1s[0].b = ssbo_140.content.m1s[0].b;
     ssbo_scalar.content.m2s[0].a[0] = ssbo_140.content.m2s[0].a[0];
     ssbo_scalar.content.m2s[0].b = ssbo_140.content.m2s[0].b;
-    ssbo_scalar.content.m0.a[0] = float2(ssbo_140.content.m0.a[0]);
+    ssbo_scalar.content.m0.a[0] = ssbo_140.content.m0.a[0];
     ssbo_scalar.content.m0.b = ssbo_140.content.m0.b;
     ssbo_scalar.content.m1.a = float3(ssbo_140.content.m1.a);
     ssbo_scalar.content.m1.b = ssbo_140.content.m1.b;
diff --git a/reference/shaders-msl/asm/vert/packed-bool-to-uint.asm.vert b/reference/shaders-msl/asm/vert/packed-bool-to-uint.asm.vert
@@ -29,7 +29,7 @@ vertex main0_out main0(main0_in in [[stage_in]], constant defaultUniformsVS& _9
 {
     main0_out out = {};
     out.gl_Position = _9.umatrix * float4(_9.uquad[int(gl_VertexIndex)].x, _9.uquad[int(gl_VertexIndex)].y, in.a_position.z, in.a_position.w);
-    if (uint(_9.flags.flags[0]) != 0u)
+    if (_9.flags.flags[0] != 0u)
     {
         out.gl_Position.z = 0.0;
     }
diff --git a/reference/shaders-msl/asm/vert/packed-bool2-to-packed_uint2.asm.vert b/reference/shaders-msl/asm/vert/packed-bool2-to-packed_uint2.asm.vert
@@ -5,7 +5,7 @@ using namespace metal;
 
 struct Struct
 {
-    packed_uint2 flags[1];
+    uint2 flags[1];
 };
 
 struct defaultUniformsVS
@@ -29,7 +29,7 @@ vertex main0_out main0(main0_in in [[stage_in]], constant defaultUniformsVS& _9
 {
     main0_out out = {};
     out.gl_Position = _9.umatrix * float4(_9.uquad[int(gl_VertexIndex)].x, _9.uquad[int(gl_VertexIndex)].y, in.a_position.z, in.a_position.w);
-    if (_9.flags.flags[0][0u] != 0u)
+    if (_9.flags.flags[0].x != 0u)
     {
         out.gl_Position.z = 0.0;
     }
diff --git a/reference/shaders-msl/comp/struct-packing.comp b/reference/shaders-msl/comp/struct-packing.comp
@@ -66,7 +66,7 @@ struct SSBO1
 
 struct S0_1
 {
-    packed_float2 a[1];
+    float2 a[1];
     char _m1_pad[8];
     float b;
     char _m0_final_padding[12];
@@ -124,13 +124,13 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
 kernel void main0(device SSBO1& ssbo_430 [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]])
 {
     Content_1 _60 = ssbo_140.content;
-    ssbo_430.content.m0s[0].a[0] = float2(_60.m0s[0].a[0]);
+    ssbo_430.content.m0s[0].a[0] = _60.m0s[0].a[0];
     ssbo_430.content.m0s[0].b = _60.m0s[0].b;
     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] = float2(_60.m0.a[0]);
+    ssbo_430.content.m0.a[0] = _60.m0.a[0];
     ssbo_430.content.m0.b = _60.m0.b;
     ssbo_430.content.m1.a = float3(_60.m1.a);
     ssbo_430.content.m1.b = _60.m1.b;
diff --git a/reference/shaders-msl/vert/packed-bool-to-uint.vert b/reference/shaders-msl/vert/packed-bool-to-uint.vert
@@ -29,7 +29,7 @@ vertex main0_out main0(main0_in in [[stage_in]], constant defaultUniformsVS& _24
 {
     main0_out out = {};
     out.gl_Position = _24.umatrix * float4(_24.uquad[int(gl_VertexIndex)].x, _24.uquad[int(gl_VertexIndex)].y, in.a_position.z, in.a_position.w);
-    if (uint(_24.flags.flags[0]) != 0u)
+    if (_24.flags.flags[0] != 0u)
     {
         out.gl_Position.z = 0.0;
     }
diff --git a/reference/shaders-msl/vert/packed-bool2-to-packed_uint2.vert b/reference/shaders-msl/vert/packed-bool2-to-packed_uint2.vert
@@ -5,7 +5,7 @@ using namespace metal;
 
 struct Struct
 {
-    packed_uint2 flags[1];
+    uint2 flags[1];
 };
 
 struct defaultUniformsVS
@@ -29,7 +29,7 @@ vertex main0_out main0(main0_in in [[stage_in]], constant defaultUniformsVS& _25
 {
     main0_out out = {};
     out.gl_Position = _25.umatrix * float4(_25.uquad[int(gl_VertexIndex)].x, _25.uquad[int(gl_VertexIndex)].y, in.a_position.z, in.a_position.w);
-    if (_25.flags.flags[0][0u] != 0u)
+    if (_25.flags.flags[0].x != 0u)
     {
         out.gl_Position.z = 0.0;
     }
diff --git a/spirv_msl.cpp b/spirv_msl.cpp
@@ -3087,11 +3087,11 @@ bool CompilerMSL::validate_member_packing_rules_msl(const SPIRType &type, uint32
 	{
 		// If we have an array type, array stride must match exactly with SPIR-V.
 
-		// An exception to this requirement is if we have one array element and a packed decoration.
+		// An exception to this requirement is if we have one array element.
 		// This comes from DX scalar layout workaround.
 		// If app tries to be cheeky and access the member out of bounds, this will not work, but this is the best we can do.
-		bool relax_array_stride = has_extended_member_decoration(type.self, index, SPIRVCrossDecorationPhysicalTypePacked) &&
-		                          mbr_type.array.back() == 1 && mbr_type.array_size_literal.back();
+		// In OpAccessChain with logical memory models, access chains must be in-bounds in SPIR-V specification.
+		bool relax_array_stride = mbr_type.array.back() == 1 && mbr_type.array_size_literal.back();
 
 		if (!relax_array_stride)
 		{
@@ -3137,7 +3137,9 @@ void CompilerMSL::ensure_member_packing_rules_msl(SPIRType &ib_type, uint32_t in
 		SPIRV_CROSS_THROW("Cannot perform any repacking for structs when it is used as a member of another struct.");
 
 	// Perform remapping here.
-	set_extended_member_decoration(ib_type.self, index, SPIRVCrossDecorationPhysicalTypePacked);
+	// There is nothing to be gained by using packed scalars, so don't attempt it.
+	if (!is_scalar(ib_type))
+		set_extended_member_decoration(ib_type.self, index, SPIRVCrossDecorationPhysicalTypePacked);
 
 	// Try validating again, now with packed.
 	if (validate_member_packing_rules_msl(ib_type, index))

Original file line number	Diff line number	Diff line change
`@@ -29,7 +29,7 @@ vertex main0_out main0(main0_in in [[stage_in]], constant defaultUniformsVS& _9`
`29`	`29`	`{`
`30`	`30`	`main0_out out = {};`
`31`	`31`	`out.gl_Position = _9.umatrix * float4(_9.uquad[int(gl_VertexIndex)].x, _9.uquad[int(gl_VertexIndex)].y, in.a_position.z, in.a_position.w);`
`32`		`- if (uint(_9.flags.flags[0]) != 0u)`
	`32`	`+ if (_9.flags.flags[0] != 0u)`
`33`	`33`	`{`
`34`	`34`	`out.gl_Position.z = 0.0;`
`35`	`35`	`}`
Original file line number	Diff line number	Diff line change
`@@ -5,7 +5,7 @@ using namespace metal;`
`5`	`5`
`6`	`6`	`struct Struct`
`7`	`7`	`{`
`8`		`- packed_uint2 flags[1];`
	`8`	`+ uint2 flags[1];`
`9`	`9`	`};`
`10`	`10`
`11`	`11`	`struct defaultUniformsVS`
`@@ -29,7 +29,7 @@ vertex main0_out main0(main0_in in [[stage_in]], constant defaultUniformsVS& _9`
`29`	`29`	`{`
`30`	`30`	`main0_out out = {};`
`31`	`31`	`out.gl_Position = _9.umatrix * float4(_9.uquad[int(gl_VertexIndex)].x, _9.uquad[int(gl_VertexIndex)].y, in.a_position.z, in.a_position.w);`
`32`		`- if (_9.flags.flags[0][0u] != 0u)`
	`32`	`+ if (_9.flags.flags[0].x != 0u)`
`33`	`33`	`{`
`34`	`34`	`out.gl_Position.z = 0.0;`
`35`	`35`	`}`