Merge pull request KhronosGroup#2215 from KhronosGroup/fix-2210

GLSL: Fix and workaround some awkward code patterns

reference/opt/shaders-hlsl/comp/ssbo-store-array.comp

@@ -0,0 +1,24 @@
+struct Data
+{
+    uint arr[3];
+};
+
+static uint _34;
+
+RWByteAddressBuffer _13 : register(u0);
+
+void comp_main()
+{
+    uint _33[3] = { _34, _34, _34 };
+    [unroll]
+    for (int _0ident = 0; _0ident < 3; _0ident++)
+    {
+        _13.Store(_0ident * 4 + 0, _33[_0ident]);
+    }
+}
+
+[numthreads(1, 1, 1)]
+void main()
+{
+    comp_main();
+}

reference/shaders-no-opt/asm/comp/block-undef.noeliminate.invalid.asm.comp

@@ -0,0 +1,13 @@
+#version 460
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+struct FHitGroupRootConstants
+{
+    uint BaseInstanceIndex;
+    uint UserData;
+};
+
+void main()
+{
+}
+

reference/shaders-no-opt/asm/comp/empty-ubo-block.noeliminate.asm.comp

@@ -0,0 +1,12 @@
+#version 460
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(binding = 15, std140) uniform type_Primitive
+{
+    int empty_struct_member;
+} Primitive;
+
+void main()
+{
+}
+

shaders-no-opt/asm/comp/block-undef.noeliminate.invalid.asm.comp

@@ -0,0 +1,34 @@
+               OpCapability Shader
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpSource GLSL 460
+               OpName %main "main"
+               OpName %FHitGroupRootConstants "FHitGroupRootConstants"
+               OpMemberName %FHitGroupRootConstants 0 "BaseInstanceIndex"
+               OpMemberName %FHitGroupRootConstants 1 "UserData"
+               OpName %type_StructuredBuffer_FHitGroupRootConstants "type_StructuredBuffer_FHitGroupRootConstants"
+               OpMemberName %type_StructuredBuffer_FHitGroupRootConstants 0 "_m0"
+               OpName %HitGroupData "HitGroupData"
+               OpMemberDecorate %FHitGroupRootConstants 0 Offset 0
+               OpMemberDecorate %FHitGroupRootConstants 1 Offset 4
+               OpDecorate %_runtimearr_FHitGroupRootConstants ArrayStride 8
+               OpMemberDecorate %type_StructuredBuffer_FHitGroupRootConstants 0 NonWritable
+               OpMemberDecorate %type_StructuredBuffer_FHitGroupRootConstants 0 Offset 0
+               OpDecorate %type_StructuredBuffer_FHitGroupRootConstants BufferBlock
+               OpDecorate %HitGroupData DescriptorSet 0
+               OpDecorate %HitGroupData Binding 20
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+%FHitGroupRootConstants = OpTypeStruct %uint %uint
+%_runtimearr_FHitGroupRootConstants = OpTypeRuntimeArray %FHitGroupRootConstants
+%type_StructuredBuffer_FHitGroupRootConstants = OpTypeStruct %_runtimearr_FHitGroupRootConstants
+%_ptr_StorageBuffer_type_StructuredBuffer_FHitGroupRootConstants = OpTypePointer StorageBuffer %type_StructuredBuffer_FHitGroupRootConstants
+%HitGroupData = OpVariable %_ptr_StorageBuffer_type_StructuredBuffer_FHitGroupRootConstants StorageBuffer
+%728 = OpUndef %_ptr_StorageBuffer_type_StructuredBuffer_FHitGroupRootConstants
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+               OpReturn
+               OpFunctionEnd

shaders-no-opt/asm/comp/empty-ubo-block.noeliminate.asm.comp

@@ -0,0 +1,21 @@
+               OpCapability Shader
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpSource GLSL 460
+               OpName %main "main"
+               OpName %type_Primitive "type_Primitive"
+               OpName %Primitive "Primitive"
+               OpDecorate %type_Primitive Block
+               OpDecorate %Primitive DescriptorSet 0
+               OpDecorate %Primitive Binding 15
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+%type_Primitive = OpTypeStruct
+%_ptr_Uniform_type_Primitive = OpTypePointer Uniform %type_Primitive
+  %Primitive = OpVariable %_ptr_Uniform_type_Primitive Uniform
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+               OpReturn
+               OpFunctionEnd

spirv_cross.cpp

@@ -1048,20 +1048,25 @@ ShaderResources Compiler::get_shader_resources(const unordered_set<VariableID> *
 	return res;
 }
 
-bool Compiler::type_is_block_like(const SPIRType &type) const
+bool Compiler::type_is_top_level_block(const spirv_cross::SPIRType &type) const
 {
 	if (type.basetype != SPIRType::Struct)
 		return false;
+	return has_decoration(type.self, DecorationBlock) || has_decoration(type.self, DecorationBufferBlock);
+}
 
-	if (has_decoration(type.self, DecorationBlock) || has_decoration(type.self, DecorationBufferBlock))
-	{
+bool Compiler::type_is_block_like(const SPIRType &type) const
+{
+	if (type_is_top_level_block(type))
 		return true;
-	}
 
-	// Block-like types may have Offset decorations.
-	for (uint32_t i = 0; i < uint32_t(type.member_types.size()); i++)
-		if (has_member_decoration(type.self, i, DecorationOffset))
-			return true;
+	if (type.basetype == SPIRType::Struct)
+	{
+		// Block-like types may have Offset decorations.
+		for (uint32_t i = 0; i < uint32_t(type.member_types.size()); i++)
+			if (has_member_decoration(type.self, i, DecorationOffset))
+				return true;
+	}
 
 	return false;
 }

spirv_cross.hpp

@@ -1150,6 +1150,7 @@ class Compiler
 	bool type_is_top_level_pointer(const SPIRType &type) const;
 	bool type_is_top_level_array(const SPIRType &type) const;
 	bool type_is_block_like(const SPIRType &type) const;
+	bool type_is_top_level_block(const SPIRType &type) const;
 	bool type_is_opaque_value(const SPIRType &type) const;
 
 	bool reflection_ssbo_instance_name_is_significant() const;

spirv_glsl.cpp

@@ -2465,6 +2465,10 @@ void CompilerGLSL::emit_buffer_block_native(const SPIRVariable &var)
 		i++;
 	}
 
+	// Don't declare empty blocks in GLSL, this is not allowed.
+	if (type_is_empty(type) && !backend.supports_empty_struct)
+		statement("int empty_struct_member;");
+
 	// var.self can be used as a backup name for the block name,
 	// so we need to make sure we don't disturb the name here on a recompile.
 	// It will need to be reset if we have to recompile.
@@ -2803,6 +2807,9 @@ string CompilerGLSL::constant_value_macro_name(uint32_t id)
 void CompilerGLSL::emit_specialization_constant_op(const SPIRConstantOp &constant)
 {
 	auto &type = get<SPIRType>(constant.basetype);
+	// This will break. It is bogus and should not be legal.
+	if (type_is_top_level_block(type))
+		return;
 	add_resource_name(constant.self);
 	auto name = to_name(constant.self);
 	statement("const ", variable_decl(type, name), " = ", constant_op_expression(constant), ";");
@@ -2832,6 +2839,10 @@ void CompilerGLSL::emit_constant(const SPIRConstant &constant)
 {
 	auto &type = get<SPIRType>(constant.constant_type);
 
+	// This will break. It is bogus and should not be legal.
+	if (type_is_top_level_block(type))
+		return;
+
 	SpecializationConstant wg_x, wg_y, wg_z;
 	ID workgroup_size_id = get_work_group_size_specialization_constants(wg_x, wg_y, wg_z);
 
@@ -3581,6 +3592,10 @@ void CompilerGLSL::emit_resources()
 		{
 			auto &id = ir.ids[id_];
 
+			// Skip declaring any bogus constants or undefs which use block types.
+			// We don't declare block types directly, so this will never work.
+			// Should not be legal SPIR-V, so this is considered a workaround.
+
 			if (id.get_type() == TypeConstant)
 			{
 				auto &c = id.get<SPIRConstant>();
@@ -3638,6 +3653,10 @@ void CompilerGLSL::emit_resources()
 				if (type.basetype == SPIRType::Void)
 					return;
 
+				// This will break. It is bogus and should not be legal.
+				if (type_is_top_level_block(type))
+					return;
+
 				string initializer;
 				if (options.force_zero_initialized_variables && type_can_zero_initialize(type))
 					initializer = join(" = ", to_zero_initialized_expression(undef.basetype));