AuroraMiddleware/SPIRV-Cross
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Merge pull request #2215 from KhronosGroup/fix-2210

Browse Source 
GLSL: Fix and workaround some awkward code patterns
This commit is contained in:
Hans-Kristian Arntzen 2023-10-11 14:20:11 +02:00 committed by GitHub
commit 082b1df25f
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
8 changed files with 137 additions and 8 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

24
reference/opt/shaders-hlsl/comp/ssbo-store-array.comp Normal file
View File

@ -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();
}

13
reference/shaders-no-opt/asm/comp/block-undef.noeliminate.invalid.asm.comp Normal file
View File

@ -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()
{
}

12
reference/shaders-no-opt/asm/comp/empty-ubo-block.noeliminate.asm.comp Normal file
View File

@ -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()
{
}

34
shaders-no-opt/asm/comp/block-undef.noeliminate.invalid.asm.comp Normal file
View File

@ -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

21
shaders-no-opt/asm/comp/empty-ubo-block.noeliminate.asm.comp Normal file
View File

@ -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

21
spirv_cross.cpp
View File

@ -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;
}

1
spirv_cross.hpp
View File

@ -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;

19
spirv_glsl.cpp
View File

@ -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));