diff --git a/src/BUILD.gn b/src/BUILD.gn
index 1af8434ab6..4fa57467a3 100644
--- a/src/BUILD.gn
+++ b/src/BUILD.gn
@@ -281,6 +281,8 @@ source_set("libtint_core_src") {
     "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 b3ae095921..53810cb7d7 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -96,6 +96,8 @@ set(TINT_LIB_SRCS
   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 6e438d814a..db0cffd76a 100644
--- a/src/ast/function.cc
+++ b/src/ast/function.cc
@@ -40,7 +40,6 @@ Function::Function(const Source& source,
   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 eb856c351b..1352a71ec2 100644
--- a/src/ast/function.h
+++ b/src/ast/function.h
@@ -63,6 +63,18 @@ class Function : public Castable<Function, Node> {
   /// @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 0000000000..d00d4b0180
--- /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 0000000000..ce242c64ee
--- /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 7011e468bf..46bc613df8 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 @@ bool Resolver::ValidateFunction(const ast::Function* func) {
   }
 
   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 @@ bool Resolver::Function(ast::Function* func) {
     }
   }
 
-  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 40916db361..32c7e81775 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 @@ bool GeneratorImpl::EmitFunction(std::ostream& out, ast::Function* func) {
 
   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 f894c4e63b..7187eb821d 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 @@ bool GeneratorImpl::EmitFunction(ast::Function* func) {
     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 @@ bool GeneratorImpl::EmitFunction(ast::Function* func) {
     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) {