diff --git a/reference/shaders-no-opt/asm/comp/block-undef.noeliminate.invalid.asm.comp b/reference/shaders-no-opt/asm/comp/block-undef.noeliminate.invalid.asm.comp new file mode 100644 index 00000000..dba99a04 --- /dev/null +++ b/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() +{ +} + diff --git a/shaders-no-opt/asm/comp/block-undef.noeliminate.invalid.asm.comp b/shaders-no-opt/asm/comp/block-undef.noeliminate.invalid.asm.comp new file mode 100644 index 00000000..9abef684 --- /dev/null +++ b/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 diff --git a/spirv_cross.cpp b/spirv_cross.cpp index c1c6a11d..7d3296f0 100644 --- a/spirv_cross.cpp +++ b/spirv_cross.cpp @@ -1048,20 +1048,25 @@ ShaderResources Compiler::get_shader_resources(const unordered_set * 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; } diff --git a/spirv_cross.hpp b/spirv_cross.hpp index 8f7ba447..8b85f7c5 100644 --- a/spirv_cross.hpp +++ b/spirv_cross.hpp @@ -1150,6 +1150,7 @@ protected: 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; diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index 83434e5a..6fac25d8 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -2807,6 +2807,9 @@ string CompilerGLSL::constant_value_macro_name(uint32_t id) void CompilerGLSL::emit_specialization_constant_op(const SPIRConstantOp &constant) { auto &type = get(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), ";"); @@ -2836,6 +2839,10 @@ void CompilerGLSL::emit_constant(const SPIRConstant &constant) { auto &type = get(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); @@ -3585,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(); @@ -3642,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));