MSL: Fixes from review for SPV_KHR_physical_storage_buffer extension.

- For physical storage pointers, emit forward reference
  declarations in addition to definitions.
- Minor syntax corrections.

reference/opt/shaders-msl/comp/buffer_device_address-recursive-struct-pointers.msl2.comp

@@ -3,7 +3,9 @@
 
 using namespace metal;
 
-struct t21
+struct t21;
+
+struct t24
 {
     int4 m0[2];
     int m1;
@@ -12,7 +14,7 @@ struct t21
     float2x4 m4;
 };
 
-struct t24
+struct t21
 {
     int4 m0[2];
     int m1;

reference/opt/shaders-msl/comp/buffer_device_address.msl2.comp

@@ -3,22 +3,25 @@
 
 using namespace metal;
 
+struct Position;
+struct PositionReferences;
+
 struct Position
 {
     float2 positions[1];
 };
 
-struct PositionReferences
-{
-    device Position* buffers[1];
-};
-
 struct Registers
 {
     device PositionReferences* references;
     float fract_time;
 };
 
+struct PositionReferences
+{
+    device Position* buffers[1];
+};
+
 constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(8u, 8u, 1u);
 
 kernel void main0(constant Registers& registers [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_NumWorkGroups [[threadgroups_per_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])

reference/opt/shaders-msl/vert/buffer_device_address.msl2.vert

@@ -3,22 +3,25 @@
 
 using namespace metal;
 
+struct Position;
+struct PositionReferences;
+
 struct Position
 {
     float2 positions[1];
 };
 
-struct PositionReferences
-{
-    device Position* buffers[1];
-};
-
 struct Registers
 {
     float4x4 view_projection;
     device PositionReferences* references;
 };
 
+struct PositionReferences
+{
+    device Position* buffers[1];
+};
+
 struct main0_out
 {
     float4 out_color [[user(locn0)]];

reference/shaders-msl/comp/buffer_device_address-recursive-struct-pointers.msl2.comp

@@ -3,7 +3,9 @@
 
 using namespace metal;
 
-struct t21
+struct t21;
+
+struct t24
 {
     int4 m0[2];
     int m1;
@@ -12,7 +14,7 @@ struct t21
     float2x4 m4;
 };
 
-struct t24
+struct t21
 {
     int4 m0[2];
     int m1;

reference/shaders-msl/comp/buffer_device_address.msl2.comp

@@ -3,22 +3,25 @@
 
 using namespace metal;
 
+struct Position;
+struct PositionReferences;
+
 struct Position
 {
     float2 positions[1];
 };
 
-struct PositionReferences
-{
-    device Position* buffers[1];
-};
-
 struct Registers
 {
     device PositionReferences* references;
     float fract_time;
 };
 
+struct PositionReferences
+{
+    device Position* buffers[1];
+};
+
 constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(8u, 8u, 1u);
 
 kernel void main0(constant Registers& registers [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_NumWorkGroups [[threadgroups_per_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])

reference/shaders-msl/vert/buffer_device_address.msl2.vert

@@ -3,22 +3,25 @@
 
 using namespace metal;
 
+struct Position;
+struct PositionReferences;
+
 struct Position
 {
     float2 positions[1];
 };
 
-struct PositionReferences
-{
-    device Position* buffers[1];
-};
-
 struct Registers
 {
     float4x4 view_projection;
     device PositionReferences* references;
 };
 
+struct PositionReferences
+{
+    device Position* buffers[1];
+};
+
 struct main0_out
 {
     float4 out_color [[user(locn0)]];

spirv_common.hpp

@@ -591,7 +591,6 @@ struct SPIRType : IVariant
 	uint32_t pointer_depth = 0;
 	bool pointer = false;
 	bool forward_pointer = false;
-	bool was_forward_referenced = false;
 
 	spv::StorageClass storage = spv::StorageClassGeneric;
 

spirv_msl.cpp

@@ -6804,6 +6804,25 @@ void CompilerMSL::emit_specialization_constants_and_structs()
 	// these types for purpose of iterating over them in ir.ids_for_type and friends.
 	auto loop_lock = ir.create_loop_soft_lock();
 
+	// Physical storage buffer pointers can have cyclical references,
+	// so emit forward declarations of them before other structs.
+	// Ignore type_id because we want the underlying struct type from the pointer.
+	ir.for_each_typed_id<SPIRType>([&](uint32_t /* type_id */, const SPIRType &type) {
+		if (type.basetype == SPIRType::Struct &&
+			type.pointer && type.storage == StorageClassPhysicalStorageBuffer &&
+			declared_structs.count(type.self) == 0)
+		{
+			statement("struct ", to_name(type.self), ";");
+			declared_structs.insert(type.self);
+			emitted = true;
+		}
+	});
+	if (emitted)
+		statement("");
+
+	emitted = false;
+	declared_structs.clear();
+
 	for (auto &id_ : ir.ids_for_constant_or_type)
 	{
 		auto &id = ir.ids[id_];
@@ -6880,9 +6899,7 @@ void CompilerMSL::emit_specialization_constants_and_structs()
 			auto &type = id.get<SPIRType>();
 			TypeID type_id = type.self;
 
-			// Forward references can change the declaration dependency order.
-			// Ensure that change is considered when emitting MSL structs, which must be emitted in dependency order.
-			bool is_struct = (type.basetype == SPIRType::Struct) && type.array.empty() && (!type.pointer || type.was_forward_referenced);
+			bool is_struct = (type.basetype == SPIRType::Struct) && type.array.empty() && !type.pointer;
 			bool is_block =
 			    has_decoration(type.self, DecorationBlock) || has_decoration(type.self, DecorationBufferBlock);
 
@@ -13543,7 +13560,7 @@ string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id)
 		const char *restrict_kw;
 
 		auto type_address_space = get_type_address_space(type, id);
-		const auto* p_parent_type = &get<SPIRType>(type.parent_type);
+		const auto *p_parent_type = &get<SPIRType>(type.parent_type);
 
 		// Work around C pointer qualifier rules. If glsl_type is a pointer type as well
 		// we'll need to emit the address space to the right.
@@ -13556,7 +13573,7 @@ string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id)
 		{
 			// Since this is not a pointer-to-pointer, ensure we've dug down to the base type.
 			// Some situations chain pointers even though they are not formally pointers-of-pointers.
-			while(type_is_pointer(*p_parent_type))
+			while (type_is_pointer(*p_parent_type))
 				p_parent_type = &get<SPIRType>(p_parent_type->parent_type);
 
 			type_name = join(type_address_space, " ", type_to_glsl(*p_parent_type, id));

spirv_parser.cpp

@@ -659,16 +659,6 @@ void Parser::parse(const Instruction &instruction)
 	{
 		uint32_t id = ops[0];
 
-		// The type may have previously been declared as an OpTypeForwardPointer.
-		// If so, we want to retain that knowledge, in case content needs to be emitted
-		// relative to when the forward pointer was declared, not when the pointer is
-		// declared. SPIR-V allows composites to contain a type that has only been
-		// declared as a forward pointer when the composite is defined, but in many
-		// languages, it is necessary to emit the component before the composite.
-		// Retrieve this info *before* overwriting the type info for the id below.
-		auto *fwd_ptr = maybe_get<SPIRType>(id);
-		bool was_fwd_ptr = fwd_ptr && fwd_ptr->forward_pointer;
-
 		// Very rarely, we might receive a FunctionPrototype here.
 		// We won't be able to compile it, but we shouldn't crash when parsing.
 		// We should be able to reflect.
@@ -681,7 +671,6 @@ void Parser::parse(const Instruction &instruction)
 		ptrbase.pointer = true;
 		ptrbase.pointer_depth++;
 		ptrbase.storage = static_cast<StorageClass>(ops[1]);
-		ptrbase.was_forward_referenced = was_fwd_ptr;
 
 		if (ptrbase.storage == StorageClassAtomicCounter)
 			ptrbase.basetype = SPIRType::AtomicCounter;