Enable Integer Range Analysis on `Access` in Robustness Transform
This patch adds the support of Integer Range Analysis in Robustness
transform behind the `RobustnessConfig.use_integer_range_analysis`
option and uses Integer Range Analysis on the `Access` expression to
an array.
Bug: 348701956
Test: tint_unittests
Change-Id: Id6976e7fc04704324d776fd184238fb88036db62
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/242975
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Jiawei Shao <jiawei.shao@intel.com>
diff --git a/src/tint/lang/core/ir/transform/robustness.cc b/src/tint/lang/core/ir/transform/robustness.cc
index 73c41ec..9d716c9 100644
--- a/src/tint/lang/core/ir/transform/robustness.cc
+++ b/src/tint/lang/core/ir/transform/robustness.cc
@@ -28,7 +28,9 @@
#include "src/tint/lang/core/ir/transform/robustness.h"
#include <algorithm>
+#include <optional>
+#include "src/tint/lang/core/ir/analysis/integer_range_analysis.h"
#include "src/tint/lang/core/ir/builder.h"
#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/core/ir/validator.h"
@@ -57,6 +59,9 @@
/// The type manager.
core::type::Manager& ty{ir.Types()};
+ /// For integer range analysis when needed
+ std::optional<ir::analysis::IntegerRangeAnalysis> integer_range_analysis = {};
+
/// Process the module.
void Process() {
// Find the access instructions that may need to be clamped.
@@ -65,6 +70,11 @@
Vector<ir::StoreVectorElement*, 64> vector_stores;
Vector<ir::CoreBuiltinCall*, 64> subgroup_matrix_calls;
Vector<ir::CoreBuiltinCall*, 64> texture_calls;
+
+ if (config.use_integer_range_analysis) {
+ integer_range_analysis.emplace(&ir);
+ }
+
for (auto* inst : ir.Instructions()) {
tint::Switch(
inst, //
@@ -215,6 +225,49 @@
inst->SetOperand(op_idx, clamped_idx);
}
+ /// Check if operand @p idx may be less than 0 or greater than @p limit with integer range
+ /// analysis algorithm.
+ /// @param idx the index to check
+ /// @param limit the upper limit @idx to compare with.
+ /// @returns true when @idx may be out of bound, false otherwise
+ bool IndexMayOutOfBound(ir::Value* idx, ir::Value* limit) {
+ // Return true when integer range analysis is disabled.
+ if (!integer_range_analysis.has_value()) {
+ return true;
+ }
+
+ // Return true when `limit` is not a constant value.
+ auto* const_limit = limit->As<ir::Constant>();
+ if (!const_limit) {
+ return true;
+ }
+
+ // Return true when we cannot get a valid range for `idx`.
+ const auto* integer_range = integer_range_analysis->GetInfo(idx);
+ if (!integer_range) {
+ return true;
+ }
+
+ TINT_ASSERT(const_limit->Value()->Type()->Is<type::U32>());
+ uint32_t const_limit_value = const_limit->Value()->ValueAs<uint32_t>();
+
+ using SignedIntegerRange = ir::analysis::IntegerRangeInfo::SignedIntegerRange;
+ using UnsignedIntegerRange = ir::analysis::IntegerRangeInfo::UnsignedIntegerRange;
+
+ // Return true when `idx` may be negative or the upper bound of `idx` is greater than
+ // `limit`.
+ if (std::holds_alternative<UnsignedIntegerRange>(integer_range->range)) {
+ UnsignedIntegerRange range = std::get<UnsignedIntegerRange>(integer_range->range);
+ return range.max_bound > static_cast<uint64_t>(const_limit_value);
+ } else {
+ SignedIntegerRange range = std::get<SignedIntegerRange>(integer_range->range);
+ if (range.min_bound < 0) {
+ return true;
+ }
+ return range.max_bound > static_cast<int64_t>(const_limit_value);
+ }
+ }
+
/// Clamp the indices of an access instruction to ensure they are within the limits of the types
/// that they are indexing into.
/// @param access the access instruction
@@ -262,7 +315,7 @@
});
// If there's a dynamic limit that needs enforced, clamp the index operand.
- if (limit) {
+ if (limit && IndexMayOutOfBound(idx, limit)) {
ClampOperand(access, ir::Access::kIndicesOperandOffset + i, limit);
}
diff --git a/src/tint/lang/core/ir/transform/robustness.h b/src/tint/lang/core/ir/transform/robustness.h
index 4371a80..2f0044a 100644
--- a/src/tint/lang/core/ir/transform/robustness.h
+++ b/src/tint/lang/core/ir/transform/robustness.h
@@ -72,6 +72,9 @@
/// Should the transform skip index clamping on runtime-sized arrays?
bool disable_runtime_sized_array_index_clamping = false;
+ /// Should the integer range analysis be used before doing index clamping?
+ bool use_integer_range_analysis = false;
+
/// Reflection for this class
TINT_REFLECT(RobustnessConfig,
clamp_value,
@@ -84,7 +87,8 @@
clamp_workgroup,
predicate_subgroup_matrix,
bindings_ignored,
- disable_runtime_sized_array_index_clamping);
+ disable_runtime_sized_array_index_clamping,
+ use_integer_range_analysis);
};
/// Robustness is a transform that prevents out-of-bounds memory accesses.
diff --git a/src/tint/lang/core/ir/transform/robustness_test.cc b/src/tint/lang/core/ir/transform/robustness_test.cc
index 120bc54..f07d30e 100644
--- a/src/tint/lang/core/ir/transform/robustness_test.cc
+++ b/src/tint/lang/core/ir/transform/robustness_test.cc
@@ -61,6 +61,8 @@
}
using IR_BindingVariableRobustnessTest = TransformTestWithParam<BindingVariableCase>;
+using IR_RobustnessWithIntegerRangeAnalysisTest = TransformTest;
+
////////////////////////////////////////////////////////////////
// These tests use the function address space.
// Test clamping of vectors, matrices, and fixed-size arrays.
@@ -4171,6 +4173,808 @@
EXPECT_EQ(expect, str());
}
+TEST_F(IR_RobustnessWithIntegerRangeAnalysisTest, AccessArrayWithIndex_MaxBound_Equal_Limit) {
+ auto* func = b.Function("func", ty.void_());
+ b.Append(func->Block(), [&] {
+ Var* idx = nullptr;
+ // arr = array<u32, 20>
+ auto* arr = b.Var("arr", ty.ptr(function, ty.array<u32, 20>()));
+ auto* loop = b.Loop();
+ b.Append(loop->Initializer(), [&] {
+ // idx = 0u
+ idx = b.Var("idx", 0_u);
+ b.NextIteration(loop);
+ });
+ b.Append(loop->Body(), [&] {
+ // idx < 20u
+ auto* binary = b.LessThan<bool>(b.Load(idx), 20_u);
+ auto* ifelse = b.If(binary);
+ b.Append(ifelse->True(), [&] { b.ExitIf(ifelse); });
+ b.Append(ifelse->False(), [&] { b.ExitLoop(loop); });
+ // access_arr = arr[idx]
+ // idx: [0u, 19u]
+ auto* load_idx = b.Load(idx);
+ auto* access_arr = b.Access(ty.ptr<function, u32>(), arr, load_idx);
+ b.Load(access_arr);
+ b.Continue(loop);
+ });
+ b.Append(loop->Continuing(), [&] {
+ // idx++
+ b.Store(idx, b.Add<u32>(b.Load(idx), 1_u));
+ b.NextIteration(loop);
+ });
+ b.Return(func);
+ });
+
+ auto* expect = R"(
+%func = func():void {
+ $B1: {
+ %arr:ptr<function, array<u32, 20>, read_write> = var undef
+ loop [i: $B2, b: $B3, c: $B4] { # loop_1
+ $B2: { # initializer
+ %idx:ptr<function, u32, read_write> = var 0u
+ next_iteration # -> $B3
+ }
+ $B3: { # body
+ %4:u32 = load %idx
+ %5:bool = lt %4, 20u
+ if %5 [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ exit_if # if_1
+ }
+ $B6: { # false
+ exit_loop # loop_1
+ }
+ }
+ %6:u32 = load %idx
+ %7:ptr<function, u32, read_write> = access %arr, %6
+ %8:u32 = load %7
+ continue # -> $B4
+ }
+ $B4: { # continuing
+ %9:u32 = load %idx
+ %10:u32 = add %9, 1u
+ store %idx, %10
+ next_iteration # -> $B3
+ }
+ }
+ ret
+ }
+}
+)";
+
+ RobustnessConfig cfg;
+ cfg.clamp_function = true;
+ cfg.use_integer_range_analysis = true;
+ Run(Robustness, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_RobustnessWithIntegerRangeAnalysisTest, AccessArrayWithIndex_MaxBound_LessThan_Limit) {
+ auto* func = b.Function("func", ty.void_());
+ b.Append(func->Block(), [&] {
+ Var* idx = nullptr;
+ // arr = array<u32, 20>
+ auto* arr = b.Var("arr", ty.ptr(function, ty.array<u32, 20>()));
+ auto* loop = b.Loop();
+ b.Append(loop->Initializer(), [&] {
+ // idx = 0u
+ idx = b.Var("idx", 0_u);
+ b.NextIteration(loop);
+ });
+ b.Append(loop->Body(), [&] {
+ // idx < 19u
+ auto* binary = b.LessThan<bool>(b.Load(idx), 19_u);
+ auto* ifelse = b.If(binary);
+ b.Append(ifelse->True(), [&] { b.ExitIf(ifelse); });
+ b.Append(ifelse->False(), [&] { b.ExitLoop(loop); });
+ // access_arr = arr[idx]
+ // idx: [0u, 18u]
+ auto* load_idx = b.Load(idx);
+ auto* access_arr = b.Access(ty.ptr<function, u32>(), arr, load_idx);
+ b.Load(access_arr);
+ b.Continue(loop);
+ });
+ b.Append(loop->Continuing(), [&] {
+ // idx++
+ b.Store(idx, b.Add<u32>(b.Load(idx), 1_u));
+ b.NextIteration(loop);
+ });
+ b.Return(func);
+ });
+
+ auto* expect = R"(
+%func = func():void {
+ $B1: {
+ %arr:ptr<function, array<u32, 20>, read_write> = var undef
+ loop [i: $B2, b: $B3, c: $B4] { # loop_1
+ $B2: { # initializer
+ %idx:ptr<function, u32, read_write> = var 0u
+ next_iteration # -> $B3
+ }
+ $B3: { # body
+ %4:u32 = load %idx
+ %5:bool = lt %4, 19u
+ if %5 [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ exit_if # if_1
+ }
+ $B6: { # false
+ exit_loop # loop_1
+ }
+ }
+ %6:u32 = load %idx
+ %7:ptr<function, u32, read_write> = access %arr, %6
+ %8:u32 = load %7
+ continue # -> $B4
+ }
+ $B4: { # continuing
+ %9:u32 = load %idx
+ %10:u32 = add %9, 1u
+ store %idx, %10
+ next_iteration # -> $B3
+ }
+ }
+ ret
+ }
+}
+)";
+
+ RobustnessConfig cfg;
+ cfg.clamp_function = true;
+ cfg.use_integer_range_analysis = true;
+ Run(Robustness, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_RobustnessWithIntegerRangeAnalysisTest, AccessArrayWithIndex_MaxBound_GreaterThan_Limit) {
+ auto* func = b.Function("func", ty.void_());
+ b.Append(func->Block(), [&] {
+ Var* idx = nullptr;
+ // arr = array<u32, 20>
+ auto* arr = b.Var("arr", ty.ptr(function, ty.array<u32, 20>()));
+ auto* loop = b.Loop();
+ b.Append(loop->Initializer(), [&] {
+ // idx = 0u
+ idx = b.Var("idx", 0_u);
+ b.NextIteration(loop);
+ });
+ b.Append(loop->Body(), [&] {
+ // idx < 21u
+ auto* binary = b.LessThan<bool>(b.Load(idx), 21_u);
+ auto* ifelse = b.If(binary);
+ b.Append(ifelse->True(), [&] { b.ExitIf(ifelse); });
+ b.Append(ifelse->False(), [&] { b.ExitLoop(loop); });
+ // access_arr = arr[idx]
+ // idx: [0, 20u]
+ auto* load_idx = b.Load(idx);
+ auto* access_arr = b.Access(ty.ptr<function, u32>(), arr, load_idx);
+ b.Load(access_arr);
+ b.Continue(loop);
+ });
+ b.Append(loop->Continuing(), [&] {
+ // idx++
+ b.Store(idx, b.Add<u32>(b.Load(idx), 1_u));
+ b.NextIteration(loop);
+ });
+ b.Return(func);
+ });
+
+ auto* expect = R"(
+%func = func():void {
+ $B1: {
+ %arr:ptr<function, array<u32, 20>, read_write> = var undef
+ loop [i: $B2, b: $B3, c: $B4] { # loop_1
+ $B2: { # initializer
+ %idx:ptr<function, u32, read_write> = var 0u
+ next_iteration # -> $B3
+ }
+ $B3: { # body
+ %4:u32 = load %idx
+ %5:bool = lt %4, 21u
+ if %5 [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ exit_if # if_1
+ }
+ $B6: { # false
+ exit_loop # loop_1
+ }
+ }
+ %6:u32 = load %idx
+ %7:u32 = min %6, 19u
+ %8:ptr<function, u32, read_write> = access %arr, %7
+ %9:u32 = load %8
+ continue # -> $B4
+ }
+ $B4: { # continuing
+ %10:u32 = load %idx
+ %11:u32 = add %10, 1u
+ store %idx, %11
+ next_iteration # -> $B3
+ }
+ }
+ ret
+ }
+}
+)";
+
+ RobustnessConfig cfg;
+ cfg.clamp_function = true;
+ cfg.use_integer_range_analysis = true;
+ Run(Robustness, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_RobustnessWithIntegerRangeAnalysisTest, AccessArrayWithIndex_IndexNoRange) {
+ auto* func = b.Function("func", ty.void_());
+ auto* param = b.FunctionParam("param", ty.u32());
+ func->AppendParam(param);
+ b.Append(func->Block(), [&] {
+ // arr = array<u32, 20>
+ auto* arr = b.Var("arr", ty.ptr(function, ty.array<u32, 20>()));
+ auto* access_arr = b.Access(ty.ptr<function, u32>(), arr, param);
+ b.Load(access_arr);
+ b.Return(func);
+ });
+
+ auto* expect = R"(
+%func = func(%param:u32):void {
+ $B1: {
+ %arr:ptr<function, array<u32, 20>, read_write> = var undef
+ %4:u32 = min %param, 19u
+ %5:ptr<function, u32, read_write> = access %arr, %4
+ %6:u32 = load %5
+ ret
+ }
+}
+)";
+
+ RobustnessConfig cfg;
+ cfg.clamp_function = true;
+ cfg.use_integer_range_analysis = true;
+ Run(Robustness, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_RobustnessWithIntegerRangeAnalysisTest, AccessArrayWithExpression_MaxBound_Equal_Limit) {
+ auto* func = b.ComputeFunction("my_func", 4_u, 1_u, 1_u);
+ auto* local_invocation_id = b.FunctionParam("local_id", mod.Types().vec3<u32>());
+ local_invocation_id->SetBuiltin(tint::core::BuiltinValue::kLocalInvocationId);
+ func->SetParams({local_invocation_id});
+
+ b.Append(func->Block(), [&] {
+ Var* idx = nullptr;
+
+ // factor = 4
+ auto* factor = b.Constant(4_u);
+ // arr = array<u32, 32>
+ auto* arr = b.Var("arr", ty.ptr(function, ty.array<u32, 32>()));
+ auto* loop = b.Loop();
+ b.Append(loop->Initializer(), [&] {
+ // idx = 0
+ idx = b.Var("idx", 0_u);
+ b.NextIteration(loop);
+ });
+ b.Append(loop->Body(), [&] {
+ // idx < 8
+ auto* binary = b.LessThan<bool>(b.Load(idx), 8_u);
+ auto* ifelse = b.If(binary);
+ b.Append(ifelse->True(), [&] { b.ExitIf(ifelse); });
+ b.Append(ifelse->False(), [&] { b.ExitLoop(loop); });
+
+ // access_idx = idx * factor + local_id.x
+ // access_idx: [0, 31] (idx: [0, 7], factor = 4, local_id.x: [0, 3])
+ auto* load_idx = b.Load(idx);
+ auto* multiply = b.Multiply<u32>(load_idx, factor);
+ auto* access_local_id_x = b.Access(ty.u32(), local_invocation_id, 0_u);
+ auto* access_idx = b.Add<u32>(multiply, access_local_id_x);
+
+ // access_arr = arr[access_idx]
+ auto* access_arr = b.Access(ty.ptr<function, u32>(), arr, access_idx);
+ b.Load(access_arr);
+
+ b.Continue(loop);
+ });
+ b.Append(loop->Continuing(), [&] {
+ // idx++
+ b.Store(idx, b.Add<u32>(b.Load(idx), 1_u));
+ b.NextIteration(loop);
+ });
+
+ b.Return(func);
+ });
+
+ auto* expect = R"(
+%my_func = @compute @workgroup_size(4u, 1u, 1u) func(%local_id:vec3<u32> [@local_invocation_id]):void {
+ $B1: {
+ %arr:ptr<function, array<u32, 32>, read_write> = var undef
+ loop [i: $B2, b: $B3, c: $B4] { # loop_1
+ $B2: { # initializer
+ %idx:ptr<function, u32, read_write> = var 0u
+ next_iteration # -> $B3
+ }
+ $B3: { # body
+ %5:u32 = load %idx
+ %6:bool = lt %5, 8u
+ if %6 [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ exit_if # if_1
+ }
+ $B6: { # false
+ exit_loop # loop_1
+ }
+ }
+ %7:u32 = load %idx
+ %8:u32 = mul %7, 4u
+ %9:u32 = access %local_id, 0u
+ %10:u32 = add %8, %9
+ %11:ptr<function, u32, read_write> = access %arr, %10
+ %12:u32 = load %11
+ continue # -> $B4
+ }
+ $B4: { # continuing
+ %13:u32 = load %idx
+ %14:u32 = add %13, 1u
+ store %idx, %14
+ next_iteration # -> $B3
+ }
+ }
+ ret
+ }
+}
+)";
+
+ RobustnessConfig cfg;
+ cfg.clamp_function = true;
+ cfg.use_integer_range_analysis = true;
+ Run(Robustness, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_RobustnessWithIntegerRangeAnalysisTest, AccessArrayWithIndex_I32_NegativeMinBound) {
+ auto* func = b.Function("func", ty.void_());
+ b.Append(func->Block(), [&] {
+ Var* idx = nullptr;
+ // arr = array<u32, 20>
+ auto* arr = b.Var("arr", ty.ptr(function, ty.array<u32, 20>()));
+ auto* loop = b.Loop();
+ b.Append(loop->Initializer(), [&] {
+ // idx = -1
+ idx = b.Var("idx", -1_i);
+ b.NextIteration(loop);
+ });
+ b.Append(loop->Body(), [&] {
+ // idx < 20
+ auto* binary = b.LessThan<bool>(b.Load(idx), 20_i);
+ auto* ifelse = b.If(binary);
+ b.Append(ifelse->True(), [&] { b.ExitIf(ifelse); });
+ b.Append(ifelse->False(), [&] { b.ExitLoop(loop); });
+ // access_arr = arr[idx]
+ // idx: [-1, 19]
+ auto* load_idx = b.Load(idx);
+ auto* access_arr = b.Access(ty.ptr<function, u32>(), arr, load_idx);
+ b.Load(access_arr);
+ b.Continue(loop);
+ });
+ b.Append(loop->Continuing(), [&] {
+ // idx++
+ b.Store(idx, b.Add<i32>(b.Load(idx), 1_i));
+ b.NextIteration(loop);
+ });
+ b.Return(func);
+ });
+
+ auto* expect = R"(
+%func = func():void {
+ $B1: {
+ %arr:ptr<function, array<u32, 20>, read_write> = var undef
+ loop [i: $B2, b: $B3, c: $B4] { # loop_1
+ $B2: { # initializer
+ %idx:ptr<function, i32, read_write> = var -1i
+ next_iteration # -> $B3
+ }
+ $B3: { # body
+ %4:i32 = load %idx
+ %5:bool = lt %4, 20i
+ if %5 [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ exit_if # if_1
+ }
+ $B6: { # false
+ exit_loop # loop_1
+ }
+ }
+ %6:i32 = load %idx
+ %7:u32 = convert %6
+ %8:u32 = min %7, 19u
+ %9:ptr<function, u32, read_write> = access %arr, %8
+ %10:u32 = load %9
+ continue # -> $B4
+ }
+ $B4: { # continuing
+ %11:i32 = load %idx
+ %12:i32 = add %11, 1i
+ store %idx, %12
+ next_iteration # -> $B3
+ }
+ }
+ ret
+ }
+}
+)";
+
+ RobustnessConfig cfg;
+ cfg.clamp_function = true;
+ cfg.use_integer_range_analysis = true;
+ Run(Robustness, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_RobustnessWithIntegerRangeAnalysisTest, AccessArrayWithIndex_MaxBound_Equal_Limit_I32) {
+ auto* func = b.Function("func", ty.void_());
+ b.Append(func->Block(), [&] {
+ Var* idx = nullptr;
+ // arr = array<u32, 20>
+ auto* arr = b.Var("arr", ty.ptr(function, ty.array<u32, 20>()));
+ auto* loop = b.Loop();
+ b.Append(loop->Initializer(), [&] {
+ // idx = 0
+ idx = b.Var("idx", 0_i);
+ b.NextIteration(loop);
+ });
+ b.Append(loop->Body(), [&] {
+ // idx < 20
+ auto* binary = b.LessThan<bool>(b.Load(idx), 20_i);
+ auto* ifelse = b.If(binary);
+ b.Append(ifelse->True(), [&] { b.ExitIf(ifelse); });
+ b.Append(ifelse->False(), [&] { b.ExitLoop(loop); });
+ // access_arr = arr[idx]
+ // idx: [0, 19]
+ auto* load_idx = b.Load(idx);
+ auto* access_arr = b.Access(ty.ptr<function, u32>(), arr, load_idx);
+ b.Load(access_arr);
+ b.Continue(loop);
+ });
+ b.Append(loop->Continuing(), [&] {
+ // idx++
+ b.Store(idx, b.Add<i32>(b.Load(idx), 1_i));
+ b.NextIteration(loop);
+ });
+ b.Return(func);
+ });
+
+ auto* expect = R"(
+%func = func():void {
+ $B1: {
+ %arr:ptr<function, array<u32, 20>, read_write> = var undef
+ loop [i: $B2, b: $B3, c: $B4] { # loop_1
+ $B2: { # initializer
+ %idx:ptr<function, i32, read_write> = var 0i
+ next_iteration # -> $B3
+ }
+ $B3: { # body
+ %4:i32 = load %idx
+ %5:bool = lt %4, 20i
+ if %5 [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ exit_if # if_1
+ }
+ $B6: { # false
+ exit_loop # loop_1
+ }
+ }
+ %6:i32 = load %idx
+ %7:ptr<function, u32, read_write> = access %arr, %6
+ %8:u32 = load %7
+ continue # -> $B4
+ }
+ $B4: { # continuing
+ %9:i32 = load %idx
+ %10:i32 = add %9, 1i
+ store %idx, %10
+ next_iteration # -> $B3
+ }
+ }
+ ret
+ }
+}
+)";
+
+ RobustnessConfig cfg;
+ cfg.clamp_function = true;
+ cfg.use_integer_range_analysis = true;
+ Run(Robustness, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_RobustnessWithIntegerRangeAnalysisTest,
+ AccessArrayWithIndex_MaxBound_LessThan_Limit_I32) {
+ auto* func = b.Function("func", ty.void_());
+ b.Append(func->Block(), [&] {
+ Var* idx = nullptr;
+ // arr = array<u32, 20>
+ auto* arr = b.Var("arr", ty.ptr(function, ty.array<u32, 20>()));
+ auto* loop = b.Loop();
+ b.Append(loop->Initializer(), [&] {
+ // idx = 0
+ idx = b.Var("idx", 0_i);
+ b.NextIteration(loop);
+ });
+ b.Append(loop->Body(), [&] {
+ // idx < 19
+ auto* binary = b.LessThan<bool>(b.Load(idx), 19_i);
+ auto* ifelse = b.If(binary);
+ b.Append(ifelse->True(), [&] { b.ExitIf(ifelse); });
+ b.Append(ifelse->False(), [&] { b.ExitLoop(loop); });
+ // access_arr = arr[idx]
+ // idx: [0, 18]
+ auto* load_idx = b.Load(idx);
+ auto* access_arr = b.Access(ty.ptr<function, u32>(), arr, load_idx);
+ b.Load(access_arr);
+ b.Continue(loop);
+ });
+ b.Append(loop->Continuing(), [&] {
+ // idx++
+ b.Store(idx, b.Add<i32>(b.Load(idx), 1_i));
+ b.NextIteration(loop);
+ });
+ b.Return(func);
+ });
+
+ auto* expect = R"(
+%func = func():void {
+ $B1: {
+ %arr:ptr<function, array<u32, 20>, read_write> = var undef
+ loop [i: $B2, b: $B3, c: $B4] { # loop_1
+ $B2: { # initializer
+ %idx:ptr<function, i32, read_write> = var 0i
+ next_iteration # -> $B3
+ }
+ $B3: { # body
+ %4:i32 = load %idx
+ %5:bool = lt %4, 19i
+ if %5 [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ exit_if # if_1
+ }
+ $B6: { # false
+ exit_loop # loop_1
+ }
+ }
+ %6:i32 = load %idx
+ %7:ptr<function, u32, read_write> = access %arr, %6
+ %8:u32 = load %7
+ continue # -> $B4
+ }
+ $B4: { # continuing
+ %9:i32 = load %idx
+ %10:i32 = add %9, 1i
+ store %idx, %10
+ next_iteration # -> $B3
+ }
+ }
+ ret
+ }
+}
+)";
+
+ RobustnessConfig cfg;
+ cfg.clamp_function = true;
+ cfg.use_integer_range_analysis = true;
+ Run(Robustness, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_RobustnessWithIntegerRangeAnalysisTest,
+ AccessArrayWithIndex_MaxBound_GreaterThan_Limit_I32) {
+ auto* func = b.Function("func", ty.void_());
+ b.Append(func->Block(), [&] {
+ Var* idx = nullptr;
+ // arr = array<u32, 20>
+ auto* arr = b.Var("arr", ty.ptr(function, ty.array<u32, 20>()));
+ auto* loop = b.Loop();
+ b.Append(loop->Initializer(), [&] {
+ // idx = 0
+ idx = b.Var("idx", 0_i);
+ b.NextIteration(loop);
+ });
+ b.Append(loop->Body(), [&] {
+ // idx < 21
+ auto* binary = b.LessThan<bool>(b.Load(idx), 21_i);
+ auto* ifelse = b.If(binary);
+ b.Append(ifelse->True(), [&] { b.ExitIf(ifelse); });
+ b.Append(ifelse->False(), [&] { b.ExitLoop(loop); });
+ // access_arr = arr[idx]
+ // idx: [0, 20]
+ auto* load_idx = b.Load(idx);
+ auto* access_arr = b.Access(ty.ptr<function, u32>(), arr, load_idx);
+ b.Load(access_arr);
+ b.Continue(loop);
+ });
+ b.Append(loop->Continuing(), [&] {
+ // idx++
+ b.Store(idx, b.Add<i32>(b.Load(idx), 1_i));
+ b.NextIteration(loop);
+ });
+ b.Return(func);
+ });
+
+ auto* expect = R"(
+%func = func():void {
+ $B1: {
+ %arr:ptr<function, array<u32, 20>, read_write> = var undef
+ loop [i: $B2, b: $B3, c: $B4] { # loop_1
+ $B2: { # initializer
+ %idx:ptr<function, i32, read_write> = var 0i
+ next_iteration # -> $B3
+ }
+ $B3: { # body
+ %4:i32 = load %idx
+ %5:bool = lt %4, 21i
+ if %5 [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ exit_if # if_1
+ }
+ $B6: { # false
+ exit_loop # loop_1
+ }
+ }
+ %6:i32 = load %idx
+ %7:u32 = convert %6
+ %8:u32 = min %7, 19u
+ %9:ptr<function, u32, read_write> = access %arr, %8
+ %10:u32 = load %9
+ continue # -> $B4
+ }
+ $B4: { # continuing
+ %11:i32 = load %idx
+ %12:i32 = add %11, 1i
+ store %idx, %12
+ next_iteration # -> $B3
+ }
+ }
+ ret
+ }
+}
+)";
+
+ RobustnessConfig cfg;
+ cfg.clamp_function = true;
+ cfg.use_integer_range_analysis = true;
+ Run(Robustness, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_RobustnessWithIntegerRangeAnalysisTest, AccessArrayWithIndex_IndexNoRange_I32) {
+ auto* func = b.Function("func", ty.void_());
+ auto* param = b.FunctionParam("param", ty.i32());
+ func->AppendParam(param);
+ b.Append(func->Block(), [&] {
+ // arr = array<u32, 20>
+ auto* arr = b.Var("arr", ty.ptr(function, ty.array<u32, 20>()));
+ auto* access_arr = b.Access(ty.ptr<function, u32>(), arr, param);
+ b.Load(access_arr);
+ b.Return(func);
+ });
+
+ auto* expect = R"(
+%func = func(%param:i32):void {
+ $B1: {
+ %arr:ptr<function, array<u32, 20>, read_write> = var undef
+ %4:u32 = convert %param
+ %5:u32 = min %4, 19u
+ %6:ptr<function, u32, read_write> = access %arr, %5
+ %7:u32 = load %6
+ ret
+ }
+}
+)";
+
+ RobustnessConfig cfg;
+ cfg.clamp_function = true;
+ cfg.use_integer_range_analysis = true;
+ Run(Robustness, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_RobustnessWithIntegerRangeAnalysisTest, AccessArrayWithIndex_DynamicSizedArray) {
+ // arr = array<u32>
+ auto* arr = b.Var("arr", ty.ptr(storage, ty.array<u32>()));
+ arr->SetBindingPoint(0, 0);
+ mod.root_block->Append(arr);
+
+ auto* func = b.Function("func", ty.void_());
+ b.Append(func->Block(), [&] {
+ Var* idx = nullptr;
+ auto* loop = b.Loop();
+ b.Append(loop->Initializer(), [&] {
+ // idx = 0u
+ idx = b.Var("idx", 0_u);
+ b.NextIteration(loop);
+ });
+ b.Append(loop->Body(), [&] {
+ // idx < 20u
+ auto* binary = b.LessThan<bool>(b.Load(idx), 20_u);
+ auto* ifelse = b.If(binary);
+ b.Append(ifelse->True(), [&] { b.ExitIf(ifelse); });
+ b.Append(ifelse->False(), [&] { b.ExitLoop(loop); });
+ // access_arr = arr[idx]
+ // idx: [0u, 19u]
+ auto* load_idx = b.Load(idx);
+ auto* access_arr = b.Access(ty.ptr<storage, u32>(), arr, load_idx);
+ b.Load(access_arr);
+ b.Continue(loop);
+ });
+ b.Append(loop->Continuing(), [&] {
+ // idx++
+ b.Store(idx, b.Add<u32>(b.Load(idx), 1_u));
+ b.NextIteration(loop);
+ });
+ b.Return(func);
+ });
+
+ auto* expect = R"(
+$B1: { # root
+ %arr:ptr<storage, array<u32>, read_write> = var undef @binding_point(0, 0)
+}
+
+%func = func():void {
+ $B2: {
+ loop [i: $B3, b: $B4, c: $B5] { # loop_1
+ $B3: { # initializer
+ %idx:ptr<function, u32, read_write> = var 0u
+ next_iteration # -> $B4
+ }
+ $B4: { # body
+ %4:u32 = load %idx
+ %5:bool = lt %4, 20u
+ if %5 [t: $B6, f: $B7] { # if_1
+ $B6: { # true
+ exit_if # if_1
+ }
+ $B7: { # false
+ exit_loop # loop_1
+ }
+ }
+ %6:u32 = load %idx
+ %7:u32 = arrayLength %arr
+ %8:u32 = sub %7, 1u
+ %9:u32 = min %6, %8
+ %10:ptr<storage, u32, read_write> = access %arr, %9
+ %11:u32 = load %10
+ continue # -> $B5
+ }
+ $B5: { # continuing
+ %12:u32 = load %idx
+ %13:u32 = add %12, 1u
+ store %idx, %13
+ next_iteration # -> $B4
+ }
+ }
+ ret
+ }
+}
+)";
+
+ RobustnessConfig cfg;
+ cfg.clamp_function = true;
+ cfg.clamp_storage = true;
+ cfg.use_integer_range_analysis = true;
+ Run(Robustness, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
INSTANTIATE_TEST_SUITE_P(, IR_RobustnessTest, testing::Values(false, true));
INSTANTIATE_TEST_SUITE_P(,
