Add ast::InternalDecoration
An tint-internal decoration used to add metadata between a sanitizer transform and a backend.
Will be used for declaring backend-specific intrinsic calls.
Change-Id: Ia05ba7dada0148de2d490605ba4d15c593075356
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/46868
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/BUILD.gn b/src/BUILD.gn
index 1af8434..4fa5746 100644
--- a/src/BUILD.gn
+++ b/src/BUILD.gn
@@ -281,6 +281,8 @@
"ast/if_statement.h",
"ast/int_literal.cc",
"ast/int_literal.h",
+ "ast/internal_decoration.cc",
+ "ast/internal_decoration.h",
"ast/literal.cc",
"ast/literal.h",
"ast/location_decoration.cc",
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index b3ae095..53810cb 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -96,6 +96,8 @@
ast/if_statement.h
ast/int_literal.cc
ast/int_literal.h
+ ast/internal_decoration.cc
+ ast/internal_decoration.h
ast/literal.cc
ast/literal.h
ast/location_decoration.cc
diff --git a/src/ast/function.cc b/src/ast/function.cc
index 6e438d8..db0cffd 100644
--- a/src/ast/function.cc
+++ b/src/ast/function.cc
@@ -40,7 +40,6 @@
for (auto* param : params_) {
TINT_ASSERT(param);
}
- TINT_ASSERT(body_);
TINT_ASSERT(symbol_.IsValid());
TINT_ASSERT(return_type_);
}
diff --git a/src/ast/function.h b/src/ast/function.h
index eb856c3..1352a71 100644
--- a/src/ast/function.h
+++ b/src/ast/function.h
@@ -63,6 +63,18 @@
/// @returns the decorations attached to this function
const DecorationList& decorations() const { return decorations_; }
+ /// @returns the decoration with the type `T` or nullptr if this function does
+ /// not contain a decoration with the given type
+ template <typename T>
+ const T* find_decoration() const {
+ for (auto* deco : decorations()) {
+ if (auto* d = deco->As<T>()) {
+ return d;
+ }
+ }
+ return nullptr;
+ }
+
/// @returns the workgroup size {x, y, z} for the function. {1, 1, 1} will be
/// return if no workgroup size was set.
std::tuple<uint32_t, uint32_t, uint32_t> workgroup_size() const;
diff --git a/src/ast/internal_decoration.cc b/src/ast/internal_decoration.cc
new file mode 100644
index 0000000..d00d4b0
--- /dev/null
+++ b/src/ast/internal_decoration.cc
@@ -0,0 +1,34 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "src/ast/internal_decoration.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ast::InternalDecoration);
+
+namespace tint {
+namespace ast {
+
+InternalDecoration::InternalDecoration() : Base(Source{}) {}
+
+InternalDecoration::~InternalDecoration() = default;
+
+void InternalDecoration::to_str(const semantic::Info&,
+ std::ostream& out,
+ size_t indent) const {
+ make_indent(out, indent);
+ out << "tint_internal(" << Name() << ")" << std::endl;
+}
+
+} // namespace ast
+} // namespace tint
diff --git a/src/ast/internal_decoration.h b/src/ast/internal_decoration.h
new file mode 100644
index 0000000..ce242c6
--- /dev/null
+++ b/src/ast/internal_decoration.h
@@ -0,0 +1,53 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SRC_AST_INTERNAL_DECORATION_H_
+#define SRC_AST_INTERNAL_DECORATION_H_
+
+#include <string>
+
+#include "src/ast/decoration.h"
+
+namespace tint {
+namespace ast {
+
+/// A decoration used to indicate that a function is tint-internal.
+/// These decorations are not produced by generators, but instead are usually
+/// created by transforms for consumption by a particular backend.
+/// Functions annotated with this decoration will have relaxed validation.
+class InternalDecoration : public Castable<InternalDecoration, Decoration> {
+ public:
+ /// Constructor
+ InternalDecoration();
+
+ /// Destructor
+ ~InternalDecoration() override;
+
+ /// @return a short description of the internal decoration which will be
+ /// displayed in WGSL as `[[internal(<name>)]]` (but is not parsable).
+ virtual std::string Name() const = 0;
+
+ /// Writes a representation of the node to the output stream
+ /// @param sem the semantic info for the program
+ /// @param out the stream to write to
+ /// @param indent number of spaces to indent the node when writing
+ void to_str(const semantic::Info& sem,
+ std::ostream& out,
+ size_t indent) const override;
+};
+
+} // namespace ast
+} // namespace tint
+
+#endif // SRC_AST_INTERNAL_DECORATION_H_
diff --git a/src/resolver/resolver.cc b/src/resolver/resolver.cc
index 7011e46..46bc613 100644
--- a/src/resolver/resolver.cc
+++ b/src/resolver/resolver.cc
@@ -26,6 +26,7 @@
#include "src/ast/discard_statement.h"
#include "src/ast/fallthrough_statement.h"
#include "src/ast/if_statement.h"
+#include "src/ast/internal_decoration.h"
#include "src/ast/loop_statement.h"
#include "src/ast/return_statement.h"
#include "src/ast/struct_block_decoration.h"
@@ -301,12 +302,18 @@
}
if (!func->return_type()->Is<type::Void>()) {
- if (!func->get_last_statement() ||
- !func->get_last_statement()->Is<ast::ReturnStatement>()) {
- diagnostics_.add_error(
- "v-0002", "non-void function must end with a return statement",
- func->source());
- return false;
+ if (func->body()) {
+ if (!func->get_last_statement() ||
+ !func->get_last_statement()->Is<ast::ReturnStatement>()) {
+ diagnostics_.add_error(
+ "v-0002", "non-void function must end with a return statement",
+ func->source());
+ return false;
+ }
+ } else if (!func->find_decoration<ast::InternalDecoration>()) {
+ TINT_ICE(diagnostics_)
+ << "Function " << builder_->Symbols().NameFor(func->symbol())
+ << " has no body and does not have the [[internal]] decoration";
}
for (auto* deco : func->return_type_decorations()) {
@@ -594,8 +601,10 @@
}
}
- if (!BlockStatement(func->body())) {
- return false;
+ if (func->body()) {
+ if (!BlockStatement(func->body())) {
+ return false;
+ }
}
variable_stack_.pop_scope();
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 40916db..32c7e81 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -20,6 +20,7 @@
#include "src/ast/call_statement.h"
#include "src/ast/constant_id_decoration.h"
#include "src/ast/fallthrough_statement.h"
+#include "src/ast/internal_decoration.h"
#include "src/ast/variable_decl_statement.h"
#include "src/semantic/array.h"
#include "src/semantic/call.h"
@@ -1464,6 +1465,11 @@
auto* func_sem = builder_.Sem().Get(func);
+ if (func->find_decoration<ast::InternalDecoration>()) {
+ // An internal function. Do not emit.
+ return true;
+ }
+
// TODO(dsinclair): This could be smarter. If the input/outputs for multiple
// entry points are the same we could generate a single struct and then have
// this determine it's the same struct and just emit once.
diff --git a/src/writer/wgsl/generator_impl.cc b/src/writer/wgsl/generator_impl.cc
index f894c4e..7187eb8 100644
--- a/src/writer/wgsl/generator_impl.cc
+++ b/src/writer/wgsl/generator_impl.cc
@@ -21,6 +21,7 @@
#include "src/ast/call_statement.h"
#include "src/ast/constant_id_decoration.h"
#include "src/ast/float_literal.h"
+#include "src/ast/internal_decoration.h"
#include "src/ast/module.h"
#include "src/ast/sint_literal.h"
#include "src/ast/stage_decoration.h"
@@ -299,6 +300,9 @@
if (auto* stage = deco->As<ast::StageDecoration>()) {
out_ << "stage(" << stage->value() << ")";
}
+ if (auto* internal = deco->As<ast::InternalDecoration>()) {
+ out_ << "internal(" << internal->Name() << ")";
+ }
out_ << "]]" << std::endl;
}
@@ -339,8 +343,14 @@
return false;
}
- out_ << " ";
- return EmitBlockAndNewline(func->body());
+ if (func->body()) {
+ out_ << " ";
+ return EmitBlockAndNewline(func->body());
+ } else {
+ out_ << std::endl;
+ }
+
+ return true;
}
bool GeneratorImpl::EmitImageFormat(const type::ImageFormat fmt) {
