[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 <dneto@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Auto-Submit: dan sinclair <dsinclair@chromium.org>

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
 )");
 }

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
 )");
 }

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<ast::BlockStatement>());
 
   ASSERT_TRUE(b.GenerateFunction(&func));
-  EXPECT_EQ(DumpInstructions(b.debug()), R"(OpName %3 "tint_615f66756e63"
+  EXPECT_EQ(DumpBuilder(b), R"(OpName %3 "tint_615f66756e63"
-)");
+%2 = OpTypeVoid
-  EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeVoid
 %1 = OpTypeFunction %2
+%3 = OpFunction %2 None %1
+%4 = OpLabel
+OpReturn
+OpFunctionEnd
 )");
+}
 
-  ASSERT_GE(b.functions().size(), 1u);
+TEST_F(BuilderTest, Function_Terminator_Return) {
-  const auto& ret = b.functions()[0];
+  ast::type::VoidType void_type;
-  EXPECT_EQ(DumpInstruction(ret.declaration()), R"(%3 = OpFunction %2 None %1
+
+  auto* body = create<ast::BlockStatement>();
+  body->append(create<ast::ReturnStatement>());
+
+  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<ast::Variable>("a", ast::StorageClass::kPrivate, &f32);
+  td.RegisterVariableForTesting(var_a);
+
+  auto* body = create<ast::BlockStatement>();
+  body->append(
+      create<ast::ReturnStatement>(create<ast::IdentifierExpression>("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<ast::BlockStatement>();
+  body->append(create<ast::DiscardStatement>());
+
+  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
 )");
 }
 

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
 )");
 }

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
 )");
 }

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<void(const Instruction&)> 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, {}});
 }
 

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<f32>(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<f32>(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;
 }

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

test/function.wgsl

@@ -19,5 +19,4 @@ fn main() -> f32 {
 [[stage(compute)]]
 [[workgroup_size(2)]]
 fn ep() -> void {
-  return;
 }

test/simple.wgsl

@@ -15,7 +15,6 @@
 [[location(0)]] var<out> gl_FragColor : vec4<f32>;
 
 fn bar() -> void {
-  return;
 }
 
 [[stage(fragment)]]
@@ -23,5 +22,4 @@ fn main() -> void {
     var a : vec2<f32> = vec2<f32>();
     gl_FragColor = vec4<f32>(0.4, 0.4, 0.8, 1.0);
     bar();
-    return;
 }

test/triangle.wgsl

@@ -24,7 +24,6 @@ const pos : array<vec2<f32>, 3> = array<vec2<f32>, 3>(
 [[stage(vertex)]]
 fn vtx_main() -> void {
   Position = vec4<f32>(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<f32>(1.0, 0.0, 0.0, 1.0);
-  return;
 }