Range Analysis: Compute range of `Access` from local invocation ID
This patch computes the range of all the `Access` instructions that
accesses the built-in function parameter local invocation ID and
returns nullptr when computing the range of any other `Access`
instructions.
Bug: 348701956
Test: tint_unittests
Change-Id: Idf0f40858b5a6b4542407792936277195616a39e
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/239484
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Jiawei Shao <jiawei.shao@intel.com>
diff --git a/src/tint/lang/core/ir/analysis/integer_range_analysis.cc b/src/tint/lang/core/ir/analysis/integer_range_analysis.cc
index 56306a3..b6401f1 100644
--- a/src/tint/lang/core/ir/analysis/integer_range_analysis.cc
+++ b/src/tint/lang/core/ir/analysis/integer_range_analysis.cc
@@ -30,6 +30,7 @@
#include <limits>
#include "src/tint/lang/core/constant/scalar.h"
+#include "src/tint/lang/core/ir/access.h"
#include "src/tint/lang/core/ir/binary.h"
#include "src/tint/lang/core/ir/exit_if.h"
#include "src/tint/lang/core/ir/exit_loop.h"
@@ -105,41 +106,50 @@
return nullptr;
}
+ // Currently we only support the query of ranges of `local_invocation_id` or
+ // `local_invocation_index`.
+ if (!param->Builtin()) {
+ return nullptr;
+ }
+
+ switch (*param->Builtin()) {
+ case BuiltinValue::kLocalInvocationIndex:
+ case BuiltinValue::kLocalInvocationId:
+ break;
+ default:
+ return nullptr;
+ }
+
const auto& info = integer_function_param_range_info_map_.GetOrAdd(
param, [&]() -> Vector<IntegerRangeInfo, 3> {
- if (param->Builtin() == core::BuiltinValue::kLocalInvocationIndex) {
- // We shouldn't be trying to use range analysis on a module that has
- // non-constant workgroup sizes, since we will always have replaced pipeline
- // overrides with constant values early in the pipeline.
- TINT_ASSERT(function_->WorkgroupSizeAsConst().has_value());
- std::array<uint32_t, 3> workgroup_size =
- function_->WorkgroupSizeAsConst().value();
- uint64_t max_bound =
- workgroup_size[0] * workgroup_size[1] * workgroup_size[2] - 1u;
- constexpr uint64_t kMinBound = 0;
+ switch (*param->Builtin()) {
+ case BuiltinValue::kLocalInvocationIndex: {
+ // We shouldn't be trying to use range analysis on a module that has
+ // non-constant workgroup sizes, since we will always have replaced pipeline
+ // overrides with constant values early in the pipeline.
+ TINT_ASSERT(function_->WorkgroupSizeAsConst().has_value());
+ std::array<uint32_t, 3> workgroup_size =
+ function_->WorkgroupSizeAsConst().value();
+ uint64_t max_bound =
+ workgroup_size[0] * workgroup_size[1] * workgroup_size[2] - 1u;
+ constexpr uint64_t kMinBound = 0;
- return {IntegerRangeInfo(kMinBound, max_bound)};
- }
-
- if (param->Builtin() == core::BuiltinValue::kLocalInvocationId) {
- TINT_ASSERT(function_->WorkgroupSizeAsConst().has_value());
- std::array<uint32_t, 3> workgroup_size =
- function_->WorkgroupSizeAsConst().value();
-
- constexpr uint64_t kMinBound = 0;
- Vector<IntegerRangeInfo, 3> integerRanges;
- for (uint32_t size_x_y_z : workgroup_size) {
- integerRanges.Push({kMinBound, size_x_y_z - 1u});
+ return {IntegerRangeInfo(kMinBound, max_bound)};
}
- return integerRanges;
- }
+ case BuiltinValue::kLocalInvocationId: {
+ TINT_ASSERT(function_->WorkgroupSizeAsConst().has_value());
+ std::array<uint32_t, 3> workgroup_size =
+ function_->WorkgroupSizeAsConst().value();
- if (param->Type()->IsUnsignedIntegerScalar()) {
- return {IntegerRangeInfo(0, std::numeric_limits<uint64_t>::max())};
- } else {
- TINT_ASSERT(param->Type()->IsSignedIntegerScalar());
- return {IntegerRangeInfo(std::numeric_limits<int64_t>::min(),
- std::numeric_limits<int64_t>::max())};
+ constexpr uint64_t kMinBound = 0;
+ Vector<IntegerRangeInfo, 3> integerRanges;
+ for (uint32_t size_x_y_z : workgroup_size) {
+ integerRanges.Push({kMinBound, size_x_y_z - 1u});
+ }
+ return integerRanges;
+ }
+ default:
+ TINT_UNREACHABLE();
}
});
@@ -163,6 +173,26 @@
return GetInfo(load_from_var);
}
+ const IntegerRangeInfo* GetInfo(const Access* access) {
+ const Value* obj = access->Object();
+
+ // Currently we only support the access to `local_invocation_id` or `local_invocation_index`
+ // as a function parameter.
+ const FunctionParam* function_param = obj->As<FunctionParam>();
+ if (!function_param) {
+ return nullptr;
+ }
+ if (access->Indices().Length() > 1) {
+ return nullptr;
+ }
+ if (!access->Indices()[0]->As<Constant>()) {
+ return nullptr;
+ }
+ uint32_t index =
+ static_cast<uint32_t>(GetValueFromConstant(access->Indices()[0]->As<Constant>()));
+ return GetInfo(function_param, index);
+ }
+
/// Analyze a loop to compute the range of the loop control variable if possible.
void AnalyzeLoop(const Loop* loop) {
const Var* index = GetLoopControlVariableFromConstantInitializer(loop);
@@ -682,4 +712,8 @@
return impl_->GetInfo(load);
}
+const IntegerRangeInfo* IntegerRangeAnalysis::GetInfo(const Access* access) {
+ return impl_->GetInfo(access);
+}
+
} // namespace tint::core::ir::analysis
diff --git a/src/tint/lang/core/ir/analysis/integer_range_analysis.h b/src/tint/lang/core/ir/analysis/integer_range_analysis.h
index b6805a4..9c14509 100644
--- a/src/tint/lang/core/ir/analysis/integer_range_analysis.h
+++ b/src/tint/lang/core/ir/analysis/integer_range_analysis.h
@@ -33,6 +33,7 @@
#include <variant>
namespace tint::core::ir {
+class Access;
class Binary;
class Function;
class FunctionParam;
@@ -93,6 +94,10 @@
/// it has a meaningful range. Returns nullptr otherwise.
const IntegerRangeInfo* GetInfo(const Load* load_var);
+ /// Returns the integer range info of a given `Access` variable if it is an integer variable and
+ /// it has a meaningful range. Returns nullptr otherwise.
+ const IntegerRangeInfo* GetInfo(const Access* access);
+
/// Note: This function is only for tests.
/// Returns the pointer of the loop control variable in the given loop when its initializer
/// meets the below requirements.
diff --git a/src/tint/lang/core/ir/analysis/integer_range_analysis_test.cc b/src/tint/lang/core/ir/analysis/integer_range_analysis_test.cc
index 7b8e872..814193f 100644
--- a/src/tint/lang/core/ir/analysis/integer_range_analysis_test.cc
+++ b/src/tint/lang/core/ir/analysis/integer_range_analysis_test.cc
@@ -8616,5 +8616,184 @@
EXPECT_EQ(nullptr, analysis.GetInfo(load_a));
}
+TEST_F(IR_IntegerRangeAnalysisTest, AccessToLocalInvocationID) {
+ auto* func = b.ComputeFunction("my_func", 4_u, 3_u, 2_u);
+ auto* local_invocation_id = b.FunctionParam("localInvocationId", mod.Types().vec3<u32>());
+ local_invocation_id->SetBuiltin(tint::core::BuiltinValue::kLocalInvocationId);
+ func->SetParams({local_invocation_id});
+
+ Access* access_x = nullptr;
+ Access* access_y = nullptr;
+ Access* access_z = nullptr;
+ b.Append(func->Block(), [&] {
+ access_x = b.Access(ty.u32(), local_invocation_id, 0_u);
+ access_y = b.Access(ty.u32(), local_invocation_id, 1_u);
+ access_z = b.Access(ty.u32(), local_invocation_id, 2_u);
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%my_func = @compute @workgroup_size(4u, 3u, 2u) func(%localInvocationId:vec3<u32> [@local_invocation_id]):void {
+ $B1: {
+ %3:u32 = access %localInvocationId, 0u
+ %4:u32 = access %localInvocationId, 1u
+ %5:u32 = access %localInvocationId, 2u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ EXPECT_EQ(Validate(mod), Success);
+
+ IntegerRangeAnalysis analysis(func);
+
+ auto* access_x_info = analysis.GetInfo(access_x);
+ ASSERT_NE(nullptr, access_x_info);
+ ASSERT_TRUE(
+ std::holds_alternative<IntegerRangeInfo::UnsignedIntegerRange>(access_x_info->range));
+ const auto& range_access_x =
+ std::get<IntegerRangeInfo::UnsignedIntegerRange>(access_x_info->range);
+ EXPECT_EQ(0u, range_access_x.min_bound);
+ EXPECT_EQ(3u, range_access_x.max_bound);
+
+ auto* access_y_info = analysis.GetInfo(access_y);
+ ASSERT_NE(nullptr, access_x_info);
+ ASSERT_TRUE(
+ std::holds_alternative<IntegerRangeInfo::UnsignedIntegerRange>(access_y_info->range));
+ const auto& range_access_y =
+ std::get<IntegerRangeInfo::UnsignedIntegerRange>(access_y_info->range);
+ EXPECT_EQ(0u, range_access_y.min_bound);
+ EXPECT_EQ(2u, range_access_y.max_bound);
+
+ auto* access_z_info = analysis.GetInfo(access_z);
+ ASSERT_NE(nullptr, access_x_info);
+ ASSERT_TRUE(
+ std::holds_alternative<IntegerRangeInfo::UnsignedIntegerRange>(access_z_info->range));
+ const auto& range_access_z =
+ std::get<IntegerRangeInfo::UnsignedIntegerRange>(access_z_info->range);
+ EXPECT_EQ(0u, range_access_z.min_bound);
+ EXPECT_EQ(1u, range_access_z.max_bound);
+}
+
+TEST_F(IR_IntegerRangeAnalysisTest, NotAccessToFunctionParam) {
+ auto* func = b.Function("func", ty.void_());
+ Access* access = nullptr;
+ b.Append(func->Block(), [&] {
+ auto* dst = b.Var(ty.ptr<function, array<u32, 24u>>());
+ access = b.Access(ty.ptr<function, u32>(), dst, 0_u);
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%func = func():void {
+ $B1: {
+ %2:ptr<function, array<u32, 24>, read_write> = var undef
+ %3:ptr<function, u32, read_write> = access %2, 0u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ EXPECT_EQ(Validate(mod), Success);
+
+ IntegerRangeAnalysis analysis(func);
+ auto* info = analysis.GetInfo(access);
+ ASSERT_EQ(nullptr, info);
+}
+
+TEST_F(IR_IntegerRangeAnalysisTest, AccessToFunctionParamNoRange) {
+ auto* func = b.ComputeFunction("my_func", 4_u, 3_u, 2_u);
+ auto* global_invocation_id = b.FunctionParam("globalId", mod.Types().vec3<u32>());
+ global_invocation_id->SetBuiltin(tint::core::BuiltinValue::kGlobalInvocationId);
+ func->SetParams({global_invocation_id});
+
+ Access* access_x = nullptr;
+ Access* access_y = nullptr;
+ Access* access_z = nullptr;
+ b.Append(func->Block(), [&] {
+ access_x = b.Access(ty.u32(), global_invocation_id, 0_u);
+ access_y = b.Access(ty.u32(), global_invocation_id, 1_u);
+ access_z = b.Access(ty.u32(), global_invocation_id, 2_u);
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%my_func = @compute @workgroup_size(4u, 3u, 2u) func(%globalId:vec3<u32> [@global_invocation_id]):void {
+ $B1: {
+ %3:u32 = access %globalId, 0u
+ %4:u32 = access %globalId, 1u
+ %5:u32 = access %globalId, 2u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ EXPECT_EQ(Validate(mod), Success);
+
+ IntegerRangeAnalysis analysis(func);
+ ASSERT_EQ(nullptr, analysis.GetInfo(access_x));
+ ASSERT_EQ(nullptr, analysis.GetInfo(access_y));
+ ASSERT_EQ(nullptr, analysis.GetInfo(access_z));
+}
+
+TEST_F(IR_IntegerRangeAnalysisTest, AccessToNonIntegerFunctionParam) {
+ auto* func = b.Function("func", ty.void_());
+ Access* access = nullptr;
+ auto* param = b.FunctionParam("param", ty.vec4<f32>());
+ func->SetParams({param});
+ b.Append(func->Block(), [&] {
+ access = b.Access(ty.f32(), param, 0_u);
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%func = func(%param:vec4<f32>):void {
+ $B1: {
+ %3:f32 = access %param, 0u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ EXPECT_EQ(Validate(mod), Success);
+
+ IntegerRangeAnalysis analysis(func);
+ auto* info = analysis.GetInfo(access);
+ ASSERT_EQ(nullptr, info);
+}
+
+TEST_F(IR_IntegerRangeAnalysisTest, NonConstantAccessIndex) {
+ auto* func = b.ComputeFunction("my_func", 4_u, 3_u, 2_u);
+ auto* local_invocation_id = b.FunctionParam("localInvocationId", mod.Types().vec3<u32>());
+ local_invocation_id->SetBuiltin(tint::core::BuiltinValue::kLocalInvocationId);
+ func->SetParams({local_invocation_id});
+
+ Access* access_x = nullptr;
+ b.Append(func->Block(), [&] {
+ auto* var = b.Var(ty.ptr<function, u32>());
+ auto* index = b.Load(var);
+ access_x = b.Access(ty.u32(), local_invocation_id, index);
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%my_func = @compute @workgroup_size(4u, 3u, 2u) func(%localInvocationId:vec3<u32> [@local_invocation_id]):void {
+ $B1: {
+ %3:ptr<function, u32, read_write> = var undef
+ %4:u32 = load %3
+ %5:u32 = access %localInvocationId, %4
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ EXPECT_EQ(Validate(mod), Success);
+
+ IntegerRangeAnalysis analysis(func);
+
+ auto* access_x_info = analysis.GetInfo(access_x);
+ ASSERT_EQ(nullptr, access_x_info);
+}
+
} // namespace
} // namespace tint::core::ir::analysis
