From 3bf3e9b20c6e06b57c0966e37ba990cc01d39365 Mon Sep 17 00:00:00 2001 From: Jiawei Shao Date: Wed, 18 Dec 2024 17:21:11 -0800 Subject: [PATCH] Tint: Implement range analysis on `LocalInvocationIndex` This patch adds the basic infrastructure of integer range analysis and implements the computation of the range on `LocalInvocationIndex`. Bug: chromium:348701956 Test: tint_unittests Change-Id: I93320c451aefcdc0bb60fe436cebdef172d9d046 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/219174 Reviewed-by: James Price Commit-Queue: Jiawei Shao --- src/tint/lang/core/ir/analysis/BUILD.bazel | 3 + src/tint/lang/core/ir/analysis/BUILD.cmake | 3 + src/tint/lang/core/ir/analysis/BUILD.gn | 7 +- .../ir/analysis/integer_range_analysis.cc | 93 +++++++++ .../core/ir/analysis/integer_range_analysis.h | 84 ++++++++ .../analysis/integer_range_analysis_test.cc | 188 ++++++++++++++++++ 6 files changed, 377 insertions(+), 1 deletion(-) create mode 100644 src/tint/lang/core/ir/analysis/integer_range_analysis.cc create mode 100644 src/tint/lang/core/ir/analysis/integer_range_analysis.h create mode 100644 src/tint/lang/core/ir/analysis/integer_range_analysis_test.cc diff --git a/src/tint/lang/core/ir/analysis/BUILD.bazel b/src/tint/lang/core/ir/analysis/BUILD.bazel index 135191b7cb..6e8e286a5a 100644 --- a/src/tint/lang/core/ir/analysis/BUILD.bazel +++ b/src/tint/lang/core/ir/analysis/BUILD.bazel @@ -39,9 +39,11 @@ load("@bazel_skylib//lib:selects.bzl", "selects") cc_library( name = "analysis", srcs = [ + "integer_range_analysis.cc", "loop_analysis.cc", ], hdrs = [ + "integer_range_analysis.h", "loop_analysis.h", ], deps = [ @@ -69,6 +71,7 @@ cc_library( name = "test", alwayslink = True, srcs = [ + "integer_range_analysis_test.cc", "loop_analysis_test.cc", ], deps = [ diff --git a/src/tint/lang/core/ir/analysis/BUILD.cmake b/src/tint/lang/core/ir/analysis/BUILD.cmake index f29f3f3941..d2896ad447 100644 --- a/src/tint/lang/core/ir/analysis/BUILD.cmake +++ b/src/tint/lang/core/ir/analysis/BUILD.cmake @@ -39,6 +39,8 @@ # Kind: lib ################################################################################ tint_add_target(tint_lang_core_ir_analysis lib + lang/core/ir/analysis/integer_range_analysis.cc + lang/core/ir/analysis/integer_range_analysis.h lang/core/ir/analysis/loop_analysis.cc lang/core/ir/analysis/loop_analysis.h ) @@ -70,6 +72,7 @@ tint_target_add_external_dependencies(tint_lang_core_ir_analysis lib # Kind: test ################################################################################ tint_add_target(tint_lang_core_ir_analysis_test test + lang/core/ir/analysis/integer_range_analysis_test.cc lang/core/ir/analysis/loop_analysis_test.cc ) diff --git a/src/tint/lang/core/ir/analysis/BUILD.gn b/src/tint/lang/core/ir/analysis/BUILD.gn index c1f76eab62..ba4efb07ae 100644 --- a/src/tint/lang/core/ir/analysis/BUILD.gn +++ b/src/tint/lang/core/ir/analysis/BUILD.gn @@ -45,6 +45,8 @@ if (tint_build_unittests || tint_build_benchmarks) { libtint_source_set("analysis") { sources = [ + "integer_range_analysis.cc", + "integer_range_analysis.h", "loop_analysis.cc", "loop_analysis.h", ] @@ -69,7 +71,10 @@ libtint_source_set("analysis") { } if (tint_build_unittests) { tint_unittests_source_set("unittests") { - sources = [ "loop_analysis_test.cc" ] + sources = [ + "integer_range_analysis_test.cc", + "loop_analysis_test.cc", + ] deps = [ "${dawn_root}/src/utils:utils", "${tint_src_dir}:gmock_and_gtest", diff --git a/src/tint/lang/core/ir/analysis/integer_range_analysis.cc b/src/tint/lang/core/ir/analysis/integer_range_analysis.cc new file mode 100644 index 0000000000..ed339f54fa --- /dev/null +++ b/src/tint/lang/core/ir/analysis/integer_range_analysis.cc @@ -0,0 +1,93 @@ +// 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/analysis/integer_range_analysis.h" + +#include + +#include "src/tint/lang/core/ir/traverse.h" + +namespace tint::core::ir::analysis { + +IntegerRangeInfo::IntegerRangeInfo(int64_t min_bound, int64_t max_bound) { + TINT_ASSERT(min_bound <= max_bound); + range = SignedIntegerRange{min_bound, max_bound}; +} + +IntegerRangeInfo::IntegerRangeInfo(uint64_t min_bound, uint64_t max_bound) { + TINT_ASSERT(min_bound <= max_bound); + range = UnsignedIntegerRange{min_bound, max_bound}; +} + +struct IntegerRangeAnalysisImpl { + explicit IntegerRangeAnalysisImpl(Function* func) : function_(func) {} + + const IntegerRangeInfo* GetInfo(const FunctionParam* param) { + if (!param->Type()->IsIntegerScalar()) { + return nullptr; + } + + const auto& info = + integer_function_param_range_info_map_.GetOrAdd(param, [&]() -> IntegerRangeInfo { + 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 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->Type()->IsUnsignedIntegerScalar()) { + return IntegerRangeInfo(0, std::numeric_limits::max()); + } else { + TINT_ASSERT(param->Type()->IsSignedIntegerScalar()); + return IntegerRangeInfo(std::numeric_limits::min(), + std::numeric_limits::max()); + } + }); + + return &info; + } + + private: + Function* function_; + Hashmap integer_function_param_range_info_map_; +}; + +IntegerRangeAnalysis::IntegerRangeAnalysis(Function* func) + : impl_(new IntegerRangeAnalysisImpl(func)) {} +IntegerRangeAnalysis::~IntegerRangeAnalysis() = default; + +const IntegerRangeInfo* IntegerRangeAnalysis::GetInfo(const FunctionParam* param) { + return impl_->GetInfo(param); +} +} // 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 new file mode 100644 index 0000000000..03bd633008 --- /dev/null +++ b/src/tint/lang/core/ir/analysis/integer_range_analysis.h @@ -0,0 +1,84 @@ +// 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_ANALYSIS_INTEGER_RANGE_ANALYSIS_H_ +#define SRC_TINT_LANG_CORE_IR_ANALYSIS_INTEGER_RANGE_ANALYSIS_H_ + +#include +#include + +#include "src/tint/lang/core/ir/function.h" +#include "src/tint/lang/core/ir/var.h" + +namespace tint::core::ir::analysis { + +/// The result of a integer range analysis: the upper and lower bound of a given integer variable. +/// The bound is inclusive, which means the value x being bound satisfies: +/// min_bound <= x <= max_bound. +struct IntegerRangeInfo { + IntegerRangeInfo() = default; + IntegerRangeInfo(int64_t min_bound, int64_t max_bound); + IntegerRangeInfo(uint64_t min_bound, uint64_t max_bound); + + struct SignedIntegerRange { + int64_t min_bound; + int64_t max_bound; + }; + struct UnsignedIntegerRange { + uint64_t min_bound; + uint64_t max_bound; + }; + std::variant range; +}; + +struct IntegerRangeAnalysisImpl; + +/// IntegerRangeAnalysis is a helper used to analyze integer ranges. +class IntegerRangeAnalysis { + public: + /// Constructor + /// @param func the function to cache analyses for + explicit IntegerRangeAnalysis(ir::Function* func); + ~IntegerRangeAnalysis(); + + /// Returns the integer range info of a given parameter, if it is an integer parameter. + /// Otherwise is not analyzable and returns nullptr. If it is the first time to query the info, + /// the result will also be stored into a cache for future queries. + /// @param param the variable to get information about + /// @returns the integer range info + const IntegerRangeInfo* GetInfo(const FunctionParam* param); + + private: + IntegerRangeAnalysis(const IntegerRangeAnalysis&) = delete; + IntegerRangeAnalysis(IntegerRangeAnalysis&&) = delete; + + std::unique_ptr impl_; +}; + +} // namespace tint::core::ir::analysis + +#endif // SRC_TINT_LANG_CORE_IR_ANALYSIS_INTEGER_RANGE_ANALYSIS_H_ 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 new file mode 100644 index 0000000000..c01ba5853f --- /dev/null +++ b/src/tint/lang/core/ir/analysis/integer_range_analysis_test.cc @@ -0,0 +1,188 @@ +// 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/analysis/integer_range_analysis.h" + +#include + +#include "src/tint/lang/core/ir/ir_helper_test.h" +#include "src/tint/lang/core/ir/validator.h" + +namespace tint::core::ir::analysis { +namespace { + +using namespace tint::core::fluent_types; // NOLINT +using namespace tint::core::number_suffixes; // NOLINT + +class IR_IntegerRangeAnalysisTest : public IRTestHelper {}; + +TEST_F(IR_IntegerRangeAnalysisTest, LocalInvocationIndex_u32_XYZ) { + auto* func = b.ComputeFunction("my_func", 4_u, 3_u, 2_u); + auto* localInvocationIndex = b.FunctionParam("localInvocationIndex", mod.Types().u32()); + localInvocationIndex->SetBuiltin(tint::core::BuiltinValue::kLocalInvocationIndex); + func->SetParams({localInvocationIndex}); + + b.Append(func->Block(), [&] { + auto* dst = b.Var(ty.ptr>()); + auto* access_dst = b.Access(ty.ptr(), dst, localInvocationIndex); + b.Store(access_dst, localInvocationIndex); + b.Return(func); + }); + + auto* src = R"( +%my_func = @compute @workgroup_size(4u, 3u, 2u) func(%localInvocationIndex:u32 [@local_invocation_index]):void { + $B1: { + %3:ptr, read_write> = var + %4:ptr = access %3, %localInvocationIndex + store %4, %localInvocationIndex + ret + } +} +)"; + EXPECT_EQ(src, str()); + EXPECT_EQ(Validate(mod), Success); + + IntegerRangeAnalysis analysis(func); + auto* info = analysis.GetInfo(localInvocationIndex); + ASSERT_NE(nullptr, info); + ASSERT_TRUE(std::holds_alternative(info->range)); + + const auto& range = std::get(info->range); + EXPECT_EQ(0u, range.min_bound); + EXPECT_EQ(23u, range.max_bound); +} + +TEST_F(IR_IntegerRangeAnalysisTest, LocalInvocationIndex_i32_XYZ) { + auto* func = b.ComputeFunction("my_func", 5_i, 4_i, 3_i); + auto* localInvocationIndex = b.FunctionParam("localInvocationIndex", mod.Types().u32()); + localInvocationIndex->SetBuiltin(tint::core::BuiltinValue::kLocalInvocationIndex); + func->SetParams({localInvocationIndex}); + + b.Append(func->Block(), [&] { + auto* dst = b.Var(ty.ptr>()); + auto* access_dst = b.Access(ty.ptr(), dst, localInvocationIndex); + b.Store(access_dst, localInvocationIndex); + b.Return(func); + }); + + auto* src = R"( +%my_func = @compute @workgroup_size(5i, 4i, 3i) func(%localInvocationIndex:u32 [@local_invocation_index]):void { + $B1: { + %3:ptr, read_write> = var + %4:ptr = access %3, %localInvocationIndex + store %4, %localInvocationIndex + ret + } +} +)"; + EXPECT_EQ(src, str()); + EXPECT_EQ(Validate(mod), Success); + + IntegerRangeAnalysis analysis(func); + auto* info = analysis.GetInfo(localInvocationIndex); + ASSERT_NE(nullptr, info); + ASSERT_TRUE(std::holds_alternative(info->range)); + + const auto& range = std::get(info->range); + EXPECT_EQ(0u, range.min_bound); + EXPECT_EQ(59u, range.max_bound); +} + +TEST_F(IR_IntegerRangeAnalysisTest, LocalInvocationIndex_1_Y_1) { + auto* func = b.ComputeFunction("my_func", 1_u, 8_u, 1_u); + auto* localInvocationIndex = b.FunctionParam("localInvocationIndex", mod.Types().u32()); + localInvocationIndex->SetBuiltin(tint::core::BuiltinValue::kLocalInvocationIndex); + func->SetParams({localInvocationIndex}); + + b.Append(func->Block(), [&] { + auto* dst = b.Var(ty.ptr>()); + auto* access_dst = b.Access(ty.ptr(), dst, localInvocationIndex); + b.Store(access_dst, localInvocationIndex); + b.Return(func); + }); + + auto* src = R"( +%my_func = @compute @workgroup_size(1u, 8u, 1u) func(%localInvocationIndex:u32 [@local_invocation_index]):void { + $B1: { + %3:ptr, read_write> = var + %4:ptr = access %3, %localInvocationIndex + store %4, %localInvocationIndex + ret + } +} +)"; + EXPECT_EQ(src, str()); + EXPECT_EQ(Validate(mod), Success); + + IntegerRangeAnalysis analysis(func); + auto* info = analysis.GetInfo(localInvocationIndex); + ASSERT_NE(nullptr, info); + ASSERT_TRUE(std::holds_alternative(info->range)); + + const auto& range = std::get(info->range); + EXPECT_EQ(0u, range.min_bound); + EXPECT_EQ(7u, range.max_bound); +} + +TEST_F(IR_IntegerRangeAnalysisTest, LocalInvocationIndex_1_1_Z) { + auto* func = b.ComputeFunction("my_func", 1_u, 1_u, 16_u); + auto* localInvocationIndex = b.FunctionParam("localInvocationIndex", mod.Types().u32()); + localInvocationIndex->SetBuiltin(tint::core::BuiltinValue::kLocalInvocationIndex); + func->SetParams({localInvocationIndex}); + + b.Append(func->Block(), [&] { + auto* dst = b.Var(ty.ptr>()); + auto* access_dst = b.Access(ty.ptr(), dst, localInvocationIndex); + b.Store(access_dst, localInvocationIndex); + b.Return(func); + }); + + auto* src = R"( +%my_func = @compute @workgroup_size(1u, 1u, 16u) func(%localInvocationIndex:u32 [@local_invocation_index]):void { + $B1: { + %3:ptr, read_write> = var + %4:ptr = access %3, %localInvocationIndex + store %4, %localInvocationIndex + ret + } +} +)"; + EXPECT_EQ(src, str()); + EXPECT_EQ(Validate(mod), Success); + + IntegerRangeAnalysis analysis(func); + auto* info = analysis.GetInfo(localInvocationIndex); + ASSERT_NE(nullptr, info); + ASSERT_TRUE(std::holds_alternative(info->range)); + + const auto& range = std::get(info->range); + EXPECT_EQ(0u, range.min_bound); + EXPECT_EQ(15u, range.max_bound); +} + +} // namespace +} // namespace tint::core::ir::analysis