Merge pull request #2521 from KhronosGroup/fix-2515

HansKristian-Work · web-flow · commit df3b1cb0d3b2 · 2025-08-07T13:08:32.000+02:00
MSL: Fix edge case where a reference is taken of packed vector element.
diff --git a/reference/shaders-msl-no-opt/comp/atomic-cmpxchg-packed-vector.comp b/reference/shaders-msl-no-opt/comp/atomic-cmpxchg-packed-vector.comp
@@ -0,0 +1,37 @@
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+using namespace metal;
+
+struct AttData0
+{
+    packed_uint3 att0[1];
+};
+
+kernel void main0(device AttData0& __restrict _22 [[buffer(0)]])
+{
+    uint newVal = 432u;
+    uint prevVal = 0u;
+    uint curVal = 0u;
+    for (;;)
+    {
+        uint _30;
+        do
+        {
+            _30 = prevVal;
+        } while (!atomic_compare_exchange_weak_explicit((device atomic_uint*)&((device uint*)&_22.att0[0])[0u], &_30, newVal, memory_order_relaxed, memory_order_relaxed) && _30 == prevVal);
+        curVal = _30;
+        if (_30 != prevVal)
+        {
+            continue;
+        }
+        else
+        {
+            break;
+        }
+    }
+}
+
diff --git a/reference/shaders-msl-no-opt/packing/isolated-scalar-access.comp b/reference/shaders-msl-no-opt/packing/isolated-scalar-access.comp
@@ -17,7 +17,7 @@ kernel void main0(device SSBO& _12 [[buffer(0)]])
     threadgroup float4 shared_vec4;
     threadgroup float3 shared_vec3;
     ((device float*)&_12.v)[0u] = 10.0;
-    _12.v3[1u] = 40.0;
+    ((device float*)&_12.v3)[1u] = 40.0;
     ((device float*)&_12.cm[1])[2u] = 20.0;
     ((device float*)&_12.rm[1u])[3] = 30.0;
     ((threadgroup float*)&shared_vec4)[2u] = 40.0;
diff --git a/reference/shaders-msl-no-opt/packing/matrix-2x3-scalar.comp b/reference/shaders-msl-no-opt/packing/matrix-2x3-scalar.comp
@@ -74,8 +74,8 @@ void copy_columns(device SSBOCol& _29, device SSBORow& _41)
 static inline __attribute__((always_inline))
 void copy_elements(device SSBOCol& _29, device SSBORow& _41)
 {
-    _29.col_major0[0][1u] = ((device float*)&_41.row_major0[1u])[0];
-    ((device float*)&_41.row_major0[1u])[0] = _29.col_major0[0][1u];
+    ((device float*)&_29.col_major0[0])[1u] = ((device float*)&_41.row_major0[1u])[0];
+    ((device float*)&_41.row_major0[1u])[0] = ((device float*)&_29.col_major0[0])[1u];
 }
 
 kernel void main0(device SSBOCol& _29 [[buffer(0)]], device SSBORow& _41 [[buffer(1)]])
diff --git a/reference/shaders-msl-no-opt/packing/matrix-3x2-scalar.comp b/reference/shaders-msl-no-opt/packing/matrix-3x2-scalar.comp
@@ -73,8 +73,8 @@ void copy_columns(device SSBOCol& _29, device SSBORow& _41)
 static inline __attribute__((always_inline))
 void copy_elements(device SSBOCol& _29, device SSBORow& _41)
 {
-    ((device float*)&_29.col_major0[0])[1u] = _41.row_major0[1u][0];
-    _41.row_major0[1u][0] = ((device float*)&_29.col_major0[0])[1u];
+    ((device float*)&_29.col_major0[0])[1u] = ((device float*)&_41.row_major0[1u])[0];
+    ((device float*)&_41.row_major0[1u])[0] = ((device float*)&_29.col_major0[0])[1u];
 }
 
 kernel void main0(device SSBOCol& _29 [[buffer(0)]], device SSBORow& _41 [[buffer(1)]])
diff --git a/reference/shaders-msl-no-opt/packing/matrix-3x3-scalar.comp b/reference/shaders-msl-no-opt/packing/matrix-3x3-scalar.comp
@@ -84,8 +84,8 @@ void copy_columns(device SSBOCol& _29, device SSBORow& _41)
 static inline __attribute__((always_inline))
 void copy_elements(device SSBOCol& _29, device SSBORow& _41)
 {
-    _29.col_major0[0][1u] = _41.row_major0[1u][0];
-    _41.row_major0[1u][0] = _29.col_major0[0][1u];
+    ((device float*)&_29.col_major0[0])[1u] = ((device float*)&_41.row_major0[1u])[0];
+    ((device float*)&_41.row_major0[1u])[0] = ((device float*)&_29.col_major0[0])[1u];
 }
 
 kernel void main0(device SSBOCol& _29 [[buffer(0)]], device SSBORow& _41 [[buffer(1)]])
diff --git a/reference/shaders-msl-no-opt/packing/matrix-3x4-scalar.comp b/reference/shaders-msl-no-opt/packing/matrix-3x4-scalar.comp
@@ -81,8 +81,8 @@ void copy_columns(device SSBOCol& _29, device SSBORow& _41)
 static inline __attribute__((always_inline))
 void copy_elements(device SSBOCol& _29, device SSBORow& _41)
 {
-    ((device float*)&_29.col_major0[0])[1u] = _41.row_major0[1u][0];
-    _41.row_major0[1u][0] = ((device float*)&_29.col_major0[0])[1u];
+    ((device float*)&_29.col_major0[0])[1u] = ((device float*)&_41.row_major0[1u])[0];
+    ((device float*)&_41.row_major0[1u])[0] = ((device float*)&_29.col_major0[0])[1u];
 }
 
 kernel void main0(device SSBOCol& _29 [[buffer(0)]], device SSBORow& _41 [[buffer(1)]])
diff --git a/reference/shaders-msl-no-opt/packing/matrix-4x3-scalar.comp b/reference/shaders-msl-no-opt/packing/matrix-4x3-scalar.comp
@@ -80,8 +80,8 @@ void copy_columns(device SSBOCol& _29, device SSBORow& _41)
 static inline __attribute__((always_inline))
 void copy_elements(device SSBOCol& _29, device SSBORow& _41)
 {
-    _29.col_major0[0][1u] = ((device float*)&_41.row_major0[1u])[0];
-    ((device float*)&_41.row_major0[1u])[0] = _29.col_major0[0][1u];
+    ((device float*)&_29.col_major0[0])[1u] = ((device float*)&_41.row_major0[1u])[0];
+    ((device float*)&_41.row_major0[1u])[0] = ((device float*)&_29.col_major0[0])[1u];
 }
 
 kernel void main0(device SSBOCol& _29 [[buffer(0)]], device SSBORow& _41 [[buffer(1)]])
diff --git a/reference/shaders-msl-no-opt/packing/struct-packing-array-of-scalar.comp b/reference/shaders-msl-no-opt/packing/struct-packing-array-of-scalar.comp
@@ -17,6 +17,6 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
 
 kernel void main0(device SSBOScalar& buffer_scalar [[buffer(0)]])
 {
-    buffer_scalar.v[1].a[1u] = 1.0;
+    ((device float*)&buffer_scalar.v[1].a)[1u] = 1.0;
 }
 
diff --git a/reference/shaders-msl-no-opt/packing/struct-packing-recursive.comp b/reference/shaders-msl-no-opt/packing/struct-packing-recursive.comp
@@ -28,6 +28,6 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
 
 kernel void main0(device SSBOScalar& buffer_scalar [[buffer(0)]])
 {
-    buffer_scalar.baz.a.a.a[3u] = 10.0;
+    ((device float*)&buffer_scalar.baz.a.a.a)[3u] = 10.0;
 }
 
diff --git a/reference/shaders-msl-no-opt/packing/struct-packing.comp b/reference/shaders-msl-no-opt/packing/struct-packing.comp
@@ -23,7 +23,7 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
 
 kernel void main0(device SSBOScalar& buffer_scalar [[buffer(0)]])
 {
-    buffer_scalar.foo.a[0u] = 10.0;
-    buffer_scalar.bar.a[0u] = 20.0;
+    ((device float*)&buffer_scalar.foo.a)[0u] = 10.0;
+    ((device float*)&buffer_scalar.bar.a)[0u] = 20.0;
 }
 
diff --git a/shaders-msl-no-opt/comp/atomic-cmpxchg-packed-vector.comp b/shaders-msl-no-opt/comp/atomic-cmpxchg-packed-vector.comp
@@ -0,0 +1,17 @@
+#version 460
+#extension GL_EXT_nonuniform_qualifier : enable
+#extension GL_EXT_scalar_block_layout : require
+
+layout(scalar, binding=1) restrict buffer AttData0 {
+    uvec3 att0[];
+};
+
+void main() {
+    uint newVal = 432;
+    uint prevVal = 0;
+    uint curVal = 0;
+    
+    while ( (curVal = atomicCompSwap(att0[0].x, prevVal, newVal)) != prevVal)
+    {
+    }
+};
diff --git a/spirv_msl.cpp b/spirv_msl.cpp
@@ -9219,7 +9219,12 @@ bool CompilerMSL::prepare_access_chain_for_scalar_access(std::string &expr, cons
 	// and there is a risk of concurrent write access to other components,
 	// we must cast the access chain to a plain pointer to ensure we only access the exact scalars we expect.
 	// The MSL compiler refuses to allow component-level access for any non-packed vector types.
-	if (!is_packed && (storage == StorageClassStorageBuffer || storage == StorageClassWorkgroup))
+	// MSL refuses to take address or reference to vector component, even for packed types, so just force
+	// through the pointer cast. No much we can do sadly.
+	// For packed types, we could technically omit this if we know the reference does not have to turn into a pointer
+	// of some kind, but that requires external analysis passes to figure out, and
+	// this case is likely rare enough that we don't need to bother.
+	if (storage == StorageClassStorageBuffer || storage == StorageClassWorkgroup)
 	{
 		const char *addr_space = storage == StorageClassWorkgroup ? "threadgroup" : "device";
 		expr = join("((", addr_space, " ", type_to_glsl(type), "*)&", enclose_expression(expr), ")");

Original file line number	Diff line number	Diff line change
`@@ -74,8 +74,8 @@ void copy_columns(device SSBOCol& _29, device SSBORow& _41)`
`74`	`74`	`static inline __attribute__((always_inline))`
`75`	`75`	`void copy_elements(device SSBOCol& _29, device SSBORow& _41)`
`76`	`76`	`{`
`77`		`- _29.col_major0[0][1u] = ((device float*)&_41.row_major0[1u])[0];`
`78`		`- ((device float*)&_41.row_major0[1u])[0] = _29.col_major0[0][1u];`
	`77`	`+ ((device float)&_29.col_major0[0])[1u] = ((device float)&_41.row_major0[1u])[0];`
	`78`	`+ ((device float)&_41.row_major0[1u])[0] = ((device float)&_29.col_major0[0])[1u];`
`79`	`79`	`}`
`80`	`80`
`81`	`81`	`kernel void main0(device SSBOCol& _29 [[buffer(0)]], device SSBORow& _41 [[buffer(1)]])`
Original file line number	Diff line number	Diff line change
`@@ -73,8 +73,8 @@ void copy_columns(device SSBOCol& _29, device SSBORow& _41)`
`73`	`73`	`static inline __attribute__((always_inline))`
`74`	`74`	`void copy_elements(device SSBOCol& _29, device SSBORow& _41)`
`75`	`75`	`{`
`76`		`- ((device float*)&_29.col_major0[0])[1u] = _41.row_major0[1u][0];`
`77`		`- _41.row_major0[1u][0] = ((device float*)&_29.col_major0[0])[1u];`
	`76`	`+ ((device float)&_29.col_major0[0])[1u] = ((device float)&_41.row_major0[1u])[0];`
	`77`	`+ ((device float)&_41.row_major0[1u])[0] = ((device float)&_29.col_major0[0])[1u];`
`78`	`78`	`}`
`79`	`79`
`80`	`80`	`kernel void main0(device SSBOCol& _29 [[buffer(0)]], device SSBORow& _41 [[buffer(1)]])`
Original file line number	Diff line number	Diff line change
`@@ -84,8 +84,8 @@ void copy_columns(device SSBOCol& _29, device SSBORow& _41)`
`84`	`84`	`static inline __attribute__((always_inline))`
`85`	`85`	`void copy_elements(device SSBOCol& _29, device SSBORow& _41)`
`86`	`86`	`{`
`87`		`- _29.col_major0[0][1u] = _41.row_major0[1u][0];`
`88`		`- _41.row_major0[1u][0] = _29.col_major0[0][1u];`
	`87`	`+ ((device float)&_29.col_major0[0])[1u] = ((device float)&_41.row_major0[1u])[0];`
	`88`	`+ ((device float)&_41.row_major0[1u])[0] = ((device float)&_29.col_major0[0])[1u];`
`89`	`89`	`}`
`90`	`90`
`91`	`91`	`kernel void main0(device SSBOCol& _29 [[buffer(0)]], device SSBORow& _41 [[buffer(1)]])`
Original file line number	Diff line number	Diff line change
`@@ -81,8 +81,8 @@ void copy_columns(device SSBOCol& _29, device SSBORow& _41)`
`81`	`81`	`static inline __attribute__((always_inline))`
`82`	`82`	`void copy_elements(device SSBOCol& _29, device SSBORow& _41)`
`83`	`83`	`{`
`84`		`- ((device float*)&_29.col_major0[0])[1u] = _41.row_major0[1u][0];`
`85`		`- _41.row_major0[1u][0] = ((device float*)&_29.col_major0[0])[1u];`
	`84`	`+ ((device float)&_29.col_major0[0])[1u] = ((device float)&_41.row_major0[1u])[0];`
	`85`	`+ ((device float)&_41.row_major0[1u])[0] = ((device float)&_29.col_major0[0])[1u];`
`86`	`86`	`}`
`87`	`87`
`88`	`88`	`kernel void main0(device SSBOCol& _29 [[buffer(0)]], device SSBORow& _41 [[buffer(1)]])`
Original file line number	Diff line number	Diff line change
`@@ -80,8 +80,8 @@ void copy_columns(device SSBOCol& _29, device SSBORow& _41)`
`80`	`80`	`static inline __attribute__((always_inline))`
`81`	`81`	`void copy_elements(device SSBOCol& _29, device SSBORow& _41)`
`82`	`82`	`{`
`83`		`- _29.col_major0[0][1u] = ((device float*)&_41.row_major0[1u])[0];`
`84`		`- ((device float*)&_41.row_major0[1u])[0] = _29.col_major0[0][1u];`
	`83`	`+ ((device float)&_29.col_major0[0])[1u] = ((device float)&_41.row_major0[1u])[0];`
	`84`	`+ ((device float)&_41.row_major0[1u])[0] = ((device float)&_29.col_major0[0])[1u];`
`85`	`85`	`}`
`86`	`86`
`87`	`87`	`kernel void main0(device SSBOCol& _29 [[buffer(0)]], device SSBORow& _41 [[buffer(1)]])`
Original file line number	Diff line number	Diff line change
`@@ -17,6 +17,6 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);`
`17`	`17`
`18`	`18`	`kernel void main0(device SSBOScalar& buffer_scalar [[buffer(0)]])`
`19`	`19`	`{`
`20`		`- buffer_scalar.v[1].a[1u] = 1.0;`
	`20`	`+ ((device float*)&buffer_scalar.v[1].a)[1u] = 1.0;`
`21`	`21`	`}`
`22`	`22`
Original file line number	Diff line number	Diff line change
`@@ -28,6 +28,6 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);`
`28`	`28`
`29`	`29`	`kernel void main0(device SSBOScalar& buffer_scalar [[buffer(0)]])`
`30`	`30`	`{`
`31`		`- buffer_scalar.baz.a.a.a[3u] = 10.0;`
	`31`	`+ ((device float*)&buffer_scalar.baz.a.a.a)[3u] = 10.0;`
`32`	`32`	`}`
`33`	`33`
Original file line number	Diff line number	Diff line change
`@@ -23,7 +23,7 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);`
`23`	`23`
`24`	`24`	`kernel void main0(device SSBOScalar& buffer_scalar [[buffer(0)]])`
`25`	`25`	`{`
`26`		`- buffer_scalar.foo.a[0u] = 10.0;`
`27`		`- buffer_scalar.bar.a[0u] = 20.0;`
	`26`	`+ ((device float*)&buffer_scalar.foo.a)[0u] = 10.0;`
	`27`	`+ ((device float*)&buffer_scalar.bar.a)[0u] = 20.0;`
`28`	`28`	`}`
`29`	`29`