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

Deal with illegal names in types as well.

Browse Source 
- Fixes issue with clip_distance flattening in MSL where member to
  flatten from would come from to_member_name, where it should have used
  the builtin name directly. This member name was modified by this patch
  and broke clip distance test shaders.

- Some cleanups with ir.meta, use ir.find_meta instead to not create
  unnecessary hashmap nodes.
This commit is contained in:
Hans-Kristian Arntzen 2020-01-16 10:28:54 +01:00
parent 79700d5412
commit f79c1e2fed
10 changed files with 317 additions and 24 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

22
reference/shaders-hlsl-no-opt/comp/illegal-struct-name.asm.comp Normal file
View File

@ -0,0 +1,22 @@
struct Foo
{
float _abs;
};
RWByteAddressBuffer _7 : register(u0);
void comp_main()
{
Foo _24;
_24._abs = asfloat(_7.Load(0));
Foo f;
f._abs = _24._abs;
int _abs = 10;
_7.Store(4, asuint(f._abs));
}
[numthreads(1, 1, 1)]
void main()
{
comp_main();
}

29
reference/shaders-msl-no-opt/comp/illegal-struct-name.asm.comp Normal file
View File

@ -0,0 +1,29 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct Foo
{
float _abs;
};
struct Foo_1
{
float _abs;
};
struct SSBO
{
Foo_1 foo;
Foo_1 foo2;
};
kernel void main0(device SSBO& _7 [[buffer(0)]])
{
Foo f;
f._abs = _7.foo._abs;
int _abs = 10;
_7.foo2._abs = f._abs;
}

22
reference/shaders-no-opt/comp/illegal-struct-name.asm.comp Normal file
View File

@ -0,0 +1,22 @@
#version 450
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
struct Foo
{
float _abs;
};
layout(binding = 0, std430) buffer SSBO
{
Foo foo;
Foo foo2;
} _7;
void main()
{
Foo f;
f._abs = _7.foo._abs;
int _abs = 10;
_7.foo2._abs = f._abs;
}

62
shaders-hlsl-no-opt/comp/illegal-struct-name.asm.comp Normal file
View File

@ -0,0 +1,62 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 8
; Bound: 31
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
OpSource GLSL 450
OpName %main "main"
OpName %Foo "Foo"
OpMemberName %Foo 0 "abs"
OpName %f "f"
OpName %Foo_0 "Foo"
OpMemberName %Foo_0 0 "abs"
OpName %SSBO "SSBO"
OpMemberName %SSBO 0 "foo"
OpMemberName %SSBO 1 "foo2"
OpName %_ ""
OpName %linear "abs"
OpMemberDecorate %Foo_0 0 Offset 0
OpMemberDecorate %SSBO 0 Offset 0
OpMemberDecorate %SSBO 1 Offset 4
OpDecorate %SSBO BufferBlock
OpDecorate %_ DescriptorSet 0
OpDecorate %_ Binding 0
%void = OpTypeVoid
%3 = OpTypeFunction %void
%float = OpTypeFloat 32
%Foo = OpTypeStruct %float
%_ptr_Function_Foo = OpTypePointer Function %Foo
%Foo_0 = OpTypeStruct %float
%SSBO = OpTypeStruct %Foo_0 %Foo_0
%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
%_ = OpVariable %_ptr_Uniform_SSBO Uniform
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%_ptr_Uniform_Foo_0 = OpTypePointer Uniform %Foo_0
%_ptr_Function_float = OpTypePointer Function %float
%_ptr_Function_int = OpTypePointer Function %int
%int_10 = OpConstant %int 10
%int_1 = OpConstant %int 1
%_ptr_Uniform_float = OpTypePointer Uniform %float
%main = OpFunction %void None %3
%5 = OpLabel
%f = OpVariable %_ptr_Function_Foo Function
%linear = OpVariable %_ptr_Function_int Function
%17 = OpAccessChain %_ptr_Uniform_Foo_0 %_ %int_0
%18 = OpLoad %Foo_0 %17
%19 = OpCompositeExtract %float %18 0
%21 = OpAccessChain %_ptr_Function_float %f %int_0
OpStore %21 %19
OpStore %linear %int_10
%26 = OpLoad %Foo %f
%27 = OpAccessChain %_ptr_Uniform_Foo_0 %_ %int_1
%28 = OpCompositeExtract %float %26 0
%30 = OpAccessChain %_ptr_Uniform_float %27 %int_0
OpStore %30 %28
OpReturn
OpFunctionEnd

62
shaders-msl-no-opt/comp/illegal-struct-name.asm.comp Normal file
View File

@ -0,0 +1,62 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 8
; Bound: 31
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
OpSource GLSL 450
OpName %main "main"
OpName %Foo "Foo"
OpMemberName %Foo 0 "abs"
OpName %f "f"
OpName %Foo_0 "Foo"
OpMemberName %Foo_0 0 "abs"
OpName %SSBO "SSBO"
OpMemberName %SSBO 0 "foo"
OpMemberName %SSBO 1 "foo2"
OpName %_ ""
OpName %linear "abs"
OpMemberDecorate %Foo_0 0 Offset 0
OpMemberDecorate %SSBO 0 Offset 0
OpMemberDecorate %SSBO 1 Offset 4
OpDecorate %SSBO BufferBlock
OpDecorate %_ DescriptorSet 0
OpDecorate %_ Binding 0
%void = OpTypeVoid
%3 = OpTypeFunction %void
%float = OpTypeFloat 32
%Foo = OpTypeStruct %float
%_ptr_Function_Foo = OpTypePointer Function %Foo
%Foo_0 = OpTypeStruct %float
%SSBO = OpTypeStruct %Foo_0 %Foo_0
%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
%_ = OpVariable %_ptr_Uniform_SSBO Uniform
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%_ptr_Uniform_Foo_0 = OpTypePointer Uniform %Foo_0
%_ptr_Function_float = OpTypePointer Function %float
%_ptr_Function_int = OpTypePointer Function %int
%int_10 = OpConstant %int 10
%int_1 = OpConstant %int 1
%_ptr_Uniform_float = OpTypePointer Uniform %float
%main = OpFunction %void None %3
%5 = OpLabel
%f = OpVariable %_ptr_Function_Foo Function
%linear = OpVariable %_ptr_Function_int Function
%17 = OpAccessChain %_ptr_Uniform_Foo_0 %_ %int_0
%18 = OpLoad %Foo_0 %17
%19 = OpCompositeExtract %float %18 0
%21 = OpAccessChain %_ptr_Function_float %f %int_0
OpStore %21 %19
OpStore %linear %int_10
%26 = OpLoad %Foo %f
%27 = OpAccessChain %_ptr_Uniform_Foo_0 %_ %int_1
%28 = OpCompositeExtract %float %26 0
%30 = OpAccessChain %_ptr_Uniform_float %27 %int_0
OpStore %30 %28
OpReturn
OpFunctionEnd

62
shaders-no-opt/comp/illegal-struct-name.asm.comp Normal file
View File

@ -0,0 +1,62 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 8
; Bound: 31
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
OpSource GLSL 450
OpName %main "main"
OpName %Foo "Foo"
OpMemberName %Foo 0 "abs"
OpName %f "f"
OpName %Foo_0 "Foo"
OpMemberName %Foo_0 0 "abs"
OpName %SSBO "SSBO"
OpMemberName %SSBO 0 "foo"
OpMemberName %SSBO 1 "foo2"
OpName %_ ""
OpName %linear "abs"
OpMemberDecorate %Foo_0 0 Offset 0
OpMemberDecorate %SSBO 0 Offset 0
OpMemberDecorate %SSBO 1 Offset 4
OpDecorate %SSBO BufferBlock
OpDecorate %_ DescriptorSet 0
OpDecorate %_ Binding 0
%void = OpTypeVoid
%3 = OpTypeFunction %void
%float = OpTypeFloat 32
%Foo = OpTypeStruct %float
%_ptr_Function_Foo = OpTypePointer Function %Foo
%Foo_0 = OpTypeStruct %float
%SSBO = OpTypeStruct %Foo_0 %Foo_0
%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
%_ = OpVariable %_ptr_Uniform_SSBO Uniform
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%_ptr_Uniform_Foo_0 = OpTypePointer Uniform %Foo_0
%_ptr_Function_float = OpTypePointer Function %float
%_ptr_Function_int = OpTypePointer Function %int
%int_10 = OpConstant %int 10
%int_1 = OpConstant %int 1
%_ptr_Uniform_float = OpTypePointer Uniform %float
%main = OpFunction %void None %3
%5 = OpLabel
%f = OpVariable %_ptr_Function_Foo Function
%linear = OpVariable %_ptr_Function_int Function
%17 = OpAccessChain %_ptr_Uniform_Foo_0 %_ %int_0
%18 = OpLoad %Foo_0 %17
%19 = OpCompositeExtract %float %18 0
%21 = OpAccessChain %_ptr_Function_float %f %int_0
OpStore %21 %19
OpStore %linear %int_10
%26 = OpLoad %Foo %f
%27 = OpAccessChain %_ptr_Uniform_Foo_0 %_ %int_1
%28 = OpCompositeExtract %float %26 0
%30 = OpAccessChain %_ptr_Uniform_float %27 %int_0
OpStore %30 %28
OpReturn
OpFunctionEnd

39
spirv_glsl.cpp
View File

@ -2241,6 +2241,36 @@ void CompilerGLSL::emit_entry_point_declarations()
{
}
void CompilerGLSL::replace_illegal_names(const unordered_set<string> &keywords)
{
ir.for_each_typed_id<SPIRVariable>([&](uint32_t, const SPIRVariable &var) {
if (is_hidden_variable(var))
return;
auto *meta = ir.find_meta(var.self);
if (!meta)
return;
auto &m = meta->decoration;
if (m.alias.compare(0, 3, "gl_") == 0 || keywords.find(m.alias) != end(keywords))
m.alias = join("_", m.alias);
});
ir.for_each_typed_id<SPIRType>([&](uint32_t, const SPIRType &type) {
auto *meta = ir.find_meta(type.self);
if (!meta)
return;
auto &m = meta->decoration;
if (m.alias.compare(0, 3, "gl_") == 0 || keywords.find(m.alias) != end(keywords))
m.alias = join("_", m.alias);
for (auto &memb : meta->members)
if (memb.alias.compare(0, 3, "gl_") == 0 || keywords.find(memb.alias) != end(keywords))
memb.alias = join("_", memb.alias);
});
}
void CompilerGLSL::replace_illegal_names()
{
// clang-format off
@ -2295,14 +2325,7 @@ void CompilerGLSL::replace_illegal_names()
};
// clang-format on
ir.for_each_typed_id<SPIRVariable>([&](uint32_t, const SPIRVariable &var) {
if (!is_hidden_variable(var))
{
auto &m = ir.meta[var.self].decoration;
if (m.alias.compare(0, 3, "gl_") == 0 || keywords.find(m.alias) != end(keywords))
m.alias = join("_", m.alias);
}
});
replace_illegal_names(keywords);
}
void CompilerGLSL::replace_fragment_output(SPIRVariable &var)

1
spirv_glsl.hpp
View File

@ -590,6 +590,7 @@ protected:
bool check_atomic_image(uint32_t id);
virtual void replace_illegal_names();
void replace_illegal_names(const std::unordered_set<std::string> &keywords);
virtual void emit_entry_point_declarations();
void replace_fragment_output(SPIRVariable &var);

10
spirv_hlsl.cpp
View File

@ -1113,15 +1113,7 @@ void CompilerHLSL::replace_illegal_names()
"line", "linear", "matrix", "point", "row_major", "sampler",
};
ir.for_each_typed_id<SPIRVariable>([&](uint32_t, SPIRVariable &var) {
if (!is_hidden_variable(var))
{
auto &m = ir.meta[var.self].decoration;
if (keywords.find(m.alias) != end(keywords))
m.alias = join("_", m.alias);
}
});
CompilerGLSL::replace_illegal_names(keywords);
CompilerGLSL::replace_illegal_names();
}

32
spirv_msl.cpp
View File

@ -1701,6 +1701,7 @@ void CompilerMSL::add_composite_variable_to_interface_block(StorageClass storage
set_name(var.self, builtin_to_glsl(builtin, StorageClassFunction));
bool flatten_from_ib_var = false;
string flatten_from_ib_mbr_name;
if (storage == StorageClassOutput && is_builtin && builtin == BuiltInClipDistance)
{
@ -1708,7 +1709,9 @@ void CompilerMSL::add_composite_variable_to_interface_block(StorageClass storage
uint32_t clip_array_mbr_idx = uint32_t(ib_type.member_types.size());
ib_type.member_types.push_back(get_variable_data_type_id(var));
set_member_decoration(ib_type.self, clip_array_mbr_idx, DecorationBuiltIn, BuiltInClipDistance);
set_member_name(ib_type.self, clip_array_mbr_idx, builtin_to_glsl(BuiltInClipDistance, StorageClassOutput));
flatten_from_ib_mbr_name = builtin_to_glsl(BuiltInClipDistance, StorageClassOutput);
set_member_name(ib_type.self, clip_array_mbr_idx, flatten_from_ib_mbr_name);
// When we flatten, we flatten directly from the "out" struct,
// not from a function variable.
@ -1817,7 +1820,7 @@ void CompilerMSL::add_composite_variable_to_interface_block(StorageClass storage
";");
}
else if (flatten_from_ib_var)
statement(ib_var_ref, ".", mbr_name, " = ", ib_var_ref, ".", to_name(var.self), "[", i, "];");
statement(ib_var_ref, ".", mbr_name, " = ", ib_var_ref, ".", flatten_from_ib_mbr_name, "[", i, "];");
else
statement(ib_var_ref, ".", mbr_name, " = ", to_name(var.self), "[", i, "];");
});
@ -1902,6 +1905,7 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass
usable_type = &get<SPIRType>(usable_type->parent_type);
bool flatten_from_ib_var = false;
string flatten_from_ib_mbr_name;
if (storage == StorageClassOutput && is_builtin && builtin == BuiltInClipDistance)
{
@ -1909,7 +1913,9 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass
uint32_t clip_array_mbr_idx = uint32_t(ib_type.member_types.size());
ib_type.member_types.push_back(mbr_type_id);
set_member_decoration(ib_type.self, clip_array_mbr_idx, DecorationBuiltIn, BuiltInClipDistance);
set_member_name(ib_type.self, clip_array_mbr_idx, builtin_to_glsl(BuiltInClipDistance, StorageClassOutput));
flatten_from_ib_mbr_name = builtin_to_glsl(BuiltInClipDistance, StorageClassOutput);
set_member_name(ib_type.self, clip_array_mbr_idx, flatten_from_ib_mbr_name);
// When we flatten, we flatten directly from the "out" struct,
// not from a function variable.
@ -1983,7 +1989,7 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass
entry_func.fixup_hooks_out.push_back([=, &var, &var_type]() {
if (flatten_from_ib_var)
{
statement(ib_var_ref, ".", mbr_name, " = ", ib_var_ref, ".", to_member_name(var_type, mbr_idx),
statement(ib_var_ref, ".", mbr_name, " = ", ib_var_ref, ".", flatten_from_ib_mbr_name,
"[", i, "];");
}
else
@ -10395,19 +10401,31 @@ void CompilerMSL::replace_illegal_names()
};
ir.for_each_typed_id<SPIRVariable>([&](uint32_t self, SPIRVariable &) {
auto &dec = ir.meta[self].decoration;
auto *meta = ir.find_meta(self);
if (!meta)
return;
auto &dec = meta->decoration;
if (keywords.find(dec.alias) != end(keywords))
dec.alias += "0";
});
ir.for_each_typed_id<SPIRFunction>([&](uint32_t self, SPIRFunction &) {
auto &dec = ir.meta[self].decoration;
auto *meta = ir.find_meta(self);
if (!meta)
return;
auto &dec = meta->decoration;
if (illegal_func_names.find(dec.alias) != end(illegal_func_names))
dec.alias += "0";
});
ir.for_each_typed_id<SPIRType>([&](uint32_t self, SPIRType &) {
for (auto &mbr_dec : ir.meta[self].members)
auto *meta = ir.find_meta(self);
if (!meta)
return;
for (auto &mbr_dec : meta->members)
if (keywords.find(mbr_dec.alias) != end(keywords))
mbr_dec.alias += "0";
});