### Configure Ascend Toolkit Environment

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/README.md

Source the environment setup script provided by the Ascend Toolkit installation.

```shell
source /path/to/install/Ascend/ascend-toolkit/set_env.sh
```

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

### Install Ascend Toolkit

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/README.md

Execute the Ascend Toolkit installation script. Ensure you have the correct version and path.

```shell
chmod +x Ascend-cann-toolkit_{ascend-cann-toolkit version}_linux-aarch64.run
./Ascend-cann-toolkit_{ascend-cann-toolkit version}_linux-aarch64.run --install
```

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

### Install TileLang dependencies

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/CONTRIBUTING.md

Run this command after cloning the repository to install the necessary project dependencies.

```bash
python setup.py install
```

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

### Full Example of tilelang.jit.compile for NPU

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/developer/npu runtime.md

A comprehensive example of using tilelang.jit.compile with NPU target, including output index, execution backend, verbosity, and pass configurations.

```python
# Full example

 kernel = tilelang.jit.compile(
     func=my_prim_func,
     out_idx=0,
     execution_backend="cython",
     target="npuir",  # 必须设置为 "npuir"
     verbose=True,
     pass_configs={
         "tir.disable_vectorize": True,
         "tl.dynamic_vectorize_size_bits": 256
     }
 )
```

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

### Install PyTorch NPU Support

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/README.md

Install the pybind11 and torch_npu Python packages.

```shell
pip install pybind11 torch_npu
```

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

### Tilelang VEXP2 Usage Example

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/数学操作/T.vexp2.md

Example demonstrating the use of T.vexp2 for element-wise base-2 exponentiation on tensors. This example is JIT-compiled for the 'npuir' target.

```python
@tilelang.jit(target="npuir")
def vec_exp2(M, N, block_M, block_N):
    m_num = M // block_M
    n_num = N // block_N
    dtype = "float16"
    BLOCK_SIZE = 20

    @T.prim_func
    def main(
        A: T.Tensor((M, N), dtype),
        B: T.Tensor((M, N), dtype),
        C: T.Tensor((M, N), dtype),
    ):
        with T.Kernel(BLOCK_SIZE, is_npu=True) as (cid, _):
            A_VEC = T.alloc_ub((block_M, block_N), dtype)
            B_VEC = T.alloc_ub((block_M, block_N), dtype)
            C_VEC = T.alloc_ub((block_M, block_N), dtype)
            for i in T.serial(T.ceildiv(m_num * n_num, BLOCK_SIZE)):
                block_id = i * BLOCK_SIZE + cid
                if block_id < m_num * n_num:
                    block_id_m = block_id // n_num
                    block_id_n = block_id % n_num
                    bx = block_id_m * block_M
                    by = block_id_n * block_N
                    T.copy(A[bx, by], A_VEC)
                    T.copy(B[bx, by], B_VEC)
                    T.vexp2(A_VEC, B_VEC, C_VEC)
                    T.copy(C_VEC, C[bx, by])

    return main

```

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

### Install Static Libraries

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/CMakeLists.txt

Installs static libraries ('tilelang_static', 'tvm_runtime') to the library destination if 'TILE_LANG_INSTALL_STATIC_LIB' is enabled. This is for deployment scenarios requiring static linking.

```cmake
if(TILE_LANG_INSTALL_STATIC_LIB)
  install(TARGETS tilelang_static tvm_runtime
    LIBRARY DESTINATION lib${LIB_SUFFIX}
  )
else()
  if(DEFINED TVM_PREBUILD_PATH)
    install(TARGETS tilelang tilelang_module
      RUNTIME DESTINATION bin
      LIBRARY DESTINATION lib${LIB_SUFFIX}
    )
  else()
    install(TARGETS tvm_runtime tilelang tilelang_module
      RUNTIME DESTINATION bin
      LIBRARY DESTINATION lib${LIB_SUFFIX}
    )
  endif()
endif()
```

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

### T.flip Usage Example

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/索引与元素操作/T.flip.md

An example demonstrating how to use the T.flip operation within a Tilelang JIT function.

```APIDOC
## Usage Example

This example shows how to use `T.flip` in a JIT-compiled function for the NPU.

```python
@tilelang.jit(target="npuir")
def vec_flip(block_M, block_N, dtype="float16"):
    BLOCK_SIZE = 1

    @T.prim_func
    def main(
        A: T.Tensor((block_M, block_N), dtype),
        C: T.Tensor((block_M, block_N), dtype),
    ):
        with T.Kernel(BLOCK_SIZE, is_npu=True) as (cid, _):
            A_VEC = T.alloc_ub((block_M, block_N), dtype)
            C_VEC = T.alloc_ub((block_M, block_N), dtype)

            T.copy(A, A_VEC)
            T.flip(A_VEC[:block_M, :block_N], C_VEC[:block_M, :block_N], 1)
            T.copy(C_VEC, C)

    return main
```
```

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

### Tilelang vrelu Implementation Example

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/数学操作/T.vrelu.md

Demonstrates how to implement the T.vrelu operation within a Tilelang JIT function. This example computes ReLU values for a tensor and copies data between host and device buffers.

```python
@tilelang.jit(target="npuir")
def vecrelu(M, N, block_M, block_N, dtype="float16"):
    @T.prim_func
    def vecrelu_(A: T.Tensor((M, N), dtype), B: T.Tensor((M, N), dtype)):
        with T.Kernel(T.ceildiv(N, block_N) * T.ceildiv(M, block_M), is_npu=True) as (
            cid,
            _,
        ):
            by = cid // T.ceildiv(N, block_N)
            bx = cid % T.ceildiv(N, block_N)

            A_BUF = T.alloc_shared((block_M, block_N), dtype)
            B_BUF = T.alloc_shared((block_M, block_N), dtype)

            T.copy(A[by * block_M, bx * block_N], A_BUF)
            T.vrelu(A_BUF, B_BUF)
            T.copy(B_BUF, B[by * block_M, bx * block_N])

    return vecrelu_
```

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

### Tilelang VEXP Implementation Example

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/数学操作/T.vexp.md

Example demonstrating how to implement the T.vexp operation within a Tilelang JIT function for Ascend NPUs. It shows tensor allocation, data copying, the vexp call, and output copying.

```python
@tilelang.jit(target="npuir")
def vec_exp(M, N, block_M, block_N, dtype="float16"):
    m_num = M // block_M
    n_num = N // block_N

    @T.prim_func
    def main(
        Input: T.Tensor((M, N), dtype),
        Output: T.Tensor((M, N), dtype),
    ):
        with T.Kernel(m_num * n_num, is_npu=True) as (cid, _):
            bx = cid // n_num
            by = cid % n_num

            ub_input = T.alloc_ub((block_M, block_N), dtype)
            ub_output = T.alloc_ub((block_M, block_N), dtype)

            T.copy(Input[bx * block_M, by * block_N], ub_input)
            T.vexp(ub_input, ub_output)
            T.copy(ub_output, Output[bx * block_M, by * block_N])

    return main
```

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

### Build AscendNPU-IR

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/README.md

Build the AscendNPU-IR using the provided installation script. An alternative allows specifying a local path for bishengir-compile.

```shell
cd tilelang-ascend
bash install_npuir.sh
bash install_npuir.sh --bishengir-path=/path/to/bishengir-compile
```

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

### Tilelang vsin Implementation Example

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/数学操作/T.vsin.md

An example demonstrating the implementation of the T.vsin operation using Tilelang's JIT compilation for Ascend NPUs. This snippet shows how to define a kernel for computing sine values on a tensor.

```python
import tilelang as T

@tilelang.jit(target="npuir")
def vecsin(M, N, block_M, block_N, dtype="float16"):
    @T.prim_func
    def vecsin_(A: T.Tensor((M, N), dtype), B: T.Tensor((M, N), dtype)):
        with T.Kernel(T.ceildiv(N, block_N) * T.ceildiv(M, block_M), is_npu=True) as (
            cid,
            _,
        ):
            by = cid // T.ceildiv(N, block_N)
            bx = cid % T.ceildiv(N, block_N)

            A_BUF = T.alloc_shared((block_M, block_N), dtype)
            B_BUF = T.alloc_shared((block_M, block_N), dtype)

            T.copy(A[by * block_M, bx * block_N], A_BUF)
            T.vsin(A_BUF, B_BUF)
            T.copy(B_BUF, B[by * block_M, bx * block_N])

    return vecsin_
```

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

### vabs_kernel Example

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/数学操作/T.vabs.md

An example demonstrating the usage of the T.vabs operator within a Tilelang JIT kernel for processing a 2D tensor in Expert Mode.

```APIDOC
## vabs_kernel Example

### Description
This example shows how to implement the `vabs` operation for a 2D tensor using Tilelang's JIT compilation for the NPU. It includes allocating memory, performing the `vabs` operation, and copying data back.

### Code
```python
@tilelang.jit(target="npuir")
def vabs_kernel(M, N, dtype="float16", out_dtype="float16"):
    @T.prim_func
    def main(src: T.Tensor((M, N), dtype), dst: T.Tensor((M, N), out_dtype)):
        # Use a single block to process the whole tensor
        with T.Kernel(1, is_npu=True) as (bid, _):
            # Allocate UB memory for input and output
            src_ub = T.alloc_ub((M, N), dtype)
            dst_ub = T.alloc_ub((M, N), out_dtype)
            # Copy data from GM to UB
            T.copy(src, src_ub)
            T.vabs(src_ub, dst_ub)
            # Copy results back to GM
            T.copy(dst_ub, dst)

    return main
```
```

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

### Element-wise Power Operation Example

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/数学操作/T.vpow.md

This example demonstrates how to implement an element-wise power operation using T.vpow within a Tilelang JIT function. It defines a kernel for NPU execution, copying data to shared buffers, performing the vpow operation, and copying results back.

```python
@tilelang.jit(target="npuir")
def vecpow(M, N, block_M, block_N, dtype="int32"):
    @T.prim_func
    def vecpow_(
        A: T.Tensor((M, N), dtype),
        B: T.Tensor((M, N), dtype),
        C: T.Tensor((M, N), dtype),
    ):
        with T.Kernel(T.ceildiv(N, block_N) * T.ceildiv(M, block_M), is_npu=True) as (
            cid,
            _,
        ):
            by = cid // T.ceildiv(N, block_N)
            bx = cid % T.ceildiv(N, block_N)

            A_BUF = T.alloc_shared((block_M, block_N), dtype)
            B_BUF = T.alloc_shared((block_M, block_N), dtype)
            C_BUF = T.alloc_shared((block_M, block_N), dtype)

            T.copy(A[by * block_M, bx * block_N], A_BUF)
            T.copy(B[by * block_M, bx * block_N], B_BUF)
            T.vpow(A_BUF, B_BUF, C_BUF)
            T.copy(C_BUF, C[by * block_M, bx * block_N])

    return vecpow_
```

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

### Vec Flip Example

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/索引与元素操作/T.flip.md

This example demonstrates how to use T.flip within a Tilelang JIT function for NPU execution. It defines a kernel to copy data, flip it along axis 1, and then copy the result.

```python
@tilelang.jit(target="npuir")
def vec_flip(block_M, block_N, dtype="float16"):
    BLOCK_SIZE = 1

    @T.prim_func
    def main(
        A: T.Tensor((block_M, block_N), dtype),
        C: T.Tensor((block_M, block_N), dtype),
    ):
        with T.Kernel(BLOCK_SIZE, is_npu=True) as (cid, _):
            A_VEC = T.alloc_ub((block_M, block_N), dtype)
            C_VEC = T.alloc_ub((block_M, block_N), dtype)

            T.copy(A, A_VEC)
            T.flip(A_VEC[:block_M, :block_N], C_VEC[:block_M, :block_N], 1)
            T.copy(C_VEC, C)

    return main
```

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

### Tilelang VAND Operation Example

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/逻辑操作/T.vand.md

This example demonstrates how to use the T.vand operation within a Tilelang JIT function targeting an NPU. It includes tensor allocation, copying data, performing the VAND operation, and copying the result back.

```python
@tilelang.jit(target="npuir")
def vec_and(block_M, block_N, dtype="float16"):

    BLOCK_SIZE = 1

    @T.prim_func
    def main(
        A: T.Tensor((block_M, block_N), dtype),
        B: T.Tensor((block_M, block_N), dtype),
        C: T.Tensor((block_M, block_N), dtype),
    ):
        with T.Kernel(BLOCK_SIZE, is_npu=True) as (cid, _):
            A_VEC = T.alloc_ub((block_M, block_N), dtype)
            B_VEC = T.alloc_ub((block_M, block_N), dtype)
            C_VEC = T.alloc_ub((block_M, block_N), dtype)

            T.copy(A, A_VEC)
            T.copy(B, B_VEC)
            T.vand(A_VEC, B_VEC, C_VEC)
            T.copy(C_VEC, C)

    return main

```

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

### 2D Tensor Atomic Add Example

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/原子操作/T.atomic_add.md

This example demonstrates an atomic add operation on a 2D tensor using Tilelang's JIT compilation for the NPUIR target. It includes kernel setup, data copying to shared memory, and the atomic add operation itself.

```python
@tilelang.jit(target="npuir")
def vec_atomic_add_2d(M, N, block_M, block_N, dtype="float32"):
    m_num = M // block_M
    n_num = N // block_N

    @T.prim_func
    def vecAtomicAdd2D(
        A: T.Tensor((M, N), dtype),
        B: T.Tensor((M, N), dtype),
        shape_M: T.int32,
        shape_N: T.int32,
    ):
        with T.Kernel(m_num * n_num, is_npu=True) as (cid, _):
            blockx = cid // n_num
            bx = blockx * block_M
            blocky = cid % n_num
            by = blocky * block_N
            A_VEC = T.alloc_shared((block_M, block_N), dtype)
            t0 = shape_M - bx
            tile_size_M = T.min(block_M, t0)
            t0 = shape_N - by
            tile_size_N = T.min(block_N, t0)
            T.copy(A[bx, by], A_VEC, [tile_size_M, tile_size_N])
            T.atomic_add(B[bx, by], A_VEC, [tile_size_M, tile_size_N])

    return vecAtomicAdd2D
```

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

### Element-wise XOR with Tilelang JIT

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/逻辑操作/T.vxor.md

This example demonstrates how to use T.vxor within a Tilelang JIT function for NPU execution. It includes tensor copying and the XOR operation on sub-tensors.

```python
@tilelang.jit(target="npuir")
def vec_exp(M, N, block_M, block_N, dtype="float16"):
    m_num = M // block_M
    n_num = N // block_N

    @T.prim_func
    def main(
        Input1: T.Tensor((M, N), dtype),
        Input2: T.Tensor((M, N), dtype),
        Output: T.Tensor((M, N), dtype),
    ):
        with T.Kernel(m_num * n_num, is_npu=True) as (cid, _):
            bx = cid // n_num
            by = cid % n_num
            ub_input1 = T.alloc_ub((block_M, block_N), dtype)
            ub_input2 = T.alloc_ub((block_M, block_N), dtype)
            ub_output = T.alloc_ub((block_M, block_N), dtype)

            T.copy(Input1[bx * block_M, by * block_N], ub_input1)
            T.copy(Input2[bx * block_M, by * block_N], ub_input2)
            T.vxor(ub_input1, ub_input2, ub_output)
            T.copy(ub_output, Output[bx * block_M, by * block_N])

    return main
```

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

### Tilelang Expression-Level Log2 Example

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/数学操作/T.vlog2.md

Demonstrates the expression-level T.log2 operation within a Tilelang prim_func. This example shows how to apply log2 to individual elements of a tensor in TIR.

```Python
@tilelang.jit(target="npuir")
def log2_expr_example(block_M):
    @T.prim_func
    def update_logsum(
        logsum: T.Tensor((block_M,), "float32"),
        scores_max: T.Tensor((block_M,), "float32"),
    ):
        scale = 0.5
        for i in T.Parallel(block_M):
            logsum[i] = T.log2(logsum[i]) + scores_max[i] * scale

    return update_logsum
```

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

### Example: Implementing arange for a Tensor

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/创建操作/T.arange.md

This example demonstrates how to use the arange operation within a Tilelang JIT function to fill tensors with sequential data. It shows the creation of two tensors, A and B, both filled using arange with different offsets and strides.

```python
@tilelang.jit(target="npuir")
def vec_arange(M, N, block_M, block_N, src_dtype="float32", dst_dtype="float16"):
    m_num = M // block_M
    n_num = N // block_N

    @T.prim_func
    def main(
        A: T.Tensor((M, N), dst_dtype),
        B: T.Tensor((M, N), dst_dtype),
    ):
        with T.Kernel(m_num * n_num, is_npu=True) as (cid, _):
            bx_ = cid // n_num
            bx = bx_ * block_M
            by_ = cid % n_num
            by = by_ * block_N

            A_VEC = T.alloc_ub((block_M, block_N), dst_dtype)
            B_VEC = T.alloc_ub((block_M, block_N), dst_dtype)
            strides = [1, 2]
            T.arange(A_VEC, strides, offset=1)
            T.arange(B_VEC, strides)
            T.copy(A_VEC, A[bx, by])
            T.copy(B_VEC, B[bx, by])

    return main

```

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

### Install Python Dependencies for Ascend

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/README.md

Install required Python packages for Ascend Toolkit development, ensuring compatible NumPy version.

```shell
pip3 install attrs cython 'numpy>=1.19.2,<=1.24.0' decorator sympy cffi pyyaml pathlib2 psutil protobuf==3.20.0 scipy requests absl-py
```

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

### Tilelang atomic_addx4 Example

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/原子操作/T.atomic_addx4.md

Demonstrates a 2D tensor atomic add calculation using tilelang.jit. This function defines a Tilelang program for atomic addition on the NPU.

```python
import tilelang as T

@T.jit(target="npuir")
def run_atomic_addx4(M, N, block_M, block_N, dtype="float32"):
    m_num = M // block_M
    n_num = N // block_N

    @T.prim_func
    def atomicAddx4Program(
        A: T.Tensor((M, N), dtype),
        B: T.Tensor((M, N), dtype),
        shape_M: T.int32,
        shape_N: T.int32,
    ):
        with T.Kernel(m_num * n_num, is_npu=True) as (cid, _):
            blockx = cid // n_num
            blocky = cid % n_num
            A_VEC = T.alloc_ub((1, 4), dtype)

            for i, j in T.Parallel(block_M, block_N // 4):
                bx = blockx * block_M + i
                by = blocky * block_N + j * 4

                T.copy(A[bx, by], A_VEC, [1, 4])
                T.atomic_addx4(B[bx, by], A_VEC, [1, 4])

    return atomicAddx4Program
```

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

### Vector Element-wise OR Operation Example

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/逻辑操作/T.vor.md

This example demonstrates how to implement a vector element-wise OR operation using Tilelang's T.vor. It defines a JIT-compiled function for NPU execution, handling tensor copying and the core OR operation within a kernel.

```python
@tilelang.jit(target="npuir")
def vec_exp(M, N, block_M, block_N, dtype="float16"):
    m_num = M // block_M
    n_num = N // block_N

    @T.prim_func
    def main(
        Input1: T.Tensor((M, N), dtype),
        Input2: T.Tensor((M, N), dtype),
        Output: T.Tensor((M, N), dtype),
    ):
        with T.Kernel(m_num * n_num, is_npu=True) as (cid, _):
            bx = cid // n_num
            by = cid % n_num
            ub_input1 = T.alloc_ub((block_M, block_N), dtype)
            ub_input2 = T.alloc_ub((block_M, block_N), dtype)
            ub_output = T.alloc_ub((block_M, block_N), dtype)

            T.copy(Input1[bx * block_M, by * block_N], ub_input1)
            T.copy(Input2[bx * block_M, by * block_N], ub_input2)
            T.vor(ub_input1, ub_input2, ub_output)
            T.copy(ub_output, Output[bx * block_M, by * block_N])

    return main
```

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

### Tilelang NPUIR Vlog2 Kernel Example

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/数学操作/T.vlog2.md

Example of a Tilelang NPUIR kernel for the vlog2 operation. It defines a prim_func that processes tensor tiles, performing log2 using an intermediate buffer.

```Python
@tilelang.jit(target="npuir")
def vlog2_kernel(M, N, block_M, block_N, dtype="float16"):
    m_num = M // block_M
    n_num = N // block_N
    block_size = 8

    @T.prim_func
    def main(
        src: T.Tensor((M, N), dtype),
        dst: T.Tensor((M, N), dtype),
    ):
        with T.Kernel(block_size, is_npu=True) as (cid, _):
            src_ub = T.alloc_ub((block_M, block_N), dtype)
            dst_ub = T.alloc_ub((block_M, block_N), dtype)
            tmp_ub = T.alloc_ub((block_M, block_N), dtype)

            for i in T.serial(T.ceildiv(m_num * n_num, block_size)):
                block_id = i * block_size + cid
                if block_id < m_num * n_num:
                    block_id_m = block_id // n_num
                    block_id_n = block_id % n_num
                    bx = block_id_m * block_M
                    by = block_id_n * block_N

                    T.copy(src[bx, by], src_ub)
                    T.vlog2(src_ub, dst_ub, tmp_ub)
                    T.copy(dst_ub, dst[bx, by])

    return main
```

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

### TileLang Vector Add Kernel Example

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/README.md

This Python code initializes TileLang, clears the cache, and imports necessary libraries for NPU development. It sets up the environment for defining and executing NPU kernels.

```python
# Copyright (c) Tile-AI Corporation.
# Licensed under the MIT License.

import os

import tilelang
import tilelang.language as T  # Import TileLang DSL for kernel definition

import torch
import torch_npu  # Import NPU (Neural Processing Unit) backend support for PyTorch

# Clear any previously cached compiled kernels to ensure a clean run
tilelang.cache.clear_cache()

```

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

### General Reduction Output Format

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/tilelang/carver/README.md

Example of the dictionary-based output structure returned by the recommendation engine.

```python
{
  'block': [1, 128],
  'thread': [1, 128],
  'rstep': [64],
  ...
},
{
  'block': [2, 64],
  'thread': [2, 64],
  'rstep': [64],
  ...
},
...
{
  'block': [1, 16],
  'thread': [1, 16],
  'rstep': [512],
  'reduce_thread': [8],
  ...
}
```

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

### Tilelang vmul Example with Tensor Multiplication

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/数学操作/T.vmul.md

This example demonstrates how to use T.vmul for element-wise multiplication of two tensors with the same shape. Ensure operands are allocated on UB and the operation is within a T.Kernel context.

```python
@tilelang.jit(target="npuir")
def vec_mul(M, N, dtype):
    BLOCK_SIZE = 20

    @T.prim_func
    def main(
        A: T.Tensor((M, N), dtype),
        B: T.Tensor((M, N), dtype),
        C: T.Tensor((M, N), dtype),
    ):
        with T.Kernel(BLOCK_SIZE, is_npu=True) as (cid, _):
            A_VEC = T.alloc_ub((M, N), dtype)
            B_VEC = T.alloc_ub((M, N), dtype)
            C_VEC = T.alloc_ub((M, N), dtype)
            T.copy(A, A_VEC)
            T.copy(B, B_VEC)
            T.vmul(A_VEC, B_VEC, C_VEC)
            T.copy(C_VEC, C)

    return main
```

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

### Tilelang vselect Kernel Implementation

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/条件操作/T.vselect.md

This example demonstrates a Tilelang kernel that utilizes the vselect operation. It performs a comparison and then selects values based on the condition, copying data to and from shared memory.

```python
@tilelang.jit(target="npuir")
def select_kernel(N, block_M, dtype="float16"):
    grid = (N + block_M - 1) // block_M

    @T.prim_func
    def main(
        A: T.Tensor((N,), dtype),
        B: T.Tensor((N,), dtype),
        Out: T.Tensor((N,), dtype),
    ):
        with T.Kernel(grid, is_npu=True) as (bx, _):
            cond_ub = T.alloc_shared((block_M,), "bool")
            acc_A = T.alloc_shared((block_M,), dtype)
            acc_B = T.alloc_shared((block_M,), dtype)
            out_ub = T.alloc_shared((block_M,), dtype)

            for i in T.serial(grid):
                start = i * block_M
                end = T.min(start + block_M, N)
                cur_size = end - start

                T.copy(A[start:end], acc_A[0:cur_size])
                T.copy(B[start:end], acc_B[0:cur_size])

                T.vcmp(acc_A[0:cur_size], acc_B[0:cur_size], cond_ub[0:cur_size], "ge")
                T.vselect(
                    cond_ub[0:cur_size],
                    acc_A[0:cur_size],
                    acc_B[0:cur_size],
                    out_ub[0:cur_size],
                )
                T.copy(out_ub[0:cur_size], Out[start:end])

    return main

```

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

### Tilelang vcmp Kernel Example

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/比较操作/T.vcmp.md

This example demonstrates how to implement a comparison kernel using T.vcmp for element-wise greater than or equal to comparison between two tensors. Ensure input tensors A and B have the same shape and the output tensor C is of boolean type.

```python
@tilelang.jit(target="npuir")
def vcmp_kernel(M, N, dtype):
    BLOCK_SIZE = 1

    @T.prim_func
    def main(
        A: T.Tensor((M, N), dtype),
        B: T.Tensor((M, N), dtype),
        C: T.Tensor((M, N), "bool"),
    ):

        with T.Kernel(BLOCK_SIZE, is_npu=True) as (cid, _):
            A_ub = T.alloc_shared((M, N), dtype)
            B_ub = T.alloc_shared((M, N), dtype)
            C_ub = T.alloc_shared((M, N), "bool")

            T.copy(A, A_ub)
            T.copy(B, B_ub)
            T.vcmp(A_ub, B_ub, C_ub, "ge")
            T.copy(C_ub, C)

    return main
```

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

### vabs Kernel Implementation (Expert Mode)

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/数学操作/T.vabs.md

Example of implementing the T.vabs operator within a Tilelang JIT kernel for NPU execution. This demonstrates allocating UB memory, performing the vabs operation, and copying data back.

```python
@tilelang.jit(target="npuir")
def vabs_kernel(M, N, dtype="float16", out_dtype="float16"):
    @T.prim_func
    def main(src: T.Tensor((M, N), dtype), dst: T.Tensor((M, N), out_dtype)):
        # Use a single block to process the whole tensor
        with T.Kernel(1, is_npu=True) as (bid, _):
            # Allocate UB memory for input and output
            src_ub = T.alloc_ub((M, N), dtype)
            dst_ub = T.alloc_ub((M, N), out_dtype)
            # Copy data from GM to UB
            T.copy(src, src_ub)
            T.vabs(src_ub, dst_ub)
            # Copy results back to GM
            T.copy(dst_ub, dst)

    return main
```

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

### Tilelang cumsum Kernel Implementation

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/排序操作/T.vcumsum.md

This example demonstrates how to implement a cumulative sum kernel using Tilelang. It includes setting up the kernel, allocating memory, copying data, performing the cumulative sum, and copying results back. The `dim` and `reverse` parameters control the operation.

```python
@tilelang.jit(target="npuir")
def cumsum_kernel(M, N, dim, reverse, dtype="float16", accum_dtype="float16"):
    BLOCK_SIZE = 1

    @T.prim_func
    def main(src: T.Tensor((M, N), dtype), dst: T.Tensor((M, N), accum_dtype)):

        with T.Kernel(BLOCK_SIZE, is_npu=True) as (cid, _):
            # Allocate UB memory
            src_ub = T.alloc_ub((M, N), dtype)
            dst_ub = T.alloc_ub((M, N), accum_dtype)

            # Copy data from GM to UB
            T.copy(src, src_ub)

            # Perform cumulative sum
            T.cumsum(src_ub, dst_ub, dim=dim, reverse=reverse)

            # Copy results back to GM
            T.copy(dst_ub, dst)

    return main
```

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

### Two-Dimensional Tile Padding Example

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/shape操作/T.pad.md

Example demonstrating symmetric padding on the row dimension of a two-dimensional tile. This snippet shows how to pad a tensor using T.pad, including padding with a specific value and padding based on loop variables.

```python
@T.prim_func
def main(
    A: T.Tensor((M, N), src_dtype),
    B: T.Tensor((2 * M, 2 * N), dst_dtype),
    C: T.Tensor((M, N), dst_dtype),
):
    m_num = M // block_M
    n_num = N // block_N

    with T.Kernel(m_num * n_num, is_npu=True) as (cid, _):
        bx_ = cid // n_num
        bx = bx_ * block_M
        by_ = cid % n_num
        by = by_ * block_N

        A_VEC = T.alloc_ub((block_M, block_N), src_dtype)
        B_VEC = T.alloc_ub((2 * block_M, block_N), dst_dtype)
        C_VEC = T.alloc_ub((block_M + 2 * m_num * n_num, block_N), dst_dtype)

        T.copy(A[bx, by], A_VEC)

        # 1）在第 0 维两端各 pad block_M/2 行，pad_value = 0.0
        T.pad(A_VEC, B_VEC, 0.0, [block_M / 2, 0], [block_M / 2, 0])

        # 2）在第 0 维前后各 pad cid 行，作为示例
        T.pad(A_VEC, C_VEC, 0.0, [cid, 0], [cid, 0])

        T.copy(B_VEC, B[2 * bx, by])
        T.copy(C_VEC, C[bx, by])
```

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

### Matmul Template Output Format

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/tilelang/carver/README.md

Example output for matrix multiplication hints, including warp and tensor core configuration flags.

```python
{
  'block': [32, 64],
  'warp': [16, 32],
  'rstep': [128],
  'use_tc': True,
  ...
},
{
  'block': [64, 32],
  'warp': [32, 16],
  'rstep': [128],
  'use_tc': True,
  ...
},
...
{
  'block': [256, 32],
  'warp': [128, 16],
  'rstep': [32],
  'use_tc': True,
  ...
}
```

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

### vclamp with Tensor Bounds

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/比较操作/T.vclamp.md

Illustrates using vclamp with element-wise minimum and maximum bounds provided as tensors. The bounds tensors must be broadcastable to the source and destination tensors. This example is for float16 and float32.

```Python
@tilelang.jit(target="npuir")
def vclamp_tensor_kernel(M, N, dtype="float16"):
    block_size = 1

    @T.prim_func
    def main(
        src: T.Tensor((M, N), dtype),
        dst: T.Tensor((M, N), dtype),
        min_val: T.Tensor((M, N), dtype),
        max_val: T.Tensor((M, N), dtype),
    ):
        with T.Kernel(block_size, is_npu=True) as (cid, _):
            src_ub = T.alloc_shared((M, N), dtype)
            dst_ub = T.alloc_fragment((M, N), dtype)
            min_ub = T.alloc_shared((M, N), dtype)
            max_ub = T.alloc_shared((M, N), dtype)

            T.copy(src, src_ub)
            T.copy(min_val, min_ub)
            T.copy(max_val, max_ub)
            T.vclamp(src_ub, dst_ub, min_ub, max_ub)
            T.copy(dst_ub, dst)

    return main
```

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

### Using ceildiv in Kernel Definition

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/数学操作/T.ceildiv.md

This example demonstrates how to use T.ceildiv within a Tilelang JIT function to define kernel dimensions, specifically for calculating the number of blocks needed for N and M dimensions based on block sizes.

```python
import tilelang

@tilelang.jit(target="npuir")
def vec_add(block_M, block_N, dtype="float32"):
    M = T.symbolic("M")
    N = T.symbolic("N")

    @T.prim_func
    def main(
        A: T.Tensor((M, N), dtype),
        B: T.Tensor((M, N), dtype),
        C: T.Tensor((M, N), dtype),
    ):
        with T.Kernel(T.ceildiv(N, block_N) * T.ceildiv(M, block_M), is_npu=True) as (
            cid,
            _,
        ):
            # 代码段省略
            pass

    return main
```

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

### Sigmoid Kernel Implementation in Tilelang

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/数学操作/T.sigmoid.md

This example demonstrates how to implement a sigmoid kernel for a tensor of shape (M, N) using Tilelang. It utilizes `tilelang.jit` for NPU compilation and defines a primary function `main` that performs the sigmoid computation.

```python
@tilelang.jit(target="npuir")
def sigmoid_kernel(M, N, dtype):
    BLOCK_SIZE = 1

    @T.prim_func
    def main(src: T.Tensor((M, N), dtype), dst: T.Tensor((M, N), dtype)):

        with T.Kernel(BLOCK_SIZE, is_npu=True) as (cid, _):
            src_ub = T.alloc_shared((M, N), dtype)
            dst_ub = T.alloc_shared((M, N), dtype)

            T.copy(src, src_ub)
            T.vsigmoid(src_ub, dst_ub)
            T.copy(dst_ub, dst)

    return main
```

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

### Tilelang vtanh OP Usage Example

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/数学操作/T.vtanh.md

Demonstrates how to use the vtanh operation within a Tilelang JIT function for NPU execution. It includes tensor allocation, data copying, vtanh computation, and copying results back.

```python
@tilelang.jit(target="npuir")
def vec_tanh(M, N, dtype):
    BLOCK_SIZE = 8

    @T.prim_func
    def main(
        src: T.Tensor((M, N), dtype),
        dst: T.Tensor((M, N), dtype),
    ):
        with T.Kernel(BLOCK_SIZE, is_npu=True) as (cid, _):
            src_ub = T.alloc_ub((M, N), dtype)
            dst_ub = T.alloc_ub((M, N), dtype)

            T.copy(src, src_ub)
            T.vtanh(src_ub, dst_ub)
            T.copy(src_ub, dst)

    return main
```

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

### Vector Natural Logarithm Implementation

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/数学操作/T.vLn.md

This example demonstrates how to implement a vector natural logarithm operation using Tilelang's JIT compilation for NPUIR. It defines a prim_func that copies data to a temporary buffer, applies T.vLn, and then copies the result back.

```python
@tilelang.jit(target="npuir")
def vec_ln(M, N, block_M, block_N, dtype="float16"):
    m_num = M // block_M
    n_num = N // block_N

    @T.prim_func
    def main(
        Input: T.Tensor((M, N), dtype),
        Output: T.Tensor((M, N), dtype),
    ):
        with T.Kernel(m_num * n_num, is_npu=True) as (cid, _):
            bx = cid // n_num
            by = cid % n_num

            ub_input = T.alloc_ub((block_M, block_N), dtype)
            ub_output = T.alloc_ub((block_M, block_N), dtype)

            T.copy(Input[bx * block_M, by * block_N], ub_input)
            T.vLn(ub_input, ub_output)
            T.copy(ub_output, Output[bx * block_M, by * block_N])

    return main
```

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

### Matrix Multiplication Kernel with alloc_L1

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/内存操作/T.alloc_L1.md

An example demonstrating the use of T.alloc_L1 within a matrix multiplication kernel. This snippet shows how L1 buffers are allocated for matrices A, B, and C, and how they are used in conjunction with other NPUIR operations.

```python
@tilelang.jit(target="npuir")
def matmul(M, N, K, block_M, block_N, K_L1, dtype="float16", accum_dtype="float32"):
    m_num = M // block_M
    n_num = N // block_N

    @T.prim_func
    def main(
        A: T.Tensor((M, K), dtype),
        B: T.Tensor((K, N), dtype),
        C: T.Tensor((M, N), dtype),
    ):
        with T.Kernel(m_num * n_num, is_npu=True) as (cid, _), T.Scope("Cube"):
            bx = cid // n_num * block_M
            by = cid % n_num * block_N
            A_BUF = T.alloc_L1([block_M, K_L1], dtype)
            B_BUF = T.alloc_L1([K_L1, block_N], dtype)
            C_BUF = T.alloc_L0C([block_M, block_N], accum_dtype)

            for i in T.serial(T.ceildiv(K, K_L1)):
                T.npuir_load_nd2nz(A[bx, i * K_L1], A_BUF, [block_M, K_L1])
                T.npuir_load_nd2nz(B[i * K_L1, by], B_BUF, [K_L1, block_N])

                if i == 0:
                    T.npuir_dot(
                        A_BUF,
                        B_BUF,
                        C_BUF,
                        initC=True,
                        size=[block_M, K_L1, block_N],
                    )
                else:
                    T.npuir_dot(
                        A_BUF,
                        B_BUF,
                        C_BUF,
                        initC=False,
                        size=[block_M, K_L1, block_N],
                    )

                T.npuir_store_fixpipe(
                    C_BUF, C[bx, by], size=[block_M, block_N], enable_nz2nd=True
                )

    return main
```

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

### Element-wise Addition of Tensors with Broadcasting

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/数学操作/T.vadd.md

This example demonstrates element-wise addition of two tensors of shape (M, N) using T.vadd. It includes copying data to shared memory buffers and back to the destination tensor.

```python
T.vadd(src0, src1, dst)
```

```python
@tilelang.jit(target="npuir")
def sub_kernel(M, N, dtype):
    BLOCK_SIZE = 1

    @T.prim_func
    def main(
        src0: T.Tensor((M, N), dtype),
        src1: T.Tensor((M, N), dtype),
        dst: T.Tensor((M, N), dtype),
    ):

        with T.Kernel(BLOCK_SIZE, is_npu=True) as (cid, _):
            src0_ub = T.alloc_shared((M, N), dtype)
            src1_ub = T.alloc_shared((M, N), dtype)
            dst_ub = T.alloc_shared((M, N), dtype)

            T.copy(src0, src0_ub)
            T.copy(src1, src1_ub)
            T.vadd(src0_ub, src1_ub, dst_ub)
            T.copy(dst_ub, dst)

    return main
```

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

### Vector Subtraction Kernel using T.alloc_ub

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/内存操作/T.alloc_ub.md

Example demonstrating the use of `T.alloc_ub` to allocate Ascend UB memory for a vector subtraction operation. This snippet shows how to define a Tilelang JIT function for NPU execution, including buffer allocation and data copying.

```python
@tilelang.jit(target="npuir")
def vecsub(M, N, block_M, block_N, dtype="float16"):
    @T.prim_func
    def vecsub_(
        A: T.Tensor((M, N), dtype),
        B: T.Tensor((M, N), dtype),
        C: T.Tensor((M, N), dtype),
    ):
        with T.Kernel(T.ceildiv(N, block_N) * T.ceildiv(M, block_M), is_npu=True) as (
            cid,
            _,
        ):
            by = cid // T.ceildiv(N, block_N)
            bx = cid % T.ceildiv(N, block_N)

            A_BUF = T.alloc_ub((block_M, block_N), dtype)
            B_BUF = T.alloc_ub((block_M, block_N), dtype)
            C_BUF = T.alloc_ub((block_M, block_N), dtype)

            T.copy(A[by * block_M, bx * block_N], A_BUF)
            T.copy(B[by * block_M, bx * block_N], B_BUF)
            T.vsub(A_BUF, B_BUF, C_BUF)
            T.copy(C_BUF, C[by * block_M, bx * block_N])

    return vecsub_
```

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

### Matrix Multiplication with store_fixpipe

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/Tilelang.language/内存操作/T.store_fixpipe.md

This example demonstrates using T.store_fixpipe to move data from L0C to GM after a matrix multiplication. It includes loading data into buffers, performing GEMM, and then storing the result using store_fixpipe with nz2nd enabled.

```python
@tilelang.jit(target="npuir")
def kernel_mha_qk_matmul(b, n, s, d, block_d, dtype="float16", accum_dtype="float32"):

    @T.prim_func
    def main(
        Q: T.Tensor((b, n, s, d), dtype),
        K: T.Tensor((b, n, s, d), dtype),
        A: T.Tensor((b, n, s, s), accum_dtype),
    ):
        with T.Kernel(b * n, is_npu=True) as (cid, _), T.Scope("Cube"):
            b_id = cid // n
            n_id = cid % n

            Q_BUF = T.alloc_L1([s, block_d], dtype)
            K_BUF = T.alloc_L1([s, block_d], dtype)
            A_BUF = T.alloc_L0C([s, s], accum_dtype)

            for i in T.serial(T.ceildiv(d, block_d)):
                real_block_d = d - i * block_d
                real_block_d = T.min(real_block_d, block_d)
                T.load_nd2nz(Q[b_id, n_id, 0, i * block_d], Q_BUF, [s, real_block_d])
                T.load_nd2nz(K[b_id, n_id, 0, i * block_d], K_BUF, [s, real_block_d])

                T.gemm(
                    Q_BUF,
                    K_BUF,
                    A_BUF,
                    initC=(i == 0),
                    b_transpose=True,
                    size=[s, real_block_d, s],
                )

            T.store_fixpipe(A_BUF, A[b_id, n_id, 0, 0], size=[s, s], enable_nz2nd=True)

    return main

```

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

### Basic Usage of tilelang.jit.compile for NPU

Source: https://github.com/happyhappyhappy-arch/tilelang-ascend/blob/npuir/docs/developer/npu runtime.md

Demonstrates the basic usage of tilelang.jit.compile, emphasizing the requirement to specify target="npuir".

```python
# Basic Usage - The target="npuir" must be specified. 
 kernel = tilelang.jit.compile(func=my_func, target="npuir")
```