[ir] Handle array types in substitute overrides
This CL adds support to Substitute Overrides to handle array types which
use an override as the array count.
Bug: 374971092
Change-Id: Icb64955d44d6733423cb60a5b4af201e7c26d076
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/213916
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/core/ir/BUILD.bazel b/src/tint/lang/core/ir/BUILD.bazel
index 74c159a..5f6dbd9 100644
--- a/src/tint/lang/core/ir/BUILD.bazel
+++ b/src/tint/lang/core/ir/BUILD.bazel
@@ -159,6 +159,7 @@
"//src/tint/lang/core",
"//src/tint/lang/core/constant",
"//src/tint/lang/core/intrinsic",
+ "//src/tint/lang/core/ir/type",
"//src/tint/lang/core/type",
"//src/tint/utils/containers",
"//src/tint/utils/diagnostic",
diff --git a/src/tint/lang/core/ir/BUILD.cmake b/src/tint/lang/core/ir/BUILD.cmake
index 3af8189..ae3de43 100644
--- a/src/tint/lang/core/ir/BUILD.cmake
+++ b/src/tint/lang/core/ir/BUILD.cmake
@@ -36,6 +36,7 @@
include(lang/core/ir/binary/BUILD.cmake)
include(lang/core/ir/transform/BUILD.cmake)
+include(lang/core/ir/type/BUILD.cmake)
################################################################################
# Target: tint_lang_core_ir
@@ -161,6 +162,7 @@
tint_lang_core
tint_lang_core_constant
tint_lang_core_intrinsic
+ tint_lang_core_ir_type
tint_lang_core_type
tint_utils_containers
tint_utils_diagnostic
diff --git a/src/tint/lang/core/ir/BUILD.gn b/src/tint/lang/core/ir/BUILD.gn
index 15a78d0..1fa9bdd 100644
--- a/src/tint/lang/core/ir/BUILD.gn
+++ b/src/tint/lang/core/ir/BUILD.gn
@@ -164,6 +164,7 @@
"${tint_src_dir}/lang/core",
"${tint_src_dir}/lang/core/constant",
"${tint_src_dir}/lang/core/intrinsic",
+ "${tint_src_dir}/lang/core/ir/type",
"${tint_src_dir}/lang/core/type",
"${tint_src_dir}/utils/containers",
"${tint_src_dir}/utils/diagnostic",
diff --git a/src/tint/lang/core/ir/disassembler.cc b/src/tint/lang/core/ir/disassembler.cc
index 4d9091a..59c6cb1 100644
--- a/src/tint/lang/core/ir/disassembler.cc
+++ b/src/tint/lang/core/ir/disassembler.cc
@@ -32,7 +32,6 @@
#include <optional>
#include <string_view>
-#include "src//tint/lang/core/ir/unary.h"
#include "src/tint/lang/core/binary_op.h"
#include "src/tint/lang/core/constant/composite.h"
#include "src/tint/lang/core/constant/scalar.h"
@@ -60,10 +59,14 @@
#include "src/tint/lang/core/ir/switch.h"
#include "src/tint/lang/core/ir/swizzle.h"
#include "src/tint/lang/core/ir/terminate_invocation.h"
+#include "src/tint/lang/core/ir/type/array_count.h"
+#include "src/tint/lang/core/ir/unary.h"
#include "src/tint/lang/core/ir/unreachable.h"
#include "src/tint/lang/core/ir/unused.h"
#include "src/tint/lang/core/ir/user_call.h"
#include "src/tint/lang/core/ir/var.h"
+#include "src/tint/lang/core/type/array.h"
+#include "src/tint/lang/core/type/pointer.h"
#include "src/tint/lang/core/type/struct.h"
#include "src/tint/lang/core/type/type.h"
#include "src/tint/utils/ice/ice.h"
@@ -961,11 +964,37 @@
EmitLine();
}
-StyledText Disassembler::NameOf(const type::Type* ty) {
+StyledText Disassembler::NameOf(const core::type::Type* ty) {
if (!ty) {
return StyledText{} << StyleError("undef");
}
+ auto ary_emit = [&](StyledText& out, const core::type::Array* ary,
+ const core::ir::type::ValueArrayCount* cnt) -> void {
+ out << "array<" << ary->ElemType()->FriendlyName() << ", " << NameOf(cnt->value) << ">";
+ };
+
+ if (auto* ptr = ty->As<core::type::Pointer>()) {
+ if (auto* ary = ty->UnwrapPtr()->As<core::type::Array>()) {
+ if (auto* cnt = ary->Count()->As<core::ir::type::ValueArrayCount>()) {
+ auto out = StyledText{} << "ptr<";
+ if (ptr->AddressSpace() != core::AddressSpace::kUndefined) {
+ out << ptr->AddressSpace() << ", ";
+ }
+ ary_emit(out, ary, cnt);
+ out << ", " << ptr->Access() << ">";
+
+ return out;
+ }
+ }
+ } else if (auto* ary = ty->UnwrapPtr()->As<core::type::Array>()) {
+ if (auto* cnt = ary->Count()->As<core::ir::type::ValueArrayCount>()) {
+ auto out = StyledText{};
+ ary_emit(out, ary, cnt);
+ return out;
+ }
+ }
+
return StyledText{} << StyleType(ty->FriendlyName());
}
diff --git a/src/tint/lang/core/ir/transform/BUILD.bazel b/src/tint/lang/core/ir/transform/BUILD.bazel
index 450fc75..b0ede84 100644
--- a/src/tint/lang/core/ir/transform/BUILD.bazel
+++ b/src/tint/lang/core/ir/transform/BUILD.bazel
@@ -99,6 +99,7 @@
"//src/tint/lang/core/constant",
"//src/tint/lang/core/intrinsic",
"//src/tint/lang/core/ir",
+ "//src/tint/lang/core/ir/type",
"//src/tint/lang/core/type",
"//src/tint/utils/containers",
"//src/tint/utils/diagnostic",
@@ -165,6 +166,7 @@
"//src/tint/lang/core/intrinsic",
"//src/tint/lang/core/ir",
"//src/tint/lang/core/ir/transform",
+ "//src/tint/lang/core/ir/type",
"//src/tint/lang/core/type",
"//src/tint/lang/wgsl",
"//src/tint/lang/wgsl/ast",
diff --git a/src/tint/lang/core/ir/transform/BUILD.cmake b/src/tint/lang/core/ir/transform/BUILD.cmake
index 9fa86ab..2a72ef5 100644
--- a/src/tint/lang/core/ir/transform/BUILD.cmake
+++ b/src/tint/lang/core/ir/transform/BUILD.cmake
@@ -98,6 +98,7 @@
tint_lang_core_constant
tint_lang_core_intrinsic
tint_lang_core_ir
+ tint_lang_core_ir_type
tint_lang_core_type
tint_utils_containers
tint_utils_diagnostic
@@ -158,6 +159,7 @@
tint_lang_core_intrinsic
tint_lang_core_ir
tint_lang_core_ir_transform
+ tint_lang_core_ir_type
tint_lang_core_type
tint_lang_wgsl
tint_lang_wgsl_ast
diff --git a/src/tint/lang/core/ir/transform/BUILD.gn b/src/tint/lang/core/ir/transform/BUILD.gn
index 945ce31..412163b 100644
--- a/src/tint/lang/core/ir/transform/BUILD.gn
+++ b/src/tint/lang/core/ir/transform/BUILD.gn
@@ -104,6 +104,7 @@
"${tint_src_dir}/lang/core/constant",
"${tint_src_dir}/lang/core/intrinsic",
"${tint_src_dir}/lang/core/ir",
+ "${tint_src_dir}/lang/core/ir/type",
"${tint_src_dir}/lang/core/type",
"${tint_src_dir}/utils/containers",
"${tint_src_dir}/utils/diagnostic",
@@ -159,6 +160,7 @@
"${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/type",
"${tint_src_dir}/lang/core/type",
"${tint_src_dir}/lang/wgsl",
"${tint_src_dir}/lang/wgsl/ast",
diff --git a/src/tint/lang/core/ir/transform/substitute_overrides.cc b/src/tint/lang/core/ir/transform/substitute_overrides.cc
index 943fdaa..c4b17c9 100644
--- a/src/tint/lang/core/ir/transform/substitute_overrides.cc
+++ b/src/tint/lang/core/ir/transform/substitute_overrides.cc
@@ -31,7 +31,9 @@
#include "src/tint/lang/core/ir/builder.h"
#include "src/tint/lang/core/ir/evaluator.h"
+#include "src/tint/lang/core/ir/type/array_count.h"
#include "src/tint/lang/core/ir/validator.h"
+#include "src/tint/lang/core/ir/value.h"
#include "src/tint/utils/result/result.h"
using namespace tint::core::fluent_types; // NOLINT
@@ -55,12 +57,13 @@
core::type::Manager& ty{ir.Types()};
/// Map of override id to value
- Hashmap<OverrideId, Value*, 8> override_id_to_value_{};
+ Hashmap<OverrideId, Constant*, 8> override_id_to_value_{};
/// Process the module.
Result<SuccessType> Process() {
Vector<Instruction*, 8> to_remove;
- Vector<Value*, 8> values_to_propagate;
+ Vector<Constant*, 8> values_to_propagate;
+ Vector<core::ir::Var*, 4> vars_with_value_array_count;
// 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
@@ -70,8 +73,14 @@
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>()) {
+ if (auto* var = inst->As<core::ir::Var>()) {
+ if (auto* ary = var->Result(0)->Type()->UnwrapPtr()->As<core::type::Array>()) {
+ if (ary->Count()->Is<core::ir::type::ValueArrayCount>()) {
+ vars_with_value_array_count.Push(var);
+ }
+ }
+ } else {
+ // Gather all the non-var instructions which we'll remove
to_remove.Push(inst);
}
continue;
@@ -80,7 +89,7 @@
// 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);
+ auto* replacement = CreateConstant(override->Result(0)->Type(), iter->second);
ReplaceOverride(override, replacement);
values_to_propagate.Push(replacement);
to_remove.Push(override);
@@ -96,9 +105,8 @@
return Failure(error);
}
- core::ir::Value* replacement = nullptr;
- if (override->Initializer()->Is<core::ir::Constant>()) {
- replacement = override->Initializer();
+ core::ir::Constant* replacement = override->Initializer()->As<core::ir::Constant>();
+ if (replacement) {
// Remove the initializer such that we don't find the override as a usage when we
// try to propagate the replacement.
override->SetInitializer(nullptr);
@@ -131,28 +139,56 @@
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();
+ auto new_value = CalculateOverride(val);
+ if (!new_value.Get()) {
+ return new_value.Failure();
}
-
- new_wg[i] = new_value;
+ new_wg[i] = new_value.Get();
}
func->SetWorkgroupSize(new_wg);
}
- // TODO(dsinclair): Replace array type
+ // Replace array types using overrides
+ for (auto var : vars_with_value_array_count) {
+ auto* old_ptr = var->Result(0)->Type()->As<core::type::Pointer>();
+ TINT_ASSERT(old_ptr);
+
+ auto* old_ty = old_ptr->UnwrapPtr()->As<core::type::Array>();
+ auto* cnt = old_ty->Count()->As<core::ir::type::ValueArrayCount>();
+ TINT_ASSERT(cnt);
+
+ auto new_value = CalculateOverride(cnt->value);
+ if (!new_value.Get()) {
+ return new_value.Failure();
+ }
+
+ uint32_t num_elements = new_value.Get()->Value()->ValueAs<uint32_t>();
+ auto* new_cnt = ty.Get<core::type::ConstantArrayCount>(num_elements);
+ auto* new_ty = ty.Get<core::type::Array>(old_ty->ElemType(), new_cnt, old_ty->Align(),
+ num_elements * old_ty->Stride(),
+ old_ty->Stride(), old_ty->ImplicitStride());
+
+ auto* new_ptr = ty.ptr(old_ptr->AddressSpace(), new_ty, old_ptr->Access());
+ var->Result(0)->SetType(new_ptr);
+
+ // The `Var` type needs to propagate to certain usages.
+ Vector<core::ir::Instruction*, 2> to_replace;
+ to_replace.Push(var);
+
+ while (!to_replace.IsEmpty()) {
+ auto* inst = to_replace.Pop();
+
+ for (auto usage : inst->Result(0)->UsagesUnsorted()) {
+ if (!usage->instruction->Is<core::ir::Let>()) {
+ continue;
+ }
+
+ usage->instruction->Result(0)->SetType(new_ptr);
+ to_replace.Push(usage->instruction);
+ }
+ }
+ }
// Remove any non-var instruction in the root block
for (auto* inst : to_remove) {
@@ -170,12 +206,31 @@
return Success;
}
- void ReplaceOverride(core::ir::Override* override, core::ir::Value* replacement) {
+ Result<core::ir::Constant*> CalculateOverride(core::ir::Value* val) {
+ auto* count_value = val->As<core::ir::InstructionResult>();
+ TINT_ASSERT(count_value);
+
+ if (auto* override = count_value->Instruction()->As<core::ir::Override>()) {
+ auto replacement = override_id_to_value_.Get(override->OverrideId());
+ TINT_ASSERT(replacement);
+ return *replacement;
+ }
+ auto r = eval::Eval(b, count_value);
+ if (r != Success) {
+ return r.Failure();
+ }
+ // Must be able to evaluate the constant.
+ TINT_ASSERT(r.Get());
+
+ return r;
+ }
+
+ void ReplaceOverride(core::ir::Override* override, core::ir::Constant* 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) {
+ Result<SuccessType> Propagate(Vector<core::ir::Constant*, 8>& values_to_propagate) {
while (!values_to_propagate.IsEmpty()) {
auto* value = values_to_propagate.Pop();
for (auto usage : value->UsagesSorted()) {
@@ -224,7 +279,7 @@
[&](Default) { return false; });
}
- Value* CreateValue(const core::type::Type* type, double val) {
+ Constant* CreateConstant(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)); },
diff --git a/src/tint/lang/core/ir/transform/substitute_overrides_test.cc b/src/tint/lang/core/ir/transform/substitute_overrides_test.cc
index db0c5ce..e93841a 100644
--- a/src/tint/lang/core/ir/transform/substitute_overrides_test.cc
+++ b/src/tint/lang/core/ir/transform/substitute_overrides_test.cc
@@ -32,6 +32,7 @@
#include "gtest/gtest.h"
#include "src/tint/lang/core/ir/transform/helper_test.h"
+#include "src/tint/lang/core/ir/type/array_count.h"
namespace tint::core::ir::transform {
namespace {
@@ -770,9 +771,156 @@
EXPECT_EQ(expect, str());
}
-// TODO(dsinclair): Support array type overrides
-TEST_F(IR_SubstituteOverridesTest, DISABLED_OverrideArraySize) {
- FAIL();
+TEST_F(IR_SubstituteOverridesTest, OverrideArraySize) {
+ b.Append(mod.root_block, [&] {
+ auto* x = b.Override("x", ty.u32());
+ x->SetOverrideId({2});
+
+ auto* cnt = ty.Get<core::ir::type::ValueArrayCount>(x->Result(0));
+ auto* ary = ty.Get<core::type::Array>(ty.i32(), cnt, 4_u, 4_u, 4_u, 4_u);
+ b.Var("v", ty.ptr(core::AddressSpace::kWorkgroup, ary, core::Access::kReadWrite));
+ });
+
+ auto* func = b.ComputeFunction("foo");
+ b.Append(func->Block(), [&] { b.Return(func); });
+
+ auto* src = R"(
+$B1: { # root
+ %x:u32 = override @id(2)
+ %v:ptr<workgroup, array<i32, %x>, read_write> = var
+}
+
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ ret
+ }
+}
+)";
+
+ auto* expect = R"(
+$B1: { # root
+ %v:ptr<workgroup, array<i32, 5>, read_write> = var
+}
+
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ SubstituteOverridesConfig cfg{};
+ cfg.map[OverrideId{2}] = 5;
+ Run(SubstituteOverrides, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_SubstituteOverridesTest, OverrideArraySizeExpression) {
+ b.Append(mod.root_block, [&] {
+ auto* x = b.Override("x", ty.u32());
+ x->SetOverrideId({2});
+
+ auto* inst = b.Multiply(ty.u32(), x, 2_u);
+ auto* cnt = ty.Get<core::ir::type::ValueArrayCount>(inst->Result(0));
+ auto* ary = ty.Get<core::type::Array>(ty.i32(), cnt, 4_u, 4_u, 4_u, 4_u);
+ b.Var("v", ty.ptr(core::AddressSpace::kWorkgroup, ary, core::Access::kReadWrite));
+ });
+
+ auto* func = b.ComputeFunction("foo");
+ b.Append(func->Block(), [&] { b.Return(func); });
+
+ auto* src = R"(
+$B1: { # root
+ %x:u32 = override @id(2)
+ %2:u32 = mul %x, 2u
+ %v:ptr<workgroup, array<i32, %2>, read_write> = var
+}
+
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ ret
+ }
+}
+)";
+
+ auto* expect = R"(
+$B1: { # root
+ %v:ptr<workgroup, array<i32, 10>, read_write> = var
+}
+
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ SubstituteOverridesConfig cfg{};
+ cfg.map[OverrideId{2}] = 5;
+ Run(SubstituteOverrides, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_SubstituteOverridesTest, OverrideArraySizeIntoLet) {
+ core::ir::Var* v = nullptr;
+ b.Append(mod.root_block, [&] {
+ auto* x = b.Override("x", ty.u32());
+ x->SetOverrideId({2});
+
+ auto* cnt = ty.Get<core::ir::type::ValueArrayCount>(x->Result(0));
+ auto* ary = ty.Get<core::type::Array>(ty.i32(), cnt, 4_u, 4_u, 4_u, 4_u);
+ v = b.Var("v", ty.ptr(core::AddressSpace::kWorkgroup, ary, core::Access::kReadWrite));
+ });
+
+ auto* func = b.ComputeFunction("foo");
+ b.Append(func->Block(), [&] {
+ auto* y = b.Let("y", v);
+ b.Let("z", y);
+ b.Return(func);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %x:u32 = override @id(2)
+ %v:ptr<workgroup, array<i32, %x>, read_write> = var
+}
+
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %y:ptr<workgroup, array<i32, %x>, read_write> = let %v
+ %z:ptr<workgroup, array<i32, %x>, read_write> = let %y
+ ret
+ }
+}
+)";
+
+ auto* expect = R"(
+$B1: { # root
+ %v:ptr<workgroup, array<i32, 5>, read_write> = var
+}
+
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %y:ptr<workgroup, array<i32, 5>, read_write> = let %v
+ %z:ptr<workgroup, array<i32, 5>, read_write> = let %y
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ SubstituteOverridesConfig cfg{};
+ cfg.map[OverrideId{2}] = 5;
+ Run(SubstituteOverrides, cfg);
+
+ EXPECT_EQ(expect, str());
}
} // namespace
diff --git a/src/tint/lang/core/ir/type/BUILD.bazel b/src/tint/lang/core/ir/type/BUILD.bazel
new file mode 100644
index 0000000..b107203
--- /dev/null
+++ b/src/tint/lang/core/ir/type/BUILD.bazel
@@ -0,0 +1,64 @@
+# 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 = "type",
+ srcs = [
+ "array_count.cc",
+ ],
+ hdrs = [
+ "array_count.h",
+ ],
+ deps = [
+ "//src/tint/lang/core/type",
+ "//src/tint/utils/containers",
+ "//src/tint/utils/ice",
+ "//src/tint/utils/id",
+ "//src/tint/utils/macros",
+ "//src/tint/utils/math",
+ "//src/tint/utils/memory",
+ "//src/tint/utils/rtti",
+ "//src/tint/utils/symbol",
+ "//src/tint/utils/text",
+ "//src/tint/utils/traits",
+ "//src/utils",
+ ],
+ copts = COPTS,
+ visibility = ["//visibility:public"],
+)
+
diff --git a/src/tint/lang/core/ir/type/BUILD.cmake b/src/tint/lang/core/ir/type/BUILD.cmake
new file mode 100644
index 0000000..4178a14
--- /dev/null
+++ b/src/tint/lang/core/ir/type/BUILD.cmake
@@ -0,0 +1,62 @@
+# 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_type
+# Kind: lib
+################################################################################
+tint_add_target(tint_lang_core_ir_type lib
+ lang/core/ir/type/array_count.cc
+ lang/core/ir/type/array_count.h
+)
+
+tint_target_add_dependencies(tint_lang_core_ir_type lib
+ tint_lang_core_type
+ tint_utils_containers
+ tint_utils_ice
+ tint_utils_id
+ tint_utils_macros
+ tint_utils_math
+ tint_utils_memory
+ tint_utils_rtti
+ tint_utils_symbol
+ tint_utils_text
+ tint_utils_traits
+)
+
+tint_target_add_external_dependencies(tint_lang_core_ir_type lib
+ "src_utils"
+)
diff --git a/src/tint/lang/core/ir/type/BUILD.gn b/src/tint/lang/core/ir/type/BUILD.gn
new file mode 100644
index 0000000..f9aa2c7
--- /dev/null
+++ b/src/tint/lang/core/ir/type/BUILD.gn
@@ -0,0 +1,61 @@
+# 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")
+
+libtint_source_set("type") {
+ sources = [
+ "array_count.cc",
+ "array_count.h",
+ ]
+ deps = [
+ "${dawn_root}/src/utils:utils",
+ "${tint_src_dir}/lang/core/type",
+ "${tint_src_dir}/utils/containers",
+ "${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/rtti",
+ "${tint_src_dir}/utils/symbol",
+ "${tint_src_dir}/utils/text",
+ "${tint_src_dir}/utils/traits",
+ ]
+}
diff --git a/src/tint/lang/core/ir/type/array_count.cc b/src/tint/lang/core/ir/type/array_count.cc
new file mode 100644
index 0000000..d82925b
--- /dev/null
+++ b/src/tint/lang/core/ir/type/array_count.cc
@@ -0,0 +1,54 @@
+// 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/type/array_count.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::core::ir::type::ValueArrayCount);
+
+namespace tint::core::ir::type {
+
+ValueArrayCount::ValueArrayCount(Value* val)
+ : Base(static_cast<size_t>(tint::TypeCode::Of<ValueArrayCount>().bits)), value(val) {}
+
+ValueArrayCount::~ValueArrayCount() = default;
+
+bool ValueArrayCount::Equals(const UniqueNode& other) const {
+ if (auto* v = other.As<ValueArrayCount>()) {
+ return value == v->value;
+ }
+ return false;
+}
+
+std::string ValueArrayCount::FriendlyName() const {
+ return "<value>";
+}
+
+core::type::ArrayCount* ValueArrayCount::Clone(core::type::CloneContext&) const {
+ TINT_UNREACHABLE() << "Value array count clone not available";
+}
+
+} // namespace tint::core::ir::type
diff --git a/src/tint/lang/core/ir/type/array_count.h b/src/tint/lang/core/ir/type/array_count.h
new file mode 100644
index 0000000..4f2a650
--- /dev/null
+++ b/src/tint/lang/core/ir/type/array_count.h
@@ -0,0 +1,67 @@
+// 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_TYPE_ARRAY_COUNT_H_
+#define SRC_TINT_LANG_CORE_IR_TYPE_ARRAY_COUNT_H_
+
+#include <string>
+
+#include "src/tint/lang/core/type/array_count.h"
+
+// Predeclarations
+namespace tint::core::ir {
+class Value;
+}
+
+namespace tint::core::ir::type {
+
+/// The variant of an ArrayCount when the count is an ir `Value`.
+class ValueArrayCount final : public Castable<ValueArrayCount, core::type::ArrayCount> {
+ public:
+ /// Constructor
+ /// @param val the value
+ explicit ValueArrayCount(Value* val);
+ ~ValueArrayCount() override;
+
+ /// @param other the other node
+ /// @returns true if this array count is equal @p other
+ bool Equals(const core::type::UniqueNode& other) const override;
+
+ /// @returns the friendly name for this array count
+ std::string FriendlyName() const override;
+
+ /// @param ctx the clone context
+ /// @returns a clone of this type
+ core::type::ArrayCount* Clone(core::type::CloneContext& ctx) const override;
+
+ /// The value
+ Value* value;
+};
+
+} // namespace tint::core::ir::type
+
+#endif // SRC_TINT_LANG_CORE_IR_TYPE_ARRAY_COUNT_H_
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/BUILD.bazel b/src/tint/lang/wgsl/reader/program_to_ir/BUILD.bazel
index 456ac24..17f5563 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/BUILD.bazel
+++ b/src/tint/lang/wgsl/reader/program_to_ir/BUILD.bazel
@@ -50,6 +50,7 @@
"//src/tint/lang/core/constant",
"//src/tint/lang/core/intrinsic",
"//src/tint/lang/core/ir",
+ "//src/tint/lang/core/ir/type",
"//src/tint/lang/core/type",
"//src/tint/lang/wgsl",
"//src/tint/lang/wgsl/ast",
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/BUILD.cmake b/src/tint/lang/wgsl/reader/program_to_ir/BUILD.cmake
index a8161ee..783dc58 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/BUILD.cmake
+++ b/src/tint/lang/wgsl/reader/program_to_ir/BUILD.cmake
@@ -51,6 +51,7 @@
tint_lang_core_constant
tint_lang_core_intrinsic
tint_lang_core_ir
+ tint_lang_core_ir_type
tint_lang_core_type
tint_lang_wgsl
tint_lang_wgsl_ast
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/BUILD.gn b/src/tint/lang/wgsl/reader/program_to_ir/BUILD.gn
index a6c64f9..6d5b319 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/BUILD.gn
+++ b/src/tint/lang/wgsl/reader/program_to_ir/BUILD.gn
@@ -55,6 +55,7 @@
"${tint_src_dir}/lang/core/constant",
"${tint_src_dir}/lang/core/intrinsic",
"${tint_src_dir}/lang/core/ir",
+ "${tint_src_dir}/lang/core/ir/type",
"${tint_src_dir}/lang/core/type",
"${tint_src_dir}/lang/wgsl",
"${tint_src_dir}/lang/wgsl/ast",
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/ir_program_test.h b/src/tint/lang/wgsl/reader/program_to_ir/ir_program_test.h
index 2ce447c..0761c34 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/ir_program_test.h
+++ b/src/tint/lang/wgsl/reader/program_to_ir/ir_program_test.h
@@ -66,8 +66,7 @@
return lower.Failure();
}
- auto validate = core::ir::Validate(
- result.Get(), core::ir::Capabilities{core::ir::Capability::kAllowOverrides});
+ auto validate = core::ir::Validate(result.Get(), kCapabilities);
if (validate != Success) {
return validate.Failure();
}
@@ -81,13 +80,16 @@
Source::File file("test.wgsl", std::move(wgsl));
auto result = wgsl::reader::WgslToIR(&file);
if (result == Success) {
- auto validated = core::ir::Validate(result.Get());
+ auto validated = core::ir::Validate(result.Get(), kCapabilities);
if (validated != Success) {
return validated.Failure();
}
}
return result;
}
+
+ core::ir::Capabilities kCapabilities =
+ core::ir::Capabilities{core::ir::Capability::kAllowOverrides};
};
using IRProgramTest = IRProgramTestBase<testing::Test>;
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
index 66ae2d4..83ed512 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
@@ -40,6 +40,7 @@
#include "src/tint/lang/core/ir/loop.h"
#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/core/ir/switch.h"
+#include "src/tint/lang/core/ir/type/array_count.h"
#include "src/tint/lang/core/ir/value.h"
#include "src/tint/lang/core/type/pointer.h"
#include "src/tint/lang/core/type/reference.h"
@@ -87,6 +88,7 @@
#include "src/tint/lang/wgsl/ast/while_statement.h"
#include "src/tint/lang/wgsl/ir/builtin_call.h"
#include "src/tint/lang/wgsl/program/program.h"
+#include "src/tint/lang/wgsl/sem/array_count.h"
#include "src/tint/lang/wgsl/sem/builtin_enum_expression.h"
#include "src/tint/lang/wgsl/sem/builtin_fn.h"
#include "src/tint/lang/wgsl/sem/call.h"
@@ -914,9 +916,8 @@
if (!val) {
return;
}
- auto* sem = impl.program_.Sem().Get(expr);
- auto* ty = sem->Type()->Clone(impl.clone_ctx_.type_ctx);
core::ir::Instruction* inst = nullptr;
+ auto* sem = impl.program_.Sem().Get(expr);
switch (expr->op) {
case core::UnaryOp::kAddressOf:
case core::UnaryOp::kIndirection:
@@ -924,15 +925,21 @@
// pointer.
Bind(expr, val);
return;
- case core::UnaryOp::kComplement:
+ case core::UnaryOp::kComplement: {
+ auto* ty = sem->Type()->Clone(impl.clone_ctx_.type_ctx);
inst = impl.builder_.Complement(ty, val);
break;
- case core::UnaryOp::kNegation:
+ }
+ case core::UnaryOp::kNegation: {
+ auto* ty = sem->Type()->Clone(impl.clone_ctx_.type_ctx);
inst = impl.builder_.Negation(ty, val);
break;
- case core::UnaryOp::kNot:
+ }
+ case core::UnaryOp::kNot: {
+ auto* ty = sem->Type()->Clone(impl.clone_ctx_.type_ctx);
inst = impl.builder_.Not(ty, val);
break;
+ }
}
impl.current_block_->Append(inst);
Bind(expr, inst->Result(0));
@@ -1173,9 +1180,38 @@
var,
[&](const ast::Var* v) {
auto* ref = sem->Type()->As<core::type::Reference>();
- auto* ty = builder_.ir.Types().Get<core::type::Pointer>(
- ref->AddressSpace(), ref->StoreType()->Clone(clone_ctx_.type_ctx),
- ref->Access());
+ const core::type::Type* store_ty = nullptr;
+
+ const auto* ary = ref->StoreType()->As<core::type::Array>();
+ // If the array has an override count
+ if (ary && !ary->Count()
+ ->IsAnyOf<core::type::RuntimeArrayCount,
+ core::type::ConstantArrayCount>()) {
+ core::ir::Value* count = tint::Switch(
+ ary->Count(), //
+ [&](const sem::UnnamedOverrideArrayCount* u) {
+ return EmitValueExpression(u->expr->Declaration());
+ },
+ [&](const sem::NamedOverrideArrayCount* n) {
+ return scopes_.Get(n->variable->Declaration()->name->symbol);
+ },
+ TINT_ICE_ON_NO_MATCH);
+
+ if (!count) {
+ return;
+ }
+
+ auto* ary_count =
+ builder_.ir.Types().Get<core::ir::type::ValueArrayCount>(count);
+ store_ty = builder_.ir.Types().Get<core::type::Array>(
+ ary->ElemType()->Clone(clone_ctx_.type_ctx), ary_count, ary->Align(),
+ ary->Size(), ary->Stride(), ary->ImplicitStride());
+ } else {
+ store_ty = ref->StoreType()->Clone(clone_ctx_.type_ctx);
+ }
+
+ auto* ty = builder_.ir.Types().Get<core::type::Pointer>(ref->AddressSpace(),
+ store_ty, ref->Access());
auto* val = builder_.Var(ty);
if (v->initializer) {
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
index 60722d7..07af5e2 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
@@ -1264,5 +1264,32 @@
)");
}
+TEST_F(IR_FromProgramTest, OverrideWithLetAddressOf) {
+ auto* src = R"(
+override x = 1;
+var<workgroup> arr : array<u32, x>;
+
+fn a() {
+ let y = &arr;
+}
+)";
+ auto res = Build(src);
+ ASSERT_EQ(res, Success);
+
+ auto m = res.Move();
+ EXPECT_EQ(core::ir::Disassembler(m).Plain(), R"($B1: { # root
+ %x:i32 = override, 1i @id(0)
+ %arr:ptr<workgroup, array<u32, %x>, read_write> = var
+}
+
+%a = func():void {
+ $B2: {
+ %y:ptr<workgroup, array<u32, %x>, read_write> = let %arr
+ ret
+ }
+}
+)");
+}
+
} // namespace
} // namespace tint::wgsl::reader
