From a7c9391dc62549e78ce14941b09c76120f339a84 Mon Sep 17 00:00:00 2001 From: dan sinclair Date: Wed, 18 Nov 2020 18:25:30 +0000 Subject: [PATCH] [spirv-writer] Support optional trailing return. This CL updates the SPIRV-Writer to inject an OpReturn as the trailing statement in a function if the function does not end with a `discard` or a `return` statement. R=bclayton@google.com, dneto@google.com Fixes: tint:302 Change-Id: I2e7c7beff15ad30c779c591bb75cf97fc0960bf7 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/33160 Reviewed-by: David Neto Commit-Queue: dan sinclair Auto-Submit: dan sinclair --- src/writer/spirv/builder_call_test.cc | 2 + .../spirv/builder_function_decoration_test.cc | 2 + src/writer/spirv/builder_function_test.cc | 81 +++++++++++++++++-- src/writer/spirv/builder_intrinsic_test.cc | 26 ++++++ src/writer/spirv/builder_switch_test.cc | 5 ++ src/writer/spirv/function.cc | 20 +++++ test/compute_boids.wgsl | 4 - test/cube.wgsl | 2 - test/function.wgsl | 1 - test/simple.wgsl | 2 - test/triangle.wgsl | 2 - 11 files changed, 130 insertions(+), 17 deletions(-) diff --git a/src/writer/spirv/builder_call_test.cc b/src/writer/spirv/builder_call_test.cc index 1ca5c0e88f..1fc8aa455d 100644 --- a/src/writer/spirv/builder_call_test.cc +++ b/src/writer/spirv/builder_call_test.cc @@ -95,6 +95,7 @@ OpFunctionEnd %12 = OpFunction %11 None %10 %13 = OpLabel %14 = OpFunctionCall %2 %3 %15 %15 +OpReturn OpFunctionEnd )"); } @@ -156,6 +157,7 @@ OpFunctionEnd %12 = OpFunction %2 None %11 %13 = OpLabel %14 = OpFunctionCall %2 %4 %15 %15 +OpReturn OpFunctionEnd )"); } diff --git a/src/writer/spirv/builder_function_decoration_test.cc b/src/writer/spirv/builder_function_decoration_test.cc index a68f7804ca..8b59b2ed43 100644 --- a/src/writer/spirv/builder_function_decoration_test.cc +++ b/src/writer/spirv/builder_function_decoration_test.cc @@ -258,9 +258,11 @@ OpName %5 "tint_6d61696e32" %1 = OpTypeFunction %2 %3 = OpFunction %2 None %1 %4 = OpLabel +OpReturn OpFunctionEnd %5 = OpFunction %2 None %1 %6 = OpLabel +OpReturn OpFunctionEnd )"); } diff --git a/src/writer/spirv/builder_function_test.cc b/src/writer/spirv/builder_function_test.cc index 528ee18d88..1a6de922ca 100644 --- a/src/writer/spirv/builder_function_test.cc +++ b/src/writer/spirv/builder_function_test.cc @@ -18,6 +18,7 @@ #include "spirv/unified1/spirv.h" #include "spirv/unified1/spirv.hpp11" #include "src/ast/decorated_variable.h" +#include "src/ast/discard_statement.h" #include "src/ast/function.h" #include "src/ast/identifier_expression.h" #include "src/ast/member_accessor_expression.h" @@ -51,15 +52,83 @@ TEST_F(BuilderTest, Function_Empty) { ast::Function func("a_func", {}, &void_type, create()); ASSERT_TRUE(b.GenerateFunction(&func)); - EXPECT_EQ(DumpInstructions(b.debug()), R"(OpName %3 "tint_615f66756e63" -)"); - EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeVoid + EXPECT_EQ(DumpBuilder(b), R"(OpName %3 "tint_615f66756e63" +%2 = OpTypeVoid %1 = OpTypeFunction %2 +%3 = OpFunction %2 None %1 +%4 = OpLabel +OpReturn +OpFunctionEnd )"); +} - ASSERT_GE(b.functions().size(), 1u); - const auto& ret = b.functions()[0]; - EXPECT_EQ(DumpInstruction(ret.declaration()), R"(%3 = OpFunction %2 None %1 +TEST_F(BuilderTest, Function_Terminator_Return) { + ast::type::VoidType void_type; + + auto* body = create(); + body->append(create()); + + ast::Function func("a_func", {}, &void_type, body); + + ASSERT_TRUE(b.GenerateFunction(&func)); + EXPECT_EQ(DumpBuilder(b), R"(OpName %3 "tint_615f66756e63" +%2 = OpTypeVoid +%1 = OpTypeFunction %2 +%3 = OpFunction %2 None %1 +%4 = OpLabel +OpReturn +OpFunctionEnd +)"); +} + +TEST_F(BuilderTest, Function_Terminator_ReturnValue) { + ast::type::VoidType void_type; + ast::type::F32Type f32; + + auto* var_a = create("a", ast::StorageClass::kPrivate, &f32); + td.RegisterVariableForTesting(var_a); + + auto* body = create(); + body->append( + create(create("a"))); + ASSERT_TRUE(td.DetermineResultType(body)) << td.error(); + + ast::Function func("a_func", {}, &void_type, body); + + ASSERT_TRUE(b.GenerateGlobalVariable(var_a)) << b.error(); + ASSERT_TRUE(b.GenerateFunction(&func)) << b.error(); + EXPECT_EQ(DumpBuilder(b), R"(OpName %1 "tint_61" +OpName %7 "tint_615f66756e63" +%3 = OpTypeFloat 32 +%2 = OpTypePointer Private %3 +%4 = OpConstantNull %3 +%1 = OpVariable %2 Private %4 +%6 = OpTypeVoid +%5 = OpTypeFunction %6 +%7 = OpFunction %6 None %5 +%8 = OpLabel +%9 = OpLoad %3 %1 +OpReturnValue %9 +OpFunctionEnd +)"); +} + +TEST_F(BuilderTest, Function_Terminator_Discard) { + ast::type::VoidType void_type; + + auto* body = create(); + body->append(create()); + + ast::Function func("a_func", {}, &void_type, body); + + ASSERT_TRUE(b.GenerateFunction(&func)); + EXPECT_EQ(DumpBuilder(b), R"(OpName %3 "tint_615f66756e63" +%2 = OpTypeVoid +%1 = OpTypeFunction %2 +%3 = OpFunction %2 None %1 +%4 = OpLabel +OpKill +OpFunctionEnd )"); } diff --git a/src/writer/spirv/builder_intrinsic_test.cc b/src/writer/spirv/builder_intrinsic_test.cc index de8cecfa56..c877885dc8 100644 --- a/src/writer/spirv/builder_intrinsic_test.cc +++ b/src/writer/spirv/builder_intrinsic_test.cc @@ -940,6 +940,7 @@ OpName %7 "tint_615f66756e63" %8 = OpLabel %11 = OpLoad %3 %1 %9 = OpExtInst %3 %10 Round %11 +OpReturn OpFunctionEnd )"); } @@ -967,6 +968,7 @@ OpName %3 "tint_615f66756e63" %4 = OpLabel %5 = OpExtInst %6 %7 )" + param.op + R"( %8 +OpReturn OpFunctionEnd )"); } @@ -994,6 +996,7 @@ OpName %3 "tint_615f66756e63" %4 = OpLabel %5 = OpExtInst %6 %8 )" + param.op + R"( %10 +OpReturn OpFunctionEnd )"); } @@ -1042,6 +1045,7 @@ OpName %3 "tint_615f66756e63" %3 = OpFunction %2 None %1 %4 = OpLabel %5 = OpExtInst %6 %7 Length %8 +OpReturn OpFunctionEnd )"); } @@ -1066,6 +1070,7 @@ OpName %3 "tint_615f66756e63" %3 = OpFunction %2 None %1 %4 = OpLabel %5 = OpExtInst %6 %7 Length %10 +OpReturn OpFunctionEnd )"); } @@ -1090,6 +1095,7 @@ OpName %3 "tint_615f66756e63" %3 = OpFunction %2 None %1 %4 = OpLabel %5 = OpExtInst %6 %8 Normalize %10 +OpReturn OpFunctionEnd )"); } @@ -1118,6 +1124,7 @@ OpName %3 "tint_615f66756e63" %4 = OpLabel %5 = OpExtInst %6 %7 )" + param.op + R"( %8 %8 +OpReturn OpFunctionEnd )"); } @@ -1146,6 +1153,7 @@ OpName %3 "tint_615f66756e63" %4 = OpLabel %5 = OpExtInst %6 %8 )" + param.op + R"( %10 %10 +OpReturn OpFunctionEnd )"); } @@ -1177,6 +1185,7 @@ OpName %3 "tint_615f66756e63" %3 = OpFunction %2 None %1 %4 = OpLabel %5 = OpExtInst %6 %7 Distance %8 %8 +OpReturn OpFunctionEnd )"); } @@ -1202,6 +1211,7 @@ OpName %3 "tint_615f66756e63" %3 = OpFunction %2 None %1 %4 = OpLabel %5 = OpExtInst %6 %7 Distance %10 %10 +OpReturn OpFunctionEnd )"); } @@ -1228,6 +1238,7 @@ OpName %3 "tint_615f66756e63" %3 = OpFunction %2 None %1 %4 = OpLabel %5 = OpExtInst %6 %8 Cross %10 %10 +OpReturn OpFunctionEnd )"); } @@ -1255,6 +1266,7 @@ OpName %3 "tint_615f66756e63" %4 = OpLabel %5 = OpExtInst %6 %7 )" + param.op + R"( %8 %8 %8 +OpReturn OpFunctionEnd )"); } @@ -1284,6 +1296,7 @@ OpName %3 "tint_615f66756e63" %4 = OpLabel %5 = OpExtInst %6 %8 )" + param.op + R"( %10 %10 %10 +OpReturn OpFunctionEnd )"); } @@ -1320,6 +1333,7 @@ OpName %3 "tint_615f66756e63" %4 = OpLabel %5 = OpExtInst %6 %7 )" + param.op + R"( %8 +OpReturn OpFunctionEnd )"); } @@ -1347,6 +1361,7 @@ OpName %3 "tint_615f66756e63" %4 = OpLabel %5 = OpExtInst %6 %8 )" + param.op + R"( %10 +OpReturn OpFunctionEnd )"); } @@ -1377,6 +1392,7 @@ OpName %3 "tint_615f66756e63" %4 = OpLabel %5 = OpExtInst %6 %7 )" + param.op + R"( %8 +OpReturn OpFunctionEnd )"); } @@ -1404,6 +1420,7 @@ OpName %3 "tint_615f66756e63" %4 = OpLabel %5 = OpExtInst %6 %8 )" + param.op + R"( %10 +OpReturn OpFunctionEnd )"); } @@ -1434,6 +1451,7 @@ OpName %3 "tint_615f66756e63" %4 = OpLabel %5 = OpExtInst %6 %7 )" + param.op + R"( %8 %8 +OpReturn OpFunctionEnd )"); } @@ -1461,6 +1479,7 @@ OpName %3 "tint_615f66756e63" %4 = OpLabel %5 = OpExtInst %6 %8 )" + param.op + R"( %10 %10 +OpReturn OpFunctionEnd )"); } @@ -1492,6 +1511,7 @@ OpName %3 "tint_615f66756e63" %4 = OpLabel %5 = OpExtInst %6 %7 )" + param.op + R"( %8 %8 +OpReturn OpFunctionEnd )"); } @@ -1519,6 +1539,7 @@ OpName %3 "tint_615f66756e63" %4 = OpLabel %5 = OpExtInst %6 %8 )" + param.op + R"( %10 %10 +OpReturn OpFunctionEnd )"); } @@ -1550,6 +1571,7 @@ OpName %3 "tint_615f66756e63" %4 = OpLabel %5 = OpExtInst %6 %7 )" + param.op + R"( %8 %8 %8 +OpReturn OpFunctionEnd )"); } @@ -1579,6 +1601,7 @@ OpName %3 "tint_615f66756e63" %4 = OpLabel %5 = OpExtInst %6 %8 )" + param.op + R"( %10 %10 %10 +OpReturn OpFunctionEnd )"); } @@ -1609,6 +1632,7 @@ OpName %3 "tint_615f66756e63" %4 = OpLabel %5 = OpExtInst %6 %7 )" + param.op + R"( %8 %8 %8 +OpReturn OpFunctionEnd )"); } @@ -1638,6 +1662,7 @@ OpName %3 "tint_615f66756e63" %4 = OpLabel %5 = OpExtInst %6 %8 )" + param.op + R"( %10 %10 %10 +OpReturn OpFunctionEnd )"); } @@ -1673,6 +1698,7 @@ OpName %5 "tint_766172" %4 = OpLabel %13 = OpLoad %7 %5 %11 = OpExtInst %9 %12 Determinant %13 +OpReturn OpFunctionEnd )"); } diff --git a/src/writer/spirv/builder_switch_test.cc b/src/writer/spirv/builder_switch_test.cc index 2738b28189..b4b110b292 100644 --- a/src/writer/spirv/builder_switch_test.cc +++ b/src/writer/spirv/builder_switch_test.cc @@ -141,6 +141,7 @@ OpBranch %9 %11 = OpLabel OpBranch %9 %9 = OpLabel +OpReturn OpFunctionEnd )"); } @@ -198,6 +199,7 @@ OpSwitch %10 %11 OpStore %1 %12 OpBranch %9 %9 = OpLabel +OpReturn OpFunctionEnd )"); } @@ -288,6 +290,7 @@ OpBranch %9 OpStore %1 %16 OpBranch %9 %9 = OpLabel +OpReturn OpFunctionEnd )"); } @@ -379,6 +382,7 @@ OpBranch %9 OpStore %1 %16 OpBranch %9 %9 = OpLabel +OpReturn OpFunctionEnd )"); } @@ -501,6 +505,7 @@ OpBranch %9 %11 = OpLabel OpBranch %9 %9 = OpLabel +OpReturn OpFunctionEnd )"); } diff --git a/src/writer/spirv/function.cc b/src/writer/spirv/function.cc index 01da96ce9d..5c5c726afa 100644 --- a/src/writer/spirv/function.cc +++ b/src/writer/spirv/function.cc @@ -17,6 +17,15 @@ namespace tint { namespace writer { namespace spirv { +namespace { + +// Returns true if the given Op is a function terminator +bool OpIsFunctionTerminator(spv::Op op) { + return op == spv::Op::OpReturn || op == spv::Op::OpReturnValue || + op == spv::Op::OpKill; +} + +} // namespace Function::Function() : declaration_(Instruction{spv::Op::OpNop, {}}), @@ -47,6 +56,17 @@ void Function::iterate(std::function cb) const { cb(inst); } + bool needs_terminator = false; + if (instructions_.empty()) { + needs_terminator = true; + } else { + const auto& last = instructions_.back(); + needs_terminator = !OpIsFunctionTerminator(last.opcode()); + } + if (needs_terminator) { + cb(Instruction{spv::Op::OpReturn, {}}); + } + cb(Instruction{spv::Op::OpFunctionEnd, {}}); } diff --git a/test/compute_boids.wgsl b/test/compute_boids.wgsl index 1d32673598..6bee6d3a9a 100644 --- a/test/compute_boids.wgsl +++ b/test/compute_boids.wgsl @@ -26,7 +26,6 @@ fn vert_main() -> void { (a_pos.x * cos(angle)) - (a_pos.y * sin(angle)), (a_pos.x * sin(angle)) + (a_pos.y * cos(angle))); gl_Position = vec4(pos + a_particlePos, 0.0, 1.0); - return; } # fragment shader @@ -35,7 +34,6 @@ fn vert_main() -> void { [[stage(fragment)]] fn frag_main() -> void { fragColor = vec4(1.0, 1.0, 1.0, 1.0); - return; } # compute shader @@ -137,6 +135,4 @@ fn comp_main() -> void { # Write back particlesB.particles[index].pos = vPos; particlesB.particles[index].vel = vVel; - - return; } diff --git a/test/cube.wgsl b/test/cube.wgsl index eb67254dae..d56f32f1d1 100644 --- a/test/cube.wgsl +++ b/test/cube.wgsl @@ -28,7 +28,6 @@ fn vtx_main() -> void { Position = uniforms.modelViewProjectionMatrix * cur_position; vtxFragColor = color; - return; } # Fragment shader @@ -38,5 +37,4 @@ fn vtx_main() -> void { [[stage(fragment)]] fn frag_main() -> void { outColor = fragColor; - return; } diff --git a/test/function.wgsl b/test/function.wgsl index f2b2e2281d..604b8ffef3 100644 --- a/test/function.wgsl +++ b/test/function.wgsl @@ -19,5 +19,4 @@ fn main() -> f32 { [[stage(compute)]] [[workgroup_size(2)]] fn ep() -> void { - return; } diff --git a/test/simple.wgsl b/test/simple.wgsl index ff5f38e6a2..844bfcbbf4 100644 --- a/test/simple.wgsl +++ b/test/simple.wgsl @@ -15,7 +15,6 @@ [[location(0)]] var gl_FragColor : vec4; fn bar() -> void { - return; } [[stage(fragment)]] @@ -23,5 +22,4 @@ fn main() -> void { var a : vec2 = vec2(); gl_FragColor = vec4(0.4, 0.4, 0.8, 1.0); bar(); - return; } diff --git a/test/triangle.wgsl b/test/triangle.wgsl index 0159eb6d5a..d612c94ee2 100644 --- a/test/triangle.wgsl +++ b/test/triangle.wgsl @@ -24,7 +24,6 @@ const pos : array, 3> = array, 3>( [[stage(vertex)]] fn vtx_main() -> void { Position = vec4(pos[VertexIndex], 0.0, 1.0); - return; } # Fragment shader @@ -33,5 +32,4 @@ fn vtx_main() -> void { [[stage(fragment)]] fn frag_main() -> void { outColor = vec4(1.0, 0.0, 0.0, 1.0); - return; }