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

Handle inout properly with split access chains.

Browse Source 
Found some other issues. Had some bugs with variable writes not properly
invalidating if writes came from split access chains.
This commit is contained in:
Hans-Kristian Arntzen 2018-05-11 10:14:20 +02:00
parent 97e38bcd38
commit 7eba247864
11 changed files with 201 additions and 55 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

47
reference/opt/shaders-msl/comp/struct-packing.comp
View File

@ -97,29 +97,30 @@ struct SSBO0
kernel void main0(device SSBO0& ssbo_140 [[buffer(0)]], device SSBO1& ssbo_430 [[buffer(1)]])
{
ssbo_430.content.m0s[0].a[0] = ssbo_140.content.m0s[0].a[0];
ssbo_430.content.m0s[0].b = ssbo_140.content.m0s[0].b;
ssbo_430.content.m1s[0].a = ssbo_140.content.m1s[0].a;
ssbo_430.content.m1s[0].b = ssbo_140.content.m1s[0].b;
ssbo_430.content.m2s[0].a[0] = ssbo_140.content.m2s[0].a[0];
ssbo_430.content.m2s[0].b = ssbo_140.content.m2s[0].b;
ssbo_430.content.m0.a[0] = ssbo_140.content.m0.a[0];
ssbo_430.content.m0.b = ssbo_140.content.m0.b;
ssbo_430.content.m1.a = ssbo_140.content.m1.a;
ssbo_430.content.m1.b = ssbo_140.content.m1.b;
ssbo_430.content.m2.a[0] = ssbo_140.content.m2.a[0];
ssbo_430.content.m2.b = ssbo_140.content.m2.b;
ssbo_430.content.m3.a = ssbo_140.content.m3.a;
ssbo_430.content.m3.b = ssbo_140.content.m3.b;
ssbo_430.content.m4 = ssbo_140.content.m4;
ssbo_430.content.m3s[0].c = ssbo_140.content.m3s[0].c;
ssbo_430.content.m3s[1].c = ssbo_140.content.m3s[1].c;
ssbo_430.content.m3s[2].c = ssbo_140.content.m3s[2].c;
ssbo_430.content.m3s[3].c = ssbo_140.content.m3s[3].c;
ssbo_430.content.m3s[4].c = ssbo_140.content.m3s[4].c;
ssbo_430.content.m3s[5].c = ssbo_140.content.m3s[5].c;
ssbo_430.content.m3s[6].c = ssbo_140.content.m3s[6].c;
ssbo_430.content.m3s[7].c = ssbo_140.content.m3s[7].c;
Content_1 _60 = ssbo_140.content;
ssbo_430.content.m0s[0].a[0] = _60.m0s[0].a[0];
ssbo_430.content.m0s[0].b = _60.m0s[0].b;
ssbo_430.content.m1s[0].a = _60.m1s[0].a;
ssbo_430.content.m1s[0].b = _60.m1s[0].b;
ssbo_430.content.m2s[0].a[0] = _60.m2s[0].a[0];
ssbo_430.content.m2s[0].b = _60.m2s[0].b;
ssbo_430.content.m0.a[0] = _60.m0.a[0];
ssbo_430.content.m0.b = _60.m0.b;
ssbo_430.content.m1.a = _60.m1.a;
ssbo_430.content.m1.b = _60.m1.b;
ssbo_430.content.m2.a[0] = _60.m2.a[0];
ssbo_430.content.m2.b = _60.m2.b;
ssbo_430.content.m3.a = _60.m3.a;
ssbo_430.content.m3.b = _60.m3.b;
ssbo_430.content.m4 = _60.m4;
ssbo_430.content.m3s[0].c = _60.m3s[0].c;
ssbo_430.content.m3s[1].c = _60.m3s[1].c;
ssbo_430.content.m3s[2].c = _60.m3s[2].c;
ssbo_430.content.m3s[3].c = _60.m3s[3].c;
ssbo_430.content.m3s[4].c = _60.m3s[4].c;
ssbo_430.content.m3s[5].c = _60.m3s[5].c;
ssbo_430.content.m3s[6].c = _60.m3s[6].c;
ssbo_430.content.m3s[7].c = _60.m3s[7].c;
ssbo_430.content.m1.a = ssbo_430.content.m3.a * ssbo_430.m6[1][1];
}

9
reference/opt/shaders/asm/geom/inout-split-access-chain-handle.asm.geom Normal file
View File

@ -0,0 +1,9 @@
#version 440
layout(triangles) in;
layout(max_vertices = 5, triangle_strip) out;
void main()
{
gl_Position = gl_in[0].gl_Position;
}

4
reference/opt/shaders/comp/struct-packing.comp
View File

@ -62,7 +62,7 @@ struct Content_1
S4 m3s[8];
};
layout(binding = 1, std430) buffer SSBO1
layout(binding = 1, std430) restrict buffer SSBO1
{
Content content;
Content content1[2];
@ -78,7 +78,7 @@ layout(binding = 1, std430) buffer SSBO1
float array[];
} ssbo_430;
layout(binding = 0, std140) buffer SSBO0
layout(binding = 0, std140) restrict buffer SSBO0
{
Content_1 content;
Content_1 content1[2];

47
reference/shaders-msl/comp/struct-packing.comp
View File

@ -120,29 +120,30 @@ struct SSBO0
kernel void main0(device SSBO0& ssbo_140 [[buffer(0)]], device SSBO1& ssbo_430 [[buffer(1)]])
{
ssbo_430.content.m0s[0].a[0] = ssbo_140.content.m0s[0].a[0];
ssbo_430.content.m0s[0].b = ssbo_140.content.m0s[0].b;
ssbo_430.content.m1s[0].a = ssbo_140.content.m1s[0].a;
ssbo_430.content.m1s[0].b = ssbo_140.content.m1s[0].b;
ssbo_430.content.m2s[0].a[0] = ssbo_140.content.m2s[0].a[0];
ssbo_430.content.m2s[0].b = ssbo_140.content.m2s[0].b;
ssbo_430.content.m0.a[0] = ssbo_140.content.m0.a[0];
ssbo_430.content.m0.b = ssbo_140.content.m0.b;
ssbo_430.content.m1.a = ssbo_140.content.m1.a;
ssbo_430.content.m1.b = ssbo_140.content.m1.b;
ssbo_430.content.m2.a[0] = ssbo_140.content.m2.a[0];
ssbo_430.content.m2.b = ssbo_140.content.m2.b;
ssbo_430.content.m3.a = ssbo_140.content.m3.a;
ssbo_430.content.m3.b = ssbo_140.content.m3.b;
ssbo_430.content.m4 = ssbo_140.content.m4;
ssbo_430.content.m3s[0].c = ssbo_140.content.m3s[0].c;
ssbo_430.content.m3s[1].c = ssbo_140.content.m3s[1].c;
ssbo_430.content.m3s[2].c = ssbo_140.content.m3s[2].c;
ssbo_430.content.m3s[3].c = ssbo_140.content.m3s[3].c;
ssbo_430.content.m3s[4].c = ssbo_140.content.m3s[4].c;
ssbo_430.content.m3s[5].c = ssbo_140.content.m3s[5].c;
ssbo_430.content.m3s[6].c = ssbo_140.content.m3s[6].c;
ssbo_430.content.m3s[7].c = ssbo_140.content.m3s[7].c;
Content_1 _60 = ssbo_140.content;
ssbo_430.content.m0s[0].a[0] = _60.m0s[0].a[0];
ssbo_430.content.m0s[0].b = _60.m0s[0].b;
ssbo_430.content.m1s[0].a = _60.m1s[0].a;
ssbo_430.content.m1s[0].b = _60.m1s[0].b;
ssbo_430.content.m2s[0].a[0] = _60.m2s[0].a[0];
ssbo_430.content.m2s[0].b = _60.m2s[0].b;
ssbo_430.content.m0.a[0] = _60.m0.a[0];
ssbo_430.content.m0.b = _60.m0.b;
ssbo_430.content.m1.a = _60.m1.a;
ssbo_430.content.m1.b = _60.m1.b;
ssbo_430.content.m2.a[0] = _60.m2.a[0];
ssbo_430.content.m2.b = _60.m2.b;
ssbo_430.content.m3.a = _60.m3.a;
ssbo_430.content.m3.b = _60.m3.b;
ssbo_430.content.m4 = _60.m4;
ssbo_430.content.m3s[0].c = _60.m3s[0].c;
ssbo_430.content.m3s[1].c = _60.m3s[1].c;
ssbo_430.content.m3s[2].c = _60.m3s[2].c;
ssbo_430.content.m3s[3].c = _60.m3s[3].c;
ssbo_430.content.m3s[4].c = _60.m3s[4].c;
ssbo_430.content.m3s[5].c = _60.m3s[5].c;
ssbo_430.content.m3s[6].c = _60.m3s[6].c;
ssbo_430.content.m3s[7].c = _60.m3s[7].c;
ssbo_430.content.m1.a = ssbo_430.content.m3.a * ssbo_430.m6[1][1];
}

23
reference/shaders/asm/geom/inout-split-access-chain-handle.asm.geom Normal file
View File

@ -0,0 +1,23 @@
#version 440
layout(triangles) in;
layout(max_vertices = 5, triangle_strip) out;
struct Data
{
vec4 ApiPerspectivePosition;
};
void Copy(inout Data inputStream[3])
{
inputStream[0].ApiPerspectivePosition = gl_in[0].gl_Position;
}
void main()
{
Data inputStream[3];
Data param[3] = inputStream;
Copy(param);
inputStream = param;
gl_Position = inputStream[0].ApiPerspectivePosition;
}

4
reference/shaders/comp/struct-packing.comp
View File

@ -85,7 +85,7 @@ struct Content_1
S4_1 m3s[8];
};
layout(binding = 1, std430) buffer SSBO1
layout(binding = 1, std430) restrict buffer SSBO1
{
Content content;
Content content1[2];
@ -101,7 +101,7 @@ layout(binding = 1, std430) buffer SSBO1
float array[];
} ssbo_430;
layout(binding = 0, std140) buffer SSBO0
layout(binding = 0, std140) restrict buffer SSBO0
{
Content_1 content;
Content_1 content1[2];

90
shaders/asm/geom/inout-split-access-chain-handle.asm.geom Normal file
View File

@ -0,0 +1,90 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 3
; Bound: 42
; Schema: 0
OpCapability Geometry
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Geometry %main "main" %gl_in %_
OpExecutionMode %main Triangles
OpExecutionMode %main Invocations 1
OpExecutionMode %main OutputTriangleStrip
OpExecutionMode %main OutputVertices 5
OpSource GLSL 440
OpName %main "main"
OpName %Data "Data"
OpMemberName %Data 0 "ApiPerspectivePosition"
OpName %Copy_struct_Data_vf41_3__ "Copy(struct-Data-vf41[3];"
OpName %inputStream "inputStream"
OpName %gl_PerVertex "gl_PerVertex"
OpMemberName %gl_PerVertex 0 "gl_Position"
OpMemberName %gl_PerVertex 1 "gl_PointSize"
OpMemberName %gl_PerVertex 2 "gl_ClipDistance"
OpName %gl_in "gl_in"
OpName %inputStream_0 "inputStream"
OpName %param "param"
OpName %gl_PerVertex_0 "gl_PerVertex"
OpMemberName %gl_PerVertex_0 0 "gl_Position"
OpMemberName %gl_PerVertex_0 1 "gl_PointSize"
OpMemberName %gl_PerVertex_0 2 "gl_ClipDistance"
OpName %_ ""
OpMemberDecorate %gl_PerVertex 0 BuiltIn Position
OpMemberDecorate %gl_PerVertex 1 BuiltIn PointSize
OpMemberDecorate %gl_PerVertex 2 BuiltIn ClipDistance
OpDecorate %gl_PerVertex Block
OpMemberDecorate %gl_PerVertex_0 0 BuiltIn Position
OpMemberDecorate %gl_PerVertex_0 1 BuiltIn PointSize
OpMemberDecorate %gl_PerVertex_0 2 BuiltIn ClipDistance
OpDecorate %gl_PerVertex_0 Block
%void = OpTypeVoid
%3 = OpTypeFunction %void
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%Data = OpTypeStruct %v4float
%uint = OpTypeInt 32 0
%uint_3 = OpConstant %uint 3
%_arr_Data_uint_3 = OpTypeArray %Data %uint_3
%_ptr_Function__Data = OpTypePointer Function %Data
%_ptr_Function__arr_Data_uint_3 = OpTypePointer Function %_arr_Data_uint_3
%13 = OpTypeFunction %void %_ptr_Function__arr_Data_uint_3
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%uint_1 = OpConstant %uint 1
%_arr_float_uint_1 = OpTypeArray %float %uint_1
%gl_PerVertex = OpTypeStruct %v4float %float %_arr_float_uint_1
%_arr_gl_PerVertex_uint_3 = OpTypeArray %gl_PerVertex %uint_3
%_ptr_Input__arr_gl_PerVertex_uint_3 = OpTypePointer Input %_arr_gl_PerVertex_uint_3
%gl_in = OpVariable %_ptr_Input__arr_gl_PerVertex_uint_3 Input
%_ptr_Input_v4float = OpTypePointer Input %v4float
%_ptr_Function_v4float = OpTypePointer Function %v4float
%gl_PerVertex_0 = OpTypeStruct %v4float %float %_arr_float_uint_1
%_ptr_Output_gl_PerVertex_0 = OpTypePointer Output %gl_PerVertex_0
%_ = OpVariable %_ptr_Output_gl_PerVertex_0 Output
%_ptr_Output_v4float = OpTypePointer Output %v4float
%main = OpFunction %void None %3
%5 = OpLabel
%inputStream_0 = OpVariable %_ptr_Function__arr_Data_uint_3 Function
%param = OpVariable %_ptr_Function__arr_Data_uint_3 Function
%32 = OpLoad %_arr_Data_uint_3 %inputStream_0
OpStore %param %32
%33 = OpFunctionCall %void %Copy_struct_Data_vf41_3__ %param
%34 = OpLoad %_arr_Data_uint_3 %param
OpStore %inputStream_0 %34
%59 = OpAccessChain %_ptr_Function__Data %inputStream_0 %int_0
%38 = OpAccessChain %_ptr_Function_v4float %59 %int_0
%39 = OpLoad %v4float %38
%41 = OpAccessChain %_ptr_Output_v4float %_ %int_0
OpStore %41 %39
OpReturn
OpFunctionEnd
%Copy_struct_Data_vf41_3__ = OpFunction %void None %13
%inputStream = OpFunctionParameter %_ptr_Function__arr_Data_uint_3
%16 = OpLabel
%26 = OpAccessChain %_ptr_Input_v4float %gl_in %int_0 %int_0
%27 = OpLoad %v4float %26
%28 = OpAccessChain %_ptr_Function__Data %inputStream %int_0
%29 = OpAccessChain %_ptr_Function_v4float %28 %int_0
OpStore %29 %27
OpReturn
OpFunctionEnd

4
shaders/comp/struct-packing.comp
View File

@ -44,7 +44,7 @@ struct Content
S4 m3s[8];
};
layout(binding = 1, std430) buffer SSBO1
layout(binding = 1, std430) restrict buffer SSBO1
{
Content content;
Content content1[2];
@ -61,7 +61,7 @@ layout(binding = 1, std430) buffer SSBO1
float array[];
} ssbo_430;
layout(binding = 0, std140) buffer SSBO0
layout(binding = 0, std140) restrict buffer SSBO0
{
Content content;
Content content1[2];

9
spirv_common.hpp
View File

@ -1111,9 +1111,10 @@ public:
void set(std::unique_ptr<IVariant> val, uint32_t new_type)
{
holder = std::move(val);
if (type != TypeNone && type != new_type)
if (!allow_type_rewrite && type != TypeNone && type != new_type)
SPIRV_CROSS_THROW("Overwriting a variant with new type.");
type = new_type;
allow_type_rewrite = false;
}
template <typename T>
@ -1154,9 +1155,15 @@ public:
type = TypeNone;
}
void set_allow_type_rewrite()
{
allow_type_rewrite = true;
}
private:
std::unique_ptr<IVariant> holder;
uint32_t type = TypeNone;
bool allow_type_rewrite = false;
};
template <typename T>

15
spirv_cross.cpp
View File

@ -104,7 +104,13 @@ bool Compiler::variable_storage_is_aliased(const SPIRVariable &v)
meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock);
bool image = type.basetype == SPIRType::Image;
bool counter = type.basetype == SPIRType::AtomicCounter;
bool is_restrict = meta[v.self].decoration.decoration_flags.get(DecorationRestrict);
bool is_restrict;
if (ssbo)
is_restrict = get_buffer_block_flags(v).get(DecorationRestrict);
else
is_restrict = has_decoration(v.self, DecorationRestrict);
return !is_restrict && (ssbo || image || counter);
}
@ -3191,6 +3197,9 @@ bool Compiler::DummySamplerForCombinedImageHandler::handle(Op opcode, const uint
uint32_t ptr = args[2];
compiler.set<SPIRExpression>(id, "", result_type, true);
compiler.register_read(id, ptr, true);
// Other backends might use SPIRAccessChain for this later.
compiler.ids[id].set_allow_type_rewrite();
break;
}
@ -3696,6 +3705,10 @@ void Compiler::analyze_variable_scope(SPIRFunction &entry)
notify_variable_access(args[i], current_block->self);
// The result of an access chain is a fixed expression and is not really considered a temporary.
auto &e = compiler.set<SPIRExpression>(args[1], "", args[0], true);
auto *backing_variable = compiler.maybe_get_backing_variable(ptr);
e.loaded_from = backing_variable ? backing_variable->self : 0;
compiler.ids[args[1]].set_allow_type_rewrite();
break;
}

4
spirv_glsl.cpp
View File

@ -6073,7 +6073,9 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
bool need_transpose, result_is_packed;
auto e = access_chain(ops[2], &ops[3], length - 3, get<SPIRType>(ops[0]), &need_transpose, &result_is_packed);
auto &expr = set<SPIRExpression>(ops[1], move(e), ops[0], should_forward(ops[2]));
expr.loaded_from = ops[2];
auto *backing_variable = maybe_get_backing_variable(ops[2]);
expr.loaded_from = backing_variable ? backing_variable->self : ops[2];
expr.need_transpose = need_transpose;
// Mark the result as being packed. Some platforms handled packed vectors differently than non-packed.