MSL: Don't dereference forwarded copies of OpVariable pointers.

Fixes two more CTS tests under
`dEQP-VK.spirv_assembly.instruction.compute.empty_struct.copying`.

reference/opt/shaders-msl/asm/comp/copy-object-ssbo-to-ssbo.asm.comp

@@ -0,0 +1,23 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct _19
+{
+};
+struct _5
+{
+    int _m0;
+    _19 _m1;
+    char _m2_pad[4];
+    _19 _m2;
+    char _m3_pad[4];
+    int _m3;
+};
+
+kernel void main0(device _5& _3 [[buffer(0)]], device _5& _4 [[buffer(1)]])
+{
+    _4 = _3;
+}
+

reference/opt/shaders-msl/asm/comp/copy-object-ubo-to-ssbo.asm.comp

@@ -0,0 +1,24 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct _19
+{
+};
+struct _5
+{
+    int _m0;
+    char _m1_pad[12];
+    _19 _m1;
+    char _m2_pad[16];
+    _19 _m2;
+    char _m3_pad[16];
+    int _m3;
+};
+
+kernel void main0(constant _5& _3 [[buffer(0)]], device _5& _4 [[buffer(1)]])
+{
+    _4 = _3;
+}
+

reference/shaders-msl/asm/comp/copy-object-ssbo-to-ssbo.asm.comp

@@ -0,0 +1,23 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct _19
+{
+};
+struct _5
+{
+    int _m0;
+    _19 _m1;
+    char _m2_pad[4];
+    _19 _m2;
+    char _m3_pad[4];
+    int _m3;
+};
+
+kernel void main0(device _5& _3 [[buffer(0)]], device _5& _4 [[buffer(1)]])
+{
+    _4 = _3;
+}
+

reference/shaders-msl/asm/comp/copy-object-ubo-to-ssbo.asm.comp

@@ -0,0 +1,24 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct _19
+{
+};
+struct _5
+{
+    int _m0;
+    char _m1_pad[12];
+    _19 _m1;
+    char _m2_pad[16];
+    _19 _m2;
+    char _m3_pad[16];
+    int _m3;
+};
+
+kernel void main0(constant _5& _3 [[buffer(0)]], device _5& _4 [[buffer(1)]])
+{
+    _4 = _3;
+}
+

shaders-msl/asm/comp/copy-object-ssbo-to-ssbo.asm.comp

@@ -0,0 +1,43 @@
+OpCapability Shader
+OpExtension "SPV_KHR_storage_buffer_storage_class"
+OpMemoryModel Logical GLSL450
+OpEntryPoint GLCompute %main "main" %var_id
+OpExecutionMode %main LocalSize 1 1 1
+OpDecorate %var_id BuiltIn GlobalInvocationId
+OpDecorate %var_input Binding 0
+OpDecorate %var_input DescriptorSet 0
+OpDecorate %var_outdata Binding 1
+OpDecorate %var_outdata DescriptorSet 0
+OpMemberDecorate %type_container_struct 0 Offset 0
+OpMemberDecorate %type_container_struct 1 Offset 4
+OpMemberDecorate %type_container_struct 2 Offset 8
+OpMemberDecorate %type_container_struct 3 Offset 12
+OpDecorate %type_container_struct Block
+%bool      = OpTypeBool
+%void      = OpTypeVoid
+%voidf     = OpTypeFunction %void
+%u32       = OpTypeInt 32 0
+%i32       = OpTypeInt 32 1
+%f32       = OpTypeFloat 32
+%uvec3     = OpTypeVector %u32 3
+%fvec3     = OpTypeVector %f32 3
+%uvec3ptr  = OpTypePointer Input %uvec3
+%i32ptr    = OpTypePointer Uniform %i32
+%f32ptr    = OpTypePointer Uniform %f32
+%i32arr    = OpTypeRuntimeArray %i32
+%f32arr    = OpTypeRuntimeArray %f32
+%type_empty_struct					= OpTypeStruct
+%type_container_struct				= OpTypeStruct %i32 %type_empty_struct %type_empty_struct %i32
+%type_container_struct_ubo_ptr		= OpTypePointer Uniform %type_container_struct
+%type_container_struct_ssbo_ptr	= OpTypePointer StorageBuffer %type_container_struct
+%var_id							= OpVariable %uvec3ptr Input
+%var_input						= OpVariable %type_container_struct_ssbo_ptr StorageBuffer
+%var_outdata					= OpVariable %type_container_struct_ssbo_ptr StorageBuffer
+
+%main								= OpFunction %void None %voidf
+%label								= OpLabel
+%input_copy					= OpCopyObject %type_container_struct_ssbo_ptr %var_input
+%result						= OpLoad %type_container_struct %input_copy
+OpStore %var_outdata %result
+OpReturn
+OpFunctionEnd

shaders-msl/asm/comp/copy-object-ubo-to-ssbo.asm.comp

@@ -0,0 +1,43 @@
+OpCapability Shader
+OpExtension "SPV_KHR_storage_buffer_storage_class"
+OpMemoryModel Logical GLSL450
+OpEntryPoint GLCompute %main "main" %var_id
+OpExecutionMode %main LocalSize 1 1 1
+OpDecorate %var_id BuiltIn GlobalInvocationId
+OpDecorate %var_input Binding 0
+OpDecorate %var_input DescriptorSet 0
+OpDecorate %var_outdata Binding 1
+OpDecorate %var_outdata DescriptorSet 0
+OpMemberDecorate %type_container_struct 0 Offset 0
+OpMemberDecorate %type_container_struct 1 Offset 16
+OpMemberDecorate %type_container_struct 2 Offset 32
+OpMemberDecorate %type_container_struct 3 Offset 48
+OpDecorate %type_container_struct Block
+%bool      = OpTypeBool
+%void      = OpTypeVoid
+%voidf     = OpTypeFunction %void
+%u32       = OpTypeInt 32 0
+%i32       = OpTypeInt 32 1
+%f32       = OpTypeFloat 32
+%uvec3     = OpTypeVector %u32 3
+%fvec3     = OpTypeVector %f32 3
+%uvec3ptr  = OpTypePointer Input %uvec3
+%i32ptr    = OpTypePointer Uniform %i32
+%f32ptr    = OpTypePointer Uniform %f32
+%i32arr    = OpTypeRuntimeArray %i32
+%f32arr    = OpTypeRuntimeArray %f32
+%type_empty_struct					= OpTypeStruct
+%type_container_struct				= OpTypeStruct %i32 %type_empty_struct %type_empty_struct %i32
+%type_container_struct_ubo_ptr		= OpTypePointer Uniform %type_container_struct
+%type_container_struct_ssbo_ptr	= OpTypePointer StorageBuffer %type_container_struct
+%var_id							= OpVariable %uvec3ptr Input
+%var_input						= OpVariable %type_container_struct_ubo_ptr Uniform
+%var_outdata					= OpVariable %type_container_struct_ssbo_ptr StorageBuffer
+
+%main								= OpFunction %void None %voidf
+%label								= OpLabel
+%input_copy					= OpCopyObject %type_container_struct_ubo_ptr %var_input
+%result						= OpLoad %type_container_struct %input_copy
+OpStore %var_outdata %result
+OpReturn
+OpFunctionEnd

spirv_glsl.cpp

@@ -10061,9 +10061,32 @@ bool CompilerGLSL::should_dereference(uint32_t id)
 	if (auto *var = maybe_get<SPIRVariable>(id))
 		return var->phi_variable;
 
-	// If id is an access chain, we should not dereference it.
 	if (auto *expr = maybe_get<SPIRExpression>(id))
-		return !expr->access_chain;
+	{
+		// If id is an access chain, we should not dereference it.
+		if (expr->access_chain)
+			return false;
+
+		// If id is a forwarded copy of a variable pointer, we should not dereference it.
+		SPIRVariable *var = nullptr;
+		while (expr->loaded_from && expression_is_forwarded(expr->self))
+		{
+			auto &src_type = expression_type(expr->loaded_from);
+			// To be a copy, the pointer and its source expression must be the
+			// same type. Can't check type.self, because for some reason that's
+			// usually the base type with pointers stripped off. This check is
+			// complex enough that I've hoisted it out of the while condition.
+			if (src_type.pointer != type.pointer || src_type.pointer_depth != type.pointer ||
+			    src_type.parent_type != type.parent_type)
+				break;
+			if ((var = maybe_get<SPIRVariable>(expr->loaded_from)))
+				break;
+			if (!(expr = maybe_get<SPIRExpression>(expr->loaded_from)))
+				break;
+		}
+
+		return !var || var->phi_variable;
+	}
 
 	// Otherwise, we should dereference this pointer expression.
 	return true;
@@ -11660,7 +11683,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 			// RHS expression is immutable, so just forward it.
 			// Copying these things really make no sense, but
 			// seems to be allowed anyways.
-			auto &e = set<SPIRExpression>(id, to_expression(rhs), result_type, true);
+			auto &e = emit_op(result_type, id, to_expression(rhs), true, true);
 			if (pointer)
 			{
 				auto *var = maybe_get_backing_variable(rhs);