This is an automated email from the ASF dual-hosted git repository.
spectrometerHBH pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm.git
The following commit(s) were added to refs/heads/main by this push:
new 88d2de8f0b [TEST] Simplify and organize native C++ tests (#20676)
88d2de8f0b is described below
commit 88d2de8f0b55575ba63e97a243a84d7d8b7f4362
Author: Tianqi Chen <[email protected]>
AuthorDate: Mon Oct 12 14:13:32 2026 -0400
[TEST] Simplify and organize native C++ tests (#20676)
Remove redundant native tests and organize the retained C++ coverage by
subsystem. Consolidate existing target and schema fixtures, retain
focused native extension and lifetime checks, and make the logging
exception helper fail when no exception is thrown.
---
tests/cpp/.gitignore | 3 -
tests/cpp/README.md | 4 +
tests/cpp/cuda_launch_test.cc | 199 ------
tests/cpp/expr_test.cc | 201 ------
tests/cpp/{ => ir}/expr_functor_test.cc | 0
.../op_signature_test.cc} | 37 -
tests/cpp/ir_functor_test.cc | 758 ---------------------
tests/cpp/ndarray_test.cc | 73 --
tests/cpp/nested_msg_test.cc | 343 ----------
tests/cpp/object_protocol_test.cc | 85 ---
tests/cpp/random_engine_test.cc | 65 --
tests/cpp/runtime/logging_test.cc | 19 +-
tests/cpp/runtime/thread_storage_scope_test.cc | 42 --
.../cpp/{support_test.cc => support/arena_test.cc} | 24 +-
tests/cpp/support/ring_buffer_test.cc | 68 --
tests/cpp/{ => sym}/pattern_match_test.cc | 2 +-
tests/cpp/sym_simplify_test.cc | 156 -----
tests/cpp/{ => target}/config_schema_test.cc | 34 +-
tests/cpp/target/target_test.cc | 141 ++++
tests/cpp/target_test.cc | 400 -----------
tests/cpp/te_compute_test.cc | 59 --
tests/cpp/threading_backend_test.cc | 196 ------
tests/cpp/tir_analysis_side_effect.cc | 38 --
tests/cpp/tir_scalable_datatype.cc | 213 ------
tests/cpp/{ => tirx}/stmt_functor_test.cc | 0
tests/cpp/topi_ewise_test.cc | 33 -
26 files changed, 162 insertions(+), 3031 deletions(-)
diff --git a/tests/cpp/.gitignore b/tests/cpp/.gitignore
deleted file mode 100644
index de6dea3ebb..0000000000
--- a/tests/cpp/.gitignore
+++ /dev/null
@@ -1,3 +0,0 @@
-unittest
-*.d
-*_test
diff --git a/tests/cpp/README.md b/tests/cpp/README.md
index 884054809e..1956893cc8 100644
--- a/tests/cpp/README.md
+++ b/tests/cpp/README.md
@@ -21,3 +21,7 @@ This folder contains some unit tests for C++ utilities in the
codebase.
In principle we aim to do most compiler related tests in the python to
bring more development velocity, and only use this folder for low-level
unit-tests.
All tests should finish fast and not dependent on a presence of an accelerator
device.
+
+Tests are grouped by subsystem in `ir/`, `runtime/`, `support/`, `sym/`,
+`target/`, and `tirx/`. CMake discovers `.cc` files in these folders
recursively.
+Configure and build out of source in `build/`, then run `build/cpptest`.
diff --git a/tests/cpp/cuda_launch_test.cc b/tests/cpp/cuda_launch_test.cc
deleted file mode 100644
index 53168357e5..0000000000
--- a/tests/cpp/cuda_launch_test.cc
+++ /dev/null
@@ -1,199 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-
-#include <algorithm>
-#include <cstring>
-#include <string>
-
-#include "../../src/backend/cuda/launch/launch_plan.h"
-#include "../../src/runtime/metadata.h"
-
-namespace {
-using namespace tvm_cuda_launch;
-
-TEST(CudaLaunch, LegacyTagsAreDecodedOnlyAtTheBoundary) {
- PackedLaunchPlan plan({},
- {"blockIdx.x", "threadIdx.x", "clusterCtaIdx.x",
"preferredClusterCtaIdx.x",
- "tirx.use_programtic_dependent_launch",
"tirx.use_dyn_shared_memory"},
- {});
- EXPECT_EQ(plan.num_values, 5);
- EXPECT_EQ(plan.operands[4].field, Field::kProgrammaticStreamSerialization);
- EXPECT_EQ(plan.operands[4].index, -1);
- LaunchValues legacy;
- legacy.SetInt(Field::kGridX, 0);
- legacy.SetInt(Field::kClusterX, 1);
- legacy.SetInt(Field::kPreferredClusterX, 1);
- plan.NormalizeLegacyPresence(&legacy);
- EXPECT_EQ(legacy.Int(Field::kGridX), 1);
- EXPECT_TRUE(legacy.Has(Option::kCluster));
- EXPECT_FALSE(legacy.Has(Option::kPreferredCluster));
- EXPECT_THROW(PackedLaunchPlan({}, {"tirx.use_dyn_shared_memory",
"blockIdx.x"}, {}),
- std::invalid_argument);
- EXPECT_THROW(PackedLaunchPlan({"grid.x"}, {"blockIdx.x"}, {}),
std::invalid_argument);
- EXPECT_THROW(PackedLaunchPlan({"grid.x", "grid.x"}, {}, {}),
std::invalid_argument);
-}
-
-class MetadataStream : public tvm::support::Stream {
- public:
- using Stream::Read;
- using Stream::Write;
- size_t Read(void* ptr, size_t size) override {
- size = std::min(size, data.size() - offset);
- std::memcpy(ptr, data.data() + offset, size);
- offset += size;
- return size;
- }
- size_t Write(const void* ptr, size_t size) override {
- data.append(static_cast<const char*>(ptr), size);
- return size;
- }
- std::string data;
- size_t offset = 0;
-};
-
-TEST(CudaLaunch, NativeOperandsRetainTheirTypes) {
- PackedLaunchPlan plan({"grid.x", "grid.y", "grid.z", "block.x", "block.y",
"block.z", "stream",
- "access_policy_window.hit_ratio"},
- {}, {});
- void* stream = reinterpret_cast<void*>(0x20000);
- auto decode = tvm::ffi::Function::FromPacked([&](tvm::ffi::PackedArgs args,
tvm::ffi::Any*) {
- LaunchValues values = plan.Extract(args, 1);
- EXPECT_EQ(values.Int(Field::kGridX), 8);
- EXPECT_EQ(values.Int(Field::kBlockX), 32);
- EXPECT_EQ(values.Get(Field::kStream).handle, stream);
- EXPECT_DOUBLE_EQ(values.Get(Field::kAccessPolicyWindowHitRatio).real,
0.25);
- EXPECT_FALSE(values.Has(Option::kCluster));
- });
- decode(999, 8, 1, 1, 32, 1, 1, stream, 0.25);
- EXPECT_THROW(decode(999, 8), std::invalid_argument);
-}
-
-TEST(CudaLaunch, MetadataReadsLegacyAndRoundTripsNative) {
- using namespace tvm;
- runtime::FunctionInfo old("kernel", {{kDLInt, 32, 1}}, {"blockIdx.x",
"threadIdx.x"}, {});
- MetadataStream bytes;
- bytes.Write(old->name);
- bytes.Write(old->arg_types);
- bytes.Write(old->launch_param_tags);
- bytes.Write(old->arg_extra_tags);
- runtime::FunctionInfo decoded(ffi::UnsafeInit{});
- ASSERT_TRUE(bytes.Read(&decoded));
- EXPECT_EQ(decoded->name, "kernel");
- EXPECT_EQ(decoded->launch_param_tags.size(), 2);
- EXPECT_TRUE(decoded->cuda_launch_fields.empty());
-
- runtime::FunctionInfo native("native", {{kDLInt, 32, 1}}, {}, {},
- {"grid.x", "grid.y", "grid.z", "block.x",
"block.y", "block.z"},
- {{"min_blocks_per_sm", 2}});
- MetadataStream native_bytes;
- native_bytes.Write(native);
- ASSERT_TRUE(native_bytes.Read(&decoded));
- EXPECT_EQ(decoded->cuda_launch_fields.size(), 6);
- EXPECT_EQ(decoded->cuda_kernel_attrs.at("min_blocks_per_sm"), 2);
- auto json_decoded = ffi::make_object<runtime::FunctionInfoObj>();
- json_decoded->LoadFromJSON(native->SaveToJSON().cast<ffi::json::Object>());
- EXPECT_EQ(json_decoded->cuda_launch_fields.size(), 6);
- EXPECT_EQ(json_decoded->cuda_kernel_attrs.at("min_blocks_per_sm"), 2);
-}
-} // namespace
-
-#if __has_include(<cuda.h>) && __has_include(<cuda_runtime.h>)
-#include <cuda.h>
-#include <cuda_runtime.h>
-#if CUDA_VERSION >= 12000 && CUDART_VERSION >= 12000
-#include "../../src/backend/cuda/launch/launch_config.h"
-
-namespace {
-TEST(CudaLaunch, DriverAndRuntimeEncodeTheSameLaunch) {
- LaunchValues values;
- values.SetInt(Field::kGridX, 8);
- values.SetInt(Field::kBlockX, 128);
- values.SetInt(Field::kClusterX, 2);
- values.SetInt(Field::kCooperative, 1);
- values.SetInt(Field::kProgrammaticStreamSerialization, 1);
- values.SetInt(Field::kPriority, -2);
- values.SetInt(Field::kDynamicSmemBytes, 65536);
- values.SetFloat(Field::kAccessPolicyWindowHitRatio, 0.5);
- values.SetInt(Field::kAccessPolicyWindowNumBytes, 4096);
- values.SetHandle(Field::kAccessPolicyWindowBasePtr,
reinterpret_cast<void*>(0x10000));
- values.SetHandle(Field::kStream, reinterpret_cast<void*>(0x20000));
- values = Normalize(values, {});
- CUlaunchConfig driver{};
- cudaLaunchConfig_t runtime{};
- std::array<CUlaunchAttribute, kMaxAttributes> driver_attrs{};
- std::array<cudaLaunchAttribute, kMaxAttributes> runtime_attrs{};
- EncodeDriver(values, {}, nullptr, &driver, &driver_attrs);
- EncodeRuntime(values, {}, nullptr, &runtime, &runtime_attrs);
- EXPECT_EQ(driver.gridDimX, runtime.gridDim.x);
- EXPECT_EQ(driver.blockDimX, runtime.blockDim.x);
- EXPECT_EQ(driver.sharedMemBytes, runtime.dynamicSmemBytes);
- EXPECT_EQ(driver.hStream, runtime.stream);
- ASSERT_EQ(driver.numAttrs, runtime.numAttrs);
- for (unsigned i = 0; i < driver.numAttrs; ++i) {
- EXPECT_EQ(static_cast<int>(driver_attrs[i].id),
static_cast<int>(runtime_attrs[i].id));
- EXPECT_EQ(
- std::memcmp(&driver_attrs[i].value, &runtime_attrs[i].val,
sizeof(driver_attrs[i].value)),
- 0);
- }
- values.SetInt(Field::kGridX, 7);
- EXPECT_THROW(Normalize(values, {}), std::invalid_argument);
- values.SetInt(Field::kGridX, 0);
- EXPECT_THROW(Normalize(values, {}), std::invalid_argument);
-}
-
-struct FakeResourceAPI {
- using Kernel = int;
- static inline int initializations = 0;
- static void Initialize(Kernel, int, ResourceState* state) {
- ++initializations;
- state->portable_bytes = 49152;
- state->optin_bytes = 229376;
- }
- static void AllowCluster(Kernel) {}
- static void SetDynamicSmem(Kernel, size_t) {}
-};
-
-TEST(CudaLaunch, ResourceCachingHandlesIncreasingSizesAndExplicitModes) {
- ResourceState state;
- FakeResourceAPI::initializations = 0;
- for (int64_t bytes : {1024, 65536, 131072, 1024}) {
- LaunchValues values;
- values.SetInt(Field::kDynamicSmemBytes, bytes);
- PrepareResources<FakeResourceAPI>(0, 0, &values, &state);
-#if CUDA_VERSION >= 13020 && CUDART_VERSION >= 13020
- if (bytes > 49152) {
- EXPECT_EQ(values.Int(Field::kSharedMemoryMode), 2);
- }
-#else
- if (bytes > 49152) {
- EXPECT_GE(state.configured_bytes, bytes);
- }
-#endif
- }
- EXPECT_EQ(FakeResourceAPI::initializations, 1);
- LaunchValues explicit_mode;
- explicit_mode.SetInt(Field::kSharedMemoryMode, 1);
- PrepareResources<FakeResourceAPI>(0, 0, &explicit_mode, &state);
- EXPECT_EQ(explicit_mode.Int(Field::kSharedMemoryMode), 1);
-}
-} // namespace
-#endif
-#endif
diff --git a/tests/cpp/expr_test.cc b/tests/cpp/expr_test.cc
deleted file mode 100644
index c373d0905b..0000000000
--- a/tests/cpp/expr_test.cc
+++ /dev/null
@@ -1,201 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-#include <tvm/ffi/cast.h>
-#include <tvm/ffi/extra/structural_equal.h>
-#include <tvm/ir/location.h>
-#include <tvm/ir/prim/expr.h>
-#include <tvm/relax/expr.h>
-#include <tvm/runtime/logging.h>
-#include <tvm/s_tir/stmt.h>
-#include <tvm/te/operation.h>
-#include <tvm/tirx/function.h>
-
-#include <type_traits>
-
-namespace {
-using namespace tvm;
-template <typename Ref>
-void CheckRequiredIRReference() {
- static_assert(!Ref::_type_is_nullable);
- static_assert(!std::is_default_constructible_v<Ref>);
- ffi::Any none;
- EXPECT_FALSE(none.as<Ref>().has_value());
- EXPECT_FALSE(none.try_cast<Ref>().has_value());
- EXPECT_THROW(none.as_or_throw<Ref>(), ffi::Error);
- EXPECT_THROW(none.cast<Ref>(), ffi::Error);
- EXPECT_FALSE(none.cast<ffi::Optional<Ref>>().has_value());
- ffi::Any missing_element = ffi::Array<ffi::Any>({none});
- EXPECT_THROW(missing_element.cast<ffi::Array<Ref>>(), ffi::Error);
- ffi::TypedFunction<Ref()> missing_result =
- ffi::Function::FromTyped([]() -> ffi::Any { return {}; });
- EXPECT_THROW(missing_result(), ffi::Error);
- ffi::TypedFunction<Ref(Ref)> required = [](Ref value) { return value; };
- EXPECT_THROW(required.packed()(none), ffi::Error);
- ffi::TypedFunction<ffi::Optional<Ref>(ffi::Optional<Ref>)> optional =
- [](ffi::Optional<Ref> value) { return value; };
- EXPECT_EQ(optional.packed()(none).type_index(), ffi::TypeIndex::kTVMFFINone);
-}
-} // namespace
-
-TEST(Expr, RequiredIRReferences) {
- using namespace tvm;
- CheckRequiredIRReference<Expr>();
- CheckRequiredIRReference<PrimExpr>();
- CheckRequiredIRReference<TypedExpr<PrimType>>();
- CheckRequiredIRReference<IntExpr>();
- CheckRequiredIRReference<PrimVar>();
- CheckRequiredIRReference<tirx::TensorVar>();
- CheckRequiredIRReference<Stmt>();
- CheckRequiredIRReference<Evaluate>();
- CheckRequiredIRReference<s_tir::SBlock>();
- CheckRequiredIRReference<tirx::Function>();
- CheckRequiredIRReference<relax::Function>();
- CheckRequiredIRReference<Type>();
- CheckRequiredIRReference<PrimType>();
- EXPECT_TRUE(Type::Missing().as<MissingType>().has_value());
-
EXPECT_TRUE(ffi::Any(Type::Missing()).cast<Type>().as<MissingType>().has_value());
-
EXPECT_THROW(ffi::Array<ffi::Any>({ffi::Any()}).as_or_throw<ffi::Array<Expr>>(),
ffi::Error);
-}
-
-TEST(Expr, NonNullablePrimitiveFallbacks) {
- using namespace tvm;
- EXPECT_EQ(ffi::Any(false).cast<PrimExpr>().as<IntImmNode>()->value, 0);
- EXPECT_EQ(ffi::Any(0).cast<PrimExpr>().as<IntImmNode>()->value, 0);
- EXPECT_EQ(ffi::Any(1.25).cast<PrimExpr>().as<FloatImmNode>()->value, 1.25);
- EXPECT_TRUE(ffi::Any(0).cast<Expr>().defined());
- Var primitive("i", PrimType::Int(32));
- EXPECT_TRUE(ffi::Any(primitive).cast<PrimExpr>().same_as(primitive));
-
EXPECT_TRUE(ffi::Any(primitive).cast<TypedExpr<PrimType>>().same_as(primitive));
- Var wrong_type("x", AnyType());
- EXPECT_FALSE(ffi::Any(wrong_type).try_cast<PrimExpr>().has_value());
-
EXPECT_FALSE(ffi::Any(wrong_type).try_cast<TypedExpr<PrimType>>().has_value());
- auto tensor = te::placeholder({1}, PrimType::Float(32), "input");
- EXPECT_TRUE(ffi::Any(tensor(0)).cast<PrimExpr>().defined());
-}
-
-TEST(Expr, Basic) {
- using namespace tvm;
- using namespace tvm::tirx;
- PrimVar x("x");
- auto z = max(x + 1 + 2, 100);
- ffi::ObjectRef tmp = z;
- PrimExpr zz = tmp.as_or_throw<PrimExpr>();
- std::ostringstream os;
- os << z;
- TVM_FFI_ICHECK(zz.same_as(z));
- TVM_FFI_ICHECK(os.str() == "T.max(x + 1 + 2, 100)");
-}
-
-TEST(Expr, VarTypeAnnotation) {
- using namespace tvm;
- using namespace tvm::tirx;
- PrimVar x("x", PrimType::Float(32));
- PrimVar y("y", PrimType::Float(32));
- tvm::ffi::StructuralEqual checker;
- TVM_FFI_ICHECK(checker(x.ty(), y.ty()));
- TVM_FFI_ICHECK(checker(x->ty, y->ty));
-}
-
-TEST(Expr, VarCopyHelpers) {
- using namespace tvm;
- using namespace tvm::tirx;
-
- SourceLoc loc(SourceName::Get("test.cc"), 1, 1, 1, 10);
- Type pointer_type = PtrType(PrimType::Float(32), "global");
- Var var("x", pointer_type, loc);
-
- Var renamed = var.CopyWithName("y");
- EXPECT_FALSE(renamed.same_as(var));
- EXPECT_EQ(renamed->name, "y");
- EXPECT_TRUE(renamed->ty.same_as(pointer_type));
- EXPECT_TRUE(renamed->loc.same_as(loc));
-
- PrimType dtype = PrimType::Int(64);
- Var retyped = var.CopyWithDType(dtype);
- EXPECT_FALSE(retyped.same_as(var));
- EXPECT_EQ(retyped->name, "x");
- EXPECT_TRUE(retyped->ty.same_as(dtype));
- EXPECT_TRUE(retyped->loc.same_as(loc));
-
- PrimVar prim_var("i", PrimType::Int(32), loc);
-
static_assert(std::is_same_v<decltype(prim_var.CopyWithDType(PrimType::Float(32))),
PrimVar>);
- PrimType prim_dtype = PrimType::Float(32);
- PrimVar retyped_prim_var = prim_var.CopyWithDType(prim_dtype);
- EXPECT_FALSE(retyped_prim_var.same_as(prim_var));
- EXPECT_EQ(retyped_prim_var->name, "i");
- EXPECT_TRUE(retyped_prim_var.ty().same_as(prim_dtype));
- EXPECT_TRUE(retyped_prim_var->loc.same_as(loc));
-}
-
-TEST(Expr, PrimTypeBoolLanes) {
- using namespace tvm;
- PrimType boolx4 = PrimType::Bool(4);
- TVM_FFI_ICHECK(boolx4.IsFixedLengthVector());
- TVM_FFI_ICHECK(boolx4.MatchesCode(DLDataTypeCode::kDLBool));
- TVM_FFI_ICHECK_EQ(boolx4.lanes(), 4);
- TVM_FFI_ICHECK(boolx4.MatchesElementType(DLDataTypeCode::kDLBool, 8));
-}
-
-TEST(Expr, IntegerView) {
- using namespace tvm;
- for (PrimType ty : {PrimType::Int(8), PrimType::UInt(64),
PrimType::Int(128)}) {
- ffi::Any value = Var("i", ty);
- EXPECT_TRUE(value.cast<IntExpr>().same_as(value.cast<PrimExpr>()));
- EXPECT_TRUE(value.cast<TypedExpr<PrimType>>().ty() == ty);
- }
- for (PrimType ty : {PrimType::Float(32), PrimType::Bool(), PrimType::Int(32,
4),
- PrimType::ScalableVector(kDLInt, 32, 4)}) {
- ffi::Any value = Var("x", ty);
- EXPECT_FALSE(value.try_cast<IntExpr>());
- EXPECT_TRUE(value.try_cast<PrimExpr>());
- }
- ffi::Any nonprimitive = Var("x", AnyType());
- EXPECT_FALSE(nonprimitive.try_cast<IntExpr>());
- EXPECT_FALSE(nonprimitive.try_cast<PrimExpr>());
- EXPECT_FALSE(nonprimitive.try_cast<TypedExpr<PrimType>>());
- EXPECT_EQ(ffi::Any(1).cast<IntExpr>().ty().bits(), 32);
- EXPECT_EQ(ffi::Any(int64_t{1} << 40).cast<IntExpr>().ty().bits(), 64);
-}
-
-TEST(ExprNodeRef, Basic) {
- using namespace tvm;
- using namespace tvm::tirx;
- PrimVar x("x");
- PrimExpr z = max(x + 1 + 2, 100);
- const prim::MaxNode* op = z.as<prim::MaxNode>();
- TVM_FFI_ICHECK(ffi::GetRef<ffi::ObjectRef>(op).same_as(z));
-}
-
-TEST(Expr, DeepEqualTensorLoadSourceIdentity) {
- using namespace tvm;
- Var source("source", PtrType(PrimType::Float(32)));
- Var other_source("source", PtrType(PrimType::Float(32)));
- auto load = [](Expr source, PrimExpr index) {
- auto node = ffi::make_object<TensorLoadNode>(source);
- node->ty = PrimType::Float(32);
- node->indices = {index};
- return TensorLoad(node);
- };
- prim::ExprDeepEqual equal;
- EXPECT_TRUE(equal(load(source, 0), load(source, 0)));
- EXPECT_FALSE(equal(load(source, 0), load(other_source, 0)));
- EXPECT_FALSE(equal(load(source, 0), load(source, 1)));
-}
diff --git a/tests/cpp/expr_functor_test.cc b/tests/cpp/ir/expr_functor_test.cc
similarity index 100%
rename from tests/cpp/expr_functor_test.cc
rename to tests/cpp/ir/expr_functor_test.cc
diff --git a/tests/cpp/ir_op_signature_test.cc
b/tests/cpp/ir/op_signature_test.cc
similarity index 79%
rename from tests/cpp/ir_op_signature_test.cc
rename to tests/cpp/ir/op_signature_test.cc
index c97dcfd653..2efa3ebc78 100644
--- a/tests/cpp/ir_op_signature_test.cc
+++ b/tests/cpp/ir/op_signature_test.cc
@@ -40,18 +40,6 @@ TEST(OpSignature, MetadataAndValidator) {
sig::var_args<PrimExpr>("rest", "Remaining indices."),
sig::ty_arg<PrimType>("T"),
sig::var_ty_args<PrimType>("Ts", "More types."),
sig::call_attrs<DictAttrsNode>());
Op op = def.op();
- EXPECT_EQ(op->args_info.size(), 2);
- EXPECT_EQ(op->args_info[0]->name, "value");
- EXPECT_EQ(op->args_info[0]->doc, "The input.");
- EXPECT_EQ(op->args_info[1]->name, "index");
- ASSERT_TRUE(op->var_args_info.has_value());
- EXPECT_EQ(op->var_args_info.value()->name, "rest");
- EXPECT_EQ(op->var_args_info.value()->doc, "Remaining indices.");
- EXPECT_EQ(op->ty_args_info.size(), 1);
- EXPECT_EQ(op->ty_args_info[0]->name, "T");
- ASSERT_TRUE(op->var_ty_args_info.has_value());
- EXPECT_EQ(op->var_ty_args_info.value()->name, "Ts");
- EXPECT_EQ(op->var_ty_args_info.value()->doc, "More types.");
EXPECT_EQ(op->attrs_type_key, DictAttrsNode::_type_key);
EXPECT_EQ(op->attrs_type_index, DictAttrsNode::RuntimeTypeIndex());
@@ -99,15 +87,7 @@ TEST(OpSignature, MetadataAndValidator) {
ffi::TypeTraitsBase::GetMismatchTypeInfo(&raw));
OpDef(op->name).set_attr<bool>("FSignatureTestMarker", true);
- EXPECT_EQ(op->args_info.size(), 2);
- EXPECT_TRUE(op->var_args_info.has_value());
EXPECT_NO_THROW(OpDef(op->name).signature(sig::arg("replacement")));
- ASSERT_EQ(op->args_info.size(), 1);
- EXPECT_EQ(op->args_info[0]->name, "replacement");
- EXPECT_FALSE(op->var_args_info.has_value());
- EXPECT_TRUE(op->ty_args_info.empty());
- EXPECT_FALSE(op->var_ty_args_info.has_value());
- EXPECT_TRUE(op->attrs_type_key.empty());
EXPECT_NO_THROW(op.Validate(minimum.get()));
EXPECT_NO_THROW(OpDef(op->name).add_arg("legacy", ""));
EXPECT_NO_THROW(OpDef(op->name).add_ty_arg("legacy_type", ""));
@@ -118,8 +98,6 @@ TEST(OpSignature, DefaultsAndRepeatedRegistration) {
OpDef plain("test.op_signature.defaults");
plain.signature(sig::arg("value"), sig::ty_arg("T"));
Op op = plain.op();
- EXPECT_FALSE(op->var_args_info.has_value());
- EXPECT_FALSE(op->var_ty_args_info.has_value());
EXPECT_NO_THROW(Call(AnyType(), op, {Var("x", AnyType())}, std::nullopt,
{AnyType()}));
EXPECT_NO_THROW(Call(AnyType(), op, {Var("x", AnyType())}, DictAttrs(),
{AnyType()}));
EXPECT_EQ(ValidationError(op, Call(AnyType(), op, {Var("x", AnyType()),
Var("y", AnyType())},
@@ -134,10 +112,6 @@ TEST(OpSignature, DefaultsAndRepeatedRegistration) {
OpDef tails("test.op_signature.default_tails");
tails.signature(sig::var_args("values"), sig::var_ty_args("types"));
- EXPECT_TRUE(tails.op()->args_info.empty());
- EXPECT_TRUE(tails.op()->ty_args_info.empty());
- EXPECT_EQ(tails.op()->var_args_info.value()->name, "values");
- EXPECT_EQ(tails.op()->var_ty_args_info.value()->name, "types");
EXPECT_NO_THROW(Call(AnyType(), tails.op(), {}, std::nullopt, {}));
EXPECT_NO_THROW(Call(AnyType(), tails.op(), {Var("x", AnyType())},
std::nullopt, {AnyType()}));
@@ -145,24 +119,14 @@ TEST(OpSignature, DefaultsAndRepeatedRegistration) {
EXPECT_NO_THROW(repeated.signature(sig::arg("", "value doc"),
sig::var_args("", "tail doc"),
sig::ty_arg("", "type doc"),
sig::var_ty_args("", "type tail doc")));
- EXPECT_EQ(repeated.op()->args_info[0]->name, "");
- EXPECT_EQ(repeated.op()->var_args_info.value()->name, "");
- EXPECT_EQ(repeated.op()->ty_args_info[0]->name, "");
- EXPECT_EQ(repeated.op()->var_ty_args_info.value()->name, "");
EXPECT_TRUE(repeated.op()->validator != nullptr);
EXPECT_NO_THROW(repeated.op().Validate(
Call(AnyType(), repeated.op(), {Var("x", AnyType())}, std::nullopt,
{AnyType()}).get()));
EXPECT_NO_THROW(repeated.signature(sig::arg("changed"), sig::arg("")));
- EXPECT_EQ(repeated.op()->args_info.size(), 2);
- EXPECT_EQ(repeated.op()->args_info[0]->name, "changed");
- EXPECT_EQ(repeated.op()->args_info[1]->name, "");
OpDef legacy("test.op_signature.legacy");
legacy.add_arg("old", "Existing metadata.");
EXPECT_NO_THROW(legacy.signature(sig::arg("new")));
- EXPECT_EQ(legacy.op()->args_info.size(), 1);
- EXPECT_EQ(legacy.op()->args_info[0]->name, "new");
- EXPECT_FALSE(legacy.op()->var_args_info.has_value());
EXPECT_TRUE(legacy.op()->validator != nullptr);
OpDef manual("test.op_signature.manual_hook");
@@ -176,7 +140,6 @@ TEST(OpSignature, DefaultsAndRepeatedRegistration) {
EXPECT_THROW(manual.set_validator(manual_hook), ffi::Error);
EXPECT_NO_THROW(manual.set_validator(manual_hook, true));
EXPECT_NO_THROW(manual.signature(sig::arg("value")));
- EXPECT_EQ(manual.op()->args_info[0]->name, "value");
Call manual_call(AnyType(), manual.op(), {Var("x", AnyType())});
manual.op().Validate(manual_call.get());
EXPECT_EQ(manual_calls, 1);
diff --git a/tests/cpp/ir_functor_test.cc b/tests/cpp/ir_functor_test.cc
deleted file mode 100644
index 5644ef9251..0000000000
--- a/tests/cpp/ir_functor_test.cc
+++ /dev/null
@@ -1,758 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-#include <tvm/ffi/extra/structural_equal.h>
-#include <tvm/ffi/extra/structural_mutate.h>
-#include <tvm/ffi/extra/structural_visit.h>
-#include <tvm/ir/expr_functor.h>
-#include <tvm/ir/module.h>
-#include <tvm/ir/object_functor.h>
-#include <tvm/ir/prim/expr.h>
-#include <tvm/ir/prim/op.h>
-#include <tvm/runtime/logging.h>
-#include <tvm/s_tir/stmt.h>
-#include <tvm/s_tir/stmt_functor.h>
-#include <tvm/tirx/analysis.h>
-#include <tvm/tirx/function.h>
-#include <tvm/tirx/op/abi.h>
-#include <tvm/tirx/op/memory.h>
-
-#include <initializer_list>
-#include <unordered_set>
-
-TEST(IRF, Basic) {
- using namespace tvm;
- using namespace tvm::tirx;
- PrimVar x("x");
- auto z = x + 1;
-
- ObjectFunctor<int(const ffi::ObjectRef& n, int b)> f;
- f.SetDispatch<VarNode>([](const ffi::ObjectRef& n, int b) { return b; });
- f.SetDispatch<prim::AddNode>([](const ffi::ObjectRef& n, int b) { return b +
2; });
- TVM_FFI_ICHECK_EQ(f(x, 2), 2);
- TVM_FFI_ICHECK_EQ(f(z, 2), 4);
-}
-
-TEST(IRF, ObjectFunctorDispatch) {
- using namespace tvm;
- PrimVar x("x");
- ObjectFunctor<int(const ffi::ObjectRef&)> f;
-
- EXPECT_FALSE(f.CanDispatch(x));
- EXPECT_THROW(f(x), ffi::Error);
- f.SetDispatch<ffi::Object>([](const ffi::ObjectRef&) {
- return 1;
- }).SetDispatch<ExprNode>([](const ffi::ObjectRef&) { return 2; });
- // Invocation uses the nearest registered ancestor; CanDispatch checks only
the exact type.
- EXPECT_FALSE(f.CanDispatch(x));
- EXPECT_EQ(f(x), 2);
- f.SetDispatch<VarNode>([](const ffi::ObjectRef&) { return 3; });
- EXPECT_TRUE(f.CanDispatch(x));
- EXPECT_EQ(f(x), 3);
- EXPECT_THROW(f.SetDispatch<VarNode>([](const ffi::ObjectRef&) { return 4;
}), ffi::Error);
-
- f.ClearDispatch<VarNode>();
- EXPECT_FALSE(f.CanDispatch(x));
- EXPECT_EQ(f(x), 2);
- f.ClearDispatch<ExprNode>();
- EXPECT_EQ(f(x), 1);
- f.SetDispatch<VarNode>([](const ffi::ObjectRef&) { return 4; });
- EXPECT_TRUE(f.CanDispatch(x));
- EXPECT_EQ(f(x), 4);
-}
-
-TEST(IRF, ObjectFunctorFinalize) {
- using namespace tvm;
- PrimVar x("x");
- PrimExpr z = x + 1;
- ObjectFunctor<int(const ffi::ObjectRef&, int)> f;
- f.SetDispatch<ExprNode>([](const ffi::ObjectRef&, int b) {
- return b;
- }).SetDispatch<prim::AddNode>([](const ffi::ObjectRef&, int b) { return b +
2; });
-
- f.Finalize();
- EXPECT_FALSE(f.CanDispatch(x));
- EXPECT_TRUE(f.CanDispatch(z));
- EXPECT_EQ(f(x, 2), 2);
- EXPECT_EQ(f(z, 2), 4);
- EXPECT_THROW(f.Finalize(), ffi::Error);
- EXPECT_THROW(f.SetDispatch<VarNode>([](const ffi::ObjectRef&, int b) {
return b; }), ffi::Error);
- EXPECT_THROW(f.ClearDispatch<ExprNode>(), ffi::Error);
-}
-
-TEST(IRF, CountVar) {
- using namespace tvm;
- using namespace tvm::tirx;
- int n_var = 0;
- PrimVar x("x"), y("y");
-
- auto z = x + 1 + y + y;
- std::unordered_set<const ffi::Object*> visited;
- auto walk_fn = [&n_var, &visited](const Var& var) ->
ffi::Expected<ffi::WalkResult> {
- if (visited.insert(var.get()).second) {
- ++n_var;
- }
- return ffi::WalkResult::Advance();
- };
- ffi::StructuralWalk<ffi::WalkOrder::kPostOrder>(z, walk_fn);
- TVM_FFI_ICHECK_EQ(n_var, 2);
-}
-
-TEST(IRF, PreOrderStructuralWalk) {
- using namespace tvm;
- using namespace tvm::tirx;
- Stmt init =
- If(IntImm::Bool(true), Evaluate(IntImm::Int32(2)),
SeqStmt(Evaluate(IntImm::Int32(2))));
- Stmt body = Evaluate(IntImm::Int32(1));
- s_tir::SBlock block(/*iter_vars=*/{}, /*reads=*/{},
- /*writes=*/{}, /*name_hint=*/"block", /*body=*/body,
- /*init=*/SeqStmt(init));
- bool init_visited = false;
- bool stopped_at_if = true;
- bool body_visited = false;
- auto visit_if = [&](const If&) -> ffi::Expected<ffi::WalkResult> {
- init_visited = true;
- return ffi::WalkResult::Skip();
- };
- auto visit_evaluate = [&](const Evaluate& eval) ->
ffi::Expected<ffi::WalkResult> {
- if (const auto* int_imm = eval->value.as<IntImmNode>()) {
- if (int_imm->value == 2) {
- stopped_at_if = false;
- } else if (int_imm->value == 1) {
- body_visited = true;
- } else {
- TVM_FFI_THROW(InternalError) << "Unreachable";
- }
- }
- return ffi::WalkResult::Advance();
- };
- ffi::StructuralWalk<ffi::WalkOrder::kPreOrder>(block, visit_if,
visit_evaluate);
- ASSERT_EQ(init_visited, true);
- ASSERT_EQ(stopped_at_if, true);
- ASSERT_EQ(body_visited, true);
-}
-
-TEST(IRF, ExprTransform) {
- using namespace tvm;
- using namespace tvm::tirx;
- PrimVar x("x");
- auto z = x + 1;
-
- class MyExprFunctor : public tvm::ExprFunctor<int(const Expr&, int)> {
- public:
- int Dispatch_(const VarNode* op, int b) final { return b; }
- int Dispatch_(const IntImmNode* op, int b) final { return
op->value.as<int>().value(); }
- int Dispatch_(const prim::AddNode* op, int b) final {
- return Dispatch(op->a, b) + Dispatch(op->b, b);
- }
- };
- MyExprFunctor f;
- TVM_FFI_ICHECK_EQ(f(x, 2), 2);
- TVM_FFI_ICHECK_EQ(f(z, 2), 3);
- try {
- f(z - 1, 2);
- TVM_FFI_THROW(InternalError) << "should fail";
- } catch (tvm::ffi::Error&) {
- }
-}
-
-TEST(IRF, ExprVisit) {
- using namespace tvm;
- using namespace tvm::tirx;
- PrimVar x("x");
- auto z = x + 1;
-
- class MyVisitor : public tvm::ExprFunctor<void(const Expr&)>,
- public tirx::StmtFunctor<void(const Stmt&)> {
- public:
- using tvm::ExprFunctor<void(const Expr&)>::Dispatch;
- using tirx::StmtFunctor<void(const Stmt&)>::Dispatch;
- int count = 0;
- // implementation
- void Dispatch_(const VarNode* op) final { ++count; }
- void Dispatch_(const IntImmNode* op) final {}
- void Dispatch_(const prim::AddNode* op) final {
- Dispatch(op->a);
- Dispatch(op->b);
- }
- void Dispatch_(const EvaluateNode* op) final { Dispatch(op->value); }
- };
- MyVisitor v;
- v.Dispatch(Evaluate(z));
- TVM_FFI_ICHECK_EQ(v.count, 1);
-}
-
-TEST(IRF, StmtVisitor) {
- using namespace tvm;
- using namespace tvm::tirx;
- PrimVar x("x");
- class MyVisitor : public s_tir::StmtExprVisitor {
- public:
- int count = 0;
- // implementation
- ffi::Optional<VisitInterrupt> Visit_(const VarNode* op) final {
- // Buffer variables now share this hook; this fixture counts other Var
operands.
- if (!op->ty.as<TensorTypeNode>()) ++count;
- return std::nullopt;
- }
- };
- auto v = ffi::make_object<MyVisitor>();
- auto fmaketest = [&]() {
- auto z = x + 1;
- Stmt eval_body = Evaluate(z);
- PrimType dtype = PrimType::Float(32);
- TensorVar buf("b", TensorType("global", dtype, {z, z}, {}, std::nullopt,
0, 0));
- // AllocTensor is flat (no body). Return as SeqStmt with eval.
- return SeqStmt(
- {Bind(buf.var(), Call(buf.type(), tirx::alloc_tensor_op(), {},
DictAttrs(), {buf.type()})),
- eval_body});
- };
- v->Visit(fmaketest());
- // AllocTensor visits buffer shape at its definition site.
- // shape = {z, z} where z = x + 1, so x is visited twice from shape + once
from eval = 3
- TVM_FFI_ICHECK_EQ(v->count, 3);
-
- {
- // tests for block and block_realize
- Stmt body = fmaketest();
- PrimType dtype = PrimType::Float(32);
- tvm::Var buf_var("b", PtrType(dtype));
- TensorVar buffer = decl_tensor({16});
- body = SeqStmt({Bind(buffer, Call(buffer.type(),
tvm::tirx::decl_tensor_op(), {buf_var}, {},
- {buffer.type()})),
- std::move(body)});
- TensorRegion buffer_region =
- TensorRegion(buffer, {Range::FromMinExtent(x + 1, 1)},
TensorRegionType());
- s_tir::MatchTensorRegion match_buffer_region(decl_tensor({1}),
buffer_region);
-
- // construct block and block_realize
- s_tir::SBlock block = s_tir::SBlock({}, {buffer_region}, {buffer_region},
"block", body,
- SeqStmt(body), {},
{match_buffer_region});
- Stmt block_realize = s_tir::SBlockRealize({}, IntImm::Bool(true), block);
-
- v->count = 0;
- v->Visit(block_realize);
- // x visited in: reads range (1), writes range (1), match_tensors range
(1).
- // init: DeclTensor data b(1) + AllocTensor shape x,x(2) + Evaluate x(1) =
4.
- // body: DeclTensor data b(1) + AllocTensor shape x,x(2) + Evaluate x(1) =
4.
- // Total: 1 + 1 + 1 + 4 + 4 = 11.
- TVM_FFI_ICHECK_EQ(v->count, 11);
- }
-}
-
-TEST(IRF, StmtExprMutator) {
- using namespace tvm;
- using namespace tvm::tirx;
- PrimVar x("x");
-
- class MyMutator : public s_tir::StmtExprMutator {
- public:
- using s_tir::StmtExprMutator::Mutate_;
- UnchangedOr<PrimExpr> Mutate_(const prim::AddNode* op, InplaceMode) final
{ return op->a; }
- };
- auto fmakealloc = [&]() {
- auto z = x + 1;
- PrimType dtype = PrimType::Float(32);
- TensorVar buf("b", TensorType("global", dtype, {1, z}, {}, std::nullopt,
0, 0));
- return Bind(buf.var(),
- Call(buf.type(), tirx::alloc_tensor_op(), {}, DictAttrs(),
{buf.type()}));
- };
-
- auto fmakeif = [&]() {
- auto z = x + 1;
- Stmt body = Evaluate(z);
- return If(x, Evaluate(0), SeqStmt(body));
- };
-
- auto v = ffi::make_object<MyMutator>();
- {
- auto alloc = fmakealloc();
- Stmt body2 = Evaluate(1);
- auto* bufptr = alloc.as<BindNode>()->var.get();
- ffi::Array<Stmt> arr{std::move(alloc), body2, body2};
- auto* arrptr = arr.get();
- arr.MutateByApply([&](Stmt s) { return
v->Mutate(s).ValueOrUnchanged(std::move(s)); });
- TVM_FFI_ICHECK(arr.get() == arrptr);
- // buffer IS mutated now (AllocTensor mutator visits buffer shape at the
buffer definition)
- // shape was {1, x+1}, mutator transforms x+1 -> x, so buffer changes
- TVM_FFI_ICHECK(arr[0].as<BindNode>()->var.get() != bufptr);
- }
- {
- ffi::Array<Stmt> arr{fmakealloc()};
- // mutate array get reference by another one, trigger copy.
- ffi::Array<Stmt> arr2 = arr;
- auto* arrptr = arr.get();
- arr.MutateByApply([&](Stmt s) { return
v->Mutate(s).ValueOrUnchanged(std::move(s)); });
- TVM_FFI_ICHECK(arr.get() != arrptr);
- // buffer is mutated in arr but not in arr2
- TVM_FFI_ICHECK(arr[0].as<BindNode>()->var.get() !=
arr2[0].as<BindNode>()->var.get());
- // mutate but no content change.
- arr2 = arr;
- arr.MutateByApply([&](Stmt s) { return
v->Mutate(s).ValueOrUnchanged(std::move(s)); });
- TVM_FFI_ICHECK(arr2.get() == arr.get());
- }
- {
- ffi::Array<Stmt> arr{fmakeif()};
- arr.MutateByApply([&](Stmt s) { return
v->Mutate(s).ValueOrUnchanged(std::move(s)); });
- TVM_FFI_ICHECK(
-
arr[0].as<IfNode>()->else_case.value()->seq[0].as<EvaluateNode>()->value.same_as(x));
- // mutate but no content change.
- auto arr2 = arr;
- arr.MutateByApply([&](Stmt s) { return
v->Mutate(s).ValueOrUnchanged(std::move(s)); });
- TVM_FFI_ICHECK(arr2.get() == arr.get());
- }
-
- {
- auto body =
- Evaluate(Call(PrimType::Int(32), tirx::call_extern_op(),
{::tvm::StringImm("xyz"), x + 1}));
- auto res = v->Mutate(body).ValueOrUnchanged(std::move(body));
-
TVM_FFI_ICHECK(res.as<EvaluateNode>()->value.as<CallNode>()->args[1].same_as(x));
- }
- {
- Stmt body = fmakealloc();
- Stmt body2 = Evaluate(1);
- auto* ref2 = body2.get();
- auto* bufptr = body.as<BindNode>()->var.get();
- // construct a recursive SeqStmt.
- body = SeqStmt({body, body2});
- body = SeqStmt({body, body2});
- body = v->Mutate(body).ValueOrUnchanged(std::move(body));
- // the seq get flattened
- TVM_FFI_ICHECK(body.as<SeqStmtNode>()->size() == 3);
- // buffer is now mutated (shape x+1 -> x at the buffer definition)
- TVM_FFI_ICHECK(body.as<SeqStmtNode>()->seq[0].as<BindNode>()->var.get() !=
bufptr);
- TVM_FFI_ICHECK(body.as<SeqStmtNode>()->seq[1].get() == ref2);
- }
-
- {
- // Cannot cow because of bref
- Stmt body = fmakealloc();
- Stmt body2 = Evaluate(1);
- // construct a recursive SeqStmt.
- body = SeqStmt({body, body2});
- auto bref = body;
- body = SeqStmt({body, body2});
- body = v->Mutate(body).ValueOrUnchanged(std::move(body));
- // the seq get flattened
- TVM_FFI_ICHECK(body.as<SeqStmtNode>()->size() == 3);
- // buffer is mutated (shape x+1 -> x at the buffer definition)
- auto* alloc_node = body.as<SeqStmtNode>()->seq[0].as<BindNode>();
- TVM_FFI_ICHECK(alloc_node != nullptr);
- auto* alloc_call = alloc_node->value.as<CallNode>();
- TVM_FFI_ICHECK(alloc_call &&
alloc_call->op.same_as(tirx::alloc_tensor_op()));
- // bref still holds the old SeqStmt (not shared with new one due to copy)
- TVM_FFI_ICHECK(!bref.same_as(body));
- }
-
- {
- // tests for block and block_realize
- // AllocTensor and DeclTensor are flat (no body), placed as siblings in
SeqStmt
- Stmt eval_body = Evaluate(x + 1);
- TensorVar buffer = decl_tensor({16});
- tvm::Var buffer_data("buffer_data", buffer.type()->DataPtrType());
- Stmt decl = Bind(buffer, Call(buffer.type(), tvm::tirx::decl_tensor_op(),
{buffer_data}, {},
- {buffer.type()}));
- Stmt alloc = fmakealloc();
- // body is: DeclTensor, AllocTensor, Evaluate
- Stmt body = SeqStmt({decl, alloc, eval_body});
- TensorRegion buffer_region =
- TensorRegion(buffer, {Range::FromMinExtent(x + 1, 1)},
TensorRegionType());
- s_tir::MatchTensorRegion match_buffer_region(decl_tensor({1}),
buffer_region);
- // construct block and block_realize
- s_tir::SBlock block = s_tir::SBlock({}, {buffer_region}, {buffer_region},
"block", body,
- SeqStmt(body), {},
{match_buffer_region});
- Stmt block_realize = s_tir::SBlockRealize({}, IntImm::Bool(true), block);
- body = v->Mutate(block_realize).ValueOrUnchanged(std::move(block_realize));
- // the body should be changed
- s_tir::SBlock new_block = body.as<s_tir::SBlockRealizeNode>()->block;
- // body is a SeqStmt; the Evaluate(x+1) -> Evaluate(x)
- auto* seq = new_block->body.as<SeqStmtNode>();
- TVM_FFI_ICHECK(seq != nullptr);
- TVM_FFI_ICHECK(seq->seq[2].as<EvaluateNode>()->value.same_as(x));
- auto* init_seq = new_block->init.value().as<SeqStmtNode>();
- TVM_FFI_ICHECK(init_seq != nullptr);
- TVM_FFI_ICHECK(init_seq->seq[2].as<EvaluateNode>()->value.same_as(x));
- // buffer region min is mutated: x+1 -> x
- TVM_FFI_ICHECK(new_block->reads[0]->region[0]->min.same_as(x));
- TVM_FFI_ICHECK(new_block->writes[0]->region[0]->min.same_as(x));
-
TVM_FFI_ICHECK(new_block->match_tensors[0]->source->region[0]->min.same_as(x));
- }
-}
-
-TEST(IRF, StructuralMapSplicesMappedSeqStmtChild) {
- using namespace tvm;
- using namespace tvm::tirx;
-
- auto make_input = []() -> Stmt {
- return SeqStmt(
- {Evaluate(IntImm::Int32(5)), Evaluate(IntImm::Int32(1)),
Evaluate(IntImm::Int32(4))});
- };
- auto expand_one = [](const Evaluate& evaluate) -> Stmt {
- const auto* value = evaluate->value.as<IntImmNode>();
- if (value != nullptr && value->value == 1) {
- return SeqStmt({Evaluate(IntImm::Int32(2)), Evaluate(IntImm::Int32(3))});
- }
- return evaluate;
- };
- auto check_values = [](const Stmt& stmt, std::initializer_list<int64_t>
expected) {
- const auto* seq = stmt.as<SeqStmtNode>();
- ASSERT_NE(seq, nullptr);
- ASSERT_EQ(seq->seq.size(), expected.size());
- size_t i = 0;
- for (int64_t expected_value : expected) {
- const auto* evaluate = seq->seq[i].as<EvaluateNode>();
- ASSERT_NE(evaluate, nullptr);
- const auto* value = evaluate->value.as<IntImmNode>();
- ASSERT_NE(value, nullptr);
- EXPECT_EQ(value->value, expected_value);
- ++i;
- }
- };
- {
- Stmt input = make_input();
- Stmt shared = input;
- Stmt mapped =
- ffi::StructuralMap<ffi::WalkOrder::kPostOrder>(input,
expand_one).as_or_throw<Stmt>();
- EXPECT_FALSE(mapped.same_as(input));
- EXPECT_EQ(shared.as<SeqStmtNode>()->seq.size(), 3);
- check_values(mapped, {5, 2, 3, 4});
- }
-
- {
- Stmt input = make_input();
- const auto* original = input.get();
- const auto* original_array = input.as<SeqStmtNode>()->seq.GetArrayObj();
- Stmt mapped =
ffi::StructuralMap<ffi::WalkOrder::kPostOrder>(std::move(input), expand_one)
- .as_or_throw<Stmt>();
- EXPECT_EQ(mapped.get(), original);
- EXPECT_NE(mapped.as<SeqStmtNode>()->seq.GetArrayObj(), original_array);
- check_values(mapped, {5, 2, 3, 4});
- }
-
- auto make_boundary_input = []() -> Stmt {
- return SeqStmt({Evaluate(IntImm::Int32(1)), Evaluate(IntImm::Int32(2)),
- Evaluate(IntImm::Int32(3)), Evaluate(IntImm::Int32(4))});
- };
- auto check_differential = [&](const auto& transform,
std::initializer_list<int64_t> expected,
- bool expect_array_reuse) {
- Stmt ordinary_input = make_boundary_input();
- Stmt ordinary =
ffi::StructuralMap<ffi::WalkOrder::kPostOrder>(ordinary_input, transform)
- .template as_or_throw<Stmt>();
-
- Stmt inplace_input = make_boundary_input();
- const auto* original_root = inplace_input.get();
- const auto* original_array =
inplace_input.as<SeqStmtNode>()->seq.GetArrayObj();
- Stmt inplace =
-
ffi::StructuralMap<ffi::WalkOrder::kPostOrder>(std::move(inplace_input),
transform)
- .template as_or_throw<Stmt>();
-
- EXPECT_EQ(inplace.get(), original_root);
- if (expect_array_reuse) {
- EXPECT_EQ(inplace.as<SeqStmtNode>()->seq.GetArrayObj(), original_array);
- } else {
- EXPECT_NE(inplace.as<SeqStmtNode>()->seq.GetArrayObj(), original_array);
- }
- EXPECT_TRUE(ffi::StructuralEqual()(ordinary, inplace));
- check_values(inplace, expected);
- };
-
- auto shrink_first_grow_last = [](const Evaluate& evaluate) -> Stmt {
- const auto* value = evaluate->value.as<IntImmNode>();
- if (value != nullptr && value->value == 1) {
- return Evaluate(0);
- }
- if (value != nullptr && value->value == 4) {
- return SeqStmt({Evaluate(IntImm::Int32(30)),
Evaluate(IntImm::Int32(31))});
- }
- return evaluate;
- };
- check_differential(shrink_first_grow_last, {2, 3, 30, 31}, true);
-
- auto empty_first_grow_last = [](const Evaluate& evaluate) -> Stmt {
- const auto* value = evaluate->value.as<IntImmNode>();
- if (value != nullptr && value->value == 1) {
- auto empty = ffi::make_object<SeqStmtNode>();
- empty->seq = {};
- return Stmt(std::move(empty));
- }
- if (value != nullptr && value->value == 4) {
- return SeqStmt({Evaluate(IntImm::Int32(30)),
Evaluate(IntImm::Int32(31))});
- }
- return evaluate;
- };
- check_differential(empty_first_grow_last, {2, 3, 30, 31}, true);
-
- int overflow_callback_count = 0;
- auto grow_first_shrink_last = [&overflow_callback_count](const Evaluate&
evaluate) -> Stmt {
- ++overflow_callback_count;
- const auto* value = evaluate->value.as<IntImmNode>();
- if (value != nullptr && value->value == 1) {
- return SeqStmt({Evaluate(IntImm::Int32(10)),
Evaluate(IntImm::Int32(11))});
- }
- if (value != nullptr && value->value == 4) {
- return Evaluate(0);
- }
- return evaluate;
- };
- check_differential(grow_first_shrink_last, {10, 11, 2, 3}, true);
- EXPECT_EQ(overflow_callback_count, 8);
-
- auto grow_last_over_capacity = [](const Evaluate& evaluate) -> Stmt {
- const auto* value = evaluate->value.as<IntImmNode>();
- if (value != nullptr && value->value == 4) {
- return SeqStmt({Evaluate(IntImm::Int32(40)),
Evaluate(IntImm::Int32(41))});
- }
- return evaluate;
- };
- check_differential(grow_last_over_capacity, {1, 2, 3, 40, 41}, false);
-
- auto replace_middle_with_one = [](const Evaluate& evaluate) -> Stmt {
- const auto* value = evaluate->value.as<IntImmNode>();
- if (value != nullptr && value->value == 2) {
- return Evaluate(IntImm::Int32(10));
- }
- return evaluate;
- };
- check_differential(replace_middle_with_one, {1, 10, 3, 4}, true);
-
- auto remove_first = [](const Evaluate& evaluate) -> Stmt {
- const auto* value = evaluate->value.as<IntImmNode>();
- return value != nullptr && value->value == 1 ? Evaluate(0) :
Stmt(evaluate);
- };
- check_differential(remove_first, {2, 3, 4}, true);
-
- auto remove_all = [](const Evaluate&) -> Stmt { return Evaluate(0); };
- Stmt ordinary_input = make_boundary_input();
- Stmt ordinary =
ffi::StructuralMap<ffi::WalkOrder::kPostOrder>(ordinary_input, remove_all)
- .as_or_throw<Stmt>();
- Stmt inplace_input = make_boundary_input();
- Stmt inplace =
- ffi::StructuralMap<ffi::WalkOrder::kPostOrder>(std::move(inplace_input),
remove_all)
- .as_or_throw<Stmt>();
- EXPECT_TRUE(ffi::StructuralEqual()(ordinary, inplace));
- for (const Stmt& result : {ordinary, inplace}) {
- const auto* sequence = result.as<SeqStmtNode>();
- ASSERT_NE(sequence, nullptr);
- EXPECT_TRUE(sequence->seq.empty());
- }
-
- auto keep_last = [](const Evaluate& evaluate) -> Stmt {
- const auto* value = evaluate->value.as<IntImmNode>();
- return value != nullptr && value->value == 4 ? Stmt(evaluate) :
Stmt(Evaluate(0));
- };
- ordinary_input = make_boundary_input();
- ordinary =
- ffi::StructuralMap<ffi::WalkOrder::kPostOrder>(ordinary_input,
keep_last).as_or_throw<Stmt>();
- inplace_input = make_boundary_input();
- inplace =
ffi::StructuralMap<ffi::WalkOrder::kPostOrder>(std::move(inplace_input),
keep_last)
- .as_or_throw<Stmt>();
- EXPECT_TRUE(ffi::StructuralEqual()(ordinary, inplace));
- for (const Stmt& result : {ordinary, inplace}) {
- const auto* sequence = result.as<SeqStmtNode>();
- ASSERT_NE(sequence, nullptr);
- ASSERT_EQ(sequence->seq.size(), 1);
- const auto* evaluate = sequence->seq[0].as<EvaluateNode>();
- ASSERT_NE(evaluate, nullptr);
- const auto* value = evaluate->value.as<IntImmNode>();
- ASSERT_NE(value, nullptr);
- EXPECT_EQ(value->value, 4);
- }
-}
-
-TEST(IRF, StructuralMapPreservesSeqStmtElementUniqueness) {
- using namespace tvm;
- using namespace tvm::tirx;
-
- auto replace_one = [](const IntImm& value) -> PrimExpr {
- return value->value == 1 ? IntImm::Int32(2) : PrimExpr(value);
- };
-
- {
- Stmt input = SeqStmt({Evaluate(IntImm::Int32(1)),
Evaluate(IntImm::Int32(3))});
- const auto* original_root = input.get();
- const auto* original_first =
input.as<SeqStmtNode>()->seq[0].as<EvaluateNode>();
- Stmt mapped =
ffi::StructuralMap<ffi::WalkOrder::kPostOrder>(std::move(input), replace_one)
- .as_or_throw<Stmt>();
-
- const auto* mapped_seq = mapped.as<SeqStmtNode>();
- ASSERT_NE(mapped_seq, nullptr);
- EXPECT_EQ(mapped.get(), original_root);
- EXPECT_EQ(mapped_seq->seq[0].get(), original_first);
-
EXPECT_EQ(mapped_seq->seq[0].as<EvaluateNode>()->value.as<IntImmNode>()->value,
2);
- }
-
- {
- ffi::Array<Stmt> shared_seq = {Evaluate(IntImm::Int32(1)),
Evaluate(IntImm::Int32(3))};
- const auto* shared_first = shared_seq[0].as<EvaluateNode>();
- Stmt input = SeqStmt(shared_seq);
- Stmt mapped =
ffi::StructuralMap<ffi::WalkOrder::kPostOrder>(std::move(input), replace_one)
- .as_or_throw<Stmt>();
-
- const auto* mapped_seq = mapped.as<SeqStmtNode>();
- ASSERT_NE(mapped_seq, nullptr);
- EXPECT_FALSE(mapped_seq->seq.same_as(shared_seq));
- EXPECT_EQ(shared_seq[0].get(), shared_first);
- EXPECT_EQ(shared_seq[0].as<EvaluateNode>()->value.as<IntImmNode>()->value,
1);
-
EXPECT_EQ(mapped_seq->seq[0].as<EvaluateNode>()->value.as<IntImmNode>()->value,
2);
-
- Stmt unchanged_input = SeqStmt(shared_seq);
- const auto* unchanged_root = unchanged_input.get();
- auto no_float_match = [](const FloatImm& value) -> PrimExpr { return
value; };
- Stmt unchanged =
-
ffi::StructuralMap<ffi::WalkOrder::kPostOrder>(std::move(unchanged_input),
no_float_match)
- .as_or_throw<Stmt>();
- EXPECT_EQ(unchanged.get(), unchanged_root);
- EXPECT_TRUE(unchanged.as<SeqStmtNode>()->seq.same_as(shared_seq));
- }
-}
-
-TEST(IRF, StructuralHooksPreserveIndexBindRegions) {
- using namespace tvm;
- using namespace tvm::tirx;
-
- auto make_input = []() -> Stmt {
- return Bind(PrimVar("binder"), PrimVar("extent") + PrimVar("preferred"));
- };
- auto check_kinds = [](int binder, int extent, int preferred) {
- EXPECT_EQ(binder, kTVMFFIDefRegionKindSimple);
- EXPECT_EQ(extent, kTVMFFIDefRegionKindNone);
- EXPECT_EQ(preferred, kTVMFFIDefRegionKindNone);
- };
-
- {
- int binder = -1;
- int extent = -1;
- int preferred = -1;
- auto walk_fn = [&](const Var& var, TVMFFIDefRegionKind kind) ->
ffi::Expected<ffi::WalkResult> {
- if (var->name == "binder") binder = kind;
- if (var->name == "extent") extent = kind;
- if (var->name == "preferred") preferred = kind;
- return ffi::WalkResult::Advance();
- };
- ffi::StructuralWalk<ffi::WalkOrder::kPostOrder>(make_input(), walk_fn);
- check_kinds(binder, extent, preferred);
- }
-
- auto check_map = [&](Stmt input) {
- int binder = -1;
- int extent = -1;
- int preferred = -1;
- ffi::StructuralMap<ffi::WalkOrder::kPostOrder>(
- std::move(input), [&](const Var& var, TVMFFIDefRegionKind kind) -> Var
{
- if (var->name == "binder") binder = kind;
- if (var->name == "extent") extent = kind;
- if (var->name == "preferred") preferred = kind;
- return var;
- });
- check_kinds(binder, extent, preferred);
- };
-
- {
- Stmt input = make_input();
- Stmt shared = input;
- check_map(input);
- }
- check_map(make_input());
-}
-
-TEST(IRF, StructuralMapBufferDefinition) {
- using namespace tvm;
- using namespace tvm::tirx;
- PrimType dtype = PrimType::Float(32);
- tvm::Var x("x", PtrType(dtype, ""));
- PrimVar n("n", PrimType::Int(32));
-
- auto fmakebuffer = [&]() {
- return TensorVar("buf", TensorType(/*storage_scope=*/"global",
- /*dtype=*/PrimType::Float(32),
- /*shape=*/{n},
- /*strides=*/{},
- /*elem_offset=*/std::nullopt,
- /*data_alignment=*/1,
- /*offset_factor=*/1));
- };
-
- {
- // Test substitution of an explicit DeclTensor source and a dependent
- // TensorType shape. Changing the type creates one fresh Var identity
- // that is shared by the declaration and every use.
- tvm::Var y = x.CopyWithSuffix("subst");
- PrimVar m("m", PrimType::Int(32));
- TensorVar buffer = fmakebuffer();
- Stmt store = TensorStore(buffer, {IntImm::Int32(0)}, FloatImm(dtype, 0));
- Stmt decl = SeqStmt(
- {Bind(buffer, Call(buffer.type(), tvm::tirx::decl_tensor_op(), {x},
{}, {buffer.type()})),
- store});
- auto f_subst = [&](const tvm::Var& var) ->
ffi::Expected<ffi::UnchangedOr<ffi::Any>> {
- if (var.same_as(x)) return ffi::Any(y);
- if (var.same_as(n)) return ffi::Any(m);
- return ffi::Unchanged();
- };
- Stmt new_decl =
- ffi::StructuralMap<ffi::WalkOrder::kPreOrder>(decl,
f_subst).as_or_throw<Stmt>();
- auto* seq_node = new_decl.as<SeqStmtNode>();
- TVM_FFI_ICHECK(seq_node != nullptr);
- auto* decl_node = seq_node->seq[0].as<BindNode>();
- TVM_FFI_ICHECK(decl_node != nullptr);
- auto* decl_call = decl_node->value.as<CallNode>();
- TVM_FFI_ICHECK(decl_call && decl_call->op.same_as(tirx::decl_tensor_op()));
- TVM_FFI_ICHECK(decl_call->args[0].same_as(y));
-
TVM_FFI_ICHECK(decl_node->var.as_or_throw<TensorVar>()->shape[0].same_as(m));
- TVM_FFI_ICHECK(!decl_node->var.same_as(buffer));
- auto* store_node = seq_node->seq[1].as<TensorStoreNode>();
- TVM_FFI_ICHECK(store_node != nullptr);
- TVM_FFI_ICHECK(store_node->dest.same_as(decl_node->var));
- }
-
- {
- // test identity substitution on expression
- TensorVar buffer = fmakebuffer();
- PrimExpr expr = MakeTensorLoad(buffer, {IntImm::Int32(0)});
- auto f_subst = [&](const tvm::Var& var) ->
ffi::Expected<ffi::UnchangedOr<ffi::Any>> {
- return ffi::Any(var);
- };
- PrimExpr new_expr =
- ffi::StructuralMap<ffi::WalkOrder::kPreOrder>(expr,
f_subst).as_or_throw<PrimExpr>();
- // the expression is not changed
- TVM_FFI_ICHECK(new_expr.same_as(expr));
- }
-}
-
-TEST(IRF, SubstituteWithDataTypeLegalizationPreservesShiftAmounts) {
- using namespace tvm::prim;
- using namespace tvm;
- using namespace tvm::tirx;
-
- PrimVar x("x", PrimType::Int(64));
- PrimVar y("y", PrimType::Int(32));
- auto f_subst = [&](const tvm::Var& var) -> ffi::Optional<PrimExpr> {
- if (var.same_as(x)) return PrimExpr(y);
- return std::nullopt;
- };
-
- PrimExpr shift_amount = IntImm::Int64(40);
- PrimExpr widened_y = cast(PrimType::Int(64), y);
- PrimExpr actual_left = SubstituteWithDataTypeLegalization(x << shift_amount,
f_subst);
- PrimExpr actual_right = SubstituteWithDataTypeLegalization(x >>
shift_amount, f_subst);
-
- ffi::StructuralEqual structural_equal;
- EXPECT_TRUE(structural_equal(actual_left, widened_y << shift_amount));
- EXPECT_TRUE(structural_equal(actual_right, widened_y >> shift_amount));
-}
diff --git a/tests/cpp/ndarray_test.cc b/tests/cpp/ndarray_test.cc
deleted file mode 100644
index 42b81c051e..0000000000
--- a/tests/cpp/ndarray_test.cc
+++ /dev/null
@@ -1,73 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-#include <tvm/runtime/logging.h>
-#include <tvm/runtime/tensor.h>
-
-using namespace tvm;
-
-TEST(TensorTest, IsContiguous_ContiguousStride) {
- auto array = runtime::Empty({5, 10}, DLDataType{kDLFloat, 32, 1}, {kDLCPU});
- DLManagedTensor* managed_tensor = array.ToDLPack();
-
- int64_t strides[] = {10, 1};
- managed_tensor->dl_tensor.strides = strides;
-
- TVM_FFI_ICHECK(ffi::IsContiguous(managed_tensor->dl_tensor));
-
- managed_tensor->deleter(managed_tensor);
-}
-
-TEST(TensorTest, IsContiguous_NullStride) {
- auto array = runtime::Empty({5, 10}, DLDataType{kDLFloat, 32, 1}, {kDLCPU});
- DLManagedTensor* managed_tensor = array.ToDLPack();
-
- managed_tensor->dl_tensor.strides = nullptr;
-
- TVM_FFI_ICHECK(ffi::IsContiguous(managed_tensor->dl_tensor));
-
- managed_tensor->deleter(managed_tensor);
-}
-
-TEST(TensorTest, IsContiguous_AnyStrideForSingular) {
- auto array = runtime::Empty({5, 1, 10}, DLDataType{kDLFloat, 32, 1},
{kDLCPU});
- DLManagedTensor* managed_tensor = array.ToDLPack();
-
- int64_t strides[] = {10, 1, 1}; // strides[1] is normalized to 1 because
shape[1] == 1.
- managed_tensor->dl_tensor.strides = strides;
-
- TVM_FFI_ICHECK(ffi::IsContiguous(managed_tensor->dl_tensor));
-
- managed_tensor->dl_tensor.strides = nullptr;
- managed_tensor->deleter(managed_tensor);
-}
-
-TEST(TensorTest, IsContiguous_UncontiguousStride) {
- auto array = runtime::Empty({5, 1, 10}, DLDataType{kDLFloat, 32, 1},
{kDLCPU});
- DLManagedTensor* managed_tensor = array.ToDLPack();
-
- int64_t strides[] = {1, 1, 1};
- managed_tensor->dl_tensor.strides = strides;
-
- TVM_FFI_ICHECK(!ffi::IsContiguous(managed_tensor->dl_tensor));
-
- managed_tensor->dl_tensor.strides = nullptr;
- managed_tensor->deleter(managed_tensor);
-}
diff --git a/tests/cpp/nested_msg_test.cc b/tests/cpp/nested_msg_test.cc
deleted file mode 100644
index 2c732dd1de..0000000000
--- a/tests/cpp/nested_msg_test.cc
+++ /dev/null
@@ -1,343 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-#include <tvm/ffi/dtype.h>
-#include <tvm/ffi/extra/structural_equal.h>
-#include <tvm/ir/prim/expr.h>
-#include <tvm/relax/block_builder.h>
-#include <tvm/relax/nested_msg.h>
-#include <tvm/relax/type.h>
-#include <tvm/runtime/logging.h>
-
-#include <algorithm>
-#include <array>
-#include <cstring>
-#include <functional>
-#include <iterator>
-#include <new>
-#include <string>
-#include <unordered_map>
-#include <utility>
-#include <vector>
-
-using namespace tvm;
-using namespace tvm::relax;
-
-namespace {
-
-TensorType ScalarTensorType(PrimType dtype) {
- auto n = tvm::ffi::make_object<TensorTypeNode>();
- n->dtype = std::move(dtype);
- n->ndim = 0;
- return TensorType(n);
-}
-
-} // namespace
-
-TEST(NestedMsg, Basic) {
- // start with no annotation
- tvm::Var x("x", std::nullopt), y("y", std::nullopt);
-
- // constructor from array, T and nullopt.
- NestedMsg<relax::Expr> msg({x, std::nullopt, x});
-
- EXPECT_TRUE(msg.IsNested());
- EXPECT_FALSE(msg.IsLeaf());
- EXPECT_TRUE(msg != nullptr);
-
- EXPECT_ANY_THROW(msg.LeafValue());
-
- auto arr = msg.NestedArray();
- EXPECT_TRUE(arr[0].LeafValue().same_as(x));
- EXPECT_TRUE(arr[1] == nullptr);
- EXPECT_TRUE(arr[1].IsNull());
-
- EXPECT_TRUE(arr[2].LeafValue().same_as(x));
-
- auto a0 = arr[0];
- EXPECT_TRUE(a0.IsLeaf());
-
- // assignment
- // assign null
- a0 = std::nullopt;
- EXPECT_TRUE(a0 == nullptr);
-
- // assign array
- a0 = {x, {x, std::nullopt, y}};
- EXPECT_TRUE(a0.IsNested());
- auto t0 = a0.NestedArray()[1];
- EXPECT_TRUE(t0.IsNested());
- EXPECT_TRUE(t0.NestedArray()[2].LeafValue().same_as(y));
-
- // assign leaf
- a0 = x;
-
- EXPECT_TRUE(a0.IsLeaf());
- EXPECT_TRUE(a0.LeafValue().same_as(x));
-}
-
-TEST(NestedMsg, IntAndAny) {
- NestedMsg<int64_t> msg({1, std::nullopt, 2});
- Any any_msg = msg;
- NestedMsg<int64_t> msg2 = any_msg.cast<NestedMsg<int64_t>>();
-
- EXPECT_TRUE(msg2.IsNested());
- EXPECT_EQ(msg2.NestedArray()[0].LeafValue(), 1);
- EXPECT_TRUE(msg2.NestedArray()[1].IsNull());
- EXPECT_EQ(msg2.NestedArray()[2].LeafValue(), 2);
-}
-
-TEST(NestedMsg, ForEachLeaf) {
- tvm::Var x("x", std::nullopt), y("y", std::nullopt);
- NestedMsg<Expr> msg = {x, {x, y}, std::nullopt, {x, {x, y}}};
-
- int x_count = 0, y_count = 0;
-
- ForEachLeaf(msg, [&](const Expr& v) {
- if (v.same_as(x)) ++x_count;
- if (v.same_as(y)) ++y_count;
- });
- EXPECT_EQ(x_count, 4);
- EXPECT_EQ(y_count, 2);
-}
-
-TEST(NestedMsg, Equal) {
- tvm::Var x("x", std::nullopt), y("y", std::nullopt);
- tvm::Var z("z", std::nullopt);
-
- auto fequal = [](Expr lhs, Expr rhs) { return lhs.same_as(rhs); };
-
- using M = NestedMsg<relax::Expr>;
-
- EXPECT_TRUE(Equal(M(std::nullopt), M(std::nullopt), fequal));
-
- EXPECT_TRUE(Equal(M(x), M(x), fequal));
-
- EXPECT_TRUE(Equal(M({x, y}), M({x, y}), fequal));
-
- EXPECT_TRUE(Equal(M({x, std::nullopt}), M({x, std::nullopt}), fequal));
-
- EXPECT_TRUE(Equal(M({x, {std::nullopt, y}}), M({x, {std::nullopt, y}}),
fequal));
-
- EXPECT_TRUE(Equal(M({x, {std::nullopt, y}, {x, z}}), M({x, {std::nullopt,
y}, {x, z}}), fequal));
-
- // type mismatch
- EXPECT_FALSE(Equal(M({x, {std::nullopt, y}, x}), M({x, {std::nullopt, y},
{x, z}}), fequal));
-
- EXPECT_FALSE(Equal(M({x, {std::nullopt, y}, {x, std::nullopt}}),
- M({x, {std::nullopt, y}, {x, z}}), fequal));
-
- EXPECT_FALSE(Equal(M({x, {std::nullopt, y}}), M({x, {std::nullopt, y}, {x,
z}}), fequal));
-
- EXPECT_FALSE(Equal(M(x), M(std::nullopt), fequal));
-
- EXPECT_FALSE(Equal(M(std::nullopt), M(x), fequal));
-
- EXPECT_FALSE(Equal(M(x), M(ffi::Array<M>({x})), fequal));
-
- EXPECT_FALSE(Equal(M(ffi::Array<M>({x})), M(x), fequal));
-}
-
-TEST(NestedMsg, MapAndDecompose) {
- tvm::Var x("x", PrimType::Int(16));
- tvm::Var y("y", PrimType::Int(32));
- tvm::Var z("z", PrimType::Int(64));
-
- BlockBuilder bb = BlockBuilder::Create(std::nullopt);
- relax::Expr t0 = bb->Normalize(Tuple({x, y}));
- relax::Expr t1 = bb->Normalize(Tuple({t0, x, z, t0}));
-
- auto c0 = IntImm::Int32(0);
- auto c1 = IntImm::Int32(1);
- auto c2 = IntImm::Int32(2);
-
- auto output = MapToNestedMsg<IntImm>(t1, [&](Expr value) {
- if (value.same_as(x)) return c0;
- if (value.same_as(y)) return c1;
- return c2;
- });
-
- NestedMsg<IntImm> expected = {{c0, c1}, c0, c2, {c0, c1}};
-
- EXPECT_TRUE(Equal(output, expected,
- [](IntImm lhs, IntImm rhs) -> bool { return lhs->value ==
rhs->value; }));
-
- auto output2 = MapToNestedMsg<IntImm>(GetType(t1), [&](Type ty) ->
NestedMsg<IntImm> {
- const auto* prim_ty = ty.as<PrimTypeNode>();
- if (prim_ty == nullptr) return std::nullopt;
- int bits = prim_ty->dtype.bits;
- if (bits == 16) return c0;
- if (bits == 32) return c1;
- if (bits == 64) return c2;
- return std::nullopt;
- });
-
- EXPECT_TRUE(Equal(output2, expected,
- [](IntImm lhs, IntImm rhs) -> bool { return lhs->value ==
rhs->value; }));
-
- int x_count = 0, y_count = 0, z_count = 0;
-
- DecomposeNestedMsg(t1, expected, [&](Expr value, NestedMsg<IntImm> msg) {
- if (value.same_as(x)) {
- EXPECT_TRUE(msg.LeafValue().same_as(c0));
- ++x_count;
- } else if (value.same_as(y)) {
- EXPECT_TRUE(msg.LeafValue().same_as(c1));
- ++y_count;
- } else {
- EXPECT_TRUE(msg.LeafValue().same_as(c2));
- ++z_count;
- }
- });
- EXPECT_EQ(x_count, 3);
- EXPECT_EQ(y_count, 2);
- EXPECT_EQ(z_count, 1);
-}
-
-TEST(NestedMsg, MapToNestedMsgByType) {
- auto sf0 = ScalarTensorType(PrimType::Float(32));
- auto sf1 = TupleType({sf0, sf0});
- auto sf2 = TupleType({sf0, sf0});
- auto x = tvm::Var("x", TupleType({sf1, sf2, sf0}));
-
- auto msg = MapToNestedMsgByType<Expr>(x, [](Expr value) { return value; });
-
- EXPECT_TRUE(msg.IsNested());
- auto arr = msg.NestedArray();
-
- EXPECT_TRUE(arr[1].IsNested());
- auto arr1 = arr[1].NestedArray();
-
- EXPECT_TRUE(arr1[0].IsLeaf());
- EXPECT_TRUE(
- tvm::ffi::StructuralEqual()(arr1[0].LeafValue(),
TupleGetItem(TupleGetItem(x, 1), 0)));
-
- EXPECT_TRUE(arr[2].IsLeaf());
- EXPECT_TRUE(tvm::ffi::StructuralEqual()(arr[2].LeafValue(), TupleGetItem(x,
2)));
-}
-
-TEST(NestedMsg, NestedMsgToExpr) {
- auto sf0 = ScalarTensorType(PrimType::Float(32));
- auto sf1 = TupleType({sf0, sf0});
-
- auto c0 = IntImm::Int32(0);
- auto c1 = IntImm::Int32(1);
- auto c2 = IntImm::Int32(2);
-
- tvm::Var x("x", sf0), y("y", sf0), z("z", sf0);
-
- NestedMsg<IntImm> msg = {c0, {c0, c1}, {c0, {c1, c2}}};
- auto expr = NestedMsgToExpr<IntImm>(msg, [&](ffi::Optional<IntImm> leaf) {
- TVM_FFI_ICHECK(leaf.has_value());
- int value = leaf.value()->value.as<int>().value();
- switch (value) {
- case 0:
- return x;
- case 1:
- return y;
- default:
- return z;
- }
- });
-
- Expr expected = Tuple({x, Tuple({x, y}), Tuple({x, Tuple({y, z})})});
- EXPECT_TRUE(tvm::ffi::StructuralEqual()(expr, expected));
-
- // test simplified
- tvm::Var t("t", sf1);
- NestedMsg<Expr> msg1 = {TupleGetItem(t, 0), TupleGetItem(t, 1)};
- auto expr1 = NestedMsgToExpr<Expr>(msg1, [](ffi::Optional<Expr> leaf) {
return leaf.value(); });
- EXPECT_TRUE(tvm::ffi::StructuralEqual()(expr1, t));
-}
-
-TEST(NestedMsg, CombineNestedMsg) {
- auto c0 = IntImm::Int32(0);
- auto c1 = IntImm::Int32(1);
- auto c2 = IntImm::Int32(2);
-
- NestedMsg<IntImm> lhs = {c0, {c0, c1}, std::nullopt, {c0, {c1, c2}}};
- NestedMsg<IntImm> rhs = {c1, {c2, std::nullopt}, std::nullopt, {c1, {c2,
c2}}};
- NestedMsg<IntImm> expected = {c1, {c2, c1}, std::nullopt, {c1, {c2, c2}}};
-
- auto output = CombineNestedMsg(lhs, rhs, [](IntImm x, IntImm y) {
- if (x->value > y->value) return x;
- return y;
- });
-
- EXPECT_TRUE(Equal(output, expected,
- [](IntImm lhs, IntImm rhs) -> bool { return lhs->value ==
rhs->value; }));
-}
-
-TEST(NestedMsg, MapNestedMsg) {
- auto c0 = IntImm::Int32(0);
- auto c1 = IntImm::Int32(1);
- auto c2 = IntImm::Int32(2);
- auto c3 = IntImm::Int32(3);
-
- NestedMsg<IntImm> msg = {c0, {c0, c1}, std::nullopt, {c0, {c2, c1}}};
- NestedMsg<IntImm> expected = {c3, {c3, std::nullopt}, std::nullopt, {c3,
{c2, std::nullopt}}};
-
- auto output = MapNestedMsg(msg, [](IntImm x) {
- if (x->value == 0) {
- return NestedMsg<IntImm>(IntImm::Int32(3));
- } else if (x->value == 1) {
- return NestedMsg<IntImm>();
- } else {
- return NestedMsg<IntImm>(x);
- }
- });
-
- EXPECT_TRUE(Equal(output, expected,
- [](IntImm lhs, IntImm rhs) -> bool { return lhs->value ==
rhs->value; }));
-}
-
-TEST(NestedMsg, TransformTupleLeaf) {
- auto c0 = IntImm::Int32(0);
- auto c1 = IntImm::Int32(1);
- auto c2 = IntImm::Int32(2);
- using NInt = NestedMsg<IntImm>;
-
- NInt msg1 = {c0, {c0, c1}, c2, {c0, {c1, c2}}};
- NInt msg2 = {c1, {c2, c0}, c2, {c1, {c2, c0}}};
-
- PrimType s = PrimType::Int(32);
- tvm::Var x("x", s), y("y", s), z("z", s);
- BlockBuilder bb = BlockBuilder::Create(std::nullopt);
- Expr expr = bb->Normalize(Tuple({x, Tuple({x, x}), x, Tuple({x, Tuple({x,
x})})}));
-
- auto ftransleaf = [&](Expr value, std::array<NInt, 2> msgs) -> Expr {
- int lhs =
msgs[0].LeafValue().as_or_throw<IntImm>()->value.as<int>().value();
- int rhs =
msgs[1].LeafValue().as_or_throw<IntImm>()->value.as<int>().value();
- if (lhs > rhs)
- return z;
- else if (lhs == rhs)
- return value;
- else
- return y;
- };
-
- Expr expected = Tuple({y, Tuple({y, z}), x, Tuple({y, Tuple({y, z})})});
-
- EXPECT_TRUE(tvm::ffi::StructuralEqual()(
- TransformTupleLeaf(expr, std::array<NInt, 2>({msg1, msg2}), ftransleaf),
expected));
-
- EXPECT_TRUE(
- expr.same_as(TransformTupleLeaf(expr, std::array<NInt, 2>({msg1, msg1}),
ftransleaf)));
-}
diff --git a/tests/cpp/object_protocol_test.cc
b/tests/cpp/object_protocol_test.cc
deleted file mode 100644
index 992ed0e784..0000000000
--- a/tests/cpp/object_protocol_test.cc
+++ /dev/null
@@ -1,85 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-#include <tvm/ffi/memory.h>
-#include <tvm/runtime/logging.h>
-
-namespace tvm {
-namespace test {
-
-using namespace tvm::runtime;
-
-class ObjBase : public ffi::Object {
- public:
- // dynamically allocate slow
- static constexpr const uint32_t _type_child_slots = 1;
- TVM_FFI_DECLARE_OBJECT_INFO("test.ObjBase", ObjBase, ffi::Object);
-};
-
-class ObjA : public ObjBase {
- public:
- static constexpr const uint32_t _type_child_slots = 0;
- TVM_FFI_DECLARE_OBJECT_INFO("test.ObjA", ObjA, ObjBase);
-};
-
-class ObjB : public ObjBase {
- public:
- static constexpr const uint32_t _type_child_slots = 0;
- TVM_FFI_DECLARE_OBJECT_INFO("test.ObjB", ObjB, ObjBase);
-};
-
-class ObjAA : public ObjA {
- public:
- TVM_FFI_DECLARE_OBJECT_INFO_FINAL("test.ObjAA", ObjAA, ObjA);
-};
-
-} // namespace test
-} // namespace tvm
-
-TEST(ObjectHierachy, Basic) {
- using namespace tvm;
- using namespace tvm::runtime;
- using namespace tvm::test;
- using namespace tvm::ffi;
-
- ffi::ObjectRef refA(ffi::make_object<ObjA>());
- TVM_FFI_ICHECK_EQ(refA->type_index(), ObjA::RuntimeTypeIndex());
- TVM_FFI_ICHECK(refA.as<ffi::Object>() != nullptr);
- TVM_FFI_ICHECK(refA.as<ObjA>() != nullptr);
- TVM_FFI_ICHECK(refA.as<ObjBase>() != nullptr);
- TVM_FFI_ICHECK(refA.as<ObjB>() == nullptr);
- TVM_FFI_ICHECK(refA.as<ObjAA>() == nullptr);
-
- ffi::ObjectRef refAA(ffi::make_object<ObjAA>());
- TVM_FFI_ICHECK_EQ(refAA->type_index(), ObjAA::RuntimeTypeIndex());
- TVM_FFI_ICHECK(refAA.as<ffi::Object>() != nullptr);
- TVM_FFI_ICHECK(refAA.as<ObjBase>() != nullptr);
- TVM_FFI_ICHECK(refAA.as<ObjA>() != nullptr);
- TVM_FFI_ICHECK(refAA.as<ObjAA>() != nullptr);
- TVM_FFI_ICHECK(refAA.as<ObjB>() == nullptr);
-
- ffi::ObjectRef refB(ffi::make_object<ObjB>());
- TVM_FFI_ICHECK_EQ(refB->type_index(), ObjB::RuntimeTypeIndex());
- TVM_FFI_ICHECK(refB.as<ffi::Object>() != nullptr);
- TVM_FFI_ICHECK(refB.as<ObjBase>() != nullptr);
- TVM_FFI_ICHECK(refB.as<ObjA>() == nullptr);
- TVM_FFI_ICHECK(refB.as<ObjAA>() == nullptr);
- TVM_FFI_ICHECK(refB.as<ObjB>() != nullptr);
-}
diff --git a/tests/cpp/random_engine_test.cc b/tests/cpp/random_engine_test.cc
deleted file mode 100644
index afd02c09b2..0000000000
--- a/tests/cpp/random_engine_test.cc
+++ /dev/null
@@ -1,65 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-#include <tvm/runtime/logging.h>
-#include <tvm/s_tir/random_engine.h>
-
-TEST(RandomEngine, Randomness) {
- int64_t rand_state = 0;
-
- tvm::s_tir::LinearCongruentialEngine rng(&rand_state);
- rng.Seed(0x114514);
-
- bool covered[100];
- memset(covered, 0, sizeof(covered));
- for (int i = 0; i < 100000; i++) {
- covered[rng() % 100] = true;
- }
- for (int i = 0; i < 100; i++) {
- TVM_FFI_ICHECK(covered[i]);
- }
-}
-
-TEST(RandomEngine, Reproducibility) {
- int64_t rand_state_a = 0, rand_state_b = 0;
- tvm::s_tir::LinearCongruentialEngine rng_a(&rand_state_a),
rng_b(&rand_state_b);
-
- rng_a.Seed(0x23456789);
- rng_b.Seed(0x23456789);
-
- for (int i = 0; i < 100000; i++) {
- TVM_FFI_ICHECK_EQ(rng_a(), rng_b());
- }
-}
-
-TEST(RandomEngine, Serialization) {
- int64_t rand_state_a = 0, rand_state_b = 0;
- tvm::s_tir::LinearCongruentialEngine rng_a(&rand_state_a),
rng_b(&rand_state_b);
-
- rng_a.Seed(0x56728);
-
- rand_state_b = rand_state_a;
- for (int i = 0; i < 100000; i++) TVM_FFI_ICHECK_EQ(rng_a(), rng_b());
-
- for (int i = 0; i < 123456; i++) rng_a();
-
- rand_state_b = rand_state_a;
- for (int i = 0; i < 100000; i++) TVM_FFI_ICHECK_EQ(rng_a(), rng_b());
-}
diff --git a/tests/cpp/runtime/logging_test.cc
b/tests/cpp/runtime/logging_test.cc
index ab52db5fd0..3df3034025 100644
--- a/tests/cpp/runtime/logging_test.cc
+++ b/tests/cpp/runtime/logging_test.cc
@@ -61,14 +61,17 @@ TEST(TvmLogDebugSettings, VLogEnabledComplex) {
EXPECT_FALSE(settings.VerboseEnabled("my/filesystem/src/baz.cc", 0));
}
-#define MATCH_THROW(stmt, err_type, matcher) \
- try { \
- stmt; \
- } catch (const err_type& e) { \
- EXPECT_THAT(e.what(), matcher); \
- } catch (...) { \
- EXPECT_FALSE("stmt threw an unexpected exception"); \
- }
+#define MATCH_THROW(stmt, err_type, matcher) \
+ do { \
+ try { \
+ stmt; \
+ ADD_FAILURE() << "Expected " #stmt " to throw " #err_type; \
+ } catch (const err_type& e) { \
+ EXPECT_THAT(e.what(), matcher); \
+ } catch (...) { \
+ ADD_FAILURE() << #stmt " threw an unexpected exception"; \
+ } \
+ } while (false)
TEST(TvmLogDebugSettings, IllFormed) {
MATCH_THROW(
diff --git a/tests/cpp/runtime/thread_storage_scope_test.cc
b/tests/cpp/runtime/thread_storage_scope_test.cc
deleted file mode 100644
index 50579cf286..0000000000
--- a/tests/cpp/runtime/thread_storage_scope_test.cc
+++ /dev/null
@@ -1,42 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include "../../../src/runtime/thread_storage_scope.h"
-
-#include <gtest/gtest.h>
-
-namespace tvm {
-namespace runtime {
-namespace {
-
-TEST(LaunchParamConfigTest, DerivesClusterLaunchFromTagPresence) {
- LaunchParamConfig config;
- config.Init(0, {"blockIdx.x", "threadIdx.x"});
- EXPECT_FALSE(config.use_cluster_launch());
-
- config.Init(0, {"blockIdx.x", "clusterCtaIdx.x", "threadIdx.x"});
- EXPECT_TRUE(config.use_cluster_launch());
-
- config.Init(0, {"blockIdx.x", "preferredClusterCtaIdx.x", "threadIdx.x"});
- EXPECT_FALSE(config.use_cluster_launch());
-}
-
-} // namespace
-} // namespace runtime
-} // namespace tvm
diff --git a/tests/cpp/support_test.cc b/tests/cpp/support/arena_test.cc
similarity index 81%
rename from tests/cpp/support_test.cc
rename to tests/cpp/support/arena_test.cc
index 243b3f848c..c597b9028e 100644
--- a/tests/cpp/support_test.cc
+++ b/tests/cpp/support/arena_test.cc
@@ -17,35 +17,15 @@
* under the License.
*/
+#include "../../../src/support/arena.h"
+
#include <gtest/gtest.h>
-#include <tvm/runtime/logging.h>
#include <memory>
-#include "../../src/support/arena.h"
-#include "../../src/support/utils.h"
-
namespace tvm {
namespace test {
-TEST(HashTests, HashStability) {
- size_t a = 345292;
- int b = 795620;
- EXPECT_EQ(::tvm::support::HashCombine(a, b), 2677237020);
- uint64_t c = 12345;
- int d = 987654432;
- EXPECT_EQ(::tvm::support::HashCombine(c, d), 3642871070);
- size_t e = 1010101;
- size_t f = 3030303;
- EXPECT_EQ(::tvm::support::HashCombine(e, f), 2722928432);
-}
-
-TEST(StartsWithTests, Basic) {
- EXPECT_TRUE(::tvm::support::StartsWith("abc", "abc"));
- EXPECT_TRUE(::tvm::support::StartsWith("abcd", "abc"));
- EXPECT_FALSE(::tvm::support::StartsWith("abc", "abcd"));
-}
-
namespace {
// A non-trivially-destructible type: destructing it has an observable
// side effect (incrementing a counter), unlike a plain-data struct.
diff --git a/tests/cpp/support/ring_buffer_test.cc
b/tests/cpp/support/ring_buffer_test.cc
deleted file mode 100644
index 43a921e0d1..0000000000
--- a/tests/cpp/support/ring_buffer_test.cc
+++ /dev/null
@@ -1,68 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include "../../../src/support/ring_buffer.h"
-
-#include <gtest/gtest.h>
-
-namespace tvm {
-namespace support {
-namespace {
-
-TEST(RingBuffer, ReadWrite) {
- RingBuffer buffer;
- std::vector<int> data = {1, 2, 3, 4};
- std::vector<int> output;
-
- buffer.Write(data.data(), data.size() * 4);
- ASSERT_EQ(buffer.bytes_available(), data.size() * 4);
-
- output.resize(4);
- buffer.Read(output.data(), data.size() * 4);
-
- for (size_t i = 0; i < output.size(); ++i) {
- ASSERT_EQ(output[i], data[i]);
- }
-}
-
-TEST(RingBuffer, ReadWithCallback) {
- RingBuffer buffer;
- std::vector<int> data = {1, 2, 3, 4};
- std::vector<int> output;
-
- buffer.Write(data.data(), data.size() * 4);
-
- auto callback0 = [](const char* data, size_t size) -> size_t {
- const int* iptr = reinterpret_cast<const int*>(data);
- TVM_FFI_ICHECK_EQ(iptr[0], 1);
- TVM_FFI_ICHECK_EQ(iptr[1], 2);
- return size;
- };
- buffer.ReadWithCallback(callback0, 2 * sizeof(int));
- auto callback1 = [](const char* data, size_t size) -> size_t {
- const int* iptr = reinterpret_cast<const int*>(data);
- TVM_FFI_ICHECK_EQ(iptr[0], 3);
- TVM_FFI_ICHECK_EQ(iptr[1], 4);
- return size;
- };
- buffer.ReadWithCallback(callback1, 2 * sizeof(int));
-}
-} // namespace
-} // namespace support
-} // namespace tvm
diff --git a/tests/cpp/pattern_match_test.cc
b/tests/cpp/sym/pattern_match_test.cc
similarity index 99%
rename from tests/cpp/pattern_match_test.cc
rename to tests/cpp/sym/pattern_match_test.cc
index cc3aa2f5f1..ee38efd76d 100644
--- a/tests/cpp/pattern_match_test.cc
+++ b/tests/cpp/sym/pattern_match_test.cc
@@ -17,7 +17,7 @@
* under the License.
*/
-#include "../src/sym/pattern_match.h"
+#include "../../../src/sym/pattern_match.h"
#include <gtest/gtest.h>
#include <tvm/ir/prim/expr.h>
diff --git a/tests/cpp/sym_simplify_test.cc b/tests/cpp/sym_simplify_test.cc
deleted file mode 100644
index 0e5ea4717c..0000000000
--- a/tests/cpp/sym_simplify_test.cc
+++ /dev/null
@@ -1,156 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-#include <tvm/ffi/extra/structural_equal.h>
-#include <tvm/runtime/logging.h>
-#include <tvm/sym/analyzer.h>
-#include <tvm/te/operation.h>
-#include <tvm/tirx/expr.h>
-
-TEST(Simplify, MinMax) {
- tvm::sym::Analyzer ana;
- auto x = tvm::te::var("x");
- auto e1 = (tvm::max(x, 1) - tvm::max(x, 1));
- auto e1s = ana->canonical_simplify(e1);
- TVM_FFI_ICHECK(tvm::prim::IsZero(e1s));
-
- auto e2 = (x * tvm::min(x, 1)) - (x * tvm::min(x, 1));
- auto e2s = ana->canonical_simplify(e2);
- TVM_FFI_ICHECK(tvm::prim::IsZero(e2s));
-}
-
-TEST(Simplify, Mul) {
- tvm::sym::Analyzer ana;
- auto x = tvm::te::var("x");
- auto e = (x * x) - (x * x);
- auto es = ana->canonical_simplify(e);
- TVM_FFI_ICHECK(tvm::prim::IsZero(es));
-}
-
-TEST(Simplify, Mod) {
- tvm::sym::Analyzer ana;
- auto x = tvm::IntImm::Int32(10);
- auto y = tvm::IntImm::Int32(12);
- // Mod::make is used instead of % to avoid constant folding during
- // calling operator%(x,y). Mod::make doesn't try constant folding,
- // and therefore, the constant folding will be attempted in CanonicalSimplify
- auto mod = ana->canonical_simplify(tvm::prim::Mod(x, y));
- auto es = ana->canonical_simplify(mod - x);
- TVM_FFI_ICHECK(tvm::prim::IsZero(es));
-}
-
-TEST(AnalyzerObjectRef, CopySharesMutableState) {
- tvm::sym::Analyzer analyzer;
- tvm::sym::Analyzer copy = analyzer;
- auto x = tvm::te::var("x");
-
- copy->Bind(x, tvm::Range::FromMinExtent(0, 8));
-
- TVM_FFI_ICHECK(analyzer->CanProve(x < 8));
-}
-
-TEST(AnalyzerObjectRef, ConstHandleRefCanMutateAnalyzerState) {
- tvm::sym::Analyzer analyzer;
- const tvm::sym::Analyzer& analyzer_ref = analyzer;
- auto x = tvm::te::var("x");
-
- analyzer_ref->Bind(x, tvm::Range::FromMinExtent(0, 8));
-
- TVM_FFI_ICHECK(analyzer->CanProve(x < 8));
-}
-
-TEST(AnalyzerObjectRef, CloneIsIndependent) {
- tvm::sym::Analyzer analyzer;
- auto x = tvm::te::var("x");
- auto y = tvm::te::var("y");
-
- analyzer->Bind(x, tvm::Range::FromMinExtent(0, 8));
- analyzer->modular_set.Update(x, tvm::sym::ModularSet(4, 0));
-
- tvm::sym::Analyzer clone = analyzer->Clone();
- TVM_FFI_ICHECK(clone->CanProve(x < 8));
- TVM_FFI_ICHECK(clone->modular_set(x)->coeff == 4);
-
- clone->Bind(y, tvm::Range::FromMinExtent(0, 4));
- clone->modular_set.Update(x, tvm::sym::ModularSet(8, 0), true);
- TVM_FFI_ICHECK(clone->CanProve(y < 4));
- TVM_FFI_ICHECK(!analyzer->CanProve(y < 4));
- TVM_FFI_ICHECK(analyzer->CanProve(x < 8));
- TVM_FFI_ICHECK(analyzer->modular_set(x)->coeff == 4);
- TVM_FFI_ICHECK(clone->modular_set(x)->coeff == 8);
-}
-
-TEST(Simplify, AssumeConstraintKeepsTensorLoadStable) {
- using namespace tvm;
-
- sym::Analyzer analyzer;
- tirx::TensorVar buffer = tirx::decl_tensor({1}, PrimType::Int(32));
- PrimExpr load = tirx::MakeTensorLoad(buffer, {IntImm::Int32(0)});
- PrimExpr constraint = load > 0;
-
- {
- auto exit_scope = analyzer->rewrite_simplify.EnterConstraint(constraint);
- EXPECT_FALSE(tvm::prim::IsOne(analyzer->rewrite_simplify(constraint)));
- exit_scope();
- }
- {
- auto exit_scope = analyzer->rewrite_simplify.EnterConstraint(constraint,
true);
- EXPECT_TRUE(tvm::prim::IsOne(analyzer->rewrite_simplify(constraint)));
- exit_scope();
- }
-
- {
- With<sym::ConstraintContext> scope(analyzer, constraint, true);
- if (analyzer->z3_prover.IsEnabled()) {
- EXPECT_TRUE(analyzer->z3_prover.CanProve(constraint));
- }
- }
-
- if (analyzer->z3_prover.IsEnabled()) {
- EXPECT_FALSE(analyzer->z3_prover.CanProve(constraint));
- }
-
- {
- With<sym::ConstraintContext> scope(analyzer, constraint);
- if (analyzer->z3_prover.IsEnabled()) {
- EXPECT_FALSE(analyzer->z3_prover.CanProve(constraint));
- }
- }
-}
-
-TEST(ConstantFold, Broadcast) {
- tvm::ffi::StructuralEqual checker;
- auto i32x4 = tvm::prim::Broadcast(tvm::IntImm::Int32(10), 4);
- auto i64x4 = tvm::prim::cast(i32x4.ty().WithBits(64), i32x4);
- auto i64x4_expected = tvm::prim::Broadcast(tvm::IntImm::Int64(10), 4);
- ASSERT_TRUE(checker(i64x4, i64x4_expected));
-}
-
-TEST(ConstantFold, Ramp) {
- tvm::ffi::StructuralEqual checker;
- auto i32x4 = tvm::prim::Ramp(tvm::IntImm::Int32(10), tvm::IntImm::Int32(1),
4);
- auto i64x4 = tvm::prim::cast(i32x4.ty().WithBits(64), i32x4);
- auto i64x4_expected = tvm::prim::Ramp(tvm::IntImm::Int64(10),
tvm::IntImm::Int64(1), 4);
- ASSERT_TRUE(checker(i64x4, i64x4_expected));
-
- auto f32x4 = tvm::prim::cast(tvm::PrimType::Float(32, 4), i32x4);
- auto f32x4_expected = tvm::prim::Cast(tvm::PrimType::Float(32, 4), i32x4);
- ASSERT_TRUE(checker(f32x4, f32x4_expected));
-}
diff --git a/tests/cpp/config_schema_test.cc
b/tests/cpp/target/config_schema_test.cc
similarity index 91%
rename from tests/cpp/config_schema_test.cc
rename to tests/cpp/target/config_schema_test.cc
index 8d80099d96..827ce218f9 100644
--- a/tests/cpp/config_schema_test.cc
+++ b/tests/cpp/target/config_schema_test.cc
@@ -45,16 +45,6 @@ TEST(ConfigSchema, BasicTypes) {
ASSERT_EQ(resolved.at("name").cast<ffi::String>(), "test");
}
-// Type coercion: correct type passes
-TEST(ConfigSchema, TypeCoercionPasses) {
- ConfigSchema schema;
- schema.def_option<int64_t>("count");
-
- ffi::Map<ffi::String, ffi::Any> config = {{"count", int64_t(42)}};
- auto resolved = schema.Resolve(config);
- ASSERT_EQ(resolved.at("count").cast<int64_t>(), 42);
-}
-
// Type mismatch rejection (TypeError)
TEST(ConfigSchema, TypeMismatchRejection) {
ConfigSchema schema;
@@ -99,16 +89,6 @@ TEST(ConfigSchema, DefaultValueMaterialization) {
ASSERT_EQ(resolved.at("debug").cast<bool>(), false);
}
-// Default value materialization via refl::DefaultValue
-TEST(ConfigSchema, DefaultValueViaTrait) {
- ConfigSchema schema;
- schema.def_option<int64_t>("max_threads", refl::DefaultValue(int64_t(128)));
-
- ffi::Map<ffi::String, ffi::Any> config = {};
- auto resolved = schema.Resolve(config);
- ASSERT_EQ(resolved.at("max_threads").cast<int64_t>(), 128);
-}
-
// Explicit value overrides default
TEST(ConfigSchema, ExplicitOverridesDefault) {
ConfigSchema schema;
@@ -196,14 +176,16 @@ TEST(ConfigSchema, DeclarationOrderPreserved) {
ASSERT_EQ(std::string(options[2].key), "middle");
}
-// Doc string trait accepted without affecting behavior
+// Documentation traits work with explicit values and defaults
TEST(ConfigSchema, DocStringTrait) {
ConfigSchema schema;
schema.def_option<int64_t>("count", "Number of items to process");
+ schema.def_option<int64_t>("num_cores", "CPU core count",
refl::DefaultValue(int64_t(4)));
ffi::Map<ffi::String, ffi::Any> config = {{"count", int64_t(5)}};
auto resolved = schema.Resolve(config);
ASSERT_EQ(resolved.at("count").cast<int64_t>(), 5);
+ ASSERT_EQ(resolved.at("num_cores").cast<int64_t>(), 4);
}
// Array option type
@@ -310,13 +292,3 @@ TEST(ConfigSchema, CanonicalizerWithDefaults) {
auto resolved = schema.Resolve(config);
ASSERT_EQ(resolved.at("threads").cast<int64_t>(), 8);
}
-
-// Combined doc + default traits
-TEST(ConfigSchema, DocAndDefaultTraits) {
- ConfigSchema schema;
- schema.def_option<int64_t>("num_cores", "CPU core count",
refl::DefaultValue(int64_t(4)));
-
- ffi::Map<ffi::String, ffi::Any> config = {};
- auto resolved = schema.Resolve(config);
- ASSERT_EQ(resolved.at("num_cores").cast<int64_t>(), 4);
-}
diff --git a/tests/cpp/target/target_test.cc b/tests/cpp/target/target_test.cc
new file mode 100644
index 0000000000..773ddc311e
--- /dev/null
+++ b/tests/cpp/target/target_test.cc
@@ -0,0 +1,141 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include <gtest/gtest.h>
+#include <tvm/runtime/logging.h>
+#include <tvm/target/target.h>
+
+using namespace tvm;
+
+TVM_FFI_STATIC_INIT_BLOCK() {
+
TargetKindDef("TestTargetKind").set_default_device_type(kDLCPU).def_option<bool>("my_bool");
+}
+
+ffi::Map<ffi::String, ffi::Any> TestTargetParser(ffi::Map<ffi::String,
ffi::Any> target) {
+ ffi::String mcpu = target.at("mcpu").as_or_throw<ffi::String>();
+ target.Set("mcpu", ffi::String("super_") + mcpu);
+ target.Set("keys", ffi::Array<ffi::String>({"super"}));
+ target.Set("feature.test", true);
+ return target;
+}
+
+TVM_FFI_STATIC_INIT_BLOCK() {
+ TargetKindDef("TestTargetParser")
+ .set_default_device_type(kDLCPU)
+ .def_option<ffi::String>("mattr")
+ .def_option<ffi::String>("mcpu")
+ .set_default_keys({"cpu"})
+ .set_target_canonicalizer(TestTargetParser);
+}
+
+TEST(TargetCreationFail, TargetKindNotFound) {
+ ffi::Map<ffi::String, ffi::Any> config = {};
+ ASSERT_THROW({ Target{config}; }, ffi::Error);
+}
+
+TEST(TargetCreation, TargetParser) {
+ Target test_target(ffi::Map<ffi::String, ffi::Any>{
+ {"kind", ffi::String("TestTargetParser")},
+ {"mcpu", ffi::String("woof")},
+ {"mattr", ffi::String("cake")},
+ });
+ ASSERT_EQ(test_target->GetAttr<ffi::String>("mcpu").value(), "super_woof");
+ ASSERT_EQ(test_target->keys.size(), 1);
+ ASSERT_EQ(test_target->keys[0], "super");
+ ASSERT_EQ(test_target->GetAttr<ffi::String>("mattr").value(), "cake");
+ ASSERT_EQ(test_target->GetAttr<bool>("feature.test").value(), true);
+ ASSERT_EQ(test_target->GetAttr<bool>("feature.nonexistent"), std::nullopt);
+
+ Target test_target_no_parser("TestTargetKind");
+ ASSERT_EQ(test_target_no_parser->GetAttr<bool>("feature.test"),
std::nullopt);
+ ASSERT_EQ(test_target_no_parser->GetAttr<bool>("feature.test",
true).value(), true);
+}
+
+TEST(TargetCreation, RoundTripCanonicalizerFeatures) {
+ // Construct a target whose canonicalizer sets feature.test and transforms
mcpu
+ Target original(ffi::Map<ffi::String, ffi::Any>{
+ {"kind", ffi::String("TestTargetParser")},
+ {"mcpu", ffi::String("woof")},
+ });
+ ASSERT_EQ(original->GetAttr<ffi::String>("mcpu").value(), "super_woof");
+ ASSERT_EQ(original->GetAttr<bool>("feature.test").value(), true);
+
+ // Export to config and reconstruct
+ ffi::Map<ffi::String, ffi::Any> exported = original->ToConfig();
+ Target reconstructed(exported);
+
+ // Canonicalized attrs must survive the round-trip
+ // Note: mcpu gets canonicalized again (super_super_woof) because the
canonicalizer runs
+ ASSERT_TRUE(reconstructed->GetAttr<ffi::String>("mcpu").has_value());
+ ASSERT_EQ(reconstructed->GetAttr<bool>("feature.test").value(), true);
+ ASSERT_EQ(reconstructed->keys.size(), 1);
+ ASSERT_EQ(reconstructed->keys[0], "super");
+}
+
+TEST(TargetCreation, RoundTripCanonicalizerFeaturesNestedHost) {
+ // Construct a host target whose canonicalizer sets feature.test
+ Target host(ffi::Map<ffi::String, ffi::Any>{
+ {"kind", ffi::String("TestTargetParser")},
+ {"mcpu", ffi::String("woof")},
+ });
+ ASSERT_EQ(host->GetAttr<bool>("feature.test").value(), true);
+
+ // Attach it as host to another target
+ Target outer(ffi::Map<ffi::String, ffi::Any>{
+ {"kind", ffi::String("TestTargetKind")},
+ {"my_bool", true},
+ });
+ Target combined(outer, host);
+
+ // Export the outer target (includes nested host) and reconstruct
+ ffi::Map<ffi::String, ffi::Any> exported = combined->ToConfig();
+ Target reconstructed(exported);
+
+ // The nested host must reconstruct successfully with feature.* preserved
+ ffi::Optional<Target> reconstructed_host = reconstructed->GetHost();
+ ASSERT_TRUE(reconstructed_host.has_value());
+ ASSERT_EQ(reconstructed_host.value()->GetAttr<bool>("feature.test").value(),
true);
+
ASSERT_TRUE(reconstructed_host.value()->GetAttr<ffi::String>("mcpu").has_value());
+}
+
+TEST(TargetCreationFail, UnknownNonFeatureKeyStillFails) {
+ // Verify that unknown non-feature.* keys still fail schema validation
+ ffi::Map<ffi::String, ffi::Any> config = {
+ {"kind", ffi::String("TestTargetParser")},
+ {"mcpu", ffi::String("woof")},
+ {"unknown_key", ffi::String("bad")},
+ };
+ ASSERT_THROW({ Target{config}; }, tvm::ffi::Error);
+}
+
+TEST(TargetCreation, DeduplicateKeys) {
+ ffi::Map<ffi::String, ffi::Any> config = {
+ {"kind", ffi::String("llvm")},
+ {"keys", ffi::Array<ffi::String>{"cpu", "arm_cpu"}},
+ {"device", ffi::String("arm_cpu")},
+ };
+ Target target = Target(config);
+ TVM_FFI_ICHECK_EQ(target->kind, TargetKind::Get("llvm").value());
+ TVM_FFI_ICHECK_EQ(target->tag, "");
+ TVM_FFI_ICHECK_EQ(target->keys.size(), 2U);
+ TVM_FFI_ICHECK_EQ(target->keys[0], "cpu");
+ TVM_FFI_ICHECK_EQ(target->keys[1], "arm_cpu");
+ TVM_FFI_ICHECK_EQ(target->attrs.count("keys"), 0U);
+ TVM_FFI_ICHECK_EQ(target->GetAttr<ffi::String>("device"), "arm_cpu");
+}
diff --git a/tests/cpp/target_test.cc b/tests/cpp/target_test.cc
deleted file mode 100644
index 0bf728425f..0000000000
--- a/tests/cpp/target_test.cc
+++ /dev/null
@@ -1,400 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-#include <tvm/ir/expr.h>
-#include <tvm/runtime/logging.h>
-#include <tvm/target/target.h>
-
-#include <cmath>
-#include <string>
-
-using namespace tvm;
-
-TVM_FFI_STATIC_INIT_BLOCK() {
- TargetKindDef("TestTargetKind")
- .set_default_device_type(kDLCPU)
- .def_option<bool>("my_bool")
- .def_option<ffi::Array<ffi::String>>("your_names")
- .def_option<ffi::Map<ffi::String, int64_t>>("her_maps");
-}
-
-ffi::Map<ffi::String, ffi::Any> TestTargetParser(ffi::Map<ffi::String,
ffi::Any> target) {
- ffi::String mcpu = target.at("mcpu").as_or_throw<ffi::String>();
- target.Set("mcpu", ffi::String("super_") + mcpu);
- target.Set("keys", ffi::Array<ffi::String>({"super"}));
- target.Set("feature.test", true);
- return target;
-}
-
-ffi::Map<ffi::String, ffi::Any> TestAttrsPreProcessor(ffi::Map<ffi::String,
ffi::Any> attrs) {
- attrs.Set("mattr", ffi::String("woof"));
- return attrs;
-}
-
-TVM_FFI_STATIC_INIT_BLOCK() {
- TargetKindDef("TestTargetParser")
- .set_default_device_type(kDLCPU)
- .def_option<ffi::String>("mattr")
- .def_option<ffi::String>("mcpu")
- .set_default_keys({"cpu"})
- .set_target_canonicalizer(TestTargetParser);
-
- TargetKindDef("TestAttrsPreprocessor")
- .set_default_device_type(kDLCPU)
- .def_option<ffi::String>("mattr")
- .set_default_keys({"cpu"})
- .set_target_canonicalizer(TestAttrsPreProcessor);
-
- TargetKindDef("TestClashingPreprocessor")
- .set_default_device_type(kDLCPU)
- .def_option<ffi::String>("mattr")
- .def_option<ffi::String>("mcpu")
- .set_default_keys({"cpu"})
- .set_target_canonicalizer(TestTargetParser);
-}
-
-TEST(TargetCreation, NestedConfig) {
- ffi::Map<ffi::String, ffi::Any> config = {
- {"my_bool", true},
- {"your_names", ffi::Array<ffi::String>{"junru", "jian"}},
- {"kind", ffi::String("TestTargetKind")},
- {
- "her_maps",
- ffi::Map<ffi::String, int64_t>{
- {"a", 1},
- {"b", 2},
- },
- },
- };
- Target target = Target(config);
- TVM_FFI_ICHECK_EQ(target->kind, TargetKind::Get("TestTargetKind").value());
- TVM_FFI_ICHECK_EQ(target->tag, "");
- TVM_FFI_ICHECK(target->keys.empty());
- bool my_bool = target->GetAttr<bool>("my_bool").value();
- TVM_FFI_ICHECK_EQ(my_bool, true);
- ffi::Array<ffi::String> your_names =
- target->GetAttr<ffi::Array<ffi::String>>("your_names").value();
- TVM_FFI_ICHECK_EQ(your_names.size(), 2U);
- TVM_FFI_ICHECK_EQ(your_names[0], "junru");
- TVM_FFI_ICHECK_EQ(your_names[1], "jian");
- ffi::Map<ffi::String, int64_t> her_maps =
- target->GetAttr<ffi::Map<ffi::String, int64_t>>("her_maps").value();
- TVM_FFI_ICHECK_EQ(her_maps.size(), 2U);
- TVM_FFI_ICHECK_EQ(her_maps["a"], 1);
- TVM_FFI_ICHECK_EQ(her_maps["b"], 2);
-}
-
-TEST(TargetCreationFail, UnrecognizedConfigOption) {
- ffi::Map<ffi::String, ffi::Any> config = {
- {"my_bool", true},
- {"your_names", ffi::Array<ffi::String>{"junru", "jian"}},
- {"kind", ffi::String("TestTargetKind")},
- {"bad", ffi::ObjectRef(nullptr)},
- {
- "her_maps",
- ffi::Map<ffi::String, int64_t>{
- {"a", 1},
- {"b", 2},
- },
- },
- };
- bool failed = false;
- try {
- Target tgt(config);
- } catch (...) {
- failed = true;
- }
- ASSERT_EQ(failed, true);
-}
-
-TEST(TargetCreationFail, TypeMismatch) {
- ffi::Map<ffi::String, ffi::Any> config = {
- {"my_bool", ffi::String("true")},
- {"your_names", ffi::Array<ffi::String>{"junru", "jian"}},
- {"kind", ffi::String("TestTargetKind")},
- {
- "her_maps",
- ffi::Map<ffi::String, int64_t>{
- {"a", 1},
- {"b", 2},
- },
- },
- };
- bool failed = false;
- try {
- Target tgt(config);
- } catch (...) {
- failed = true;
- }
- ASSERT_EQ(failed, true);
-}
-
-TEST(TargetCreationFail, TargetKindNotFound) {
- ffi::Map<ffi::String, ffi::Any> config = {
- {"my_bool", "true"},
- {"your_names", ffi::Array<ffi::String>{"junru", "jian"}},
- {
- "her_maps",
- ffi::Map<ffi::String, int64_t>{
- {"a", 1},
- {"b", 2},
- },
- },
- };
- bool failed = false;
- try {
- Target tgt(config);
- } catch (...) {
- failed = true;
- }
- ASSERT_EQ(failed, true);
-}
-
-TEST(TargetCreation, TargetParser) {
- Target test_target(ffi::Map<ffi::String, ffi::Any>{
- {"kind", ffi::String("TestTargetParser")},
- {"mcpu", ffi::String("woof")},
- });
- ASSERT_EQ(test_target->GetAttr<ffi::String>("mcpu").value(), "super_woof");
- ASSERT_EQ(test_target->keys.size(), 1);
- ASSERT_EQ(test_target->keys[0], "super");
-}
-
-TEST(TargetCreation, TargetFeatures) {
- Target test_target_with_parser(ffi::Map<ffi::String, ffi::Any>{
- {"kind", ffi::String("TestTargetParser")},
- {"mcpu", ffi::String("woof")},
- });
- // Features are stored as "feature.xxx" keys in attrs
- ASSERT_EQ(test_target_with_parser->GetAttr<bool>("feature.test").value(),
true);
-
- Target test_target_no_parser("TestTargetKind");
- ASSERT_EQ(test_target_no_parser->GetAttr<bool>("feature.test"),
std::nullopt);
- ASSERT_EQ(test_target_no_parser->GetAttr<bool>("feature.test",
true).value(), true);
-}
-
-TEST(TargetCreation, TargetFeaturesSetByCanonicalizer) {
- // feature.* keys are set by the canonicalizer, not by user input.
- Target test_target(ffi::Map<ffi::String, ffi::Any>{
- {"kind", ffi::String("TestTargetParser")},
- {"mcpu", ffi::String("woof")},
- });
- // TestTargetParser sets "feature.test" = true
- ASSERT_EQ(test_target->GetAttr<bool>("feature.test").value(), true);
- // Non-existent features return nullopt
- ASSERT_EQ(test_target->GetAttr<bool>("feature.nonexistent"), std::nullopt);
-}
-
-TEST(TargetCreation, TargetAttrsPreProcessor) {
- Target test_target(ffi::Map<ffi::String, ffi::Any>{
- {"kind", ffi::String("TestAttrsPreprocessor")},
- {"mattr", ffi::String("cake")},
- });
- ASSERT_EQ(test_target->GetAttr<ffi::String>("mattr").value(), "woof");
-}
-
-TEST(TargetCreation, TargetParserProcessing) {
- Target test_target(ffi::Map<ffi::String, ffi::Any>{
- {"kind", ffi::String("TestClashingPreprocessor")},
- {"mcpu", ffi::String("woof")},
- {"mattr", ffi::String("cake")},
- });
- ASSERT_EQ(test_target->GetAttr<ffi::String>("mcpu").value(), "super_woof");
- ASSERT_EQ(test_target->GetAttr<ffi::String>("mattr").value(), "cake");
-}
-
-TEST(TargetCreation, RoundTripCanonicalizerFeatures) {
- // Construct a target whose canonicalizer sets feature.test and transforms
mcpu
- Target original(ffi::Map<ffi::String, ffi::Any>{
- {"kind", ffi::String("TestTargetParser")},
- {"mcpu", ffi::String("woof")},
- });
- ASSERT_EQ(original->GetAttr<ffi::String>("mcpu").value(), "super_woof");
- ASSERT_EQ(original->GetAttr<bool>("feature.test").value(), true);
-
- // Export to config and reconstruct
- ffi::Map<ffi::String, ffi::Any> exported = original->ToConfig();
- Target reconstructed(exported);
-
- // Canonicalized attrs must survive the round-trip
- // Note: mcpu gets canonicalized again (super_super_woof) because the
canonicalizer runs
- ASSERT_TRUE(reconstructed->GetAttr<ffi::String>("mcpu").has_value());
- ASSERT_EQ(reconstructed->GetAttr<bool>("feature.test").value(), true);
- ASSERT_EQ(reconstructed->keys.size(), 1);
- ASSERT_EQ(reconstructed->keys[0], "super");
-}
-
-TEST(TargetCreation, RoundTripCanonicalizerFeaturesNestedHost) {
- // Construct a host target whose canonicalizer sets feature.test
- Target host(ffi::Map<ffi::String, ffi::Any>{
- {"kind", ffi::String("TestTargetParser")},
- {"mcpu", ffi::String("woof")},
- });
- ASSERT_EQ(host->GetAttr<bool>("feature.test").value(), true);
-
- // Attach it as host to another target
- Target outer(ffi::Map<ffi::String, ffi::Any>{
- {"kind", ffi::String("TestTargetKind")},
- {"my_bool", true},
- });
- Target combined(outer, host);
-
- // Export the outer target (includes nested host) and reconstruct
- ffi::Map<ffi::String, ffi::Any> exported = combined->ToConfig();
- Target reconstructed(exported);
-
- // The nested host must reconstruct successfully with feature.* preserved
- ffi::Optional<Target> reconstructed_host = reconstructed->GetHost();
- ASSERT_TRUE(reconstructed_host.has_value());
- ASSERT_EQ(reconstructed_host.value()->GetAttr<bool>("feature.test").value(),
true);
-
ASSERT_TRUE(reconstructed_host.value()->GetAttr<ffi::String>("mcpu").has_value());
-}
-
-TEST(TargetCreationFail, UnknownNonFeatureKeyStillFails) {
- // Verify that unknown non-feature.* keys still fail schema validation
- ffi::Map<ffi::String, ffi::Any> config = {
- {"kind", ffi::String("TestTargetParser")},
- {"mcpu", ffi::String("woof")},
- {"unknown_key", ffi::String("bad")},
- };
- ASSERT_THROW({ Target{config}; }, tvm::ffi::Error);
-}
-
-TVM_FFI_STATIC_INIT_BLOCK() {
- TargetKindDef("TestStringKind")
- .set_default_device_type(kDLCPU)
- .def_option<ffi::String>("single")
- .def_option<ffi::Array<ffi::String>>("array")
- .def_option<ffi::Array<ffi::Array<ffi::String>>>("nested-array")
-
.def_option<ffi::Array<ffi::Array<ffi::Array<ffi::String>>>>("nested2-array");
-}
-
-TEST(TargetCreation, ProcessStrings) {
- // Test single string attribute
- Target test_target1(ffi::Map<ffi::String, ffi::Any>{
- {"kind", ffi::String("TestStringKind")},
- {"single", ffi::String("'string with single quote")},
- });
- ASSERT_TRUE(test_target1->GetAttr<ffi::String>("single"));
- ffi::String string1 = test_target1->GetAttr<ffi::String>("single").value();
- ASSERT_EQ(string1, "'string with single quote");
-
- // Test array of strings
- Target test_target3(ffi::Map<ffi::String, ffi::Any>{
- {"kind", ffi::String("TestStringKind")},
- {"array", ffi::Array<ffi::String>{"-danny", "-sammy=1", "-kirby=string
with space"}},
- });
- ASSERT_TRUE(test_target3->GetAttr<ffi::Array<ffi::String>>("array"));
- ffi::Array<ffi::String> array3 =
test_target3->GetAttr<ffi::Array<ffi::String>>("array").value();
- ASSERT_EQ(array3[0], "-danny");
- ASSERT_EQ(array3[1], "-sammy=1");
- ASSERT_EQ(array3[2], "-kirby=string with space");
-
- // Test nested array of strings
- Target test_target6(ffi::Map<ffi::String, ffi::Any>{
- {"kind", ffi::String("TestStringKind")},
- {"nested-array",
- ffi::Array<ffi::Array<ffi::String>>{
- ffi::Array<ffi::String>{"foo0", "foo1", "foo2"},
- ffi::Array<ffi::String>{"bar0", "bar1", "bar2"},
- ffi::Array<ffi::String>{"baz0", "baz1"},
- }},
- });
-
ASSERT_TRUE(test_target6->GetAttr<ffi::Array<ffi::Array<ffi::String>>>("nested-array"));
- ffi::Array<ffi::Array<ffi::String>> array6 =
-
test_target6->GetAttr<ffi::Array<ffi::Array<ffi::String>>>("nested-array").value();
- ASSERT_EQ(array6[0][0], "foo0");
- ASSERT_EQ(array6[0][1], "foo1");
- ASSERT_EQ(array6[0][2], "foo2");
- ASSERT_EQ(array6[1][0], "bar0");
- ASSERT_EQ(array6[1][1], "bar1");
- ASSERT_EQ(array6[1][2], "bar2");
- ASSERT_EQ(array6[2][0], "baz0");
- ASSERT_EQ(array6[2][1], "baz1");
-
- // Test doubly-nested array of strings
- Target test_target7(ffi::Map<ffi::String, ffi::Any>{
- {"kind", ffi::String("TestStringKind")},
- {"nested2-array",
- ffi::Array<ffi::Array<ffi::Array<ffi::String>>>{
- ffi::Array<ffi::Array<ffi::String>>{
- ffi::Array<ffi::String>{"foo0", "foo1"},
- ffi::Array<ffi::String>{"bar0", "bar1"},
- ffi::Array<ffi::String>{"baz0", "baz1"},
- },
- ffi::Array<ffi::Array<ffi::String>>{
- ffi::Array<ffi::String>{"zing0", "zing1"},
- ffi::Array<ffi::String>{"fred"},
- },
- }},
- });
- ASSERT_TRUE(
-
test_target7->GetAttr<ffi::Array<ffi::Array<ffi::Array<ffi::String>>>>("nested2-array"));
- ffi::Array<ffi::Array<ffi::Array<ffi::String>>> array7 =
-
test_target7->GetAttr<ffi::Array<ffi::Array<ffi::Array<ffi::String>>>>("nested2-array")
- .value();
- ASSERT_EQ(array7.size(), 2);
- ASSERT_EQ(array7[0].size(), 3);
- ASSERT_EQ(array7[0][0].size(), 2);
- ASSERT_EQ(array7[0][1].size(), 2);
- ASSERT_EQ(array7[0][2].size(), 2);
- ASSERT_EQ(array7[1].size(), 2);
- ASSERT_EQ(array7[1][0].size(), 2);
- ASSERT_EQ(array7[1][1].size(), 1);
-
- ASSERT_EQ(array7[0][0][0], "foo0");
- ASSERT_EQ(array7[0][0][1], "foo1");
- ASSERT_EQ(array7[0][1][0], "bar0");
- ASSERT_EQ(array7[0][1][1], "bar1");
- ASSERT_EQ(array7[0][2][0], "baz0");
- ASSERT_EQ(array7[0][2][1], "baz1");
- ASSERT_EQ(array7[1][0][0], "zing0");
- ASSERT_EQ(array7[1][0][1], "zing1");
- ASSERT_EQ(array7[1][1][0], "fred");
-}
-
-TEST(TargetCreation, DeduplicateKeys) {
- ffi::Map<ffi::String, ffi::Any> config = {
- {"kind", ffi::String("llvm")},
- {"keys", ffi::Array<ffi::String>{"cpu", "arm_cpu"}},
- {"device", ffi::String("arm_cpu")},
- };
- Target target = Target(config);
- TVM_FFI_ICHECK_EQ(target->kind, TargetKind::Get("llvm").value());
- TVM_FFI_ICHECK_EQ(target->tag, "");
- TVM_FFI_ICHECK_EQ(target->keys.size(), 2U);
- TVM_FFI_ICHECK_EQ(target->keys[0], "cpu");
- TVM_FFI_ICHECK_EQ(target->keys[1], "arm_cpu");
- TVM_FFI_ICHECK_EQ(target->attrs.count("keys"), 0U);
- TVM_FFI_ICHECK_EQ(target->GetAttr<ffi::String>("device"), "arm_cpu");
-}
-
-TEST(TargetKindRegistry, ListTargetKinds) {
- ffi::Array<ffi::String> names =
TargetKindRegistry::Global()->ListTargetKinds();
- TVM_FFI_ICHECK_EQ(names.empty(), false);
- TVM_FFI_ICHECK_EQ(std::count(std::begin(names), std::end(names), "llvm"), 1);
-}
-
-TEST(TargetKindRegistry, ListTargetOptions) {
- TargetKind llvm = TargetKind::Get("llvm").value();
- ffi::Map<ffi::String, ffi::String> attrs =
- TargetKindRegistry::Global()->ListTargetKindOptions(llvm);
- TVM_FFI_ICHECK_EQ(attrs.empty(), false);
-}
diff --git a/tests/cpp/te_compute_test.cc b/tests/cpp/te_compute_test.cc
deleted file mode 100644
index b4cde7bc58..0000000000
--- a/tests/cpp/te_compute_test.cc
+++ /dev/null
@@ -1,59 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-#include <tvm/runtime/logging.h>
-#include <tvm/te/operation.h>
-
-TEST(Tensor, Basic) {
- using namespace tvm;
- using namespace tvm::te;
-
- PrimVar m("m"), n("n"), l("l");
-
- Tensor A = placeholder({m, l}, PrimType::Float(32), "A");
- Tensor B = placeholder({n, l}, PrimType::Float(32), "B");
-
- auto C = compute({m, n}, [&](PrimVar i, PrimVar j) { return A[i][j]; }, "C");
-
- Tensor::Slice x = A[n];
-}
-
-TEST(Tensor, Reduce) {
- using namespace tvm::prim;
- using namespace tvm;
- using namespace tvm::te;
-
- PrimVar m("m"), n("n"), l("l");
- te::Tensor A = te::placeholder({m, l}, PrimType::Float(32), "A");
- te::Tensor B = te::placeholder({n, l}, PrimType::Float(32), "B");
- IterVar rv = reduce_axis(Range{0, l}, "k");
-
- auto C = te::compute(
- {m, n}, [&](PrimVar i, PrimVar j) { return sum(max(1 + A[i][rv] + 1,
B[j][rv]), {rv}); },
- "C");
-}
-
-TEST(Tensor, Indexing) {
- using namespace tvm;
- using namespace tvm::te;
-
- PrimVar x("x"), y("y");
- te::Tensor A = te::placeholder({x, y}, PrimType::Float(32), "A");
-}
diff --git a/tests/cpp/threading_backend_test.cc
b/tests/cpp/threading_backend_test.cc
deleted file mode 100644
index e68e0fdba8..0000000000
--- a/tests/cpp/threading_backend_test.cc
+++ /dev/null
@@ -1,196 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include "../../src/runtime/threading_backend.h"
-
-#include <gtest/gtest.h>
-#include <tvm/runtime/c_backend_api.h>
-#include <tvm/runtime/logging.h>
-
-#include <atomic>
-#include <memory>
-#include <mutex>
-#include <sstream>
-#include <thread>
-#include <unordered_map>
-#include <unordered_set>
-
-constexpr size_t N = 128;
-void AtomicCompute(int task_id, size_t n, std::atomic<size_t>* acc,
TVMParallelGroupEnv* penv) {
- const size_t N_per_task = (n + penv->num_task - 1) / penv->num_task;
- for (size_t i = task_id * N_per_task; i < n && i < (task_id + 1) *
N_per_task; ++i) {
- acc->fetch_add(i, std::memory_order_relaxed);
- }
- return;
-}
-
-class AffinityCheck {
- public:
- AffinityCheck(uint32_t parent_id, int max_concurrency, std::atomic<size_t>*
acc)
- : id_(parent_id), max_concurrency_(max_concurrency), acc_(acc) {}
-
- void Compute(int task_id, size_t n, TVMParallelGroupEnv* penv) {
- AtomicCompute(task_id, n, acc_, penv);
- }
-
- int GetComputeResult() { return acc_->load(std::memory_order_relaxed); }
-
- void GetAffinity(int task_id) {
-#if defined(__linux__)
- cpu_set_t cpuset;
- CPU_ZERO(&cpuset);
- pthread_getaffinity_np(pthread_self(), sizeof(cpu_set_t), &cpuset);
- std::lock_guard<std::mutex> lock(mutex_);
- thread_affinity_[task_id] = cpuset;
- // Printing the current thread CPU affinity.
- std::ostringstream str;
- for (int i = 0; i < max_concurrency_; i++) {
- if (CPU_ISSET(i, &cpuset)) {
- str << i << ",";
- }
- }
-#endif
- }
-
- bool VerifyAffinity(const std::vector<uint32_t>& cpus) {
-#if defined(__linux__)
- std::unordered_set<uint32_t> uset;
- cpu_set_t cpu_mask;
- CPU_ZERO(&cpu_mask);
- for (auto x : cpus) {
- CPU_SET(x, &cpu_mask);
- uset.insert(x);
- }
-
- for (auto x : thread_affinity_) {
- if (!CPU_EQUAL(&cpu_mask, &x.second)) {
- bool cpu_find = false;
- for (auto cpu : uset) {
- CPU_ISSET(cpu, &x.second);
- uset.erase(cpu);
- cpu_find = true;
- break;
- }
- if (!cpu_find) return false;
- }
- }
-#endif
- return true;
- }
-
- private:
- uint32_t id_;
- int max_concurrency_;
- std::atomic<size_t>* acc_;
- std::mutex mutex_;
-#if defined(__linux__)
- std::unordered_map<int, cpu_set_t> thread_affinity_;
-#endif
-};
-
-static FTVMParallelLambda atomic_add_task_id = [](int task_id,
TVMParallelGroupEnv* penv,
- void* cdata) -> int {
- auto* data = reinterpret_cast<std::atomic<size_t>*>(cdata);
- AtomicCompute(task_id, N, data, penv);
- return 0;
-};
-
-static FTVMParallelLambda affinity_check_task_id = [](int task_id,
TVMParallelGroupEnv* penv,
- void* cdata) -> int {
- auto* data = reinterpret_cast<AffinityCheck*>(cdata);
- data->Compute(task_id, N, penv);
- data->GetAffinity(task_id);
- return 0;
-};
-
-TEST(ThreadingBackend, TVMBackendParallelLaunch) {
- std::atomic<size_t> acc(0);
- TVMBackendParallelLaunch(atomic_add_task_id, &acc, 0);
- EXPECT_EQ(acc.load(std::memory_order_relaxed), N * (N - 1) / 2);
-}
-
-TEST(ThreadingBackend, TVMBackendParallelLaunchMultipleThreads) {
- // TODO(tulloch) use parameterised tests when available.
- size_t num_jobs_per_thread = 3;
- size_t max_num_threads = 2;
-
- for (size_t num_threads = 1; num_threads < max_num_threads; ++num_threads) {
- std::vector<std::unique_ptr<std::thread>> ts;
- for (size_t i = 0; i < num_threads; ++i) {
- ts.emplace_back(new std::thread([&]() {
- for (size_t j = 0; j < num_jobs_per_thread; ++j) {
- std::atomic<size_t> acc(0);
- TVMBackendParallelLaunch(atomic_add_task_id, &acc, 0);
- EXPECT_EQ(acc.load(std::memory_order_relaxed), N * (N - 1) / 2);
- }
- }));
- }
- for (auto& t : ts) {
- t->join();
- }
- }
-}
-
-TEST(ThreadingBackend, TVMBackendAffinityConfigure) {
- int max_concurrency = tvm::runtime::threading::MaxConcurrency();
- std::vector<std::unique_ptr<std::thread>> ts;
- // Returning as there is only one CPU available.
- if (max_concurrency <= 1) {
- return;
- }
- // Creating two threads to test the 'CPU list affinity' feature.
- const int threads_num = 2;
- // Getting the maximum number of CPUs which are available to each thread.
- const int cpus_num_per_thread = max_concurrency / threads_num;
- // Testing two mode of affinity.,
- std::vector<tvm::runtime::threading::ThreadGroup::AffinityMode> modes = {
- tvm::runtime::threading::ThreadGroup::kSpecifyOneCorePerThread,
- tvm::runtime::threading::ThreadGroup::kSpecifyThreadShareAllCore};
- for (auto mode : modes) {
- for (int thread_pool_idx = 0; thread_pool_idx < threads_num;
thread_pool_idx++) {
- ts.emplace_back(new std::thread(
- [&](int thread_pool_index, int sys_max_concurrency,
- tvm::runtime::threading::ThreadGroup::AffinityMode
affinity_mode) {
- std::atomic<size_t> acc(0);
- AffinityCheck ac(thread_pool_index, sys_max_concurrency, &acc);
- std::vector<unsigned int> cpus;
- for (int k = 0; k < cpus_num_per_thread; k++) {
- cpus.push_back(thread_pool_index * cpus_num_per_thread + k);
- }
- tvm::runtime::threading ::Configure(affinity_mode, 0, cpus);
- TVMBackendParallelLaunch(affinity_check_task_id, &ac, 0);
- EXPECT_EQ(ac.GetComputeResult(), N * (N - 1) / 2);
- EXPECT_EQ(ac.VerifyAffinity(cpus), true);
- },
- thread_pool_idx, max_concurrency, mode));
- }
- }
- for (auto& t : ts) {
- t->join();
- }
-}
-
-TEST(ThreadingBackend, TVMBackendParallelForWithThreadingBackend) {
- int n = 100;
- std::vector<int> vec(/*size=*/n, /*value=*/0);
- tvm::runtime::parallel_for_with_threading_backend([&vec](int i) { vec[i] =
i; }, 0, n);
- for (int i = 0; i < n; ++i) {
- EXPECT_EQ(vec[i], i);
- }
-}
diff --git a/tests/cpp/tir_analysis_side_effect.cc
b/tests/cpp/tir_analysis_side_effect.cc
deleted file mode 100644
index 1b14bfc83e..0000000000
--- a/tests/cpp/tir_analysis_side_effect.cc
+++ /dev/null
@@ -1,38 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-#include <tvm/ir/expr.h>
-#include <tvm/ir/prim/op.h>
-#include <tvm/runtime/logging.h>
-#include <tvm/te/operation.h>
-#include <tvm/tirx/expr.h>
-#include <tvm/tirx/op/gpu.h>
-
-TEST(SimplePasses, SideEffect) {
- using namespace tvm::prim;
- using namespace tvm;
- auto buf = tirx::decl_tensor({16}, PrimType::Float(32));
- auto i = PrimVar("i", PrimType::Int(32));
- TVM_FFI_ICHECK(SideEffect(tirx::MakeTensorLoad(buf, {i})) ==
CallEffectKind::kReadState);
- TVM_FFI_ICHECK(SideEffect(exp(prim::Cast(PrimType::Float(32), i + 1))) ==
CallEffectKind::kPure);
- TVM_FFI_ICHECK(
- SideEffect(tvm::Call(PrimType::Void(), tirx::gpu_storage_sync_op(),
{StringImm("shared")})
- .as_or_throw<PrimExpr>()) ==
CallEffectKind::kUpdateState);
-}
diff --git a/tests/cpp/tir_scalable_datatype.cc
b/tests/cpp/tir_scalable_datatype.cc
deleted file mode 100644
index 531db6f467..0000000000
--- a/tests/cpp/tir_scalable_datatype.cc
+++ /dev/null
@@ -1,213 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gmock/gmock.h>
-#include <gtest/gtest.h>
-#include <tvm/ffi/dtype.h>
-#include <tvm/ir/prim/expr.h>
-#include <tvm/ir/prim/op.h>
-#include <tvm/script/printer/printer.h>
-#include <tvm/tirx/op/abi.h>
-
-#ifdef TVM_LLVM_VERSION
-#include <llvm/IR/Intrinsics.h>
-#endif
-
-#include "../../src/script/printer/utils.h"
-
-using ::testing::HasSubstr;
-
-// ---------
-// Prim Type
-// ---------
-TEST(ScalablePrimType, TestCreateScalableType) {
- tvm::PrimType scalable_type = tvm::PrimType::ScalableVector(kDLInt, 32, 4);
- ASSERT_EQ(scalable_type.code(), kDLInt);
- ASSERT_EQ(scalable_type.bits(), 32);
- ASSERT_EQ(scalable_type.VScaleFactor(), 4);
- ASSERT_TRUE(scalable_type.IsScalableVector());
- ASSERT_TRUE(scalable_type.IsScalableVector() ||
scalable_type.IsFixedLengthVector());
-}
-
-TEST(ScalablePrimType, TestScalableWithBits) {
- tvm::PrimType scalable_type = tvm::PrimType::ScalableVector(kDLInt, 1, 8);
- scalable_type = scalable_type.WithBits(32);
- ASSERT_EQ(scalable_type.bits(), 32);
- ASSERT_TRUE(scalable_type.IsScalableVector());
- ASSERT_TRUE(scalable_type.IsScalableVector() ||
scalable_type.IsFixedLengthVector());
-}
-
-TEST(ScalablePrimType, TestScalableWithVscaleFactor) {
- tvm::PrimType type = tvm::PrimType::Int(32);
- tvm::PrimType scalable_type = tvm::PrimType::ScalableVector(type.code(),
type.bits(), 4);
- ASSERT_EQ(scalable_type.VScaleFactor(), 4);
- ASSERT_TRUE(scalable_type.IsScalableVector());
- ASSERT_TRUE(scalable_type.IsScalableVector() ||
scalable_type.IsFixedLengthVector());
-}
-
-TEST(ScalablePrimType, TestAssignScalablePrimType) {
- tvm::PrimType scalable_type = tvm::PrimType::ScalableVector(kDLInt, 32, 2);
- tvm::PrimType scalable_type_copy = scalable_type;
- ASSERT_TRUE(scalable_type_copy.IsScalableVector());
- ASSERT_TRUE(scalable_type_copy.IsScalableVector() ||
scalable_type_copy.IsFixedLengthVector());
-}
-
-TEST(ScalablePrimType, TestScalablePrimTypeEquality) {
- ASSERT_TRUE(tvm::PrimType::ScalableVector(kDLInt, 32, 4) ==
- tvm::PrimType::ScalableVector(kDLInt, 32, 4));
-}
-
-TEST(ScalablePrimType, TestScalablePrimTypeAndNonScalablePrimTypeInequality) {
- ASSERT_FALSE(tvm::PrimType::ScalableVector(kDLInt, 32, 4) ==
tvm::PrimType::Int(32, 4));
-}
-
-TEST(ScalablePrimType, TestIsScalar) {
- ASSERT_FALSE(tvm::PrimType::ScalableVector(kDLInt, 32, 4).IsScalar());
- ASSERT_TRUE(tvm::PrimType::Int(32).IsScalar());
- ASSERT_FALSE(tvm::PrimType::Int(32, 4).IsScalar());
- ASSERT_FALSE(tvm::PrimType::Void().IsScalar());
-}
-
-TEST(ScalablePrimType, TestScalablePrimTypeToString) {
- tvm::PrimType scalable_type = tvm::PrimType::ScalableVector(kDLInt, 32, 4);
- EXPECT_EQ(tvm::ffi::DLDataTypeToString(scalable_type->dtype),
"int32xvscalex4");
-}
-
-TEST(ScalablePrimType, TestStringToScalablePrimType) {
- std::string scalable_type_str = "int32xvscalex4";
- EXPECT_EQ(tvm::PrimType(tvm::ffi::StringToDLDataType(scalable_type_str)),
- tvm::PrimType::ScalableVector(kDLInt, 32, 4));
-}
-
-TEST(ScalablePrimType, TestInvalidStringToScalablePrimType) {
- std::string scalable_type_str = "int32x4xvscale";
- EXPECT_THROW(
- {
- try {
- tvm::ffi::StringToDLDataType(scalable_type_str);
- } catch (const tvm::ffi::Error& e) {
- EXPECT_THAT(e.what(), HasSubstr("unknown dtype `int32x4xvscale`"));
- throw;
- }
- },
- tvm::ffi::Error);
-}
-
-TEST(ScalablePrimType, TestGetScalableVectorBytes) {
- tvm::PrimType scalable_type = tvm::PrimType::ScalableVector(kDLInt, 32, 4);
- EXPECT_THROW(
- {
- try {
- int bytes = (scalable_type.bits() * scalable_type.lanes() + 7) / 8;
- static_cast<void>(bytes);
- } catch (const tvm::ffi::Error& e) {
- EXPECT_THAT(e.what(),
- HasSubstr("Can't fetch the lanes of a scalable vector at
a compile time"));
- throw;
- }
- },
- tvm::ffi::Error);
-}
-
-TEST(ScalablePrimType, TestScalablePrimTypeInvalidLanesError) {
- EXPECT_THROW(
- {
- try {
- tvm::PrimType::ScalableVector(kDLFloat, 62, 1);
- } catch (const tvm::ffi::Error& e) {
- EXPECT_THAT(e.what(), HasSubstr("Invalid value for vscale factor"));
- throw;
- }
- },
- tvm::ffi::Error);
-}
-
-TEST(ScalablePrimType, TestScalablePrimTypeInvalidVscaleFactorAccess) {
- tvm::PrimType fixed_length_type = tvm::PrimType::Float(32, 4);
- ASSERT_TRUE(fixed_length_type.IsFixedLengthVector());
- ASSERT_TRUE(fixed_length_type.IsScalableVector() ||
fixed_length_type.IsFixedLengthVector());
- EXPECT_THROW(
- {
- try {
- fixed_length_type.VScaleFactor();
- } catch (const tvm::ffi::Error& e) {
- EXPECT_THAT(e.what(), HasSubstr("A fixed length vector doesn't have
a vscale factor"));
- throw;
- }
- },
- tvm::ffi::Error);
-}
-
-TEST(ScalablePrimType, TestScalablePrimTypeInvalidLanesAccess) {
- tvm::PrimType scalable_type = tvm::PrimType::ScalableVector(kDLFloat, 32, 4);
- EXPECT_THROW(
- {
- try {
- scalable_type.lanes();
- } catch (const tvm::ffi::Error& e) {
- EXPECT_THAT(e.what(),
- HasSubstr("Can't fetch the lanes of a scalable vector at
a compile time"));
- throw;
- }
- },
- tvm::ffi::Error);
-}
-
-TEST(ScalablePrimType, TestScalableBool) {
- tvm::PrimType scalable_type = tvm::PrimType::ScalableVector(kDLBool, 8, 4);
- ASSERT_EQ(scalable_type.code(), kDLBool);
- ASSERT_EQ(scalable_type.bits(), 8);
- ASSERT_EQ(scalable_type.VScaleFactor(), 4);
- ASSERT_TRUE(scalable_type.IsScalableVector());
-}
-
-TEST(ScalablePrimType, TestScalableUInt) {
- tvm::PrimType scalable_type = tvm::PrimType::ScalableVector(kDLUInt, 1, 4);
- ASSERT_EQ(scalable_type.code(), kDLUInt);
- ASSERT_EQ(scalable_type.bits(), 1);
- ASSERT_EQ(scalable_type.VScaleFactor(), 4);
- ASSERT_TRUE(scalable_type.IsScalableVector());
-}
-
-// -----------
-// Integration
-// -----------
-#ifdef TVM_LLVM_VERSION
-TEST(ScalablePrimType, TestScalableIntrinCall) {
- tvm::PrimType scalable_type = tvm::PrimType::ScalableVector(kDLInt, 32, 4);
- tvm::Call call = tvm::Call(scalable_type, tvm::tirx::call_llvm_intrin_op(),
-#if TVM_LLVM_VERSION >= 200
-
{tvm::IntImm::Int32(::llvm::Intrinsic::stepvector)});
-#else
-
{tvm::IntImm::Int32(::llvm::Intrinsic::experimental_stepvector)});
-#endif
- ASSERT_EQ(call->ty.as_or_throw<tvm::PrimType>(), scalable_type);
- ASSERT_EQ(tvm::Script(call),
-#if TVM_LLVM_VERSION >= 200
- "T.call_llvm_intrin(\"llvm.stepvector\", ty=\"int32xvscalex4\")");
-#else
- "T.call_llvm_intrin(\"llvm.experimental.stepvector\",
ty=\"int32xvscalex4\")");
-#endif
-}
-#endif
-
-TEST(ScalablePrimType, TestTIRScriptScalableDtype2Str) {
- tvm::PrimType scalable_type = tvm::PrimType::ScalableVector(kDLInt, 32, 4);
- ASSERT_EQ(tvm::script::printer::DType2Str(scalable_type->dtype),
"int32xvscalex4");
-}
diff --git a/tests/cpp/stmt_functor_test.cc
b/tests/cpp/tirx/stmt_functor_test.cc
similarity index 100%
rename from tests/cpp/stmt_functor_test.cc
rename to tests/cpp/tirx/stmt_functor_test.cc
diff --git a/tests/cpp/topi_ewise_test.cc b/tests/cpp/topi_ewise_test.cc
deleted file mode 100644
index 4375fa2a40..0000000000
--- a/tests/cpp/topi_ewise_test.cc
+++ /dev/null
@@ -1,33 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-#include <tvm/te/operation.h>
-#include <tvm/topi/elemwise.h>
-
-namespace tvm {
-namespace topi {
-TEST(Tensor, Basic) {
- using namespace tvm;
- PrimVar m("m"), l("l");
- Tensor A = placeholder({m, l}, PrimType::Float(32), "A");
- auto C = topi::exp(A);
-}
-} // namespace topi
-} // namespace tvm
