add mov op

lib/TileOps/tmatmul_bias_template.py

@@ -0,0 +1,37 @@
+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+# CANN Open Software License Agreement Version 2.0 (the "License").
+# Please refer to the License for details. You may not use this file except in compliance with the License.
+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+# See LICENSE in the root of the software repository for the full text of the License.
+
+"""TileLang DSL template for pto.tmatmul.bias - Matmul with bias addition."""
+
+import tilelang_dsl as pto
+
+
+@pto.ckernel(
+    target="a5",
+    op="pto.tmatmul.bias",
+    dtypes=[
+        (pto.f16, pto.f16, pto.f32, pto.f32),  # (lhs, rhs, bias, acc) - matches IR operand order
+        (pto.bf16, pto.bf16, pto.f32, pto.f32),
+        (pto.f32, pto.f32, pto.f32, pto.f32),
+    ],
+)
+def template_tmatmul_bias(lhs: pto.Tile, rhs: pto.Tile, bias: pto.Tile, acc: pto.Tile):
+    """Matmul with bias addition.
+
+    Args:
+        lhs: Left matrix tile (L0A buffer)
+        rhs: Right matrix tile (L0B buffer)
+        bias: Bias tile (Bias Table buffer)
+        acc: Accumulator tile (L0C buffer, output)
+
+    Computes: acc = lhs @ rhs + bias
+    """
+    m, k = lhs.valid_shape
+    n, _ = rhs.valid_shape
+    pto.mad_bias(lhs.as_ptr(), rhs.as_ptr(), acc.as_ptr(), bias.as_ptr(), m, n, k, disable_gemv=True)
+    return None

lib/TileOps/tmov2bias_template.py

@@ -0,0 +1,89 @@
+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+# CANN Open Software License Agreement Version 2.0 (the "License").
+# Please refer to the License for details. You may not use this file except in compliance with the License.
+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+# See LICENSE in the root of the software repository for the full text of the License.
+
+"""TileLang DSL template for pto.tmov - Mat to Bias buffer movement.
+
+This template implements the TMOV_M2B scenario for cube kernels:
+  - Source: L1 Mat buffer (memory_space="mat", 1xN row-major layout)
+  - Destination: L0 Bias Table buffer (memory_space="bias")
+  - Uses bias_load (mte_l1_bt) intrinsic operation
+
+The Bias Table is a special 4KB buffer in L0 used for bias addition
+in matmul operations. Requirements:
+  - Row dimension must be 1
+  - Column dimension * sizeof(dtype) must be aligned to 64 bits
+  - Total size must not exceed 4KB (4096 bytes)
+
+Constraint: This template is selected when dst.memory_space == BIAS.
+"""
+
+import tilelang_dsl as pto
+
+
+def _tmov_m2b_constraint(src: pto.Tile, dst: pto.Tile) -> bool:
+    """Constraint: Mat to Bias transfer scenario.
+
+    Supported scenario:
+      - src.memory_space == MAT
+      - dst.memory_space == BIAS
+    """
+    src_ms = src.memory_space
+    dst_ms = dst.memory_space
+
+    # Check src is MAT
+    if isinstance(src_ms, str):
+        src_is_mat = src_ms == "mat"
+    elif isinstance(src_ms, pto.MemorySpace):
+        src_is_mat = src_ms == pto.MemorySpace.MAT
+    else:
+        src_is_mat = hasattr(src_ms, "value") and src_ms.value == "mat"
+
+    # Check dst is BIAS
+    if isinstance(dst_ms, str):
+        dst_is_bias = dst_ms == "bias"
+    elif isinstance(dst_ms, pto.MemorySpace):
+        dst_is_bias = dst_ms == pto.MemorySpace.BIAS
+    else:
+        dst_is_bias = hasattr(dst_ms, "value") and dst_ms.value == "bias"
+
+    return src_is_mat and dst_is_bias
+
+
+@pto.ckernel(
+    target="a5",
+    op="pto.tmov",
+    constraints=[_tmov_m2b_constraint],
+    dtypes=[
+        (pto.f32, pto.f32),
+        (pto.f16, pto.f32),
+        (pto.bf16, pto.f32),
+        (pto.i32, pto.i32),
+    ],
+)
+def template_tmov_m2b(src: pto.Tile, dst: pto.Tile):
+    """Move data from Mat buffer to Bias Table buffer.
+
+    Args:
+        src: Source tile in L1 Mat location (1xN row-major)
+        dst: Destination tile in Bias Table location
+
+    The bias data is moved using burst transfer to the Bias Table.
+    """
+    # Bias has shape 1xN, we derive N from the valid shape
+    _, n = dst.valid_shape
+    dtype = dst.element_type
+    # len_burst = ceil(N * sizeof(dtype) / 32)
+    # Use integer arithmetic: (n * dtype_size + 31) // 32
+    dtype_size = pto.bytewidth(dtype)
+    len_burst = (n * dtype_size + 31) // 32
+    # nburst = (n_burst, src_gap, dst_gap) - single burst with no gaps
+    n_burst = 1
+    src_gap = 0
+    dst_gap = 0
+    pto.mte_l1_bt(src.as_ptr(), dst.as_ptr(), len_burst, nburst=(n_burst, src_gap, dst_gap))
+    return

lib/TileOps/tmov2left_template.py

@@ -0,0 +1,78 @@
+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+# CANN Open Software License Agreement Version 2.0 (the "License").
+# Please refer to the License for details. You may not use this file except in compliance with the License.
+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+# See LICENSE in the root of the software repository for the full text of the License.
+
+"""TileLang DSL template for pto.tmov - Mat to Left buffer movement.
+
+This template implements the TMOV_M2L scenario for cube kernels:
+  - Source: L1 Mat buffer (memory_space="mat")
+  - Destination: L0A Left buffer (memory_space="left")
+  - Uses left_load (mte_l1_l0a) intrinsic operation
+
+The operation is part of the cube matmul data flow where:
+  1. Data is loaded from GM to L1 Mat via TLOAD
+  2. TMOV moves data from L1 Mat to L0A Left
+  3. The Left buffer is then used as input for TMATMUL
+
+Constraint: This template is selected when dst.memory_space == LEFT.
+"""
+
+import tilelang_dsl as pto
+
+
+def _tmov_m2l_constraint(src: pto.Tile, dst: pto.Tile) -> bool:
+    """Constraint: Mat to Left transfer scenario.
+
+    Supported scenario:
+      - src.memory_space == MAT
+      - dst.memory_space == LEFT
+    """
+    src_ms = src.memory_space
+    dst_ms = dst.memory_space
+
+    # Check src is MAT
+    if isinstance(src_ms, str):
+        src_is_mat = src_ms == "mat"
+    elif isinstance(src_ms, pto.MemorySpace):
+        src_is_mat = src_ms == pto.MemorySpace.MAT
+    else:
+        src_is_mat = hasattr(src_ms, "value") and src_ms.value == "mat"
+
+    # Check dst is LEFT
+    if isinstance(dst_ms, str):
+        dst_is_left = dst_ms == "left"
+    elif isinstance(dst_ms, pto.MemorySpace):
+        dst_is_left = dst_ms == pto.MemorySpace.LEFT
+    else:
+        dst_is_left = hasattr(dst_ms, "value") and dst_ms.value == "left"
+
+    return src_is_mat and dst_is_left
+
+
+@pto.ckernel(
+    target="a5",
+    op="pto.tmov",
+    constraints=[_tmov_m2l_constraint],
+    dtypes=[
+        (pto.f16, pto.f16),
+        (pto.bf16, pto.bf16),
+        (pto.f32, pto.f32),
+        (pto.i8, pto.i8),
+    ],
+)
+def template_tmov_m2l(src: pto.Tile, dst: pto.Tile):
+    """Move data from Mat buffer to Left buffer.
+
+    Args:
+        src: Source tile in L1 Mat location
+        dst: Destination tile in L0A Left location
+
+    The m, k dimensions are derived from the tile shapes.
+    """
+    m, k = dst.valid_shape
+    pto.mte_l1_l0a(src.as_ptr(), dst.as_ptr(), m, k)
+    return

lib/TileOps/tmov2right_template.py

@@ -0,0 +1,79 @@
+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+# CANN Open Software License Agreement Version 2.0 (the "License").
+# Please refer to the License for details. You may not use this file except in compliance with the License.
+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+# See LICENSE in the root of the software repository for the full text of the License.
+
+"""TileLang DSL template for pto.tmov - Mat to Right buffer movement.
+
+This template implements the TMOV_M2R scenario for cube kernels:
+  - Source: L1 Mat buffer (memory_space="mat")
+  - Destination: L0B Right buffer (memory_space="right")
+  - Uses right_load (mte_l1_l0b) intrinsic operation
+
+The operation is part of the cube matmul data flow where:
+  1. Data is loaded from GM to L1 Mat via TLOAD
+  2. TMOV moves data from L1 Mat to L0B Right
+  3. The Right buffer is then used as input for TMATMUL
+
+Constraint: This template is selected when dst.memory_space == RIGHT.
+"""
+
+import tilelang_dsl as pto
+
+
+def _tmov_m2r_constraint(src: pto.Tile, dst: pto.Tile) -> bool:
+    """Constraint: Mat to Right transfer scenario.
+
+    Supported scenario:
+      - src.memory_space == MAT
+      - dst.memory_space == RIGHT
+    """
+    src_ms = src.memory_space
+    dst_ms = dst.memory_space
+
+    # Check src is MAT
+    if isinstance(src_ms, str):
+        src_is_mat = src_ms == "mat"
+    elif isinstance(src_ms, pto.MemorySpace):
+        src_is_mat = src_ms == pto.MemorySpace.MAT
+    else:
+        src_is_mat = hasattr(src_ms, "value") and src_ms.value == "mat"
+
+    # Check dst is RIGHT
+    if isinstance(dst_ms, str):
+        dst_is_right = dst_ms == "right"
+    elif isinstance(dst_ms, pto.MemorySpace):
+        dst_is_right = dst_ms == pto.MemorySpace.RIGHT
+    else:
+        dst_is_right = hasattr(dst_ms, "value") and dst_ms.value == "right"
+
+    return src_is_mat and dst_is_right
+
+
+@pto.ckernel(
+    target="a5",
+    op="pto.tmov",
+    constraints=[_tmov_m2r_constraint],
+    dtypes=[
+        (pto.f16, pto.f16),
+        (pto.bf16, pto.bf16),
+        (pto.f32, pto.f32),
+        (pto.i8, pto.i8),
+    ],
+)
+def template_tmov_m2r(src: pto.Tile, dst: pto.Tile):
+    """Move data from Mat buffer to Right buffer.
+
+    Args:
+        src: Source tile in L1 Mat location
+        dst: Destination tile in L0B Right location
+
+    The k, n dimensions are derived from the tile shapes.
+    Transpose is typically enabled for Right buffer layout.
+    """
+    k, n = dst.valid_shape
+    pto.mte_l1_l0b(src.as_ptr(), dst.as_ptr(), k, n, transpose=True)
+    return

lib/TileOps/tmov2scale_template.py

@@ -0,0 +1,98 @@
+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+# CANN Open Software License Agreement Version 2.0 (the "License").
+# Please refer to the License for details. You may not use this file except in compliance with the License.
+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+# See LICENSE in the root of the software repository for the full text of the License.
+
+"""TileLang DSL template for pto.tmov - Mat to Scaling/FB buffer movement.
+
+This template implements the TMOV_M2S scenario for cube kernels:
+  - Source: L1 Mat buffer (memory_space="mat", 1xN row-major layout)
+  - Destination: L0 Fixpipe Buffer (memory_space="scaling")
+  - Uses mte_l1_fb (fixpipe buffer load) intrinsic operation
+
+The Fixpipe Buffer (FB) is a 4KB buffer used for storing quantization
+scale parameters in the fixpipe quantization flow. Requirements:
+  - Row dimension must be 1
+  - Column dimension * sizeof(dtype) must be aligned to 128 bits (16 bytes)
+  - Total size must not exceed 4KB (4096 bytes)
+
+Constraint: This template is selected when dst.memory_space == SCALING.
+
+Data format:
+  - Each f32 scale value is stored as ui64 (f32 bits in lower 32 bits, upper 32 bits = 0)
+  - Hardware interprets each ui64's lower 32 bits as one f32 scale value
+  - Total bytes = N * 8 (N ui64 elements)
+  - len_burst = N (number of ui64 elements to transfer)
+"""
+
+import tilelang_dsl as pto
+
+
+def _tmov_m2s_constraint(src: pto.Tile, dst: pto.Tile) -> bool:
+    """Constraint: Mat to Scaling transfer scenario.
+
+    Supported scenario:
+      - src.memory_space == MAT
+      - dst.memory_space == SCALING
+    """
+    src_ms = src.memory_space
+    dst_ms = dst.memory_space
+
+    # Check src is MAT
+    if isinstance(src_ms, str):
+        src_is_mat = src_ms == "mat"
+    elif isinstance(src_ms, pto.MemorySpace):
+        src_is_mat = src_ms == pto.MemorySpace.MAT
+    else:
+        src_is_mat = hasattr(src_ms, "value") and src_ms.value == "mat"
+
+    # Check dst is SCALING
+    if isinstance(dst_ms, str):
+        dst_is_scaling = dst_ms == "scaling"
+    elif isinstance(dst_ms, pto.MemorySpace):
+        dst_is_scaling = dst_ms == pto.MemorySpace.SCALING
+    else:
+        dst_is_scaling = hasattr(dst_ms, "value") and dst_ms.value == "scaling"
+
+    return src_is_mat and dst_is_scaling
+
+
+@pto.ckernel(
+    target="a5",
+    op="pto.tmov",
+    constraints=[_tmov_m2s_constraint],
+    dtypes=[
+        (pto.f32, pto.f32),
+    ],
+)
+def template_tmov_m2s(src: pto.Tile, dst: pto.Tile):
+    """Move data from Mat buffer to Fixpipe Buffer (Scaling).
+
+    Args:
+        src: Source tile in L1 Mat location (1xN row-major)
+        dst: Destination tile in Scaling/FB location
+
+    The scale parameters are loaded into FB for fixpipe quantization.
+
+    Data format: Each f32 scale is stored as ui64 (f32 bits in lower 32 bits).
+    Hardware reads each ui64 and interprets lower 32 bits as f32 scale for column[i].
+
+    mte_l1_fb uses ui64 (8-byte) burst unit:
+      - len_burst = N (number of ui64 elements, where N = column count)
+      - Total bytes = N * 8
+      - n_burst = 1 (single burst configuration)
+    """
+    # Scale has shape 1xN
+    _, n = dst.valid_shape
+    # mte_l1_fb uses ui64 (8-byte) burst unit
+    # Each ui64 contains one f32 scale value (lower 32 bits)
+    # len_burst = N (number of ui64 elements)
+    len_burst = n
+    n_burst = 1
+    src_gap = 0
+    dst_gap = 0
+    pto.mte_l1_fb(src.as_ptr(), dst.as_ptr(), len_burst, nburst=(n_burst, src_gap, dst_gap))
+    return