Merge pull request #2572 from KhronosGroup/fix-2564

HansKristian-Work · web-flow · commit 8d838b4e2e1e · 2025-11-26T17:33:42.000+01:00
MSL: Fix handling of variable pointer arguments which are not BDA.
diff --git a/reference/shaders-msl-no-opt/asm/comp/storage-buffer-pointer-argument.asm.comp b/reference/shaders-msl-no-opt/asm/comp/storage-buffer-pointer-argument.asm.comp
@@ -16,13 +16,13 @@ struct SSBORead
 };
 
 static inline __attribute__((always_inline))
-void copy_out(device float& A, device float& B)
+void copy_out(device float* A, device float* B)
 {
-    A = B;
+    *A = *B;
 }
 
 kernel void main0(device SSBO& _10 [[buffer(0)]], device SSBORead& _14 [[buffer(1)]])
 {
-    copy_out(_10.a, _14.b);
+    copy_out(&_10.a, &_14.b);
 }
 
diff --git a/reference/shaders-msl-no-opt/asm/comp/variable-ssbo-argument.spv16.asm.comp b/reference/shaders-msl-no-opt/asm/comp/variable-ssbo-argument.spv16.asm.comp
@@ -0,0 +1,23 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct _3
+{
+    uchar _m0[1];
+};
+
+static inline __attribute__((always_inline))
+void _20(device uchar* _21)
+{
+    _21[2u] = uchar(0);
+}
+
+kernel void main0(device _3& _2 [[buffer(0)]])
+{
+    _20(&_2._m0[1u]);
+}
+
diff --git a/reference/shaders-msl-no-opt/asm/comp/variable-ssbo-array-argument.spv16.asm.comp b/reference/shaders-msl-no-opt/asm/comp/variable-ssbo-array-argument.spv16.asm.comp
@@ -0,0 +1,23 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct _3
+{
+    uchar _m0[16];
+};
+
+static inline __attribute__((always_inline))
+void _20(device uchar (*_21)[16])
+{
+    (*_21)[2u] = uchar(0);
+}
+
+kernel void main0(device _3& _2 [[buffer(0)]])
+{
+    _20(&_2._m0);
+}
+
diff --git a/reference/shaders-msl/asm/comp/variable-pointers-2.asm.comp b/reference/shaders-msl/asm/comp/variable-pointers-2.asm.comp
@@ -18,9 +18,9 @@ struct bar
 };
 
 static inline __attribute__((always_inline))
-device foo* select_buffer(device foo& a, constant bar& cb)
+device foo* select_buffer(device foo* a, constant bar& cb)
 {
-    return (cb.d != 0) ? &a : nullptr;
+    return (cb.d != 0) ? a : nullptr;
 }
 
 static inline __attribute__((always_inline))
@@ -31,7 +31,7 @@ thread uint3* select_input(thread uint3& gl_GlobalInvocationID, thread uint3& gl
 
 kernel void main0(device foo& buf [[buffer(0)]], constant bar& cb [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
 {
-    device foo* _44 = select_buffer(buf, cb);
+    device foo* _44 = select_buffer(&buf, cb);
     device foo* _65 = _44;
     thread uint3* _45 = select_input(gl_GlobalInvocationID, gl_LocalInvocationID, cb);
     device foo* _66 = _65;
diff --git a/reference/shaders-msl/asm/comp/variable-pointers-store-forwarding.asm.comp b/reference/shaders-msl/asm/comp/variable-pointers-store-forwarding.asm.comp
@@ -16,14 +16,14 @@ struct bar
 };
 
 static inline __attribute__((always_inline))
-device int* _24(device foo& a, device bar& b, thread uint3& gl_GlobalInvocationID)
+device int* _24(device foo* a, device bar* b, thread uint3& gl_GlobalInvocationID)
 {
-    return (gl_GlobalInvocationID.x != 0u) ? &a.a : &b.b;
+    return (gl_GlobalInvocationID.x != 0u) ? &a->a : &b->b;
 }
 
 kernel void main0(device foo& x [[buffer(0)]], device bar& y [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
-    device int* _34 = _24(x, y, gl_GlobalInvocationID);
+    device int* _34 = _24(&x, &y, gl_GlobalInvocationID);
     device int* _33 = _34;
     int _37 = x.a;
     *_33 = 0;
diff --git a/shaders-msl-no-opt/asm/comp/variable-ssbo-argument.spv16.asm.comp b/shaders-msl-no-opt/asm/comp/variable-ssbo-argument.spv16.asm.comp
@@ -0,0 +1,44 @@
+; SPIR-V
+; Version: 1.6
+; Generator: Khronos SPIR-V Tools Assembler; 0
+; Bound: 24
+; Schema: 0
+               OpCapability VariablePointersStorageBuffer
+               OpCapability Int8
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %1 "main" %2
+               OpExecutionMode %1 LocalSize 16 1 1
+               OpDecorate %2 DescriptorSet 0
+               OpDecorate %2 Binding 0
+               OpDecorate %_struct_3 Block
+               OpMemberDecorate %_struct_3 0 Offset 0
+               OpDecorate %_runtimearr_uchar ArrayStride 1
+               OpDecorate %_ptr_StorageBuffer_uchar ArrayStride 1
+       %void = OpTypeVoid
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+     %uint_3 = OpConstant %uint 3
+      %uchar = OpTypeInt 8 0
+    %uchar_0 = OpConstant %uchar 0
+%_runtimearr_uchar = OpTypeRuntimeArray %uchar
+  %_struct_3 = OpTypeStruct %_runtimearr_uchar
+%_ptr_StorageBuffer_uchar = OpTypePointer StorageBuffer %uchar
+%_ptr_StorageBuffer__struct_3 = OpTypePointer StorageBuffer %_struct_3
+          %2 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
+         %15 = OpTypeFunction %void %_ptr_StorageBuffer_uchar
+         %16 = OpTypeFunction %void
+          %1 = OpFunction %void None %16
+         %17 = OpLabel
+         %18 = OpAccessChain %_ptr_StorageBuffer_uchar %2 %uint_0 %uint_1
+         %19 = OpFunctionCall %void %20 %18
+               OpReturn
+               OpFunctionEnd
+         %20 = OpFunction %void None %15
+         %21 = OpFunctionParameter %_ptr_StorageBuffer_uchar
+         %22 = OpLabel
+         %23 = OpPtrAccessChain %_ptr_StorageBuffer_uchar %21 %uint_2
+               OpStore %23 %uchar_0
+               OpReturn
+               OpFunctionEnd
diff --git a/shaders-msl-no-opt/asm/comp/variable-ssbo-array-argument.spv16.asm.comp b/shaders-msl-no-opt/asm/comp/variable-ssbo-array-argument.spv16.asm.comp
@@ -0,0 +1,45 @@
+; SPIR-V
+; Version: 1.6
+; Generator: Khronos SPIR-V Tools Assembler; 0
+; Bound: 24
+; Schema: 0
+               OpCapability VariablePointersStorageBuffer
+               OpCapability Int8
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %1 "main" %2
+               OpExecutionMode %1 LocalSize 16 1 1
+               OpDecorate %2 DescriptorSet 0
+               OpDecorate %2 Binding 0
+               OpDecorate %_struct_3 Block
+               OpMemberDecorate %_struct_3 0 Offset 0
+               OpDecorate %uchar_array ArrayStride 1
+       %void = OpTypeVoid
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+     %uint_3 = OpConstant %uint 3
+     %uint_16 = OpConstant %uint 16
+      %uchar = OpTypeInt 8 0
+    %uchar_0 = OpConstant %uchar 0
+%uchar_array = OpTypeArray %uchar %uint_16
+  %_struct_3 = OpTypeStruct %uchar_array
+%_ptr_StorageBuffer_uchar = OpTypePointer StorageBuffer %uchar
+%_ptr_StorageBuffer_uchar_array = OpTypePointer StorageBuffer %uchar_array
+%_ptr_StorageBuffer__struct_3 = OpTypePointer StorageBuffer %_struct_3
+          %2 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
+         %15 = OpTypeFunction %void %_ptr_StorageBuffer_uchar_array
+         %16 = OpTypeFunction %void
+          %1 = OpFunction %void None %16
+         %17 = OpLabel
+         %18 = OpAccessChain %_ptr_StorageBuffer_uchar_array %2 %uint_0
+         %19 = OpFunctionCall %void %20 %18
+               OpReturn
+               OpFunctionEnd
+         %20 = OpFunction %void None %15
+         %21 = OpFunctionParameter %_ptr_StorageBuffer_uchar_array
+         %22 = OpLabel
+         %23 = OpAccessChain %_ptr_StorageBuffer_uchar %21 %uint_2
+               OpStore %23 %uchar_0
+               OpReturn
+               OpFunctionEnd
diff --git a/spirv_cross.cpp b/spirv_cross.cpp
@@ -742,6 +742,14 @@ bool Compiler::is_physical_pointer(const SPIRType &type) const
 	return type.op == OpTypePointer && type.storage == StorageClassPhysicalStorageBuffer;
 }
 
+bool Compiler::is_physical_or_buffer_pointer(const SPIRType &type) const
+{
+	return type.op == OpTypePointer &&
+	       (type.storage == StorageClassPhysicalStorageBuffer || type.storage == StorageClassUniform ||
+	        type.storage == StorageClassStorageBuffer || type.storage == StorageClassWorkgroup ||
+	        type.storage == StorageClassPushConstant);
+}
+
 bool Compiler::is_physical_pointer_to_buffer_block(const SPIRType &type) const
 {
 	return is_physical_pointer(type) && get_pointee_type(type).self == type.parent_type &&
diff --git a/spirv_cross.hpp b/spirv_cross.hpp
@@ -694,6 +694,7 @@ class Compiler
 	bool is_array(const SPIRType &type) const;
 	bool is_pointer(const SPIRType &type) const;
 	bool is_physical_pointer(const SPIRType &type) const;
+	bool is_physical_or_buffer_pointer(const SPIRType &type) const;
 	bool is_physical_pointer_to_buffer_block(const SPIRType &type) const;
 	static bool is_runtime_size_array(const SPIRType &type);
 	uint32_t expression_type_id(uint32_t id) const;
diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp
@@ -5051,10 +5051,10 @@ void CompilerGLSL::emit_polyfills(uint32_t polyfills, bool relaxed)
 // Returns a string representation of the ID, usable as a function arg.
 // Default is to simply return the expression representation fo the arg ID.
 // Subclasses may override to modify the return value.
-string CompilerGLSL::to_func_call_arg(const SPIRFunction::Parameter &, uint32_t id)
+string CompilerGLSL::to_func_call_arg(const SPIRFunction::Parameter &arg, uint32_t id)
 {
 	// BDA expects pointers through function interface.
-	if (is_physical_pointer(expression_type(id)))
+	if (!arg.alias_global_variable && is_physical_or_buffer_pointer(expression_type(id)))
 		return to_pointer_expression(id);
 
 	// Make sure that we use the name of the original variable, and not the parameter alias.
@@ -11581,7 +11581,7 @@ bool CompilerGLSL::should_dereference(uint32_t id)
 	// If id is a variable but not a phi variable, we should not dereference it.
 	// BDA passed around as parameters are always pointers.
 	if (auto *var = maybe_get<SPIRVariable>(id))
-		return (var->parameter && is_physical_pointer(type)) || var->phi_variable;
+		return (var->parameter && is_physical_or_buffer_pointer(type)) || var->phi_variable;
 
 	if (auto *expr = maybe_get<SPIRExpression>(id))
 	{
@@ -11617,8 +11617,8 @@ bool CompilerGLSL::should_dereference(uint32_t id)
 bool CompilerGLSL::should_dereference_caller_param(uint32_t id)
 {
 	const auto &type = expression_type(id);
-	// BDA is always passed around as pointers.
-	if (is_physical_pointer(type))
+	// BDA is always passed around as pointers. Similarly, we need to pass variable buffer pointers as pointers.
+	if (is_physical_or_buffer_pointer(type))
 		return false;
 
 	return should_dereference(id);
diff --git a/spirv_msl.cpp b/spirv_msl.cpp
diff --git a/spirv_msl.hpp b/spirv_msl.hpp

Original file line number	Diff line number	Diff line change
`@@ -16,13 +16,13 @@ struct SSBORead`
`16`	`16`	`};`
`17`	`17`
`18`	`18`	`static inline __attribute__((always_inline))`
`19`		`-void copy_out(device float& A, device float& B)`
	`19`	`+void copy_out(device float* A, device float* B)`
`20`	`20`	`{`
`21`		`- A = B;`
	`21`	`+ A = B;`
`22`	`22`	`}`
`23`	`23`
`24`	`24`	`kernel void main0(device SSBO& _10 [[buffer(0)]], device SSBORead& _14 [[buffer(1)]])`
`25`	`25`	`{`
`26`		`- copy_out(_10.a, _14.b);`
	`26`	`+ copy_out(&_10.a, &_14.b);`
`27`	`27`	`}`
`28`	`28`
Original file line number	Diff line number	Diff line change
`@@ -18,9 +18,9 @@ struct bar`
`18`	`18`	`};`
`19`	`19`
`20`	`20`	`static inline __attribute__((always_inline))`
`21`		`-device foo* select_buffer(device foo& a, constant bar& cb)`
	`21`	`+device foo* select_buffer(device foo* a, constant bar& cb)`
`22`	`22`	`{`
`23`		`- return (cb.d != 0) ? &a : nullptr;`
	`23`	`+ return (cb.d != 0) ? a : nullptr;`
`24`	`24`	`}`
`25`	`25`
`26`	`26`	`static inline __attribute__((always_inline))`
`@@ -31,7 +31,7 @@ thread uint3* select_input(thread uint3& gl_GlobalInvocationID, thread uint3& gl`
`31`	`31`
`32`	`32`	`kernel void main0(device foo& buf [[buffer(0)]], constant bar& cb [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])`
`33`	`33`	`{`
`34`		`- device foo* _44 = select_buffer(buf, cb);`
	`34`	`+ device foo* _44 = select_buffer(&buf, cb);`
`35`	`35`	`device foo* _65 = _44;`
`36`	`36`	`thread uint3* _45 = select_input(gl_GlobalInvocationID, gl_LocalInvocationID, cb);`
`37`	`37`	`device foo* _66 = _65;`
Original file line number	Diff line number	Diff line change
`@@ -16,14 +16,14 @@ struct bar`
`16`	`16`	`};`
`17`	`17`
`18`	`18`	`static inline __attribute__((always_inline))`
`19`		`-device int* _24(device foo& a, device bar& b, thread uint3& gl_GlobalInvocationID)`
	`19`	`+device int* _24(device foo* a, device bar* b, thread uint3& gl_GlobalInvocationID)`
`20`	`20`	`{`
`21`		`- return (gl_GlobalInvocationID.x != 0u) ? &a.a : &b.b;`
	`21`	`+ return (gl_GlobalInvocationID.x != 0u) ? &a->a : &b->b;`
`22`	`22`	`}`
`23`	`23`
`24`	`24`	`kernel void main0(device foo& x [[buffer(0)]], device bar& y [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])`
`25`	`25`	`{`
`26`		`- device int* _34 = _24(x, y, gl_GlobalInvocationID);`
	`26`	`+ device int* _34 = _24(&x, &y, gl_GlobalInvocationID);`
`27`	`27`	`device int* _33 = _34;`
`28`	`28`	`int _37 = x.a;`
`29`	`29`	`*_33 = 0;`
Original file line number	Diff line number	Diff line change
`@@ -742,6 +742,14 @@ bool Compiler::is_physical_pointer(const SPIRType &type) const`
`742`	`742`	`return type.op == OpTypePointer && type.storage == StorageClassPhysicalStorageBuffer;`
`743`	`743`	`}`
`744`	`744`
	`745`	`+bool Compiler::is_physical_or_buffer_pointer(const SPIRType &type) const`
	`746`	`+{`
	`747`	`+ return type.op == OpTypePointer &&`
	`748`	`+ (type.storage == StorageClassPhysicalStorageBuffer \|\| type.storage == StorageClassUniform \|\|`
	`749`	`+ type.storage == StorageClassStorageBuffer \|\| type.storage == StorageClassWorkgroup \|\|`
	`750`	`+ type.storage == StorageClassPushConstant);`
	`751`	`+}`
	`752`	`+`
`745`	`753`	`bool Compiler::is_physical_pointer_to_buffer_block(const SPIRType &type) const`
`746`	`754`	`{`
`747`	`755`	`return is_physical_pointer(type) && get_pointee_type(type).self == type.parent_type &&`