### Run the vector-add example

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/installation_guide.md

This snippet demonstrates setting CANN environment variables and running the vector-add tutorial.

```bash
# Set the CANN environment variables (for example, as the root user and with the default installation path /usr/local/Ascend).
source /usr/local/Ascend/ascend-toolkit/set_env.sh
# Run the tutorials example.
python3 ./third_party/ascend/tutorials/01-vector-add.py
```

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

### Install System Library Dependencies (Ubuntu Example)

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/installation_guide.md

Installs essential system libraries like zlib1g-dev, Clang, and LLD required for building Triton-Ascend from source on Ubuntu. Includes an optional step to install ccache for build acceleration.

```bash
sudo apt update
sudo apt install zlib1g-dev clang-15 lld-15
sudo apt install ccache # optional
```

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

### Install Clang, LLD, and ccache

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/installation_guide.md

Installs necessary build tools for LLVM.

```bash
apt-get install -y clang-15 lld-15 ccache
```

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

### Install System Library Dependencies (Yum Example)

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/installation_guide.md

Installs the zlib-devel package using yum, which is a dependency for Triton-Ascend, particularly when using yum-based package managers.

```bash
sudo yum install -y zlib-devel
```

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

### Install torch_npu

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/installation_guide.md

This command installs the specified version of torch_npu.

```bash
pip install torch_npu==2.7.1
```

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

### Install runtime dependencies

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/installation_guide.md

This snippet shows how to clone the triton-ascend repository and install development requirements.

```bash
# Pull the triton-ascend source code repository and examples (optional; required to pull the source code repository when running examples without source code compilation and installation).
git clone https://gitcode.com/Ascend/triton-ascend.git
cd triton-ascend && pip install -r requirements_dev.txt
```

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

### Run Triton-Ascend Docker Container

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/installation_guide.md

Starts a Docker container from the triton-ascend-image with specified configurations and mounts.

```bash
docker run -u 0 -dit --shm-size=512g --name=triton-ascend_container --net=host --privileged \
--security-opt seccomp=unconfined \
--device=/dev/davinci0 \
--device=/dev/davinci1 \
--device=/dev/davinci2 \
--device=/dev/davinci3 \
--device=/dev/davinci4 \
--device=/dev/davinci5 \
--device=/dev/davinci6 \
--device=/dev/davinci7 \
--device=/dev/davinci_manager \
--device=/dev/devmm_svm \
--device=/dev/hisi_hdc \
-v /usr/local/dcmi:/usr/local/dcmi \
-v /usr/local/bin/npu-smi:/usr/local/bin/npu-smi \
-v /usr/local/sbin/npu-smi:/usr/local/sbin/npu-smi \
-v /usr/local/Ascend/driver:/usr/local/Ascend/driver \
-v /home:/home \
-v /etc/ascend_install.info:/etc/ascend_install.info 
triton-ascend-image:latest \
/bin/bash

# Enter the container
docker exec -u root -it triton-ascend_container /bin/bash
```

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

### Quick Installation from Source

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/installation_guide.md

Clones the Triton-Ascend repository, checks out the main branch, optionally sets the LLVM system path, and installs the package using pip.

```bash
git clone https://gitcode.com/Ascend/triton-ascend.git
cd triton-ascend
git checkout main

# Optional: If a pre-compiled LLVM is available locally, you can specify the path to avoid downloading the pre-built LLVM package.
# Skip this command if no local LLVM exists and execute the installation command directly.
export LLVM_SYSPATH=/path/to/LLVM

# Run the installation command
pip install -e python
```

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

### Build Triton-Ascend from Source

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/installation_guide.md

Installs Triton-Ascend using setup.py with specific build configurations.

```bash
LLVM_SYSPATH=${LLVM_INSTALL_PREFIX} \
TRITON_BUILD_WITH_CCACHE=true \
TRITON_BUILD_WITH_CLANG_LLD=true \
TRITON_BUILD_PROTON=OFF \
TRITON_WHEEL_NAME="triton-ascend" \
TRITON_APPEND_CMAKE_ARGS="-DTRITON_BUILD_UT=OFF" \
python3 setup.py install
```

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

### LLVM Code Preparation for Manual Installation

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/installation_guide.md

Prepares the LLVM source code for manual installation by cloning the repository, checking out a specific commit, downloading a patch, and applying it.

```bash
git clone --no-checkout https://github.com/llvm/llvm-project.git
cd llvm-project
git checkout fad3272286528b8a491085183434c5ad4b59ab92
wget https://raw.gitcode.com/Ascend/triton-ascend/blobs/2b0a06eb21438359d6d0576b622e3bb5e0292d17/fad3272.patch
git apply fad3272.patch
```

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

### Build and Install LLVM

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/installation_guide.md

Configures and builds LLVM using CMake, Ninja, and specified compilers and targets.

```bash
cd $HOME/llvm-project # Path to the LLVM code pulled by git clone
mkdir build
cd build
cmake ../llvm \
  -G Ninja \
  -DCMAKE_C_COMPILER=/usr/bin/clang-15 \
  -DCMAKE_CXX_COMPILER=/usr/bin/clang++-15 \
  -DCMAKE_LINKER=/usr/bin/lld-15 \
  -DCMAKE_BUILD_TYPE=Release \
  -DLLVM_ENABLE_ASSERTIONS=ON \
  -DLLVM_ENABLE_PROJECTS="mlir;llvm;lld" \
  -DLLVM_TARGETS_TO_BUILD="host;NVPTX;AMDGPU" \
  -DLLVM_ENABLE_LLD=ON \
  -DCMAKE_INSTALL_PREFIX=${LLVM_INSTALL_PREFIX}
ninja install
```

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

### Set LLVM Installation Prefix

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/installation_guide.md

Sets the environment variable for the LLVM installation path.

```bash
export LLVM_INSTALL_PREFIX=/path/to/llvm-install
```

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

### Install requirements.txt

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/quick_start.md

Command to install Python dependencies from requirements.txt and requirements_dev.txt.

```shell
pip install -r requirements.txt -r requirements_dev.txt
```

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

### Copy FileCheck

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/installation_guide.md

Copies the FileCheck executable to the LLVM installation's bin directory.

```bash
cp  {PATH_TO}/llvm_project/build/bin/FileCheck ${LLVM_INSTALL_PREFIX}/bin/FileCheck
```

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

### Manually install Torch and torch_npu

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/installation_guide.md

This command is used when the direct installation of torch_npu fails, requiring a manual installation of Torch with CPU support first.

```bash
pip install torch==2.7.1+cpu --index-url https://download.pytorch.org/whl/cpu
```

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

### Install Historical Stable Version

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/installation_guide.md

Installs a specific historical stable version of Triton-Ascend using pip.

```shell
pip install triton-ascend==3.2.0
```

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

### Run Triton Example

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/quick_start.md

Sets up CANN environment variables and runs a sample Triton Ascend tutorial (vector-add.py).

```bash
# Set CANN environment variables (using the default root installation path `/usr/local/Ascend` as an example)
source /usr/local/Ascend/ascend-toolkit/set_env.sh
# Clone the triton-ascend repository and examples (optional; required for running examples if not installed from source)
git clone https://gitcode.com/Ascend/triton-ascend.git
# Run the tutorials example:
python3 ./triton-ascend/third_party/ascend/tutorials/01-vector-add.py
```

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

### Clone Triton-Ascend Repository

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/installation_guide.md

Clones the Triton-Ascend Git repository.

```bash
git clone https://gitcode.com/Ascend/triton-ascend.git && cd triton-ascend
```

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

### Run Example with Debugging Flags

Source: https://github.com/ascend/triton-ascend/blob/main/docs/zh/debug_guide/debugging.md

Command to run a demonstration test case with debugging flags enabled.

```bash
TRITON_DEBUG=1 TRITON_DISABLE_CACHE=1 python 01-vector-add.py
```

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

### SIMT Compilation Mode Example

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/libdevice/simt/libdevice_simt_developer_guide.md

Triton kernel example demonstrating the use of libdevice SIMT compilation mode.

```python
# Enable libdevice SIMT compilation
import os
os.environ['TRITON_ENABLE_LIBDEVICE_SIMT'] = '1'

import triton
import triton.language as tl
import triton.language.extra.cann.libdevice as libdevice
import torch

@triton.jit
def triton_kernel(input, output, XBLOCK: tl.constexpr, XBLOCK_SUB: tl.constexpr):
    offset = tl.program_id(0) * XBLOCK
    base = tl.arange(0, XBLOCK_SUB)
    loops: tl.constexpr = XBLOCK // XBLOCK_SUB
    for loop in range(loops):
        x0 = offset + (loop * XBLOCK_SUB) + base
        x = tl.load(input + (x0), None)
        y = libdevice.abs(x)
        tl.store(output + (x0), y, None)

dtype, shape, ncore, xblock, xblock_sub = ['int32', (128, 4096), 512, 1024, 1024]
input = torch.randn(shape, dtype=dtype).npu()
output = torch.randn(shape, dtype=dtype).npu()
triton_kernel[ncore, 1, 1](input, output, xblock, xblock_sub, force_simt_only=True)
```

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

### Install Latest Stable Version

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/installation_guide.md

Installs the latest stable version of Triton-Ascend using pip, including a specific extra index URL.

```shell
pip install triton-ascend==3.2.1 --extra-index-url=https://triton-ascend.osinfra.cn/pypi/simple
```

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

### Target-Specific Adapter Representation (TTAdapter IR) Example

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/debug_guide/debugging.md

An example of the kernel.ttadapter.mlir file, showcasing the TTAdapter IR.

```mlir
module {
  func.func @add_kernel(%arg0: memref<?xi8>, %arg1: memref<?xi8>, %arg2: memref<?xf32> {tt.divisibility = 16 : i32, tt.tensor_kind = 0 : i32}, %arg3: memref<?xf32> {tt.divisibility = 16 : i32, tt.tensor_kind = 0 : i32}, %arg4: memref<?xf32> {tt.divisibility = 16 : i32, tt.tensor_kind = 1 : i32}, %arg5: i32 {tt.divisibility = 16 : i32}, %arg6: i32, %arg7: i32, %arg8: i32, %arg9: i32, %arg10: i32, %arg11: i32) attributes {SyncBlockLockArgIdx = 0 : i64, WorkspaceArgIdx = 1 : i64, global_kernel = "local", mix_mode = "aiv", parallel_mode = "simd"} {
    %cst = arith.constant 0.000000e+00 : f32
    %c1024 = arith.constant 1024 : index
    %c1024_i32 = arith.constant 1024 : i32
    %0 = arith.muli %arg9, %c1024_i32 : i32
    %1 = arith.index_cast %0 : i32 to index
    %reinterpret_cast = memref.reinterpret_cast %arg2 to offset: [%1], sizes: [1024], strides: [1] : memref<?xf32> to memref<1024xf32, strided<[1], offset: ?>>
    %alloc = memref.alloc() : memref<1024xf32>
    %2 = arith.addi %1, %c1024 : index
    %3 = arith.index_cast %arg5 : i32 to index
    %4 = arith.maxsi %1, %3 : index
    %5 = arith.minsi %2, %4 : index
    %6 = arith.subi %5, %1 : index
    %7 = arith.cmpi slt, %6, %c1024 : index
    scf.if %7 {
      linalg.fill ins(%cst : f32) outs(%alloc : memref<1024xf32>)
    } {hivm.unlikely_condition}
    %subview = memref.subview %reinterpret_cast[0] [%6] [1] : memref<1024xf32, strided<[1], offset: ?>> to memref<?xf32, strided<[1], offset: ?>>
    %subview_0 = memref.subview %alloc[0] [%6] [1] : memref<1024xf32> to memref<?xf32, strided<[1]>>
    memref.copy %subview, %subview_0 : memref<?xf32, strided<[1], offset: ?>> to memref<?xf32, strided<[1]>>
    %8 = bufferization.to_tensor %alloc restrict writable : memref<1024xf32>
    %reinterpret_cast_1 = memref.reinterpret_cast %arg3 to offset: [%1], sizes: [1024], strides: [1] : memref<?xf32> to memref<1024xf32, strided<[1], offset: ?>>
    %alloc_2 = memref.alloc() : memref<1024xf32>
    scf.if %7 {
      linalg.fill ins(%cst : f32) outs(%alloc_2 : memref<1024xf32>)
    } {hivm.unlikely_condition}
    %subview_3 = memref.subview %reinterpret_cast_1[0] [%6] [1] : memref<1024xf32, strided<[1], offset: ?>> to memref<?xf32, strided<[1], offset: ?>>
    %subview_4 = memref.subview %alloc_2[0] [%6] [1] : memref<1024xf32> to memref<?xf32, strided<[1]>>
    memref.copy %subview_3, %subview_4 : memref<?xf32, strided<[1], offset: ?>> to memref<?xf32, strided<[1]>>
    %9 = bufferization.to_tensor %alloc_2 restrict writable : memref<1024xf32>
    %10 = arith.addf %8, %9 : tensor<1024xf32>
    %reinterpret_cast_5 = memref.reinterpret_cast %arg4 to offset: [%1], sizes: [1024], strides: [1] : memref<?xf32> to memref<1024xf32, strided<[1], offset: ?>>
    %extracted_slice = tensor.extract_slice %10[0] [%6] [1] : tensor<1024xf32> to tensor<?xf32>
    %subview_6 = memref.subview %reinterpret_cast_5[0] [%6] [1] : memref<1024xf32, strided<[1], offset: ?>> to memref<?xf32, strided<[1], offset: ?>>
    bufferization.materialize_in_destination %extracted_slice in writable %subview_6 : (tensor<?xf32>, memref<?xf32, strided<[1], offset: ?>>) -> ()
    return
  }
}
```

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

### Handwritten triton.Config Example

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/autotune_guide.md

Example demonstrating the standard community-style handwritten configuration path using triton.Config objects.

```python
import triton

@triton.autotune(
    configs=[
        triton.Config({"BLOCK_M": 128, "BLOCK_N": 128}),
        triton.Config({"BLOCK_M": 64, "BLOCK_N": 256}),
    ],
    key=["M", "N"],
)
@triton.jit
def kernel(...):
    ...
```

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

### Simplified Example for PT File Generation

Source: https://github.com/ascend/triton-ascend/blob/main/third_party/ascend/unittest/kernels/README.md

This example demonstrates the process of generating a PT file, which includes constructing GPU input, running the GPU kernel to get output, and saving the input, grid, and output.

```python
import copy
import torch

DEVICE = torch.device("cuda:0")
batch_size = 2
grid = (batch_size,)

input_data = {
"output_token_ids_ptr": torch.zeros((batch_size, 4), dtype=torch.int32, device=DEVICE),
"cu_num_draft_tokens_ptr": torch.tensor([2, 1], dtype=torch.int32, device=DEVICE),
# ... 其它字段
}

# 保存输入副本到 CPU
input_data_before = {
 k: (v.clone().cpu() if isinstance(v, torch.Tensor) else copy.deepcopy(v))
 for k, v in input_data.items()
}
# 预处理 input_data_before 符合 NPU kernel 输入
input_data_before["npu_need_param_key"] = NPU_NEED_PARAMS_VALUE
# 运行 kernel（在 GPU / 参考实现上）并收集输出
triton_kernel[grid](**input_data)
# 这里用 input_data 作为示例，实际应调用对应的 triton/pytorch 函数
gpu_output = {k: (v.cpu() if isinstance(v, torch.Tensor) else v) for k, v in input_data.items()}

save_obj = {"input_data": input_data_before, "grid": grid, "gpu_output": gpu_output}
torch.save(save_obj, "<kernel_name>.pt")
# 多组用例场景：torch.save([save_obj1, save_obj2], "<kernel_name>.pt")
```

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

### Build Triton-Ascend using Docker

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/installation_guide.md

Builds a Docker image for Triton-Ascend, specifying chip type and CANN version.

```bash
git clone https://gitcode.com/Ascend/triton-ascend.git && cd triton-ascend
docker build \
--build-arg CHIP_TYPE=A3 \
--build-arg CANN_VERSION=9.0.0 \
-t triton-ascend-image:latest -f ./docker/Dockerfile .
```

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

### Add Kernel Example

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/programming_guide/index.md

An example of a Triton-Ascend kernel for element-wise addition, demonstrating the use of BLOCK_SIZE for managing on-chip memory usage.

```python
import triton.language as tl

@triton.jit
def add_kernel(x_ptr, 
               y_ptr, 
               out_ptr, 
               n,  # Total number of elements. 
               BLOCK_SIZE: tl.constexpr,  # Number of block elements. 
               ): 
    pid = tl.program_id(0) 
    NUM_CORE = tl.num_programs(0)
    NUM_BLOCKS = tl.cdiv(n, BLOCK_SIZE)
    for block_idx in range(pid, NUM_BLOCKS, NUM_CORE):
        block_start = block_idx * BLOCK_SIZE
        # The block size is BLOCK_SIZE.
        offsets = block_start + tl.arange(0, BLOCK_SIZE)
        mask = offsets < n
        # Load data of x and y to the on-chip memory.
        x = tl.load(x_ptr + offsets, mask=mask)
        y = tl.load(y_ptr + offsets, mask=mask)

        output = x + y

        tl.store(out_ptr + offsets, output, mask=mask)
```

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

### Install Python Dependencies

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/installation_guide.md

Installs necessary Python build-time dependencies for Triton-Ascend, including ninja, cmake, wheel, and pybind11.

```bash
pip install ninja cmake wheel pybind11 # build-time dependencies
```

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

### Run Docker Container

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/quick_start.md

Starts a Docker container from the built image with necessary device and volume mounts for Ascend hardware.

```bash
docker run -u 0 -dit --shm-size=512g --name=triton-ascend_container --net=host --privileged \
--security-opt seccomp=unconfined \
--device=/dev/davinci0 \
--device=/dev/davinci1 \
--device=/dev/davinci2 \
--device=/dev/davinci3 \
--device=/dev/davinci4 \
--device=/dev/davinci5 \
--device=/dev/davinci6 \
--device=/dev/davinci7 \
--device=/dev/davinci_manager \
--device=/dev/devmm_svm \
--device=/dev/hisi_hdc \
-v /usr/local/dcmi:/usr/local/dcmi \
-v /usr/local/bin/npu-smi:/usr/local/bin/npu-smi \
-v /usr/local/sbin/npu-smi:/usr/local/sbin/npu-smi \
-v /usr/local/Ascend/driver:/usr/local/Ascend/driver \
-v /home:/home \
-v /etc/ascend_install.info:/etc/ascend_install.info \
triton-ascend-image:latest \
/bin/bash

# Enter the container
docker exec -u root -it triton-ascend_container /bin/bash
```

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

### Pdb Example

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/debug_guide/debugging.md

Example of using the Pdb debugger to print intermediate variables.

```python
(Pdb) p tmp0  # Print the value of variable tmp0.
```

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

### Install Tutorial Dependencies

Source: https://github.com/ascend/triton-ascend/blob/main/python/tutorials/README.rst

Command to install the necessary dependencies for the Triton tutorials.

```bash
cd triton
pip install -e './python[tutorials]'
```

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

### Runtime Debugging Example

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/debug_guide/debugging.md

Example of using tl.device_print to print tensor values during kernel execution.

```python
import triton.language as tl

@triton.jit
def triton_kernel(out_ptr0, in_ptr0, in_ptr1, XBLOCK: tl.constexpr):
    idx = tl.arange(0, XBLOCK)
    tmp0 = tl.load(in_ptr0 + idx)
    tmp1 = tl.load(in_ptr1 + idx)
    tmp2 = tmp0 + tmp1
    tl.device_print("tmp2 after addition = ", tmp2)  # Print the intermediate result.
    tl.store(out_ptr0 + idx, tmp2)
```

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

### Example Usage

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/triton_api_extention/al/al.copy.md

Example demonstrating how to compile and use the copy kernel.

```python
import os
import triton
import triton.language as tl
import triton.extension.buffer.language as bl
import triton.language.extra.cann.extension as al
from triton.compiler.compiler import ASTSource
from triton.compiler.code_generator import ast_to_ttir
from triton._C.libtriton import ir, buffer_ir
from triton._C.libtriton.ascend import ir as ascend_ir

os.environ["TORCH_DEVICE_BACKEND_AUTOLOAD"] = "0"


class Options:
    num_warps = 4
    num_stages = 3
    num_ctas = 1
    cluster_dims = (1, 1, 1)
    enable_fp_fusion = True
    debug = False
    arch = "Ascend910_95"


def compile_kernel(kernel, signature, constants):
    """Helper to compile a kernel to MLIR."""
    src = ASTSource(kernel, signature, constants)
    context = ir.context()
    ir.load_dialects(context)
    buffer_ir.load_dialects(context)
    ascend_ir.load_dialects(context)
    module = ast_to_ttir(kernel, src, context, Options(), {}, {})
    return str(module)


@triton.jit
def copy(
    A_ptr,
    A1_ptr,
    M: tl.constexpr,
    N: tl.constexpr,
):
    offs_a = tl.arange(0, M)[:, None]
    offs_b = tl.arange(0, N)[None, :]

    offs_c = (offs_a) * M + (offs_b)
    a_ptr = A_ptr + offs_c
    a_val = tl.load(a_ptr)
    a1_ptr = A1_ptr + offs_c
    a1_val = tl.load(a1_ptr)

    add = tl.add(a_val, a1_val)
    add_ub = bl.to_buffer(add, al.ascend_address_space.UB)

    A_l1 = bl.alloc(tl.float32, [M, N], al.ascend_address_space.L1)
    al.copy_from_ub_to_l1(add_ub, A_l1)

    A_ub = bl.alloc(tl.float32, [M, N], al.ascend_address_space.UB)
    al.copy(add_ub, A_ub)


def test_copy():
    print("=" * 60)
    print("Test 1: copy ")
    print("=" * 60)
    mlir = compile_kernel(
        copy,
        {"A_ptr": "*fp32", "A1_ptr": "*fp32"},
        {"M": 16, "N": 16},
    )
    print(f"Generated MLIR ({len(mlir)} chars):\n")
    print(mlir)


if __name__ == "__main__":
    test_copy()

```

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

### Static Printing Debugging Example

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/debug_guide/debugging.md

Example of using tl.static_print to debug constant parameters during compilation.

```python
import triton.language as tl

@triton.jit
def triton_kernel(
    out_ptr0, 
    in_ptr0, 
    in_ptr1, 
    XBLOCK: tl.constexpr,  # Constant parameter during compilation
    USE_FP16: tl.constexpr  # Constant parameter during compilation
):
    # Print constant parameters during compilation.
    tl.static_print("XBLOCK = ", XBLOCK)
    tl.static_print("USE_FP16 = ", USE_FP16) 
    
    idx = tl.arange(0, XBLOCK)
    tmp0 = tl.load(in_ptr0 + idx)
    tmp1 = tl.load(in_ptr1 + idx)
    
    # Print the constant calculation result.
    elements_per_thread = XBLOCK // 32
    tl.static_print("Elements per thread = ", elements_per_thread)
    
    tmp2 = tmp0 + tmp1
    tl.store(out_ptr0 + idx, tmp2)
```

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

### Autotune Configuration Example

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/migration_guide/architecture_difference.md

Example of how to enable the 'multibuffer' option during the autotune configuration phase by passing it to triton.Config.

```python
    def get_autotune_config():
        return [
            triton.Config({'XS': 1 * 128, 'multibuffer': True}),]

```

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

### Example Usage

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/triton_api_extention/bl/alloc.md

Demonstrates how to use the bl.alloc API to allocate buffers with different data types, shapes, and address spaces on the Ascend platform. It includes kernel definitions and a main block for testing and compiling the kernel to MLIR.

```Python
import os
import triton
import triton.language as tl
from triton.compiler.compiler import ASTSource
from triton.compiler.code_generator import ast_to_ttir
import triton.extension.buffer.language as bl
import triton.language.extra.cann.extension as al
from triton._C.libtriton import ir, buffer_ir
from triton._C.libtriton.ascend import ir as ascend_ir


os.environ["TORCH_DEVICE_BACKEND_AUTOLOAD"] = "0"


class Options:
    num_warps = 4
    num_stages = 3
    num_ctas = 1
    cluster_dims = (1, 1, 1)
    enable_fp_fusion = True
    debug = False


def compile_kernel(kernel, signature, constants):
    """Helper to compile a kernel to MLIR."""
    src = ASTSource(kernel, signature, constants)
    context = ir.context()
    ir.load_dialects(context)
    buffer_ir.load_dialects(context)
    ascend_ir.load_dialects(context)
    module = ast_to_ttir(kernel, src, context, Options(), {"create_address_space": al.semantic.create_address_space}, {})
    return str(module)


# ============== Kernel definitions ==============


@triton.jit
def allocate_local_buffer(XBLOCK: tl.constexpr):
    # this statement has no effect, just to test the builder
    bl.alloc(tl.float32, [XBLOCK])
    bl.alloc(tl.float32, [XBLOCK, XBLOCK], al.ascend_address_space.UB)
    bl.alloc(tl.float32, [XBLOCK, XBLOCK], al.ascend_address_space.L1)
    bl.alloc(tl.float32, [XBLOCK, XBLOCK], al.ascend_address_space.L0A)
    bl.alloc(tl.float32, [XBLOCK, XBLOCK], al.ascend_address_space.L0B)
    bl.alloc(tl.float32, [XBLOCK, XBLOCK], al.ascend_address_space.L0C)
    bl.alloc(
        tl.float32, [XBLOCK, XBLOCK], al.ascend_address_space.UB, is_mem_unique=True
    )


# ============== Main for manual testing ==============

if __name__ == "__main__":
    print("=" * 60)
    print("Test 1: Nested Scopes")
    print("=" * 60)
    mlir = compile_kernel(
        allocate_local_buffer, {}, {"XBLOCK": 256}
    )
    print(f"✅ Generated MLIR ({len(mlir)} chars):\n")
    print(mlir)
```

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

### TTIR Example

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/debug_guide/debugging.md

An example of a Triton Intermediate Representation (TTIR) file in MLIR format, representing an add kernel.

```python
module {
  tt.func public @add_kernel(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32} , %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32} , %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32} , %arg3: i32 {tt.divisibility = 16 : i32} ) attributes {noinline = false} {
    %cst = arith.constant dense<0.000000e+00> : tensor<1024xf32> loc(#loc1)
    %c1024_i32 = arith.constant 1024 : i32 loc(#loc1)
    %0 = tt.get_program_id x : i32 loc(#loc2)
    %1 = arith.muli %0, %c1024_i32 : i32 loc(#loc3)
    %2 = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32> loc(#loc4)
    %3 = tt.splat %1 : i32 -> tensor<1024xi32> loc(#loc5)
    %4 = arith.addi %3, %2 : tensor<1024xi32> loc(#loc5)
    %5 = tt.splat %arg3 : i32 -> tensor<1024xi32> loc(#loc6)
    %6 = arith.cmpi slt, %4, %5 : tensor<1024xi32> loc(#loc6)
    %7 = tt.splat %arg0 : !tt.ptr<f32> -> tensor<1024x!tt.ptr<f32>> loc(#loc7)
    %8 = tt.addptr %7, %4 : tensor<1024x!tt.ptr<f32>>, tensor<1024xi32> loc(#loc7)
    %9 = tt.load %8, %6, %cst : tensor<1024x!tt.ptr<f32>> loc(#loc8)
    %10 = tt.splat %arg1 : !tt.ptr<f32> -> tensor<1024x!tt.ptr<f32>> loc(#loc9)
    %11 = tt.addptr %10, %4 : tensor<1024x!tt.ptr<f32>>, tensor<1024xi32> loc(#loc9)
    %12 = tt.load %11, %6, %cst : tensor<1024x!tt.ptr<f32>> loc(#loc10)
    %13 = arith.addf %9, %12 : tensor<1024xf32> loc(#loc11)
    %14 = tt.splat %arg2 : !tt.ptr<f32> -> tensor<1024x!tt.ptr<f32>> loc(#loc12)
    %15 = tt.addptr %14, %4 : tensor<1024x!tt.ptr<f32>>, tensor<1024xi32> loc(#loc12)
    tt.store %15, %13, %6 : tensor<1024x!tt.ptr<f32>> loc(#loc13)
    tt.return loc(#loc14)}
}

```

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

### Example Usage

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/triton_api_extention/al/scope.md

Demonstrates how to compile a kernel with and without auto sync enabled.

```python
)

    print(f"✅ Generated MLIR ({len(mlir)} chars):
")

    print(mlir)

    print("\n" + "=" * 60)

    print("Test 5: Disable Auto Sync")

    print("=" * 60)

    mlir = compile_kernel(

        kernel_scope_disable_auto_sync,

        {"x_ptr": "*fp32", "y_ptr": "*fp32", "out_ptr": "*fp32", "n": "i32"},

        {"BLOCK": 256},

    )

    print(f"✅ Generated MLIR ({len(mlir)} chars):
")

    print(mlir)
```

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

### Simple example of Triton Autotune

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/programming_guide/index.md

This example demonstrates how to use `triton.autotune` to find the optimal `BLOCK_SIZE` for an addition kernel.

```python
import triton.language as tl

@triton.autotune(
configs=[ # List of parameter configurations to be tested. The candidate parameter values must be powers of 2.
            triton.Config({'BLOCK_SIZE': 128}),
            triton.Config({'BLOCK_SIZE': 256}),
            triton.Config({'BLOCK_SIZE': 512}),
        ],
        key=['n_elements'], # Tune dimension: input dimension on which the parameter value depends.
    )
@triton.jit
def add_kernel(x_ptr, y_ptr, output_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(axis=0)
    block_start = pid * BLOCK_SIZE
    offsets = block_start + tl.arange(0, BLOCK_SIZE)
    mask = offsets < n_elements

    x = tl.load(x_ptr + offsets, mask=mask)
    y = tl.load(y_ptr + offsets, mask=mask)
    output = x + y
    tl.store(output_ptr + offsets, output, mask=mask)
```

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

### PDB interactive debugging example

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/debug_guide/debugging.md

Example of using pdb commands to inspect code context, variables, and step through execution.

```python
(Pdb) l  # View the current code context.
118     def compile_fn(ttir):
120         import pdb; pdb.set_trace()
121         # Check the input parameter.
122         print(f"ttir type: {type(ttir)}")
123         result = lower_function(ttir)  # <-- The current suspension position.

(Pdb) p ttir  # Check the input parameter.
(Pdb) n # Execute the next line of code.
(Pdb) p result  # View the result.
```

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

### Example Usage

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/triton_api_extention/al/sync_block_wait.md

Example code demonstrating sync block wait functionality.

```Python
oc6)
    } {hivm.tcore_type = #hivm.tcore_type<VECTOR>, noinline} loc(#loc4)
    tt.return loc(#loc7)
  } loc(#loc)
}
loc = loc("/home/ganpengfei/workspace/triton-test/sync_block_set_wait.py":52:0)
#loc1 = loc("/home/ganpengfei/workspace/triton-test/sync_block_set_wait.py":53:9)
#loc2 = loc("/home/ganpengfei/workspace/triton-test/sync_block_set_wait.py":54:66)
#loc3 = loc("/home/ganpengfei/workspace/triton-test/sync_block_set_wait.py":55:66)
#loc4 = loc("/home/ganpengfei/workspace/triton-test/sync_block_set_wait.py":56:9)
#loc5 = loc("/home/ganpengfei/workspace/triton-test/sync_block_set_wait.py":57:67)
#loc6 = loc("/home/ganpengfei/workspace/triton-test/sync_block_set_wait.py":58:67)
#loc7 = loc("/home/ganpengfei/workspace/triton-test/sync_block_set_wait.py":56:4)
```

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

### Input Example

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/triton_api_extention/bl/bind_buffer.md

Example demonstrating how to use the to_buffer function to bind a tensor to a buffer.

```python
import os
import triton
import triton.language as tl
import triton.extension.buffer.language as bl
import triton.language.extra.cann.extension as al
from triton.compiler.compiler import ASTSource
from triton.compiler.code_generator import ast_to_ttir
from triton._C.libtriton import ir
from triton._C.libtriton.ascend import ir as ascend_ir

os.environ["TORCH_DEVICE_BACKEND_AUTOLOAD"] = "0"

class Options:
    num_warps = 4
    num_stages = 3
    num_ctas = 1
    cluster_dims = (1, 1, 1)
    enable_fp_fusion = True
    debug = False

def compile_kernel(kernel, signature, constants):
    """Helper to compile a kernel to MLIR."""
    src = ASTSource(kernel, signature, constants)
    context = ir.context()
    ir.load_dialects(context)
    ascend_ir.load_dialects(context)
    module = ast_to_ttir(kernel, src, context, Options(), {}, {})
    return str(module)

@triton.jit
def bind_buffer():
    alloc = bl.alloc(tl.float32, [32, 32], al.ascend_address_space.UB)
    tensor = tl.full((32, 32), 0, dtype=tl.float32)
    bl.to_buffer(tensor, bind_buffer=alloc)

# ============== Main for manual testing ==============

if __name__ == "__main__":
    mlir = compile_kernel(bind_buffer, {}, {})
    assert len(mlir) > 0
    print(mlir)
```

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

### Install Proton

Source: https://github.com/ascend/triton-ascend/blob/main/third_party/proton/README.md

Clones the Triton repository, navigates to the Python directory, and installs Proton.

```bash
git clone https://github.com/triton-lang/triton
cd triton/python
pip install .
```

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

### Start Pipeline Test

Source: https://github.com/ascend/triton-ascend/blob/main/CONTRIBUTING.md

Comment '/compile' to start the pipeline test. If the test fails, modify the code and comment '/compile' again. After the test passes, the tag 'ci-pipeline-passed' is added.

```shell
/compile
```

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

### Calling the Vector Addition Kernel

Source: https://github.com/ascend/triton-ascend/blob/main/docs/en/programming_guide/index.md

Example of how to call the Triton add_kernel function.

```python
def add(x: torch.Tensor, y: torch.Tensor):
    output = torch.empty_like(x)
    n_elements = output.numel()
    grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']), )
    add_kernel[grid](x, y, output, n_elements, BLOCK_SIZE=1024)
    return output

```