### MoE Gather/Scatter Usage Example

Source: https://context7.com/ascend/triton-ascend-ops/llms.txt

Demonstrates the usage of the gather and scatter functions for a Mixture-of-Experts (MoE) scenario. This example sets up input tensors and calls the gather and scatter operations.

```python
# Usage example for MoE
sl, hs, ne, top_k = 1024, 1536, 64, 4
device = "npu"

x = torch.randn((sl, hs), device=device, dtype=torch.half)
top_expert = torch.randint(0, ne, (sl * top_k,), device=device, dtype=torch.int32)
bin_ids, indices = torch.sort(top_expert)
weights = torch.rand((sl * top_k,), device=device, dtype=torch.half)

# Gather tokens for experts
gathered = gather(x, indices, top_k)

# Process with expert networks (placeholder)
expert_output = gathered  # Replace with actual expert processing

# Scatter back and aggregate
output = scatter(expert_output, indices, weights, top_k)
torch.npu.synchronize()
```

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

### Scatter and Aggregate MoE Outputs using Triton

Source: https://context7.com/ascend/triton-ascend-ops/llms.txt

Implements a scatter operation to aggregate expert outputs, with optional weighting, for MoE workloads. It calculates tiling parameters and optimizes for Ascend NPUs. Requires 'torch' and 'triton'.

```python
import torch
import triton
import triton.language as tl

def scatter(x, indices, weights, top_k):
    """Scatter and aggregate expert outputs with optional weighting"""
    tokens = indices.shape[0] // top_k
    out = torch.zeros((tokens, top_k, x.shape[1]), dtype=x.dtype, device=x.device)

    from utils import get_npu_properties
    num_core = get_npu_properties()["num_vectorcore"]
    indices_length = indices.shape[0]
    block_size = (indices_length - 1) // num_core + 1
    num_columns = x.shape[1]

    max_block_x = 6144
    block_x = (min(num_columns, max_block_x) + 15) // 16 * 16
    sub_block_size = max((80 * 1024 - block_x * 12) // (block_x * 2 + 4), 1)

    scale = weights is not None
    _scatter_kernel[(num_core,)]( 
        x, out, weights, indices, indices_length,
        block_size, sub_block_size, num_columns, block_x,
        TOP_K=top_k, SCALE=scale, multibuffer=True
    )
    return out.sum(dim=1) if top_k > 1 else out.view(tokens, x.shape[1])
```

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

### Gather Tokens for MoE using Triton

Source: https://context7.com/ascend/triton-ascend-ops/llms.txt

Implements a gather operation to collect tokens based on expert assignment indices for MoE workloads. It auto-calculates tiling parameters and utilizes UB-aware tiling for performance. Requires 'torch' and 'triton'.

```python
import torch
import triton
import triton.language as tl

def gather(x, indices, top_k):
    """Gather tokens according to expert assignment indices"""
    out = torch.empty((indices.shape[0], x.shape[1]), dtype=x.dtype, device=x.device)

    # Auto-calculate tiling parameters based on NPU properties
    from utils import get_npu_properties
    num_core = get_npu_properties()["num_vectorcore"]
    indices_length = indices.shape[0]
    block_size = (indices_length - 1) // num_core + 1
    num_columns = x.shape[1]

    # UB-aware tiling (192KB UB, 96KB per pipeline stage)
    max_block_x = 20480
    block_x = (min(num_columns, max_block_x) + 15) // 16 * 16
    sub_block_size = max((80 * 1024 - block_x * 2) // (block_x * 2 + 4), 1)

    _gather_kernel[(num_core,)]( 
        x, out, indices, indices_length, block_size,
        sub_block_size, num_columns, block_x, top_k,
        multibuffer=True
    )
    return out
```

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

### Load and Process K Buffer Data in Triton

Source: https://github.com/ascend/triton-ascend-ops/blob/main/tutorial/best_practice/002-decode_grouped_attention.zh.md

This snippet demonstrates loading data from the K buffer in Triton, incorporating masking and zero initialization. It handles high-dimensional continuous to low-dimensional discrete transformations. Dependencies include Triton's tensor operations (`tl`) and predefined constants like `BLOCK_N`, `BLOCK_DPE`, `K_Buffer`, etc.

```python
kpe = tl.zeros([BLOCK_N, BLOCK_DPE], dtype=qpe.dtype)
for i in range(start_n, min(BLOCK_N + start_n, split_kv_end)):
    ind = i - start_n
    offs_buf_kpe = (
        tl.get_element(kv_loc, (ind, ))  * stride_buf_kbs
        + cur_kv_head * stride_buf_kh
        + offs_dpe[None, :]
    )
    kpe_tmp = tl.load(K_Buffer + offs_buf_kpe, mask=(mask_dpe[None, :]), other=0.0)
    kpe = tl.insert_slice(kpe, kpe_tmp, (ind, 0), (1, BLOCK_DPE), (1, 1))
kpe = tl.trans(kpe, (1, 0))
```

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

### NPU Token Rearrangement with tl.insert_slice

Source: https://context7.com/ascend/triton-ascend-ops/llms.txt

Demonstrates NPU DSL extension `tl.insert_slice` for token rearrangement. It loads data by index and inserts slices into an output tensor within a loop, then stores the result to global memory. This kernel is designed for efficient tensor manipulation on NPUs.

```python
import torch
import triton
import triton.language as tl

@triton.jit
def npu_token_rearrangement_kernel(
    x_ptr, indices, output_ptr, n_elements,
    S: tl.constexpr, D: tl.constexpr, BLOCK_SIZE: tl.constexpr
):
    pid = tl.program_id(axis=0)
    dtype = output_ptr.type.element_ty
    out_start = pid * BLOCK_SIZE * D

    # Prepare output tensor
    output = tl.full((BLOCK_SIZE, D), 0, dtype=dtype)

    # Batch load rearrangement indices
    idx_offset = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    idx_mask = idx_offset < S
    idx = tl.load(indices + idx_offset, idx_mask)

    # Load data by index and insert into output tensor in loop
    for i in tl.range(0, BLOCK_SIZE):
        data_offset = D * tl.get_element(idx, (i,)) + tl.arange(0, D)[None, :]
        data_mask = data_offset < n_elements
        data = tl.load(x_ptr + data_offset, data_mask)
        output = tl.insert_slice(output, data, [i, D], [1, D], [1, 1])

    # Batch store to global memory
    out_offset = out_start + tl.arange(0, BLOCK_SIZE)[:, None] + tl.arange(0, D)[None, :]
    out_mask = out_offset < n_elements
    tl.store(output_ptr + out_offset, output, out_mask)

# Usage example
S = 1024  # sequence length
D = 32    # dimension
BLOCK_SIZE = 22

x = torch.rand(S, D, device="npu")
indices = torch.randperm(S).to(device="npu")
output = torch.empty_like(x)

npu_token_rearrangement_kernel[(48, 1, 1)](
    x, indices, output, x.numel(), S, D, BLOCK_SIZE=BLOCK_SIZE
)
torch.npu.synchronize()
```

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

### Prevent UB Overflow with Loop-based Loading (Python/Triton)

Source: https://context7.com/ascend/triton-ascend-ops/llms.txt

Demonstrates a Triton kernel that uses loop-based memory loading to prevent unified buffer overflow on NPU when processing large data structures. It calculates necessary offsets and utilizes a loop to load page information, storing results into the output indices. Dependencies include PyTorch and Triton.

```python
import torch
import triton
import triton.language as tl

@triton.jit
def npu_alloc_extend_kernel(
    pre_lens_ptr,
    seq_lens_ptr,
    free_page_ptr,
    out_indices,
    bs_upper: tl.constexpr,
    page_size: tl.constexpr,
    max_num_extend_tokens: tl.constexpr,
    BLOCK_SIZE: tl.constexpr = 1024,
):
    pid = tl.program_id(0)

    load_offset = tl.arange(0, bs_upper)
    seq_lens = tl.load(seq_lens_ptr + load_offset, mask=load_offset <= pid)
    pre_lens = tl.load(pre_lens_ptr + load_offset, mask=load_offset <= pid)
    extend_lens = seq_lens - pre_lens

    seq_len = tl.load(seq_lens_ptr + pid)
    pre_len = tl.load(pre_lens_ptr + pid)

    sum_extend_lens = tl.sum(extend_lens)
    output_start_loc = sum_extend_lens - (seq_len - pre_len)

    num_pages_after = (seq_lens + page_size - 1) // page_size
    num_pages_before = (pre_lens + page_size - 1) // page_size
    num_new_pages = num_pages_after - num_pages_before

    sum_num_new_pages = tl.sum(num_new_pages)
    new_page_start_loc = sum_num_new_pages - ((seq_len + page_size - 1) // page_size - 
                                                (pre_len + page_size - 1) // page_size)

    num_part2 = (seq_len // page_size * page_size -
                 (pre_len + page_size - 1) // page_size * page_size)

    # Loop-based approach to prevent UB overflow
    num_loop = tl.cdiv(max_num_extend_tokens, BLOCK_SIZE)
    blk_offset = tl.arange(0, BLOCK_SIZE)
    for i in range(num_loop):
        offset_many_page = blk_offset + i * BLOCK_SIZE
        page_start = tl.load(
            free_page_ptr + new_page_start_loc + offset_many_page // page_size,
            mask=offset_many_page < num_part2,
        )
        tl.store(
            out_indices + output_start_loc + offset_many_page,
            page_start * page_size + offset_many_page % page_size,
            mask=offset_many_page < num_part2,
        )

# Usage example
batch_size = 2
page_size = 128
max_context_len = 2000
max_free_page = 2048

free_pages = torch.arange(max_free_page, dtype=torch.int32, device="npu")
prefix_lens = torch.zeros(batch_size, dtype=torch.int32, device="npu")
seq_lens = torch.tensor([1000, 1500], dtype=torch.int32, device="npu")
extend_num_tokens = torch.sum(seq_lens).item()
out_indices = torch.empty((extend_num_tokens,), dtype=torch.int64, device="npu")

max_num_extend_tokens = 1 << (extend_num_tokens - 1).bit_length()

npu_alloc_extend_kernel[(batch_size,)יי](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](npu_alloc_extend_kernel[(batch_size,)](n
```

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

### Load Value Buffer Data in Triton

Source: https://github.com/ascend/triton-ascend-ops/blob/main/tutorial/best_practice/002-decode_grouped_attention.zh.md

This snippet demonstrates loading data from the V buffer in Triton, applying masks for valid ranges. It handles high-dimensional discrete to low-dimensional continuous transformations. It relies on Triton's tensor operations and pre-defined constants like `V_Buffer`, `offs_dv`, etc.

```python
offs_buf_v = (
    kv_loc[:, None] * stride_buf_vbs
    + cur_kv_head * stride_buf_vh
    + offs_dv[None, :]
)
v = tl.load(
    V_Buffer + offs_buf_v,
    mask=(offs_n[:, None] < split_kv_end) & (mask_dv[None, :]),
    other=0.0,
)
```

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

### Vector Addition Kernel for Ascend NPU

Source: https://context7.com/ascend/triton-ascend-ops/llms.txt

A Triton kernel for parallel vector addition optimized for Ascend NPUs. It supports configurable block sizes and handles boundary conditions with automatic masking. The function takes input tensors, an output tensor, vector length, and block size as parameters.

```python
import torch
import triton
import triton.language as tl

@triton.jit
def npu_vector_add_kernel(
    x,                          # [Tensor] input tensor (1 x col)
    y,                          # [Tensor] input tensor (1 x col)
    z,                          # [Tensor] output tensor (1 x col)
    vector_len: tl.constexpr,   # len of the vector
    BLOCK_SIZE: tl.constexpr
):
    pid = tl.program_id(axis=0)
    offset = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    len_mask = offset < vector_len
    x1 = tl.load(x + offset, mask=len_mask, other=0)
    y1 = tl.load(y + offset, mask=len_mask, other=0)
    z1 = x1 + y1
    tl.store(z + offset, z1, mask=len_mask)

# Usage example
vector_len = 16384
BLOCK_SIZE = 512
BLOCK_DIM = 32

x = torch.randint(0, 100, (1, vector_len), device="npu", dtype=torch.int32)
y = torch.randint(0, 100, (1, vector_len), device="npu", dtype=torch.int32)
z = torch.zeros((1, vector_len), device="npu", dtype=torch.int32)

npu_vector_add_kernel[(BLOCK_DIM,)](x, y, z, vector_len, BLOCK_SIZE)
torch.npu.synchronize()
print(f"Result: {z}")
```

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

### Grouped Attention Decode Stage 1 Kernel

Source: https://context7.com/ascend/triton-ascend-ops/llms.txt

An optimized grouped query attention kernel for the decode phase, supporting GQA/MQA/MLA with NPU-specific memory access patterns. It includes logic to determine block sizes based on input dimensions and configures grid and kernel launch parameters for efficient execution.

```python
import torch
import triton
import triton.language as tl

def decode_grouped_attention_stage1(
    q, k_buffer, v_buffer, att_out, att_lse,
    kv_indptr, kv_indices, num_kv_splits, max_kv_splits, sm_scale
):
    Lk = k_buffer.shape[-1]
    Lv = v_buffer.shape[-1]

    if Lk == 576:
        BLOCK_DMODEL, BLOCK_DPE = 512, 64
    elif Lk == 288:
        BLOCK_DMODEL, BLOCK_DPE = 256, 32
    else:
        BLOCK_DMODEL, BLOCK_DPE = triton.next_power_of_2(Lk), 0

    BLOCK_DV = triton.next_power_of_2(Lv)
    batch, head_num = q.shape[0], q.shape[1]
    kv_group_num = q.shape[1] // k_buffer.shape[1]

    BLOCK_N = 16
    BLOCK_H = 16
    MIN_BLOCK_KV = 32

    grid = (batch, triton.cdiv(head_num, min(BLOCK_H, kv_group_num)), max_kv_splits)

    grouped_attention_kernel_stage1[grid](
        q, k_buffer, v_buffer, sm_scale, kv_indptr, kv_indices,
        att_out, att_lse, num_kv_splits,
        q.stride(0), q.stride(1),
        k_buffer.stride(0), k_buffer.stride(1),
        v_buffer.stride(0), v_buffer.stride(1),
        att_out.stride(0), att_out.stride(1), att_out.stride(2),
        kv_group_num=kv_group_num, q_head_num=head_num,
        BLOCK_DMODEL=BLOCK_DMODEL, BLOCK_DPE=BLOCK_DPE, BLOCK_DV=BLOCK_DV,
        BLOCK_N=BLOCK_N, BLOCK_H=BLOCK_H, MIN_BLOCK_KV=MIN_BLOCK_KV,
        num_warps=4, num_stages=2, Lk=Lk, Lv=Lv
    )

# Usage example
B, S, H_Q, H_KV, D, D_V = 1, 128, 32, 1, 576, 512
seq_len = S
total_tokens = B * seq_len
sm_scale = 1.0 / (D ** 0.5)
max_kv_splits = 4

q = torch.randn(B, H_Q, D, dtype=torch.bfloat16, device="npu")
k_buffer = torch.randn(total_tokens, H_KV, D, dtype=torch.bfloat16, device="npu")
v_buffer = torch.randn(total_tokens, H_KV, D_V, dtype=torch.bfloat16, device="npu")

b_seq_len = torch.full((B,), seq_len, device="npu")
kv_indptr = torch.zeros((B + 1,), dtype=torch.int32, device="npu")
kv_indptr[1:B+1] = torch.cumsum(b_seq_len[:B], dim=0)
kv_indices = torch.arange(total_tokens, device="npu")
num_kv_splits = torch.full((B,), 4, dtype=torch.int32, device="npu")

attn_logits = torch.empty((B, H_Q, max_kv_splits, D_V), dtype=torch.float32, device="npu")
attn_lse = torch.empty((B, H_Q, max_kv_splits), dtype=torch.float32, device="npu")

decode_grouped_attention_stage1(
    q, k_buffer, v_buffer, attn_logits, attn_lse,
    kv_indptr, kv_indices, num_kv_splits, max_kv_splits, sm_scale
)
torch.npu.synchronize()
```

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

### LayerNorm Kernel with NPU Vector Comparison

Source: https://context7.com/ascend/triton-ascend-ops/llms.txt

A Triton kernel for Layer Normalization that includes an NPU-specific workaround for integer comparison operations. It addresses scenarios where Ascend hardware might fall back to scalar mode. The function computes mean, standard deviation, and normalized output for input tensors.

```python
import torch
import triton
import triton.language as tl

@triton.jit
def npu_vector_cmp_kernel(
    X,                 # [Tensor] input tensor (row x col)
    Out,               # [Tensor] output tensor (row x col)
    Mean,              # [Vector] mean tensor (row, ) of X
    Rstd,              # [Vector] std tensor (row, ) of X
    stride_x_row,      # [Scalar] stride of row of x
    stride_out_row,    # [Scalar] stride of row of out
    M,                 # [Scalar] row number
    N,                 # [Scalar] col number
    eps,               # [Scalar] epsilon to avoid division by zeros
    BLOCK_M: tl.constexpr,
    BLOCK_N: tl.constexpr
):
    group_m = tl.program_id(0)
    group_n = tl.program_id(1)
    row = group_m

    Mean = Mean + group_n * M
    Rstd = Rstd + group_n * M
    X = X + row * stride_x_row + group_n * N
    Out = Out + row * stride_out_row + group_n * N

    cols = tl.arange(0, BLOCK_N)  # cols is int64
    x = tl.load(X + cols, mask=cols < N, other=0.0).to(tl.float32)

    mean = tl.sum(x, axis=0) / N
    tl.store(Mean + row, mean)

    # Convert to float32 for comparison to avoid scalar fallback on NPU
    cols_cmp = cols.to(tl.float32)
    xbar = tl.where(cols_cmp < N, x - mean, 0.0)

    var = tl.sum(xbar * xbar, axis=0) / N
    rstd = 1 / tl.sqrt(var + eps)
    tl.store(Rstd + row, rstd)

    mask = cols < N
    out = (x - mean) * rstd
    tl.store(Out + cols, out, mask=mask)

# Usage example
batch_size = 256
feature_dim = 128
eps = 1e-6

X = torch.rand(batch_size, feature_dim, device="npu", dtype=torch.float32)
Out = torch.empty_like(X)
Mean = torch.empty(batch_size, device="npu", dtype=torch.float32)
Rstd = torch.empty(batch_size, device="npu", dtype=torch.float32)

BLK_M = 1
BLK_N = triton.next_power_of_2(feature_dim)
num_warps = min(max(BLK_N // 256, 1), 8)

npu_vector_cmp_kernel[(batch_size // BLK_M, 1)](
    X, Out, Mean, Rstd, X.stride(0), Out.stride(0),
    batch_size, feature_dim, eps, BLK_M, BLK_N, num_warps=num_warps)
torch.npu.synchronize()
```

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

### Compute Attention Scores with Triton

Source: https://github.com/ascend/triton-ascend-ops/blob/main/tutorial/best_practice/002-decode_grouped_attention.zh.md

This code calculates scaled attention scores (qk) using a dot product between query (qpe) and key (kpe) projections in Triton. It applies scaling and masking to ensure attention is only computed within valid ranges. The result is stored back into `qk`.

```python
qk += tl.dot(qpe, kpe.to(qpe.dtype))
qk *= sm_scale

qk = tl.where(
    mask_h[:, None] & (offs_n[None, :] < split_kv_end), qk, float("-inf")
)
```