MSL: Fix calculation of atomic image buffer address.

Fix reversed coordinates: `y` should be used to calculate the row
address. Align row address to the row stride.

I've made the row alignment a function constant; this makes it possible
to override it at pipeline compile time.

Honestly, I don't know how this worked at all for Epic. It definitely
didn't work in the CTS prior to this.

4 files changed, 16 insertions, 4 deletions

diff --git a/reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.argument.msl2.comp b/reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.argument.msl2.comp
index b3bf2b84..7dea8b71 100644
--- a/reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.argument.msl2.comp
+++ b/reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.argument.msl2.comp
@@ -23,8 +23,11 @@ struct spvDescriptorSetBuffer0
     sampler uTextureSmplr [[id(4)]];
 };
 
+// The required alignment of a linear texture of R32Uint format.
+constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
+constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
 // Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
-#define spvImage2DAtomicCoord(tc, tex) (((tex).get_width() * (tc).x) + (tc).y)
+#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() +  spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
 
 kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
diff --git a/reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.comp b/reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.comp
index a8ade54b..a2846474 100644
--- a/reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.comp
+++ b/reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.comp
@@ -14,8 +14,11 @@ struct SSBO
 
 constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
 
+// The required alignment of a linear texture of R32Uint format.
+constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
+constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
 // Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
-#define spvImage2DAtomicCoord(tc, tex) (((tex).get_width() * (tc).x) + (tc).y)
+#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() +  spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
 
 kernel void main0(device SSBO& _31 [[buffer(1)]], texture2d<uint> uImage [[texture(0)]], device atomic_uint* uImage_atomic [[buffer(0)]], texture2d<float> uTexture [[texture(1)]], sampler uTextureSmplr [[sampler(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
diff --git a/reference/opt/shaders-msl/frag/pixel-interlock-ordered.msl2.argument.frag b/reference/opt/shaders-msl/frag/pixel-interlock-ordered.msl2.argument.frag
index dfef91c8..adea453d 100644
--- a/reference/opt/shaders-msl/frag/pixel-interlock-ordered.msl2.argument.frag
+++ b/reference/opt/shaders-msl/frag/pixel-interlock-ordered.msl2.argument.frag
@@ -35,8 +35,11 @@ struct spvDescriptorSetBuffer0
     device Buffer2* m_52 [[id(7), raster_order_group(0)]];
 };
 
+// The required alignment of a linear texture of R32Uint format.
+constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
+constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
 // Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
-#define spvImage2DAtomicCoord(tc, tex) (((tex).get_width() * (tc).x) + (tc).y)
+#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() +  spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
 
 fragment void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]])
 {
diff --git a/reference/opt/shaders-msl/frag/pixel-interlock-ordered.msl2.frag b/reference/opt/shaders-msl/frag/pixel-interlock-ordered.msl2.frag
index 67b79d9d..e409ea06 100644
--- a/reference/opt/shaders-msl/frag/pixel-interlock-ordered.msl2.frag
+++ b/reference/opt/shaders-msl/frag/pixel-interlock-ordered.msl2.frag
@@ -23,8 +23,11 @@ struct Buffer2
     uint quux;
 };
 
+// The required alignment of a linear texture of R32Uint format.
+constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
+constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
 // Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
-#define spvImage2DAtomicCoord(tc, tex) (((tex).get_width() * (tc).x) + (tc).y)
+#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() +  spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
 
 fragment void main0(device Buffer3& _9 [[buffer(0)]], volatile device Buffer& _42 [[buffer(2), raster_order_group(0)]], device Buffer2& _52 [[buffer(3), raster_order_group(0)]], texture2d<float, access::write> img4 [[texture(0)]], texture2d<float, access::write> img [[texture(1), raster_order_group(0)]], texture2d<float> img3 [[texture(2), raster_order_group(0)]], texture2d<uint> img2 [[texture(3), raster_order_group(0)]], device atomic_uint* img2_atomic [[buffer(1), raster_order_group(0)]])
 {