[tint][ir][val] Check vertex output builtin rules
- Moves ReferenceModuleVars to core/ir/ and modifies it to support
const and non-const operations
- Some rewrites of existing tests whose shaders were not following
the vertex rules.
- Disables some tests where it is known getting them working will
require additional work.
Fixes: 367764845
Change-Id: I6decc0253b349de0e544b428c4180a5a1863540d
WIP: Initial constification
Change-Id: I69b33b28f6b9da3b0d281f78f9a3066bd3a54110
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/207374
Commit-Queue: Ryan Harrison <rharrison@chromium.org>
Auto-Submit: Ryan Harrison <rharrison@chromium.org>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/cmd/test/BUILD.bazel b/src/tint/cmd/test/BUILD.bazel
index 3d8f21d..7720586 100644
--- a/src/tint/cmd/test/BUILD.bazel
+++ b/src/tint/cmd/test/BUILD.bazel
@@ -46,7 +46,6 @@
"//src/tint/api/common:test",
"//src/tint/lang/core/constant:test",
"//src/tint/lang/core/intrinsic:test",
- "//src/tint/lang/core/ir/transform/common:test",
"//src/tint/lang/core/ir/transform:test",
"//src/tint/lang/core/ir:test",
"//src/tint/lang/core/type:test",
diff --git a/src/tint/cmd/test/BUILD.cmake b/src/tint/cmd/test/BUILD.cmake
index 36633be..a039469 100644
--- a/src/tint/cmd/test/BUILD.cmake
+++ b/src/tint/cmd/test/BUILD.cmake
@@ -47,7 +47,6 @@
tint_api_common_test
tint_lang_core_constant_test
tint_lang_core_intrinsic_test
- tint_lang_core_ir_transform_common_test
tint_lang_core_ir_transform_test
tint_lang_core_ir_test
tint_lang_core_type_test
diff --git a/src/tint/cmd/test/BUILD.gn b/src/tint/cmd/test/BUILD.gn
index c346355..6152cfc 100644
--- a/src/tint/cmd/test/BUILD.gn
+++ b/src/tint/cmd/test/BUILD.gn
@@ -56,7 +56,6 @@
"${tint_src_dir}/lang/core/intrinsic:unittests",
"${tint_src_dir}/lang/core/ir:unittests",
"${tint_src_dir}/lang/core/ir/transform:unittests",
- "${tint_src_dir}/lang/core/ir/transform/common:unittests",
"${tint_src_dir}/lang/core/type:unittests",
"${tint_src_dir}/lang/glsl/ir:unittests",
"${tint_src_dir}/lang/hlsl/ir:unittests",
diff --git a/src/tint/lang/core/ir/BUILD.bazel b/src/tint/lang/core/ir/BUILD.bazel
index b525147..e924736 100644
--- a/src/tint/lang/core/ir/BUILD.bazel
+++ b/src/tint/lang/core/ir/BUILD.bazel
@@ -132,6 +132,7 @@
"multi_in_block.h",
"next_iteration.h",
"operand_instruction.h",
+ "referenced_module_vars.h",
"return.h",
"store.h",
"store_vector_element.h",
@@ -207,6 +208,7 @@
"multi_in_block_test.cc",
"next_iteration_test.cc",
"operand_instruction_test.cc",
+ "referenced_module_vars_test.cc",
"return_test.cc",
"store_test.cc",
"store_vector_element_test.cc",
diff --git a/src/tint/lang/core/ir/BUILD.cmake b/src/tint/lang/core/ir/BUILD.cmake
index b4d6e1d..7eb4265 100644
--- a/src/tint/lang/core/ir/BUILD.cmake
+++ b/src/tint/lang/core/ir/BUILD.cmake
@@ -119,6 +119,7 @@
lang/core/ir/next_iteration.h
lang/core/ir/operand_instruction.cc
lang/core/ir/operand_instruction.h
+ lang/core/ir/referenced_module_vars.h
lang/core/ir/return.cc
lang/core/ir/return.h
lang/core/ir/store.cc
@@ -211,6 +212,7 @@
lang/core/ir/multi_in_block_test.cc
lang/core/ir/next_iteration_test.cc
lang/core/ir/operand_instruction_test.cc
+ lang/core/ir/referenced_module_vars_test.cc
lang/core/ir/return_test.cc
lang/core/ir/store_test.cc
lang/core/ir/store_vector_element_test.cc
diff --git a/src/tint/lang/core/ir/BUILD.gn b/src/tint/lang/core/ir/BUILD.gn
index 86ec6bd..26f700b 100644
--- a/src/tint/lang/core/ir/BUILD.gn
+++ b/src/tint/lang/core/ir/BUILD.gn
@@ -122,6 +122,7 @@
"next_iteration.h",
"operand_instruction.cc",
"operand_instruction.h",
+ "referenced_module_vars.h",
"return.cc",
"return.h",
"store.cc",
@@ -208,6 +209,7 @@
"multi_in_block_test.cc",
"next_iteration_test.cc",
"operand_instruction_test.cc",
+ "referenced_module_vars_test.cc",
"return_test.cc",
"store_test.cc",
"store_vector_element_test.cc",
diff --git a/src/tint/lang/core/ir/referenced_module_vars.h b/src/tint/lang/core/ir/referenced_module_vars.h
new file mode 100644
index 0000000..34401c1
--- /dev/null
+++ b/src/tint/lang/core/ir/referenced_module_vars.h
@@ -0,0 +1,159 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#ifndef SRC_TINT_LANG_CORE_IR_REFERENCED_MODULE_VARS_H_
+#define SRC_TINT_LANG_CORE_IR_REFERENCED_MODULE_VARS_H_
+
+#include <functional>
+
+#include "src/tint/lang/core/ir/module.h"
+#include "src/tint/lang/core/ir/user_call.h"
+#include "src/tint/lang/core/ir/var.h"
+#include "src/tint/utils/containers/hashmap.h"
+#include "src/tint/utils/containers/unique_vector.h"
+#include "src/tint/utils/rtti/switch.h"
+
+// Forward declarations.
+namespace tint::core::ir {
+class Block;
+class Function;
+} // namespace tint::core::ir
+
+/// Utility that helps guarantee makes sure the same const-ness is applied to both type
+template <class Src, class Dst>
+using TranscribeConst = std::conditional_t<std::is_const<Src>{}, std::add_const_t<Dst>, Dst>;
+
+namespace tint::core::ir {
+
+/// ReferencedModuleVars is a helper to determine the set of module-scope variables that are
+/// transitively referenced by functions in a module.
+/// References are determined lazily and cached for future requests.
+///
+/// Note:
+/// The template param M is used to ensure that inputs and outputs of this class have the same
+/// const-ness. If 'Module' is supplied then the internal operations and output will not be
+/// const, which is needed for transforms. Whereas if the param is 'const Module' the internals
+/// and outputs will be const, which is needed for the IR validator.
+/// Note:
+/// Changes to the module can invalidate the cached data. This is intended to be created by
+/// operations that need this information, and discarded when they complete. Tracking this
+/// information inside the IR module would add overhead any time an instruction is added or
+/// removed from the module. Since only a few operations need this information, it is expected
+/// to be more efficient to generate it on demand.
+template <typename M>
+class ReferencedModuleVars {
+ // Replace this with concepts when C++20 is available
+ static_assert(std::is_same<std::remove_cv_t<M>, Module>());
+
+ public:
+ /// Short form aliases for types that have the same constant-ness as M.
+ /// (The single use types are not aliased)
+ using BlockT = TranscribeConst<M, Block>;
+ using VarT = TranscribeConst<M, Var>;
+ using FunctionT = TranscribeConst<M, Function>;
+
+ /// A set of a variables referenced by a function (in declaration order).
+ using VarSet = UniqueVector<VarT*, 16>;
+
+ /// Constructor.
+ /// @param ir the module
+ /// @param pred an predicate function for filtering variables
+ /// Note: @p pred is not stored by the class, so can be a lambda that captures by reference.
+ template <typename Predicate>
+ ReferencedModuleVars(M& ir, Predicate&& pred) {
+ // Loop over module-scope variables, recording the blocks that they are referenced from.
+ BlockT* root_block = ir.root_block;
+ for (auto* inst : *root_block) {
+ if (auto* var = inst->template As<VarT>()) {
+ if (pred(var)) {
+ if (!var->Result(0)) {
+ continue;
+ }
+ var->Result(0)->ForEachUseUnsorted([&](const Usage& use) {
+ block_to_direct_vars_.GetOrAddZero(use.instruction->Block()).Add(var);
+ });
+ }
+ }
+ }
+ }
+
+ /// Constructor.
+ /// Provided default predicate that accepts all variables.
+ explicit ReferencedModuleVars(M& ir) : ReferencedModuleVars(ir, [](VarT*) { return true; }) {}
+
+ /// Get the set of transitively referenced module-scope variables for a function, filtered by
+ /// the predicate function if provided.
+ /// @param func the function
+ /// @returns the set of (possibly filtered) transitively reference module-scope variables
+ VarSet& TransitiveReferences(FunctionT* func) {
+ return transitive_references_.GetOrAdd(func, [&] {
+ VarSet vars;
+ GetTransitiveReferences(func->Block(), vars);
+ return vars;
+ });
+ }
+
+ private:
+ /// A map from blocks to their directly referenced variables.
+ Hashmap<BlockT*, VarSet, 64> block_to_direct_vars_{};
+
+ /// A map from functions to their transitively referenced variables.
+ Hashmap<FunctionT*, VarSet, 8> transitive_references_;
+
+ /// Get the set of transitively referenced module-scope variables for a block.
+ /// @param block the block
+ /// @param vars the set of transitively reference module-scope variables to populate
+ void GetTransitiveReferences(BlockT* block, VarSet& vars) {
+ // Add directly referenced vars.
+ if (auto itr = block_to_direct_vars_.Get(block)) {
+ for (auto& var : *itr) {
+ vars.Add(var);
+ }
+ }
+
+ // Loop over instructions in the block to find indirectly referenced vars.
+ for (auto* inst : *block) {
+ tint::Switch(
+ inst,
+ [&](TranscribeConst<M, UserCall>* call) {
+ // Get variables referenced by a function called from this block.
+ const auto& callee_vars = TransitiveReferences(call->Target());
+ for (auto* var : callee_vars) {
+ vars.Add(var);
+ }
+ },
+ [&](TranscribeConst<M, ControlInstruction>* ctrl) {
+ // Recurse into control instructions and gather their referenced vars.
+ ctrl->ForeachBlock([&](BlockT* blk) { GetTransitiveReferences(blk, vars); });
+ });
+ }
+ }
+};
+
+} // namespace tint::core::ir
+
+#endif // SRC_TINT_LANG_CORE_IR_REFERENCED_MODULE_VARS_H_
diff --git a/src/tint/lang/core/ir/transform/common/referenced_module_vars_test.cc b/src/tint/lang/core/ir/referenced_module_vars_test.cc
similarity index 96%
rename from src/tint/lang/core/ir/transform/common/referenced_module_vars_test.cc
rename to src/tint/lang/core/ir/referenced_module_vars_test.cc
index a8ee411..dce9c18 100644
--- a/src/tint/lang/core/ir/transform/common/referenced_module_vars_test.cc
+++ b/src/tint/lang/core/ir/referenced_module_vars_test.cc
@@ -25,7 +25,7 @@
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-#include "src/tint/lang/core/ir/transform/common/referenced_module_vars.h"
+#include "src/tint/lang/core/ir/referenced_module_vars.h"
#include <string>
@@ -62,7 +62,7 @@
)";
EXPECT_EQ(src, Disassemble());
- ReferencedModuleVars vars(mod);
+ ReferencedModuleVars<Module> vars(mod);
auto& foo_vars = vars.TransitiveReferences(foo);
EXPECT_TRUE(foo_vars.IsEmpty());
}
@@ -98,7 +98,7 @@
)";
EXPECT_EQ(src, Disassemble());
- ReferencedModuleVars vars(mod);
+ ReferencedModuleVars<Module> vars(mod);
EXPECT_THAT(vars.TransitiveReferences(foo), ElementsAre(var_a, var_b));
}
@@ -141,7 +141,7 @@
)";
EXPECT_EQ(src, Disassemble());
- ReferencedModuleVars vars(mod);
+ ReferencedModuleVars<Module> vars(mod);
EXPECT_THAT(vars.TransitiveReferences(foo), ElementsAre(var_a, var_b, var_c, var_d, var_e));
}
@@ -198,7 +198,7 @@
)";
EXPECT_EQ(src, Disassemble());
- ReferencedModuleVars vars(mod);
+ ReferencedModuleVars<Module> vars(mod);
EXPECT_THAT(vars.TransitiveReferences(foo), ElementsAre(var_a, var_b));
EXPECT_THAT(vars.TransitiveReferences(bar), ElementsAre(var_a, var_c));
EXPECT_TRUE(vars.TransitiveReferences(zoo).IsEmpty());
@@ -275,7 +275,7 @@
)";
EXPECT_EQ(src, Disassemble());
- ReferencedModuleVars vars(mod);
+ ReferencedModuleVars<Module> vars(mod);
EXPECT_THAT(vars.TransitiveReferences(foo), ElementsAre(var_a, var_b, var_c, var_d));
}
@@ -335,7 +335,7 @@
)";
EXPECT_EQ(src, Disassemble());
- ReferencedModuleVars vars(mod);
+ ReferencedModuleVars<Module> vars(mod);
EXPECT_THAT(vars.TransitiveReferences(bar), ElementsAre(var_b));
EXPECT_THAT(vars.TransitiveReferences(zoo), ElementsAre(var_b));
EXPECT_THAT(vars.TransitiveReferences(foo), ElementsAre(var_a, var_b));
@@ -368,7 +368,7 @@
)";
EXPECT_EQ(src, Disassemble());
- ReferencedModuleVars vars(mod);
+ ReferencedModuleVars<Module> vars(mod);
EXPECT_THAT(vars.TransitiveReferences(foo), ElementsAre(var_a));
}
@@ -411,7 +411,7 @@
)";
EXPECT_EQ(src, Disassemble());
- ReferencedModuleVars vars(mod, [](const Var* var) {
+ ReferencedModuleVars<Module> vars(mod, [](const Var* var) {
auto* view = var->Result(0)->Type()->As<type::MemoryView>();
return view->AddressSpace() == AddressSpace::kPrivate;
});
diff --git a/src/tint/lang/core/ir/transform/BUILD.bazel b/src/tint/lang/core/ir/transform/BUILD.bazel
index b90c54f..7dcf606 100644
--- a/src/tint/lang/core/ir/transform/BUILD.bazel
+++ b/src/tint/lang/core/ir/transform/BUILD.bazel
@@ -93,7 +93,6 @@
"//src/tint/lang/core/constant",
"//src/tint/lang/core/intrinsic",
"//src/tint/lang/core/ir",
- "//src/tint/lang/core/ir/transform/common",
"//src/tint/lang/core/type",
"//src/tint/utils/containers",
"//src/tint/utils/diagnostic",
diff --git a/src/tint/lang/core/ir/transform/BUILD.cmake b/src/tint/lang/core/ir/transform/BUILD.cmake
index 1687918..fe36cb8 100644
--- a/src/tint/lang/core/ir/transform/BUILD.cmake
+++ b/src/tint/lang/core/ir/transform/BUILD.cmake
@@ -34,8 +34,6 @@
# Do not modify this file directly
################################################################################
-include(lang/core/ir/transform/common/BUILD.cmake)
-
################################################################################
# Target: tint_lang_core_ir_transform
# Kind: lib
@@ -94,7 +92,6 @@
tint_lang_core_constant
tint_lang_core_intrinsic
tint_lang_core_ir
- tint_lang_core_ir_transform_common
tint_lang_core_type
tint_utils_containers
tint_utils_diagnostic
diff --git a/src/tint/lang/core/ir/transform/BUILD.gn b/src/tint/lang/core/ir/transform/BUILD.gn
index 9387529..d4e8158 100644
--- a/src/tint/lang/core/ir/transform/BUILD.gn
+++ b/src/tint/lang/core/ir/transform/BUILD.gn
@@ -98,7 +98,6 @@
"${tint_src_dir}/lang/core/constant",
"${tint_src_dir}/lang/core/intrinsic",
"${tint_src_dir}/lang/core/ir",
- "${tint_src_dir}/lang/core/ir/transform/common",
"${tint_src_dir}/lang/core/type",
"${tint_src_dir}/utils/containers",
"${tint_src_dir}/utils/diagnostic",
diff --git a/src/tint/lang/core/ir/transform/common/BUILD.bazel b/src/tint/lang/core/ir/transform/common/BUILD.bazel
deleted file mode 100644
index c02990e..0000000
--- a/src/tint/lang/core/ir/transform/common/BUILD.bazel
+++ /dev/null
@@ -1,105 +0,0 @@
-# Copyright 2024 The Dawn & Tint Authors
-#
-# Redistribution and use in source and binary forms, with or without
-# modification, are permitted provided that the following conditions are met:
-#
-# 1. Redistributions of source code must retain the above copyright notice, this
-# list of conditions and the following disclaimer.
-#
-# 2. Redistributions in binary form must reproduce the above copyright notice,
-# this list of conditions and the following disclaimer in the documentation
-# and/or other materials provided with the distribution.
-#
-# 3. Neither the name of the copyright holder nor the names of its
-# contributors may be used to endorse or promote products derived from
-# this software without specific prior written permission.
-#
-# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
-# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
-# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-# DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
-# FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
-# DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
-# SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
-# CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
-# OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
-# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-################################################################################
-# File generated by 'tools/src/cmd/gen' using the template:
-# tools/src/cmd/gen/build/BUILD.bazel.tmpl
-#
-# To regenerate run: './tools/run gen'
-#
-# Do not modify this file directly
-################################################################################
-
-load("//src/tint:flags.bzl", "COPTS")
-load("@bazel_skylib//lib:selects.bzl", "selects")
-cc_library(
- name = "common",
- srcs = [
- "referenced_module_vars.cc",
- ],
- hdrs = [
- "referenced_module_vars.h",
- ],
- deps = [
- "//src/tint/api/common",
- "//src/tint/lang/core",
- "//src/tint/lang/core/constant",
- "//src/tint/lang/core/ir",
- "//src/tint/lang/core/type",
- "//src/tint/utils/containers",
- "//src/tint/utils/diagnostic",
- "//src/tint/utils/ice",
- "//src/tint/utils/id",
- "//src/tint/utils/macros",
- "//src/tint/utils/math",
- "//src/tint/utils/memory",
- "//src/tint/utils/reflection",
- "//src/tint/utils/result",
- "//src/tint/utils/rtti",
- "//src/tint/utils/symbol",
- "//src/tint/utils/text",
- "//src/tint/utils/traits",
- "//src/utils",
- ],
- copts = COPTS,
- visibility = ["//visibility:public"],
-)
-cc_library(
- name = "test",
- alwayslink = True,
- srcs = [
- "referenced_module_vars_test.cc",
- ],
- deps = [
- "//src/tint/api/common",
- "//src/tint/lang/core",
- "//src/tint/lang/core/constant",
- "//src/tint/lang/core/intrinsic",
- "//src/tint/lang/core/ir",
- "//src/tint/lang/core/ir/transform/common",
- "//src/tint/lang/core/ir:test",
- "//src/tint/lang/core/type",
- "//src/tint/utils/containers",
- "//src/tint/utils/diagnostic",
- "//src/tint/utils/ice",
- "//src/tint/utils/id",
- "//src/tint/utils/macros",
- "//src/tint/utils/math",
- "//src/tint/utils/memory",
- "//src/tint/utils/reflection",
- "//src/tint/utils/result",
- "//src/tint/utils/rtti",
- "//src/tint/utils/symbol",
- "//src/tint/utils/text",
- "//src/tint/utils/traits",
- "@gtest",
- "//src/utils",
- ],
- copts = COPTS,
- visibility = ["//visibility:public"],
-)
-
diff --git a/src/tint/lang/core/ir/transform/common/BUILD.cmake b/src/tint/lang/core/ir/transform/common/BUILD.cmake
deleted file mode 100644
index 35a95e0..0000000
--- a/src/tint/lang/core/ir/transform/common/BUILD.cmake
+++ /dev/null
@@ -1,106 +0,0 @@
-# Copyright 2024 The Dawn & Tint Authors
-#
-# Redistribution and use in source and binary forms, with or without
-# modification, are permitted provided that the following conditions are met:
-#
-# 1. Redistributions of source code must retain the above copyright notice, this
-# list of conditions and the following disclaimer.
-#
-# 2. Redistributions in binary form must reproduce the above copyright notice,
-# this list of conditions and the following disclaimer in the documentation
-# and/or other materials provided with the distribution.
-#
-# 3. Neither the name of the copyright holder nor the names of its
-# contributors may be used to endorse or promote products derived from
-# this software without specific prior written permission.
-#
-# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
-# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
-# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-# DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
-# FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
-# DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
-# SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
-# CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
-# OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
-# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-################################################################################
-# File generated by 'tools/src/cmd/gen' using the template:
-# tools/src/cmd/gen/build/BUILD.cmake.tmpl
-#
-# To regenerate run: './tools/run gen'
-#
-# Do not modify this file directly
-################################################################################
-
-################################################################################
-# Target: tint_lang_core_ir_transform_common
-# Kind: lib
-################################################################################
-tint_add_target(tint_lang_core_ir_transform_common lib
- lang/core/ir/transform/common/referenced_module_vars.cc
- lang/core/ir/transform/common/referenced_module_vars.h
-)
-
-tint_target_add_dependencies(tint_lang_core_ir_transform_common lib
- tint_api_common
- tint_lang_core
- tint_lang_core_constant
- tint_lang_core_ir
- tint_lang_core_type
- tint_utils_containers
- tint_utils_diagnostic
- tint_utils_ice
- tint_utils_id
- tint_utils_macros
- tint_utils_math
- tint_utils_memory
- tint_utils_reflection
- tint_utils_result
- tint_utils_rtti
- tint_utils_symbol
- tint_utils_text
- tint_utils_traits
-)
-
-tint_target_add_external_dependencies(tint_lang_core_ir_transform_common lib
- "src_utils"
-)
-
-################################################################################
-# Target: tint_lang_core_ir_transform_common_test
-# Kind: test
-################################################################################
-tint_add_target(tint_lang_core_ir_transform_common_test test
- lang/core/ir/transform/common/referenced_module_vars_test.cc
-)
-
-tint_target_add_dependencies(tint_lang_core_ir_transform_common_test test
- tint_api_common
- tint_lang_core
- tint_lang_core_constant
- tint_lang_core_intrinsic
- tint_lang_core_ir
- tint_lang_core_ir_transform_common
- tint_lang_core_ir_test
- tint_lang_core_type
- tint_utils_containers
- tint_utils_diagnostic
- tint_utils_ice
- tint_utils_id
- tint_utils_macros
- tint_utils_math
- tint_utils_memory
- tint_utils_reflection
- tint_utils_result
- tint_utils_rtti
- tint_utils_symbol
- tint_utils_text
- tint_utils_traits
-)
-
-tint_target_add_external_dependencies(tint_lang_core_ir_transform_common_test test
- "gtest"
- "src_utils"
-)
diff --git a/src/tint/lang/core/ir/transform/common/BUILD.gn b/src/tint/lang/core/ir/transform/common/BUILD.gn
deleted file mode 100644
index 0dff5dae..0000000
--- a/src/tint/lang/core/ir/transform/common/BUILD.gn
+++ /dev/null
@@ -1,102 +0,0 @@
-# Copyright 2024 The Dawn & Tint Authors
-#
-# Redistribution and use in source and binary forms, with or without
-# modification, are permitted provided that the following conditions are met:
-#
-# 1. Redistributions of source code must retain the above copyright notice, this
-# list of conditions and the following disclaimer.
-#
-# 2. Redistributions in binary form must reproduce the above copyright notice,
-# this list of conditions and the following disclaimer in the documentation
-# and/or other materials provided with the distribution.
-#
-# 3. Neither the name of the copyright holder nor the names of its
-# contributors may be used to endorse or promote products derived from
-# this software without specific prior written permission.
-#
-# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
-# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
-# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-# DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
-# FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
-# DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
-# SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
-# CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
-# OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
-# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-################################################################################
-# File generated by 'tools/src/cmd/gen' using the template:
-# tools/src/cmd/gen/build/BUILD.gn.tmpl
-#
-# To regenerate run: './tools/run gen'
-#
-# Do not modify this file directly
-################################################################################
-
-import("../../../../../../../scripts/dawn_overrides_with_defaults.gni")
-import("../../../../../../../scripts/tint_overrides_with_defaults.gni")
-
-import("${tint_src_dir}/tint.gni")
-
-if (tint_build_unittests || tint_build_benchmarks) {
- import("//testing/test.gni")
-}
-
-libtint_source_set("common") {
- sources = [
- "referenced_module_vars.cc",
- "referenced_module_vars.h",
- ]
- deps = [
- "${dawn_root}/src/utils:utils",
- "${tint_src_dir}/api/common",
- "${tint_src_dir}/lang/core",
- "${tint_src_dir}/lang/core/constant",
- "${tint_src_dir}/lang/core/ir",
- "${tint_src_dir}/lang/core/type",
- "${tint_src_dir}/utils/containers",
- "${tint_src_dir}/utils/diagnostic",
- "${tint_src_dir}/utils/ice",
- "${tint_src_dir}/utils/id",
- "${tint_src_dir}/utils/macros",
- "${tint_src_dir}/utils/math",
- "${tint_src_dir}/utils/memory",
- "${tint_src_dir}/utils/reflection",
- "${tint_src_dir}/utils/result",
- "${tint_src_dir}/utils/rtti",
- "${tint_src_dir}/utils/symbol",
- "${tint_src_dir}/utils/text",
- "${tint_src_dir}/utils/traits",
- ]
-}
-if (tint_build_unittests) {
- tint_unittests_source_set("unittests") {
- sources = [ "referenced_module_vars_test.cc" ]
- deps = [
- "${dawn_root}/src/utils:utils",
- "${tint_src_dir}:gmock_and_gtest",
- "${tint_src_dir}/api/common",
- "${tint_src_dir}/lang/core",
- "${tint_src_dir}/lang/core/constant",
- "${tint_src_dir}/lang/core/intrinsic",
- "${tint_src_dir}/lang/core/ir",
- "${tint_src_dir}/lang/core/ir:unittests",
- "${tint_src_dir}/lang/core/ir/transform/common",
- "${tint_src_dir}/lang/core/type",
- "${tint_src_dir}/utils/containers",
- "${tint_src_dir}/utils/diagnostic",
- "${tint_src_dir}/utils/ice",
- "${tint_src_dir}/utils/id",
- "${tint_src_dir}/utils/macros",
- "${tint_src_dir}/utils/math",
- "${tint_src_dir}/utils/memory",
- "${tint_src_dir}/utils/reflection",
- "${tint_src_dir}/utils/result",
- "${tint_src_dir}/utils/rtti",
- "${tint_src_dir}/utils/symbol",
- "${tint_src_dir}/utils/text",
- "${tint_src_dir}/utils/traits",
- ]
- }
-}
diff --git a/src/tint/lang/core/ir/transform/common/referenced_module_vars.cc b/src/tint/lang/core/ir/transform/common/referenced_module_vars.cc
deleted file mode 100644
index b97b9dd..0000000
--- a/src/tint/lang/core/ir/transform/common/referenced_module_vars.cc
+++ /dev/null
@@ -1,75 +0,0 @@
-// Copyright 2024 The Dawn & Tint Authors
-//
-// Redistribution and use in source and binary forms, with or without
-// modification, are permitted provided that the following conditions are met:
-//
-// 1. Redistributions of source code must retain the above copyright notice, this
-// list of conditions and the following disclaimer.
-//
-// 2. Redistributions in binary form must reproduce the above copyright notice,
-// this list of conditions and the following disclaimer in the documentation
-// and/or other materials provided with the distribution.
-//
-// 3. Neither the name of the copyright holder nor the names of its
-// contributors may be used to endorse or promote products derived from
-// this software without specific prior written permission.
-//
-// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
-// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
-// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
-// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
-// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
-// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
-// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
-// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
-// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-#include "src/tint/lang/core/ir/transform/common/referenced_module_vars.h"
-
-#include "src/tint/lang/core/ir/control_instruction.h"
-#include "src/tint/lang/core/ir/module.h"
-#include "src/tint/lang/core/ir/user_call.h"
-#include "src/tint/lang/core/ir/var.h"
-#include "src/tint/utils/rtti/switch.h"
-
-namespace tint::core::ir {
-
-const ReferencedModuleVars::VarSet& ReferencedModuleVars::TransitiveReferences(Function* func) {
- return transitive_references_.GetOrAdd(func, [&] {
- VarSet vars;
- GetTransitiveReferences(func->Block(), vars);
- return vars;
- });
-}
-
-/// Get the set of variables transitively referenced by @p block.
-/// @param block the block
-/// @param vars the set of transitively referenced variables to populate
-void ReferencedModuleVars::GetTransitiveReferences(Block* block, VarSet& vars) {
- // Add directly referenced vars.
- if (auto itr = block_to_direct_vars_.Get(block)) {
- for (auto& var : *itr) {
- vars.Add(var);
- }
- }
-
- // Loop over instructions in the block to find indirectly referenced vars.
- for (auto* inst : *block) {
- tint::Switch(
- inst,
- [&](UserCall* call) {
- // Get variables referenced by a function called from this block.
- const auto& callee_vars = TransitiveReferences(call->Target());
- for (auto* var : callee_vars) {
- vars.Add(var);
- }
- },
- [&](ControlInstruction* ctrl) {
- // Recurse into control instructions and gather their referenced vars.
- ctrl->ForeachBlock([&](Block* blk) { GetTransitiveReferences(blk, vars); });
- });
- }
-}
-
-} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/transform/common/referenced_module_vars.h b/src/tint/lang/core/ir/transform/common/referenced_module_vars.h
deleted file mode 100644
index 41e2dd9..0000000
--- a/src/tint/lang/core/ir/transform/common/referenced_module_vars.h
+++ /dev/null
@@ -1,105 +0,0 @@
-// Copyright 2024 The Dawn & Tint Authors
-//
-// Redistribution and use in source and binary forms, with or without
-// modification, are permitted provided that the following conditions are met:
-//
-// 1. Redistributions of source code must retain the above copyright notice, this
-// list of conditions and the following disclaimer.
-//
-// 2. Redistributions in binary form must reproduce the above copyright notice,
-// this list of conditions and the following disclaimer in the documentation
-// and/or other materials provided with the distribution.
-//
-// 3. Neither the name of the copyright holder nor the names of its
-// contributors may be used to endorse or promote products derived from
-// this software without specific prior written permission.
-//
-// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
-// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
-// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
-// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
-// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
-// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
-// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
-// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
-// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-#ifndef SRC_TINT_LANG_CORE_IR_TRANSFORM_COMMON_REFERENCED_MODULE_VARS_H_
-#define SRC_TINT_LANG_CORE_IR_TRANSFORM_COMMON_REFERENCED_MODULE_VARS_H_
-
-#include <functional>
-
-#include "src/tint/lang/core/ir/module.h"
-#include "src/tint/lang/core/ir/var.h"
-#include "src/tint/utils/containers/hashmap.h"
-#include "src/tint/utils/containers/unique_vector.h"
-
-// Forward declarations.
-namespace tint::core::ir {
-class Block;
-class Function;
-} // namespace tint::core::ir
-
-namespace tint::core::ir {
-
-/// ReferencedModuleVars is a helper to determine the set of module-scope variables that are
-/// transitively referenced by functions in a module.
-/// References are determined lazily and cached for future requests.
-///
-/// Note: changes to the module can invalidate the cached data. This is intended to be created by
-/// a transform that need this information, and discarded when that transform completes. Tracking
-/// this information inside the IR module would add overhead any time an instruction is added or
-/// removed from the module. Since only a few transforms need this information, we expect it to be
-/// more efficient to generate it as and when needed instead.
-class ReferencedModuleVars {
- public:
- /// The signature of a predicate used to filter variables.
- /// A predicate function should return `true` when the variable should be added to the set.
- using Predicate = std::function<bool(const Var*)>;
-
- /// A set of a variables referenced by a function (in declaration order).
- using VarSet = UniqueVector<Var*, 16>;
-
- /// Constructor.
- /// @param ir the module
- /// @param pred an optional predicate function for filtering variables
- /// Note: @p pred is not stored by the class, so can be a lambda that captures by reference.
- explicit ReferencedModuleVars(Module& ir, Predicate&& pred = {}) : ir_(ir) {
- // Loop over module-scope variables, recording the blocks that they are referenced from.
- for (auto inst : *ir_.root_block) {
- if (auto* var = inst->As<Var>()) {
- if (!pred || pred(var)) {
- var->Result(0)->ForEachUseUnsorted([&](const Usage& use) {
- block_to_direct_vars_.GetOrAddZero(use.instruction->Block()).Add(var);
- });
- }
- }
- }
- }
-
- /// Get the set of transitively referenced module-scope variables for a function, filtered by
- /// the predicate function if provided.
- /// @param func the function
- /// @returns the set of (possibly filtered) transitively reference module-scope variables
- const VarSet& TransitiveReferences(Function* func);
-
- private:
- /// The module.
- Module& ir_;
-
- /// A map from blocks to their directly referenced variables.
- Hashmap<Block*, VarSet, 64> block_to_direct_vars_{};
-
- /// A map from functions to their transitively referenced variables.
- Hashmap<Function*, VarSet, 8> transitive_references_;
-
- /// Get the set of transitively referenced module-scope variables for a block.
- /// @param block the block
- /// @param vars the set of transitively reference module-scope variables to populate
- void GetTransitiveReferences(Block* block, VarSet& vars);
-};
-
-} // namespace tint::core::ir
-
-#endif // SRC_TINT_LANG_CORE_IR_TRANSFORM_COMMON_REFERENCED_MODULE_VARS_H_
diff --git a/src/tint/lang/core/ir/transform/zero_init_workgroup_memory.cc b/src/tint/lang/core/ir/transform/zero_init_workgroup_memory.cc
index 28511b9..9566183 100644
--- a/src/tint/lang/core/ir/transform/zero_init_workgroup_memory.cc
+++ b/src/tint/lang/core/ir/transform/zero_init_workgroup_memory.cc
@@ -32,7 +32,7 @@
#include "src/tint/lang/core/ir/builder.h"
#include "src/tint/lang/core/ir/module.h"
-#include "src/tint/lang/core/ir/transform/common/referenced_module_vars.h"
+#include "src/tint/lang/core/ir/referenced_module_vars.h"
#include "src/tint/lang/core/ir/validator.h"
#include "src/tint/utils/containers/reverse.h"
@@ -63,7 +63,7 @@
core::type::Manager& ty{ir.Types()};
/// The mapping from functions to their transitively referenced workgroup variables.
- ReferencedModuleVars referenced_module_vars_{
+ ReferencedModuleVars<Module> referenced_module_vars_{
ir, [](const Var* var) {
auto* view = var->Result(0)->Type()->As<type::MemoryView>();
return view && view->AddressSpace() == AddressSpace::kWorkgroup;
diff --git a/src/tint/lang/core/ir/validator.cc b/src/tint/lang/core/ir/validator.cc
index 12a2759..e1d81b0 100644
--- a/src/tint/lang/core/ir/validator.cc
+++ b/src/tint/lang/core/ir/validator.cc
@@ -63,6 +63,7 @@
#include "src/tint/lang/core/ir/member_builtin_call.h"
#include "src/tint/lang/core/ir/multi_in_block.h"
#include "src/tint/lang/core/ir/next_iteration.h"
+#include "src/tint/lang/core/ir/referenced_module_vars.h"
#include "src/tint/lang/core/ir/return.h"
#include "src/tint/lang/core/ir/store.h"
#include "src/tint/lang/core/ir/store_vector_element.h"
@@ -75,6 +76,7 @@
#include "src/tint/lang/core/ir/user_call.h"
#include "src/tint/lang/core/ir/var.h"
#include "src/tint/lang/core/type/bool.h"
+#include "src/tint/lang/core/type/f32.h"
#include "src/tint/lang/core/type/i8.h"
#include "src/tint/lang/core/type/memory_view.h"
#include "src/tint/lang/core/type/pointer.h"
@@ -363,9 +365,23 @@
void CheckRootBlock(const Block* blk);
/// Validates the given function
- /// @param func the function validate
+ /// @param func the function to validate
void CheckFunction(const Function* func);
+ /// Validates the specific function as a vertex entry point
+ /// @param ep the function to validate
+ void CheckVertexEntryPoint(const Function* ep);
+
+ /// Validates that the type annotated with @builtin(position) is correct
+ /// @param ep the entry point to associate errors with
+ /// @param type the type to validate
+ void CheckBuiltinPosition(const Function* ep, const core::type::Type* type);
+
+ /// Validates that the type annotated with @builtin(clip_distances) is correct
+ /// @param ep the entry point to associate errors with
+ /// @param type the type to validate
+ void CheckBuiltinClipDistances(const Function* ep, const core::type::Type* type);
+
/// Validates the given instruction
/// @param inst the instruction to validate
void CheckInstruction(const Instruction* inst);
@@ -612,10 +628,11 @@
Hashmap<const ir::Block*, const ir::Function*, 64> block_to_function_{};
Hashmap<const ir::Function*, Hashset<const ir::UserCall*, 4>, 4> user_func_calls_;
Hashset<const ir::Discard*, 4> discards_;
+ core::ir::ReferencedModuleVars<const Module> referenced_module_vars_;
};
Validator::Validator(const Module& mod, Capabilities capabilities)
- : mod_(mod), capabilities_(capabilities) {}
+ : mod_(mod), capabilities_(capabilities), referenced_module_vars_(mod) {}
Validator::~Validator() = default;
@@ -1176,10 +1193,98 @@
}
}
+ if (func->Stage() == Function::PipelineStage::kVertex) {
+ CheckVertexEntryPoint(func);
+ }
+
QueueBlock(func->Block());
ProcessTasks();
}
+void Validator::CheckVertexEntryPoint(const Function* ep) {
+ const auto* ret_struct = ep->ReturnType()->As<core::type::Struct>();
+ bool contains_position = false;
+ if (ret_struct) {
+ for (auto* mem : ret_struct->Members()) {
+ if (!mem->Attributes().builtin.has_value()) {
+ continue;
+ }
+ switch (mem->Attributes().builtin.value()) {
+ case BuiltinValue::kPosition:
+ contains_position = true;
+ CheckBuiltinPosition(ep, mem->Type());
+ break;
+ case BuiltinValue::kClipDistances:
+ CheckBuiltinClipDistances(ep, mem->Type());
+ break;
+ default:
+ break;
+ }
+ }
+ } else {
+ if (ep->ReturnBuiltin() && ep->ReturnBuiltin() == BuiltinValue::kPosition) {
+ contains_position = true;
+ CheckBuiltinPosition(ep, ep->ReturnType());
+ }
+ }
+
+ for (auto var : referenced_module_vars_.TransitiveReferences(ep)) {
+ const auto* res_type = var->Result(0)->Type()->UnwrapPtrOrRef();
+ const auto* res_struct = res_type->As<core::type::Struct>();
+ if (res_struct) {
+ for (auto* mem : res_struct->Members()) {
+ if (!mem->Attributes().builtin.has_value()) {
+ continue;
+ }
+ switch (mem->Attributes().builtin.value()) {
+ case BuiltinValue::kPosition:
+ contains_position = true;
+ CheckBuiltinPosition(ep, mem->Type()->UnwrapPtrOrRef());
+ break;
+ case BuiltinValue::kClipDistances:
+ CheckBuiltinClipDistances(ep, mem->Type()->UnwrapPtrOrRef());
+ break;
+ default:
+ break;
+ }
+ }
+ } else {
+ if (!var->Attributes().builtin.has_value()) {
+ continue;
+ }
+ switch (var->Attributes().builtin.value()) {
+ case BuiltinValue::kPosition: {
+ contains_position = true;
+ CheckBuiltinPosition(ep, res_type);
+ } break;
+ case BuiltinValue::kClipDistances:
+ CheckBuiltinClipDistances(ep, var->Result(0)->Type()->UnwrapPtrOrRef());
+ break;
+ default:
+ break;
+ }
+ }
+ }
+
+ if (DAWN_UNLIKELY(!contains_position)) {
+ AddError(ep) << "position must be declared for vertex entry point output";
+ }
+}
+
+void Validator::CheckBuiltinPosition(const Function* ep, const core::type::Type* type) {
+ auto elems = type->Elements();
+ if (!type->IsFloatVector() || !elems.type->Is<core::type::F32>() || elems.count != 4) {
+ AddError(ep) << "position must be a vec4<f32>";
+ }
+}
+
+void Validator::CheckBuiltinClipDistances(const Function* ep, const core::type::Type* type) {
+ const auto elems = type->Elements();
+ if (!elems.type || !elems.type->Is<core::type::F32>() || elems.count > 8) {
+ AddError(ep) << "clip_distances must be an array<f32, N>, where N <= 8";
+ }
+}
+
void Validator::ProcessTasks() {
while (!tasks_.IsEmpty()) {
tasks_.Pop()();
@@ -1789,7 +1894,8 @@
// Ensure that values used in the loop continuing are not from the loop body, after a
// continue instruction.
if (auto* first_continue = first_continues_.GetOr(loop, nullptr)) {
- // Find the instruction in the body block that is or holds the first continue instruction.
+ // Find the instruction in the body block that is or holds the first continue
+ // instruction.
const Instruction* holds_continue = first_continue;
while (holds_continue && holds_continue->Block() &&
holds_continue->Block() != loop->Body()) {
@@ -1803,7 +1909,8 @@
if (TransitivelyHolds(loop->Continuing(), use.instruction)) {
AddError(use.instruction, use.operand_index)
<< NameOf(result)
- << " cannot be used in continuing block as it is declared after the "
+ << " cannot be used in continuing block as it is declared after "
+ "the "
"first "
<< style::Instruction("continue") << " in the loop's body";
AddDeclarationNote(result);
@@ -1986,8 +2093,8 @@
auto* func = ret->Func();
if (func == nullptr) {
- // Func() returning nullptr after CheckResultsAndOperandRange is due to the first operand
- // being not a function
+ // Func() returning nullptr after CheckResultsAndOperandRange is due to the first
+ // operand being not a function
AddError(ret) << "expected function for first operand";
return;
}
diff --git a/src/tint/lang/core/ir/validator_test.cc b/src/tint/lang/core/ir/validator_test.cc
index 5c7235f..7b90404 100644
--- a/src/tint/lang/core/ir/validator_test.cc
+++ b/src/tint/lang/core/ir/validator_test.cc
@@ -451,6 +451,133 @@
)");
}
+TEST_F(IR_ValidatorTest, Function_VertexBasicPosition) {
+ auto* f = b.Function("my_func", ty.vec4<f32>());
+ f->SetStage(Function::PipelineStage::kVertex);
+ f->SetReturnBuiltin(BuiltinValue::kPosition);
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_EQ(res, Success);
+}
+
+TEST_F(IR_ValidatorTest, Function_VertexStructPosition) {
+ auto pos_ty = ty.vec4<f32>();
+ auto pos_attr = IOAttributes();
+ pos_attr.builtin = BuiltinValue::kPosition;
+
+ auto* str_ty =
+ ty.Struct(mod.symbols.New("MyStruct"), {
+ {mod.symbols.New("pos"), pos_ty, pos_attr},
+ });
+
+ auto* f = b.Function("my_func", str_ty);
+ f->SetStage(Function::PipelineStage::kVertex);
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_EQ(res, Success);
+}
+
+TEST_F(IR_ValidatorTest, Function_VertexStructPositionAndClipDistances) {
+ auto pos_ty = ty.vec4<f32>();
+ auto pos_attr = IOAttributes();
+ pos_attr.builtin = BuiltinValue::kPosition;
+
+ auto clip_ty = ty.array<f32, 4>();
+ auto clip_attr = IOAttributes();
+ clip_attr.builtin = BuiltinValue::kClipDistances;
+
+ auto* str_ty =
+ ty.Struct(mod.symbols.New("MyStruct"), {
+ {mod.symbols.New("pos"), pos_ty, pos_attr},
+ {mod.symbols.New("clip"), clip_ty, clip_attr},
+ });
+
+ auto* f = b.Function("my_func", str_ty);
+ f->SetStage(Function::PipelineStage::kVertex);
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_EQ(res, Success);
+}
+
+TEST_F(IR_ValidatorTest, Function_VertexStructOnlyClipDistances) {
+ auto clip_ty = ty.array<f32, 4>();
+ auto clip_attr = IOAttributes();
+ clip_attr.builtin = BuiltinValue::kClipDistances;
+
+ auto* str_ty =
+ ty.Struct(mod.symbols.New("MyStruct"), {
+ {mod.symbols.New("clip"), clip_ty, clip_attr},
+ });
+
+ auto* f = b.Function("my_func", str_ty);
+ f->SetStage(Function::PipelineStage::kVertex);
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:5:1 error: position must be declared for vertex entry point output
+%my_func = @vertex func():MyStruct {
+^^^^^^^^
+
+note: # Disassembly
+MyStruct = struct @align(4) {
+ clip:array<f32, 4> @offset(0), @builtin(clip_distances)
+}
+
+%my_func = @vertex func():MyStruct {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Function_VertexMissingPosition) {
+ auto* f = b.Function("my_func", ty.vec4<f32>());
+ f->SetStage(Function::PipelineStage::kVertex);
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:1 error: position must be declared for vertex entry point output
+%my_func = @vertex func():vec4<f32> {
+^^^^^^^^
+
+note: # Disassembly
+%my_func = @vertex func():vec4<f32> {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Function_VertexPositionWrongType) {
+ auto* f = b.Function("my_func", ty.void_());
+ f->SetStage(Function::PipelineStage::kVertex);
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:1 error: position must be declared for vertex entry point output
+%my_func = @vertex func():void {
+^^^^^^^^
+
+note: # Disassembly
+%my_func = @vertex func():void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
TEST_F(IR_ValidatorTest, CallToFunctionOutsideModule) {
auto* f = b.Function("f", ty.void_());
auto* g = b.Function("g", ty.void_());
@@ -1547,7 +1674,8 @@
});
auto* ep = b.Function("ep", ty.void_());
- ep->SetStage(Function::PipelineStage::kVertex);
+ ep->SetStage(Function::PipelineStage::kCompute);
+ ep->SetWorkgroupSize(0, 0, 0);
b.Append(ep->Block(), [&] {
b.Call(func);
b.Return(ep);
@@ -1567,7 +1695,7 @@
ret
}
}
-%ep = @vertex func():void {
+%ep = @compute @workgroup_size(0, 0, 0) func():void {
$B2: {
%3:void = call %foo
ret
diff --git a/src/tint/lang/msl/writer/raise/BUILD.bazel b/src/tint/lang/msl/writer/raise/BUILD.bazel
index e820aac..e1f73fc 100644
--- a/src/tint/lang/msl/writer/raise/BUILD.bazel
+++ b/src/tint/lang/msl/writer/raise/BUILD.bazel
@@ -66,7 +66,6 @@
"//src/tint/lang/core/intrinsic",
"//src/tint/lang/core/ir",
"//src/tint/lang/core/ir/transform",
- "//src/tint/lang/core/ir/transform/common",
"//src/tint/lang/core/type",
"//src/tint/lang/msl",
"//src/tint/lang/msl/intrinsic",
diff --git a/src/tint/lang/msl/writer/raise/BUILD.cmake b/src/tint/lang/msl/writer/raise/BUILD.cmake
index d872922..9f786f3 100644
--- a/src/tint/lang/msl/writer/raise/BUILD.cmake
+++ b/src/tint/lang/msl/writer/raise/BUILD.cmake
@@ -67,7 +67,6 @@
tint_lang_core_intrinsic
tint_lang_core_ir
tint_lang_core_ir_transform
- tint_lang_core_ir_transform_common
tint_lang_core_type
tint_lang_msl
tint_lang_msl_intrinsic
diff --git a/src/tint/lang/msl/writer/raise/BUILD.gn b/src/tint/lang/msl/writer/raise/BUILD.gn
index 33e9d8f..614d731 100644
--- a/src/tint/lang/msl/writer/raise/BUILD.gn
+++ b/src/tint/lang/msl/writer/raise/BUILD.gn
@@ -71,7 +71,6 @@
"${tint_src_dir}/lang/core/intrinsic",
"${tint_src_dir}/lang/core/ir",
"${tint_src_dir}/lang/core/ir/transform",
- "${tint_src_dir}/lang/core/ir/transform/common",
"${tint_src_dir}/lang/core/type",
"${tint_src_dir}/lang/msl",
"${tint_src_dir}/lang/msl/intrinsic",
diff --git a/src/tint/lang/msl/writer/raise/module_scope_vars.cc b/src/tint/lang/msl/writer/raise/module_scope_vars.cc
index efbde47..8727e39 100644
--- a/src/tint/lang/msl/writer/raise/module_scope_vars.cc
+++ b/src/tint/lang/msl/writer/raise/module_scope_vars.cc
@@ -30,7 +30,7 @@
#include <utility>
#include "src/tint/lang/core/ir/builder.h"
-#include "src/tint/lang/core/ir/transform/common/referenced_module_vars.h"
+#include "src/tint/lang/core/ir/referenced_module_vars.h"
#include "src/tint/lang/core/ir/validator.h"
namespace tint::msl::writer::raise {
@@ -62,7 +62,7 @@
Hashmap<core::ir::Block*, core::ir::Function*, 64> block_to_function{};
/// The mapping from functions to their transitively referenced workgroup variables.
- core::ir::ReferencedModuleVars referenced_module_vars{ir};
+ core::ir::ReferencedModuleVars<core::ir::Module> referenced_module_vars{ir};
// The name of the module-scope variables structure.
static constexpr const char* kModuleVarsName = "tint_module_vars";
@@ -172,7 +172,7 @@
/// @returns the structure that holds the module-scope variables
core::ir::Value* AddModuleVarsToEntryPoint(
core::ir::Function* func,
- const core::ir::ReferencedModuleVars::VarSet& referenced_vars) {
+ const core::ir::ReferencedModuleVars<core::ir::Module>::VarSet& referenced_vars) {
core::ir::Value* module_var_struct = nullptr;
core::ir::FunctionParam* workgroup_allocation_param = nullptr;
Vector<core::type::Manager::StructMemberDesc, 4> workgroup_struct_members;
diff --git a/src/tint/lang/msl/writer/raise/simd_ballot.cc b/src/tint/lang/msl/writer/raise/simd_ballot.cc
index 565d85d..1b4c84c 100644
--- a/src/tint/lang/msl/writer/raise/simd_ballot.cc
+++ b/src/tint/lang/msl/writer/raise/simd_ballot.cc
@@ -30,7 +30,7 @@
#include <utility>
#include "src/tint/lang/core/ir/builder.h"
-#include "src/tint/lang/core/ir/transform/common/referenced_module_vars.h"
+#include "src/tint/lang/core/ir/referenced_module_vars.h"
#include "src/tint/lang/core/ir/validator.h"
#include "src/tint/lang/msl/ir/builtin_call.h"
@@ -68,9 +68,10 @@
}
// Set the subgroup size mask value from all entry points that use it.
- core::ir::ReferencedModuleVars refs(ir, [&](const core::ir::Var* var) { //
- return var == subgroup_size_mask;
- });
+ core::ir::ReferencedModuleVars<core::ir::Module> refs(ir,
+ [&](const core::ir::Var* var) { //
+ return var == subgroup_size_mask;
+ });
for (auto func : ir.functions) {
if (func->Stage() != core::ir::Function::PipelineStage::kUndefined) {
if (refs.TransitiveReferences(func).Contains(subgroup_size_mask)) {
diff --git a/src/tint/lang/spirv/reader/lower/BUILD.bazel b/src/tint/lang/spirv/reader/lower/BUILD.bazel
index 5d99ad6..06eec9c 100644
--- a/src/tint/lang/spirv/reader/lower/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/lower/BUILD.bazel
@@ -54,7 +54,6 @@
"//src/tint/lang/core/constant",
"//src/tint/lang/core/intrinsic",
"//src/tint/lang/core/ir",
- "//src/tint/lang/core/ir/transform/common",
"//src/tint/lang/core/type",
"//src/tint/utils/containers",
"//src/tint/utils/diagnostic",
diff --git a/src/tint/lang/spirv/reader/lower/BUILD.cmake b/src/tint/lang/spirv/reader/lower/BUILD.cmake
index 9758df7..3bd17b4 100644
--- a/src/tint/lang/spirv/reader/lower/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/lower/BUILD.cmake
@@ -53,7 +53,6 @@
tint_lang_core_constant
tint_lang_core_intrinsic
tint_lang_core_ir
- tint_lang_core_ir_transform_common
tint_lang_core_type
tint_utils_containers
tint_utils_diagnostic
diff --git a/src/tint/lang/spirv/reader/lower/BUILD.gn b/src/tint/lang/spirv/reader/lower/BUILD.gn
index b9adbf6..07df190 100644
--- a/src/tint/lang/spirv/reader/lower/BUILD.gn
+++ b/src/tint/lang/spirv/reader/lower/BUILD.gn
@@ -59,7 +59,6 @@
"${tint_src_dir}/lang/core/constant",
"${tint_src_dir}/lang/core/intrinsic",
"${tint_src_dir}/lang/core/ir",
- "${tint_src_dir}/lang/core/ir/transform/common",
"${tint_src_dir}/lang/core/type",
"${tint_src_dir}/utils/containers",
"${tint_src_dir}/utils/diagnostic",
diff --git a/src/tint/lang/spirv/reader/lower/shader_io.cc b/src/tint/lang/spirv/reader/lower/shader_io.cc
index 7f402d6..7fc8d18 100644
--- a/src/tint/lang/spirv/reader/lower/shader_io.cc
+++ b/src/tint/lang/spirv/reader/lower/shader_io.cc
@@ -31,7 +31,7 @@
#include "src/tint/lang/core/ir/builder.h"
#include "src/tint/lang/core/ir/module.h"
-#include "src/tint/lang/core/ir/transform/common/referenced_module_vars.h"
+#include "src/tint/lang/core/ir/referenced_module_vars.h"
#include "src/tint/lang/core/ir/validator.h"
namespace tint::spirv::reader::lower {
@@ -62,7 +62,7 @@
Hashset<core::ir::Var*, 4> output_variables{};
/// The mapping from functions to their transitively referenced output variables.
- core::ir::ReferencedModuleVars referenced_output_vars{
+ core::ir::ReferencedModuleVars<core::ir::Module> referenced_output_vars{
ir, [](const core::ir::Var* var) {
auto* view = var->Result(0)->Type()->As<core::type::MemoryView>();
return view && view->AddressSpace() == core::AddressSpace::kOut;
diff --git a/src/tint/lang/spirv/reader/parser/function_test.cc b/src/tint/lang/spirv/reader/parser/function_test.cc
index d344063..1790734 100644
--- a/src/tint/lang/spirv/reader/parser/function_test.cc
+++ b/src/tint/lang/spirv/reader/parser/function_test.cc
@@ -130,7 +130,7 @@
)");
}
-TEST_F(SpirvParserTest, VertexShader) {
+TEST_F(SpirvParserTest, DISABLED_VertexShader) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
diff --git a/src/tint/lang/spirv/writer/function_test.cc b/src/tint/lang/spirv/writer/function_test.cc
index f566d7e..9058a18 100644
--- a/src/tint/lang/spirv/writer/function_test.cc
+++ b/src/tint/lang/spirv/writer/function_test.cc
@@ -143,25 +143,51 @@
}
TEST_F(SpirvWriterTest, Function_EntryPoint_Vertex) {
- auto* func = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kVertex);
+ auto* func = b.Function("main", ty.vec4<f32>(), core::ir::Function::PipelineStage::kVertex);
+ func->SetReturnBuiltin(core::BuiltinValue::kPosition);
b.Append(func->Block(), [&] { //
- b.Return(func);
+ b.Return(func, b.Zero<vec4<f32>>());
});
ASSERT_TRUE(Generate()) << Error() << output_;
EXPECT_INST(R"(
- OpEntryPoint Vertex %main "main"
+ OpEntryPoint Vertex %main "main" %main_position_Output %main___point_size_Output
; Debug Information
- OpName %main "main" ; id %1
+ OpName %main_position_Output "main_position_Output" ; id %1
+ OpName %main___point_size_Output "main___point_size_Output" ; id %5
+ OpName %main_inner "main_inner" ; id %7
+ OpName %main "main" ; id %11
+
+ ; Annotations
+ OpDecorate %main_position_Output BuiltIn Position
+ OpDecorate %main___point_size_Output BuiltIn PointSize
; Types, variables and constants
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+%main_position_Output = OpVariable %_ptr_Output_v4float Output ; BuiltIn Position
+%_ptr_Output_float = OpTypePointer Output %float
+%main___point_size_Output = OpVariable %_ptr_Output_float Output ; BuiltIn PointSize
+ %8 = OpTypeFunction %v4float
+ %10 = OpConstantNull %v4float
%void = OpTypeVoid
- %3 = OpTypeFunction %void
+ %13 = OpTypeFunction %void
+ %float_1 = OpConstant %float 1
+
+ ; Function main_inner
+ %main_inner = OpFunction %v4float None %8
+ %9 = OpLabel
+ OpReturnValue %10
+ OpFunctionEnd
; Function main
- %main = OpFunction %void None %3
- %4 = OpLabel
+ %main = OpFunction %void None %13
+ %14 = OpLabel
+ %15 = OpFunctionCall %v4float %main_inner
+ OpStore %main_position_Output %15 None
+ OpStore %main___point_size_Output %float_1 None
OpReturn
OpFunctionEnd
)");
