diff --git a/docs/arch/pass_infra.rst b/docs/arch/pass_infra.rst
index bf7b52229d13..c54ba18b0add 100644
--- a/docs/arch/pass_infra.rst
+++ b/docs/arch/pass_infra.rst
@@ -550,7 +550,7 @@ a certain scope.
 
 .. code:: python
 
-    @tvm._ffi.register_object("transform.PassContext")
+    @tvm.ffi.register_object("transform.PassContext")
     class PassContext(tvm.runtime.Object):
         def __enter__(self):
             _transform.EnterPassContext(self)
diff --git a/python/setup.py b/python/setup.py
index ffa8c2bbd433..679f5078d3c1 100644
--- a/python/setup.py
+++ b/python/setup.py
@@ -20,6 +20,7 @@
 import pathlib
 import shutil
 import sys
+import sys
 
 from setuptools import find_packages
 from setuptools.dist import Distribution
@@ -42,7 +43,7 @@ def get_lib_path():
     """Get library path, name and version"""
     # We can not import `libinfo.py` in setup.py directly since __init__.py
     # Will be invoked which introduces dependencies
-    libinfo_py = os.path.join(CURRENT_DIR, "./tvm/_ffi/libinfo.py")
+    libinfo_py = os.path.join(CURRENT_DIR, "./tvm/libinfo.py")
     libinfo = {"__file__": libinfo_py}
     exec(compile(open(libinfo_py, "rb").read(), libinfo_py, "exec"), libinfo, libinfo)
     version = libinfo["__version__"]
@@ -145,7 +146,15 @@ def config_cython():
     try:
         from Cython.Build import cythonize
 
-        subdir = "_cy3"
+        # for python 3.12+, use limited API for future compact
+        limited_api_kwargs = {}
+        if sys.version_info >= (3, 12):
+            limited_api_kwargs = {
+                "define_macros": [
+                    ("Py_LIMITED_API", 0x030C0000),
+                ],
+                "py_limited_api": True,
+            }
 
         ret = []
         extra_compile_args = ["-std=c++17", "-DDMLC_USE_LOGGING_LIBRARY=<tvm/runtime/logging.h>"]
@@ -179,6 +188,7 @@ def config_cython():
                     library_dirs=library_dirs,
                     libraries=libraries,
                     language="c++",
+                    **limited_api_kwargs,
                 )
             )
         return cythonize(ret, compiler_directives={"language_level": 3})
diff --git a/python/tvm/__init__.py b/python/tvm/__init__.py
index 8563c84ab398..150e5d4b1dbc 100644
--- a/python/tvm/__init__.py
+++ b/python/tvm/__init__.py
@@ -19,13 +19,12 @@
 import multiprocessing
 import sys
 import os
-import traceback
 
 # top-level alias
 # tvm._ffi
-from ._ffi.base import TVMError, __version__, _RUNTIME_ONLY
+from .base import TVMError, __version__, _RUNTIME_ONLY
 
-from ._ffi import register_object, register_func, get_global_func
+from .ffi import register_object, register_func, get_global_func
 
 # top-level alias
 # tvm.runtime
diff --git a/python/tvm/_ffi/__init__.py b/python/tvm/_ffi/__init__.py
deleted file mode 100644
index 559ca84635bd..000000000000
--- a/python/tvm/_ffi/__init__.py
+++ /dev/null
@@ -1,31 +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.
-"""C interfacing code.
-
-This namespace contains everything that interacts with C code.
-Most TVM C related object are ctypes compatible, which means
-they contains a handle field that is ctypes.c_void_p and can
-be used via ctypes function calls.
-
-Some performance critical functions are implemented by cython
-and have a ctypes fallback implementation.
-"""
-from . import _pyversion
-from . import base
-from .registry import register_object, register_func
-from .registry import _init_api, get_global_func
-from ..ffi import register_error
diff --git a/python/tvm/_ffi/_pyversion.py b/python/tvm/_ffi/_pyversion.py
deleted file mode 100644
index b661cfd875fc..000000000000
--- a/python/tvm/_ffi/_pyversion.py
+++ /dev/null
@@ -1,26 +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.
-"""Python version check
-"""
-import sys
-
-# ----------------------------
-# Python3 version.
-# ----------------------------
-if not (sys.version_info[0] >= 3 and sys.version_info[1] >= 6):
-    PY3STATEMENT = "The minimal Python requirement is Python 3.6"
-    raise Exception(PY3STATEMENT)
diff --git a/python/tvm/_ffi/registry.py b/python/tvm/_ffi/registry.py
deleted file mode 100644
index b11cab48b841..000000000000
--- a/python/tvm/_ffi/registry.py
+++ /dev/null
@@ -1,29 +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.
-
-# pylint: disable=invalid-name, unused-import, wrong-import-position
-"""FFI registry to register function and objects."""
-
-import tvm.ffi
-
-from tvm.ffi import register_object, register_func, get_global_func
-
-from tvm.ffi.registry import (
-    list_global_func_names,
-    remove_global_func,
-    _init_api,
-)
diff --git a/python/tvm/arith/_ffi_api.py b/python/tvm/arith/_ffi_api.py
index c551e5651563..e05405b0fcc6 100644
--- a/python/tvm/arith/_ffi_api.py
+++ b/python/tvm/arith/_ffi_api.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.arith"""
-import tvm._ffi
+import tvm.ffi
 
 
-tvm._ffi._init_api("arith", __name__)
+tvm.ffi._init_api("arith", __name__)
diff --git a/python/tvm/arith/analyzer.py b/python/tvm/arith/analyzer.py
index f8069a717da3..919272a2734b 100644
--- a/python/tvm/arith/analyzer.py
+++ b/python/tvm/arith/analyzer.py
@@ -19,7 +19,7 @@
 import enum
 from typing import Union
 
-import tvm._ffi
+import tvm.ffi
 from tvm import tir, ir
 from tvm.runtime import Object
 
@@ -46,7 +46,7 @@ class Extension(enum.Flag):
     ComparisonOfProductAndSum = 1 << 3
 
 
-@tvm._ffi.register_object("arith.ModularSet")
+@tvm.ffi.register_object("arith.ModularSet")
 class ModularSet(Object):
     """Represent range of (coeff * x + base) for x in Z"""
 
@@ -54,7 +54,7 @@ def __init__(self, coeff, base):
         self.__init_handle_by_constructor__(_ffi_api.ModularSet, coeff, base)
 
 
-@tvm._ffi.register_object("arith.ConstIntBound")
+@tvm.ffi.register_object("arith.ConstIntBound")
 class ConstIntBound(Object):
     """Represent constant integer bound
 
diff --git a/python/tvm/arith/int_set.py b/python/tvm/arith/int_set.py
index d38f5e805f39..f779df5d4c92 100644
--- a/python/tvm/arith/int_set.py
+++ b/python/tvm/arith/int_set.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """Integer set."""
-import tvm._ffi
+import tvm.ffi
 from tvm.runtime import Object
 from . import _ffi_api
 
@@ -64,7 +64,7 @@ def single_point(point):
         return _ffi_api.intset_single_point(point)
 
 
-@tvm._ffi.register_object("arith.IntervalSet")
+@tvm.ffi.register_object("arith.IntervalSet")
 class IntervalSet(IntSet):
     """Represent set of continuous interval [min_value, max_value]
 
@@ -81,7 +81,7 @@ def __init__(self, min_value, max_value):
         self.__init_handle_by_constructor__(_ffi_api.IntervalSet, min_value, max_value)
 
 
-@tvm._ffi.register_object("arith.PresburgerSet")
+@tvm.ffi.register_object("arith.PresburgerSet")
 class PresburgerSet(IntSet):
     """Represent of Presburger Set"""
 
diff --git a/python/tvm/arith/int_solver.py b/python/tvm/arith/int_solver.py
index 6e8a010eec16..a97cda10f8eb 100644
--- a/python/tvm/arith/int_solver.py
+++ b/python/tvm/arith/int_solver.py
@@ -15,12 +15,12 @@
 # specific language governing permissions and limitations
 # under the License.
 """integer constraints data structures and solvers"""
-import tvm._ffi
+import tvm.ffi
 from tvm.runtime import Object
 from . import _ffi_api
 
 
-@tvm._ffi.register_object("arith.IntGroupBounds")
+@tvm.ffi.register_object("arith.IntGroupBounds")
 class IntGroupBounds(Object):
     """Represent integer grouped bounds which are classified into
        lower bounds (include), upper bounds (include) and equalities.
@@ -66,7 +66,7 @@ def find_best_range(self):
         return _ffi_api.IntGroupBounds_FindBestRange(self)
 
 
-@tvm._ffi.register_object("arith.IntConstraints")
+@tvm.ffi.register_object("arith.IntConstraints")
 class IntConstraints(Object):
     """Represent a set of integer constraints including variables, their ranges and
        the relations between them (either equations or inequalities)
@@ -85,7 +85,7 @@ def __init__(self, variables, ranges, relations):
         self.__init_handle_by_constructor__(_ffi_api.IntConstraints, variables, ranges, relations)
 
 
-@tvm._ffi.register_object("arith.IntConstraintsTransform")
+@tvm.ffi.register_object("arith.IntConstraintsTransform")
 class IntConstraintsTransform(Object):
     """We can have different set of variables to represent the same integer constraints.
        For example, the following two constrains are equivalent,
diff --git a/python/tvm/arith/iter_affine_map.py b/python/tvm/arith/iter_affine_map.py
index f19dd0a1bac9..dbb4087f325f 100644
--- a/python/tvm/arith/iter_affine_map.py
+++ b/python/tvm/arith/iter_affine_map.py
@@ -14,9 +14,9 @@
 # KIND, either express or implied.  See the License for the
 # specific language governing permissions and limitations
 # under the License.
-""" Iterator (quasi)affine mapping patterns."""
+"""Iterator (quasi)affine mapping patterns."""
 from enum import IntEnum
-import tvm._ffi
+import tvm.ffi
 from tvm.runtime import Object
 from tvm.ir import PrimExpr
 from . import _ffi_api
@@ -26,7 +26,7 @@ class IterMapExpr(PrimExpr):
     """Base class of all IterMap expressions."""
 
 
-@tvm._ffi.register_object("arith.IterMark")
+@tvm.ffi.register_object("arith.IterMark")
 class IterMark(Object):
     """Mark the source as an iterator in [0, extent).
 
@@ -43,7 +43,7 @@ def __init__(self, source, extent):
         self.__init_handle_by_constructor__(_ffi_api.IterMark, source, extent)
 
 
-@tvm._ffi.register_object("arith.IterSplitExpr")
+@tvm.ffi.register_object("arith.IterSplitExpr")
 class IterSplitExpr(IterMapExpr):
     """Split of an iterator.
 
@@ -70,7 +70,7 @@ def __init__(self, source, lower_factor, extent, scale):
         )
 
 
-@tvm._ffi.register_object("arith.IterSumExpr")
+@tvm.ffi.register_object("arith.IterSumExpr")
 class IterSumExpr(IterMapExpr):
     """Fuse multiple iterators by summing them with scaling.
 
diff --git a/python/tvm/_ffi/base.py b/python/tvm/base.py
similarity index 86%
rename from python/tvm/_ffi/base.py
rename to python/tvm/base.py
index 18ed40fb4cb9..13b4fc8d443a 100644
--- a/python/tvm/_ffi/base.py
+++ b/python/tvm/base.py
@@ -16,22 +16,23 @@
 # under the License.
 # coding: utf-8
 # pylint: disable=invalid-name, import-outside-toplevel
-"""Base library for TVM FFI."""
+"""Base library for TVM."""
 import ctypes
 import os
 import sys
 
-
-import numpy as np
-
 from . import libinfo
 
+# ----------------------------
+# Python3 version.
+# ----------------------------
+if not (sys.version_info[0] >= 3 and sys.version_info[1] >= 8):
+    PY3STATEMENT = "The minimal Python requirement is Python 3.8"
+    raise Exception(PY3STATEMENT)
+
 # ----------------------------
 # library loading
 # ----------------------------
-string_types = (str,)
-integer_types = (int, np.int32)
-numeric_types = integer_types + (float, np.float16, np.float32)
 
 
 def _load_lib():
@@ -62,7 +63,7 @@ def _load_lib():
 
 
 if _RUNTIME_ONLY:
-    from ..ffi import registry as _tvm_ffi_registry
+    from .ffi import registry as _tvm_ffi_registry
 
     _tvm_ffi_registry._SKIP_UNKNOWN_OBJECTS = True
 
diff --git a/python/tvm/contrib/cc.py b/python/tvm/contrib/cc.py
index 110f80db6186..04a69baee9c1 100644
--- a/python/tvm/contrib/cc.py
+++ b/python/tvm/contrib/cc.py
@@ -23,7 +23,7 @@
 import sys
 from typing import Dict
 
-from .._ffi.base import py_str
+from ..base import py_str
 from . import tar as _tar
 from . import utils as _utils
 
diff --git a/python/tvm/contrib/clang.py b/python/tvm/contrib/clang.py
index 16c465dc22ab..4d2769436d06 100644
--- a/python/tvm/contrib/clang.py
+++ b/python/tvm/contrib/clang.py
@@ -18,7 +18,7 @@
 # pylint: disable=invalid-name
 import subprocess
 
-from tvm._ffi.base import py_str
+from tvm.base import py_str
 import tvm.target
 from . import utils
 
diff --git a/python/tvm/contrib/coreml_runtime.py b/python/tvm/contrib/coreml_runtime.py
index aa4f21279967..def5d3c2e06e 100644
--- a/python/tvm/contrib/coreml_runtime.py
+++ b/python/tvm/contrib/coreml_runtime.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """CoreML runtime that load and run coreml models."""
-import tvm._ffi
+import tvm.ffi
 from ..rpc import base as rpc_base
 
 
@@ -41,7 +41,7 @@ def create(symbol, compiled_model_path, device):
     if device_type >= rpc_base.RPC_SESS_MASK:
         fcreate = device._rpc_sess.get_function(runtime_func)
     else:
-        fcreate = tvm._ffi.get_global_func(runtime_func)
+        fcreate = tvm.ffi.get_global_func(runtime_func)
         assert fcreate, "Cannot find `tvm.coreml_runtime.create` function."
 
     return CoreMLModule(fcreate(symbol, compiled_model_path))
diff --git a/python/tvm/contrib/cudnn.py b/python/tvm/contrib/cudnn.py
index 9d3f80a5c74b..1c80d4a3b9e1 100644
--- a/python/tvm/contrib/cudnn.py
+++ b/python/tvm/contrib/cudnn.py
@@ -20,7 +20,7 @@
 import numpy as np
 import tvm
 
-import tvm._ffi
+import tvm.ffi
 from tvm import te
 
 # algos can be read from cudnn.h
@@ -349,7 +349,7 @@ def _conv_find_algo(
         dims - 2, pad, stride, dilation, x_shape, w_shape
     )
     yshape = np.array(y_shape, dtype=np.int32)
-    func = tvm._ffi.get_global_func(func_name)
+    func = tvm.ffi.get_global_func(func_name)
     return func(
         tensor_format,
         dims - 2,
diff --git a/python/tvm/contrib/cutlass/_ffi_api.py b/python/tvm/contrib/cutlass/_ffi_api.py
index e71eb8c13f19..be71b0d48f13 100644
--- a/python/tvm/contrib/cutlass/_ffi_api.py
+++ b/python/tvm/contrib/cutlass/_ffi_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI API for CUTLASS BYOC."""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("contrib.cutlass", __name__)
+tvm.ffi._init_api("contrib.cutlass", __name__)
diff --git a/python/tvm/contrib/cutlass/build.py b/python/tvm/contrib/cutlass/build.py
index e12ce93e270b..0aea5bf1416a 100644
--- a/python/tvm/contrib/cutlass/build.py
+++ b/python/tvm/contrib/cutlass/build.py
@@ -26,7 +26,7 @@
 
 import tvm
 from tvm import relax, runtime
-from tvm._ffi.registry import register_func
+from tvm.ffi.registry import register_func
 from tvm.contrib.nvcc import get_cuda_version
 from tvm.topi.utils import get_const_tuple
 
diff --git a/python/tvm/contrib/cutlass/gen_tensor_op.py b/python/tvm/contrib/cutlass/gen_tensor_op.py
index 88581b1cb4f4..6fa349b28e44 100644
--- a/python/tvm/contrib/cutlass/gen_tensor_op.py
+++ b/python/tvm/contrib/cutlass/gen_tensor_op.py
@@ -24,7 +24,7 @@
 import subprocess
 import tempfile
 
-import tvm._ffi
+import tvm.ffi
 from tvm.runtime import Object
 from tvm.tir import IntImm
 
@@ -461,7 +461,7 @@ def _get_optional_int_annotation(annotations, key, default=None):
     return int(value)
 
 
-@tvm._ffi.register_func("contrib.cutlass.instantiate_template")
+@tvm.ffi.register_func("contrib.cutlass.instantiate_template")
 def instantiate_template(func_name, annotations, func_args):
     """Return CUTLASS host code based on a template and the provided annotations.
 
diff --git a/python/tvm/contrib/emcc.py b/python/tvm/contrib/emcc.py
index 3beb096b6747..a9f0bec0cf9d 100644
--- a/python/tvm/contrib/emcc.py
+++ b/python/tvm/contrib/emcc.py
@@ -20,8 +20,8 @@
 import subprocess
 from pathlib import Path
 
-from tvm._ffi.base import py_str
-from tvm._ffi.libinfo import find_lib_path
+from tvm.base import py_str
+from tvm.libinfo import find_lib_path
 
 
 def create_tvmjs_wasm(output, objects, options=None, cc="emcc", libs=None):
diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py
index d1ca5227fd0f..f4b02ff80f73 100644
--- a/python/tvm/contrib/hexagon/build.py
+++ b/python/tvm/contrib/hexagon/build.py
@@ -35,7 +35,7 @@
 from typing import Union
 
 from tvm.contrib.hexagon.hexagon_profiler import HexagonProfiler
-from ..._ffi import libinfo
+from ...ffi import libinfo
 from .session import Session
 from .tools import HEXAGON_SIMULATOR_NAME
 
diff --git a/python/tvm/contrib/hexagon/tools.py b/python/tvm/contrib/hexagon/tools.py
index 3b129b03323b..5ee89713d9a5 100644
--- a/python/tvm/contrib/hexagon/tools.py
+++ b/python/tvm/contrib/hexagon/tools.py
@@ -29,7 +29,7 @@
 
 import tvm
 import tvm.contrib.cc as cc
-from ..._ffi.registry import register_func
+from ...ffi.registry import register_func
 
 
 # Linking Hexagon shared libraries.
diff --git a/python/tvm/contrib/miopen.py b/python/tvm/contrib/miopen.py
index 0e336c1c82b9..22b08f38ca76 100644
--- a/python/tvm/contrib/miopen.py
+++ b/python/tvm/contrib/miopen.py
@@ -19,7 +19,7 @@
 import ctypes
 import numpy as np
 import tvm
-import tvm._ffi
+import tvm.ffi
 
 from tvm import te
 
@@ -94,7 +94,7 @@ def conv2d_forward(
     oshape = np.zeros((len(x.shape)), dtype=np.int32)
     xshape = x.shape
     wshape = w.shape
-    setup_func = tvm._ffi.get_global_func("tvm.contrib.miopen.conv2d.setup")
+    setup_func = tvm.ffi.get_global_func("tvm.contrib.miopen.conv2d.setup")
     algo = setup_func(
         conv_mode,
         data_type,
diff --git a/python/tvm/contrib/mrvl.py b/python/tvm/contrib/mrvl.py
index 3cf393b34160..36c932cd1a1d 100644
--- a/python/tvm/contrib/mrvl.py
+++ b/python/tvm/contrib/mrvl.py
@@ -24,10 +24,10 @@
 import base64
 import numpy as np
 import tvm
-import tvm._ffi
+import tvm.ffi
 
 
-@tvm._ffi.register_func("tvm.mrvl.find_value_in_KV_pair")
+@tvm.ffi.register_func("tvm.mrvl.find_value_in_KV_pair")
 def find_value_in_KV_pair(json_input: str, key_to_find: str) -> str:
     """This function takes the graph_json string and key to be searched in
     the json string, using json parser routine it loads the json string
@@ -54,7 +54,7 @@ def find_value_in_KV_pair(json_input: str, key_to_find: str) -> str:
     return value
 
 
-@tvm._ffi.register_func("tvm.mrvl.GetNodesJSONString")
+@tvm.ffi.register_func("tvm.mrvl.GetNodesJSONString")
 def get_nodes_json_string(graph_json):
     """This takes the graph_json string from MrvlJSONSerializer and adds / modifies
     the json string to a form suitable for the Marvell Backend.
@@ -206,7 +206,7 @@ def get_nodes_json_string(graph_json):
     return nodes_json_string
 
 
-@tvm._ffi.register_func("tvm.mrvl.ModifyConstNames")
+@tvm.ffi.register_func("tvm.mrvl.ModifyConstNames")
 def modify_const_names(nodes_json_str, consts_json_str):
     """This takes the graph module returned by build an generates nodes and constant
        meta data suitable for compilation by the back end.
@@ -329,7 +329,7 @@ def get_working_dir():
     return os.getcwd()
 
 
-@tvm._ffi.register_func("tvm.mrvl.WriteJsonFile")
+@tvm.ffi.register_func("tvm.mrvl.WriteJsonFile")
 def write_json_file(json_string, json_filename):
     """Generate json file under working directory"""
     working_dir = get_working_dir()
@@ -351,7 +351,7 @@ def delete_temp_files(symbol_name):
         shutil.rmtree(bin_folder)
 
 
-@tvm._ffi.register_func("tvm.mrvl.CompileModel")
+@tvm.ffi.register_func("tvm.mrvl.CompileModel")
 def compile_model(
     symbol_name,
     nodes_json_string,
@@ -414,7 +414,7 @@ def compile_model(
         raise RuntimeError(error_msg)
 
 
-@tvm._ffi.register_func("tvm.mrvl.CleanUpSim")
+@tvm.ffi.register_func("tvm.mrvl.CleanUpSim")
 def clean_up_sim(bin_file, input_json, input_bin, out_bin_prefix, num_outputs):
     os.remove(bin_file)
     os.remove(input_json)
@@ -424,7 +424,7 @@ def clean_up_sim(bin_file, input_json, input_bin, out_bin_prefix, num_outputs):
         os.remove(out_bin)
 
 
-@tvm._ffi.register_func("tvm.mrvl.SearchPath")
+@tvm.ffi.register_func("tvm.mrvl.SearchPath")
 def search_path(file_name):
     path = shutil.which(file_name)
     if path is None:
@@ -432,7 +432,7 @@ def search_path(file_name):
     return os.path.dirname(path)
 
 
-@tvm._ffi.register_func("tvm.mrvl.JsonToBin")
+@tvm.ffi.register_func("tvm.mrvl.JsonToBin")
 def convert_json_to_bin(json_file, input_bin_file):
     with open(json_file) as input_json:
         data = json.load(input_json)
@@ -442,7 +442,7 @@ def convert_json_to_bin(json_file, input_bin_file):
             f.write(data_b)
 
 
-@tvm._ffi.register_func("tvm.mrvl.RunSim")
+@tvm.ffi.register_func("tvm.mrvl.RunSim")
 def run_simulation(run_command, sim_directory):
     cwd_path = get_working_dir()
     os.mkdir(sim_directory)
@@ -452,6 +452,6 @@ def run_simulation(run_command, sim_directory):
     shutil.rmtree(sim_directory)
 
 
-@tvm._ffi.register_func("tvm.mrvl.TempDir")
+@tvm.ffi.register_func("tvm.mrvl.TempDir")
 def get_temp_dir():
     return tempfile.gettempdir()
diff --git a/python/tvm/contrib/msc/core/_ffi_api.py b/python/tvm/contrib/msc/core/_ffi_api.py
index c0b0e21267ea..f7c975aff98a 100644
--- a/python/tvm/contrib/msc/core/_ffi_api.py
+++ b/python/tvm/contrib/msc/core/_ffi_api.py
@@ -16,6 +16,6 @@
 # under the License.
 """tvm.contrib.msc.core._ffi_api"""
 
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("msc.core", __name__)
+tvm.ffi._init_api("msc.core", __name__)
diff --git a/python/tvm/contrib/msc/core/ir/graph.py b/python/tvm/contrib/msc/core/ir/graph.py
index 172f40e06a31..9aa5bde93380 100644
--- a/python/tvm/contrib/msc/core/ir/graph.py
+++ b/python/tvm/contrib/msc/core/ir/graph.py
@@ -25,7 +25,7 @@
 from tvm.contrib.msc.core import utils as msc_utils
 
 
-@tvm._ffi.register_object("msc.core.MSCTensor")
+@tvm.ffi.register_object("msc.core.MSCTensor")
 class MSCTensor(Object):
     """Tensor in MSCGraph
 
@@ -198,7 +198,7 @@ class BaseJoint(Object):
     """Base class of all MSC Nodes."""
 
 
-@tvm._ffi.register_object("msc.core.MSCJoint")
+@tvm.ffi.register_object("msc.core.MSCJoint")
 class MSCJoint(BaseJoint):
     """Node in MSCGraph
 
@@ -423,7 +423,7 @@ def equal(self, other: BaseJoint) -> bool:
         return msc_utils.dict_equal(self.get_attrs(), other.get_attrs())
 
 
-@tvm._ffi.register_object("msc.core.MSCPrim")
+@tvm.ffi.register_object("msc.core.MSCPrim")
 class MSCPrim(BaseJoint):
     """Prim in MSCGraph
 
@@ -447,7 +447,7 @@ def __init__(
         self.__init_handle_by_constructor__(_ffi_api.MSCPrim, index, name, optype, attrs, parents)
 
 
-@tvm._ffi.register_object("msc.core.WeightJoint")
+@tvm.ffi.register_object("msc.core.WeightJoint")
 class WeightJoint(BaseJoint):
     """Node in WeightGraph
 
@@ -565,7 +565,7 @@ class BaseGraph(Object):
     """Base class of all MSC Graphs."""
 
 
-@tvm._ffi.register_object("msc.core.MSCGraph")
+@tvm.ffi.register_object("msc.core.MSCGraph")
 class MSCGraph(BaseGraph):
     """The MSCGraph
 
@@ -954,7 +954,7 @@ def visualize(self, path: Optional[str] = None) -> str:
         return graph_proto
 
 
-@tvm._ffi.register_object("msc.core.WeightGraph")
+@tvm.ffi.register_object("msc.core.WeightGraph")
 class WeightGraph(Object):
     """The WeightGraph
 
diff --git a/python/tvm/contrib/msc/framework/tensorflow/_ffi_api.py b/python/tvm/contrib/msc/framework/tensorflow/_ffi_api.py
index d43984b9c292..5b85e16a53ba 100644
--- a/python/tvm/contrib/msc/framework/tensorflow/_ffi_api.py
+++ b/python/tvm/contrib/msc/framework/tensorflow/_ffi_api.py
@@ -16,6 +16,6 @@
 # under the License.
 """tvm.contrib.msc.framework.tensorflow._ffi_api"""
 
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("msc.framework.tensorflow", __name__)
+tvm.ffi._init_api("msc.framework.tensorflow", __name__)
diff --git a/python/tvm/contrib/msc/framework/tensorrt/_ffi_api.py b/python/tvm/contrib/msc/framework/tensorrt/_ffi_api.py
index c0fa9c2c0559..4db71f3a19de 100644
--- a/python/tvm/contrib/msc/framework/tensorrt/_ffi_api.py
+++ b/python/tvm/contrib/msc/framework/tensorrt/_ffi_api.py
@@ -16,6 +16,6 @@
 # under the License.
 """tvm.contrib.msc.framework.tensorrt._ffi_api"""
 
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("msc.framework.tensorrt", __name__)
+tvm.ffi._init_api("msc.framework.tensorrt", __name__)
diff --git a/python/tvm/contrib/msc/framework/torch/_ffi_api.py b/python/tvm/contrib/msc/framework/torch/_ffi_api.py
index 190e7507fb07..d12fcf2e2f87 100644
--- a/python/tvm/contrib/msc/framework/torch/_ffi_api.py
+++ b/python/tvm/contrib/msc/framework/torch/_ffi_api.py
@@ -16,6 +16,6 @@
 # under the License.
 """tvm.contrib.msc.framework.torch._ffi_api"""
 
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("msc.framework.torch", __name__)
+tvm.ffi._init_api("msc.framework.torch", __name__)
diff --git a/python/tvm/contrib/msc/framework/tvm/_ffi_api.py b/python/tvm/contrib/msc/framework/tvm/_ffi_api.py
index e82612a6403f..a3683181b0e4 100644
--- a/python/tvm/contrib/msc/framework/tvm/_ffi_api.py
+++ b/python/tvm/contrib/msc/framework/tvm/_ffi_api.py
@@ -16,6 +16,6 @@
 # under the License.
 """tvm.contrib.msc.framework.tvm._ffi_api"""
 
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("msc.framework.tvm", __name__)
+tvm.ffi._init_api("msc.framework.tvm", __name__)
diff --git a/python/tvm/contrib/msc/plugin/_ffi_api.py b/python/tvm/contrib/msc/plugin/_ffi_api.py
index 0e12c29242d1..c566d3b0d332 100644
--- a/python/tvm/contrib/msc/plugin/_ffi_api.py
+++ b/python/tvm/contrib/msc/plugin/_ffi_api.py
@@ -16,6 +16,6 @@
 # under the License.
 """tvm.contrib.msc.plugin._ffi_api"""
 
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("msc.plugin", __name__)
+tvm.ffi._init_api("msc.plugin", __name__)
diff --git a/python/tvm/contrib/msc/plugin/op/_ffi_api.py b/python/tvm/contrib/msc/plugin/op/_ffi_api.py
index 2111e11227a1..0d8ad3c5e457 100644
--- a/python/tvm/contrib/msc/plugin/op/_ffi_api.py
+++ b/python/tvm/contrib/msc/plugin/op/_ffi_api.py
@@ -16,6 +16,6 @@
 # under the License.
 """tvm.contrib.msc.plugin.op._ffi_api"""
 
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("msc.plugin.op", __name__)
+tvm.ffi._init_api("msc.plugin.op", __name__)
diff --git a/python/tvm/contrib/ndk.py b/python/tvm/contrib/ndk.py
index 14820c0ca8ab..c1441c496ae8 100644
--- a/python/tvm/contrib/ndk.py
+++ b/python/tvm/contrib/ndk.py
@@ -25,8 +25,8 @@
 import tempfile
 from pathlib import Path
 
-from .._ffi import register_func
-from .._ffi.base import py_str
+from ..ffi import register_func
+from ..base import py_str
 from . import utils as _utils, tar as _tar, cc as _cc
 from .cc import get_target_by_dump_machine
 
diff --git a/python/tvm/contrib/nnpack.py b/python/tvm/contrib/nnpack.py
index 010bef533c00..1b4f51850805 100644
--- a/python/tvm/contrib/nnpack.py
+++ b/python/tvm/contrib/nnpack.py
@@ -17,7 +17,7 @@
 """External function interface to NNPACK libraries."""
 import tvm
 from tvm import te
-import tvm._ffi
+import tvm.ffi
 
 
 def is_available():
@@ -232,4 +232,4 @@ def convolution_inference_weight_transform(
     )
 
 
-tvm._ffi._init_api("tvm.contrib.nnpack")
+tvm.ffi._init_api("tvm.contrib.nnpack")
diff --git a/python/tvm/contrib/nvcc.py b/python/tvm/contrib/nvcc.py
index c8b749b36bf1..45e2793fbb6f 100644
--- a/python/tvm/contrib/nvcc.py
+++ b/python/tvm/contrib/nvcc.py
@@ -22,10 +22,10 @@
 import subprocess
 import warnings
 
-import tvm._ffi
+import tvm.ffi
 from tvm.target import Target
 
-from .._ffi.base import py_str
+from ..base import py_str
 from . import utils
 
 
@@ -198,14 +198,14 @@ def get_cuda_version(cuda_path=None):
     raise RuntimeError("Cannot read cuda version file")
 
 
-@tvm._ffi.register_func
+@tvm.ffi.register_func
 def tvm_callback_cuda_compile(code, target):  # pylint: disable=unused-argument
     """use nvcc to generate fatbin code for better optimization"""
     ptx = compile_cuda(code, target_format="fatbin")
     return ptx
 
 
-@tvm._ffi.register_func("tvm_callback_libdevice_path")
+@tvm.ffi.register_func("tvm_callback_libdevice_path")
 def find_libdevice_path(arch):
     """Utility function to find libdevice
 
@@ -270,7 +270,7 @@ def callback_libdevice_path(arch):
         return ""
 
 
-@tvm._ffi.register_func("tvm.contrib.nvcc.get_compute_version")
+@tvm.ffi.register_func("tvm.contrib.nvcc.get_compute_version")
 def get_target_compute_version(target=None):
     """Utility function to get compute capability of compilation target.
 
@@ -415,7 +415,7 @@ def have_cudagraph():
         return False
 
 
-@tvm._ffi.register_func("tvm.contrib.nvcc.supports_bf16")
+@tvm.ffi.register_func("tvm.contrib.nvcc.supports_bf16")
 def have_bf16(compute_version):
     """Either bf16 support is provided in the compute capability or not
 
@@ -431,7 +431,7 @@ def have_bf16(compute_version):
     return False
 
 
-@tvm._ffi.register_func("tvm.contrib.nvcc.supports_fp8")
+@tvm.ffi.register_func("tvm.contrib.nvcc.supports_fp8")
 def have_fp8(compute_version):
     """Whether fp8 support is provided in the specified compute capability or not
 
@@ -449,7 +449,7 @@ def have_fp8(compute_version):
     return False
 
 
-@tvm._ffi.register_func("tvm.contrib.nvcc.supports_fp4")
+@tvm.ffi.register_func("tvm.contrib.nvcc.supports_fp4")
 def have_fp4(compute_version):
     """Whether fp4 support is provided in the specified compute capability or not
 
diff --git a/python/tvm/contrib/pickle_memoize.py b/python/tvm/contrib/pickle_memoize.py
index 72efb8064312..a5d09fb9316a 100644
--- a/python/tvm/contrib/pickle_memoize.py
+++ b/python/tvm/contrib/pickle_memoize.py
@@ -23,7 +23,6 @@
 import sys
 
 import functools
-from .._ffi.base import string_types
 
 try:
     import cPickle as pickle
@@ -115,7 +114,7 @@ def memoize(key, save_at_exit=False):
 
     def _register(f):
         """Registration function"""
-        allow_types = (string_types, int, float, tuple)
+        allow_types = (str, int, float, tuple)
         fkey = key + "." + f.__name__ + ".pkl"
         if fkey not in Cache.cache_by_key:
             Cache.cache_by_key[fkey] = Cache(fkey, save_at_exit)
diff --git a/python/tvm/contrib/random.py b/python/tvm/contrib/random.py
index bbc74fccac94..6a17693b9162 100644
--- a/python/tvm/contrib/random.py
+++ b/python/tvm/contrib/random.py
@@ -17,7 +17,7 @@
 """External function interface to random library."""
 import tvm
 from tvm import te
-import tvm._ffi
+import tvm.ffi
 
 
 def randint(low, high, size, dtype="int32"):
@@ -112,4 +112,4 @@ def normal(loc, scale, size):
     )
 
 
-tvm._ffi._init_api("tvm.contrib.random")
+tvm.ffi._init_api("tvm.contrib.random")
diff --git a/python/tvm/contrib/rocm.py b/python/tvm/contrib/rocm.py
index f3427463b3e0..6e6a985c2732 100644
--- a/python/tvm/contrib/rocm.py
+++ b/python/tvm/contrib/rocm.py
@@ -20,8 +20,8 @@
 import os
 from os.path import join, exists
 
-import tvm._ffi
-from tvm._ffi.base import py_str
+import tvm.ffi
+from tvm.base import py_str
 import tvm.runtime
 import tvm.target
 
@@ -99,7 +99,7 @@ def rocm_link(in_file, out_file, lld=None):
         raise RuntimeError(msg)
 
 
-@tvm._ffi.register_func("tvm_callback_rocm_link")
+@tvm.ffi.register_func("tvm_callback_rocm_link")
 def callback_rocm_link(obj_bin):
     """Links object file generated from LLVM to HSA Code Object
 
@@ -123,7 +123,7 @@ def callback_rocm_link(obj_bin):
     return cobj_bin
 
 
-@tvm._ffi.register_func("tvm_callback_rocm_bitcode_path")
+@tvm.ffi.register_func("tvm_callback_rocm_bitcode_path")
 def callback_rocm_bitcode_path(rocdl_dir=None):
     """Utility function to find ROCm device library bitcodes
 
@@ -227,7 +227,7 @@ def have_matrixcore(compute_version=None):
     return False
 
 
-@tvm._ffi.register_func("tvm_callback_rocm_get_arch")
+@tvm.ffi.register_func("tvm_callback_rocm_get_arch")
 def get_rocm_arch(rocm_path=None):
     """Utility function to get the AMD GPU architecture
 
diff --git a/python/tvm/contrib/spirv.py b/python/tvm/contrib/spirv.py
index 94b24d0c7b09..0484562ee737 100644
--- a/python/tvm/contrib/spirv.py
+++ b/python/tvm/contrib/spirv.py
@@ -18,7 +18,7 @@
 import subprocess
 import os
 from . import utils
-from .._ffi.base import py_str
+from ..base import py_str
 
 
 def optimize(spv_bin):
diff --git a/python/tvm/contrib/tar.py b/python/tvm/contrib/tar.py
index 67175b8b278c..7322ed447197 100644
--- a/python/tvm/contrib/tar.py
+++ b/python/tvm/contrib/tar.py
@@ -22,7 +22,7 @@
 import shutil
 import subprocess
 from . import utils
-from .._ffi.base import py_str
+from ..base import py_str
 
 
 def tar(output, files):
diff --git a/python/tvm/contrib/tflite_runtime.py b/python/tvm/contrib/tflite_runtime.py
index 1558e36d51af..aceeefd248f4 100644
--- a/python/tvm/contrib/tflite_runtime.py
+++ b/python/tvm/contrib/tflite_runtime.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """TFLite runtime that load and run tflite models."""
-import tvm._ffi
+import tvm.ffi
 from ..rpc import base as rpc_base
 
 
@@ -45,7 +45,7 @@ def create(tflite_model_bytes, device, runtime_target="cpu"):
     if device_type >= rpc_base.RPC_SESS_MASK:
         fcreate = device._rpc_sess.get_function(runtime_func)
     else:
-        fcreate = tvm._ffi.get_global_func(runtime_func)
+        fcreate = tvm.ffi.get_global_func(runtime_func)
 
     return TFLiteModule(fcreate(bytearray(tflite_model_bytes), device))
 
diff --git a/python/tvm/contrib/thrust.py b/python/tvm/contrib/thrust.py
index 8f3178429589..9a05cfafbac3 100644
--- a/python/tvm/contrib/thrust.py
+++ b/python/tvm/contrib/thrust.py
@@ -17,7 +17,7 @@
 """Utilities for thrust"""
 import logging
 
-from tvm._ffi import get_global_func
+from tvm.ffi import get_global_func
 
 
 def maybe_warn(target, func_name):
diff --git a/python/tvm/contrib/tvmjs.py b/python/tvm/contrib/tvmjs.py
index d936c8a2276c..e24b88a3f8c3 100644
--- a/python/tvm/contrib/tvmjs.py
+++ b/python/tvm/contrib/tvmjs.py
@@ -34,7 +34,7 @@
     ml_dtypes = None
 
 import tvm
-from tvm._ffi.libinfo import find_lib_path
+from tvm.libinfo import find_lib_path
 from tvm.runtime import DataType
 
 from .emcc import create_tvmjs_wasm
diff --git a/python/tvm/contrib/xcode.py b/python/tvm/contrib/xcode.py
index d12367330dde..adfc2dcd8459 100644
--- a/python/tvm/contrib/xcode.py
+++ b/python/tvm/contrib/xcode.py
@@ -21,7 +21,7 @@
 import sys
 import subprocess
 import json
-from .._ffi.base import py_str
+from ..base import py_str
 from . import utils
 
 
diff --git a/python/tvm/dlight/analysis/common_analysis.py b/python/tvm/dlight/analysis/common_analysis.py
index be260b894203..a3499274e5a8 100644
--- a/python/tvm/dlight/analysis/common_analysis.py
+++ b/python/tvm/dlight/analysis/common_analysis.py
@@ -20,7 +20,7 @@
 from typing_extensions import Literal
 
 from tvm import ir, tir
-from tvm._ffi import get_global_func
+from tvm.ffi import get_global_func
 from tvm.target.target import Target
 from tvm.tir import Schedule
 from tvm.tir.schedule import BlockRV
diff --git a/python/tvm/driver/_ffi_api.py b/python/tvm/driver/_ffi_api.py
index c423656d78f5..1ceecc9c94c6 100644
--- a/python/tvm/driver/_ffi_api.py
+++ b/python/tvm/driver/_ffi_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.driver"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("driver", __name__)
+tvm.ffi._init_api("driver", __name__)
diff --git a/python/tvm/exec/disco_worker.py b/python/tvm/exec/disco_worker.py
index b1f1554b56f9..ecfeaa1ebb88 100644
--- a/python/tvm/exec/disco_worker.py
+++ b/python/tvm/exec/disco_worker.py
@@ -22,7 +22,7 @@
 from typing import Callable
 
 import tvm
-from tvm._ffi import get_global_func, register_func
+from tvm.ffi import get_global_func, register_func
 from tvm.runtime import NDArray, ShapeTuple, String
 from tvm.runtime.ndarray import array
 
diff --git a/python/tvm/generic.py b/python/tvm/generic.py
deleted file mode 100644
index 7c46312c2ea5..000000000000
--- a/python/tvm/generic.py
+++ /dev/null
@@ -1,19 +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.
-"""Generic operators."""
-# pylint:disable=unused-wildcard-import, wildcard-import
-from .tir.generic import *
diff --git a/python/tvm/ir/_ffi_analysis_api.py b/python/tvm/ir/_ffi_analysis_api.py
index 0013ec3b5026..ca38c2309f41 100644
--- a/python/tvm/ir/_ffi_analysis_api.py
+++ b/python/tvm/ir/_ffi_analysis_api.py
@@ -16,7 +16,7 @@
 # under the License.
 """FFI APIs for tvm.ir.analysis"""
 
-import tvm._ffi
+import tvm.ffi
 
 
-tvm._ffi._init_api("ir.analysis", __name__)
+tvm.ffi._init_api("ir.analysis", __name__)
diff --git a/python/tvm/ir/_ffi_api.py b/python/tvm/ir/_ffi_api.py
index d3a9505c38d0..6434a3925e98 100644
--- a/python/tvm/ir/_ffi_api.py
+++ b/python/tvm/ir/_ffi_api.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.ir"""
-import tvm._ffi
+import tvm.ffi
 
 
-tvm._ffi._init_api("ir", __name__)
+tvm.ffi._init_api("ir", __name__)
diff --git a/python/tvm/ir/_ffi_instrument_api.py b/python/tvm/ir/_ffi_instrument_api.py
index bf62caf30e5a..d88faf7fddd0 100644
--- a/python/tvm/ir/_ffi_instrument_api.py
+++ b/python/tvm/ir/_ffi_instrument_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.instrument"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("instrument", __name__)
+tvm.ffi._init_api("instrument", __name__)
diff --git a/python/tvm/ir/_ffi_transform_api.py b/python/tvm/ir/_ffi_transform_api.py
index bb01b559c3d8..1a27fc58776c 100644
--- a/python/tvm/ir/_ffi_transform_api.py
+++ b/python/tvm/ir/_ffi_transform_api.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.transform"""
-import tvm._ffi
+import tvm.ffi
 
 
-tvm._ffi._init_api("transform", __name__)
+tvm.ffi._init_api("transform", __name__)
diff --git a/python/tvm/ir/attrs.py b/python/tvm/ir/attrs.py
index fc63138043fa..6565a8de37b4 100644
--- a/python/tvm/ir/attrs.py
+++ b/python/tvm/ir/attrs.py
@@ -15,14 +15,14 @@
 # specific language governing permissions and limitations
 # under the License.
 """TVM Attribute module, which is mainly used for defining attributes of operators."""
-import tvm._ffi
+import tvm.ffi
 
 from tvm.runtime import Object
 import tvm.runtime._ffi_node_api
 from . import _ffi_api
 
 
-@tvm._ffi.register_object
+@tvm.ffi.register_object
 class Attrs(Object):
     """Attribute node, which is mainly use for defining attributes of operators.
 
@@ -93,7 +93,7 @@ def __getitem__(self, item):
         return self.__getattr__(item)
 
 
-@tvm._ffi.register_object
+@tvm.ffi.register_object
 class DictAttrs(Attrs):
     """Dictionary attributes."""
 
diff --git a/python/tvm/ir/base.py b/python/tvm/ir/base.py
index 50b2c595b33f..a31be4c40ccb 100644
--- a/python/tvm/ir/base.py
+++ b/python/tvm/ir/base.py
@@ -15,9 +15,9 @@
 # specific language governing permissions and limitations
 # under the License.
 """Common base structures."""
-import tvm._ffi
+import tvm.ffi
 import tvm.error
-from tvm._ffi import get_global_func, register_object
+from tvm.ffi import get_global_func, register_object
 from tvm.runtime import Object, _ffi_node_api
 
 from . import _ffi_api, json_compact
diff --git a/python/tvm/ir/diagnostics/__init__.py b/python/tvm/ir/diagnostics/__init__.py
index c4d4fcc57807..ac4adc3306e6 100644
--- a/python/tvm/ir/diagnostics/__init__.py
+++ b/python/tvm/ir/diagnostics/__init__.py
@@ -22,7 +22,7 @@
 and the DiagnosticRenderer.
 """
 import enum
-import tvm._ffi
+import tvm.ffi
 from . import _ffi_api
 from ... import get_global_func, register_func, Object
 
@@ -69,7 +69,7 @@ class DiagnosticLevel(enum.IntEnum):
     HELP = 50
 
 
-@tvm._ffi.register_object("Diagnostic")
+@tvm.ffi.register_object("Diagnostic")
 class Diagnostic(Object):
     """A single diagnostic object from TVM."""
 
@@ -77,7 +77,7 @@ def __init__(self, level, span, message):
         self.__init_handle_by_constructor__(_ffi_api.Diagnostic, level, span, message)
 
 
-@tvm._ffi.register_object("DiagnosticRenderer")
+@tvm.ffi.register_object("DiagnosticRenderer")
 class DiagnosticRenderer(Object):
     """
     A diagnostic renderer, which given a diagnostic context produces a "rendered"
@@ -100,7 +100,7 @@ def render(self, ctx):
 
 
 # Register the diagnostic context.
-@tvm._ffi.register_object("DiagnosticContext")
+@tvm.ffi.register_object("DiagnosticContext")
 class DiagnosticContext(Object):
     """
     A diagnostic context which records active errors
diff --git a/python/tvm/ir/diagnostics/_ffi_api.py b/python/tvm/ir/diagnostics/_ffi_api.py
index 430fd17f4d8a..fb157c977510 100644
--- a/python/tvm/ir/diagnostics/_ffi_api.py
+++ b/python/tvm/ir/diagnostics/_ffi_api.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI for TVM diagnostics."""
-import tvm._ffi
+import tvm.ffi
 
 
-tvm._ffi._init_api("diagnostics", __name__)
+tvm.ffi._init_api("diagnostics", __name__)
diff --git a/python/tvm/ir/expr.py b/python/tvm/ir/expr.py
index d140b5867c6e..197b0831bf25 100644
--- a/python/tvm/ir/expr.py
+++ b/python/tvm/ir/expr.py
@@ -18,7 +18,7 @@
 from numbers import Number
 from typing import Optional
 
-import tvm._ffi
+import tvm.ffi
 
 from ..runtime import Object, Scriptable
 from . import _ffi_api
@@ -71,7 +71,7 @@ def struct_info(self) -> Optional["tvm.relax.StructInfo"]:
         return _ffi_api.ExprStructInfo(self)
 
 
-@tvm._ffi.register_object("GlobalVar")
+@tvm.ffi.register_object("GlobalVar")
 class GlobalVar(RelaxExpr):
     """A global variable in the IR.
 
@@ -117,7 +117,7 @@ def __call__(self, *args: RelaxExpr) -> BaseExpr:
         raise RuntimeError(f"Do not know how to handle GlobalVar.__call__ for types {arg_types}")
 
 
-@tvm._ffi.register_object
+@tvm.ffi.register_object
 class Range(Node, Scriptable):
     """Represent a range in TVM.
 
diff --git a/python/tvm/ir/instrument.py b/python/tvm/ir/instrument.py
index 4aee761af698..1e1505858f50 100644
--- a/python/tvm/ir/instrument.py
+++ b/python/tvm/ir/instrument.py
@@ -19,13 +19,13 @@
 import inspect
 import functools
 
-import tvm._ffi
+import tvm.ffi
 import tvm.runtime
 
 from . import _ffi_instrument_api
 
 
-@tvm._ffi.register_object("instrument.PassInstrument")
+@tvm.ffi.register_object("instrument.PassInstrument")
 class PassInstrument(tvm.runtime.Object):
     """A pass instrument implementation.
 
@@ -225,7 +225,7 @@ def create_pass_instrument(pi_cls):
     return create_pass_instrument
 
 
-@tvm._ffi.register_object("instrument.PassInstrument")
+@tvm.ffi.register_object("instrument.PassInstrument")
 class PassTimingInstrument(tvm.runtime.Object):
     """A wrapper to create a passes time instrument that implemented in C++"""
 
diff --git a/python/tvm/ir/module.py b/python/tvm/ir/module.py
index 8347e218beb9..6033dc6f8066 100644
--- a/python/tvm/ir/module.py
+++ b/python/tvm/ir/module.py
@@ -20,8 +20,7 @@
 
 from typing import Dict, Union
 
-import tvm._ffi
-from tvm._ffi.base import string_types
+import tvm.ffi
 from tvm.runtime import Scriptable
 from tvm.runtime.object import Object
 
@@ -31,7 +30,7 @@
 from .base import Node
 
 
-@tvm._ffi.register_object("IRModule")
+@tvm.ffi.register_object("IRModule")
 class IRModule(Node, Scriptable):
     """IRModule that holds functions and type definitions.
 
@@ -49,7 +48,7 @@ def __init__(self, functions=None, attrs=None, global_infos=None):
         elif isinstance(functions, dict):
             mapped_funcs = {}
             for k, v in functions.items():
-                if isinstance(k, string_types):
+                if isinstance(k, str):
                     k = _expr.GlobalVar(k)
                 if not isinstance(k, _expr.GlobalVar):
                     raise TypeError("Expect functions to be Dict[GlobalVar, Function]")
@@ -98,7 +97,7 @@ def __setitem__(self, var, val):
 
     def _add(self, var, val, update=True):
         if isinstance(val, _expr.RelaxExpr):
-            if isinstance(var, string_types):
+            if isinstance(var, str):
                 if _ffi_api.Module_ContainGlobalVar(self, var):
                     var = _ffi_api.Module_GetGlobalVar(self, var)
                 else:
@@ -118,7 +117,7 @@ def __getitem__(self, var):
         val: Union[Function, Type]
             The definition referenced by :code:`var` (either a function or type).
         """
-        if isinstance(var, string_types):
+        if isinstance(var, str):
             return _ffi_api.Module_Lookup_str(self, var)
         assert isinstance(var, _expr.GlobalVar)
         return _ffi_api.Module_Lookup(self, var)
diff --git a/python/tvm/ir/op.py b/python/tvm/ir/op.py
index 932aef24c60d..41105c4549dd 100644
--- a/python/tvm/ir/op.py
+++ b/python/tvm/ir/op.py
@@ -16,13 +16,13 @@
 # under the License.
 # pylint: disable=invalid-name
 """Primitive operators in the TVM IR."""
-import tvm._ffi
+import tvm.ffi
 
 from . import _ffi_api
 from .expr import RelaxExpr
 
 
-@tvm._ffi.register_object("Op")
+@tvm.ffi.register_object("Op")
 class Op(RelaxExpr):
     """Primitive operator in the IR."""
 
diff --git a/python/tvm/ir/supply.py b/python/tvm/ir/supply.py
index a501e8849e03..046432edfd99 100644
--- a/python/tvm/ir/supply.py
+++ b/python/tvm/ir/supply.py
@@ -20,7 +20,7 @@
 from . import _ffi_api
 
 
-@tvm._ffi.register_object("NameSupply")
+@tvm.ffi.register_object("NameSupply")
 class NameSupply(Object):
     """NameSupply that can be used to generate unique names.
 
@@ -77,7 +77,7 @@ def contains_name(self, name, add_prefix=True):
         return _ffi_api.NameSupply_ContainsName(self, name, add_prefix)
 
 
-@tvm._ffi.register_object("GlobalVarSupply")
+@tvm.ffi.register_object("GlobalVarSupply")
 class GlobalVarSupply(Object):
     """GlobalVarSupply that holds a mapping between names and GlobalVars.
 
diff --git a/python/tvm/ir/transform.py b/python/tvm/ir/transform.py
index 644909f4d481..45050d44af0b 100644
--- a/python/tvm/ir/transform.py
+++ b/python/tvm/ir/transform.py
@@ -19,13 +19,13 @@
 import inspect
 import functools
 
-import tvm._ffi
+import tvm.ffi
 import tvm.runtime
 
 from . import _ffi_transform_api
 
 
-@tvm._ffi.register_object("transform.PassInfo")
+@tvm.ffi.register_object("transform.PassInfo")
 class PassInfo(tvm.runtime.Object):
     """The class contains the meta data required by a pass. It is the
     container of information needed by running an optimization or analysis.
@@ -50,7 +50,7 @@ def __init__(self, opt_level, name, required=None, traceable=False):
         )
 
 
-@tvm._ffi.register_object("transform.PassContext")
+@tvm.ffi.register_object("transform.PassContext")
 class PassContext(tvm.runtime.Object):
     """The basis where a TVM optimization/analysis runs on.
     Each pass context contains a number of auxiliary information that is used
@@ -209,7 +209,7 @@ def get_tuning_api_database(self):
         return _ffi_transform_api.GetTuningAPIDatabase(self)
 
 
-@tvm._ffi.register_object("transform.Pass")
+@tvm.ffi.register_object("transform.Pass")
 class Pass(tvm.runtime.Object):
     """The base class of all passes. All methods here are just simple wrappers
     that are implemented in the backend. They are defined for users to
@@ -238,7 +238,7 @@ def __call__(self, mod):
         return _ffi_transform_api.RunPass(self, mod)
 
 
-@tvm._ffi.register_object("transform.ModulePass")
+@tvm.ffi.register_object("transform.ModulePass")
 class ModulePass(Pass):
     """A pass that works on tvm.IRModule. Users don't need to interact with
     this class directly. Instead, a module pass should be created through
@@ -249,7 +249,7 @@ class ModulePass(Pass):
     """
 
 
-@tvm._ffi.register_object("transform.Sequential")
+@tvm.ffi.register_object("transform.Sequential")
 class Sequential(Pass):
     """A pass that works on a sequence of pass objects. Multiple passes can be
     executed sequentially using this class.
diff --git a/python/tvm/ir/type.py b/python/tvm/ir/type.py
index 3d372012b649..9ec8ef8fbd02 100644
--- a/python/tvm/ir/type.py
+++ b/python/tvm/ir/type.py
@@ -16,7 +16,7 @@
 # under the License.
 """Unified type system in the project."""
 import tvm
-import tvm._ffi
+import tvm.ffi
 from tvm.runtime import Scriptable
 
 from . import _ffi_api
@@ -38,7 +38,7 @@ def same_as(self, other):
         return super().__eq__(other)
 
 
-@tvm._ffi.register_object("PrimType")
+@tvm.ffi.register_object("PrimType")
 class PrimType(Type):
     """Primitive data type in the low level IR
 
@@ -52,7 +52,7 @@ def __init__(self, dtype):
         self.__init_handle_by_constructor__(_ffi_api.PrimType, dtype)
 
 
-@tvm._ffi.register_object("PointerType")
+@tvm.ffi.register_object("PointerType")
 class PointerType(Type):
     """PointerType used in the low-level TIR.
 
@@ -69,7 +69,7 @@ def __init__(self, element_type, storage_scope=""):
         self.__init_handle_by_constructor__(_ffi_api.PointerType, element_type, storage_scope)
 
 
-@tvm._ffi.register_object("TupleType")
+@tvm.ffi.register_object("TupleType")
 class TupleType(Type):
     """The type of tuple values.
 
@@ -83,7 +83,7 @@ def __init__(self, fields):
         self.__init_handle_by_constructor__(_ffi_api.TupleType, fields)
 
 
-@tvm._ffi.register_object("FuncType")
+@tvm.ffi.register_object("FuncType")
 class FuncType(Type):
     """Function type.
 
diff --git a/python/tvm/ir/type_relation.py b/python/tvm/ir/type_relation.py
index dba42dbce4a1..d0175fda5706 100644
--- a/python/tvm/ir/type_relation.py
+++ b/python/tvm/ir/type_relation.py
@@ -15,13 +15,13 @@
 # specific language governing permissions and limitations
 # under the License.
 """Type relation and function for type checking."""
-import tvm._ffi
+import tvm.ffi
 
 from .type import Type, TypeConstraint
 from . import _ffi_api
 
 
-@tvm._ffi.register_object("TypeCall")
+@tvm.ffi.register_object("TypeCall")
 class TypeCall(Type):
     """Type function application.
 
@@ -43,7 +43,7 @@ def __init__(self, func, args):
         self.__init_handle_by_constructor__(_ffi_api.TypeCall, func, args)
 
 
-@tvm._ffi.register_object("TypeRelation")
+@tvm.ffi.register_object("TypeRelation")
 class TypeRelation(TypeConstraint):
     """User defined type relation, it is an input-output relation on types.
 
diff --git a/python/tvm/_ffi/libinfo.py b/python/tvm/libinfo.py
similarity index 98%
rename from python/tvm/_ffi/libinfo.py
rename to python/tvm/libinfo.py
index 55d4d8165aee..c344315334d2 100644
--- a/python/tvm/_ffi/libinfo.py
+++ b/python/tvm/libinfo.py
@@ -47,8 +47,8 @@ def get_dll_directories():
     # An installed TVM's curr_path will look something like:
     #   $PREFIX/lib/python3.6/site-packages/tvm/_ffi
     ffi_dir = os.path.dirname(os.path.realpath(os.path.expanduser(__file__)))
-    source_dir = os.path.join(ffi_dir, "..", "..", "..")
-    install_lib_dir = os.path.join(ffi_dir, "..", "..", "..", "..")
+    source_dir = os.path.join(ffi_dir, "..", "..")
+    install_lib_dir = os.path.join(ffi_dir, "..", "..", "..")
 
     dll_path = []
 
diff --git a/python/tvm/meta_schedule/_ffi_api.py b/python/tvm/meta_schedule/_ffi_api.py
index 24022191a8b4..89b8df086001 100644
--- a/python/tvm/meta_schedule/_ffi_api.py
+++ b/python/tvm/meta_schedule/_ffi_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.meta_schedule"""
-from .._ffi import _init_api
+from ..ffi import _init_api
 
 _init_api("meta_schedule", __name__)  # pylint: disable=protected-access
diff --git a/python/tvm/meta_schedule/arg_info.py b/python/tvm/meta_schedule/arg_info.py
index 7390c544a50b..69c8d6d4c5dc 100644
--- a/python/tvm/meta_schedule/arg_info.py
+++ b/python/tvm/meta_schedule/arg_info.py
@@ -17,7 +17,7 @@
 """The argument information"""
 from typing import Any, List, Union
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.ir import IRModule
 from tvm.runtime import DataType, Object, ShapeTuple
 from tvm.tir import PrimFunc
diff --git a/python/tvm/meta_schedule/builder/builder.py b/python/tvm/meta_schedule/builder/builder.py
index 221077cfbd6c..f323e15bd532 100644
--- a/python/tvm/meta_schedule/builder/builder.py
+++ b/python/tvm/meta_schedule/builder/builder.py
@@ -21,7 +21,7 @@
 from typing_extensions import Literal
 
 # isort: on
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.ir import IRModule
 from tvm.runtime import NDArray, Object
 from tvm.target import Target
diff --git a/python/tvm/meta_schedule/builder/local_builder.py b/python/tvm/meta_schedule/builder/local_builder.py
index ae9ad6574e34..ff738c6265c3 100644
--- a/python/tvm/meta_schedule/builder/local_builder.py
+++ b/python/tvm/meta_schedule/builder/local_builder.py
@@ -19,7 +19,7 @@
 import tempfile
 from typing import Callable, Dict, List, Optional, Union
 
-from tvm._ffi import register_func
+from tvm.ffi import register_func
 from tvm.ir import IRModule
 from tvm.runtime import Module, NDArray, load_param_dict, save_param_dict
 from tvm.target import Target
diff --git a/python/tvm/meta_schedule/cost_model/cost_model.py b/python/tvm/meta_schedule/cost_model/cost_model.py
index 541154d4cc59..9abd50b94c75 100644
--- a/python/tvm/meta_schedule/cost_model/cost_model.py
+++ b/python/tvm/meta_schedule/cost_model/cost_model.py
@@ -24,7 +24,7 @@
 # isort: on
 
 import numpy as np  # type: ignore
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.runtime import Object
 
 from .. import _ffi_api
diff --git a/python/tvm/meta_schedule/cost_model/mlp_model.py b/python/tvm/meta_schedule/cost_model/mlp_model.py
index 9167d30e9008..9191eee6a68f 100644
--- a/python/tvm/meta_schedule/cost_model/mlp_model.py
+++ b/python/tvm/meta_schedule/cost_model/mlp_model.py
@@ -542,7 +542,7 @@ def load(  # pylint: disable=too-many-locals
                                 "_workload.json", "_candidates.json"
                             ),
                         )
-                    except tvm._ffi.base.TVMError:
+                    except tvm.base.TVMError:
                         continue
                     candidates, results = [], []
                     tuning_records = database.get_all_tuning_records()
diff --git a/python/tvm/meta_schedule/database/database.py b/python/tvm/meta_schedule/database/database.py
index 601571089592..7abaead68018 100644
--- a/python/tvm/meta_schedule/database/database.py
+++ b/python/tvm/meta_schedule/database/database.py
@@ -22,7 +22,7 @@
 
 # isort: on
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.ir.module import IRModule
 from tvm.runtime import Object
 from tvm.target import Target
diff --git a/python/tvm/meta_schedule/database/json_database.py b/python/tvm/meta_schedule/database/json_database.py
index 102a13b90d98..f3b188493767 100644
--- a/python/tvm/meta_schedule/database/json_database.py
+++ b/python/tvm/meta_schedule/database/json_database.py
@@ -18,7 +18,7 @@
 import os.path as osp
 from typing import Optional
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .database import Database
diff --git a/python/tvm/meta_schedule/database/memory_database.py b/python/tvm/meta_schedule/database/memory_database.py
index 34a6a141970a..53755333839c 100644
--- a/python/tvm/meta_schedule/database/memory_database.py
+++ b/python/tvm/meta_schedule/database/memory_database.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """A database that stores TuningRecords in memory"""
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .database import Database
diff --git a/python/tvm/meta_schedule/database/ordered_union_database.py b/python/tvm/meta_schedule/database/ordered_union_database.py
index 35b0a9e282c1..a451d8ee2fd1 100644
--- a/python/tvm/meta_schedule/database/ordered_union_database.py
+++ b/python/tvm/meta_schedule/database/ordered_union_database.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """A database consists of multiple databases."""
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .database import Database
diff --git a/python/tvm/meta_schedule/database/schedule_fn_database.py b/python/tvm/meta_schedule/database/schedule_fn_database.py
index c7d175cb79d3..3b7dfa79f6bf 100644
--- a/python/tvm/meta_schedule/database/schedule_fn_database.py
+++ b/python/tvm/meta_schedule/database/schedule_fn_database.py
@@ -17,7 +17,7 @@
 """A database for injecting handcrafted schedule functions."""
 from typing import Callable
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.tir import Schedule
 
 from .. import _ffi_api
diff --git a/python/tvm/meta_schedule/database/union_database.py b/python/tvm/meta_schedule/database/union_database.py
index ae55ebe79614..7f896c1da61f 100644
--- a/python/tvm/meta_schedule/database/union_database.py
+++ b/python/tvm/meta_schedule/database/union_database.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """A database consists of multiple databases."""
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .database import Database
diff --git a/python/tvm/meta_schedule/extracted_task.py b/python/tvm/meta_schedule/extracted_task.py
index b69a38ef6dc0..0cdede120b6f 100644
--- a/python/tvm/meta_schedule/extracted_task.py
+++ b/python/tvm/meta_schedule/extracted_task.py
@@ -17,7 +17,7 @@
 """Extracted tasks from high-level IR."""
 from typing import List
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.ir import IRModule
 from tvm.runtime import Object
 from tvm.target import Target
diff --git a/python/tvm/meta_schedule/feature_extractor/feature_extractor.py b/python/tvm/meta_schedule/feature_extractor/feature_extractor.py
index c14c97e0f526..bd37214db997 100644
--- a/python/tvm/meta_schedule/feature_extractor/feature_extractor.py
+++ b/python/tvm/meta_schedule/feature_extractor/feature_extractor.py
@@ -22,7 +22,7 @@
 
 # isort: on
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.runtime import Object
 from tvm.runtime.ndarray import NDArray
 
diff --git a/python/tvm/meta_schedule/feature_extractor/per_store_feature.py b/python/tvm/meta_schedule/feature_extractor/per_store_feature.py
index f6a456d707be..b1098bd4ea7c 100644
--- a/python/tvm/meta_schedule/feature_extractor/per_store_feature.py
+++ b/python/tvm/meta_schedule/feature_extractor/per_store_feature.py
@@ -18,7 +18,7 @@
 """We extract one feature vector per BufferStoreNode statement in a TIR Stmt,
 so we call this feature as "per-store" feature.
 """
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .feature_extractor import FeatureExtractor
diff --git a/python/tvm/meta_schedule/measure_callback/add_to_database.py b/python/tvm/meta_schedule/measure_callback/add_to_database.py
index ab61e87f647d..f40dffeaad44 100644
--- a/python/tvm/meta_schedule/measure_callback/add_to_database.py
+++ b/python/tvm/meta_schedule/measure_callback/add_to_database.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """A callback that adds the measurement results into the database"""
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .measure_callback import MeasureCallback
diff --git a/python/tvm/meta_schedule/measure_callback/measure_callback.py b/python/tvm/meta_schedule/measure_callback/measure_callback.py
index d4a10c1e4009..17a7f45460e9 100644
--- a/python/tvm/meta_schedule/measure_callback/measure_callback.py
+++ b/python/tvm/meta_schedule/measure_callback/measure_callback.py
@@ -23,7 +23,7 @@
 
 # isort: on
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.runtime import Object
 
 from .. import _ffi_api
diff --git a/python/tvm/meta_schedule/measure_callback/remove_build_artifact.py b/python/tvm/meta_schedule/measure_callback/remove_build_artifact.py
index 4b2e1ab7f428..82c18f8f9065 100644
--- a/python/tvm/meta_schedule/measure_callback/remove_build_artifact.py
+++ b/python/tvm/meta_schedule/measure_callback/remove_build_artifact.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """A callback that removes the build artifacts from the disk"""
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .measure_callback import MeasureCallback
diff --git a/python/tvm/meta_schedule/measure_callback/update_cost_model.py b/python/tvm/meta_schedule/measure_callback/update_cost_model.py
index c6ee1d26fe6d..5b8b0306d421 100644
--- a/python/tvm/meta_schedule/measure_callback/update_cost_model.py
+++ b/python/tvm/meta_schedule/measure_callback/update_cost_model.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """A measure callback that updates the cost model"""
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .measure_callback import MeasureCallback
diff --git a/python/tvm/meta_schedule/mutator/mutate_compute_location.py b/python/tvm/meta_schedule/mutator/mutate_compute_location.py
index bb361247bf62..5ebe04a6b13a 100644
--- a/python/tvm/meta_schedule/mutator/mutate_compute_location.py
+++ b/python/tvm/meta_schedule/mutator/mutate_compute_location.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """A mutator that mutates the compute-at location decision of SampleComputeLocation"""
-from tvm._ffi.registry import register_object
+from tvm.ffi.registry import register_object
 
 from .. import _ffi_api
 from .mutator import Mutator
diff --git a/python/tvm/meta_schedule/mutator/mutate_parallel.py b/python/tvm/meta_schedule/mutator/mutate_parallel.py
index c66dddb825f4..c7736fdcf71d 100644
--- a/python/tvm/meta_schedule/mutator/mutate_parallel.py
+++ b/python/tvm/meta_schedule/mutator/mutate_parallel.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """Mutator that mutates the parallel extent"""
-from tvm._ffi.registry import register_object
+from tvm.ffi.registry import register_object
 
 from .. import _ffi_api
 from .mutator import Mutator
diff --git a/python/tvm/meta_schedule/mutator/mutate_thread_binding.py b/python/tvm/meta_schedule/mutator/mutate_thread_binding.py
index 6a2553f94346..2225ca76c77d 100644
--- a/python/tvm/meta_schedule/mutator/mutate_thread_binding.py
+++ b/python/tvm/meta_schedule/mutator/mutate_thread_binding.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """Mutator that mutates the thread binding extent"""
-from tvm._ffi.registry import register_object
+from tvm.ffi.registry import register_object
 
 from .. import _ffi_api
 from .mutator import Mutator
diff --git a/python/tvm/meta_schedule/mutator/mutate_tile_size.py b/python/tvm/meta_schedule/mutator/mutate_tile_size.py
index ff432a6633b9..90cccdc3f5db 100644
--- a/python/tvm/meta_schedule/mutator/mutate_tile_size.py
+++ b/python/tvm/meta_schedule/mutator/mutate_tile_size.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """Mutator that mutates the decision of instruction Sample-Perfect-Tile"""
-from tvm._ffi.registry import register_object
+from tvm.ffi.registry import register_object
 
 from .. import _ffi_api
 from .mutator import Mutator
diff --git a/python/tvm/meta_schedule/mutator/mutate_unroll.py b/python/tvm/meta_schedule/mutator/mutate_unroll.py
index f81953d008d4..9575c3fc22d9 100644
--- a/python/tvm/meta_schedule/mutator/mutate_unroll.py
+++ b/python/tvm/meta_schedule/mutator/mutate_unroll.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """Mutator that mutates auto unroll step"""
-from tvm._ffi.registry import register_object
+from tvm.ffi.registry import register_object
 
 from .. import _ffi_api
 from .mutator import Mutator
diff --git a/python/tvm/meta_schedule/mutator/mutator.py b/python/tvm/meta_schedule/mutator/mutator.py
index 188cb30c5b69..6991c72bec41 100644
--- a/python/tvm/meta_schedule/mutator/mutator.py
+++ b/python/tvm/meta_schedule/mutator/mutator.py
@@ -22,7 +22,7 @@
 
 # isort: on
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.runtime import Object
 from tvm.tir.schedule import Trace
 
@@ -80,7 +80,7 @@ def create(
             "cuda",
             "cuda-tensorcore",
             "hexagon",
-        ]
+        ],
     ) -> Dict["Mutator", float]:
         """Create a list of default mutators.
 
diff --git a/python/tvm/meta_schedule/postproc/disallow_async_strided_mem_copy.py b/python/tvm/meta_schedule/postproc/disallow_async_strided_mem_copy.py
index 0dcff9bf45a3..5c50b2064426 100644
--- a/python/tvm/meta_schedule/postproc/disallow_async_strided_mem_copy.py
+++ b/python/tvm/meta_schedule/postproc/disallow_async_strided_mem_copy.py
@@ -16,7 +16,7 @@
 # under the License.
 """A postprocessor that checks if the IRModule has any strided memory copies"""
 
-from tvm._ffi.registry import register_object
+from tvm.ffi.registry import register_object
 from .. import _ffi_api
 from .postproc import Postproc
 
diff --git a/python/tvm/meta_schedule/postproc/disallow_dynamic_loop.py b/python/tvm/meta_schedule/postproc/disallow_dynamic_loop.py
index 5515d288e0e7..34c13aded935 100644
--- a/python/tvm/meta_schedule/postproc/disallow_dynamic_loop.py
+++ b/python/tvm/meta_schedule/postproc/disallow_dynamic_loop.py
@@ -16,7 +16,7 @@
 # under the License.
 """A postprocessor that checks if the IRModule has any loop with non-constant extent"""
 
-from tvm._ffi.registry import register_object
+from tvm.ffi.registry import register_object
 from .. import _ffi_api
 from .postproc import Postproc
 
diff --git a/python/tvm/meta_schedule/postproc/postproc.py b/python/tvm/meta_schedule/postproc/postproc.py
index af7fe9e9c502..33daabc3951c 100644
--- a/python/tvm/meta_schedule/postproc/postproc.py
+++ b/python/tvm/meta_schedule/postproc/postproc.py
@@ -22,7 +22,7 @@
 
 # isort: on
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.runtime import Object
 from tvm.tir.schedule import Schedule
 
diff --git a/python/tvm/meta_schedule/postproc/rewrite_cooperative_fetch.py b/python/tvm/meta_schedule/postproc/rewrite_cooperative_fetch.py
index e2d7c2212382..20c354ce601d 100644
--- a/python/tvm/meta_schedule/postproc/rewrite_cooperative_fetch.py
+++ b/python/tvm/meta_schedule/postproc/rewrite_cooperative_fetch.py
@@ -17,7 +17,7 @@
 """A postprocessor that rewrites the cooperative fetch annotation to actual
 vectorized cooperative fetching in loop bindings."""
 
-from tvm._ffi.registry import register_object
+from tvm.ffi.registry import register_object
 from .. import _ffi_api
 from .postproc import Postproc
 
diff --git a/python/tvm/meta_schedule/postproc/rewrite_layout.py b/python/tvm/meta_schedule/postproc/rewrite_layout.py
index 10addefee542..13556f1909d2 100644
--- a/python/tvm/meta_schedule/postproc/rewrite_layout.py
+++ b/python/tvm/meta_schedule/postproc/rewrite_layout.py
@@ -16,7 +16,7 @@
 # under the License.
 """A postprocessor that rewrites the layout of input tensor"""
 
-from tvm._ffi.registry import register_object
+from tvm.ffi.registry import register_object
 
 from .. import _ffi_api
 from .postproc import Postproc
diff --git a/python/tvm/meta_schedule/postproc/rewrite_parallel_vectorize_unroll.py b/python/tvm/meta_schedule/postproc/rewrite_parallel_vectorize_unroll.py
index abe7288acba9..0be7cdbe118f 100644
--- a/python/tvm/meta_schedule/postproc/rewrite_parallel_vectorize_unroll.py
+++ b/python/tvm/meta_schedule/postproc/rewrite_parallel_vectorize_unroll.py
@@ -17,7 +17,7 @@
 """A postprocessor that applies parallelization, vectorization and auto unrolling
 according to the annotation of each block"""
 
-from tvm._ffi.registry import register_object
+from tvm.ffi.registry import register_object
 from .. import _ffi_api
 from .postproc import Postproc
 
diff --git a/python/tvm/meta_schedule/postproc/rewrite_reduction_block.py b/python/tvm/meta_schedule/postproc/rewrite_reduction_block.py
index 7e15ed493ccb..30c8cf9b0699 100644
--- a/python/tvm/meta_schedule/postproc/rewrite_reduction_block.py
+++ b/python/tvm/meta_schedule/postproc/rewrite_reduction_block.py
@@ -16,7 +16,7 @@
 # under the License.
 """A postprocessor that rewrites reduction block by moving the init block out."""
 
-from tvm._ffi.registry import register_object
+from tvm.ffi.registry import register_object
 from .. import _ffi_api
 from .postproc import Postproc
 
diff --git a/python/tvm/meta_schedule/postproc/rewrite_tensorize.py b/python/tvm/meta_schedule/postproc/rewrite_tensorize.py
index 85075c41b43c..e04ddcbdf223 100644
--- a/python/tvm/meta_schedule/postproc/rewrite_tensorize.py
+++ b/python/tvm/meta_schedule/postproc/rewrite_tensorize.py
@@ -16,7 +16,7 @@
 # under the License.
 """A postprocessor that tensorize related components."""
 
-from tvm._ffi.registry import register_object
+from tvm.ffi.registry import register_object
 from .. import _ffi_api
 from .postproc import Postproc
 
diff --git a/python/tvm/meta_schedule/postproc/rewrite_unbound_block.py b/python/tvm/meta_schedule/postproc/rewrite_unbound_block.py
index aef5bca690e4..ca4c9cdcd624 100644
--- a/python/tvm/meta_schedule/postproc/rewrite_unbound_block.py
+++ b/python/tvm/meta_schedule/postproc/rewrite_unbound_block.py
@@ -16,7 +16,7 @@
 # under the License.
 """A postprocessor that adds thread binding to unbound blocks"""
 
-from tvm._ffi.registry import register_object
+from tvm.ffi.registry import register_object
 
 from .. import _ffi_api
 from .postproc import Postproc
diff --git a/python/tvm/meta_schedule/postproc/verify_gpu_code.py b/python/tvm/meta_schedule/postproc/verify_gpu_code.py
index 501e4423196c..1a74eadaa906 100644
--- a/python/tvm/meta_schedule/postproc/verify_gpu_code.py
+++ b/python/tvm/meta_schedule/postproc/verify_gpu_code.py
@@ -16,7 +16,7 @@
 # under the License.
 """A postprocessor that verifies if the GPU code is correct"""
 
-from tvm._ffi.registry import register_object
+from tvm.ffi.registry import register_object
 from .. import _ffi_api
 from .postproc import Postproc
 
diff --git a/python/tvm/meta_schedule/postproc/verify_vtcm_limit.py b/python/tvm/meta_schedule/postproc/verify_vtcm_limit.py
index 28d202d5b338..51a38624d28e 100644
--- a/python/tvm/meta_schedule/postproc/verify_vtcm_limit.py
+++ b/python/tvm/meta_schedule/postproc/verify_vtcm_limit.py
@@ -16,7 +16,7 @@
 # under the License.
 """A postprocessor that verifies the VTCM usage of a given schedule."""
 
-from tvm._ffi.registry import register_object
+from tvm.ffi.registry import register_object
 from .. import _ffi_api
 from .postproc import Postproc
 
diff --git a/python/tvm/meta_schedule/profiler.py b/python/tvm/meta_schedule/profiler.py
index 7b7bb6e6d17f..65c1079d65b0 100644
--- a/python/tvm/meta_schedule/profiler.py
+++ b/python/tvm/meta_schedule/profiler.py
@@ -19,7 +19,7 @@
 from contextlib import contextmanager
 from typing import Dict, Optional
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.runtime import Object
 
 from . import _ffi_api
diff --git a/python/tvm/meta_schedule/relax_integration.py b/python/tvm/meta_schedule/relax_integration.py
index 0e5adc4e9982..613405c8ad3b 100644
--- a/python/tvm/meta_schedule/relax_integration.py
+++ b/python/tvm/meta_schedule/relax_integration.py
@@ -23,7 +23,7 @@
 
 # isort: on
 
-from tvm._ffi import get_global_func, register_func
+from tvm.ffi import get_global_func, register_func
 from tvm.ir import IRModule
 from tvm.ir.transform import PassContext
 from tvm.runtime import NDArray
diff --git a/python/tvm/meta_schedule/runner/runner.py b/python/tvm/meta_schedule/runner/runner.py
index 1a8f78414e91..0c2609469a19 100644
--- a/python/tvm/meta_schedule/runner/runner.py
+++ b/python/tvm/meta_schedule/runner/runner.py
@@ -22,7 +22,7 @@
 
 # isort: on
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.runtime import Object
 
 from .. import _ffi_api
diff --git a/python/tvm/meta_schedule/schedule_rule/add_rfactor.py b/python/tvm/meta_schedule/schedule_rule/add_rfactor.py
index 72f9fc92f96e..ceb18a6c3aa6 100644
--- a/python/tvm/meta_schedule/schedule_rule/add_rfactor.py
+++ b/python/tvm/meta_schedule/schedule_rule/add_rfactor.py
@@ -17,7 +17,7 @@
 """Add-rfactor Rule that add-rfactor to some blocks if needed"""
 from typing import Optional
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .schedule_rule import ScheduleRule
diff --git a/python/tvm/meta_schedule/schedule_rule/apply_custom_rule.py b/python/tvm/meta_schedule/schedule_rule/apply_custom_rule.py
index 29e25f992930..26f61aa8ceb6 100644
--- a/python/tvm/meta_schedule/schedule_rule/apply_custom_rule.py
+++ b/python/tvm/meta_schedule/schedule_rule/apply_custom_rule.py
@@ -16,7 +16,7 @@
 # under the License.
 """Create a rule that applies customized rules registered using block attribute `schedule_rule`.
 The rule will be dispatched according to target keys."""
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .schedule_rule import ScheduleRule
diff --git a/python/tvm/meta_schedule/schedule_rule/auto_bind.py b/python/tvm/meta_schedule/schedule_rule/auto_bind.py
index 99a91f606e32..ef34e45061f7 100644
--- a/python/tvm/meta_schedule/schedule_rule/auto_bind.py
+++ b/python/tvm/meta_schedule/schedule_rule/auto_bind.py
@@ -17,7 +17,7 @@
 """Auto-bind Rule that binds blocks to threads if needed"""
 from typing import List, Optional
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .schedule_rule import ScheduleRule
diff --git a/python/tvm/meta_schedule/schedule_rule/auto_inline.py b/python/tvm/meta_schedule/schedule_rule/auto_inline.py
index c84dbaf89b97..8cd122ec93d3 100644
--- a/python/tvm/meta_schedule/schedule_rule/auto_inline.py
+++ b/python/tvm/meta_schedule/schedule_rule/auto_inline.py
@@ -17,7 +17,7 @@
 """Auto-Inline. Rule that inlines spatial blocks if it satisfies some conditions"""
 from typing import List, Optional
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .schedule_rule import ScheduleRule
diff --git a/python/tvm/meta_schedule/schedule_rule/cross_thread_reduction.py b/python/tvm/meta_schedule/schedule_rule/cross_thread_reduction.py
index f242e42aea4b..d2c780b72854 100644
--- a/python/tvm/meta_schedule/schedule_rule/cross_thread_reduction.py
+++ b/python/tvm/meta_schedule/schedule_rule/cross_thread_reduction.py
@@ -17,7 +17,7 @@
 """Rules which apply cross-thread reduction to some reduction blocks correspondingly when needed"""
 from typing import List
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .schedule_rule import ScheduleRule
diff --git a/python/tvm/meta_schedule/schedule_rule/multi_level_tiling.py b/python/tvm/meta_schedule/schedule_rule/multi_level_tiling.py
index 19651a2ce18e..2f389190d662 100644
--- a/python/tvm/meta_schedule/schedule_rule/multi_level_tiling.py
+++ b/python/tvm/meta_schedule/schedule_rule/multi_level_tiling.py
@@ -18,7 +18,7 @@
 from typing import Any, Dict, List, Mapping, NamedTuple, Optional, Callable
 
 from tvm.tir.schedule import Schedule, BlockRV
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .schedule_rule import ScheduleRule
diff --git a/python/tvm/meta_schedule/schedule_rule/parallel_vectorize_unroll.py b/python/tvm/meta_schedule/schedule_rule/parallel_vectorize_unroll.py
index a79ea918670e..e9626c40e39c 100644
--- a/python/tvm/meta_schedule/schedule_rule/parallel_vectorize_unroll.py
+++ b/python/tvm/meta_schedule/schedule_rule/parallel_vectorize_unroll.py
@@ -18,7 +18,7 @@
 each block in a follow-up post processor"""
 from typing import List, Optional
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .schedule_rule import ScheduleRule
diff --git a/python/tvm/meta_schedule/schedule_rule/random_compute_location.py b/python/tvm/meta_schedule/schedule_rule/random_compute_location.py
index 2355b0bfa8e5..81de07afbbed 100644
--- a/python/tvm/meta_schedule/schedule_rule/random_compute_location.py
+++ b/python/tvm/meta_schedule/schedule_rule/random_compute_location.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """Rule that randomly select a compute-at location for a free block"""
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .schedule_rule import ScheduleRule
diff --git a/python/tvm/meta_schedule/schedule_rule/schedule_rule.py b/python/tvm/meta_schedule/schedule_rule/schedule_rule.py
index b71a4c7c3538..5684e68c715f 100644
--- a/python/tvm/meta_schedule/schedule_rule/schedule_rule.py
+++ b/python/tvm/meta_schedule/schedule_rule/schedule_rule.py
@@ -25,7 +25,7 @@
 
 # isort: on
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.runtime import Object
 from tvm.tir.schedule import BlockRV, Schedule
 
diff --git a/python/tvm/meta_schedule/search_strategy/evolutionary_search.py b/python/tvm/meta_schedule/search_strategy/evolutionary_search.py
index 44f32527fad9..1833ef23bda1 100644
--- a/python/tvm/meta_schedule/search_strategy/evolutionary_search.py
+++ b/python/tvm/meta_schedule/search_strategy/evolutionary_search.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """Evolutionary Search Strategy"""
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .search_strategy import SearchStrategy
diff --git a/python/tvm/meta_schedule/search_strategy/replay_func.py b/python/tvm/meta_schedule/search_strategy/replay_func.py
index f4660014241a..09e5c58d077a 100644
--- a/python/tvm/meta_schedule/search_strategy/replay_func.py
+++ b/python/tvm/meta_schedule/search_strategy/replay_func.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """Replay Trace Search Strategy"""
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .search_strategy import SearchStrategy
diff --git a/python/tvm/meta_schedule/search_strategy/replay_trace.py b/python/tvm/meta_schedule/search_strategy/replay_trace.py
index e24ad5a5219a..a25596524451 100644
--- a/python/tvm/meta_schedule/search_strategy/replay_trace.py
+++ b/python/tvm/meta_schedule/search_strategy/replay_trace.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """Replay Trace Search Strategy"""
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .search_strategy import SearchStrategy
diff --git a/python/tvm/meta_schedule/search_strategy/search_strategy.py b/python/tvm/meta_schedule/search_strategy/search_strategy.py
index 8822b097945e..ab4a6fb7b636 100644
--- a/python/tvm/meta_schedule/search_strategy/search_strategy.py
+++ b/python/tvm/meta_schedule/search_strategy/search_strategy.py
@@ -24,7 +24,7 @@
 from typing_extensions import Literal
 
 # isort: on
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.runtime import Object
 from tvm.tir.schedule import Schedule
 
diff --git a/python/tvm/meta_schedule/space_generator/post_order_apply.py b/python/tvm/meta_schedule/space_generator/post_order_apply.py
index 930e8a51dc61..eee9ea0d0e5d 100644
--- a/python/tvm/meta_schedule/space_generator/post_order_apply.py
+++ b/python/tvm/meta_schedule/space_generator/post_order_apply.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """Post Order Apply Space Generator."""
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .space_generator import (
diff --git a/python/tvm/meta_schedule/space_generator/schedule_fn.py b/python/tvm/meta_schedule/space_generator/schedule_fn.py
index 65956e843679..2cb1538a5abc 100644
--- a/python/tvm/meta_schedule/space_generator/schedule_fn.py
+++ b/python/tvm/meta_schedule/space_generator/schedule_fn.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """Union of meta Schedule design space generators."""
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .space_generator import (
diff --git a/python/tvm/meta_schedule/space_generator/space_generator.py b/python/tvm/meta_schedule/space_generator/space_generator.py
index c1d9765067bc..8c9effa6e656 100644
--- a/python/tvm/meta_schedule/space_generator/space_generator.py
+++ b/python/tvm/meta_schedule/space_generator/space_generator.py
@@ -24,7 +24,7 @@
 from typing_extensions import Literal
 
 # isort: on
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.ir import IRModule
 from tvm.runtime import Object
 from tvm.tir.schedule import Schedule
diff --git a/python/tvm/meta_schedule/space_generator/space_generator_union.py b/python/tvm/meta_schedule/space_generator/space_generator_union.py
index e3d8f441d1ef..f512f6535550 100644
--- a/python/tvm/meta_schedule/space_generator/space_generator_union.py
+++ b/python/tvm/meta_schedule/space_generator/space_generator_union.py
@@ -17,7 +17,7 @@
 """Union of meta Schedule design space generators."""
 from typing import List
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from .space_generator import (
diff --git a/python/tvm/meta_schedule/task_scheduler/gradient_based.py b/python/tvm/meta_schedule/task_scheduler/gradient_based.py
index 963de8711e10..7bac23bb3fad 100644
--- a/python/tvm/meta_schedule/task_scheduler/gradient_based.py
+++ b/python/tvm/meta_schedule/task_scheduler/gradient_based.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """Gradient Based Task Scheduler"""
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from ..logging import get_logger, get_logging_func
diff --git a/python/tvm/meta_schedule/task_scheduler/round_robin.py b/python/tvm/meta_schedule/task_scheduler/round_robin.py
index e5c7f14af424..6475b4102a1d 100644
--- a/python/tvm/meta_schedule/task_scheduler/round_robin.py
+++ b/python/tvm/meta_schedule/task_scheduler/round_robin.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """Round Robin Task Scheduler"""
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .. import _ffi_api
 from ..logging import get_logger, get_logging_func
diff --git a/python/tvm/meta_schedule/task_scheduler/task_scheduler.py b/python/tvm/meta_schedule/task_scheduler/task_scheduler.py
index d56d944474e9..9d6fec88b63b 100644
--- a/python/tvm/meta_schedule/task_scheduler/task_scheduler.py
+++ b/python/tvm/meta_schedule/task_scheduler/task_scheduler.py
@@ -22,7 +22,7 @@
 
 # isort: on
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.runtime import Object
 
 from .. import _ffi_api
diff --git a/python/tvm/meta_schedule/testing/validate_database.py b/python/tvm/meta_schedule/testing/validate_database.py
index ec7a4237546b..14ff32c0178a 100644
--- a/python/tvm/meta_schedule/testing/validate_database.py
+++ b/python/tvm/meta_schedule/testing/validate_database.py
@@ -25,7 +25,7 @@
 
 import tvm
 from tvm import meta_schedule as ms
-from tvm._ffi import get_global_func, register_func
+from tvm.ffi import get_global_func, register_func
 from tvm.ir import IRModule
 from tvm.support import describe
 from tvm.target import Target
@@ -469,9 +469,9 @@ def f_with_args_run_evaluator_common(
             number=evaluator_config.number,
             repeat=evaluator_config.repeat,
             min_repeat_ms=evaluator_config.min_repeat_ms,
-            f_preproc="cache_flush_cpu_non_first_arg"
-            if evaluator_config.enable_cpu_cache_flush
-            else "",
+            f_preproc=(
+                "cache_flush_cpu_non_first_arg" if evaluator_config.enable_cpu_cache_flush else ""
+            ),
         )
 
         repeated_costs: List[List[float]] = []
diff --git a/python/tvm/meta_schedule/tir_integration.py b/python/tvm/meta_schedule/tir_integration.py
index bffade49a072..b171c9711802 100644
--- a/python/tvm/meta_schedule/tir_integration.py
+++ b/python/tvm/meta_schedule/tir_integration.py
@@ -22,7 +22,7 @@
 
 # isort: on
 from tvm import ir, tir
-from tvm._ffi import register_func
+from tvm.ffi import register_func
 from tvm.target import Target
 from tvm.tir.expr import IntImm
 
diff --git a/python/tvm/meta_schedule/tune_context.py b/python/tvm/meta_schedule/tune_context.py
index 6f76452a57b5..5512b7a2682b 100644
--- a/python/tvm/meta_schedule/tune_context.py
+++ b/python/tvm/meta_schedule/tune_context.py
@@ -24,7 +24,7 @@
 # isort: on
 
 from tvm import IRModule
-from tvm._ffi import register_object, register_func
+from tvm.ffi import register_object, register_func
 from tvm.runtime import Object
 from tvm.target import Target
 from tvm.tir import PrimFunc, Schedule
diff --git a/python/tvm/meta_schedule/utils.py b/python/tvm/meta_schedule/utils.py
index 8932fcdc3eaa..61b32e1e324b 100644
--- a/python/tvm/meta_schedule/utils.py
+++ b/python/tvm/meta_schedule/utils.py
@@ -22,7 +22,7 @@
 
 import numpy as np  # type: ignore
 import psutil  # type: ignore
-from tvm._ffi import get_global_func, register_func
+from tvm.ffi import get_global_func, register_func
 from tvm.error import TVMError
 from tvm.ir import Array, IRModule, Map
 from tvm.rpc import RPCSession
diff --git a/python/tvm/relax/_ffi_api.py b/python/tvm/relax/_ffi_api.py
index a127e1c81378..db1ca055865a 100644
--- a/python/tvm/relax/_ffi_api.py
+++ b/python/tvm/relax/_ffi_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI API for Relax."""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("relax", __name__)
+tvm.ffi._init_api("relax", __name__)
diff --git a/python/tvm/relax/analysis/_ffi_api.py b/python/tvm/relax/analysis/_ffi_api.py
index 40ee05c3960d..fb44606f1122 100644
--- a/python/tvm/relax/analysis/_ffi_api.py
+++ b/python/tvm/relax/analysis/_ffi_api.py
@@ -14,6 +14,6 @@
 # KIND, either express or implied.  See the License for the
 # specific language governing permissions and limitations
 """FFI APIs"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("relax.analysis", __name__)
+tvm.ffi._init_api("relax.analysis", __name__)
diff --git a/python/tvm/relax/backend/_ffi_api.py b/python/tvm/relax/backend/_ffi_api.py
index d1378b2eacc2..17d7a18a338d 100644
--- a/python/tvm/relax/backend/_ffi_api.py
+++ b/python/tvm/relax/backend/_ffi_api.py
@@ -16,6 +16,6 @@
 # under the License.
 """FFI API for Relax backend."""
 
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("relax.backend", __name__)
+tvm.ffi._init_api("relax.backend", __name__)
diff --git a/python/tvm/relax/backend/metal/coreml.py b/python/tvm/relax/backend/metal/coreml.py
index b5caa688f221..139e5cc2b997 100644
--- a/python/tvm/relax/backend/metal/coreml.py
+++ b/python/tvm/relax/backend/metal/coreml.py
@@ -19,7 +19,7 @@
 
 import os
 import shutil
-import tvm._ffi
+import tvm.ffi
 from tvm.contrib import coreml_runtime
 from tvm.contrib.xcode import compile_coreml
 
@@ -144,7 +144,7 @@ def _conv2d_pattern(pattern_name):
         *default_unary_patterns(op_name="nn.avg_pool2d"),
         *conv2d_patterns(),
         *clip_patterns(),
-        *matmul_patterns()
+        *matmul_patterns(),
         # TODO(@tvm-team): enable when relax op is implemented
         # ("coreml.nn.batch_flatten", is_op("relax.nn.batch_flatten")(wildcard())),
     ]
@@ -463,7 +463,7 @@ def compile(self, out_dir):
         compile_coreml(model, self.model_name, out_dir)
 
 
-@tvm._ffi.register_func("relax.ext.coreml")
+@tvm.ffi.register_func("relax.ext.coreml")
 def coreml_compiler(funcs, options, constant_names):
     """
     Create a CoreML runtime from a Relax module.
diff --git a/python/tvm/relax/binding_rewrite.py b/python/tvm/relax/binding_rewrite.py
index a9f6d878ad0d..22215206ac4b 100644
--- a/python/tvm/relax/binding_rewrite.py
+++ b/python/tvm/relax/binding_rewrite.py
@@ -20,13 +20,13 @@
 from typing import Optional
 
 import tvm
-import tvm._ffi
+import tvm.ffi
 from tvm.runtime import Object
 from . import Binding, DataflowBlock, Expr, Function, Var
 from . import _ffi_api
 
 
-@tvm._ffi.register_object("relax.DataflowBlockRewrite")
+@tvm.ffi.register_object("relax.DataflowBlockRewrite")
 class DataflowBlockRewrite(Object):
     """
     A binding/statement-level dataflow block rewriter.
diff --git a/python/tvm/relax/block_builder.py b/python/tvm/relax/block_builder.py
index 37866840bd68..e09a9fab263a 100644
--- a/python/tvm/relax/block_builder.py
+++ b/python/tvm/relax/block_builder.py
@@ -100,7 +100,7 @@ def __exit__(self, ptype, value, trace):
         self._bb.end_scope()
 
 
-@tvm._ffi.register_object("relax.BlockBuilder")
+@tvm.ffi.register_object("relax.BlockBuilder")
 class BlockBuilder(Object):
     """A builder to build Relax IR for testing and dev.
 
diff --git a/python/tvm/relax/distributed/_ffi_api.py b/python/tvm/relax/distributed/_ffi_api.py
index 57411e82613f..6544a8d35572 100644
--- a/python/tvm/relax/distributed/_ffi_api.py
+++ b/python/tvm/relax/distributed/_ffi_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.relax.distributed"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("relax.distributed", __name__)
+tvm.ffi._init_api("relax.distributed", __name__)
diff --git a/python/tvm/relax/distributed/global_info.py b/python/tvm/relax/distributed/global_info.py
index 9639fd885cb6..3f549ecfa37e 100644
--- a/python/tvm/relax/distributed/global_info.py
+++ b/python/tvm/relax/distributed/global_info.py
@@ -26,7 +26,7 @@
 from . import _ffi_api as ffi
 
 
-@tvm._ffi.register_object("relax.distributed.DeviceMesh")
+@tvm.ffi.register_object("relax.distributed.DeviceMesh")
 class DeviceMesh(GlobalInfo):
     """Device mesh express a view of topology of devices,
        represented by an n-d matrix of device ids.
diff --git a/python/tvm/relax/distributed/struct_info.py b/python/tvm/relax/distributed/struct_info.py
index b5e258516d05..50087b98841a 100644
--- a/python/tvm/relax/distributed/struct_info.py
+++ b/python/tvm/relax/distributed/struct_info.py
@@ -33,7 +33,7 @@ class PlacementSpecKind(enum.IntEnum):
     kReplica = 1
 
 
-@tvm._ffi.register_object("relax.distributed.PlacementSpec")
+@tvm.ffi.register_object("relax.distributed.PlacementSpec")
 class PlacementSpec(Object):
     """Describes how data is distributed in one dimension of the device mesh
 
@@ -80,7 +80,7 @@ def replica() -> "PlacementSpec":
         return _ffi_api.Replica()
 
 
-@tvm._ffi.register_object("relax.distributed.Placement")
+@tvm.ffi.register_object("relax.distributed.Placement")
 class Placement(Object):
     """Describes how data is distributed in each dimension of the device mesh
 
@@ -110,7 +110,7 @@ def from_text(text: str) -> "Placement":
         return _ffi_api.PlacementFromText(text)
 
 
-@tvm._ffi.register_object("relax.DTensorStructInfo")
+@tvm.ffi.register_object("relax.DTensorStructInfo")
 class DTensorStructInfo(StructInfo):
     """StructInfo of a Distributed Tensor value.
 
diff --git a/python/tvm/relax/distributed/transform/_ffi_api.py b/python/tvm/relax/distributed/transform/_ffi_api.py
index d064ae2bb931..b694a67116d2 100644
--- a/python/tvm/relax/distributed/transform/_ffi_api.py
+++ b/python/tvm/relax/distributed/transform/_ffi_api.py
@@ -14,6 +14,6 @@
 # KIND, either express or implied.  See the License for the
 # specific language governing permissions and limitations
 """FFI APIs for tvm.relax.distributed.transform"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("relax.distributed.transform", __name__)
+tvm.ffi._init_api("relax.distributed.transform", __name__)
diff --git a/python/tvm/relax/dpl/_ffi.py b/python/tvm/relax/dpl/_ffi.py
index 6699e42bee63..72bf073bedfc 100644
--- a/python/tvm/relax/dpl/_ffi.py
+++ b/python/tvm/relax/dpl/_ffi.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """DataFlow Pattern Language FFI bindings."""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("relax.dpl", __name__)
+tvm.ffi._init_api("relax.dpl", __name__)
diff --git a/python/tvm/relax/dpl/pattern.py b/python/tvm/relax/dpl/pattern.py
index 3d076e8fad35..42486ee28948 100644
--- a/python/tvm/relax/dpl/pattern.py
+++ b/python/tvm/relax/dpl/pattern.py
@@ -23,7 +23,7 @@
 from typing import Dict, List, Optional, Tuple, Union
 
 import tvm
-import tvm._ffi as tvm_ffi
+import tvm.ffi as tvm_ffi
 from tvm.ir.container import Array
 from tvm.ir.expr import PrimExpr
 from tvm.ir.op import Op
diff --git a/python/tvm/relax/dpl/rewrite.py b/python/tvm/relax/dpl/rewrite.py
index 96c69e9266a2..a9782057c8fb 100644
--- a/python/tvm/relax/dpl/rewrite.py
+++ b/python/tvm/relax/dpl/rewrite.py
@@ -20,7 +20,7 @@
 
 from tvm.ir import IRModule
 from tvm.runtime import Object
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 
 from .pattern import DFPattern
 from .context import PatternContext
diff --git a/python/tvm/relax/exec_builder.py b/python/tvm/relax/exec_builder.py
index 699860786072..0939dabe16ed 100644
--- a/python/tvm/relax/exec_builder.py
+++ b/python/tvm/relax/exec_builder.py
@@ -56,7 +56,7 @@ def __exit__(self, ptype, value, trace):
         self.exit_callback()
 
 
-@tvm._ffi.register_object("relax.ExecBuilder")
+@tvm.ffi.register_object("relax.ExecBuilder")
 class ExecBuilder(Object):
     """A builder to emit instructions and build executable for the virtual machine."""
 
diff --git a/python/tvm/relax/expr.py b/python/tvm/relax/expr.py
index 950de9eac022..0fa8c4df88f8 100644
--- a/python/tvm/relax/expr.py
+++ b/python/tvm/relax/expr.py
@@ -22,11 +22,10 @@
 import numpy as _np  # type: ignore
 
 import tvm
-import tvm._ffi
+import tvm.ffi
 import tvm.ir
 import tvm.relax
 from tvm import DataType
-from tvm._ffi import base as _base
 from tvm.runtime import Object
 from tvm.runtime import ndarray as _nd
 
@@ -43,7 +42,7 @@
 GlobalVar = Union[tvm.ir.GlobalVar]
 
 
-@tvm._ffi.register_object("relax.Id")
+@tvm.ffi.register_object("relax.Id")
 class Id(Object):
     """Unique identifier(name) used in Var.
     Guaranteed to be stable across all passes.
@@ -528,7 +527,7 @@ def __getitem__(self, axis: Union[int, PrimExpr, Expr]) -> Expr:
         return tvm.relax.Call(op, [self.tensor, axis])
 
 
-@tvm._ffi.register_object("relax.expr.Call")
+@tvm.ffi.register_object("relax.expr.Call")
 class Call(ExprWithOp):
     """Function call node in Relax.
 
@@ -577,7 +576,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("relax.expr.If")
+@tvm.ffi.register_object("relax.expr.If")
 class If(ExprWithOp):
     """A conditional expression in Relax.
 
@@ -609,7 +608,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("relax.expr.Tuple")
+@tvm.ffi.register_object("relax.expr.Tuple")
 class Tuple(ExprWithOp):
     """Tuple expression that groups several fields together.
 
@@ -644,7 +643,7 @@ def __len__(self) -> int:
         return len(self.fields)
 
 
-@tvm._ffi.register_object("relax.expr.TupleGetItem")
+@tvm.ffi.register_object("relax.expr.TupleGetItem")
 class TupleGetItem(ExprWithOp):
     """Get index-th item from a tuple.
 
@@ -670,7 +669,7 @@ def __init__(self, tuple_value: Expr, index: int, span: Optional[Span] = None):
         )
 
 
-@tvm._ffi.register_object("relax.expr.ShapeExpr")
+@tvm.ffi.register_object("relax.expr.ShapeExpr")
 class ShapeExpr(ExprWithOp):
     """A shape expression which allows users to construct a shape containing PrimExpr.
 
@@ -708,7 +707,7 @@ def make_shape(shape: Union[List[Any], typing.Tuple[Any, ...]]) -> ShapeExpr:
     raise ValueError("Wrong type")
 
 
-@tvm._ffi.register_object("relax.expr.Constant")
+@tvm.ffi.register_object("relax.expr.Constant")
 class Constant(ExprWithOp):
     """Constant Tensor
 
@@ -742,7 +741,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("relax.expr.Var")
+@tvm.ffi.register_object("relax.expr.Var")
 class Var(ExprWithOp):
     """The variable class for all Relax bindings.
 
@@ -789,7 +788,7 @@ def name_hint(self) -> str:
         return name
 
 
-@tvm._ffi.register_object("relax.expr.DataflowVar")
+@tvm.ffi.register_object("relax.expr.DataflowVar")
 class DataflowVar(Var):
     """A sub-type of the variable node used to mark dataflow variables from
     normal visible "function local" bindings.
@@ -838,7 +837,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("relax.expr.PrimValue")
+@tvm.ffi.register_object("relax.expr.PrimValue")
 class PrimValue(Expr, Scriptable):
     """The prim expr representing the value."""
 
@@ -850,7 +849,7 @@ def __init__(self, value: Union[PrimExpr, int], span: Optional[Span] = None) ->
         self.__init_handle_by_constructor__(_ffi_api.PrimValue, value, span)  # type: ignore
 
 
-@tvm._ffi.register_object("relax.expr.StringImm")
+@tvm.ffi.register_object("relax.expr.StringImm")
 class StringImm(Expr, Scriptable):
     """Represent a string literal constant."""
 
@@ -861,7 +860,7 @@ def __init__(self, value: str, span: Optional[Span] = None) -> None:
         self.__init_handle_by_constructor__(_ffi_api.StringImm, value, span)  # type: ignore
 
 
-@tvm._ffi.register_object("relax.expr.DataTypeImm")
+@tvm.ffi.register_object("relax.expr.DataTypeImm")
 class DataTypeImm(Expr, Scriptable):
     """Represent a data type constant."""
 
@@ -872,7 +871,7 @@ def __init__(self, value: Union[DataType, str], span: Optional[Span] = None) ->
         self.__init_handle_by_constructor__(_ffi_api.DataTypeImm, value, span)  # type: ignore
 
 
-@tvm._ffi.register_object("relax.expr.Binding")
+@tvm.ffi.register_object("relax.expr.Binding")
 class Binding(Node, Scriptable):
     """The base class of a binding in Relax."""
 
@@ -880,7 +879,7 @@ class Binding(Node, Scriptable):
     span: Optional[Span]
 
 
-@tvm._ffi.register_object("relax.expr.MatchCast")
+@tvm.ffi.register_object("relax.expr.MatchCast")
 class MatchCast(Binding):
     """Runtime-match the value to the struct info.
 
@@ -912,7 +911,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("relax.expr.VarBinding")
+@tvm.ffi.register_object("relax.expr.VarBinding")
 class VarBinding(Binding):
     """Variable binding, bind he variable of the lhs with the rhs.
 
@@ -934,7 +933,7 @@ def __init__(self, var: Var, value: Expr, span: Optional[Span] = None) -> None:
         self.__init_handle_by_constructor__(_ffi_api.VarBinding, var, value, span)  # type: ignore
 
 
-@tvm._ffi.register_object("relax.expr.BindingBlock")
+@tvm.ffi.register_object("relax.expr.BindingBlock")
 class BindingBlock(Node, Scriptable):
     """base class of binding block, bindings inside can be impure
     (with side effect or control flow)"""
@@ -946,7 +945,7 @@ def __init__(self, bindings: List[Binding], span: Optional[Span] = None) -> None
         self.__init_handle_by_constructor__(_ffi_api.BindingBlock, bindings, span)  # type: ignore
 
 
-@tvm._ffi.register_object("relax.expr.DataflowBlock")
+@tvm.ffi.register_object("relax.expr.DataflowBlock")
 class DataflowBlock(BindingBlock):
     """dataflow block, bindings inside are pure (no side effect and no control flow)"""
 
@@ -958,7 +957,7 @@ def __init__(self, bindings: List[Binding], span: Optional[Span] = None) -> None
         self.__init_handle_by_constructor__(_ffi_api.DataflowBlock, bindings, span)  # type: ignore
 
 
-@tvm._ffi.register_object("relax.expr.SeqExpr")
+@tvm.ffi.register_object("relax.expr.SeqExpr")
 class SeqExpr(ExprWithOp):
     """A sequence of binding blocks followed by an expression."""
 
@@ -970,7 +969,7 @@ def __init__(self, blocks: List[BindingBlock], body: Expr, span: Optional[Span]
         self.__init_handle_by_constructor__(_ffi_api.SeqExpr, blocks, body, span)  # type: ignore
 
 
-@tvm._ffi.register_object("relax.expr.Function")
+@tvm.ffi.register_object("relax.expr.Function")
 class Function(BaseFunc, Scriptable):
     """A Relax function."""
 
@@ -1109,7 +1108,7 @@ def inline_functions(
         return _ffi_api.FunctionInlineFunctions(self, function_map)  # type: ignore
 
 
-@tvm._ffi.register_object("relax.expr.ExternFunc")
+@tvm.ffi.register_object("relax.expr.ExternFunc")
 class ExternFunc(BaseFunc, ExprWithOp):
     """extern function, which represents a PackedFunc."""
 
@@ -1154,7 +1153,7 @@ def const(
     - bool maps to "bool"
     - other using the same default rule as numpy.
     """
-    if isinstance(value, (_base.numeric_types, (bool, list))):
+    if isinstance(value, (Number, (bool, list))):
         value = _np.array(value, dtype=dtype)
 
     if not dtype:
diff --git a/python/tvm/relax/expr_functor.py b/python/tvm/relax/expr_functor.py
index 12647733cbd7..49d3b14505ba 100644
--- a/python/tvm/relax/expr_functor.py
+++ b/python/tvm/relax/expr_functor.py
@@ -261,7 +261,7 @@ def visit_var_def(self, var: Var):
             raise TypeError("Invalid type: {0}".format(type(var)))
 
 
-@tvm._ffi.register_object("expr_functor.PyExprVisitor")
+@tvm.ffi.register_object("expr_functor.PyExprVisitor")
 class _PyExprVisitor(Object):
     """
     A TVM object to support customization of ExprVisitor on the python side.
@@ -781,7 +781,7 @@ def visit_span(self, span: Span) -> None:
         return _ffi_api.ExprVisitorVisitSpan(self._outer(), span)  # type: ignore
 
 
-@tvm._ffi.register_object("expr_functor.PyExprMutator")
+@tvm.ffi.register_object("expr_functor.PyExprMutator")
 class _PyExprMutator(Object):
     """
     A TVM object to support customization of ExprMutator on the python side.
diff --git a/python/tvm/relax/op/_ffi_api.py b/python/tvm/relax/op/_ffi_api.py
index 8dc6a1b4fbb0..1d16a024d1d4 100644
--- a/python/tvm/relax/op/_ffi_api.py
+++ b/python/tvm/relax/op/_ffi_api.py
@@ -14,6 +14,6 @@
 # KIND, either express or implied.  See the License for the
 # specific language governing permissions and limitations
 """FFI APIs for tvm.relax.op"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("relax.op", __name__)
+tvm.ffi._init_api("relax.op", __name__)
diff --git a/python/tvm/relax/op/_op_gradient.py b/python/tvm/relax/op/_op_gradient.py
index 6878f9733163..41eaa5de5008 100644
--- a/python/tvm/relax/op/_op_gradient.py
+++ b/python/tvm/relax/op/_op_gradient.py
@@ -21,7 +21,7 @@
 from typing import List
 
 from tvm import relax
-from tvm._ffi.base import TVMError
+from tvm.base import TVMError
 from tvm.arith import Analyzer
 from tvm.relax.struct_info import ShapeStructInfo
 
@@ -1090,7 +1090,7 @@ def log_softmax_grad(
         Returns `[y_grad - sum(y_grad, axis, keepdims=True) * exp(y)]`
     """
     y_exp = exp(orig_var)
-    return [(output_grad - sum(output_grad, orig_call.attrs.axis, True) * y_exp)]
+    return [output_grad - sum(output_grad, orig_call.attrs.axis, True) * y_exp]
 
 
 @register_gradient("relax.nn.cross_entropy_with_logits")
diff --git a/python/tvm/relax/op/builtin/_ffi_api.py b/python/tvm/relax/op/builtin/_ffi_api.py
index 42fe8cb65234..a7f48af57697 100644
--- a/python/tvm/relax/op/builtin/_ffi_api.py
+++ b/python/tvm/relax/op/builtin/_ffi_api.py
@@ -14,6 +14,6 @@
 # KIND, either express or implied.  See the License for the
 # specific language governing permissions and limitations
 """FFI APIs for tvm.relax.op.builtin"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("relax.op.builtin", __name__)
+tvm.ffi._init_api("relax.op.builtin", __name__)
diff --git a/python/tvm/relax/op/ccl/_ffi_api.py b/python/tvm/relax/op/ccl/_ffi_api.py
index cdf468781061..bf605aae6ab0 100644
--- a/python/tvm/relax/op/ccl/_ffi_api.py
+++ b/python/tvm/relax/op/ccl/_ffi_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """Operators serving for Collective Communications Library (CCL) operators"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("relax.op.ccl", __name__)
+tvm.ffi._init_api("relax.op.ccl", __name__)
diff --git a/python/tvm/relax/op/distributed/_ffi_api.py b/python/tvm/relax/op/distributed/_ffi_api.py
index 9b1b4d68d6da..394cb8c262b2 100644
--- a/python/tvm/relax/op/distributed/_ffi_api.py
+++ b/python/tvm/relax/op/distributed/_ffi_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.relax.op.distributed"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("relax.op.dist", __name__)
+tvm.ffi._init_api("relax.op.dist", __name__)
diff --git a/python/tvm/relax/op/grad/_ffi_api.py b/python/tvm/relax/op/grad/_ffi_api.py
index 9b819dd4df29..415d590f01f0 100644
--- a/python/tvm/relax/op/grad/_ffi_api.py
+++ b/python/tvm/relax/op/grad/_ffi_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.relax.op.grad"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("relax.op.grad", __name__)
+tvm.ffi._init_api("relax.op.grad", __name__)
diff --git a/python/tvm/relax/op/image/_ffi_api.py b/python/tvm/relax/op/image/_ffi_api.py
index e666203ae7ff..8c813231f9a0 100644
--- a/python/tvm/relax/op/image/_ffi_api.py
+++ b/python/tvm/relax/op/image/_ffi_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """Constructor APIs"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("relax.op.image", __name__)
+tvm.ffi._init_api("relax.op.image", __name__)
diff --git a/python/tvm/relax/op/memory/_ffi_api.py b/python/tvm/relax/op/memory/_ffi_api.py
index 475de481b22e..fb829b7db953 100644
--- a/python/tvm/relax/op/memory/_ffi_api.py
+++ b/python/tvm/relax/op/memory/_ffi_api.py
@@ -14,6 +14,6 @@
 # KIND, either express or implied.  See the License for the
 # specific language governing permissions and limitations
 """FFI APIs for tvm.relax.op.memory"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("relax.op.memory", __name__)
+tvm.ffi._init_api("relax.op.memory", __name__)
diff --git a/python/tvm/relax/op/nn/_ffi_api.py b/python/tvm/relax/op/nn/_ffi_api.py
index 1785345ac1b1..b5f735127ec2 100644
--- a/python/tvm/relax/op/nn/_ffi_api.py
+++ b/python/tvm/relax/op/nn/_ffi_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """Constructor APIs"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("relax.op.nn", __name__)
+tvm.ffi._init_api("relax.op.nn", __name__)
diff --git a/python/tvm/relax/op/op_attrs.py b/python/tvm/relax/op/op_attrs.py
index fe527e38e8a8..4869ff252065 100644
--- a/python/tvm/relax/op/op_attrs.py
+++ b/python/tvm/relax/op/op_attrs.py
@@ -16,179 +16,179 @@
 # under the License.
 """The attributes node used for Relax operators"""
 from tvm.ir import Attrs
-import tvm._ffi
+import tvm.ffi
 
 
-@tvm._ffi.register_object("relax.attrs.CallTIRWithGradAttrs")
+@tvm.ffi.register_object("relax.attrs.CallTIRWithGradAttrs")
 class CallTIRWithGradAttrs(Attrs):
     """Attributes used in call_tir_with_grad operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.InitAttrs")
+@tvm.ffi.register_object("relax.attrs.InitAttrs")
 class InitAttrs(Attrs):
     """Attributes used in full/full_like, ones/ones_like, and zeros/zeros_like operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.TriluAttrs")
+@tvm.ffi.register_object("relax.attrs.TriluAttrs")
 class TriluAttrs(Attrs):
     """Attributes used in tril and triu operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.AstypeAttrs")
+@tvm.ffi.register_object("relax.attrs.AstypeAttrs")
 class AstypeAttrs(Attrs):
     """Attributes used in astype operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.TakeAttrs")
+@tvm.ffi.register_object("relax.attrs.TakeAttrs")
 class TakeAttrs(Attrs):
     """Attributes used in take operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.StridedSliceAttrs")
+@tvm.ffi.register_object("relax.attrs.StridedSliceAttrs")
 class StridedSliceAttrs(Attrs):
     """Attributes used in strided_slice operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.MatmulAttrs")
+@tvm.ffi.register_object("relax.attrs.MatmulAttrs")
 class MatmulAttrs(Attrs):
     """Attributes for matmul operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.Conv2DAttrs")
+@tvm.ffi.register_object("relax.attrs.Conv2DAttrs")
 class Conv2DAttrs(Attrs):
     """Attributes for nn.conv2d"""
 
 
-@tvm._ffi.register_object("relax.attrs.Conv3DAttrs")
+@tvm.ffi.register_object("relax.attrs.Conv3DAttrs")
 class Conv3DAttrs(Attrs):
     """Attributes for nn.conv3d"""
 
 
-@tvm._ffi.register_object("relax.attrs.Conv2DTransposeAttrs")
+@tvm.ffi.register_object("relax.attrs.Conv2DTransposeAttrs")
 class Conv2DTransposeAttrs(Attrs):
     """Attributes for nn.conv2d_transpose"""
 
 
-@tvm._ffi.register_object("relax.attrs.Pool2DAttrs")
+@tvm.ffi.register_object("relax.attrs.Pool2DAttrs")
 class Pool2DAttrs(Attrs):
     """Attributes for nn.max_pool2d"""
 
 
-@tvm._ffi.register_object("relax.attrs.AdaptivePool2DAttrs")
+@tvm.ffi.register_object("relax.attrs.AdaptivePool2DAttrs")
 class AdaptivePool2DAttrs(Attrs):
     """Attributes for 2d adaptive pool operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.SoftmaxAttrs")
+@tvm.ffi.register_object("relax.attrs.SoftmaxAttrs")
 class SoftmaxAttrs(Attrs):
     """Attributes for nn.softmax"""
 
 
-@tvm._ffi.register_object("relax.attrs.BatchNormAttrs")
+@tvm.ffi.register_object("relax.attrs.BatchNormAttrs")
 class BatchNormAttrs(Attrs):
     """Attributes used in batch_norm operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.LayerNormAttrs")
+@tvm.ffi.register_object("relax.attrs.LayerNormAttrs")
 class LayerNormAttrs(Attrs):
     """Attributes used in layer_norm operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.DropoutAttrs")
+@tvm.ffi.register_object("relax.attrs.DropoutAttrs")
 class DropoutAttrs(Attrs):
     """Attributes for dropout operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.StatisticalAttrs")
+@tvm.ffi.register_object("relax.attrs.StatisticalAttrs")
 class StatisticalAttrs(Attrs):
     """Attributes used in statistical operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.ConcatAttrs")
+@tvm.ffi.register_object("relax.attrs.ConcatAttrs")
 class ConcatAttrs(Attrs):
     """Attributes for concat operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.ExpandDimsAttrs")
+@tvm.ffi.register_object("relax.attrs.ExpandDimsAttrs")
 class ExpandDimsAttrs(Attrs):
     """Attributes for expand_dims operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.PermuteDimsAttrs")
+@tvm.ffi.register_object("relax.attrs.PermuteDimsAttrs")
 class PermuteDimsAttrs(Attrs):
     """Attributes for permute_dims operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.SortAttrs")
+@tvm.ffi.register_object("relax.attrs.SortAttrs")
 class SortAttrs(Attrs):
     """Attributes for sort operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.ArgsortAttrs")
+@tvm.ffi.register_object("relax.attrs.ArgsortAttrs")
 class ArgsortAttrs(Attrs):
     """Attributes for argsort operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.SplitAttrs")
+@tvm.ffi.register_object("relax.attrs.SplitAttrs")
 class SplitAttrs(Attrs):
     """Attributes used in split operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.SqueezeAttrs")
+@tvm.ffi.register_object("relax.attrs.SqueezeAttrs")
 class SqueezeAttrs(Attrs):
     """Attributes for squeeze operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.StackAttrs")
+@tvm.ffi.register_object("relax.attrs.StackAttrs")
 class StackAttrs(Attrs):
     """Attributes for concat operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.IndexPutAttrs")
+@tvm.ffi.register_object("relax.attrs.IndexPutAttrs")
 class IndexPutAttrs(Attrs):
     """Attributes for index_put operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.LayoutTransformAttrs")
+@tvm.ffi.register_object("relax.attrs.LayoutTransformAttrs")
 class LayoutTransformAttrs(Attrs):
     """Attributes used in layout_transform operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.Resize2DAttrs")
+@tvm.ffi.register_object("relax.attrs.Resize2DAttrs")
 class Resize2DAttrs(Attrs):
     """Attributes used in image resize2d operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.ArgmaxArgminAttrs")
+@tvm.ffi.register_object("relax.attrs.ArgmaxArgminAttrs")
 class ArgmaxArgminAttrs(Attrs):
     """Attributes for argmax/argmin operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.RepeatAttrs")
+@tvm.ffi.register_object("relax.attrs.RepeatAttrs")
 class RepeatAttrs(Attrs):
     """Attributes for repeat operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.TileAttrs")
+@tvm.ffi.register_object("relax.attrs.TileAttrs")
 class TileAttrs(Attrs):
     """Attributes for tile operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.ScanopAttrs")
+@tvm.ffi.register_object("relax.attrs.ScanopAttrs")
 class ScanopAttrs(Attrs):
     """Attributes for scan operators"""
 
 
-@tvm._ffi.register_object("relax.attrs.TopKAttrs")
+@tvm.ffi.register_object("relax.attrs.TopKAttrs")
 class TopKAttrs(Attrs):
     """Attributes for topk operators"""
 
 
-@tvm._ffi.register_object("relax.attrs.EinsumAttrs")
+@tvm.ffi.register_object("relax.attrs.EinsumAttrs")
 class EinsumAttrs(Attrs):
     """Attributes for einsum operator"""
 
 
-@tvm._ffi.register_object("relax.attrs.FlipAttrs")
+@tvm.ffi.register_object("relax.attrs.FlipAttrs")
 class FlipAttrs(Attrs):
     """Attributes for flip operator"""
diff --git a/python/tvm/relax/op/vm/_ffi_api.py b/python/tvm/relax/op/vm/_ffi_api.py
index 786b73c76c64..f3b6cea13b67 100644
--- a/python/tvm/relax/op/vm/_ffi_api.py
+++ b/python/tvm/relax/op/vm/_ffi_api.py
@@ -14,6 +14,6 @@
 # KIND, either express or implied.  See the License for the
 # specific language governing permissions and limitations
 """FFI APIs for tvm.relax.op.vm"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("relax.op.vm", __name__)
+tvm.ffi._init_api("relax.op.vm", __name__)
diff --git a/python/tvm/relax/struct_info.py b/python/tvm/relax/struct_info.py
index de1b1ac3bfc3..c143f098328c 100644
--- a/python/tvm/relax/struct_info.py
+++ b/python/tvm/relax/struct_info.py
@@ -18,7 +18,7 @@
 """The struct info nodes of the Relax language."""
 from typing import List, Optional, Union
 
-import tvm._ffi
+import tvm.ffi
 import tvm
 
 from tvm.ir import Span, EnvFunc, Array, VDevice
@@ -29,7 +29,7 @@
 from . import _ffi_api, ty, expr
 
 
-@tvm._ffi.register_object("relax.ObjectStructInfo")
+@tvm.ffi.register_object("relax.ObjectStructInfo")
 class ObjectStructInfo(StructInfo):
     """StructInfo of an Object."""
 
@@ -37,7 +37,7 @@ def __init__(self, span: Span = None) -> None:
         self.__init_handle_by_constructor__(_ffi_api.ObjectStructInfo, span)  # type: ignore
 
 
-@tvm._ffi.register_object("relax.PrimStructInfo")
+@tvm.ffi.register_object("relax.PrimStructInfo")
 class PrimStructInfo(StructInfo):
     """StructInfo of a primitive POD value.
 
@@ -107,7 +107,7 @@ def __init__(
             )  # type: ignore
 
 
-@tvm._ffi.register_object("relax.ShapeStructInfo")
+@tvm.ffi.register_object("relax.ShapeStructInfo")
 class ShapeStructInfo(StructInfo):
     """StructInfo of a shape value.
 
@@ -136,7 +136,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("relax.TensorStructInfo")
+@tvm.ffi.register_object("relax.TensorStructInfo")
 class TensorStructInfo(StructInfo):
     """StructInfo of a Tensor value.
 
@@ -180,7 +180,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("relax.TupleStructInfo")
+@tvm.ffi.register_object("relax.TupleStructInfo")
 class TupleStructInfo(StructInfo):
     """StructInfo of a Tuple value.
 
@@ -197,7 +197,7 @@ def __init__(self, fields: List[StructInfo], span: Span = None) -> None:
         self.__init_handle_by_constructor__(_ffi_api.TupleStructInfo, fields, span)  # type: ignore
 
 
-@tvm._ffi.register_object("relax.FuncStructInfo")
+@tvm.ffi.register_object("relax.FuncStructInfo")
 class FuncStructInfo(StructInfo):
     """StructInfo of a function value.
 
diff --git a/python/tvm/relax/testing/transform.py b/python/tvm/relax/testing/transform.py
index 02c79bd4fa6e..198b07e51ea7 100644
--- a/python/tvm/relax/testing/transform.py
+++ b/python/tvm/relax/testing/transform.py
@@ -70,7 +70,7 @@ def dataflow_alias_analysis(
     return res_alias_sets, res_tuple_map  # type: ignore
 
 
-@tvm._ffi.register_object("relax.transform.InplaceOpportunity")
+@tvm.ffi.register_object("relax.transform.InplaceOpportunity")
 class InplaceOpportunity(Object):
     """
     Represents an opportunity to make a binding in-place. Exposed only for testing;
diff --git a/python/tvm/relax/training/_ffi_api.py b/python/tvm/relax/training/_ffi_api.py
index 70cb83fc0e6f..9b7dbcdee748 100644
--- a/python/tvm/relax/training/_ffi_api.py
+++ b/python/tvm/relax/training/_ffi_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.relax.training"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("relax.training", __name__)
+tvm.ffi._init_api("relax.training", __name__)
diff --git a/python/tvm/relax/training/utils.py b/python/tvm/relax/training/utils.py
index 4d1a32177227..dd433435e278 100644
--- a/python/tvm/relax/training/utils.py
+++ b/python/tvm/relax/training/utils.py
@@ -21,7 +21,7 @@
 
 import tvm
 from tvm import relax
-from tvm._ffi.registry import register_func
+from tvm.ffi.registry import register_func
 from tvm.relax.block_builder import BlockBuilder
 
 from ..expr import Function, Var, Call
diff --git a/python/tvm/relax/transform/_ffi_api.py b/python/tvm/relax/transform/_ffi_api.py
index 667aa62c2c95..3c4387a3cbb8 100644
--- a/python/tvm/relax/transform/_ffi_api.py
+++ b/python/tvm/relax/transform/_ffi_api.py
@@ -14,6 +14,6 @@
 # KIND, either express or implied.  See the License for the
 # specific language governing permissions and limitations
 """FFI APIs for tvm.transform"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("relax.transform", __name__)
+tvm.ffi._init_api("relax.transform", __name__)
diff --git a/python/tvm/relax/transform/transform.py b/python/tvm/relax/transform/transform.py
index 4ee53b032918..8e74f0897720 100644
--- a/python/tvm/relax/transform/transform.py
+++ b/python/tvm/relax/transform/transform.py
@@ -36,14 +36,14 @@
 from ..expr import Var
 
 
-@tvm._ffi.register_object("relax.FunctionPass")
+@tvm.ffi.register_object("relax.FunctionPass")
 class FunctionPass(tvm.ir.transform.Pass):
     """A pass that works on each tvm.relax.Function in a module. A function
     pass class should be created through `function_pass`.
     """
 
 
-@tvm._ffi.register_object("relax.DataflowBlockPass")
+@tvm.ffi.register_object("relax.DataflowBlockPass")
 class DataflowBlockPass(tvm.ir.transform.Pass):
     """A pass that works on each tvm.relax.DataflowBlock in a module."""
 
@@ -820,7 +820,7 @@ def FuseTIR() -> tvm.ir.transform.Pass:
     return _ffi_api.FuseTIR()  # type: ignore
 
 
-@tvm._ffi.register_object("relax.transform.PatternCheckContext")
+@tvm.ffi.register_object("relax.transform.PatternCheckContext")
 class PatternCheckContext(Object):
     """
     The input of check function `FusionPattern.check`.
@@ -854,7 +854,7 @@ class PatternCheckContext(Object):
     value_to_bound_var: Mapping[Expr, Var]
 
 
-@tvm._ffi.register_object("relax.transform.FusionPattern")
+@tvm.ffi.register_object("relax.transform.FusionPattern")
 class FusionPattern(Object):
     """
     The pattern used by `FuseOpsByPattern`. It's mainly DFPattern but with other
diff --git a/python/tvm/relax/transform/tuning_api/_ffi_api.py b/python/tvm/relax/transform/tuning_api/_ffi_api.py
index f31522d02595..54caece700ef 100644
--- a/python/tvm/relax/transform/tuning_api/_ffi_api.py
+++ b/python/tvm/relax/transform/tuning_api/_ffi_api.py
@@ -14,6 +14,6 @@
 # KIND, either express or implied.  See the License for the
 # specific language governing permissions and limitations
 """FFI APIs for relax.tuning_api"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("relax.tuning_api", __name__)
+tvm.ffi._init_api("relax.tuning_api", __name__)
diff --git a/python/tvm/relax/transform/tuning_api/database.py b/python/tvm/relax/transform/tuning_api/database.py
index 9477e142bad4..cbc103423b0f 100644
--- a/python/tvm/relax/transform/tuning_api/database.py
+++ b/python/tvm/relax/transform/tuning_api/database.py
@@ -24,7 +24,7 @@
 from tvm.meta_schedule.database import Workload
 from tvm.tir.schedule.trace import JSON_TYPE
 from tvm.target import Target
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from .primitives import Trace
 from . import _ffi_api
 
diff --git a/python/tvm/relax/transform/tuning_api/default_functions.py b/python/tvm/relax/transform/tuning_api/default_functions.py
index bfb73e11a06b..cbd71a06e608 100644
--- a/python/tvm/relax/transform/tuning_api/default_functions.py
+++ b/python/tvm/relax/transform/tuning_api/default_functions.py
@@ -33,7 +33,7 @@
     LocalRunner,
     RunnerInput,
 )
-from tvm._ffi.registry import register_func
+from tvm.ffi.registry import register_func
 from .primitives import Knob, Trace
 
 logger = logging.getLogger("TuningAPI")  # pylint: disable=invalid-name
diff --git a/python/tvm/relax/transform/tuning_api/primitives.py b/python/tvm/relax/transform/tuning_api/primitives.py
index 67b81ba7e99c..fdc3769f3e5a 100644
--- a/python/tvm/relax/transform/tuning_api/primitives.py
+++ b/python/tvm/relax/transform/tuning_api/primitives.py
@@ -23,7 +23,7 @@
 from tvm.ir.module import IRModule
 from tvm.relax import Expr
 from tvm.tir.schedule.trace import JSON_TYPE, _json_from_tvm
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from . import _ffi_api
 
 logger = logging.getLogger("TuningAPI")  # pylint: disable=invalid-name
diff --git a/python/tvm/relax/ty.py b/python/tvm/relax/ty.py
index b0afb069435a..426695c9f1fe 100644
--- a/python/tvm/relax/ty.py
+++ b/python/tvm/relax/ty.py
@@ -16,13 +16,13 @@
 # under the License.
 # pylint: disable=invalid-name, unused-import
 """The type nodes of the Relax language."""
-import tvm._ffi
+import tvm.ffi
 from tvm.ir import Type, TupleType, FuncType, Span
 
 from . import _ffi_api
 
 
-@tvm._ffi.register_object("relax.ShapeType")
+@tvm.ffi.register_object("relax.ShapeType")
 class ShapeType(Type):
     """The type of shape in Relax.
 
@@ -37,7 +37,7 @@ def __init__(self, ndim: int = -1, span: Span = None) -> None:
         self.__init_handle_by_constructor__(_ffi_api.ShapeType, ndim, span)  # type: ignore
 
 
-@tvm._ffi.register_object("relax.ObjectType")
+@tvm.ffi.register_object("relax.ObjectType")
 class ObjectType(Type):
     """A type that corresponds to tvm::runtime::Object, is base of all possible object
     values in TVM."""
@@ -46,7 +46,7 @@ def __init__(self, span: Span = None) -> None:
         self.__init_handle_by_constructor__(_ffi_api.ObjectType, span)  # type: ignore
 
 
-@tvm._ffi.register_object("relax.DynTensorType")
+@tvm.ffi.register_object("relax.DynTensorType")
 class TensorType(Type):
     """A dynamic tensor type in Relax.
 
@@ -65,7 +65,7 @@ def __init__(self, ndim=-1, dtype="float32", span: Span = None) -> None:
         self.__init_handle_by_constructor__(_ffi_api.TensorType, ndim, dtype, span)  # type: ignore
 
 
-@tvm._ffi.register_object("relax.PackedFuncType")
+@tvm.ffi.register_object("relax.PackedFuncType")
 class PackedFuncType(Type):
     """The type of ExternFunc in Relax."""
 
diff --git a/python/tvm/rpc/_ffi_api.py b/python/tvm/rpc/_ffi_api.py
index 1a7cc739b5c1..3b77e7a552e3 100644
--- a/python/tvm/rpc/_ffi_api.py
+++ b/python/tvm/rpc/_ffi_api.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.rpc"""
-import tvm._ffi
+import tvm.ffi
 
 
-tvm._ffi._init_api("rpc", __name__)
+tvm.ffi._init_api("rpc", __name__)
diff --git a/python/tvm/rpc/base.py b/python/tvm/rpc/base.py
index 8be74fe88c96..235934212c0e 100644
--- a/python/tvm/rpc/base.py
+++ b/python/tvm/rpc/base.py
@@ -25,7 +25,7 @@
 import random
 import logging
 
-from .._ffi.base import py_str
+from ..base import py_str
 
 # Magic header for RPC data plane
 RPC_MAGIC = 0xFF271
diff --git a/python/tvm/rpc/client.py b/python/tvm/rpc/client.py
index f9e677e49e98..ea78b0d7d418 100644
--- a/python/tvm/rpc/client.py
+++ b/python/tvm/rpc/client.py
@@ -22,8 +22,8 @@
 import struct
 import time
 
-import tvm._ffi
-from tvm._ffi.base import TVMError
+import tvm.ffi
+from tvm.base import TVMError
 from tvm.contrib import utils
 from tvm.runtime import ndarray as nd
 from tvm.runtime import Device
@@ -263,7 +263,7 @@ def __init__(self):
         RPCSession.__init__(self, _ffi_api.LocalSession())
 
 
-@tvm._ffi.register_func("rpc.PopenSession")
+@tvm.ffi.register_func("rpc.PopenSession")
 def _popen_session(binary):
     temp = utils.tempdir()
 
diff --git a/python/tvm/rpc/minrpc.py b/python/tvm/rpc/minrpc.py
index 842ba7e49814..2e46965a2050 100644
--- a/python/tvm/rpc/minrpc.py
+++ b/python/tvm/rpc/minrpc.py
@@ -16,7 +16,7 @@
 # under the License.
 """Utils to path."""
 import os
-from tvm._ffi import libinfo
+from tvm import libinfo
 from tvm.contrib import cc
 
 
diff --git a/python/tvm/rpc/proxy.py b/python/tvm/rpc/proxy.py
index 7997274d3f5a..5987d3f101ba 100644
--- a/python/tvm/rpc/proxy.py
+++ b/python/tvm/rpc/proxy.py
@@ -47,7 +47,7 @@
 from . import base
 from .base import TrackerCode
 from .server import _server_env
-from .._ffi.base import py_str
+from ..base import py_str
 
 
 class ForwardHandler(object):
diff --git a/python/tvm/rpc/server.py b/python/tvm/rpc/server.py
index d334fe0cf7ba..eb345260e300 100644
--- a/python/tvm/rpc/server.py
+++ b/python/tvm/rpc/server.py
@@ -36,10 +36,10 @@
 import time
 import errno
 import sys
-import tvm._ffi
+import tvm.ffi
 
-from tvm._ffi.base import py_str
-from tvm._ffi.libinfo import find_lib_path
+from tvm.base import py_str
+from tvm.libinfo import find_lib_path
 from tvm.runtime.module import load_module as _load_module
 from tvm.contrib import utils
 from tvm.contrib.popen_pool import PopenWorker
@@ -70,11 +70,11 @@ def _server_env(load_library, work_path=None):
         temp = utils.tempdir()
 
     # pylint: disable=unused-variable
-    @tvm._ffi.register_func("tvm.rpc.server.workpath", override=True)
+    @tvm.ffi.register_func("tvm.rpc.server.workpath", override=True)
     def get_workpath(path):
         return temp.relpath(path)
 
-    @tvm._ffi.register_func("tvm.rpc.server.load_module", override=True)
+    @tvm.ffi.register_func("tvm.rpc.server.load_module", override=True)
     def load_module(file_name):
         """Load module from remote side."""
         path = temp.relpath(file_name)
@@ -82,7 +82,7 @@ def load_module(file_name):
         logger.info("load_module %s", path)
         return m
 
-    @tvm._ffi.register_func("tvm.rpc.server.download_linked_module", override=True)
+    @tvm.ffi.register_func("tvm.rpc.server.download_linked_module", override=True)
     def download_linked_module(file_name):
         """Load module from remote side."""
         # pylint: disable=import-outside-toplevel
diff --git a/python/tvm/rpc/tracker.py b/python/tvm/rpc/tracker.py
index cd32c4a8221c..4ceb338404f7 100644
--- a/python/tvm/rpc/tracker.py
+++ b/python/tvm/rpc/tracker.py
@@ -60,7 +60,7 @@
         f"RPCTracker module requires tornado package {error_msg}. Try 'pip install tornado'."
     )
 
-from .._ffi.base import py_str
+from ..base import py_str
 from . import base
 from .base import RPC_TRACKER_MAGIC, TrackerCode
 
diff --git a/python/tvm/runtime/_ffi_api.py b/python/tvm/runtime/_ffi_api.py
index f0d1bdeb0a76..71f96983ee18 100644
--- a/python/tvm/runtime/_ffi_api.py
+++ b/python/tvm/runtime/_ffi_api.py
@@ -15,8 +15,8 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.runtime"""
-import tvm._ffi
+import tvm.ffi
 
 # Exports functions registered via TVM_FFI_REGISTER_GLOBAL with the "runtime" prefix.
 # e.g. TVM_FFI_REGISTER_GLOBAL("runtime.ModuleLoadFromFile")
-tvm._ffi._init_api("runtime", __name__)
+tvm.ffi._init_api("runtime", __name__)
diff --git a/python/tvm/runtime/_ffi_node_api.py b/python/tvm/runtime/_ffi_node_api.py
index 600206d42583..493dfceab59e 100644
--- a/python/tvm/runtime/_ffi_node_api.py
+++ b/python/tvm/runtime/_ffi_node_api.py
@@ -17,7 +17,7 @@
 
 # pylint: disable=invalid-name, unused-argument
 """FFI for tvm.node"""
-import tvm._ffi
+import tvm.ffi
 import tvm.ffi.core
 
 
@@ -47,4 +47,4 @@ def LoadJSON(json_str):
 
 # Exports functions registered via TVM_FFI_REGISTER_GLOBAL with the "node" prefix.
 # e.g. TVM_FFI_REGISTER_GLOBAL("node.AsRepr")
-tvm._ffi._init_api("node", __name__)
+tvm.ffi._init_api("node", __name__)
diff --git a/python/tvm/runtime/disco/_ffi_api.py b/python/tvm/runtime/disco/_ffi_api.py
index 340be86708db..79e1a52ad44e 100644
--- a/python/tvm/runtime/disco/_ffi_api.py
+++ b/python/tvm/runtime/disco/_ffi_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs from C++"""
-from ..._ffi import _init_api
+from ...ffi import _init_api
 
 _init_api("runtime.disco", __name__)
diff --git a/python/tvm/runtime/disco/process_pool.py b/python/tvm/runtime/disco/process_pool.py
index 95969e038e0f..8f05f28e9158 100644
--- a/python/tvm/runtime/disco/process_pool.py
+++ b/python/tvm/runtime/disco/process_pool.py
@@ -20,7 +20,7 @@
 import subprocess
 import sys
 
-from tvm._ffi import register_func
+from tvm.ffi import register_func
 from tvm.runtime import ShapeTuple
 
 
diff --git a/python/tvm/runtime/disco/session.py b/python/tvm/runtime/disco/session.py
index 719e399b597f..c551eac428b7 100644
--- a/python/tvm/runtime/disco/session.py
+++ b/python/tvm/runtime/disco/session.py
@@ -25,7 +25,7 @@
 
 import numpy as np
 
-from ..._ffi import get_global_func, register_func, register_object
+from ...ffi import get_global_func, register_func, register_object
 from ..device import Device
 from ..container import ShapeTuple
 from ..ndarray import NDArray
diff --git a/python/tvm/runtime/module.py b/python/tvm/runtime/module.py
index bb1fbb5fe3c2..3dd4de5da0a8 100644
--- a/python/tvm/runtime/module.py
+++ b/python/tvm/runtime/module.py
@@ -24,8 +24,8 @@
 import numpy as np
 
 import tvm.ffi
-from tvm._ffi.base import _RUNTIME_ONLY
-from tvm._ffi.libinfo import find_include_path
+from tvm.base import _RUNTIME_ONLY
+from tvm.libinfo import find_include_path
 
 from . import _ffi_api
 
diff --git a/python/tvm/runtime/ndarray.py b/python/tvm/runtime/ndarray.py
index 7c298050929f..9a026707cb48 100644
--- a/python/tvm/runtime/ndarray.py
+++ b/python/tvm/runtime/ndarray.py
@@ -70,7 +70,7 @@ def from_dlpack(ext_tensor):
     )
 
 
-@tvm._ffi.register_object("object.NDArray")
+@tvm.ffi.register_object("object.NDArray")
 class NDArray(tvm.ffi.core.NDArray):
     """Lightweight NDArray class of TVM runtime.
 
diff --git a/python/tvm/runtime/object_path.py b/python/tvm/runtime/object_path.py
index ff223b75998c..45e4925a3e28 100644
--- a/python/tvm/runtime/object_path.py
+++ b/python/tvm/runtime/object_path.py
@@ -22,7 +22,7 @@
 
 from typing import Optional
 
-import tvm._ffi
+import tvm.ffi
 from tvm.runtime import Object
 from . import _ffi_node_api
 
@@ -40,7 +40,7 @@
 )
 
 
-@tvm._ffi.register_object("ObjectPath")
+@tvm.ffi.register_object("ObjectPath")
 class ObjectPath(Object):
     """
     Path to an object from some root object.
@@ -94,42 +94,42 @@ def missing_map_entry(self) -> "ObjectPath":
     __hash__ = Object.__hash__
 
 
-@tvm._ffi.register_object("RootPath")
+@tvm.ffi.register_object("RootPath")
 class RootPath(ObjectPath):
     pass
 
 
-@tvm._ffi.register_object("AttributeAccessPath")
+@tvm.ffi.register_object("AttributeAccessPath")
 class AttributeAccessPath(ObjectPath):
     pass
 
 
-@tvm._ffi.register_object("UnknownAttributeAccessPath")
+@tvm.ffi.register_object("UnknownAttributeAccessPath")
 class UnknownAttributeAccessPath(ObjectPath):
     pass
 
 
-@tvm._ffi.register_object("ArrayIndexPath")
+@tvm.ffi.register_object("ArrayIndexPath")
 class ArrayIndexPath(ObjectPath):
     pass
 
 
-@tvm._ffi.register_object("MissingArrayElementPath")
+@tvm.ffi.register_object("MissingArrayElementPath")
 class MissingArrayElementPath(ObjectPath):
     pass
 
 
-@tvm._ffi.register_object("MapValuePath")
+@tvm.ffi.register_object("MapValuePath")
 class MapValuePath(ObjectPath):
     pass
 
 
-@tvm._ffi.register_object("MissingMapEntryPath")
+@tvm.ffi.register_object("MissingMapEntryPath")
 class MissingMapEntryPath(ObjectPath):
     pass
 
 
-@tvm._ffi.register_object("ObjectPathPair")
+@tvm.ffi.register_object("ObjectPathPair")
 class ObjectPathPair(Object):
     """
     Pair of ObjectPaths, one for each object being tested for structural equality.
diff --git a/python/tvm/runtime/profiling/__init__.py b/python/tvm/runtime/profiling/__init__.py
index 903cd443876f..45189a008495 100644
--- a/python/tvm/runtime/profiling/__init__.py
+++ b/python/tvm/runtime/profiling/__init__.py
@@ -17,7 +17,7 @@
 """Registration of profiling objects in python."""
 
 from typing import Dict, Sequence, Optional
-from ... import _ffi
+from ... import ffi as _ffi
 from . import _ffi_api
 from .. import Object, Device
 
diff --git a/python/tvm/runtime/profiling/_ffi_api.py b/python/tvm/runtime/profiling/_ffi_api.py
index d26b847a699f..85e5d4ca020c 100644
--- a/python/tvm/runtime/profiling/_ffi_api.py
+++ b/python/tvm/runtime/profiling/_ffi_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI for profiling"""
-from ... import _ffi
+from ...ffi import _init_api
 
-_ffi._init_api("runtime.profiling", __name__)
+_init_api("runtime.profiling", __name__)
diff --git a/python/tvm/runtime/relax_vm.py b/python/tvm/runtime/relax_vm.py
index 411835cd216f..d69c3308fad4 100644
--- a/python/tvm/runtime/relax_vm.py
+++ b/python/tvm/runtime/relax_vm.py
@@ -18,12 +18,12 @@
 """The Relax virtual machine."""
 from enum import IntEnum
 from typing import Any, Callable, Dict, List, Optional, Tuple, Union
+from numbers import Number, Integral
 
 import numpy as np  # type: ignore
 
 import tvm
-from tvm._ffi import base as _base
-from tvm._ffi import register_func
+from tvm.ffi import register_func
 from tvm.runtime import Device, Object, PackedFunc
 from tvm.runtime.profiling import Report
 
@@ -198,7 +198,7 @@ def _convert(self, arg: Any, cargs: List) -> None:
         def _gettype(arg):
             if isinstance(arg, np.float16):
                 return "float16"
-            elif isinstance(arg, (_base.integer_types, bool)):
+            elif isinstance(arg, (Integral, bool)):
                 return "int32"
             else:
                 return "float32"
@@ -215,7 +215,7 @@ def _gettype(arg):
             for field in arg:
                 self._convert(field, field_args)
             cargs.append(tuple(field_args))
-        elif isinstance(arg, (_base.numeric_types, bool)):
+        elif isinstance(arg, (Number, bool)):
             dtype = _gettype(arg)
             value = tvm.nd.array(np.array(arg, dtype=dtype), device=tvm.cpu(0))
             cargs.append(value)
diff --git a/python/tvm/runtime/script_printer.py b/python/tvm/runtime/script_printer.py
index ad3f612c4e29..ade34e1e9b85 100644
--- a/python/tvm/runtime/script_printer.py
+++ b/python/tvm/runtime/script_printer.py
@@ -18,7 +18,7 @@
 import os
 from typing import Dict, List, Optional, Sequence
 
-from tvm._ffi import get_global_func, register_object
+from tvm.ffi import get_global_func, register_object
 from tvm.runtime import Object
 
 from . import _ffi_node_api
diff --git a/python/tvm/runtime/support.py b/python/tvm/runtime/support.py
index 3716460a2709..149a66ef7b55 100644
--- a/python/tvm/runtime/support.py
+++ b/python/tvm/runtime/support.py
@@ -19,10 +19,10 @@
 
 import re
 
-import tvm._ffi
+import tvm.ffi
 
 
-@tvm._ffi.register_func("tvm.runtime.regex_match")
+@tvm.ffi.register_func("tvm.runtime.regex_match")
 def _regex_match(regex_pattern: str, match_against: str) -> bool:
     """Check if a pattern matches a regular expression
 
diff --git a/python/tvm/script/_ffi_api.py b/python/tvm/script/_ffi_api.py
index ebc638f3fd35..8ae8f7b7f9a5 100644
--- a/python/tvm/script/_ffi_api.py
+++ b/python/tvm/script/_ffi_api.py
@@ -14,7 +14,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.script"""
-import tvm._ffi
+import tvm.ffi
 
 
-tvm._ffi._init_api("script", __name__)
+tvm.ffi._init_api("script", __name__)
diff --git a/python/tvm/script/ir_builder/_ffi_api.py b/python/tvm/script/ir_builder/_ffi_api.py
index 68811c9e018c..8ee223051986 100644
--- a/python/tvm/script/ir_builder/_ffi_api.py
+++ b/python/tvm/script/ir_builder/_ffi_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.script.ir_builder"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("script.ir_builder", __name__)  # pylint: disable=protected-access
+tvm.ffi._init_api("script.ir_builder", __name__)  # pylint: disable=protected-access
diff --git a/python/tvm/script/ir_builder/base.py b/python/tvm/script/ir_builder/base.py
index 1d5d050444f7..95b5c5002558 100644
--- a/python/tvm/script/ir_builder/base.py
+++ b/python/tvm/script/ir_builder/base.py
@@ -17,7 +17,7 @@
 """A generic IRBuilder across the TVM stack"""
 from typing import Any, Callable, List
 
-from tvm._ffi import register_object as _register_object
+from tvm.ffi import register_object as _register_object
 from tvm.runtime import Object as _Object
 
 from . import _ffi_api
diff --git a/python/tvm/script/ir_builder/ir/_ffi_api.py b/python/tvm/script/ir_builder/ir/_ffi_api.py
index 874cc278af83..5b9d801a6ed3 100644
--- a/python/tvm/script/ir_builder/ir/_ffi_api.py
+++ b/python/tvm/script/ir_builder/ir/_ffi_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("script.ir_builder.ir", __name__)  # pylint: disable=protected-access
+tvm.ffi._init_api("script.ir_builder.ir", __name__)  # pylint: disable=protected-access
diff --git a/python/tvm/script/ir_builder/ir/frame.py b/python/tvm/script/ir_builder/ir/frame.py
index e16d86dc227e..d2737fde59a6 100644
--- a/python/tvm/script/ir_builder/ir/frame.py
+++ b/python/tvm/script/ir_builder/ir/frame.py
@@ -16,7 +16,7 @@
 # under the License.
 """Package tvm.script.ir_builder.ir.frame"""
 
-from tvm._ffi import register_object as _register_object
+from tvm.ffi import register_object as _register_object
 
 from ..base import IRBuilderFrame
 
diff --git a/python/tvm/script/ir_builder/relax/_ffi_api.py b/python/tvm/script/ir_builder/relax/_ffi_api.py
index 6e2098cf88af..1c767bacc4c5 100644
--- a/python/tvm/script/ir_builder/relax/_ffi_api.py
+++ b/python/tvm/script/ir_builder/relax/_ffi_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.script.ir_builder.relax"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("script.ir_builder.relax", __name__)  # pylint: disable=protected-access
+tvm.ffi._init_api("script.ir_builder.relax", __name__)  # pylint: disable=protected-access
diff --git a/python/tvm/script/ir_builder/relax/distributed/_ffi_api.py b/python/tvm/script/ir_builder/relax/distributed/_ffi_api.py
index d7121744890a..4d2ba60c2002 100644
--- a/python/tvm/script/ir_builder/relax/distributed/_ffi_api.py
+++ b/python/tvm/script/ir_builder/relax/distributed/_ffi_api.py
@@ -15,8 +15,8 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.script.ir_builder.relax.distributed"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api(
+tvm.ffi._init_api(
     "script.ir_builder.relax.distributed", __name__
 )  # pylint: disable=protected-access
diff --git a/python/tvm/script/ir_builder/relax/distributed/ir.py b/python/tvm/script/ir_builder/relax/distributed/ir.py
index aa0e2a26d34c..159ad5aea169 100644
--- a/python/tvm/script/ir_builder/relax/distributed/ir.py
+++ b/python/tvm/script/ir_builder/relax/distributed/ir.py
@@ -18,6 +18,7 @@
 
 """IRBuilder for distributed Relax dialect"""
 from typing import Union, List, Tuple, Optional
+from numbers import Number
 
 import numpy as _np  # type: ignore
 import tvm
@@ -27,7 +28,7 @@
 from tvm.relax.expr import Tuple as RxTuple
 from tvm.relax.distributed import DTensorStructInfo
 from tvm.relax.utils import args_converter
-from tvm._ffi import base as _base
+from tvm import base as _base
 from tvm.runtime import ndarray as _nd
 from tvm.relax.op.distributed import (
     redistribute as _redistribute,
@@ -114,7 +115,7 @@ def const(
     if not isinstance(struct_info, DTensorStructInfo):
         raise TypeError("struct_info needs to be an instance of DTensorStructInfo. ")
     dtype = str(struct_info.tensor_sinfo.dtype)
-    if isinstance(value, (_base.numeric_types, (bool, list))):
+    if isinstance(value, (Number, (bool, list))):
         value = _np.array(value, dtype=dtype)
 
     if isinstance(value, (_np.ndarray, _np.generic)):
diff --git a/python/tvm/script/ir_builder/relax/frame.py b/python/tvm/script/ir_builder/relax/frame.py
index 97e181fbe4be..181f62ec4f39 100644
--- a/python/tvm/script/ir_builder/relax/frame.py
+++ b/python/tvm/script/ir_builder/relax/frame.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """IR Builder Frame for Relax dialect"""
-from tvm._ffi import register_object as _register_object
+from tvm.ffi import register_object as _register_object
 
 from ..base import IRBuilderFrame
 
diff --git a/python/tvm/script/ir_builder/tir/_ffi_api.py b/python/tvm/script/ir_builder/tir/_ffi_api.py
index 876f5f3a35a0..69797f986afd 100644
--- a/python/tvm/script/ir_builder/tir/_ffi_api.py
+++ b/python/tvm/script/ir_builder/tir/_ffi_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("script.ir_builder.tir", __name__)  # pylint: disable=protected-access
+tvm.ffi._init_api("script.ir_builder.tir", __name__)  # pylint: disable=protected-access
diff --git a/python/tvm/script/ir_builder/tir/frame.py b/python/tvm/script/ir_builder/tir/frame.py
index b2229d503bfb..e3ce2e6e2eb1 100644
--- a/python/tvm/script/ir_builder/tir/frame.py
+++ b/python/tvm/script/ir_builder/tir/frame.py
@@ -17,7 +17,7 @@
 """IRBuilder for TIR"""
 from typing import List, Union
 
-from tvm._ffi import register_object as _register_object
+from tvm.ffi import register_object as _register_object
 from tvm.tir import Buffer, Var
 
 from ..base import IRBuilderFrame
diff --git a/python/tvm/script/parser/core/parser.py b/python/tvm/script/parser/core/parser.py
index f40b9a7cf6d3..78da15ca1f27 100644
--- a/python/tvm/script/parser/core/parser.py
+++ b/python/tvm/script/parser/core/parser.py
@@ -24,7 +24,7 @@
 
 import numpy as np
 
-from tvm._ffi.base import TVMError
+from tvm.base import TVMError
 from tvm.error import DiagnosticError
 from tvm.ir import GlobalVar
 
diff --git a/python/tvm/script/printer/_ffi_api.py b/python/tvm/script/printer/_ffi_api.py
index 944ecad01e08..9cbf6cfdca22 100644
--- a/python/tvm/script/printer/_ffi_api.py
+++ b/python/tvm/script/printer/_ffi_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.script.printer"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("script.printer", __name__)  # pylint: disable=protected-access
+tvm.ffi._init_api("script.printer", __name__)  # pylint: disable=protected-access
diff --git a/python/tvm/script/printer/doc.py b/python/tvm/script/printer/doc.py
index 9a6e7f1b8c8f..02a67e916bc0 100644
--- a/python/tvm/script/printer/doc.py
+++ b/python/tvm/script/printer/doc.py
@@ -19,7 +19,7 @@
 from enum import IntEnum, unique
 from typing import Dict, List, Optional, Sequence, Tuple, Union
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.runtime import Object, ObjectPath
 from tvm.tir import FloatImm, IntImm
 
diff --git a/python/tvm/support.py b/python/tvm/support.py
index a50a5e7b5732..7e0ad5875f83 100644
--- a/python/tvm/support.py
+++ b/python/tvm/support.py
@@ -22,11 +22,11 @@
 import sys
 
 import tvm
-import tvm._ffi
+import tvm.ffi
 from .runtime.module import Module
 from . import get_global_func
 
-tvm._ffi._init_api("support", __name__)
+tvm.ffi._init_api("support", __name__)
 
 
 def libinfo():
diff --git a/python/tvm/target/_ffi_api.py b/python/tvm/target/_ffi_api.py
index 3f3c4f2b8e46..489b59b4c6ae 100644
--- a/python/tvm/target/_ffi_api.py
+++ b/python/tvm/target/_ffi_api.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.target"""
-import tvm._ffi
+import tvm.ffi
 
 
-tvm._ffi._init_api("target", __name__)
+tvm.ffi._init_api("target", __name__)
diff --git a/python/tvm/target/datatype.py b/python/tvm/target/datatype.py
index aaf30afaf535..5003dbdde253 100644
--- a/python/tvm/target/datatype.py
+++ b/python/tvm/target/datatype.py
@@ -26,7 +26,7 @@
     BinaryOpExpr as _BinaryOpExpr,
 )
 from tvm.tir.op import call_pure_extern
-from tvm._ffi import register_func as _register_func
+from tvm.ffi import register_func as _register_func
 from tvm.tir import call_intrin
 
 
@@ -214,7 +214,7 @@ class name (e.g. Add, LE, Cast, Call).
         )
     else:
         lower_func_name = "tvm.datatype.lower." + target + "." + op_name + "." + src_type_name
-    tvm._ffi.register_func(lower_func_name, lower_func)
+    tvm.ffi.register_func(lower_func_name, lower_func)
 
 
 def register_min_func(func, type_name):
diff --git a/python/tvm/target/detect_target.py b/python/tvm/target/detect_target.py
index 57e39dfdb0ef..ec1875eb90a1 100644
--- a/python/tvm/target/detect_target.py
+++ b/python/tvm/target/detect_target.py
@@ -17,7 +17,7 @@
 """Detect target."""
 from typing import Union
 
-from .._ffi import get_global_func
+from ..ffi import get_global_func
 from ..runtime import Device
 from ..runtime.ndarray import device
 from . import Target
diff --git a/python/tvm/target/target.py b/python/tvm/target/target.py
index ba7b93333f1f..eb8bf1f9b807 100644
--- a/python/tvm/target/target.py
+++ b/python/tvm/target/target.py
@@ -20,8 +20,8 @@
 import warnings
 from typing import Union
 
-import tvm._ffi
-from tvm._ffi import register_func as _register_func
+import tvm.ffi
+from tvm.ffi import register_func as _register_func
 from tvm.runtime import Device
 from tvm.runtime import Object, convert
 from tvm.runtime.container import String
@@ -30,7 +30,7 @@
 from . import _ffi_api
 
 
-@tvm._ffi.register_object
+@tvm.ffi.register_object
 class TargetKind(Object):
     """Kind of a compilation target"""
 
@@ -53,7 +53,7 @@ def __getattr__(self, name: str):
         return _ffi_api.TargetGetFeature(self.target, name)
 
 
-@tvm._ffi.register_object
+@tvm.ffi.register_object
 class Target(Object):
     """Target device information, use through TVM API.
 
diff --git a/python/tvm/target/virtual_device.py b/python/tvm/target/virtual_device.py
index b2bc0bbcf5ee..3d923a4623d2 100644
--- a/python/tvm/target/virtual_device.py
+++ b/python/tvm/target/virtual_device.py
@@ -22,7 +22,7 @@
 from . import _ffi_api
 
 
-@tvm._ffi.register_object
+@tvm.ffi.register_object
 class VirtualDevice(Object):
     """A compile time representation for where data is to be stored at runtime,
     and how to compile code to compute it."""
diff --git a/python/tvm/target/x86.py b/python/tvm/target/x86.py
index c040eface808..177021f1433f 100644
--- a/python/tvm/target/x86.py
+++ b/python/tvm/target/x86.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """Common x86 related utilities"""
-from .._ffi import register_func
+from ..ffi import register_func
 from .codegen import target_has_features
 
 
diff --git a/python/tvm/te/_ffi_api.py b/python/tvm/te/_ffi_api.py
index ac814c844724..98e466e9e88c 100644
--- a/python/tvm/te/_ffi_api.py
+++ b/python/tvm/te/_ffi_api.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.te"""
-import tvm._ffi
+import tvm.ffi
 
 
-tvm._ffi._init_api("te", __name__)
+tvm.ffi._init_api("te", __name__)
diff --git a/python/tvm/te/operation.py b/python/tvm/te/operation.py
index 8d72fc794011..4a5d2425e669 100644
--- a/python/tvm/te/operation.py
+++ b/python/tvm/te/operation.py
@@ -14,18 +14,17 @@
 # KIND, either express or implied.  See the License for the
 # specific language governing permissions and limitations
 # under the License.
-""" Operation class for computation declaration."""
+"""Operation class for computation declaration."""
 import inspect
 
 # pylint: disable=invalid-name
 from numbers import Integral as _Integral
 from typing import List, Optional, Union
 
-import tvm._ffi
+import tvm.ffi
 import tvm.arith._ffi_api
 import tvm.tir
 import tvm.tir._ffi_api
-from tvm._ffi.base import string_types
 from tvm.ir import Array
 from tvm.runtime import convert
 
@@ -516,7 +515,7 @@ def thread_axis(dom=None, tag="", name="", span=None):
     axis : IterVar
         The thread itervar.
     """
-    if isinstance(dom, string_types):
+    if isinstance(dom, str):
         tag, dom = dom, None
     if not tag:
         raise ValueError("tag must be given as Positional or keyword argument")
diff --git a/python/tvm/te/tensor.py b/python/tvm/te/tensor.py
index dde92870518d..aad18a8b016c 100644
--- a/python/tvm/te/tensor.py
+++ b/python/tvm/te/tensor.py
@@ -16,7 +16,7 @@
 # under the License.
 """Tensor class for computation declaration."""
 # pylint: disable=invalid-name
-import tvm._ffi
+import tvm.ffi
 
 from tvm.runtime import Object, ObjectGeneric
 from tvm.tir import expr as _expr, DataProducer
@@ -48,7 +48,7 @@ def dtype(self):
         return self.tensor.dtype
 
 
-@tvm._ffi.register_object("te.Tensor")
+@tvm.ffi.register_object("te.Tensor")
 class Tensor(DataProducer, _expr.ExprOp):
     """Tensor object, to construct, see function.Tensor"""
 
@@ -141,12 +141,12 @@ def input_tensors(self):
         return _ffi_api.OpInputTensors(self)
 
 
-@tvm._ffi.register_object
+@tvm.ffi.register_object
 class PlaceholderOp(Operation):
     """Placeholder operation."""
 
 
-@tvm._ffi.register_object
+@tvm.ffi.register_object
 class BaseComputeOp(Operation):
     """Compute operation."""
 
@@ -161,12 +161,12 @@ def reduce_axis(self):
         return self.__getattr__("reduce_axis")
 
 
-@tvm._ffi.register_object
+@tvm.ffi.register_object
 class ComputeOp(BaseComputeOp):
     """Scalar operation."""
 
 
-@tvm._ffi.register_object
+@tvm.ffi.register_object
 class ScanOp(Operation):
     """Scan operation."""
 
@@ -176,6 +176,6 @@ def scan_axis(self):
         return self.__getattr__("scan_axis")
 
 
-@tvm._ffi.register_object
+@tvm.ffi.register_object
 class ExternOp(Operation):
     """External operation."""
diff --git a/python/tvm/testing/_ffi_api.py b/python/tvm/testing/_ffi_api.py
index 56a77223b767..e3c30d1299a1 100644
--- a/python/tvm/testing/_ffi_api.py
+++ b/python/tvm/testing/_ffi_api.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.testing"""
-import tvm._ffi
+import tvm.ffi
 
 
-tvm._ffi._init_api("testing", __name__)
+tvm.ffi._init_api("testing", __name__)
diff --git a/python/tvm/testing/popen_pool.py b/python/tvm/testing/popen_pool.py
index 42a34ccc61da..0fc3ce219030 100644
--- a/python/tvm/testing/popen_pool.py
+++ b/python/tvm/testing/popen_pool.py
@@ -36,19 +36,19 @@ def after_initializer():
     return TEST_GLOBAL_STATE_1, TEST_GLOBAL_STATE_2, TEST_GLOBAL_STATE_3
 
 
-@tvm._ffi.register_func("testing.identity_py")
+@tvm.ffi.register_func("testing.identity_py")
 def identity_py(arg):
     return arg
 
 
 def register_ffi():
-    @tvm._ffi.register_func("testing.nested_identity_py")
+    @tvm.ffi.register_func("testing.nested_identity_py")
     def _identity_py(arg):  # pylint: disable=unused-variable
         return arg
 
 
 def call_py_ffi(arg):
-    _identity_py = tvm._ffi.get_global_func("testing.nested_identity_py")
+    _identity_py = tvm.ffi.get_global_func("testing.nested_identity_py")
     return _identity_py(arg)
 
 
diff --git a/python/tvm/testing/utils.py b/python/tvm/testing/utils.py
index ec3fc28b1d8a..6b047de4460a 100644
--- a/python/tvm/testing/utils.py
+++ b/python/tvm/testing/utils.py
@@ -87,7 +87,7 @@ def test_something():
 import tvm.arith
 import tvm.tir
 import tvm.te
-import tvm._ffi
+import tvm.ffi
 
 from tvm.target import codegen
 from tvm.contrib import nvcc, cudnn, rocm
diff --git a/python/tvm/tir/_ffi_api.py b/python/tvm/tir/_ffi_api.py
index 1b60b8c81c6d..8c438557c8c1 100644
--- a/python/tvm/tir/_ffi_api.py
+++ b/python/tvm/tir/_ffi_api.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.tir"""
-import tvm._ffi
+import tvm.ffi
 
 
-tvm._ffi._init_api("tir", __name__)
+tvm.ffi._init_api("tir", __name__)
diff --git a/python/tvm/tir/analysis/_ffi_api.py b/python/tvm/tir/analysis/_ffi_api.py
index 6c1687e8a520..40a7b4caf340 100644
--- a/python/tvm/tir/analysis/_ffi_api.py
+++ b/python/tvm/tir/analysis/_ffi_api.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.tir.analysis"""
-import tvm._ffi
+import tvm.ffi
 
 
-tvm._ffi._init_api("tir.analysis", __name__)
+tvm.ffi._init_api("tir.analysis", __name__)
diff --git a/python/tvm/tir/block_dependence_info.py b/python/tvm/tir/block_dependence_info.py
index 5f1664628890..67a644967e4b 100644
--- a/python/tvm/tir/block_dependence_info.py
+++ b/python/tvm/tir/block_dependence_info.py
@@ -18,7 +18,7 @@
 to store the block level dependences"""
 
 from typing import Union, Optional
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.ir.module import IRModule
 from tvm.runtime import Object
 from tvm.tir import Block, PrimFunc
diff --git a/python/tvm/tir/block_scope.py b/python/tvm/tir/block_scope.py
index 30e047b4f78a..b24cca0707a0 100644
--- a/python/tvm/tir/block_scope.py
+++ b/python/tvm/tir/block_scope.py
@@ -18,7 +18,7 @@
 from enum import IntEnum
 from typing import List, Optional, Union
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.runtime import Object
 from tvm.tir import Block, For
 
diff --git a/python/tvm/tir/buffer.py b/python/tvm/tir/buffer.py
index 72c2a40fedd2..1f40520e55be 100644
--- a/python/tvm/tir/buffer.py
+++ b/python/tvm/tir/buffer.py
@@ -17,15 +17,14 @@
 """Abstraction for array data structures."""
 from numbers import Integral
 
-import tvm._ffi
-from tvm._ffi.base import string_types
+import tvm.ffi
 from tvm.ir import PointerType, PrimExpr, PrimType, Range
 from tvm.runtime import Object, Scriptable, convert
 
 from . import _ffi_api
 
 
-@tvm._ffi.register_object("tir.Buffer")
+@tvm.ffi.register_object("tir.Buffer")
 class Buffer(Object, Scriptable):
     """Symbolic data buffer in TVM.
 
@@ -85,7 +84,7 @@ def access_ptr(self, access_mask, ptr_type="handle", content_lanes=1, offset=0,
           # Get access ptr for read with extent
           buffer.access_ptr("r", extent = 100)
         """
-        if isinstance(access_mask, string_types):
+        if isinstance(access_mask, str):
             mask = 0
             for value in access_mask:
                 if value == "r":
@@ -350,6 +349,6 @@ def decl_buffer(
     )
 
 
-@tvm._ffi.register_object("tir.DataProducer")
+@tvm.ffi.register_object("tir.DataProducer")
 class DataProducer(Object):
     pass
diff --git a/python/tvm/tir/data_layout.py b/python/tvm/tir/data_layout.py
index 71cc404ee23b..39874640ff40 100644
--- a/python/tvm/tir/data_layout.py
+++ b/python/tvm/tir/data_layout.py
@@ -17,13 +17,13 @@
 """Data layout."""
 from typing import Union
 
-import tvm._ffi
+import tvm.ffi
 
 from tvm.runtime import Object
 from . import _ffi_api
 
 
-@tvm._ffi.register_object("tir.Layout")
+@tvm.ffi.register_object("tir.Layout")
 class Layout(Object):
     """Layout is composed of upper cases, lower cases and numbers,
     where upper case indicates a primal axis and
@@ -81,7 +81,7 @@ def factor_of(self, axis):
         return _ffi_api.LayoutFactorOf(self, axis)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.BijectiveLayout")
+@tvm.ffi.register_object("tir.BijectiveLayout")
 class BijectiveLayout(Object):
     """Bijective mapping for two layouts (src-layout and dst-layout).
     It provides shape and index conversion between each other.
diff --git a/python/tvm/tir/expr.py b/python/tvm/tir/expr.py
index b293343cae74..e57c01f23afc 100644
--- a/python/tvm/tir/expr.py
+++ b/python/tvm/tir/expr.py
@@ -29,7 +29,7 @@
 """
 from typing import List, Optional, Union
 
-import tvm._ffi
+import tvm.ffi
 import tvm.ir._ffi_api
 from tvm import ir
 from tvm.ir import Op, PrimExpr
@@ -349,7 +349,7 @@ class LogicalExpr(PrimExprWithOp):
     pass
 
 
-@tvm._ffi.register_object("tir.Var")
+@tvm.ffi.register_object("tir.Var")
 class Var(PrimExprWithOp):
     """Symbolic variable.
 
@@ -372,7 +372,7 @@ def __init__(self, name: str, dtype: Union[str, ir.Type], span: Optional[Span] =
         self.__init_handle_by_constructor__(_ffi_api.Var, name, dtype, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.SizeVar")
+@tvm.ffi.register_object("tir.SizeVar")
 class SizeVar(Var):
     """Symbolic variable to represent a tensor index size
        which is greater or equal to zero.
@@ -394,7 +394,7 @@ def __init__(self, name: str, dtype: Union[str, ir.Type], span: Optional[Span] =
         self.__init_handle_by_constructor__(_ffi_api.SizeVar, name, dtype, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.IterVar")
+@tvm.ffi.register_object("tir.IterVar")
 class IterVar(ExprOp, Object, Scriptable):
     """Represent iteration variable.
 
@@ -467,7 +467,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("tir.CommReducer")
+@tvm.ffi.register_object("tir.CommReducer")
 class CommReducer(Object, Scriptable):
     """Commutative reduce operator
 
@@ -507,7 +507,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("tir.Reduce")
+@tvm.ffi.register_object("tir.Reduce")
 class Reduce(PrimExprWithOp):
     """Reduce node.
 
@@ -558,7 +558,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object
+@tvm.ffi.register_object
 class FloatImm(ConstExpr):
     """Float constant.
 
@@ -585,7 +585,7 @@ def __float__(self) -> float:
         return self.value
 
 
-@tvm._ffi.register_object
+@tvm.ffi.register_object
 class IntImm(ConstExpr):
     """Int constant.
 
@@ -627,7 +627,7 @@ def __bool__(self) -> bool:
         return self.__nonzero__()
 
 
-@tvm._ffi.register_object("tir.StringImm")  # type: ignore
+@tvm.ffi.register_object("tir.StringImm")  # type: ignore
 class StringImm(ConstExpr):
     """String constant.
 
@@ -659,7 +659,7 @@ def __hash__(self) -> int:
         return PrimExpr.__hash__(self)
 
 
-@tvm._ffi.register_object("tir.Cast")
+@tvm.ffi.register_object("tir.Cast")
 class Cast(PrimExprWithOp):
     """Cast expression.
 
@@ -681,7 +681,7 @@ def __init__(self, dtype, value, span: Optional[Span] = None) -> None:
         self.__init_handle_by_constructor__(_ffi_api.Cast, dtype, value, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.Add")
+@tvm.ffi.register_object("tir.Add")
 class Add(BinaryOpExpr):
     """Add node.
 
@@ -701,7 +701,7 @@ def __init__(self, a: PrimExpr, b: PrimExpr, span: Optional[Span] = None) -> Non
         self.__init_handle_by_constructor__(_ffi_api.Add, a, b, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.Sub")
+@tvm.ffi.register_object("tir.Sub")
 class Sub(BinaryOpExpr):
     """Sub node.
 
@@ -721,7 +721,7 @@ def __init__(self, a: PrimExpr, b: PrimExpr, span: Optional[Span] = None) -> Non
         self.__init_handle_by_constructor__(_ffi_api.Sub, a, b, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.Mul")
+@tvm.ffi.register_object("tir.Mul")
 class Mul(BinaryOpExpr):
     """Mul node.
 
@@ -741,7 +741,7 @@ def __init__(self, a: PrimExpr, b: PrimExpr, span: Optional[Span] = None) -> Non
         self.__init_handle_by_constructor__(_ffi_api.Mul, a, b, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.Div")
+@tvm.ffi.register_object("tir.Div")
 class Div(BinaryOpExpr):
     """Div node.
 
@@ -761,7 +761,7 @@ def __init__(self, a: PrimExpr, b: PrimExpr, span: Optional[Span] = None) -> Non
         self.__init_handle_by_constructor__(_ffi_api.Div, a, b, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.Mod")
+@tvm.ffi.register_object("tir.Mod")
 class Mod(BinaryOpExpr):
     """Mod node.
 
@@ -781,7 +781,7 @@ def __init__(self, a: PrimExpr, b: PrimExpr, span: Optional[Span] = None) -> Non
         self.__init_handle_by_constructor__(_ffi_api.Mod, a, b, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.FloorDiv")
+@tvm.ffi.register_object("tir.FloorDiv")
 class FloorDiv(BinaryOpExpr):
     """FloorDiv node.
 
@@ -801,7 +801,7 @@ def __init__(self, a: PrimExpr, b: PrimExpr, span: Optional[Span] = None) -> Non
         self.__init_handle_by_constructor__(_ffi_api.FloorDiv, a, b, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.FloorMod")
+@tvm.ffi.register_object("tir.FloorMod")
 class FloorMod(BinaryOpExpr):
     """FloorMod node.
 
@@ -821,7 +821,7 @@ def __init__(self, a: PrimExpr, b: PrimExpr, span: Optional[Span] = None) -> Non
         self.__init_handle_by_constructor__(_ffi_api.FloorMod, a, b, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.Min")
+@tvm.ffi.register_object("tir.Min")
 class Min(BinaryOpExpr):
     """Min node.
 
@@ -841,7 +841,7 @@ def __init__(self, a: PrimExpr, b: PrimExpr, span: Optional[Span] = None) -> Non
         self.__init_handle_by_constructor__(_ffi_api.Min, a, b, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.Max")
+@tvm.ffi.register_object("tir.Max")
 class Max(BinaryOpExpr):
     """Max node.
 
@@ -861,7 +861,7 @@ def __init__(self, a: PrimExpr, b: PrimExpr, span: Optional[Span] = None) -> Non
         self.__init_handle_by_constructor__(_ffi_api.Max, a, b, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.EQ")
+@tvm.ffi.register_object("tir.EQ")
 class EQ(CmpExpr):
     """EQ node.
 
@@ -881,7 +881,7 @@ def __init__(self, a: PrimExpr, b: PrimExpr, span: Optional[Span] = None) -> Non
         self.__init_handle_by_constructor__(_ffi_api.EQ, a, b, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.NE")
+@tvm.ffi.register_object("tir.NE")
 class NE(CmpExpr):
     """NE node.
 
@@ -901,7 +901,7 @@ def __init__(self, a: PrimExpr, b: PrimExpr, span: Optional[Span] = None) -> Non
         self.__init_handle_by_constructor__(_ffi_api.NE, a, b, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.LT")
+@tvm.ffi.register_object("tir.LT")
 class LT(CmpExpr):
     """LT node.
 
@@ -921,7 +921,7 @@ def __init__(self, a: PrimExpr, b: PrimExpr, span: Optional[Span] = None) -> Non
         self.__init_handle_by_constructor__(_ffi_api.LT, a, b, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.LE")
+@tvm.ffi.register_object("tir.LE")
 class LE(CmpExpr):
     """LE node.
 
@@ -941,7 +941,7 @@ def __init__(self, a: PrimExpr, b: PrimExpr, span: Optional[Span] = None) -> Non
         self.__init_handle_by_constructor__(_ffi_api.LE, a, b, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.GT")
+@tvm.ffi.register_object("tir.GT")
 class GT(CmpExpr):
     """GT node.
 
@@ -961,7 +961,7 @@ def __init__(self, a: PrimExpr, b: PrimExpr, span: Optional[Span] = None) -> Non
         self.__init_handle_by_constructor__(_ffi_api.GT, a, b, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.GE")
+@tvm.ffi.register_object("tir.GE")
 class GE(CmpExpr):
     """GE node.
 
@@ -981,7 +981,7 @@ def __init__(self, a: PrimExpr, b: PrimExpr, span: Optional[Span] = None) -> Non
         self.__init_handle_by_constructor__(_ffi_api.GE, a, b, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.And")
+@tvm.ffi.register_object("tir.And")
 class And(LogicalExpr):
     """And node.
 
@@ -1001,7 +1001,7 @@ def __init__(self, a: PrimExpr, b: PrimExpr, span: Optional[Span] = None) -> Non
         self.__init_handle_by_constructor__(_ffi_api.And, a, b, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.Or")
+@tvm.ffi.register_object("tir.Or")
 class Or(LogicalExpr):
     """Or node.
 
@@ -1024,7 +1024,7 @@ def __init__(self, a: PrimExpr, b: PrimExpr, span: Optional[Span] = None) -> Non
         self.__init_handle_by_constructor__(_ffi_api.Or, a, b, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.Not")
+@tvm.ffi.register_object("tir.Not")
 class Not(LogicalExpr):
     """Not node.
 
@@ -1043,7 +1043,7 @@ def __init__(self, a: PrimExpr, span: Optional[Span] = None) -> None:
         self.__init_handle_by_constructor__(_ffi_api.Not, a, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.Select")
+@tvm.ffi.register_object("tir.Select")
 class Select(PrimExprWithOp):
     """Select node.
 
@@ -1087,7 +1087,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("tir.BufferLoad")
+@tvm.ffi.register_object("tir.BufferLoad")
 class BufferLoad(PrimExprWithOp):
     """Buffer load node.
 
@@ -1122,7 +1122,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("tir.ProducerLoad")
+@tvm.ffi.register_object("tir.ProducerLoad")
 class ProducerLoad(PrimExprWithOp):
     """Producer load node.
 
@@ -1149,7 +1149,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("tir.Ramp")
+@tvm.ffi.register_object("tir.Ramp")
 class Ramp(PrimExprWithOp):
     """Ramp node.
 
@@ -1180,7 +1180,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("tir.Broadcast")
+@tvm.ffi.register_object("tir.Broadcast")
 class Broadcast(PrimExprWithOp):
     """Broadcast node.
 
@@ -1203,7 +1203,7 @@ def __init__(self, value: PrimExpr, lanes: PrimExpr, span: Optional[Span] = None
         self.__init_handle_by_constructor__(_ffi_api.Broadcast, value, lanes, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.Shuffle")
+@tvm.ffi.register_object("tir.Shuffle")
 class Shuffle(PrimExprWithOp):
     """Shuffle node.
 
@@ -1241,7 +1241,7 @@ class CallEffectKind:
     Opaque = UpdateState
 
 
-@tvm._ffi.register_object("tir.Call")
+@tvm.ffi.register_object("tir.Call")
 class Call(PrimExprWithOp):
     """Call node.
 
@@ -1281,7 +1281,7 @@ def __init__(
         self.__init_handle_by_constructor__(_ffi_api.Call, dtype, op, args, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.Let")
+@tvm.ffi.register_object("tir.Let")
 class Let(PrimExprWithOp):
     """Let node.
 
diff --git a/python/tvm/tir/function.py b/python/tvm/tir/function.py
index eb3c50b409c8..55bae37809f0 100644
--- a/python/tvm/tir/function.py
+++ b/python/tvm/tir/function.py
@@ -22,7 +22,7 @@
 from typing import Callable, List, Mapping, Optional, Tuple, Union
 
 import tvm
-import tvm._ffi
+import tvm.ffi
 import tvm.runtime
 from tvm.ir import BaseFunc, Range
 from tvm.runtime import Object, Scriptable
@@ -33,7 +33,7 @@
 from .expr import PrimExpr, Var
 
 
-@tvm._ffi.register_object("tir.PrimFunc")
+@tvm.ffi.register_object("tir.PrimFunc")
 class PrimFunc(BaseFunc, Scriptable):
     """A function declaration expression.
 
@@ -174,7 +174,7 @@ def mem_copy_16_16(a: T.handle, b: T.handle) -> None:
         return _ffi_api.Specialize(self, param_map)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.TensorIntrin")
+@tvm.ffi.register_object("tir.TensorIntrin")
 class TensorIntrin(Object):
     """A tensor intrinsic.
 
@@ -230,7 +230,7 @@ def get(name: str, allow_missing: bool = False) -> Optional["TensorIntrin"]:
         return _ffi_api.TensorIntrinGet(name, allow_missing)  # pylint: type: ignore
 
 
-@tvm._ffi.register_object("tir.IndexMap")
+@tvm.ffi.register_object("tir.IndexMap")
 class IndexMap(Object):
     """A mapping from multi-dimensional indices to another set of multi-dimensional indices
 
diff --git a/python/tvm/tir/ir_builder.py b/python/tvm/tir/ir_builder.py
index 777d46ec7b0d..7a9708848ab4 100644
--- a/python/tvm/tir/ir_builder.py
+++ b/python/tvm/tir/ir_builder.py
@@ -16,7 +16,6 @@
 # under the License.
 """Developer API of IR node builder make function."""
 import tvm
-from tvm._ffi.base import string_types
 from tvm.runtime import ObjectGeneric, const
 from tvm.ir import container as _container
 
@@ -194,9 +193,9 @@ def scope_attr(self, node, attr_key, value):
             ib.scope_attr(x, "storage_scope", "global")
             x[i] = x[i - 1] + 1
         """
-        if isinstance(node, string_types):
+        if isinstance(node, str):
             node = _expr.StringImm(node)
-        if isinstance(value, string_types):
+        if isinstance(value, str):
             value = _expr.StringImm(value)
         # thread_extent could be zero for dynamic workloads
         if attr_key == "thread_extent":
diff --git a/python/tvm/tir/op.py b/python/tvm/tir/op.py
index 4c0b51280e3e..57aa060cd0c5 100644
--- a/python/tvm/tir/op.py
+++ b/python/tvm/tir/op.py
@@ -18,7 +18,7 @@
 """Operators used in TIR expression."""
 from typing import Any, Optional, Union
 
-import tvm._ffi
+import tvm.ffi
 from tvm import tir
 from tvm.ir import Array, Op, PrimExpr
 from tvm.ir.base import Span
@@ -1927,7 +1927,7 @@ def all(*args, span=None):
     return val
 
 
-@tvm._ffi.register_func("tvm.default_trace_action")
+@tvm.ffi.register_func("tvm.default_trace_action")
 def _tvm_default_trace_action(*args):
     print(list(args))
 
diff --git a/python/tvm/tir/schedule/_ffi_api.py b/python/tvm/tir/schedule/_ffi_api.py
index ae8bdfde54bf..b854145beb6a 100644
--- a/python/tvm/tir/schedule/_ffi_api.py
+++ b/python/tvm/tir/schedule/_ffi_api.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.tir.schedule"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("tir.schedule", __name__)  # pylint: disable=protected-access
+tvm.ffi._init_api("tir.schedule", __name__)  # pylint: disable=protected-access
diff --git a/python/tvm/tir/schedule/analysis.py b/python/tvm/tir/schedule/analysis.py
index 15748a99c81a..491a689c9309 100644
--- a/python/tvm/tir/schedule/analysis.py
+++ b/python/tvm/tir/schedule/analysis.py
@@ -17,7 +17,7 @@
 """Analysis used in TensorIR scheduling"""
 from typing import List, Optional
 
-import tvm._ffi
+import tvm.ffi
 from tvm.runtime import Object
 
 from ..buffer import Buffer
@@ -62,7 +62,7 @@ def suggest_index_map(
     )
 
 
-@tvm._ffi.register_object("tir.schedule.TensorizeInfo")
+@tvm.ffi.register_object("tir.schedule.TensorizeInfo")
 class TensorizeInfo(Object):
     """Necessary information used for tensorization."""
 
@@ -90,7 +90,7 @@ def get_tensorize_loop_mapping(
     return _ffi_api.GetTensorizeLoopMapping(sch, block, desc_func, allow_padding)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.schedule.AutoTensorizeMappingInfo")
+@tvm.ffi.register_object("tir.schedule.AutoTensorizeMappingInfo")
 class AutoTensorizeMappingInfo(Object):
     """Necessary information used to perform transformations for tensorization."""
 
diff --git a/python/tvm/tir/schedule/instruction.py b/python/tvm/tir/schedule/instruction.py
index 09b2d70dc321..5a8563e652b6 100644
--- a/python/tvm/tir/schedule/instruction.py
+++ b/python/tvm/tir/schedule/instruction.py
@@ -17,7 +17,7 @@
 """Schedule instructions each corresponds to a schedule primitive"""
 from typing import TYPE_CHECKING, Any, List, Union
 
-from tvm._ffi import register_object as _register_object
+from tvm.ffi import register_object as _register_object
 from tvm.runtime import Object
 
 from . import _ffi_api
diff --git a/python/tvm/tir/schedule/schedule.py b/python/tvm/tir/schedule/schedule.py
index d0dd427e6d5e..5325ecdc16c4 100644
--- a/python/tvm/tir/schedule/schedule.py
+++ b/python/tvm/tir/schedule/schedule.py
@@ -18,7 +18,7 @@
 import inspect
 from typing import Callable, Dict, List, Literal, Optional, Tuple, Union
 
-from tvm._ffi import register_object as _register_object
+from tvm.ffi import register_object as _register_object
 from tvm.error import TVMError, register_error
 from tvm.ir import GlobalVar, IRModule, PrimExpr
 from tvm.runtime import Object
diff --git a/python/tvm/tir/schedule/state.py b/python/tvm/tir/schedule/state.py
index df2eb534e633..f082a9e92ea7 100644
--- a/python/tvm/tir/schedule/state.py
+++ b/python/tvm/tir/schedule/state.py
@@ -20,7 +20,7 @@
 from enum import IntEnum
 from typing import Dict, Optional, Union
 
-from tvm._ffi import register_object
+from tvm.ffi import register_object
 from tvm.ir import IRModule
 from tvm.runtime import Object
 from tvm.tir import Block, BlockRealize, For, PrimFunc
diff --git a/python/tvm/tir/schedule/trace.py b/python/tvm/tir/schedule/trace.py
index 25cf5244addf..15bb201ae641 100644
--- a/python/tvm/tir/schedule/trace.py
+++ b/python/tvm/tir/schedule/trace.py
@@ -18,7 +18,7 @@
 import os
 from typing import TYPE_CHECKING, Any, Callable, Dict, List, Optional
 
-from tvm._ffi import register_object as _register_object
+from tvm.ffi import register_object as _register_object
 from tvm.runtime import Object
 
 from ...ir import Array, Map, save_json
diff --git a/python/tvm/tir/stmt.py b/python/tvm/tir/stmt.py
index aa3b17a7a12f..a04f80b55e7a 100644
--- a/python/tvm/tir/stmt.py
+++ b/python/tvm/tir/stmt.py
@@ -29,7 +29,7 @@
 from enum import IntEnum
 from typing import List, Mapping, Optional, Union
 
-import tvm._ffi
+import tvm.ffi
 from tvm.ir import PrimExpr, Range, Span
 from tvm.runtime import Object, Scriptable, const, NDArray
 
@@ -42,7 +42,7 @@ class Stmt(Object, Scriptable):
     """Base class of all the statements."""
 
 
-@tvm._ffi.register_object("tir.LetStmt")
+@tvm.ffi.register_object("tir.LetStmt")
 class LetStmt(Stmt):
     """LetStmt node.
 
@@ -72,7 +72,7 @@ def __init__(self, var: Var, value: PrimExpr, body: Stmt, span: Optional[Span] =
         )
 
 
-@tvm._ffi.register_object("tir.AssertStmt")
+@tvm.ffi.register_object("tir.AssertStmt")
 class AssertStmt(Stmt):
     """AssertStmt node.
 
@@ -120,7 +120,7 @@ class ForKind(IntEnum):
     THREAD_BINDING = 4  # pylint: disable=invalid-name
 
 
-@tvm._ffi.register_object("tir.For")
+@tvm.ffi.register_object("tir.For")
 class For(Stmt):
     """For node.
 
@@ -185,7 +185,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("tir.While")
+@tvm.ffi.register_object("tir.While")
 class While(Stmt):
     """While node.
 
@@ -209,7 +209,7 @@ def __init__(self, condition: PrimExpr, body: Stmt, span: Optional[Span] = None)
         self.__init_handle_by_constructor__(_ffi_api.While, condition, body, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.BufferStore")
+@tvm.ffi.register_object("tir.BufferStore")
 class BufferStore(Stmt):
     """Buffer store node.
 
@@ -252,7 +252,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("tir.BufferRealize")
+@tvm.ffi.register_object("tir.BufferRealize")
 class BufferRealize(Stmt):
     """Buffer realize node.
 
@@ -293,7 +293,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("tir.ProducerStore")
+@tvm.ffi.register_object("tir.ProducerStore")
 class ProducerStore(Stmt):
     """ProducerStore node.
 
@@ -329,7 +329,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("tir.Allocate")
+@tvm.ffi.register_object("tir.Allocate")
 class Allocate(Stmt):
     """Allocate node.
 
@@ -389,7 +389,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("tir.AllocateConst")
+@tvm.ffi.register_object("tir.AllocateConst")
 class AllocateConst(Stmt):
     """Allocate constant node.
 
@@ -451,7 +451,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("tir.DeclBuffer")
+@tvm.ffi.register_object("tir.DeclBuffer")
 class DeclBuffer(Stmt):
     """DeclBuffer node.
 
@@ -475,7 +475,7 @@ def __init__(self, buffer: Buffer, body: Stmt, span: Optional[Span] = None) -> N
         self.__init_handle_by_constructor__(_ffi_api.DeclBuffer, buffer, body, span)
 
 
-@tvm._ffi.register_object("tir.AttrStmt")
+@tvm.ffi.register_object("tir.AttrStmt")
 class AttrStmt(Stmt):
     """AttrStmt node.
 
@@ -511,7 +511,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("tir.ProducerRealize")
+@tvm.ffi.register_object("tir.ProducerRealize")
 class ProducerRealize(Stmt):
     """ProducerRealize node.
 
@@ -563,7 +563,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("tir.SeqStmt")
+@tvm.ffi.register_object("tir.SeqStmt")
 class SeqStmt(Stmt):
     """Sequence of statements.
 
@@ -589,7 +589,7 @@ def __len__(self):
         return len(self.seq)
 
 
-@tvm._ffi.register_object("tir.IfThenElse")
+@tvm.ffi.register_object("tir.IfThenElse")
 class IfThenElse(Stmt):
     """IfThenElse node.
 
@@ -624,7 +624,7 @@ def __init__(
         )
 
 
-@tvm._ffi.register_object("tir.Evaluate")
+@tvm.ffi.register_object("tir.Evaluate")
 class Evaluate(Stmt):
     """Evaluate node.
 
@@ -644,7 +644,7 @@ def __init__(self, value: PrimExpr, span: Optional[Span] = None) -> None:
         self.__init_handle_by_constructor__(_ffi_api.Evaluate, value, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.Prefetch")
+@tvm.ffi.register_object("tir.Prefetch")
 class Prefetch(Stmt):
     """Prefetch node.
 
@@ -668,7 +668,7 @@ def __init__(self, buffer: Buffer, bounds: List[Range], span: Optional[Span] = N
         self.__init_handle_by_constructor__(_ffi_api.Prefetch, buffer, bounds, span)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.BufferRegion")
+@tvm.ffi.register_object("tir.BufferRegion")
 class BufferRegion(Object, Scriptable):
     """BufferRegion node.
 
@@ -688,7 +688,7 @@ def __init__(self, buffer: Buffer, region: List[Range]) -> None:
         self.__init_handle_by_constructor__(_ffi_api.BufferRegion, buffer, region)  # type: ignore
 
 
-@tvm._ffi.register_object("tir.MatchBufferRegion")
+@tvm.ffi.register_object("tir.MatchBufferRegion")
 class MatchBufferRegion(Object, Scriptable):
     """MatchBufferRegion node.
 
@@ -710,7 +710,7 @@ def __init__(self, buffer: Buffer, source: BufferRegion) -> None:
         )
 
 
-@tvm._ffi.register_object("tir.Block")
+@tvm.ffi.register_object("tir.Block")
 class Block(Stmt):
     """Block node.
 
@@ -792,7 +792,7 @@ def __init__(
         )  # type: ignore
 
 
-@tvm._ffi.register_object("tir.BlockRealize")
+@tvm.ffi.register_object("tir.BlockRealize")
 class BlockRealize(Stmt):
     """BlockRealize node.
 
diff --git a/python/tvm/tir/tensor_intrin/cuda.py b/python/tvm/tir/tensor_intrin/cuda.py
index 57b1c3b873d7..6f964c94370d 100644
--- a/python/tvm/tir/tensor_intrin/cuda.py
+++ b/python/tvm/tir/tensor_intrin/cuda.py
@@ -18,7 +18,7 @@
 """Intrinsics for tensorization on NVIDIA GPU."""
 from typing import Dict, Literal, Optional, Tuple
 
-from tvm._ffi import register_func
+from tvm.ffi import register_func
 from tvm.runtime import convert
 from tvm.script import tir as T
 from tvm.tir import Cast, IntImm, TensorIntrin
diff --git a/python/tvm/tir/transform/_ffi_api.py b/python/tvm/tir/transform/_ffi_api.py
index 86f7bdf5dac3..8a6607c11af0 100644
--- a/python/tvm/tir/transform/_ffi_api.py
+++ b/python/tvm/tir/transform/_ffi_api.py
@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI APIs for tvm.tir.transform"""
-import tvm._ffi
+import tvm.ffi
 
 
-tvm._ffi._init_api("tir.transform", __name__)
+tvm.ffi._init_api("tir.transform", __name__)
diff --git a/python/tvm/tir/transform/function_pass.py b/python/tvm/tir/transform/function_pass.py
index c6c825699e9a..b679d4ab16ce 100644
--- a/python/tvm/tir/transform/function_pass.py
+++ b/python/tvm/tir/transform/function_pass.py
@@ -19,13 +19,13 @@
 import functools
 from typing import Callable, List, Optional, Union
 
-import tvm._ffi
+import tvm.ffi
 from tvm.ir.transform import Pass, PassInfo
 
 from . import _ffi_api
 
 
-@tvm._ffi.register_object("tir.PrimFuncPass")
+@tvm.ffi.register_object("tir.PrimFuncPass")
 class PrimFuncPass(Pass):
     """A pass that works on each :py:func:`tvm.tir.PrimFunc` in a module. A function
     pass class should be created through py:func:`tvm.tir.transform.function_pass`.
diff --git a/python/tvm/topi/__init__.py b/python/tvm/topi/__init__.py
index fa4e98a89a42..51417874210e 100644
--- a/python/tvm/topi/__init__.py
+++ b/python/tvm/topi/__init__.py
@@ -24,7 +24,7 @@
 Some of the schedule function may have been specially optimized for a
 specific workload.
 """
-from tvm._ffi.libinfo import __version__
+from tvm.libinfo import __version__
 
 # Ensure C++ schedules get registered first, so python schedules can
 # override them.
diff --git a/python/tvm/topi/cpp/cuda.py b/python/tvm/topi/cpp/cuda.py
index ce2efa929824..22f97293d38d 100644
--- a/python/tvm/topi/cpp/cuda.py
+++ b/python/tvm/topi/cpp/cuda.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI for CUDA TOPI ops and schedules"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("topi.cuda", "tvm.topi.cpp.cuda")
+tvm.ffi._init_api("topi.cuda", "tvm.topi.cpp.cuda")
diff --git a/python/tvm/topi/cpp/generic.py b/python/tvm/topi/cpp/generic.py
index d314eca8b22d..3230d5428bb2 100644
--- a/python/tvm/topi/cpp/generic.py
+++ b/python/tvm/topi/cpp/generic.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI for generic TOPI ops and schedules"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("topi.generic", "tvm.topi.cpp.generic")
+tvm.ffi._init_api("topi.generic", "tvm.topi.cpp.generic")
diff --git a/python/tvm/topi/cpp/impl.py b/python/tvm/topi/cpp/impl.py
index 2c877c300dc9..e5473a7e6602 100644
--- a/python/tvm/topi/cpp/impl.py
+++ b/python/tvm/topi/cpp/impl.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """Load Lib for C++ TOPI ops and schedules"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("topi", "tvm.topi.cpp")
+tvm.ffi._init_api("topi", "tvm.topi.cpp")
diff --git a/python/tvm/topi/cpp/nn.py b/python/tvm/topi/cpp/nn.py
index 0e3cee703de0..2ea1fc371404 100644
--- a/python/tvm/topi/cpp/nn.py
+++ b/python/tvm/topi/cpp/nn.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI for NN TOPI ops and schedules"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("topi.nn", "tvm.topi.cpp.nn")
+tvm.ffi._init_api("topi.nn", "tvm.topi.cpp.nn")
diff --git a/python/tvm/topi/cpp/rocm.py b/python/tvm/topi/cpp/rocm.py
index eab51107beb7..771fc3c3f0f3 100644
--- a/python/tvm/topi/cpp/rocm.py
+++ b/python/tvm/topi/cpp/rocm.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI for Rocm TOPI ops and schedules"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("topi.rocm", "tvm.topi.cpp.rocm")
+tvm.ffi._init_api("topi.rocm", "tvm.topi.cpp.rocm")
diff --git a/python/tvm/topi/cpp/utils.py b/python/tvm/topi/cpp/utils.py
index 60a2747f9abb..b78a6baa0f01 100644
--- a/python/tvm/topi/cpp/utils.py
+++ b/python/tvm/topi/cpp/utils.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI for TOPI utility functions"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("topi.utils", "tvm.topi.cpp.utils")
+tvm.ffi._init_api("topi.utils", "tvm.topi.cpp.utils")
diff --git a/python/tvm/topi/cpp/vision/__init__.py b/python/tvm/topi/cpp/vision/__init__.py
index 000602fb399d..5fdf1ac4e3a8 100644
--- a/python/tvm/topi/cpp/vision/__init__.py
+++ b/python/tvm/topi/cpp/vision/__init__.py
@@ -16,8 +16,8 @@
 # under the License.
 
 """FFI for vision TOPI ops and schedules"""
-import tvm._ffi
+import tvm.ffi
 
 from . import yolo
 
-tvm._ffi._init_api("topi.vision", "tvm.topi.cpp.vision")
+tvm.ffi._init_api("topi.vision", "tvm.topi.cpp.vision")
diff --git a/python/tvm/topi/cpp/vision/yolo.py b/python/tvm/topi/cpp/vision/yolo.py
index 17e2327295d2..5d8bdd99d24c 100644
--- a/python/tvm/topi/cpp/vision/yolo.py
+++ b/python/tvm/topi/cpp/vision/yolo.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI for Yolo TOPI ops and schedules"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("topi.vision.yolo", "tvm.topi.cpp.vision.yolo")
+tvm.ffi._init_api("topi.vision.yolo", "tvm.topi.cpp.vision.yolo")
diff --git a/python/tvm/topi/cpp/x86.py b/python/tvm/topi/cpp/x86.py
index 0034af02c572..18de30c668a3 100644
--- a/python/tvm/topi/cpp/x86.py
+++ b/python/tvm/topi/cpp/x86.py
@@ -15,6 +15,6 @@
 # specific language governing permissions and limitations
 # under the License.
 """FFI for x86 TOPI ops and schedules"""
-import tvm._ffi
+import tvm.ffi
 
-tvm._ffi._init_api("topi.x86", "tvm.topi.cpp.x86")
+tvm.ffi._init_api("topi.x86", "tvm.topi.cpp.x86")
diff --git a/python/tvm/topi/generic_op_impl.py b/python/tvm/topi/generic_op_impl.py
index 661e24d2a45c..2e0beadbc177 100644
--- a/python/tvm/topi/generic_op_impl.py
+++ b/python/tvm/topi/generic_op_impl.py
@@ -64,7 +64,7 @@ def _tensor_bop_impl(lhs, rhs):
         it performs tensor-scalar {op} operation on an element-wise basis.
 
         Otherwise, it performs default generic.{op} operation, as defined
-        in tvm.generic module.
+        in tvm.tir.generic module.
 
         Parameters
         ----------
diff --git a/tests/python/codegen/test_gpu_codegen_allreduce.py b/tests/python/codegen/test_gpu_codegen_allreduce.py
index 5a24fcb197a5..aa56411cc9e0 100644
--- a/tests/python/codegen/test_gpu_codegen_allreduce.py
+++ b/tests/python/codegen/test_gpu_codegen_allreduce.py
@@ -102,7 +102,7 @@ def compile_metal(src, target):
 
     if define_metal_compile_callback:
         if cached is None:
-            tvm._ffi.registry.remove_global_func(name)
+            tvm.ffi.registry.remove_global_func(name)
         else:
             tvm.register_func(name, cached, override=True)
 
diff --git a/tests/python/contrib/test_hexagon/test_vtcm.py b/tests/python/contrib/test_hexagon/test_vtcm.py
index ea9feb740319..2795f5630163 100644
--- a/tests/python/contrib/test_hexagon/test_vtcm.py
+++ b/tests/python/contrib/test_hexagon/test_vtcm.py
@@ -62,7 +62,7 @@ def test_vtcm_limit(vtcm_capacity, limited):
     def _raises_exception(f):
         try:
             f()
-        except tvm._ffi.base.TVMError:
+        except tvm.base.TVMError:
             return True
         return False
 
diff --git a/tests/python/disco/test_loader.py b/tests/python/disco/test_loader.py
index ba0287afc61a..5089336f09d3 100644
--- a/tests/python/disco/test_loader.py
+++ b/tests/python/disco/test_loader.py
@@ -25,7 +25,7 @@
 import tvm.testing
 from tvm import dlight as dl
 from tvm import relax as rx
-from tvm._ffi import register_func
+from tvm.ffi import register_func
 from tvm.contrib import tvmjs
 from tvm.runtime import ShapeTuple
 from tvm.runtime import disco as di
diff --git a/tests/python/meta_schedule/test_meta_schedule_builder.py b/tests/python/meta_schedule/test_meta_schedule_builder.py
index a74ac893262f..090a393fbeeb 100644
--- a/tests/python/meta_schedule/test_meta_schedule_builder.py
+++ b/tests/python/meta_schedule/test_meta_schedule_builder.py
@@ -14,7 +14,7 @@
 # KIND, either express or implied.  See the License for the
 # specific language governing permissions and limitations
 # under the License.
-""" Test Meta Schedule Builder """
+"""Test Meta Schedule Builder"""
 
 import os
 import sys
@@ -25,7 +25,7 @@
 import tvm.testing
 
 from tvm import script
-from tvm._ffi import register_func
+from tvm.ffi import register_func
 from tvm.meta_schedule.builder import (
     BuilderInput,
     BuilderResult,
diff --git a/tests/python/meta_schedule/test_meta_schedule_post_order_apply.py b/tests/python/meta_schedule/test_meta_schedule_post_order_apply.py
index 8d374c01ae19..57d9d0961088 100644
--- a/tests/python/meta_schedule/test_meta_schedule_post_order_apply.py
+++ b/tests/python/meta_schedule/test_meta_schedule_post_order_apply.py
@@ -25,7 +25,7 @@
 import tvm.testing
 from tvm import te
 from tvm.ir.module import IRModule
-from tvm._ffi import register_func
+from tvm.ffi import register_func
 from tvm.error import TVMError
 from tvm.meta_schedule import TuneContext
 from tvm.meta_schedule.schedule_rule import PyScheduleRule
diff --git a/tests/python/meta_schedule/test_meta_schedule_runner.py b/tests/python/meta_schedule/test_meta_schedule_runner.py
index 03ab8c58b48d..e5deefe7507c 100644
--- a/tests/python/meta_schedule/test_meta_schedule_runner.py
+++ b/tests/python/meta_schedule/test_meta_schedule_runner.py
@@ -25,7 +25,7 @@
 import pytest
 import tvm
 import tvm.testing
-from tvm._ffi import register_func
+from tvm.ffi import register_func
 from tvm.meta_schedule.arg_info import TensorInfo
 from tvm.meta_schedule.builder import BuilderInput, LocalBuilder
 from tvm.meta_schedule.runner import (
diff --git a/tests/python/meta_schedule/test_meta_schedule_space_generator.py b/tests/python/meta_schedule/test_meta_schedule_space_generator.py
index ef2be381c694..9457a9a40f00 100644
--- a/tests/python/meta_schedule/test_meta_schedule_space_generator.py
+++ b/tests/python/meta_schedule/test_meta_schedule_space_generator.py
@@ -14,7 +14,7 @@
 # KIND, either express or implied.  See the License for the
 # specific language governing permissions and limitations
 # under the License.
-""" Test Meta Schedule SpaceGenerator """
+"""Test Meta Schedule SpaceGenerator"""
 # pylint: disable=missing-function-docstring
 
 import math
@@ -22,7 +22,7 @@
 import pytest
 import tvm
 import tvm.testing
-from tvm._ffi.base import TVMError
+from tvm.base import TVMError
 from tvm.meta_schedule.space_generator import (
     PySpaceGenerator,
     ScheduleFn,
diff --git a/tests/python/relax/test_binding_rewrite.py b/tests/python/relax/test_binding_rewrite.py
index d0d3344eb61e..d3c78b3657a8 100644
--- a/tests/python/relax/test_binding_rewrite.py
+++ b/tests/python/relax/test_binding_rewrite.py
@@ -18,7 +18,7 @@
 import pytest
 import tvm
 import tvm.testing
-from tvm._ffi.base import TVMError
+from tvm.base import TVMError
 from tvm.relax.analysis import name_to_binding
 from tvm.relax.binding_rewrite import DataflowBlockRewrite
 from tvm.relax.expr import DataflowVar, Var
diff --git a/tests/python/relax/test_frontend_nn_extern_module.py b/tests/python/relax/test_frontend_nn_extern_module.py
index ef97cfd9056c..5aa48dd18aba 100644
--- a/tests/python/relax/test_frontend_nn_extern_module.py
+++ b/tests/python/relax/test_frontend_nn_extern_module.py
@@ -119,8 +119,8 @@ def test_sym(
 
 def _compile_cc(src: Path, dst: Path):
     # pylint: disable=import-outside-toplevel
-    from tvm._ffi.base import py_str
-    from tvm._ffi.libinfo import find_include_path
+    from tvm.base import py_str
+    from tvm.libinfo import find_include_path
 
     # pylint: enable=import-outside-toplevel
 
diff --git a/tests/python/relax/test_op_distributed.py b/tests/python/relax/test_op_distributed.py
index dc5440c0c113..2290174e464f 100644
--- a/tests/python/relax/test_op_distributed.py
+++ b/tests/python/relax/test_op_distributed.py
@@ -16,7 +16,7 @@
 # under the License.
 import pytest
 import tvm
-from tvm._ffi.base import TVMError
+from tvm.base import TVMError
 import tvm.testing
 from tvm import relax
 from tvm.script.parser import relax as R
diff --git a/tests/python/relax/test_op_grad.py b/tests/python/relax/test_op_grad.py
index 5ae30adb70ef..a8a17c7a135a 100644
--- a/tests/python/relax/test_op_grad.py
+++ b/tests/python/relax/test_op_grad.py
@@ -16,7 +16,7 @@
 # under the License.
 import pytest
 import tvm
-from tvm._ffi.base import TVMError
+from tvm.base import TVMError
 import tvm.testing
 from tvm import relax
 from tvm.ir import Op
diff --git a/tests/python/relax/test_relax_operators.py b/tests/python/relax/test_relax_operators.py
index 1fdded92393c..c94dd9f5789d 100644
--- a/tests/python/relax/test_relax_operators.py
+++ b/tests/python/relax/test_relax_operators.py
@@ -24,7 +24,7 @@
 import tvm
 import tvm.testing
 from tvm import relax
-from tvm._ffi.base import TVMError
+from tvm.base import TVMError
 from tvm.script import ir as I, relax as R, tir as T
 
 exec_mode = tvm.testing.parameter("bytecode", "compiled")
diff --git a/tests/python/relax/test_transform_gradient.py b/tests/python/relax/test_transform_gradient.py
index 072edea5c400..47c41ca108f9 100644
--- a/tests/python/relax/test_transform_gradient.py
+++ b/tests/python/relax/test_transform_gradient.py
@@ -20,7 +20,7 @@
 import tvm
 import tvm.testing
 from tvm import relax
-from tvm._ffi.base import TVMError
+from tvm.base import TVMError
 from tvm.ir.base import assert_structural_equal
 from tvm.script.parser import relax as R, tir as T, ir as I
 
diff --git a/tests/python/runtime/test_runtime_error.py b/tests/python/runtime/test_runtime_error.py
index 6f950d4ca77c..6eb6cc9b641b 100644
--- a/tests/python/runtime/test_runtime_error.py
+++ b/tests/python/runtime/test_runtime_error.py
@@ -92,7 +92,7 @@ def flevel3():
 
 @functools.lru_cache()
 def _has_debug_symbols():
-    lib = tvm._ffi.base._LIB
+    lib = tvm.base._LIB
     headers = subprocess.check_output(["objdump", "--section-headers", lib._name], encoding="utf-8")
     return ".debug" in headers
 
diff --git a/tests/python/runtime/test_runtime_rpc.py b/tests/python/runtime/test_runtime_rpc.py
index d06846400dca..6711ccf92f3f 100644
--- a/tests/python/runtime/test_runtime_rpc.py
+++ b/tests/python/runtime/test_runtime_rpc.py
@@ -103,7 +103,7 @@ def check_remote():
         assert f1(10) == 11
         f3 = client.get_function("rpc.test.except")
 
-        with pytest.raises(tvm._ffi.base.TVMError):
+        with pytest.raises(tvm.base.TVMError):
             f3("abc")
 
         f2 = client.get_function("rpc.test.strcat")
diff --git a/tests/python/te/test_te_verify_compute.py b/tests/python/te/test_te_verify_compute.py
index 7ea9321e18c9..880e960da7fa 100644
--- a/tests/python/te/test_te_verify_compute.py
+++ b/tests/python/te/test_te_verify_compute.py
@@ -35,14 +35,14 @@ def test_verify_compute():
     # Valid compute
     try:
         B = te.compute((n,), f1, name="B")
-    except tvm._ffi.base.TVMError as ex:
+    except tvm.base.TVMError as ex:
         assert False
 
     #
     # Valid compute
     try:
         B = te.compute((n,), f2, name="B")
-    except tvm._ffi.base.TVMError as ex:
+    except tvm.base.TVMError as ex:
         assert False
 
     #
@@ -50,7 +50,7 @@ def test_verify_compute():
     try:
         B = te.compute((n,), f3, name="B")
         assert False
-    except tvm._ffi.base.TVMError as ex:
+    except tvm.base.TVMError as ex:
         pass
 
     #
@@ -58,7 +58,7 @@ def test_verify_compute():
     try:
         B = te.compute((n,), f4, name="B")
         assert False
-    except tvm._ffi.base.TVMError as ex:
+    except tvm.base.TVMError as ex:
         pass
 
     #
@@ -66,7 +66,7 @@ def test_verify_compute():
     try:
         B0, B1 = te.compute((n,), f5, name="B")
         assert False
-    except tvm._ffi.base.TVMError as ex:
+    except tvm.base.TVMError as ex:
         pass
 
     #
@@ -74,7 +74,7 @@ def test_verify_compute():
     try:
         B0, B1 = te.compute((n,), f6, name="B")
         assert False
-    except tvm._ffi.base.TVMError as ex:
+    except tvm.base.TVMError as ex:
         pass
 
 
diff --git a/tests/python/tir-base/test_tir_base.py b/tests/python/tir-base/test_tir_base.py
index b33e846741cd..d204ebfb6084 100644
--- a/tests/python/tir-base/test_tir_base.py
+++ b/tests/python/tir-base/test_tir_base.py
@@ -17,7 +17,7 @@
 import tvm
 import pytest
 from tvm import tir
-from tvm._ffi.base import TVMError
+from tvm.base import TVMError
 from tvm.ir.transform import PassContext
 import itertools
 import pytest
diff --git a/tests/python/tir-transform/test_tir_transform_inject_ptx_async_copy.py b/tests/python/tir-transform/test_tir_transform_inject_ptx_async_copy.py
index 46487f969b96..da079f46e38e 100644
--- a/tests/python/tir-transform/test_tir_transform_inject_ptx_async_copy.py
+++ b/tests/python/tir-transform/test_tir_transform_inject_ptx_async_copy.py
@@ -430,7 +430,7 @@ def tvm_callback_cuda_postproc(code, _):
 
     # Restore previous postproc func to avoid impacting other tests
     if prev_postproc is None:
-        tvm._ffi.registry.remove_global_func(func_name)
+        tvm.ffi.registry.remove_global_func(func_name)
     else:
         tvm.register_func(func_name, prev_postproc, override=True)
 
diff --git a/version.py b/version.py
index a8ae77a8d5f2..bf2f0343d6c7 100644
--- a/version.py
+++ b/version.py
@@ -20,7 +20,7 @@
 This script runs and update all the locations that related to versions
 
 List of affected files:
-- tvm-root/python/tvm/_ffi/libinfo.py
+- tvm-root/python/tvm/libinfo.py
 - tvm-root/include/tvm/runtime/base.h
 - tvm-root/conda/recipe/meta.yaml
 - tvm-root/web/package.json
@@ -170,7 +170,7 @@ def sync_version(pub_ver, local_ver, dry_run):
     """Synchronize version."""
     # python uses the PEP-440: local version
     update(
-        os.path.join(PROJ_ROOT, "python", "tvm", "_ffi", "libinfo.py"),
+        os.path.join(PROJ_ROOT, "python", "tvm", "libinfo.py"),
         r"(?<=__version__ = \")[.0-9a-z\+]+",
         local_ver,
         dry_run,