Apache TVM

https://tvm.apache.org/docs/
Admin
Apache TVM is a deep learning compiler that enables high-performance machine learning optimization...
Tokens:1,058,113
Snippets:9,253
Trust Score:10
Update:1 month ago
Show doc for...
Tokens:
Raw
### start()

Source: https://tvm.apache.org/docs/reference/api/javadoc/org/apache/tvm/rpc/Server.html

Starts the RPC server instance.

```APIDOC
## start()

### Description
Starts the RPC server process.
```

--------------------------------

### start(inst)

Source: https://tvm.apache.org/docs/reference/api/typedoc/classes/_internal_.Environment.html

Marks the start of a WebAssembly instance within the environment.

```APIDOC
## Method: start

### Description
Marks the start of the provided WebAssembly instance.

### Parameters
- **inst** (Instance) - Required - The WebAssembly instance to start.
```

--------------------------------

### Layout configuration for ldstmatrix

Source: https://tvm.apache.org/docs/_sources/tirx/tile_primitives/copy/ldstmatrix.rst

Example setup for loading 2 row-major matrix tiles using the m8n8 fragment layout.

```python
from tvm.tirx.layout import S, TileLayout, laneid

num = 2; M, N = 8, num * 8
r_layout = TileLayout(S[(8, 4, num, 2) : (4 @ laneid, 1 @ laneid, 2, 1)])
s_layout = TileLayout(S[(8, 4, num, 2) : (num * 8, 2, 8, 1)])     # row-major
full = (slice(0, 8), slice(0, 4), slice(0, num), slice(0, 2))
```

--------------------------------

### Install Kernel Library for Development

Source: https://tvm.apache.org/docs/_sources/tirx/install.rst

Clone the repository and install in editable mode for development.

```bash
git clone https://github.com/mlc-ai/tirx-kernels
cd tirx-kernels
pip install -e .
```

--------------------------------

### Full IRModule and Transformation Example

Source: https://tvm.apache.org/docs/deep_dive/tensor_ir/tutorials/tir_transformation.html

A complete example showing an IRModule definition followed by a transformation function.

```python
# from tvm.script import ir as I
# from tvm.script import tirx as T
# from tvm.tirx.layout import Axis

@I.ir_module
class Module:
    @T.prim_func(s_tir=True)
    def main(A: T.Buffer((128, 128), "float32"), B: T.Buffer((128, 128), "float32"), C: T.Buffer((128, 128), "float32")):
        T.func_attr({"tirx.noalias": True})
        # with T.sblock("root"):
        Y = T.sblock_alloc_buffer((128, 128))
        for i, j_0 in T.grid(128, 16):
            for j_1_init in range(8):
                with T.sblock("Y_init"):
                    vi = T.axis.spatial(128, i)
                    vj = T.axis.spatial(128, j_0 * 8 + j_1_init)
                    T.reads()
                    T.writes(Y[vi, vj])
                    Y[vi, vj] = T.float32(0.0)
            for k, j_1 in T.grid(128, 8):
                with T.sblock("Y_update"):
                    vi = T.axis.spatial(128, i)
                    vj = T.axis.spatial(128, j_0 * 8 + j_1)
                    vk = T.axis.reduce(128, k)
                    T.reads(Y[vi, vj], A[vi, vk], B[vk, vj])
                    T.writes(Y[vi, vj])
                    Y[vi, vj] = Y[vi, vj] + A[vi, vk] * B[vk, vj]
            for ax0 in range(8):
                with T.sblock("C"):
                    vi = T.axis.spatial(128, i)
                    vj = T.axis.spatial(128, j_0 * 8 + ax0)
                    T.reads(Y[vi, vj])
                    T.writes(C[vi, vj])
                    C[vi, vj] = T.max(Y[vi, vj], T.float32(0.0))

# from tvm import s_tir
def apply_trace(sch: s_tir.Schedule) -> None:
  b0 = sch.get_sblock(name="Y", func_name="main")
  l1, l2, l3 = sch.get_loops(block=b0)
  l4, l5 = sch.split(loop=l2, factors=[None, 8], preserve_unit_iters=True, disable_predication=False)
  sch.reorder(l4, l3, l5)
  b6 = sch.get_sblock(name="C", func_name="main")
  sch.reverse_compute_at(block=b6, loop=l4, preserve_unit_loops=False, index=-1)
  b7 = sch.decompose_reduction(block=b0, loop=l3)
```

--------------------------------

### Install Google Test from Source

Source: https://tvm.apache.org/docs/_sources/install/from_source.rst

Commands to clone, build, and install Google Test for running C++ tests.

```bash
git clone https://github.com/google/googletest
cd googletest
mkdir build
cd build
cmake -DBUILD_SHARED_LIBS=ON ..
make
sudo make install
```

--------------------------------

### Start()

Source: https://tvm.apache.org/docs/reference/api/doxygen/classtvm_1_1runtime_1_1TimerNode.html

Starts the timer. This function should only be called once per object.

```APIDOC
## void Start()

### Description
Starts the timer. This function should only be called once per object.
```

--------------------------------

### Install Kernel Library

Source: https://tvm.apache.org/docs/_sources/tirx/install.rst

Install the optional tirx-kernels library from PyPI.

```bash
pip install tirx-kernels
```

--------------------------------

### Model Output Example

Source: https://tvm.apache.org/docs/how_to/tutorials/customize_opt.html

Example output generated after running the optimized model on the GPU.

```text
[[25302.758 26301.07  25925.635 24935.379 25223.234 24493.31  24251.89
  24630.104 25003.176 23884.219]]
```

--------------------------------

### RFactor Schedule Example

Source: https://tvm.apache.org/docs/reference/api/python/s_tir/schedule.html

Example showing the TensorIR before rfactor and the subsequent scheduling commands to apply it.

```python
@T.prim_func(s_tir=True)
def before_rfactor(a: T.handle, b: T.handle) -> None:
    A = T.match_buffer(a, (128, 128, 128))
    B = T.match_buffer(b, (128,))
    for ii, i, j in T.grid(128, 128, 128):
    with T.sblock("B"):
        vii, vi, vj = T.axis.remap("SRR", [ii, i, j])
        with T.init():
            B[vii] = 0.0
        B[vii] = B[vii] + A[vii, vi, vj]
```

```python
sch = tvm.s_tir.Schedule(before_rfactor)
_, _, k = sch.get_loops(sch.get_sblock("B"))
sch.rfactor(k, 0)
print(sch.mod["main"].script())
```

--------------------------------

### Model Output Example

Source: https://tvm.apache.org/docs/get_started/tutorials/quick_start.html

Example output from the model execution.

```text
[[24552.506 26274.057 25085.32  25691.854 26284.488 24763.63  26721.383
  24417.2   24932.246 25592.887]]
```

--------------------------------

### Install Reference Dependencies

Source: https://tvm.apache.org/docs/_sources/tirx/install.rst

Install pinned, mutually compatible revisions for correctness tests and reference baselines.

```bash
python scripts/install_reference_dependencies.py
```

--------------------------------

### Calculate Start Index

Source: https://tvm.apache.org/docs/reference/api/doxygen/pooling_8h_source.html

Computes the starting index for a given output index and dimension ratio.

```cpp
inline PrimExpr start_index(const PrimVar& out_index, const PrimExpr& odim, const PrimExpr& idim) {
313 return indexdiv(out_index * idim, odim);
314}
```

--------------------------------

### Install TVM via Environment Variables

Source: https://tvm.apache.org/docs/_sources/install/from_source.rst

Configure the environment to use a custom TVM installation path.

```bash
export TVM_HOME=/path-to-tvm
pip install --target=$TVM_HOME/python $TVM_HOME/3rdparty/tvm-ffi
export PYTHONPATH=$TVM_HOME/python:$PYTHONPATH
```

--------------------------------

### tvm.script.ir_builder.ir.ir_module

Source: https://tvm.apache.org/docs/reference/api/python/script/ir_builder.html

Start an IRModule frame.

```APIDOC
## tvm.script.ir_builder.ir.ir_module() -> IRModuleFrame

### Description
Start an ir_module frame.

### Returns
- **frame** (IRModuleFrame) - The constructed frame.
```

--------------------------------

### Merge parallel matmuls

Source: https://tvm.apache.org/docs/_sources/deep_dive/relax/dpl.rst

Example setup for merging three parallel matmuls using rewrite_bindings.

```python
from tvm.script import relax as R
from tvm.relax.dpl import wildcard, is_op, rewrite_bindings, PatternContext
```

--------------------------------

### Initialize and Use Trainer

Source: https://tvm.apache.org/docs/reference/api/python/relax/training.html

Example usage of SetupTrainer to prepare a module and Trainer to execute training steps.

```python
setup_trainer = SetupTrainer(
    MSELoss(reduction="sum"),
    SGD(0.001),
    [pred_ty, target_ty],
)
train_mod = setup_trainer(Backbone)
ex = tvm.compile(train_mod, target)
vm = relax.VirtualMachine(ex, dev)

trainer = training.Trainer(train_mod, vm, dev, False)

trainer.xaiver_uniform_init_params()
trainer.predict(input_instances)
trainer.update([input_instances], [labels])
```

--------------------------------

### Mark device entry with T.device_entry

Source: https://tvm.apache.org/docs/tirx/api/script.html

Marks the start of a device-region within a PrimFunc, separating host-side setup from device-side execution.

```python
@T.prim_func
def kernel(...):
    A = T.match_buffer(...)
    T.device_entry()           # device region starts here
    bx = T.cta_id([SM_COUNT])  # standalone scope-id def
    ...
```

--------------------------------

### Import Dependencies and Setup Environment

Source: https://tvm.apache.org/docs/_downloads/ee6bb02238ab43f2b87505cc74484f6b/mix_python_and_tvm_with_pymodule.ipynb

Initialize necessary libraries and check for environment compatibility.

```python
import os

try:
    import torch
    import torch.nn.functional as F
except ImportError:
    torch = None

import tvm
from tvm import relax
from tvm.relax.base_py_module import BasePyModule
from tvm.script import ir as I
from tvm.script import relax as R
from tvm.script import tirx as T

IS_IN_CI = os.getenv("CI", "").lower() == "true"
HAS_TORCH = torch is not None
RUN_EXAMPLE = HAS_TORCH and not IS_IN_CI
```

--------------------------------

### Init

Source: https://tvm.apache.org/docs/reference/api/doxygen/classtvm_1_1runtime_1_1vm_1_1VirtualMachine.html

Initializes the virtual machine for a set of devices.

```APIDOC
## Init(const std::vector<Device>& devices, const std::vector<AllocatorType>& alloc_types)

### Description
Initialize the virtual machine for a set of devices.

### Parameters
- **devices** (const std::vector<Device>&) - The set of TVM devices.
- **alloc_types** (const std::vector<AllocatorType>&) - The allocator types for each device.
```

--------------------------------

### TCGen05 Asynchronous Copy Demonstration

Source: https://tvm.apache.org/docs/_sources/tirx/tile_primitives/copy_async/tcgen05_ldst.rst

Example showing the setup and usage of asynchronous copy operations between local registers and shared tensor memory.

```python
@Tx.prim_func
def copy_async_test(A_ptr: Tx.handle, B_ptr: Tx.handle):
    A = Tx.match_buffer(A_ptr, (128, WIDTH), "float16"); B = Tx.match_buffer(B_ptr, (128, WIDTH), "float16")
    Tx.device_entry()
    warp_id = Tx.warp_id([4]); wg_id = Tx.warpgroup_id([1]); tid = Tx.thread_id([128])
    tmem_addr = Tx.alloc_shared([1], "uint32")
    if wg_id == 0:
        if warp_id == 0:
            Tx.ptx["tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32"](
                Tx.address_of(tmem_addr), Tx.uint32(32))
        Tx.tvm_storage_sync("shared")
        tmem = Tx.decl_buffer((128, WIDTH), "float16", scope="tmem", allocated_addr=tmem_addr[0],
                             layout=TileLayout(S[(128, WIDTH) : (1 @ TLane, 1 @ TCol)]))
        A_reg = Tx.alloc_local((WIDTH,), "float16"); B_reg = Tx.alloc_local((WIDTH,), "float16")
        A_local = A_reg.view(128, WIDTH, layout=local_view)
        B_local = B_reg.view(128, WIDTH, layout=local_view)
        # ... load A into A_reg, zero B_reg, cta_sync ...
        Tx.tile.wg.copy_async(tmem[:, :], A_local[:, :]); Tx.ptx.tcgen05.wait__st.sync.aligned()   # store (local -> tmem)
        Tx.cuda.cta_sync()
        Tx.tile.wg.copy_async(B_local[:, :], tmem[:, :]); Tx.ptx.tcgen05.wait__ld.sync.aligned()   # load  (tmem -> local)
        # ... write B_reg out; tcgen05.dealloc ...
```

--------------------------------

### Init

Source: https://tvm.apache.org/docs/reference/api/doxygen/classtvm_1_1runtime_1_1vm_1_1VirtualMachine.html

Initializes the virtual machine with specific devices and allocator types.

```APIDOC
## virtual void Init(const std::vector< Device > &devices, const std::vector< AllocatorType > &alloc_types)

### Description
Initializes the virtual machine for a set of devices.

### Parameters
- **devices** (const std::vector< Device > &) - The list of devices to initialize.
- **alloc_types** (const std::vector< AllocatorType > &) - The list of allocator types to use.
```

--------------------------------

### Start RPC Server

Source: https://tvm.apache.org/docs/_downloads/148819f3421b8d89b1723c3e15e3f19f/cross_compilation_and_rpc.ipynb

Commands to initialize and verify the RPC server on the remote device.

```bash
python -m tvm.exec.rpc_server --host 0.0.0.0 --port=9090
```

```bash
INFO:root:RPCServer: bind to 0.0.0.0:9090
```

--------------------------------

### Clone and Prepare Website Repository

Source: https://tvm.apache.org/docs/contribute/release_process.html

Initialize the local environment for updating the TVM website documentation.

```bash
git clone https://github.com/apache/tvm-site.git
pushd tvm-site
git checkout asf-site
pushd docs
```

--------------------------------

### RPC Workflow Demonstration

Source: https://tvm.apache.org/docs/_sources/how_to/tutorials/cross_compilation_and_rpc.rst

A conceptual overview of the steps required to set up an RPC server and execute a model remotely.

```python
print("\nRPC workflow (works for any remote device):")
print("=" * 50)
print("1. Start RPC server on target device:")
print("   python -m tvm.exec.rpc_server --host 0.0.0.0 --port=9090")
print("\n2. Connect from local machine:")
print("   remote = rpc.connect('DEVICE_IP', 9090)")
print("\n3. Upload compiled library:")
print("   remote.upload('model_deployed.so')")
print("   remote.upload('model_params.npz')")
print("\n4. Load and run remotely:")
print("   lib = remote.load_module('model_deployed.so')")
print("   vm = relax.VirtualMachine(lib, remote.cpu())")
print("   result = vm['main'](input, *params)")
```

--------------------------------

### Package and Deploy RPC Runtime

Source: https://tvm.apache.org/docs/how_to/tutorials/cross_compilation_and_rpc.html

Steps to package the cross-compiled runtime and set up the environment on the target device.

```bash
rm -rf tvm_runtime_package
mkdir tvm_runtime_package
cp -a python tvm_runtime_package/
cp cross_build/lib/libtvm_ffi.so tvm_runtime_package/python/tvm/
cp cross_build/lib/libtvm_runtime*.so tvm_runtime_package/python/tvm/
tar -czf tvm_runtime.tar.gz -C tvm_runtime_package python
```

```bash
tar -xzf tvm_runtime.tar.gz
export PYTHONPATH=`pwd`/python:${PYTHONPATH}
```

--------------------------------

### start(event_type, leader=None)

Source: https://tvm.apache.org/docs/_sources/tirx/native_basics/cuda/profiling.rst

Marks the beginning of a profiled region.

```APIDOC
## start(event_type, leader=None)

### Description
Opens a region for profiling. Every thread executes this, but only the leader thread stores the record.

### Parameters
- **event_type** (int) - Required - The type of event to record.
- **leader** (bool, optional) - Optional - Override for the default leader predicate.
```

--------------------------------

### Install TVM

Source: https://tvm.apache.org/docs/_downloads/e2a9c4bfdec7a9365ef67c5335b1aaa4/tir_transformation.ipynb

Installs the latest development build of TVM from PyPI.

```shell
%%shell
# Installs the latest dev build of TVM from PyPI. If you wish to build
# from source, see https://tvm.apache.org/docs/install/from_source.html
pip install apache-tvm --pre
```

--------------------------------

### initWebGPU

Source: https://tvm.apache.org/docs/reference/api/typedoc/classes/Instance.html

Initialize WebGPU in the runtime.

```APIDOC
## initWebGPU

### Description
Initialize webgpu in the runtime.

### Parameters
- **device** (GPUDevice) - Required - The given GPU device.

### Returns
- **void**
```

--------------------------------

### One-hot encoding example

Source: https://tvm.apache.org/docs/reference/api/python/topi.html

Example demonstrating the output structure of the one_hot operation.

```python
indices = [0, 1, 2]

topi.one_hot(indices, 3) =
    [[1, 0, 0],
     [0, 1, 0],
     [0, 0, 1]]
```

--------------------------------

### Decompose Padding IR Example

Source: https://tvm.apache.org/docs/reference/api/python/s_tir/schedule.html

Example of an IR state after applying decompose-padding.

```python
@T.prim_func(s_tir=True)
def main(
    A: T.Buffer((127, 127), "float32"),
    B: T.Buffer((127, 127), "float32"),
    C: T.Buffer((127, 127), "float32"),
):
    # with T.sblock("root"):
    A_pad = T.sblock_alloc_buffer((128, 128))
    B_pad = T.sblock_alloc_buffer((128, 128))
    C_pad = T.sblock_alloc_buffer((128, 128))
    for i0, i1 in T.grid(128, 128):
        with T.sblock("A_pad"):
            v0, v1 = T.axis.remap("SS", [i0, i1])
            A_pad[v0, v1] = T.if_then_else(
                v0 < 127 and v1 < 127,
                A[v0, v1],
                T.float32(0),
            )
    for i0, i1 in T.grid(128, 128):
        with T.sblock("B_pad"):
            v0, v1 = T.axis.remap("SS", [i0, i1])
            B_pad[v0, v1] = T.if_then_else(
                v0 < 127 and v1 < 127,
                B[v0, v1],
                T.float32(0),
            )
    for i0, i1, i2 in T.grid(128, 128, 128):
        with T.sblock("C_shared"):
            i, j, k = T.axis.remap("SSR", [i0, i1, i2])
            with T.init():
                C_pad[i, j] = T.float32(0)
            C_pad[i, j] = C_pad[i, j] + A_pad[i, k] * B_pad[k, j]
    for i0, i1 in T.grid(127, 127):
        with T.sblock("C_pad"):
            v0, v1 = T.axis.remap("SS", [i0, i1])
            C[v0, v1] = C_pad[v0, v1]
```

--------------------------------

### Initialize dsmem demonstration

Source: https://tvm.apache.org/docs/_sources/tirx/tile_primitives/copy_async/dsmem.rst

Sets up the parameters for a 2-CTA cluster demonstration involving a float16 tile copy.

```python
from tvm.backend.cuda.lang import MBarrier

shape, dtype, CLUSTER_N = (128, 64), "float16", 2
src_layout = dst_layout = TileLayout(S[128, 64])
copy_bytes = 128 * 64 * 2
r = (slice(0, 128), slice(0, 64))
```

--------------------------------

### List Installed Kernels

Source: https://tvm.apache.org/docs/_sources/tirx/install.rst

Display the list of installed kernels in JSON format.

```bash
python -m tirx_kernels.registry --format json
```

--------------------------------

### Verify TIRx Installation

Source: https://tvm.apache.org/docs/_sources/tirx/install.rst

Check the installation by importing the tvm and tvm.tirx modules.

```bash
python -c "import tvm, tvm.tirx; print(tvm.__version__)"
```

--------------------------------

### InitializeWithTuneContext(const TuneContext &context)

Source: https://tvm.apache.org/docs/reference/api/doxygen/classtvm_1_1s__tir_1_1meta__schedule_1_1PyPostprocNode.html

Initializes the postprocessor with a tuning context.

```APIDOC
## InitializeWithTuneContext(const TuneContext &context)

### Description
Initializes the postprocessor with the provided tuning context. This method should be called only once before any other methods.

### Parameters
- **context** (const TuneContext &) - Required - The tuning context for initialization.
```

--------------------------------

### Initialize Build Directory

Source: https://tvm.apache.org/docs/_sources/install/from_source.rst

Prepare the build directory and copy the default CMake configuration file.

```bash
cd tvm
rm -rf build && mkdir build && cd build
# Specify the build configuration via CMake options
cp ../cmake/config.cmake .
```

--------------------------------

### Nonzero operator example

Source: https://tvm.apache.org/docs/reference/api/python/relax/op.html

Example demonstrating the output of the nonzero operator on a 2D tensor.

```python
x = [[0, 1],
     [2, 0]]
nonzero(x) = [[0, 1],
              [1, 0]]
```

--------------------------------

### Configure and enter a PassContext scope

Source: https://tvm.apache.org/docs/reference/api/doxygen/classtvm_1_1transform_1_1PassContext.html

Demonstrates creating a new PassContext, setting the optimization level, and entering a scope using the With RAII wrapper.

```cpp
auto new_ctx = PassContext::Create();
ctx->opt_level = 2;
With<PassContext> scope(ctx);
// pass context in effect.
```

--------------------------------

### Transform layout example

Source: https://tvm.apache.org/docs/reference/api/python/s_tir/schedule.html

Demonstrates the initial state, the application of transform_layout, and the resulting IR.

```python
@T.prim_func(s_tir=True)
def before_transform_layout(a: T.handle, c: T.handle) -> None:
    A = T.match_buffer(a, (128, 128), "float32")
    B = T.sblock_alloc_buffer((128, 128), "float32")
    C = T.match_buffer(c, (128, 128), "float32")
    for i, j in T.grid(128, 128):
        with T.sblock("B"):
            vi, vj = T.axis.remap("SS", [i, j])
            B[vi, vj] = A[vi, vj] * 2.0
    for i, j in T.grid(128, 128):
        with T.sblock("C"):
            vi, vj = T.axis.remap("SS", [i, j])
            C[vi, vj] = B[vi, vj] + 1.0
```

```python
sch = tvm.s_tir.Schedule(before_storage_align)
sch.transform_layout(sch.get_sblock("B"), buffer=("write",0),
                     index_map=lambda m, n: (m // 16, n // 16, m % 16, n % 16))
print(sch.mod["main"].script())
```

```python
@T.prim_func(s_tir=True)
def two_elementwise_transformed_intermediate_buffer(a: T.handle, c: T.handle) -> None:
    A = T.match_buffer(a, (128, 128), "float32")
    B = T.sblock_alloc_buffer((8, 8, 16, 16), "float32")
    C = T.match_buffer(c, (128, 128), "float32")
    for i, j in T.grid(128, 128):
        with T.sblock("B"):
            vi, vj = T.axis.remap("SS", [i, j])
            B[vi // 16, vj // 16, vi % 16, vj % 16] = A[vi, vj] * 2.0
    for i, j in T.grid(128, 128):
        with T.sblock("C"):
            vi, vj = T.axis.remap("SS", [i, j])
            C[vi, vj] = B[vi // 16, vj // 16, vi % 16, vj % 16] + 1.0
```

--------------------------------

### Install Python Dependencies

Source: https://tvm.apache.org/docs/_sources/install/from_source.rst

Install necessary and optional Python dependencies for TVM features.

```bash
pip3 install numpy cython
```

```bash
pip3 install tornado
```

```bash
pip3 install tornado psutil 'xgboost>=1.1.0' cloudpickle
```

--------------------------------

### Install tvm-ffi Package

Source: https://tvm.apache.org/docs/_sources/install/from_source.rst

Install the required tvm-ffi package after the build process is complete.

```bash
cd 3rdparty/tvm-ffi; pip install .; cd ..
```

--------------------------------

### RPC Deployment Workflow

Source: https://tvm.apache.org/docs/_downloads/148819f3421b8d89b1723c3e15e3f19f/cross_compilation_and_rpc.ipynb

Demonstrate the steps to connect to a remote device and upload compiled artifacts for inference.

```python
if local_demo:
    # For demonstration, show the code structure without execution
    print("\nRPC workflow (works for any remote device):")
    print("=" * 50)
    print("1. Start RPC server on target device:")
    print("   python -m tvm.exec.rpc_server --host 0.0.0.0 --port=9090")
    print("\n2. Connect from local machine:")
    print("   remote = rpc.connect('DEVICE_IP', 9090)")
    print("\n3. Upload compiled library:")
    print("   remote.upload('model_deployed.so')")
    print("   remote.upload('model_params.npz')")
    print("\n4.")
```

--------------------------------

### Install Python test dependencies

Source: https://tvm.apache.org/docs/_sources/contribute/pull_request.rst

Install required packages for running Python tests.

```bash
pip install --user pytest pytest-xdist Cython
```

--------------------------------

### Interact with Relax Virtual Machine in Python

Source: https://tvm.apache.org/docs/_sources/arch/relax_vm.rst

Demonstrates compiling a module, creating a VM, and executing functions using both direct and stateful interfaces.

```python
import tvm
from tvm import relax
import numpy as np

# Compile
ex = tvm.compile(MyModule, target="llvm")

# Create VM
vm = relax.VirtualMachine(ex, tvm.cpu())

# Direct invocation
inp = tvm.runtime.tensor(np.random.rand(3, 4).astype("float32"))
result = vm["main"](inp)

# Stateful interface (useful for RPC)
vm.set_input("main", inp)
vm.invoke_stateful("main")
output = vm.get_outputs("main")
```

--------------------------------

### Build an IRModule to VM Executable

Source: https://tvm.apache.org/docs/reference/api/python/relax/relax.html

Example demonstrating the compilation of a Relax IRModule into an executable using a specified target.

```python
class InputModule:
    @R.function
    def foo(x: Tensor((3, 4), "float32"), y: Tensor((3, 4), "float32")):
        z = R.add(x, y)
        return z

mod = InputModule
target = tvm.target.Target("llvm", host="llvm")
ex = tvm.compile(mod, target)
```

--------------------------------

### SetupTrainer

Source: https://tvm.apache.org/docs/reference/api/python/relax/training.html

Transforms a backbone module into a complete, legalized trainer module by appending loss functions and an optimizer.

```APIDOC
class tvm.relax.training.SetupTrainer(loss, optimizer, loss_args, legalize=True)

Parameters:
- loss (Loss) - The loss function to append.
- optimizer (Optimizer) - The optimizer function.
- loss_args (list[TensorType]) - Arguments for the loss function.
- legalize (bool) - Whether to legalize the module (default: True).
```

--------------------------------

### TVMScript Module Example

Source: https://tvm.apache.org/docs/reference/api/python/relax/block_builder.html

An example of a TVMScript module structure that can be built using the block builder.

```python
@tvm.script.ir_module
class Module:
    @T.prim_func(s_tir=True)
    def te_func(var_rxplaceholder: T.handle, var_compute: T.handle, n: T.int64) -> None:
        rxplaceholder = T.match_buffer(var_rxplaceholder, [n + T.int64(1)],
                                       dtype="float32")
        compute = T.match_buffer(var_compute, [n + T.int64(1)], dtype="float32")
        # body
        # with T.sblock("root")
        for i0 in T.serial(0, n + T.int64(1)):
            with T.sblock("compute"):
                i = T.axis.spatial(n + T.int64(1), i0)
                T.reads([rxplaceholder[i]])
                T.writes([compute[i]])
                compute[i] = rxplaceholder[i]

    @R.function
    def rx_func(x: Tensor((n,), "float32"), y: Tensor(((n + 1),), "float32"))
        -> Tensor(None, "float32", ndim=-1):
        # block 0
        gv = relax.call_tir(te_func, (y, n), R.Tensor((n + 1,), "float32"))
        return gv
```

--------------------------------

### Make

Source: https://tvm.apache.org/docs/reference/api/doxygen/classtvm_1_1VirtualDeviceCache.html

Creates or retrieves a unique VirtualDevice instance based on the provided device parameters.

```APIDOC
## VirtualDevice Make(int device_type, int virtual_device_id, Target target, MemoryScope memory_scope)

### Description
Returns the unique `VirtualDevice` representing the given fields. This method ensures that identical configurations result in the same object instance.

### Parameters
- **device_type** (int) - Optional - The device type (default: kInvalidDeviceType).
- **virtual_device_id** (int) - Optional - The virtual device ID (default: -1).
- **target** (Target) - Optional - The target configuration (default: {}).
- **memory_scope** (MemoryScope) - Optional - The memory scope (default: {}).
```

--------------------------------

### Operator Registration Example

Source: https://tvm.apache.org/docs/reference/api/doxygen/ir_2op_8h.html

Example usage of the TVM_REGISTER_OP macro to define an 'add' operator and its attributes.

```cpp
TVM_REGISTER_OP("add")
.describe("add two inputs together")
.set_num_inputs(2)
.set_attr<OpKernel>("gpu_kernel", AddKernel);
```

--------------------------------

### Generated TIRx IR Example

Source: https://tvm.apache.org/docs/_sources/tirx/tile_primitives/copy/ldstmatrix.rst

Example of the generated TIRx IR for a specific matrix configuration.

```python
for mm in Tx.unroll(1):
    Tx.ptx["ldmatrix.sync.aligned.m8n8.x2.shared.b16"](
        r_local[0], r_local[2], smem_ptr)
```

--------------------------------

### Compile and initialize the Virtual Machine

Source: https://tvm.apache.org/docs/_sources/how_to/tutorials/optimize_llm.rst

Uses the registered 'opt_llm' pipeline to compile the module and initialize the TVM Virtual Machine.

```python
with target:
    ex = tvm.compile(mod, target, relax_pipeline=relax.get_pipeline("opt_llm"))
    vm = relax.VirtualMachine(ex, dev)
```

--------------------------------

### TVMScript Syntax Sugar Examples

Source: https://tvm.apache.org/docs/_downloads/315fcda965a8d605f81705edf19ea2c6/relax_creation.ipynb

Examples of shorthand syntax versus normalized IR output.

```python
lv0 = R.matmul(data, R.permute_dims(w0)) + b0
```

```python
lv: R.Tensor((784, 128), dtype="float32") = R.permute_dims(w0, axes=None)
lv1: R.Tensor((n, 128), dtype="float32") = R.matmul(data, lv)
lv0: R.Tensor((n, 128), dtype="float32") = R.add(lv1, b0)
```

--------------------------------

### Launch and Query RPC Server

Source: https://tvm.apache.org/docs/_downloads/148819f3421b8d89b1723c3e15e3f19f/cross_compilation_and_rpc.ipynb

Commands to start the RPC server on the target and verify connectivity via the tracker.

```shell
# Through an RPC proxy.
python3 -m tvm.exec.rpc_server \
    --host RPC_PROXY_IP \
    --port RPC_PROXY_PORT \
    --through-proxy \
    --key RPC_KEY

# Directly to an RPC tracker.
python3 -m tvm.exec.rpc_server \
    --tracker RPC_TRACKER_IP:RPC_TRACKER_PORT \
    --key RPC_KEY
```

```shell
python3 -m tvm.exec.query_rpc_tracker --host RPC_TRACKER_IP --port RPC_TRACKER_PORT
```

--------------------------------

### Start RPC Tracker and Proxy

Source: https://tvm.apache.org/docs/_downloads/148819f3421b8d89b1723c3e15e3f19f/cross_compilation_and_rpc.ipynb

Commands to initialize the RPC tracker and proxy services on a host machine.

```shell
python3 -m tvm.exec.rpc_tracker --host RPC_TRACKER_IP --port 9190 --port-end 9191
```

```shell
python3 -m tvm.exec.rpc_proxy \
    --host RPC_PROXY_IP \
    --port 9090 \
    --port-end 9091 \
    --tracker RPC_TRACKER_IP:RPC_TRACKER_PORT
```

--------------------------------

### Install System Dependencies on Ubuntu/Debian

Source: https://tvm.apache.org/docs/_sources/install/from_source.rst

Install required system libraries for TVM on Debian-based distributions.

```bash
sudo apt update
sudo apt install zlib1g-dev libxml2-dev
```

--------------------------------

### Create

Source: https://tvm.apache.org/docs/reference/api/doxygen/classtvm_1_1runtime_1_1vm_1_1VirtualMachine.html

Creates a new instance of the VirtualMachine.

```APIDOC
## static ffi::ObjectPtr< VirtualMachine > Create()

### Description
Creates a specific instance of the VirtualMachine.

### Returns
- **ffi::ObjectPtr< VirtualMachine >** - A pointer to the newly created VirtualMachine instance.
```

--------------------------------

### Start Jupyter Notebook

Source: https://tvm.apache.org/docs/_sources/install/docker.rst

Launch a Jupyter notebook server inside the container environment.

```bash
jupyter notebook
```

```bash
jupyter notebook --ip=0.0.0.0
```

--------------------------------

### tvm.topi.gpu.topk

Source: https://tvm.apache.org/docs/reference/api/python/topi.html

Get the top k elements in an input tensor along the given axis.

```APIDOC
## tvm.topi.gpu.topk(data, k=1, axis=-1, ret_type='both', is_ascend=False, dtype='int64')

### Description
Get the top k elements in an input tensor along the given axis.
```

--------------------------------

### Pad Einsum example

Source: https://tvm.apache.org/docs/reference/api/python/s_tir/schedule.html

Demonstrates applying the pad_einsum primitive to a block matching the Einsum pattern.

```python
@T.prim_func(s_tir=True)
def before_pad_einsum(
    A: T.Buffer((127, 127), "float32"),
    B: T.Buffer((127, 127), "float32"),
    C: T.Buffer((127, 127), "float32"),
) -> None:
    for i0, i1, i2 in T.grid(127, 127, 127):
        with T.sblock("C_shared"):
            i, j, k = T.axis.remap("SSR", [i0, i1, i2])
            with T.init():
                C[i, j] = T.float32(0)
            C[i, j] = C[i, j] + A[i, k] * B[k, j]
```

```python
sch = tvm.s_tir.Schedule(before_pad_einsum, debug_mask="all")
block = sch.get_sblock("C_shared")
sch.pad_einsum(block, [32, 32, 32])
print(sch.mod["main"].script())
```

--------------------------------

### Install TFLite Frontend Dependency

Source: https://tvm.apache.org/docs/_sources/install/from_source.rst

Install the specific TensorFlow version required for TFLite frontend testing.

```bash
pip install tensorflow==2.19.0
```

--------------------------------

### tvm.support.libinfo()

Source: https://tvm.apache.org/docs/reference/api/python/support.html

Returns a dictionary containing compile-time information about the TVM installation.

```APIDOC
## tvm.support.libinfo()

### Description
Returns a dictionary of compile-time info. This function acts as a minimal Python fallback, synthesizing values from build-time hints.

### Returns
- **dict**: A dictionary containing compile-time configuration information.
```

--------------------------------

### tvm.topi.gpu.topk_thrust

Source: https://tvm.apache.org/docs/reference/api/python/topi.html

Get the top k elements in an input tensor along the given axis using the thrust library.

```APIDOC
## tvm.topi.gpu.topk_thrust(data, k=1, axis=-1, ret_type='both', is_ascend=False, dtype='int64', workspace=None)

### Description
Get the top k elements in an input tensor along the given axis.

### Parameters
- **data** (tvm.te.Tensor) - Required - The input tensor.
- **k** (int) - Optional - Number of top elements to select. Return all elements if k < 1.
- **axis** (int) - Optional - Axis along which to sort the input tensor.
- **ret_type** (str) - Optional - The return type [both, values, indices].
- **is_ascend** (boolean) - Optional - Whether to sort in ascending or descending order.
- **dtype** (string) - Optional - The data type of the indices output.
- **workspace** (Optional[tvm.te.Tensor]) - Optional - A buffer to store intermediate results.

### Returns
- **out** (tvm.te.Tensor or List[tvm.te.Tensor]) - The computed result.
```

--------------------------------

### Launch RPC Server

Source: https://tvm.apache.org/docs/how_to/tutorials/cross_compilation_and_rpc.html

Commands to start an RPC server either through a proxy or directly to a tracker.

```bash
# Through an RPC proxy.
python3 -m tvm.exec.rpc_server \
    --host RPC_PROXY_IP \
    --port RPC_PROXY_PORT \
    --through-proxy \
    --key RPC_KEY

# Directly to an RPC tracker.
python3 -m tvm.exec.rpc_server \
    --tracker RPC_TRACKER_IP:RPC_TRACKER_PORT \
    --key RPC_KEY
```

--------------------------------

### DSMEM Copy Primitive Example

Source: https://tvm.apache.org/docs/_sources/tirx/tile_primitives/copy_async/dsmem.rst

A complete example of using Tx.tile.copy_async with the dsmem dispatch to transfer data between CTAs.

```python
@Tx.prim_func
def dsmem_copy(A_ptr: Tx.handle, B_ptr: Tx.handle):
    A = Tx.match_buffer(A_ptr, shape, dtype); B = Tx.match_buffer(B_ptr, shape, dtype)
    Tx.device_entry()
    cbx = Tx.cta_id_in_cluster([CLUSTER_N]); Tx.cta_id([CLUSTER_N]); tid = Tx.thread_id([1])
    pool = Tx.SMEMPool()
    src_raw = pool.alloc([8192], dtype, align=128)
    src_smem = Tx.decl_buffer(list(shape), dtype, src_raw.data,
                             byte_offset=src_raw.byte_offset,
                             scope="shared.dyn", layout=src_layout)
    dst_raw = pool.alloc([8192], dtype, align=128)
    dst_smem = Tx.decl_buffer(list(shape), dtype, dst_raw.data,
                             byte_offset=dst_raw.byte_offset,
                             scope="shared.dyn", layout=dst_layout)
    mbar = MBarrier(pool, 1); pool.commit()
    mbar.init(1); Tx.ptx.fence.mbarrier_init.release.cluster(); Tx.cuda.cluster_sync()
    if tid == 0:
        if cbx == 0:                                      # source CTA
            Tx.tile.copy(src_smem[r], A[r])                    # global -> local shared
            Tx.ptx.fence.proxy.async_.shared__cta()
            Tx.tile.copy_async(dst_smem[r], src_smem[r], dispatch="dsmem",
                          mbar=mbar.ptr_to([0]), remote_cta_id=Tx.int32(1))   # -> CTA 1
        else:                                             # destination CTA
            Tx.ptx.mbarrier.arrive.expect_tx.shared.b64(mbar.ptr_to([0]), Tx.uint32(copy_bytes))
            mbar.wait(0, 0)
            Tx.tile.copy(B[r], dst_smem[r])                    # remote shared -> global
    Tx.cuda.cluster_sync()
```

--------------------------------

### Defining and Using a Context with tvm::With

Source: https://tvm.apache.org/docs/reference/api/doxygen/classtvm_1_1With.html

Shows the structure of a context class and how to use the With wrapper to manage its scope.

```cpp
// context class
class MyContext {
private:
friend class With<MyContext>;
MyContext(arguments);
void EnterWithScope();
void ExitWithScope();
};
{
With<MyContext> scope(arguments);
// effect take place.
}
```

--------------------------------

### tvm.topi.vision.get_valid_counts

Source: https://tvm.apache.org/docs/reference/api/python/topi.html

Get valid count of bounding boxes given a score threshold and rearrange valid boxes to the top.

```APIDOC
## tvm.topi.vision.get_valid_counts

### Description
Get valid count of bounding boxes given a score threshold. Also moves valid boxes to the top of input data.

### Parameters
- **data** (tvm.te.Tensor) - Required - Input data. 3-D tensor with shape [batch_size, num_anchors, elem_length].
- **score_threshold** (float) - Optional - Lower limit of score for valid bounding boxes.
- **id_index** (int) - Optional - Index of the class categories, -1 to disable.
- **score_index** (int) - Optional - Index of the scores/confidence of boxes.

### Returns
- **valid_count** (tvm.te.Tensor) - 1-D tensor for valid number of boxes, shape [batch_size].
- **out_tensor** (tvm.te.Tensor) - Rearranged data tensor, shape [batch_size, num_anchors, elem_length].
- **out_indices** (tvm.te.Tensor) - Related index in input data, shape [batch_size, num_anchors].
```

--------------------------------

### Install Apache TVM via pip

Source: https://tvm.apache.org/docs/_sources/install/pypi.rst

Use this command for a standard installation of the Apache TVM Python package.

```bash
pip install apache-tvm
```

--------------------------------

### Connect and Deploy to Remote Device

Source: https://tvm.apache.org/docs/_downloads/148819f3421b8d89b1723c3e15e3f19f/cross_compilation_and_rpc.ipynb

Establishes an RPC connection to a remote device, uploads the compiled library and parameters, and initializes the Virtual Machine.

```python
device_host = "192.168.1.100"  # Replace with your device IP
    device_port = 9090
    remote = rpc.connect(device_host, device_port)
    print(f"Connected to remote device at {device_host}:{device_port}")

    # Upload library and parameters to remote device
    remote.upload(lib_path)
    remote.upload(params_path)
    print("Uploaded files to remote device")

    # Load the library on the remote device
    lib = remote.load_module("model_deployed.so")

    # Choose device on remote machine
    # For CPU: dev = remote.cpu()
    # For CUDA GPU: dev = remote.cuda(0)
    # For OpenCL: dev = remote.cl(0)
    dev = remote.cpu()

    # Create VM and load parameters
    vm = relax.VirtualMachine(lib, dev)

    # Load parameters from the uploaded file
    # Note: In practice, you might load this from the remote filesystem
    params_npz = np.load(params_path)
    remote_params = [tvm.runtime.tensor(params_npz[f"p_{i}"], dev) for i in range(len(params_npz))]
```

--------------------------------

### Reshape operator dimension inference examples

Source: https://tvm.apache.org/docs/reference/api/python/relax/frontend.html

Examples demonstrating how the -1 dimension inference works for the reshape operator at compile-time.

```text
x.shape = (2, 3, 4), shape = (6, 1, -1), result.shape = (6, 1, 4)
x.shape = (2, 3, 4), shape = (3, -1, 8), result.shape = (3, 1, 8)
x.shape = (2, 3, 4), shape = (-1,), result.shape = (24,)
```

--------------------------------

### Initialize RPC Server with Callback

Source: https://tvm.apache.org/docs/reference/api/python/rpc.html

Use server_init_callback to register custom functions in the server environment before it starts.

```python
def server_init_callback():
    import tvm
    # must import mypackage here
    import mypackage

    tvm.register_global_func("function", mypackage.func)

server = rpc.Server(host, server_init_callback=server_init_callback)
```

--------------------------------

### Install TVM via Pip Local Project

Source: https://tvm.apache.org/docs/_sources/install/from_source.rst

Install TVM in editable mode using a local project path.

```bash
conda activate your-own-env
conda install python # make sure python is installed
export TVM_LIBRARY_PATH=/path-to-tvm/build
pip install -e /path-to-tvm
```

--------------------------------

### Join()

Source: https://tvm.apache.org/docs/reference/api/doxygen/functions_func_j.html

Joins a virtual device.

```APIDOC
## Join()

### Description
Joins a tvm::VirtualDevice instance.

### Signature
void Join()
```

--------------------------------

### Install Apache TVM with CUDA extras

Source: https://tvm.apache.org/docs/_sources/install/pypi.rst

Install the CUDA extra to include Python-side CUDA bindings for runtime compilation workflows.

```bash
pip install "apache-tvm[cuda]"
```

--------------------------------

### Causal Mask BottomRight Example (seq_len=2, seq_len_kv=4)

Source: https://tvm.apache.org/docs/reference/api/python/relax/op.html

Example matrix for BottomRight causal mask with 2 query and 4 key sequences.

```text
[[1, 1, 1, 0],
[1, 1, 1, 1]]
```

--------------------------------

### Server(String proxyHost, int proxyPort, String key)

Source: https://tvm.apache.org/docs/reference/api/javadoc/org/apache/tvm/rpc/Server.html

Constructor to initialize an RPC server that connects to a specified proxy.

```APIDOC
## Server(String proxyHost, int proxyPort, String key)

### Description
Starts an RPC server that connects to a proxy server.

### Parameters
- **proxyHost** (String) - Required - The hostname of the proxy server.
- **proxyPort** (int) - Required - The port of the proxy server.
- **key** (String) - Required - The key used to identify the server to the proxy.
```

--------------------------------

### Execution Time Summary Output

Source: https://tvm.apache.org/docs/deep_dive/tensor_ir/tutorials/tir_transformation.html

Example output showing the performance metrics of the executed module.

```text
Execution time summary:
 mean (ms)   median (ms)    max (ms)     min (ms)     std (ms)
   2.5589       2.5589       2.5589       2.5589       0.0000
```

--------------------------------

### Causal Mask TopLeft Example (seq_len=2, seq_len_kv=4)

Source: https://tvm.apache.org/docs/reference/api/python/relax/op.html

Example matrix for TopLeft causal mask with 2 query and 4 key sequences.

```text
[[1, 0, 0, 0],
[1, 1, 0, 0]]
```

--------------------------------

### Causal Mask BottomRight Example (seq_len=4, seq_len_kv=2)

Source: https://tvm.apache.org/docs/reference/api/python/relax/op.html

Example matrix for BottomRight causal mask with 4 query and 2 key sequences.

```text
[[1, 1],
[1, 1],
[1, 1],
[1, 1]]
```

--------------------------------

### PreTuning()

Source: https://tvm.apache.org/docs/reference/api/doxygen/classtvm_1_1s__tir_1_1meta__schedule_1_1PySearchStrategyNode.html

Initializes the search strategy before the tuning process begins. This method sets up the necessary parameters and state for the tuning session.

```APIDOC
## void PreTuning(int max_trials, int num_trials_per_iter, const ffi::Array< s_tir::Schedule > &design_spaces, const ffi::Optional< Database > &database, const ffi::Optional< CostModel > &cost_model)

### Description
Prepares the search strategy for the tuning process. This should be called after initialization and before tuning begins.

### Parameters
- **max_trials** (int) - The maximum number of trials.
- **num_trials_per_iter** (int) - The number of trials per iteration.
- **design_spaces** (ffi::Array< s_tir::Schedule >) - The design spaces used during the tuning process.
- **database** (ffi::Optional< Database >) - The database used during the tuning process.
- **cost_model** (ffi::Optional< CostModel >) - The cost model used during the tuning process.
```

--------------------------------

### Causal Mask TopLeft Example (seq_len=4, seq_len_kv=2)

Source: https://tvm.apache.org/docs/reference/api/python/relax/op.html

Example matrix for TopLeft causal mask with 4 query and 2 key sequences.

```text
[[1, 0],
[1, 1],
[1, 1],
[1, 1]]
```

--------------------------------

### InitializeWithTuneContext

Source: https://tvm.apache.org/docs/reference/api/doxygen/classtvm_1_1s__tir_1_1meta__schedule_1_1SearchStrategyNode.html

Initializes the search strategy with a specific tuning context.

```APIDOC
## InitializeWithTuneContext

### Description
Initialize the search strategy with tuning context. This method is supposed to be called only once before every other method.

### Parameters
- **context** (TuneContext) - The tuning context for initialization.
```

--------------------------------

### Python numpydoc Docstring Example

Source: https://tvm.apache.org/docs/_sources/contribute/document.rst

Standard format for documenting Python functions using numpydoc, including parameters, return values, and usage examples.

```python
def myfunction(arg1, arg2, arg3=3):
    """Briefly describe my function.

    Parameters
    ----------
    arg1 : Type1
        Description of arg1

    arg2 : Type2
        Description of arg2

    arg3 : Type3, optional
        Description of arg3

    Returns
    -------
    rv1 : RType1
        Description of return type one

    Examples
    --------
    .. code:: python

        # Example usage of myfunction
        x = myfunction(1, 2)
    """
    return rv1
```

--------------------------------

### Get

Source: https://tvm.apache.org/docs/reference/api/doxygen/classtvm_1_1relax_1_1ExecBuilderNode.html

Finalizes the build.

```APIDOC
### Get

**Signature**: `ffi::ObjectPtr< vm::VMExecutable > Get()`

Finalizes the build, runs formalization, and returns the final VMExecutable.
```
Tokens:
Raw