[ir] Implement Substitute Overrides
This Cl adds the Substitute overrides transform for the IR. It is not
used by any backend yet, the array type override information still needs
to be substituted.
Bug: 374971092
Change-Id: If1121836256488a029ee4d5201e2df86a5d968a3
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/213474
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/core/ir/builder.h b/src/tint/lang/core/ir/builder.h
index ec9b23a..5f6e4f2 100644
--- a/src/tint/lang/core/ir/builder.h
+++ b/src/tint/lang/core/ir/builder.h
@@ -1749,7 +1749,10 @@
/// @param name the override name
/// @param value the override value
/// @returns the instruction
- template <typename VALUE>
+ template <
+ typename VALUE,
+ typename = std::enable_if_t<
+ !traits::IsTypeOrDerived<std::remove_pointer_t<std::decay_t<VALUE>>, core::type::Type>>>
ir::Override* Override(std::string_view name, VALUE&& value) {
auto* val = Value(std::forward<VALUE>(value));
if (DAWN_UNLIKELY(!val)) {
@@ -1762,6 +1765,48 @@
return override;
}
+ /// Creates a new `override` declaration
+ /// @param src the source
+ /// @param name the override name
+ /// @param value the override value
+ /// @returns the instruction
+ template <
+ typename VALUE,
+ typename = std::enable_if_t<
+ !traits::IsTypeOrDerived<std::remove_pointer_t<std::decay_t<VALUE>>, core::type::Type>>>
+ ir::Override* Override(Source src, std::string_view name, VALUE&& value) {
+ auto* val = Value(std::forward<VALUE>(value));
+ if (DAWN_UNLIKELY(!val)) {
+ TINT_ASSERT(val);
+ return nullptr;
+ }
+ auto* override = Append(ir.CreateInstruction<ir::Override>(InstructionResult(val->Type())));
+ override->SetInitializer(val);
+ ir.SetName(override->Result(0), name);
+ ir.SetSource(override, src);
+ return override;
+ }
+
+ /// Creates a new `override` declaration, with an unassigned value
+ /// @param name the override name
+ /// @param type the override type
+ /// @returns the instruction
+ ir::Override* Override(std::string_view name, const type::Type* type) {
+ return Override(Source{}, name, type);
+ }
+
+ /// Creates a new `override` declaration, with an unassigned value
+ /// @param name the override name
+ /// @param type the override type
+ /// @returns the instruction
+ ir::Override* Override(Source src, std::string_view name, const type::Type* type) {
+ auto* override = ir.CreateInstruction<ir::Override>(InstructionResult(type));
+ ir.SetName(override->Result(0), name);
+ ir.SetSource(override, src);
+ Append(override);
+ return override;
+ }
+
/// Creates a new `override` declaration, with an unassigned value
/// @param type the override type
/// @returns the instruction
diff --git a/src/tint/lang/core/ir/evaluator.cc b/src/tint/lang/core/ir/evaluator.cc
index 397736d..1a89a68 100644
--- a/src/tint/lang/core/ir/evaluator.cc
+++ b/src/tint/lang/core/ir/evaluator.cc
@@ -86,7 +86,10 @@
[&](core::ir::CoreBuiltinCall* c) { return EvalCoreBuiltinCall(c); },
[&](core::ir::CoreUnary* u) { return EvalUnary(u); },
[&](core::ir::Swizzle* s) { return EvalSwizzle(s); }, //
- TINT_ICE_ON_NO_MATCH);
+ [&](Default) {
+ // Treat any unknown instruction as a termination point for trying to eval.
+ return nullptr;
+ });
},
TINT_ICE_ON_NO_MATCH);
}
diff --git a/src/tint/lang/core/ir/transform/BUILD.bazel b/src/tint/lang/core/ir/transform/BUILD.bazel
index e82b9c7..450fc75 100644
--- a/src/tint/lang/core/ir/transform/BUILD.bazel
+++ b/src/tint/lang/core/ir/transform/BUILD.bazel
@@ -60,6 +60,7 @@
"shader_io.cc",
"single_entry_point.cc",
"std140.cc",
+ "substitute_overrides.cc",
"value_to_let.cc",
"vectorize_scalar_matrix_constructors.cc",
"zero_init_workgroup_memory.cc",
@@ -86,6 +87,7 @@
"shader_io.h",
"single_entry_point.h",
"std140.h",
+ "substitute_overrides.h",
"value_to_let.h",
"vectorize_scalar_matrix_constructors.h",
"zero_init_workgroup_memory.h",
@@ -141,6 +143,7 @@
"robustness_test.cc",
"single_entry_point_test.cc",
"std140_test.cc",
+ "substitute_overrides_test.cc",
"value_to_let_test.cc",
"vectorize_scalar_matrix_constructors_test.cc",
"zero_init_workgroup_memory_test.cc",
diff --git a/src/tint/lang/core/ir/transform/BUILD.cmake b/src/tint/lang/core/ir/transform/BUILD.cmake
index 5496c4b..9fa86ab 100644
--- a/src/tint/lang/core/ir/transform/BUILD.cmake
+++ b/src/tint/lang/core/ir/transform/BUILD.cmake
@@ -81,6 +81,8 @@
lang/core/ir/transform/single_entry_point.h
lang/core/ir/transform/std140.cc
lang/core/ir/transform/std140.h
+ lang/core/ir/transform/substitute_overrides.cc
+ lang/core/ir/transform/substitute_overrides.h
lang/core/ir/transform/value_to_let.cc
lang/core/ir/transform/value_to_let.h
lang/core/ir/transform/vectorize_scalar_matrix_constructors.cc
@@ -142,6 +144,7 @@
lang/core/ir/transform/robustness_test.cc
lang/core/ir/transform/single_entry_point_test.cc
lang/core/ir/transform/std140_test.cc
+ lang/core/ir/transform/substitute_overrides_test.cc
lang/core/ir/transform/value_to_let_test.cc
lang/core/ir/transform/vectorize_scalar_matrix_constructors_test.cc
lang/core/ir/transform/zero_init_workgroup_memory_test.cc
@@ -225,6 +228,7 @@
lang/core/ir/transform/rename_conflicts_fuzz.cc
lang/core/ir/transform/robustness_fuzz.cc
lang/core/ir/transform/std140_fuzz.cc
+ lang/core/ir/transform/substitute_overrides_fuzz.cc
lang/core/ir/transform/value_to_let_fuzz.cc
lang/core/ir/transform/vectorize_scalar_matrix_constructors_fuzz.cc
lang/core/ir/transform/zero_init_workgroup_memory_fuzz.cc
diff --git a/src/tint/lang/core/ir/transform/BUILD.gn b/src/tint/lang/core/ir/transform/BUILD.gn
index d3cbca3..945ce31 100644
--- a/src/tint/lang/core/ir/transform/BUILD.gn
+++ b/src/tint/lang/core/ir/transform/BUILD.gn
@@ -87,6 +87,8 @@
"single_entry_point.h",
"std140.cc",
"std140.h",
+ "substitute_overrides.cc",
+ "substitute_overrides.h",
"value_to_let.cc",
"value_to_let.h",
"vectorize_scalar_matrix_constructors.cc",
@@ -142,6 +144,7 @@
"robustness_test.cc",
"single_entry_point_test.cc",
"std140_test.cc",
+ "substitute_overrides_test.cc",
"value_to_let_test.cc",
"vectorize_scalar_matrix_constructors_test.cc",
"zero_init_workgroup_memory_test.cc",
@@ -216,6 +219,7 @@
"rename_conflicts_fuzz.cc",
"robustness_fuzz.cc",
"std140_fuzz.cc",
+ "substitute_overrides_fuzz.cc",
"value_to_let_fuzz.cc",
"vectorize_scalar_matrix_constructors_fuzz.cc",
"zero_init_workgroup_memory_fuzz.cc",
diff --git a/src/tint/lang/core/ir/transform/helper_test.h b/src/tint/lang/core/ir/transform/helper_test.h
index 22ed109..1103791 100644
--- a/src/tint/lang/core/ir/transform/helper_test.h
+++ b/src/tint/lang/core/ir/transform/helper_test.h
@@ -61,6 +61,14 @@
EXPECT_EQ(ir::Validate(mod, capabilities), Success);
}
+ /// Calls the `transform` but return the result instead of validating.
+ /// @param transform_func the transform to run
+ /// @param args the arguments to the transform function
+ template <typename TRANSFORM, typename... ARGS>
+ Result<SuccessType> RunWithFailure(TRANSFORM&& transform_func, ARGS&&... args) {
+ return transform_func(mod, std::forward<ARGS>(args)...);
+ }
+
/// @returns the transformed module as a disassembled string
std::string str() { return "\n" + ir::Disassembler(mod).Plain(); }
diff --git a/src/tint/lang/core/ir/transform/substitute_overrides.cc b/src/tint/lang/core/ir/transform/substitute_overrides.cc
new file mode 100644
index 0000000..943fdaa
--- /dev/null
+++ b/src/tint/lang/core/ir/transform/substitute_overrides.cc
@@ -0,0 +1,252 @@
+// 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/substitute_overrides.h"
+
+#include <functional>
+
+#include "src/tint/lang/core/ir/builder.h"
+#include "src/tint/lang/core/ir/evaluator.h"
+#include "src/tint/lang/core/ir/validator.h"
+#include "src/tint/utils/result/result.h"
+
+using namespace tint::core::fluent_types; // NOLINT
+using namespace tint::core::number_suffixes; // NOLINT
+
+namespace tint::core::ir::transform {
+namespace {
+
+/// PIMPL state for the transform.
+struct State {
+ /// The IR module.
+ Module& ir;
+
+ /// The configuration
+ const SubstituteOverridesConfig& cfg;
+
+ /// The IR builder.
+ Builder b{ir};
+
+ /// The type manager.
+ core::type::Manager& ty{ir.Types()};
+
+ /// Map of override id to value
+ Hashmap<OverrideId, Value*, 8> override_id_to_value_{};
+
+ /// Process the module.
+ Result<SuccessType> Process() {
+ Vector<Instruction*, 8> to_remove;
+ Vector<Value*, 8> values_to_propagate;
+
+ // Note, we don't `Destroy` the overrides when we substitute them. We need them to stay
+ // alive because the `workgroup_size` and `array` usages aren't in the `Usages` list so
+ // haven't been replaced yet.
+
+ // Find all overrides in the root block and replace them
+ for (auto* inst : *ir.root_block) {
+ auto* override = inst->As<core::ir::Override>();
+ if (!override) {
+ // Gather all the non-var instructions which we'll remove
+ if (!inst->Is<core::ir::Var>()) {
+ to_remove.Push(inst);
+ }
+ continue;
+ }
+
+ // Check if the user provided an override for the given ID.
+ auto iter = cfg.map.find(override->OverrideId());
+ if (iter != cfg.map.end()) {
+ auto* replacement = CreateValue(override->Result(0)->Type(), iter->second);
+ ReplaceOverride(override, replacement);
+ values_to_propagate.Push(replacement);
+ to_remove.Push(override);
+ continue;
+ }
+
+ if (override->Initializer() == nullptr) {
+ diag::Diagnostic error{};
+ error.severity = diag::Severity::Error;
+ error.source = ir.SourceOf(override);
+ error << "Initializer not provided for override, and override not overridden.";
+
+ return Failure(error);
+ }
+
+ core::ir::Value* replacement = nullptr;
+ if (override->Initializer()->Is<core::ir::Constant>()) {
+ replacement = override->Initializer();
+ // Remove the initializer such that we don't find the override as a usage when we
+ // try to propagate the replacement.
+ override->SetInitializer(nullptr);
+ } else {
+ auto r = eval::Eval(b, override->Initializer());
+ if (r != Success) {
+ return r.Failure();
+ }
+ replacement = r.Get();
+ }
+ ReplaceOverride(override, replacement);
+ values_to_propagate.Push(replacement);
+ to_remove.Push(override);
+ }
+
+ // Find any workgroup_sizes to replace
+ for (auto func : ir.functions) {
+ if (func->Stage() != core::ir::Function::PipelineStage::kCompute) {
+ continue;
+ }
+
+ auto wgs = func->WorkgroupSize();
+ TINT_ASSERT(wgs.has_value());
+
+ std::array<ir::Value*, 3> new_wg{};
+ for (size_t i = 0; i < 3; ++i) {
+ auto* val = wgs.value()[i];
+
+ if (val->Is<core::ir::Constant>()) {
+ new_wg[i] = val;
+ continue;
+ }
+ auto* res = val->As<core::ir::InstructionResult>();
+ TINT_ASSERT(res);
+
+ core::ir::Value* new_value = nullptr;
+ if (auto* override = res->Instruction()->As<core::ir::Override>()) {
+ auto replacement = override_id_to_value_.Get(override->OverrideId());
+ TINT_ASSERT(replacement);
+ new_value = *replacement;
+ } else {
+ auto r = eval::Eval(b, val);
+ if (r != Success) {
+ return r.Failure();
+ }
+ new_value = r.Get();
+ }
+
+ new_wg[i] = new_value;
+ }
+ func->SetWorkgroupSize(new_wg);
+ }
+
+ // TODO(dsinclair): Replace array type
+
+ // Remove any non-var instruction in the root block
+ for (auto* inst : to_remove) {
+ inst->Destroy();
+ }
+
+ {
+ // Propagate any replaced override instructions up their instruction chains
+ auto res = Propagate(values_to_propagate);
+ if (res != Success) {
+ return res;
+ }
+ }
+
+ return Success;
+ }
+
+ void ReplaceOverride(core::ir::Override* override, core::ir::Value* replacement) {
+ override_id_to_value_.Add(override->OverrideId(), replacement);
+ override->Result(0)->ReplaceAllUsesWith(replacement);
+ }
+
+ Result<SuccessType> Propagate(Vector<core::ir::Value*, 8>& values_to_propagate) {
+ while (!values_to_propagate.IsEmpty()) {
+ auto* value = values_to_propagate.Pop();
+ for (auto usage : value->UsagesSorted()) {
+ // If the instruction has no results, then it was destroyed already and we can just
+ // skip it.
+ if (!usage.instruction->Result(0)) {
+ continue;
+ }
+
+ if (!NeedsEval(usage.instruction)) {
+ continue;
+ }
+
+ auto r = eval::Eval(b, usage.instruction);
+ if (r != Success) {
+ return r.Failure();
+ }
+
+ // The replacement can be a `nullptr` if we try to evaluate something like a `dpdx`
+ // builtin which doesn't have a `@const` annotation.
+ auto* replacement = r.Get();
+ if (!replacement) {
+ continue;
+ }
+
+ usage.instruction->Result(0)->ReplaceAllUsesWith(replacement);
+ values_to_propagate.Push(replacement);
+ usage.instruction->Destroy();
+ }
+ }
+
+ return Success;
+ }
+
+ bool NeedsEval(core::ir::Instruction* inst) {
+ return tint::Switch( //
+ inst, //
+ [&](core::ir::Bitcast*) { return true; }, //
+ [&](core::ir::Access*) { return true; }, //
+ [&](core::ir::Construct*) { return true; }, //
+ [&](core::ir::Convert*) { return true; }, //
+ [&](core::ir::CoreBinary*) { return true; }, //
+ [&](core::ir::CoreBuiltinCall*) { return true; }, //
+ [&](core::ir::CoreUnary*) { return true; }, //
+ [&](core::ir::Swizzle*) { return true; }, //
+ [&](Default) { return false; });
+ }
+
+ Value* CreateValue(const core::type::Type* type, double val) {
+ return tint::Switch(
+ type,
+ [&](const core::type::Bool*) { return b.Constant(!std::equal_to<double>()(val, 0.0)); },
+ [&](const core::type::I32*) { return b.Constant(i32(val)); },
+ [&](const core::type::U32*) { return b.Constant(u32(val)); },
+ [&](const core::type::F32*) { return b.Constant(f32(val)); },
+ [&](const core::type::F16*) { return b.Constant(f16(val)); }, //
+ TINT_ICE_ON_NO_MATCH);
+ }
+};
+
+} // namespace
+
+SubstituteOverridesConfig::SubstituteOverridesConfig() = default;
+
+Result<SuccessType> SubstituteOverrides(Module& ir, const SubstituteOverridesConfig& cfg) {
+ auto result =
+ ValidateAndDumpIfNeeded(ir, "core.SubstituteOverrides", kSubstituteOverridesCapabilities);
+ if (result != Success) {
+ return result;
+ }
+ return State{ir, cfg}.Process();
+}
+
+} // namespace tint::core::ir::transform
diff --git a/src/tint/lang/core/ir/transform/substitute_overrides.h b/src/tint/lang/core/ir/transform/substitute_overrides.h
new file mode 100644
index 0000000..f438181
--- /dev/null
+++ b/src/tint/lang/core/ir/transform/substitute_overrides.h
@@ -0,0 +1,81 @@
+// 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_SUBSTITUTE_OVERRIDES_H_
+#define SRC_TINT_LANG_CORE_IR_TRANSFORM_SUBSTITUTE_OVERRIDES_H_
+
+#include <unordered_map>
+
+#include "src/tint/api/common/override_id.h"
+#include "src/tint/lang/core/ir/validator.h"
+#include "src/tint/utils/reflection/reflection.h"
+#include "src/tint/utils/result/result.h"
+
+// Forward declarations.
+namespace tint::core::ir {
+class Module;
+}
+
+namespace tint::core::ir::transform {
+
+/// The capabilities that the transform can support.
+const core::ir::Capabilities kSubstituteOverridesCapabilities{
+ core::ir::Capability::kAllowOverrides,
+};
+
+/// Configuration options for the transform
+struct SubstituteOverridesConfig {
+ /// Constructor
+ SubstituteOverridesConfig();
+
+ /// Copy constructor
+ SubstituteOverridesConfig(const SubstituteOverridesConfig&) = default;
+
+ /// Destructor
+ ~SubstituteOverridesConfig() = default;
+
+ /// Assignment operator
+ /// @returns this config
+ SubstituteOverridesConfig& operator=(const SubstituteOverridesConfig&) = default;
+
+ /// The map of override identifier to the override value.
+ /// The value is always a double coming into the transform and will be
+ /// converted to the correct type through and initializer.
+ std::unordered_map<OverrideId, double> map;
+
+ /// Reflect the fields of this class so that it can be used by tint::ForeachField()
+ TINT_REFLECT(SubstituteOverridesConfig, map);
+};
+
+/// Substitute overrides to their constant values.
+/// @param module the module to transform
+/// @returns success or failure
+Result<SuccessType> SubstituteOverrides(Module& module, const SubstituteOverridesConfig& cfg);
+
+} // namespace tint::core::ir::transform
+
+#endif // SRC_TINT_LANG_CORE_IR_TRANSFORM_SUBSTITUTE_OVERRIDES_H_
diff --git a/src/tint/lang/core/ir/transform/substitute_overrides_fuzz.cc b/src/tint/lang/core/ir/transform/substitute_overrides_fuzz.cc
new file mode 100644
index 0000000..549de1a
--- /dev/null
+++ b/src/tint/lang/core/ir/transform/substitute_overrides_fuzz.cc
@@ -0,0 +1,52 @@
+// 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/substitute_overrides.h"
+
+#include "src/tint/cmd/fuzz/ir/fuzz.h"
+#include "src/tint/lang/core/ir/module.h"
+#include "src/tint/lang/core/ir/validator.h"
+
+namespace tint::core::ir::transform {
+namespace {
+
+void SubstituteOverridesFuzzer(Module& module, SubstituteOverridesConfig config) {
+ if (auto res = SubstituteOverrides(module, config); res != Success) {
+ return;
+ }
+
+ Capabilities capabilities;
+ if (auto res = Validate(module, capabilities); res != Success) {
+ TINT_ICE() << "result of SubstituteOverrides failed IR validation\n" << res.Failure();
+ }
+}
+
+} // namespace
+} // namespace tint::core::ir::transform
+
+TINT_IR_MODULE_FUZZER(tint::core::ir::transform::SubstituteOverridesFuzzer,
+ tint::core::ir::transform::kSubstituteOverridesCapabilities);
diff --git a/src/tint/lang/core/ir/transform/substitute_overrides_test.cc b/src/tint/lang/core/ir/transform/substitute_overrides_test.cc
new file mode 100644
index 0000000..db0c5ce
--- /dev/null
+++ b/src/tint/lang/core/ir/transform/substitute_overrides_test.cc
@@ -0,0 +1,779 @@
+// 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/substitute_overrides.h"
+
+#include <limits>
+#include <utility>
+
+#include "gtest/gtest.h"
+#include "src/tint/lang/core/ir/transform/helper_test.h"
+
+namespace tint::core::ir::transform {
+namespace {
+
+using namespace tint::core::fluent_types; // NOLINT
+using namespace tint::core::number_suffixes; // NOLINT
+
+using IR_SubstituteOverridesTest = TransformTest;
+
+TEST_F(IR_SubstituteOverridesTest, NoOverridesNoChange) {
+ auto* func = b.Function("foo", ty.void_());
+ func->Block()->Append(b.Return(func));
+
+ auto* expect = R"(
+%foo = func():void {
+ $B1: {
+ ret
+ }
+}
+)";
+
+ SubstituteOverridesConfig cfg{};
+ Run(SubstituteOverrides, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_SubstituteOverridesTest, UnsetOverrideTriggersError) {
+ b.Append(mod.root_block, [&] {
+ auto* o = b.Override(Source{{1, 2}}, "a", ty.i32());
+ o->SetOverrideId({1});
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %a:i32 = override @id(1)
+}
+
+)";
+ EXPECT_EQ(src, str());
+
+ SubstituteOverridesConfig cfg{};
+ auto result = RunWithFailure(SubstituteOverrides, cfg);
+ ASSERT_NE(result, Success);
+ EXPECT_EQ(result.Failure().reason.Str(),
+ R"(1:2 error: Initializer not provided for override, and override not overridden.)");
+}
+
+TEST_F(IR_SubstituteOverridesTest, OverrideWithDefault) {
+ core::ir::Override* o = nullptr;
+ b.Append(mod.root_block, [&] {
+ o = b.Override(Source{{1, 2}}, "a", 2_u);
+ o->SetOverrideId({1});
+ });
+
+ auto* func = b.Function("foo", ty.u32());
+ b.Append(func->Block(), [&] { b.Return(func, o->Result(0)); });
+
+ auto* src = R"(
+$B1: { # root
+ %a:u32 = override, 2u @id(1)
+}
+
+%foo = func():u32 {
+ $B2: {
+ ret %a
+ }
+}
+)";
+
+ auto* expect = R"(
+%foo = func():u32 {
+ $B1: {
+ ret 2u
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ SubstituteOverridesConfig cfg{};
+ Run(SubstituteOverrides, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_SubstituteOverridesTest, OverrideWithDefaultWithOverride) {
+ core::ir::Override* o = nullptr;
+ b.Append(mod.root_block, [&] {
+ o = b.Override(Source{{1, 2}}, "a", 2_u);
+ o->SetOverrideId({1});
+ });
+
+ auto* func = b.Function("foo", ty.u32());
+ b.Append(func->Block(), [&] { b.Return(func, o->Result(0)); });
+
+ auto* src = R"(
+$B1: { # root
+ %a:u32 = override, 2u @id(1)
+}
+
+%foo = func():u32 {
+ $B2: {
+ ret %a
+ }
+}
+)";
+
+ auto* expect = R"(
+%foo = func():u32 {
+ $B1: {
+ ret 55u
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ SubstituteOverridesConfig cfg{};
+ cfg.map[OverrideId{1}] = 55;
+ Run(SubstituteOverrides, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_SubstituteOverridesTest, OverrideWithoutDefaultWithOverride) {
+ core::ir::Override* o = nullptr;
+ b.Append(mod.root_block, [&] {
+ o = b.Override(Source{{1, 2}}, "a", ty.u32());
+ o->SetOverrideId({1});
+ });
+
+ auto* func = b.Function("foo", ty.u32());
+ b.Append(func->Block(), [&] { b.Return(func, o->Result(0)); });
+
+ auto* src = R"(
+$B1: { # root
+ %a:u32 = override @id(1)
+}
+
+%foo = func():u32 {
+ $B2: {
+ ret %a
+ }
+}
+)";
+
+ auto* expect = R"(
+%foo = func():u32 {
+ $B1: {
+ ret 55u
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ SubstituteOverridesConfig cfg{};
+ cfg.map[OverrideId{1}] = 55;
+ Run(SubstituteOverrides, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_SubstituteOverridesTest, OverrideWithComplexInitNoOverrides) {
+ core::ir::Override* o = nullptr;
+ b.Append(mod.root_block, [&] {
+ auto* add = b.Add(ty.u32(), 2_u, 4_u);
+
+ o = b.Override(Source{{1, 2}}, "a", ty.u32());
+ o->SetOverrideId({1});
+ o->SetInitializer(add->Result(0));
+ });
+
+ auto* func = b.Function("foo", ty.u32());
+ b.Append(func->Block(), [&] { b.Return(func, o->Result(0)); });
+
+ auto* src = R"(
+$B1: { # root
+ %1:u32 = add 2u, 4u
+ %a:u32 = override, %1 @id(1)
+}
+
+%foo = func():u32 {
+ $B2: {
+ ret %a
+ }
+}
+)";
+
+ auto* expect = R"(
+%foo = func():u32 {
+ $B1: {
+ ret 6u
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ SubstituteOverridesConfig cfg{};
+ Run(SubstituteOverrides, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_SubstituteOverridesTest, OverrideWithComplexInitComponentOverride) {
+ core::ir::Override* o = nullptr;
+ b.Append(mod.root_block, [&] {
+ auto* add = b.Add(ty.u32(), 2_u, 4_u);
+
+ o = b.Override(Source{{1, 2}}, "a", ty.u32());
+ o->SetOverrideId({1});
+ o->SetInitializer(add->Result(0));
+ });
+
+ auto* func = b.Function("foo", ty.u32());
+ b.Append(func->Block(), [&] { b.Return(func, o->Result(0)); });
+
+ auto* src = R"(
+$B1: { # root
+ %1:u32 = add 2u, 4u
+ %a:u32 = override, %1 @id(1)
+}
+
+%foo = func():u32 {
+ $B2: {
+ ret %a
+ }
+}
+)";
+
+ auto* expect = R"(
+%foo = func():u32 {
+ $B1: {
+ ret 55u
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ SubstituteOverridesConfig cfg{};
+ cfg.map[OverrideId{1}] = 55;
+ Run(SubstituteOverrides, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_SubstituteOverridesTest, OverrideWithComplexIncludingOverride) {
+ core::ir::Override* o = nullptr;
+ b.Append(mod.root_block, [&] {
+ auto* x = b.Override("x", ty.u32());
+ x->SetOverrideId({2});
+
+ auto* add = b.Add(ty.u32(), x, 4_u);
+
+ o = b.Override(Source{{1, 2}}, "a", ty.u32());
+ o->SetOverrideId({1});
+ o->SetInitializer(add->Result(0));
+ });
+
+ auto* func = b.Function("foo", ty.u32());
+ b.Append(func->Block(), [&] { b.Return(func, o->Result(0)); });
+
+ auto* src = R"(
+$B1: { # root
+ %x:u32 = override @id(2)
+ %2:u32 = add %x, 4u
+ %a:u32 = override, %2 @id(1)
+}
+
+%foo = func():u32 {
+ $B2: {
+ ret %a
+ }
+}
+)";
+
+ auto* expect = R"(
+%foo = func():u32 {
+ $B1: {
+ ret 9u
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ SubstituteOverridesConfig cfg{};
+ cfg.map[OverrideId{2}] = 5;
+ Run(SubstituteOverrides, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_SubstituteOverridesTest, OverrideWithComplexGenError) {
+ core::ir::Override* o = nullptr;
+ b.Append(mod.root_block, [&] {
+ auto* x = b.Override("x", ty.f32());
+ x->SetOverrideId({2});
+
+ auto* add = b.Add(ty.f32(), x, f32(std::numeric_limits<float>::max() - 1));
+ b.ir.SetSource(add, Source{{1, 2}});
+
+ o = b.Override("a", ty.f32());
+ o->SetOverrideId({1});
+ o->SetInitializer(add->Result(0));
+ });
+
+ auto* func = b.Function("foo", ty.f32());
+ b.Append(func->Block(), [&] { b.Return(func, o->Result(0)); });
+
+ auto* src = R"(
+$B1: { # root
+ %x:f32 = override @id(2)
+ %2:f32 = add %x, 340282346638528859811704183484516925440.0f
+ %a:f32 = override, %2 @id(1)
+}
+
+%foo = func():f32 {
+ $B2: {
+ ret %a
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ SubstituteOverridesConfig cfg{};
+ cfg.map[OverrideId{2}] = static_cast<double>(std::numeric_limits<float>::max());
+ auto result = RunWithFailure(SubstituteOverrides, cfg);
+ ASSERT_NE(result, Success);
+ EXPECT_EQ(
+ result.Failure().reason.Str(),
+ R"(1:2 error: '340282346638528859811704183484516925440.0 + 340282346638528859811704183484516925440.0' cannot be represented as 'f32')");
+}
+
+TEST_F(IR_SubstituteOverridesTest, OverrideWorkgroupSize) {
+ core::ir::Override* o = nullptr;
+ core::ir::Override* x = nullptr;
+ b.Append(mod.root_block, [&] {
+ x = b.Override("x", ty.u32());
+ x->SetOverrideId({2});
+
+ auto* add = b.Add(ty.u32(), x, 4_u);
+
+ o = b.Override(Source{{1, 2}}, "a", ty.u32());
+ o->SetOverrideId({1});
+ o->SetInitializer(add->Result(0));
+ });
+
+ auto* func = b.ComputeFunction("foo", o, x, o);
+ b.Append(func->Block(), [&] { b.Return(func); });
+
+ auto* src = R"(
+$B1: { # root
+ %x:u32 = override @id(2)
+ %2:u32 = add %x, 4u
+ %a:u32 = override, %2 @id(1)
+}
+
+%foo = @compute @workgroup_size(%a, %x, %a) func():void {
+ $B2: {
+ ret
+ }
+}
+)";
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(9u, 5u, 9u) func():void {
+ $B1: {
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ SubstituteOverridesConfig cfg{};
+ cfg.map[OverrideId{2}] = 5;
+ Run(SubstituteOverrides, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_SubstituteOverridesTest, FunctionExpression) {
+ core::ir::Override* o = nullptr;
+ core::ir::Override* x = nullptr;
+ b.Append(mod.root_block, [&] {
+ x = b.Override("x", ty.u32());
+ x->SetOverrideId({2});
+
+ auto* add = b.Add(ty.u32(), x, 4_u);
+
+ o = b.Override(Source{{1, 2}}, "a", ty.u32());
+ o->SetOverrideId({1});
+ o->SetInitializer(add->Result(0));
+ });
+
+ auto* func = b.ComputeFunction("foo");
+ b.Append(func->Block(), [&] {
+ b.Let("y", b.Divide(ty.u32(), 10_u, x));
+ b.Let("z", b.Multiply(ty.u32(), 5_u, o));
+ b.Return(func);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %x:u32 = override @id(2)
+ %2:u32 = add %x, 4u
+ %a:u32 = override, %2 @id(1)
+}
+
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %5:u32 = div 10u, %x
+ %y:u32 = let %5
+ %7:u32 = mul 5u, %a
+ %z:u32 = let %7
+ ret
+ }
+}
+)";
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %y:u32 = let 2u
+ %z:u32 = let 45u
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ SubstituteOverridesConfig cfg{};
+ cfg.map[OverrideId{2}] = 5;
+ Run(SubstituteOverrides, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_SubstituteOverridesTest, FunctionExpressionNonConstBuiltin) {
+ core::ir::Override* x = nullptr;
+ b.Append(mod.root_block, [&] {
+ x = b.Override("x", ty.f32());
+ x->SetOverrideId({2});
+ });
+
+ auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ b.Append(func->Block(), [&] {
+ b.Let("y", b.Call(ty.f32(), core::BuiltinFn::kDpdx, b.Multiply(ty.f32(), x, 4_f)));
+ b.Return(func);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %x:f32 = override @id(2)
+}
+
+%foo = @fragment func():void {
+ $B2: {
+ %3:f32 = mul %x, 4.0f
+ %4:f32 = dpdx %3
+ %y:f32 = let %4
+ ret
+ }
+}
+)";
+
+ auto* expect = R"(
+%foo = @fragment func():void {
+ $B1: {
+ %2:f32 = dpdx 20.0f
+ %y:f32 = let %2
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ SubstituteOverridesConfig cfg{};
+ cfg.map[OverrideId{2}] = 5;
+ Run(SubstituteOverrides, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_SubstituteOverridesTest, FunctionExpressionMultiOperand) {
+ core::ir::Override* o = nullptr;
+ core::ir::Override* x = nullptr;
+ b.Append(mod.root_block, [&] {
+ x = b.Override("x", ty.u32());
+ x->SetOverrideId({2});
+
+ auto* add = b.Add(ty.u32(), x, 4_u);
+
+ o = b.Override(Source{{1, 2}}, "a", ty.u32());
+ o->SetOverrideId({1});
+ o->SetInitializer(add->Result(0));
+ });
+
+ auto* func = b.ComputeFunction("foo");
+ b.Append(func->Block(), [&] {
+ b.Let("y", b.Divide(ty.u32(), 10_u, o));
+ auto* k = b.Add(ty.u32(), 1_u, b.Multiply(ty.u32(), 2_u, x));
+ b.Let("z", b.Multiply(ty.u32(), k, o));
+ b.Return(func);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %x:u32 = override @id(2)
+ %2:u32 = add %x, 4u
+ %a:u32 = override, %2 @id(1)
+}
+
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %5:u32 = div 10u, %a
+ %y:u32 = let %5
+ %7:u32 = mul 2u, %x
+ %8:u32 = add 1u, %7
+ %9:u32 = mul %8, %a
+ %z:u32 = let %9
+ ret
+ }
+}
+)";
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %y:u32 = let 1u
+ %z:u32 = let 99u
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ SubstituteOverridesConfig cfg{};
+ cfg.map[OverrideId{2}] = 5;
+ Run(SubstituteOverrides, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_SubstituteOverridesTest, FunctionExpressionMultiOperandFlipOrder) {
+ core::ir::Override* o = nullptr;
+ core::ir::Override* x = nullptr;
+ b.Append(mod.root_block, [&] {
+ x = b.Override("x", ty.u32());
+ x->SetOverrideId({2});
+
+ auto* add = b.Add(ty.u32(), x, 4_u);
+
+ o = b.Override(Source{{1, 2}}, "a", ty.u32());
+ o->SetOverrideId({1});
+ o->SetInitializer(add->Result(0));
+ });
+
+ auto* func = b.ComputeFunction("foo");
+ b.Append(func->Block(), [&] {
+ b.Let("y", b.Divide(ty.u32(), 10_u, o));
+ auto* k = b.Add(ty.u32(), 1_u, b.Multiply(ty.u32(), 2_u, o));
+ b.Let("z", b.Multiply(ty.u32(), k, x));
+ b.Return(func);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %x:u32 = override @id(2)
+ %2:u32 = add %x, 4u
+ %a:u32 = override, %2 @id(1)
+}
+
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %5:u32 = div 10u, %a
+ %y:u32 = let %5
+ %7:u32 = mul 2u, %a
+ %8:u32 = add 1u, %7
+ %9:u32 = mul %8, %x
+ %z:u32 = let %9
+ ret
+ }
+}
+)";
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %y:u32 = let 1u
+ %z:u32 = let 95u
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ SubstituteOverridesConfig cfg{};
+ cfg.map[OverrideId{2}] = 5;
+ Run(SubstituteOverrides, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_SubstituteOverridesTest, FunctionExpressionMultiOperandNonConstFn) {
+ core::ir::Override* o = nullptr;
+ core::ir::Override* x = nullptr;
+ b.Append(mod.root_block, [&] {
+ x = b.Override("x", ty.f32());
+ x->SetOverrideId({2});
+
+ auto* add = b.Add(ty.f32(), x, 4_f);
+
+ o = b.Override(Source{{1, 2}}, "a", ty.f32());
+ o->SetOverrideId({1});
+ o->SetInitializer(add->Result(0));
+ });
+
+ auto* func = b.ComputeFunction("foo");
+ b.Append(func->Block(), [&] {
+ b.Let("y", b.Divide(ty.f32(), 10_f, x));
+ auto* k = b.Call(ty.f32(), core::BuiltinFn::kDpdx, x);
+ b.Let("z", b.Multiply(ty.f32(), k, o));
+ b.Return(func);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %x:f32 = override @id(2)
+ %2:f32 = add %x, 4.0f
+ %a:f32 = override, %2 @id(1)
+}
+
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %5:f32 = div 10.0f, %x
+ %y:f32 = let %5
+ %7:f32 = dpdx %x
+ %8:f32 = mul %7, %a
+ %z:f32 = let %8
+ ret
+ }
+}
+)";
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %y:f32 = let 2.0f
+ %3:f32 = dpdx 5.0f
+ %4:f32 = mul %3, 9.0f
+ %z:f32 = let %4
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ SubstituteOverridesConfig cfg{};
+ cfg.map[OverrideId{2}] = 5;
+ Run(SubstituteOverrides, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_SubstituteOverridesTest, FunctionExpressionMultiOperandLet) {
+ core::ir::Override* o = nullptr;
+ core::ir::Override* x = nullptr;
+ b.Append(mod.root_block, [&] {
+ x = b.Override("x", ty.f32());
+ x->SetOverrideId({2});
+
+ auto* add = b.Add(ty.f32(), x, 4_f);
+
+ o = b.Override(Source{{1, 2}}, "a", ty.f32());
+ o->SetOverrideId({1});
+ o->SetInitializer(add->Result(0));
+ });
+
+ auto* func = b.ComputeFunction("foo");
+ b.Append(func->Block(), [&] {
+ b.Let("y", b.Divide(ty.f32(), 10_f, x));
+ auto* k = b.Let("k", b.Call(ty.f32(), core::BuiltinFn::kDpdx, x));
+ b.Let("z", b.Multiply(ty.f32(), k, o));
+ b.Return(func);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %x:f32 = override @id(2)
+ %2:f32 = add %x, 4.0f
+ %a:f32 = override, %2 @id(1)
+}
+
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %5:f32 = div 10.0f, %x
+ %y:f32 = let %5
+ %7:f32 = dpdx %x
+ %k:f32 = let %7
+ %9:f32 = mul %k, %a
+ %z:f32 = let %9
+ ret
+ }
+}
+)";
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %y:f32 = let 2.0f
+ %3:f32 = dpdx 5.0f
+ %k:f32 = let %3
+ %5:f32 = mul %k, 9.0f
+ %z:f32 = let %5
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ SubstituteOverridesConfig cfg{};
+ cfg.map[OverrideId{2}] = 5;
+ Run(SubstituteOverrides, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+// TODO(dsinclair): Support array type overrides
+TEST_F(IR_SubstituteOverridesTest, DISABLED_OverrideArraySize) {
+ FAIL();
+}
+
+} // namespace
+} // namespace tint::core::ir::transform
