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

Reply via email to