MSL: Fix crash where variable storage buffer pointers are passed down.

Only deal with readonly decoration for actual block types.

reference/shaders-msl-no-opt/asm/comp/storage-buffer-pointer-argument.asm.comp

@@ -0,0 +1,27 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct SSBO
+{
+    float a;
+};
+
+struct SSBORead
+{
+    float b;
+};
+
+void copy_out(device float& A, device const float& B)
+{
+    A = B;
+}
+
+kernel void main0(device SSBO& _7 [[buffer(0)]], const device SSBORead& _9 [[buffer(1)]])
+{
+    copy_out(_7.a, _9.b);
+}
+

shaders-msl-no-opt/asm/comp/storage-buffer-pointer-argument.asm.comp

@@ -0,0 +1,63 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Khronos Glslang Reference Front End; 7
+; Bound: 30
+; Schema: 0
+               OpCapability Shader
+               OpCapability VariablePointersStorageBuffer
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpSource GLSL 450
+               OpName %main "main"
+               OpName %copy_out_f1_f1_ "copy_out(f1;f1;"
+               OpName %A "A"
+               OpName %B "B"
+               OpName %SSBO "SSBO"
+               OpMemberName %SSBO 0 "a"
+               OpName %_ ""
+               OpName %SSBORead "SSBORead"
+               OpMemberName %SSBORead 0 "b"
+               OpName %__0 ""
+               OpMemberDecorate %SSBO 0 NonReadable
+               OpMemberDecorate %SSBO 0 Offset 0
+               OpDecorate %SSBO Block
+               OpDecorate %_ DescriptorSet 0
+               OpDecorate %_ Binding 0
+               OpMemberDecorate %SSBORead 0 NonWritable
+               OpMemberDecorate %SSBORead 0 Offset 0
+               OpDecorate %SSBORead Block
+               OpDecorate %__0 DescriptorSet 0
+               OpDecorate %__0 Binding 1
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+%_ptr_Function_float = OpTypePointer Function %float
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+          %8 = OpTypeFunction %void %_ptr_StorageBuffer_float %_ptr_StorageBuffer_float
+       %SSBO = OpTypeStruct %float
+%_ptr_StorageBuffer_SSBO = OpTypePointer StorageBuffer %SSBO
+          %_ = OpVariable %_ptr_StorageBuffer_SSBO StorageBuffer
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+   %SSBORead = OpTypeStruct %float
+%_ptr_StorageBuffer_SSBORead = OpTypePointer StorageBuffer %SSBORead
+        %__0 = OpVariable %_ptr_StorageBuffer_SSBORead StorageBuffer
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+      %param = OpVariable %_ptr_Function_float Function
+    %param_0 = OpVariable %_ptr_Function_float Function
+         %25 = OpAccessChain %_ptr_StorageBuffer_float %_ %int_0
+         %26 = OpAccessChain %_ptr_StorageBuffer_float %__0 %int_0
+         %27 = OpFunctionCall %void %copy_out_f1_f1_ %25 %26
+               OpReturn
+               OpFunctionEnd
+%copy_out_f1_f1_ = OpFunction %void None %8
+          %A = OpFunctionParameter %_ptr_StorageBuffer_float
+          %B = OpFunctionParameter %_ptr_StorageBuffer_float
+         %12 = OpLabel
+         %13 = OpLoad %float %B
+               OpStore %A %13
+               OpReturn
+               OpFunctionEnd

spirv_msl.cpp

@@ -5406,7 +5406,12 @@ string CompilerMSL::get_argument_address_space(const SPIRVariable &argument)
 
 	case StorageClassStorageBuffer:
 	{
-		bool readonly = ir.get_buffer_block_flags(argument).get(DecorationNonWritable);
+		// For arguments from variable pointers, we use the write count deduction, so
+		// we should not assume any constness here. Only for global SSBOs.
+		bool readonly = false;
+		if (has_decoration(type.self, DecorationBlock))
+			readonly = ir.get_buffer_block_flags(argument).get(DecorationNonWritable);
+
 		return readonly ? "const device" : "device";
 	}