diff --git a/reference/opt/shaders-msl/comp/struct-packing.comp b/reference/opt/shaders-msl/comp/struct-packing.comp index fb72a9bf..a042f7aa 100644 --- a/reference/opt/shaders-msl/comp/struct-packing.comp +++ b/reference/opt/shaders-msl/comp/struct-packing.comp @@ -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]; } diff --git a/reference/opt/shaders/asm/geom/inout-split-access-chain-handle.asm.geom b/reference/opt/shaders/asm/geom/inout-split-access-chain-handle.asm.geom new file mode 100644 index 00000000..ca1381cf --- /dev/null +++ b/reference/opt/shaders/asm/geom/inout-split-access-chain-handle.asm.geom @@ -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; +} + diff --git a/reference/opt/shaders/comp/struct-packing.comp b/reference/opt/shaders/comp/struct-packing.comp index ecd7932b..8a2ac762 100644 --- a/reference/opt/shaders/comp/struct-packing.comp +++ b/reference/opt/shaders/comp/struct-packing.comp @@ -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]; diff --git a/reference/shaders-msl/comp/struct-packing.comp b/reference/shaders-msl/comp/struct-packing.comp index 73ea6947..2b37844f 100644 --- a/reference/shaders-msl/comp/struct-packing.comp +++ b/reference/shaders-msl/comp/struct-packing.comp @@ -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]; } diff --git a/reference/shaders/asm/geom/inout-split-access-chain-handle.asm.geom b/reference/shaders/asm/geom/inout-split-access-chain-handle.asm.geom new file mode 100644 index 00000000..71082099 --- /dev/null +++ b/reference/shaders/asm/geom/inout-split-access-chain-handle.asm.geom @@ -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; +} + diff --git a/reference/shaders/comp/struct-packing.comp b/reference/shaders/comp/struct-packing.comp index d8ece843..cd1eda1b 100644 --- a/reference/shaders/comp/struct-packing.comp +++ b/reference/shaders/comp/struct-packing.comp @@ -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]; diff --git a/shaders/asm/geom/inout-split-access-chain-handle.asm.geom b/shaders/asm/geom/inout-split-access-chain-handle.asm.geom new file mode 100644 index 00000000..d011cc69 --- /dev/null +++ b/shaders/asm/geom/inout-split-access-chain-handle.asm.geom @@ -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 diff --git a/shaders/comp/struct-packing.comp b/shaders/comp/struct-packing.comp index 53a54e49..7a1be047 100644 --- a/shaders/comp/struct-packing.comp +++ b/shaders/comp/struct-packing.comp @@ -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]; diff --git a/spirv_common.hpp b/spirv_common.hpp index 8d8ba12a..5190b4e7 100644 --- a/spirv_common.hpp +++ b/spirv_common.hpp @@ -1111,9 +1111,10 @@ public: void set(std::unique_ptr 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 @@ -1154,9 +1155,15 @@ public: type = TypeNone; } + void set_allow_type_rewrite() + { + allow_type_rewrite = true; + } + private: std::unique_ptr holder; uint32_t type = TypeNone; + bool allow_type_rewrite = false; }; template diff --git a/spirv_cross.cpp b/spirv_cross.cpp index 23cd9940..1a667745 100644 --- a/spirv_cross.cpp +++ b/spirv_cross.cpp @@ -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(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(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; } diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index ac133f03..e6e44264 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -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(ops[0]), &need_transpose, &result_is_packed); auto &expr = set(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.