resolver: Validate calls to void callables

Calls to functions and intrinsics that do not return a value must only be used by a call statement.

Fixed: chromium:1241460
Change-Id: I0f940c942b55a5212367dbf9e261083beb4560ec
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/62440
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>

src/ast/intrinsic_texture_helper_test.cc

@@ -2237,6 +2237,18 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
   };
 }
 
+bool ReturnsVoid(ValidTextureOverload texture_overload) {
+  switch (texture_overload) {
+    case ValidTextureOverload::kStoreWO1dRgba32float:
+    case ValidTextureOverload::kStoreWO2dRgba32float:
+    case ValidTextureOverload::kStoreWO2dArrayRgba32float:
+    case ValidTextureOverload::kStoreWO3dRgba32float:
+      return true;
+    default:
+      return false;
+  }
+}
+
 }  // namespace test
 }  // namespace intrinsic
 }  // namespace ast

src/ast/intrinsic_texture_helper_test.h

@@ -185,6 +185,10 @@ enum class ValidTextureOverload {
   kStoreWO3dRgba32float,             // Not permutated for all texel formats
 };
 
+/// @param texture_overload the ValidTextureOverload
+/// @returns true if the ValidTextureOverload intrinsic returns no value.
+bool ReturnsVoid(ValidTextureOverload texture_overload);
+
 /// Describes a texture intrinsic overload
 struct TextureOverloadCase {
   /// Constructor for textureSample...() functions

src/resolver/call_test.cc

@@ -90,7 +90,7 @@ TEST_F(ResolverCallTest, Valid) {
     args.push_back(p.create_value(*this, 0));
   }
 
-  Func("foo", std::move(params), ty.void_(), {Return()});
+  Func("foo", std::move(params), ty.f32(), {Return(1.23f)});
   auto* call = Call("foo", std::move(args));
   WrapInFunction(call);
 

src/resolver/function_validation_test.cc

@@ -216,6 +216,20 @@ TEST_F(ResolverFunctionValidationTest,
             "type, returned 'i32', expected 'void'");
 }
 
+TEST_F(ResolverFunctionValidationTest,
+       FunctionTypeMustMatchReturnStatementType_void_fail) {
+  // fn v { return; }
+  // fn func { return v(); }
+  Func("v", {}, ty.void_(), {Return()});
+  Func("func", {}, ty.void_(),
+       {
+           Return(Call(Source{Source::Location{12, 34}}, "v")),
+       });
+
+  EXPECT_FALSE(r()->Resolve());
+  EXPECT_EQ(r()->error(), "12:34 error: function 'v' does not return a value");
+}
+
 TEST_F(ResolverFunctionValidationTest,
        FunctionTypeMustMatchReturnStatementTypeMissing_fail) {
   // fn func -> f32 { return; }

src/resolver/intrinsic_test.cc

@@ -557,7 +557,7 @@ TEST_P(ResolverIntrinsicTest_Barrier, InferType) {
   auto param = GetParam();
 
   auto* call = Call(param.name);
-  WrapInFunction(call);
+  WrapInFunction(create<ast::CallStatement>(call));
 
   EXPECT_TRUE(r()->Resolve()) << r()->error();
   ASSERT_NE(TypeOf(call), nullptr);
@@ -568,7 +568,7 @@ TEST_P(ResolverIntrinsicTest_Barrier, Error_TooManyParams) {
   auto param = GetParam();
 
   auto* call = Call(param.name, vec4<f32>(1.f, 2.f, 3.f, 4.f), 1.0f);
-  WrapInFunction(call);
+  WrapInFunction(create<ast::CallStatement>(call));
 
   EXPECT_FALSE(r()->Resolve());
 
@@ -2082,8 +2082,10 @@ TEST_P(ResolverIntrinsicTest_Texture, Call) {
   param.buildSamplerVariable(this);
 
   auto* call = Call(param.function, param.args(this));
-  Func("func", {}, ty.void_(), {Ignore(call)},
+  auto* stmt = ast::intrinsic::test::ReturnsVoid(param.overload)
-       {create<ast::StageDecoration>(ast::PipelineStage::kFragment)});
+                   ? create<ast::CallStatement>(call)
+                   : Ignore(call);
+  Func("func", {}, ty.void_(), {stmt}, {Stage(ast::PipelineStage::kFragment)});
 
   ASSERT_TRUE(r()->Resolve()) << r()->error();
 

src/resolver/intrinsic_validation_test.cc

@@ -21,6 +21,20 @@ namespace {
 
 using ResolverIntrinsicValidationTest = ResolverTest;
 
+TEST_F(ResolverIntrinsicValidationTest,
+       FunctionTypeMustMatchReturnStatementType_void_fail) {
+  // fn func { return workgroupBarrier(); }
+  Func("func", {}, ty.void_(),
+       {
+           Return(Call(Source{Source::Location{12, 34}}, "workgroupBarrier")),
+       });
+
+  EXPECT_FALSE(r()->Resolve());
+  EXPECT_EQ(
+      r()->error(),
+      "12:34 error: intrinsic 'workgroupBarrier' does not return a value");
+}
+
 TEST_F(ResolverIntrinsicValidationTest, InvalidPipelineStageDirect) {
   // [[stage(compute), workgroup_size(1)]] fn func { return dpdx(1.0); }
 
@@ -42,16 +56,16 @@ TEST_F(ResolverIntrinsicValidationTest, InvalidPipelineStageIndirect) {
 
   auto* dpdx = create<ast::CallExpression>(Source{{3, 4}}, Expr("dpdx"),
                                            ast::ExpressionList{Expr(1.0f)});
-  Func(Source{{1, 2}}, "f0", ast::VariableList{}, ty.void_(), {Ignore(dpdx)});
+  Func(Source{{1, 2}}, "f0", {}, ty.void_(), {Ignore(dpdx)});
 
-  Func(Source{{3, 4}}, "f1", ast::VariableList{}, ty.void_(),
+  Func(Source{{3, 4}}, "f1", {}, ty.void_(),
-       {Ignore(Call("f0"))});
+       {create<ast::CallStatement>(Call("f0"))});
 
-  Func(Source{{5, 6}}, "f2", ast::VariableList{}, ty.void_(),
+  Func(Source{{5, 6}}, "f2", {}, ty.void_(),
-       {Ignore(Call("f1"))});
+       {create<ast::CallStatement>(Call("f1"))});
 
-  Func(Source{{7, 8}}, "main", ast::VariableList{}, ty.void_(),
+  Func(Source{{7, 8}}, "main", {}, ty.void_(),
-       {Ignore(Call("f2"))},
+       {create<ast::CallStatement>(Call("f2"))},
        {Stage(ast::PipelineStage::kCompute), WorkgroupSize(1)});
 
   EXPECT_FALSE(r()->Resolve());

src/resolver/resolver.cc

@@ -1463,12 +1463,9 @@ bool Resolver::ValidateEntryPoint(const ast::Function* func,
   };
 
   // Inner lambda that is applied to a type and all of its members.
-  auto validate_entry_point_decorations_inner = [&](const ast::DecorationList&
+  auto validate_entry_point_decorations_inner =
-                                                        decos,
+      [&](const ast::DecorationList& decos, sem::Type* ty, Source source,
-                                                    sem::Type* ty,
+          ParamOrRetType param_or_ret, bool is_struct_member) {
-                                                    Source source,
-                                                    ParamOrRetType param_or_ret,
-                                                    bool is_struct_member) {
         // Scan decorations for pipeline IO attributes.
         // Check for overlap with attributes that have been seen previously.
         ast::Decoration* pipeline_io_attribute = nullptr;
@@ -1478,14 +1475,16 @@ bool Resolver::ValidateEntryPoint(const ast::Function* func,
           if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
             if (pipeline_io_attribute) {
               AddError("multiple entry point IO attributes", deco->source());
-          AddNote("previously consumed " + deco_to_str(pipeline_io_attribute),
+              AddNote(
+                  "previously consumed " + deco_to_str(pipeline_io_attribute),
                   pipeline_io_attribute->source());
               return false;
             }
             pipeline_io_attribute = deco;
 
             if (builtins.count(builtin->value())) {
-          AddError(deco_to_str(builtin) +
+              AddError(
+                  deco_to_str(builtin) +
                       " attribute appears multiple times as pipeline " +
                       (param_or_ret == ParamOrRetType::kParameter ? "input"
                                                                   : "output"),
@@ -1503,7 +1502,8 @@ bool Resolver::ValidateEntryPoint(const ast::Function* func,
           } else if (auto* location = deco->As<ast::LocationDecoration>()) {
             if (pipeline_io_attribute) {
               AddError("multiple entry point IO attributes", deco->source());
-          AddNote("previously consumed " + deco_to_str(pipeline_io_attribute),
+              AddNote(
+                  "previously consumed " + deco_to_str(pipeline_io_attribute),
                   pipeline_io_attribute->source());
               return false;
             }
@@ -1514,7 +1514,8 @@ bool Resolver::ValidateEntryPoint(const ast::Function* func,
                                             is_input)) {
               return false;
             }
-      } else if (auto* interpolate = deco->As<ast::InterpolateDecoration>()) {
+          } else if (auto* interpolate =
+                         deco->As<ast::InterpolateDecoration>()) {
             if (func->pipeline_stage() == ast::PipelineStage::kCompute) {
               is_invalid_compute_shader_decoration = true;
             } else if (!ValidateInterpolateDecoration(interpolate, ty)) {
@@ -1528,7 +1529,8 @@ bool Resolver::ValidateEntryPoint(const ast::Function* func,
           }
           if (is_invalid_compute_shader_decoration) {
             std::string input_or_output =
-            param_or_ret == ParamOrRetType::kParameter ? "inputs" : "output";
+                param_or_ret == ParamOrRetType::kParameter ? "inputs"
+                                                           : "output";
             AddError(
                 "decoration is not valid for compute shader " + input_or_output,
                 deco->source());
@@ -1536,13 +1538,13 @@ bool Resolver::ValidateEntryPoint(const ast::Function* func,
           }
         }
 
-    if (IsValidationEnabled(decos,
+        if (IsValidationEnabled(
-                            ast::DisabledValidation::kEntryPointParameter)) {
+                decos, ast::DisabledValidation::kEntryPointParameter)) {
           if (!ty->Is<sem::Struct>() && !pipeline_io_attribute) {
             std::string err = "missing entry point IO attribute";
             if (!is_struct_member) {
-          err +=
+              err += (param_or_ret == ParamOrRetType::kParameter
-              (param_or_ret == ParamOrRetType::kParameter ? " on parameter"
+                          ? " on parameter"
                           : " on return type");
             }
             AddError(err, source);
@@ -2338,28 +2340,42 @@ bool Resolver::Call(ast::CallExpression* call) {
     }
   }
 
+  return ValidateCall(call);
+}
+
+bool Resolver::ValidateCall(ast::CallExpression* call) {
+  if (TypeOf(call)->Is<sem::Void>()) {
+    bool is_call_statement = false;
+    if (current_statement_) {
+      if (auto* call_stmt =
+              As<ast::CallStatement>(current_statement_->Declaration())) {
+        if (call_stmt->expr() == call) {
+          is_call_statement = true;
+        }
+      }
+    }
+    if (!is_call_statement) {
+      // https://gpuweb.github.io/gpuweb/wgsl/#function-call-expr
+      // If the called function does not return a value, a function call
+      // statement should be used instead.
+      auto* ident = call->func();
+      auto name = builder_->Symbols().NameFor(ident->symbol());
+      // A function call is made to either a user declared function or an
+      // intrinsic. function_calls_ only maps CallExpression to user declared
+      // functions
+      bool is_function = function_calls_.count(call) != 0;
+      AddError((is_function ? "function" : "intrinsic") + std::string(" '") +
+                   name + "' does not return a value",
+               call->source());
+      return false;
+    }
+  }
+
   return true;
 }
 
 bool Resolver::ValidateCallStatement(ast::CallStatement* stmt) {
-  const sem::Type* return_type = nullptr;
+  const sem::Type* return_type = TypeOf(stmt->expr());
-  // A function call is made to either a user declared function or an intrinsic.
-  // function_calls_ only maps CallExpression to user declared functions
-  auto it = function_calls_.find(stmt->expr());
-  if (it != function_calls_.end()) {
-    return_type = it->second.function->return_type;
-  } else {
-    // Must be an intrinsic call
-    auto* target = builder_->Sem().Get(stmt->expr())->Target();
-    if (auto* intrinsic = target->As<sem::Intrinsic>()) {
-      return_type = intrinsic->ReturnType();
-    } else {
-      TINT_ICE(Resolver, diagnostics_)
-          << "call target was not an intrinsic, but a "
-          << intrinsic->TypeInfo().name;
-    }
-  }
-
   if (!return_type->Is<sem::Void>()) {
     // https://gpuweb.github.io/gpuweb/wgsl/#function-call-statement
     // A function call statement executes a function call where the called

src/resolver/resolver.h

@@ -276,6 +276,7 @@ class Resolver {
                                  const sem::Type* storage_type,
                                  const bool is_input,
                                  const bool is_struct_member);
+  bool ValidateCall(ast::CallExpression* call);
   bool ValidateCallStatement(ast::CallStatement* stmt);
   bool ValidateEntryPoint(const ast::Function* func, const FunctionInfo* info);
   bool ValidateFunction(const ast::Function* func, const FunctionInfo* info);

src/resolver/validation_test.cc

@@ -94,11 +94,11 @@ TEST_F(ResolverValidationTest, WorkgroupMemoryUsedInFragmentStage) {
   Global("dst", ty.vec4<f32>(), ast::StorageClass::kPrivate);
   auto* stmt = Assign(Expr("dst"), Expr(Source{{3, 4}}, "wg"));
 
-  Func(Source{{5, 6}}, "f2", ast::VariableList{}, ty.void_(), {stmt});
+  Func(Source{{5, 6}}, "f2", {}, ty.void_(), {stmt});
-  Func(Source{{7, 8}}, "f1", ast::VariableList{}, ty.void_(),
+  Func(Source{{7, 8}}, "f1", {}, ty.void_(),
-       {Ignore(Call("f2"))});
+       {create<ast::CallStatement>(Call("f2"))});
-  Func(Source{{9, 10}}, "f0", ast::VariableList{}, ty.void_(),
+  Func(Source{{9, 10}}, "f0", {}, ty.void_(),
-       {Ignore(Call("f1"))},
+       {create<ast::CallStatement>(Call("f1"))},
        ast::DecorationList{Stage(ast::PipelineStage::kFragment)});
 
   EXPECT_FALSE(r()->Resolve());

src/writer/hlsl/generator_impl_call_test.cc

@@ -23,8 +23,7 @@ namespace {
 using HlslGeneratorImplTest_Call = TestHelper;
 
 TEST_F(HlslGeneratorImplTest_Call, EmitExpression_Call_WithoutParams) {
-  Func("my_func", ast::VariableList{}, ty.void_(), ast::StatementList{},
+  Func("my_func", {}, ty.f32(), {Return(1.23f)});
-       ast::DecorationList{});
 
   auto* call = Call("my_func");
   WrapInFunction(call);
@@ -42,7 +41,7 @@ TEST_F(HlslGeneratorImplTest_Call, EmitExpression_Call_WithParams) {
            Param(Sym(), ty.f32()),
            Param(Sym(), ty.f32()),
        },
-       ty.void_(), ast::StatementList{}, ast::DecorationList{});
+       ty.f32(), {Return(1.23f)});
   Global("param1", ty.f32(), ast::StorageClass::kPrivate);
   Global("param2", ty.f32(), ast::StorageClass::kPrivate);
 

src/writer/hlsl/generator_impl_intrinsic_texture_test.cc

@@ -361,14 +361,11 @@ TEST_P(HlslGeneratorIntrinsicTextureTest, Call) {
   param.buildSamplerVariable(this);
 
   auto* call = Call(param.function, param.args(this));
+  auto* stmt = ast::intrinsic::test::ReturnsVoid(param.overload)
+                   ? create<ast::CallStatement>(call)
+                   : Ignore(call);
 
-  Func("main", ast::VariableList{}, ty.void_(),
+  Func("main", {}, ty.void_(), {stmt}, {Stage(ast::PipelineStage::kFragment)});
-       ast::StatementList{
-           Ignore(call),
-       },
-       ast::DecorationList{
-           Stage(ast::PipelineStage::kFragment),
-       });
 
   GeneratorImpl& gen = SanitizeAndBuild();
 

src/writer/msl/generator_impl_call_test.cc

@@ -23,8 +23,7 @@ namespace {
 using MslGeneratorImplTest = TestHelper;
 
 TEST_F(MslGeneratorImplTest, EmitExpression_Call_WithoutParams) {
-  Func("my_func", ast::VariableList{}, ty.void_(), ast::StatementList{},
+  Func("my_func", {}, ty.f32(), {Return(1.23f)});
-       ast::DecorationList{});
 
   auto* call = Call("my_func");
   WrapInFunction(call);
@@ -42,7 +41,7 @@ TEST_F(MslGeneratorImplTest, EmitExpression_Call_WithParams) {
            Param(Sym(), ty.f32()),
            Param(Sym(), ty.f32()),
        },
-       ty.void_(), ast::StatementList{}, ast::DecorationList{});
+       ty.f32(), {Return(1.23f)});
   Global("param1", ty.f32(), ast::StorageClass::kPrivate);
   Global("param2", ty.f32(), ast::StorageClass::kPrivate);
 

src/writer/msl/generator_impl_intrinsic_test.cc

@@ -299,7 +299,7 @@ TEST_F(MslGeneratorImplTest, Intrinsic_Call) {
 
 TEST_F(MslGeneratorImplTest, StorageBarrier) {
   auto* call = Call("storageBarrier");
-  WrapInFunction(call);
+  WrapInFunction(create<ast::CallStatement>(call));
 
   GeneratorImpl& gen = Build();
 
@@ -310,7 +310,7 @@ TEST_F(MslGeneratorImplTest, StorageBarrier) {
 
 TEST_F(MslGeneratorImplTest, WorkgroupBarrier) {
   auto* call = Call("workgroupBarrier");
-  WrapInFunction(call);
+  WrapInFunction(create<ast::CallStatement>(call));
 
   GeneratorImpl& gen = Build();
 

src/writer/msl/generator_impl_intrinsic_texture_test.cc

@@ -273,14 +273,11 @@ TEST_P(MslGeneratorIntrinsicTextureTest, Call) {
 
   auto* call =
       create<ast::CallExpression>(Expr(param.function), param.args(this));
+  auto* stmt = ast::intrinsic::test::ReturnsVoid(param.overload)
+                   ? create<ast::CallStatement>(call)
+                   : Ignore(call);
 
-  Func("main", ast::VariableList{}, ty.void_(),
+  Func("main", {}, ty.void_(), {stmt}, {Stage(ast::PipelineStage::kFragment)});
-       ast::StatementList{
-           Ignore(call),
-       },
-       ast::DecorationList{
-           Stage(ast::PipelineStage::kFragment),
-       });
 
   GeneratorImpl& gen = Build();
 

src/writer/spirv/builder_intrinsic_texture_test.cc

@@ -3691,9 +3691,10 @@ TEST_P(IntrinsicTextureTest, Call) {
   auto* sampler = param.buildSamplerVariable(this);
 
   auto* call = Call(param.function, param.args(this));
-
+  auto* stmt = ast::intrinsic::test::ReturnsVoid(param.overload)
-  Func("func", {}, ty.void_(), {Ignore(call)},
+                   ? create<ast::CallStatement>(call)
-       {create<ast::StageDecoration>(ast::PipelineStage::kFragment)});
+                   : Ignore(call);
+  Func("func", {}, ty.void_(), {stmt}, {Stage(ast::PipelineStage::kFragment)});
 
   spirv::Builder& b = Build();
 
@@ -3719,13 +3720,10 @@ TEST_P(IntrinsicTextureTest, ValidateSPIRV) {
 
   auto* call = Call(param.function, param.args(this));
 
-  Func("main", ast::VariableList{}, ty.void_(),
+  auto* stmt = ast::intrinsic::test::ReturnsVoid(param.overload)
-       ast::StatementList{
+                   ? create<ast::CallStatement>(call)
-           Ignore(call),
+                   : Ignore(call);
-       },
+  Func("main", {}, ty.void_(), {stmt}, {Stage(ast::PipelineStage::kFragment)});
-       ast::DecorationList{
-           Stage(ast::PipelineStage::kFragment),
-       });
 
   spirv::Builder& b = Build();
 
@@ -3744,7 +3742,10 @@ TEST_P(IntrinsicTextureTest, OutsideFunction_IsError) {
   auto* sampler = param.buildSamplerVariable(this);
 
   auto* call = Call(param.function, param.args(this));
-  Func("func", {}, ty.void_(), {Ignore(call)},
+  auto* stmt = ast::intrinsic::test::ReturnsVoid(param.overload)
+                   ? create<ast::CallStatement>(call)
+                   : Ignore(call);
+  Func("func", {}, ty.void_(), {stmt},
        {create<ast::StageDecoration>(ast::PipelineStage::kFragment)});
 
   spirv::Builder& b = Build();

src/writer/wgsl/generator_impl_call_test.cc

@@ -23,8 +23,7 @@ namespace {
 using WgslGeneratorImplTest = TestHelper;
 
 TEST_F(WgslGeneratorImplTest, EmitExpression_Call_WithoutParams) {
-  Func("my_func", ast::VariableList{}, ty.void_(), ast::StatementList{},
+  Func("my_func", {}, ty.f32(), {Return(1.23f)});
-       ast::DecorationList{});
 
   auto* call = Call("my_func");
   WrapInFunction(call);
@@ -42,7 +41,7 @@ TEST_F(WgslGeneratorImplTest, EmitExpression_Call_WithParams) {
            Param(Sym(), ty.f32()),
            Param(Sym(), ty.f32()),
        },
-       ty.void_(), ast::StatementList{}, ast::DecorationList{});
+       ty.f32(), {Return(1.23f)});
   Global("param1", ty.f32(), ast::StorageClass::kPrivate);
   Global("param2", ty.f32(), ast::StorageClass::kPrivate);