diff --git a/lib/TileOps/tmatmul_bias_template.py b/lib/TileOps/tmatmul_bias_template.py new file mode 100644 index 0000000000..1b28fe0a24 --- /dev/null +++ b/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 \ No newline at end of file diff --git a/lib/TileOps/tmov2bias_template.py b/lib/TileOps/tmov2bias_template.py new file mode 100644 index 0000000000..565615383a --- /dev/null +++ b/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 \ No newline at end of file diff --git a/lib/TileOps/tmov2left_template.py b/lib/TileOps/tmov2left_template.py new file mode 100644 index 0000000000..8f1863817e --- /dev/null +++ b/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 \ No newline at end of file diff --git a/lib/TileOps/tmov2right_template.py b/lib/TileOps/tmov2right_template.py new file mode 100644 index 0000000000..e72cd5bd3f --- /dev/null +++ b/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 \ No newline at end of file diff --git a/lib/TileOps/tmov2scale_template.py b/lib/TileOps/tmov2scale_template.py new file mode 100644 index 0000000000..c86e237237 --- /dev/null +++ b/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 \ No newline at end of file diff --git a/lib/TileOps/tmov2vec_template.py b/lib/TileOps/tmov2vec_template.py new file mode 100644 index 0000000000..d5f4bde2bb --- /dev/null +++ b/lib/TileOps/tmov2vec_template.py @@ -0,0 +1,85 @@ +# 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 - Acc to Vec/UB buffer movement. + +This template implements the TMOV_A2V scenario for cube kernels: + - Source: L0C Accumulator buffer (memory_space="acc") + - Destination: UB (Unified Buffer) Vec location (memory_space="ub") + - Uses acc_store_ub (mte_l0c_ub) intrinsic operation + +This is part of the fixpipe path where accumulator results are +moved from the cube unit to the vector unit for further processing. + +Constraint: This template is selected when src.memory_space == ACC +and dst.memory_space == UB. +""" + +import tilelang_dsl as pto + + +def _tmov_a2v_constraint(src: pto.Tile, dst: pto.Tile) -> bool: + """Constraint: Acc to Vec/UB transfer scenario. + + Supported scenario: + - src.memory_space == ACC + - dst.memory_space == UB + """ + src_ms = src.memory_space + dst_ms = dst.memory_space + + # Check src is ACC + if isinstance(src_ms, str): + src_is_acc = src_ms == "acc" + elif isinstance(src_ms, pto.MemorySpace): + src_is_acc = src_ms == pto.MemorySpace.ACC + else: + src_is_acc = hasattr(src_ms, "value") and src_ms.value == "acc" + + # Check dst is UB (or "vec" which maps to UB) + if isinstance(dst_ms, str): + dst_is_ub = dst_ms == "ub" or dst_ms == "vec" + elif isinstance(dst_ms, pto.MemorySpace): + dst_is_ub = dst_ms == pto.MemorySpace.UB + else: + dst_is_ub = hasattr(dst_ms, "value") and (dst_ms.value == "ub" or dst_ms.value == "vec") + + return src_is_acc and dst_is_ub + + +@pto.ckernel( + target="a5", + op="pto.tmov", + constraints=[_tmov_a2v_constraint], + dtypes=[ + (pto.f32, pto.f32), + (pto.i32, pto.i32), + ], +) +def template_tmov_a2v(src: pto.Tile, dst: pto.Tile): + """Move data from Acc buffer to Vec/UB buffer. + + Args: + src: Source tile in Acc location + dst: Destination tile in Vec/UB location + + The m, n dimensions and strides are derived from the tile shapes. + This performs NZ2ND layout conversion. + """ + m, n = dst.valid_shape + src_stride = (m + 15) // 16 * 16 # Align to 16 blocks + dst_stride = n # Row-major stride + # mte_l0c_ub takes 7 positional args: src, dst, m, n, src_stride, dst_stride, dst_mode + # dst_mode: integer for sub_blockid mode (0 or 1) + # layout: keyword arg for layout conversion ("nz2nd" for NZ to ND) + pto.mte_l0c_ub( + src.as_ptr(), dst.as_ptr(), m, n, src_stride, dst_stride, + 0, # dst_mode: sub_blockid value 0 + layout="nz2nd", + ) + return \ No newline at end of file diff --git a/lib/TileOps/tmov_fp_template.py b/lib/TileOps/tmov_fp_template.py new file mode 100644 index 0000000000..d61d6c0f92 --- /dev/null +++ b/lib/TileOps/tmov_fp_template.py @@ -0,0 +1,103 @@ +# 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 - Acc to Mat with quantization (fixpipe). + +This template implements the TMOV_FP scenario for cube kernels: + - Source: L0C Accumulator buffer (memory_space="acc") + - Scaling: FB buffer with quantization parameters (memory_space="scaling") + - Destination: L1 Mat buffer (memory_space="mat", quantized output) + - Uses fixpipe intrinsic operation + +This is part of the fixpipe quantization path where: + 1. Matmul results are accumulated in L0C (int32/float) + 2. Scale parameters are loaded into FB buffer + 3. TMOV with fp parameter performs quantization: Acc * scale -> quantized output + +Constraint: This template is selected when src.memory_space == ACC, +scale.memory_space == SCALING, and dst.memory_space == MAT. + +Supported scenarios: + - float32 accumulator -> float16/bf16 output with scale +""" + +import tilelang_dsl as pto + + +def _tmov_fp_constraint(src: pto.Tile, dst: pto.Tile, fp: pto.Tile) -> bool: + """Constraint: Fixpipe quantization scenario. + + Supported scenario: + - src.memory_space == ACC + - dst.memory_space == MAT + - fp.memory_space == SCALING + """ + src_ms = src.memory_space + dst_ms = dst.memory_space + fp_ms = fp.memory_space + + # Check src is ACC + if isinstance(src_ms, str): + src_is_acc = src_ms == "acc" + elif isinstance(src_ms, pto.MemorySpace): + src_is_acc = src_ms == pto.MemorySpace.ACC + else: + src_is_acc = hasattr(src_ms, "value") and src_ms.value == "acc" + + # Check dst is MAT + if isinstance(dst_ms, str): + dst_is_mat = dst_ms == "mat" + elif isinstance(dst_ms, pto.MemorySpace): + dst_is_mat = dst_ms == pto.MemorySpace.MAT + else: + dst_is_mat = hasattr(dst_ms, "value") and dst_ms.value == "mat" + + # Check fp is SCALING + if isinstance(fp_ms, str): + fp_is_scaling = fp_ms == "scaling" + elif isinstance(fp_ms, pto.MemorySpace): + fp_is_scaling = fp_ms == pto.MemorySpace.SCALING + else: + fp_is_scaling = hasattr(fp_ms, "value") and fp_ms.value == "scaling" + + return src_is_acc and dst_is_mat and fp_is_scaling + + +@pto.ckernel( + target="a5", + op="pto.tmov", + constraints=[_tmov_fp_constraint], + dtypes=[ + (pto.f32, pto.f16, pto.f32), # (src, dst, fp) - IR operand order + (pto.f32, pto.bf16, pto.f32), + ], +) +def template_tmov_fp(src: pto.Tile, dst: pto.Tile, fp: pto.Tile): + """Move and quantize data from Acc to Mat with scaling parameters. + + Args: + src: Source tile in Acc location (accumulator) + dst: Destination tile in Mat location (quantized output) + fp: Scaling tile in FB location (quantization params) + + The tmov with fp parameter performs fixpipe quantization using mte_l0c_l1 + with pre_quant keyword argument. + """ + # Get dimensions from destination tile + m, n = dst.valid_shape + # Strides: src is in Acc (fractal layout), dst is in Mat (row-major) + src_stride = (m + 15) // 16 * 16 # Align to 16 blocks for fractal + dst_stride = n # Row-major stride + + # Use mte_l0c_l1 with pre_quant for fixpipe quantization + # quant_mode uses _vec suffix for vector (fb/scaling buffer) payload + pto.mte_l0c_l1( + src.as_ptr(), dst.as_ptr(), m, n, src_stride, dst_stride, + pre_quant=(fp.as_ptr(), "qf322f16_pre_vec"), + ) + return \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/CMakeLists.txt b/test/tilelang_st/npu/a5/src/st/testcase/CMakeLists.txt index 22fbc3f731..9a752bad98 100644 --- a/test/tilelang_st/npu/a5/src/st/testcase/CMakeLists.txt +++ b/test/tilelang_st/npu/a5/src/st/testcase/CMakeLists.txt @@ -194,6 +194,12 @@ set(ALL_TESTCASES tfmods tcmps tmatmul + tmov2left + tmov2right + tmov2bias + tmov2scale + tmov2vec + tmov_fp ) if((TEST_CASE IN_LIST ALL_TESTCASES) OR (TEST_CASE STREQUAL "all")) diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/CMakeLists.txt b/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/CMakeLists.txt new file mode 100644 index 0000000000..aea8ec60fc --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/CMakeLists.txt @@ -0,0 +1,9 @@ +# 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. + +pto_tilelang_cube_st(tmov2bias PTO_LEVEL level3) \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/cases.py b/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/cases.py new file mode 100644 index 0000000000..d1e3c222e2 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/cases.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. + +# coding=utf-8 + +"""Single source of truth for tmov2bias ST test cases. + +Tests the TMOV2BIAS (Mat->Bias) operation with full bias integration: + - TLOAD: GM -> L1 Mat (for A, B, and Bias) + - TMOV2LEFT/TMOV2RIGHT: L1 Mat -> L0A/L0B + - TMOV2BIAS: L1 Mat -> Bias Table (mte_l1_bt) - the operation being tested + - TMATMUL_BIAS: compute Acc = Left x Right + Bias + - TSTORE: Acc -> GM +""" + +import numpy as np + + +CASES = [ + { + "name": "f16_16x16x16", + "dtype_a": np.float16, + "dtype_b": np.float16, + "dtype_c": np.float32, + "dtype_bias": np.float32, + "shape_a": (16, 16), + "shape_b": (16, 16), + "shape_c": (16, 16), + "shape_bias": (1, 16), # Bias shape: 1xN (row dimension) + "eps": 1e-3, + }, +] \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/compare.py b/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/compare.py new file mode 100644 index 0000000000..11a83503a4 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/compare.py @@ -0,0 +1,45 @@ +# 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. + +# coding=utf-8 + +import os +import sys +import numpy as np + +from cases import CASES +from st_common import result_cmp, style_fail, style_pass + + +def main(): + case_filter = sys.argv[1] if len(sys.argv) > 1 else None + + all_passed = True + for case in CASES: + if case_filter is not None and case["name"] != case_filter: + continue + + case_dir = case["name"] + shape_c = case["shape_c"] + golden = np.fromfile(os.path.join(case_dir, "golden.bin"), dtype=np.float32).reshape(shape_c) + output = np.fromfile(os.path.join(case_dir, "output.bin"), dtype=np.float32).reshape(shape_c) + + ok = result_cmp(golden, output, case["eps"]) + if ok: + print(style_pass(f"[INFO] {case['name']}: compare passed")) + else: + print(style_fail(f"[ERROR] {case['name']}: compare failed")) + all_passed = False + + if not all_passed: + sys.exit(2) + print(style_pass("[INFO] all cases passed")) + + +if __name__ == "__main__": + main() \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/gen_data.py b/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/gen_data.py new file mode 100644 index 0000000000..bc09cb2a41 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/gen_data.py @@ -0,0 +1,41 @@ +# 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. + +# coding=utf-8 + +import numpy as np + +from cases import CASES +from st_common import setup_case_rng, save_case_data + + +for case in CASES: + setup_case_rng(case) + + shape_a = case["shape_a"] + shape_b = case["shape_b"] + shape_bias = case["shape_bias"] + dtype_a = case["dtype_a"] + dtype_b = case["dtype_b"] + dtype_bias = case["dtype_bias"] + + lhs = np.random.uniform(-1.0, 1.0, size=shape_a).astype(dtype_a) + rhs = np.random.uniform(-1.0, 1.0, size=shape_b).astype(dtype_b) + bias = np.random.uniform(-0.5, 0.5, size=shape_bias).astype(dtype_bias) + + # Compute golden: matmul + bias as float32 + # Acc = lhs @ rhs, then add bias (broadcast along column dimension) + matmul_result = np.matmul(lhs.astype(np.float32), rhs.astype(np.float32)) + golden = matmul_result + bias.astype(np.float32) # bias broadcasts to (1, N) -> (M, N) + golden = golden.astype(np.float32) + + save_case_data(case["name"], {"input1": lhs, "input2": rhs, "bias": bias, "golden": golden}) + print( + f"[INFO] gen_data: {case['name']} " + f"lhs={shape_a} rhs={shape_b} bias={shape_bias} out={case['shape_c']} dtype=float16" + ) \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/launch.cpp b/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/launch.cpp new file mode 100644 index 0000000000..710e41493a --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/launch.cpp @@ -0,0 +1,19 @@ +// 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. + +#include + +#ifndef AICORE +#define AICORE [aicore] +#endif + +extern "C" __global__ AICORE void TMOV2BIAS_f16_16x16x16(__gm__ uint16_t *a, __gm__ uint16_t *b, __gm__ float *bias, __gm__ float *c); + +void LaunchTMOV2BIAS_f16_16x16x16(uint16_t *a, uint16_t *b, float *bias, float *c, void *stream) { + TMOV2BIAS_f16_16x16x16<<<1, nullptr, stream>>>((__gm__ uint16_t *)a, (__gm__ uint16_t *)b, (__gm__ float *)bias, (__gm__ float *)c); +} \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/main.cpp b/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/main.cpp new file mode 100644 index 0000000000..5e73ec670f --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/main.cpp @@ -0,0 +1,127 @@ +// 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. + +#include "acl/acl.h" +#include "test_common.h" +#include +#include +#include +#include +#include + +using namespace PtoTestCommon; + +void LaunchTMOV2BIAS_f16_16x16x16(uint16_t *a, uint16_t *b, float *bias, float *c, void *stream); + +struct TestCase { + const char *name; + void (*launch)(uint16_t *, uint16_t *, float *, float *, void *); + size_t lhsRows; + size_t lhsCols; + size_t rhsRows; + size_t rhsCols; + size_t biasRows; + size_t biasCols; + size_t outRows; + size_t outCols; +}; + +static const TestCase kCases[] = { + {"f16_16x16x16", LaunchTMOV2BIAS_f16_16x16x16, 16, 16, 16, 16, 1, 16, 16, 16}, +}; +static constexpr size_t kNumCases = sizeof(kCases) / sizeof(kCases[0]); + +static int RunCase(const TestCase &tc, int deviceId, aclrtStream stream) { + (void)deviceId; + int rc = 0; + size_t lhsBytes = tc.lhsRows * tc.lhsCols * sizeof(uint16_t); + size_t rhsBytes = tc.rhsRows * tc.rhsCols * sizeof(uint16_t); + size_t biasBytes = tc.biasRows * tc.biasCols * sizeof(float); + size_t outBytes = tc.outRows * tc.outCols * sizeof(float); + + std::printf("[INFO] === case: %s (lhs=%zux%zu, rhs=%zux%zu, bias=%zux%zu, out=%zux%zu) ===\n", + tc.name, tc.lhsRows, tc.lhsCols, tc.rhsRows, tc.rhsCols, tc.biasRows, tc.biasCols, tc.outRows, tc.outCols); + + std::string caseDir = std::string("./") + tc.name; + + void *lhsHost = nullptr, *rhsHost = nullptr, *biasHost = nullptr, *outHost = nullptr; + void *lhsDevice = nullptr, *rhsDevice = nullptr, *biasDevice = nullptr, *outDevice = nullptr; + + aclrtMallocHost(&lhsHost, lhsBytes); + aclrtMallocHost(&rhsHost, rhsBytes); + aclrtMallocHost(&biasHost, biasBytes); + aclrtMallocHost(&outHost, outBytes); + + aclrtMalloc(&lhsDevice, lhsBytes, ACL_MEM_MALLOC_HUGE_FIRST); + aclrtMalloc(&rhsDevice, rhsBytes, ACL_MEM_MALLOC_HUGE_FIRST); + aclrtMalloc(&biasDevice, biasBytes, ACL_MEM_MALLOC_HUGE_FIRST); + aclrtMalloc(&outDevice, outBytes, ACL_MEM_MALLOC_HUGE_FIRST); + + if (!ReadFile((caseDir + "/input1.bin").c_str(), lhsBytes, lhsHost, lhsBytes)) { + std::fprintf(stderr, "[ERROR] failed to read %s/input1.bin\n", caseDir.c_str()); + rc = 1; + } + if (rc == 0 && !ReadFile((caseDir + "/input2.bin").c_str(), rhsBytes, rhsHost, rhsBytes)) { + std::fprintf(stderr, "[ERROR] failed to read %s/input2.bin\n", caseDir.c_str()); + rc = 1; + } + if (rc == 0 && !ReadFile((caseDir + "/bias.bin").c_str(), biasBytes, biasHost, biasBytes)) { + std::fprintf(stderr, "[ERROR] failed to read %s/bias.bin\n", caseDir.c_str()); + rc = 1; + } + + if (rc == 0) { + aclrtMemcpy(lhsDevice, lhsBytes, lhsHost, lhsBytes, ACL_MEMCPY_HOST_TO_DEVICE); + aclrtMemcpy(rhsDevice, rhsBytes, rhsHost, rhsBytes, ACL_MEMCPY_HOST_TO_DEVICE); + aclrtMemcpy(biasDevice, biasBytes, biasHost, biasBytes, ACL_MEMCPY_HOST_TO_DEVICE); + + tc.launch(static_cast(lhsDevice), static_cast(rhsDevice), + static_cast(biasDevice), static_cast(outDevice), stream); + + aclrtSynchronizeStream(stream); + aclrtMemcpy(outHost, outBytes, outDevice, outBytes, ACL_MEMCPY_DEVICE_TO_HOST); + } + + if (rc == 0 && !WriteFile((caseDir + "/output.bin").c_str(), outHost, outBytes)) { + std::fprintf(stderr, "[ERROR] failed to write %s/output.bin\n", caseDir.c_str()); + rc = 1; + } + + if (lhsDevice) aclrtFree(lhsDevice); + if (rhsDevice) aclrtFree(rhsDevice); + if (biasDevice) aclrtFree(biasDevice); + if (outDevice) aclrtFree(outDevice); + if (lhsHost) aclrtFreeHost(lhsHost); + if (rhsHost) aclrtFreeHost(rhsHost); + if (biasHost) aclrtFreeHost(biasHost); + if (outHost) aclrtFreeHost(outHost); + + if (rc == 0) std::printf("[INFO] case %s done\n", tc.name); + return rc; +} + +int main(int argc, char *argv[]) { + const char *caseFilter = (argc > 1) ? argv[1] : nullptr; + int rc = 0, deviceId = 0; + aclrtStream stream = nullptr; + + aclInit(nullptr); + if (const char *envDevice = std::getenv("ACL_DEVICE_ID")) deviceId = std::atoi(envDevice); + aclrtSetDevice(deviceId); + aclrtCreateStream(&stream); + + for (size_t i = 0; i < kNumCases; ++i) { + if (caseFilter && std::strcmp(kCases[i].name, caseFilter) != 0) continue; + if (RunCase(kCases[i], deviceId, stream) != 0) { rc = 1; break; } + } + + if (stream) aclrtDestroyStream(stream); + aclrtResetDevice(deviceId); + aclFinalize(); + return rc; +} \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/tmov2bias.pto b/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/tmov2bias.pto new file mode 100644 index 0000000000..42ed1db6e2 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2bias/tmov2bias.pto @@ -0,0 +1,139 @@ +// 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 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 ST kernel for cube tmov (Mat->Bias): tests template expansion path with full bias integration. +// Uses pto.tmov OP which will be expanded via --enable-tile-op-expand. +// The template_tmov_m2b template should be selected when dst memory_space is "bias". +// Compiled by ptoas --enable-insert-sync --enable-tile-op-expand --pto-backend=vpto --pto-level=level3 + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @TMOV2BIAS_f16_16x16x16(%a_gm: !pto.ptr, %b_gm: !pto.ptr, %bias_gm: !pto.ptr, %c_gm: !pto.ptr) attributes {pto.aicore} { + %c0_i64 = arith.constant 0 : i64 + %c1_i64 = arith.constant 1 : i64 + %c16_i64 = arith.constant 16 : i64 + %c32_i64 = arith.constant 32 : i64 + %c64_i64 = arith.constant 64 : i64 // 16 * 4 bytes for bias (1x16 f32) + %c256_i64 = arith.constant 256 : i64 + %c512_i64 = arith.constant 512 : i64 + %c1024_i64 = arith.constant 1024 : i64 + %false = arith.constant false + + // L1 Mat tiles (source for TMOV) + %l1_a_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + %l1_b_tile = pto.alloc_tile addr = %c512_i64 + : !pto.tile_buf + + // L1 Bias Mat tile (holding bias data from GM, will be moved to Bias Table) + %l1_bias_tile = pto.alloc_tile addr = %c1024_i64 + : !pto.tile_buf + + // L0A/L0B tiles + %l0a_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + %l0b_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + + // Bias tile (L0/BT) - destination for TMOV Mat->Bias + %bias_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + + // Acc tile + %l0c_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + + // Get addresses + %l1_a = pto.tile_buf_addr %l1_a_tile + : !pto.tile_buf + -> !pto.ptr + %l1_b = pto.tile_buf_addr %l1_b_tile + : !pto.tile_buf + -> !pto.ptr + %l1_bias = pto.tile_buf_addr %l1_bias_tile + : !pto.tile_buf + -> !pto.ptr + %l0a = pto.tile_buf_addr %l0a_tile + : !pto.tile_buf + -> !pto.ptr + %l0b = pto.tile_buf_addr %l0b_tile + : !pto.tile_buf + -> !pto.ptr + %l0c = pto.tile_buf_addr %l0c_tile + : !pto.tile_buf + -> !pto.ptr + + // TLOAD A: GM -> L1 Mat + pto.mte_gm_l1_frac %a_gm, %l1_a, nd2nz, + shape(%c16_i64, %c16_i64), + src_layout(%c32_i64), + dst_group(%c1_i64, %c1_i64, %c16_i64, %c0_i64), + ctrl(%c0_i64, %false) + : !pto.ptr, !pto.ptr, nd2nz, + shape i64, i64, src_layout(i64), + dst_group i64, i64, i64, i64, ctrl i64, i1 + + pto.set_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID0"] + pto.wait_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID0"] + + // TMOV Mat->Left: L1 Mat -> L0A Left (using pto.tmov for template expansion) + pto.tmov ins(%l1_a_tile : !pto.tile_buf) + outs(%l0a_tile : !pto.tile_buf) + + // TLOAD B: GM -> L1 Mat + pto.mte_gm_l1_frac %b_gm, %l1_b, nd2nz, + shape(%c16_i64, %c16_i64), + src_layout(%c32_i64), + dst_group(%c1_i64, %c1_i64, %c16_i64, %c0_i64), + ctrl(%c0_i64, %false) + : !pto.ptr, !pto.ptr, nd2nz, + shape i64, i64, src_layout(i64), + dst_group i64, i64, i64, i64, ctrl i64, i1 + + pto.set_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID1"] + pto.wait_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID1"] + + // TMOV Mat->Right: L1 Mat -> L0B Right (using pto.tmov for template expansion) + pto.tmov ins(%l1_b_tile : !pto.tile_buf) + outs(%l0b_tile : !pto.tile_buf) + + // TLOAD Bias: GM -> L1 Mat (bias data is 1x16 f32 = 64 bytes) + // Use mte_gm_l1 for simple linear load (bias is row-major 1x16) + pto.mte_gm_l1 %bias_gm, %l1_bias, %c64_i64 + nburst(%c1_i64, %c0_i64, %c0_i64) + : !pto.ptr, !pto.ptr, i64, i64, i64, i64 + + pto.set_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID2"] + pto.wait_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID2"] + + // TMOV Mat->Bias: L1 Mat -> Bias Table (using pto.tmov for template expansion) + // This tests the template_tmov_m2b template with constraint dst.memory_space == BIAS + pto.tmov ins(%l1_bias_tile : !pto.tile_buf) + outs(%bias_tile : !pto.tile_buf) + + pto.set_flag["PIPE_MTE1", "PIPE_M", "EVENT_ID0"] + pto.wait_flag["PIPE_MTE1", "PIPE_M", "EVENT_ID0"] + + // TMATMUL_BIAS: compute Acc = Left x Right + Bias + // Uses bias_tile loaded via TMOV Mat->Bias - validates bias_tile content + pto.tmatmul.bias ins(%l0a_tile, %l0b_tile, %bias_tile : !pto.tile_buf, + !pto.tile_buf, + !pto.tile_buf) + outs(%l0c_tile : !pto.tile_buf) + + pto.set_flag["PIPE_M", "PIPE_FIX", "EVENT_ID1"] + pto.wait_flag["PIPE_M", "PIPE_FIX", "EVENT_ID1"] + + // TSTORE: Acc -> GM (output includes bias contribution) + pto.mte_l0c_gm %l0c, %c_gm, %c16_i64, %c16_i64, %c16_i64, %c16_i64, %c0_i64, %c0_i64, + nz2nd + : !pto.ptr, !pto.ptr, i64, i64, i64, i64, i64, i64 + + pto.barrier #pto.pipe + return + } +} \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/CMakeLists.txt b/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/CMakeLists.txt new file mode 100644 index 0000000000..1523762ca5 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/CMakeLists.txt @@ -0,0 +1,9 @@ +# 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. + +pto_tilelang_cube_st(tmov2left PTO_LEVEL level3) \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/cases.py b/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/cases.py new file mode 100644 index 0000000000..a9c00503ab --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/cases.py @@ -0,0 +1,36 @@ +# 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. + +# coding=utf-8 + +"""Single source of truth for tmov2left ST test cases. + +Tests the TMOV2LEFT (Mat->Left) operation explicitly through a matmul flow: + - TLOAD: GM -> L1 Mat + - TMOV2LEFT: L1 Mat -> L0A Left (mte_l1_l0a) - the operation being tested + - TMOV2RIGHT: L1 Mat -> L0B Right (mte_l1_l0b) + - TMATMUL: compute Acc = Left x Right + - TSTORE: Acc -> GM + +The correctness of TMOV2LEFT is verified by comparing the matmul output +with the expected golden result. +""" + +import numpy as np + + +CASES = [ + { + "name": "f16_16x16x16", + "dtype": np.float16, + "shape_a": (16, 16), + "shape_b": (16, 16), + "shape_c": (16, 16), + "eps": 1e-2, + }, +] \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/compare.py b/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/compare.py new file mode 100644 index 0000000000..11a83503a4 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/compare.py @@ -0,0 +1,45 @@ +# 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. + +# coding=utf-8 + +import os +import sys +import numpy as np + +from cases import CASES +from st_common import result_cmp, style_fail, style_pass + + +def main(): + case_filter = sys.argv[1] if len(sys.argv) > 1 else None + + all_passed = True + for case in CASES: + if case_filter is not None and case["name"] != case_filter: + continue + + case_dir = case["name"] + shape_c = case["shape_c"] + golden = np.fromfile(os.path.join(case_dir, "golden.bin"), dtype=np.float32).reshape(shape_c) + output = np.fromfile(os.path.join(case_dir, "output.bin"), dtype=np.float32).reshape(shape_c) + + ok = result_cmp(golden, output, case["eps"]) + if ok: + print(style_pass(f"[INFO] {case['name']}: compare passed")) + else: + print(style_fail(f"[ERROR] {case['name']}: compare failed")) + all_passed = False + + if not all_passed: + sys.exit(2) + print(style_pass("[INFO] all cases passed")) + + +if __name__ == "__main__": + main() \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/gen_data.py b/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/gen_data.py new file mode 100644 index 0000000000..3e79bc8707 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/gen_data.py @@ -0,0 +1,32 @@ +# 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. + +# coding=utf-8 + +import numpy as np + +from cases import CASES +from st_common import setup_case_rng, save_case_data + + +for case in CASES: + setup_case_rng(case) + + shape_a = case["shape_a"] + shape_b = case["shape_b"] + dtype = case["dtype"] + + lhs = np.random.uniform(-1.0, 1.0, size=shape_a).astype(dtype) + rhs = np.random.uniform(-1.0, 1.0, size=shape_b).astype(dtype) + golden = np.matmul(lhs.astype(np.float32), rhs.astype(np.float32)).astype(np.float32) + + save_case_data(case["name"], {"input1": lhs, "input2": rhs, "golden": golden}) + print( + f"[INFO] gen_data: {case['name']} " + f"lhs={shape_a} rhs={shape_b} out={case['shape_c']} dtype={dtype.__name__}" + ) \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/launch.cpp b/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/launch.cpp new file mode 100644 index 0000000000..b1d2b2a05e --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/launch.cpp @@ -0,0 +1,19 @@ +// 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. + +#include + +#ifndef AICORE +#define AICORE [aicore] +#endif + +extern "C" __global__ AICORE void TMOV2LEFT_f16_16x16x16(__gm__ uint16_t *a, __gm__ uint16_t *b, __gm__ float *c); + +void LaunchTMOV2LEFT_f16_16x16x16(uint16_t *a, uint16_t *b, float *c, void *stream) { + TMOV2LEFT_f16_16x16x16<<<1, nullptr, stream>>>((__gm__ uint16_t *)a, (__gm__ uint16_t *)b, (__gm__ float *)c); +} \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/main.cpp b/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/main.cpp new file mode 100644 index 0000000000..222be97084 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/main.cpp @@ -0,0 +1,158 @@ +// 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. + +#include "acl/acl.h" +#include "test_common.h" +#include +#include +#include +#include +#include + +using namespace PtoTestCommon; + +void LaunchTMOV2LEFT_f16_16x16x16(uint16_t *a, uint16_t *b, float *c, void *stream); + +using LaunchFn = void (*)(uint16_t *, uint16_t *, float *, void *); + +struct TestCase { + const char *name; + LaunchFn launch; + size_t lhsRows; + size_t lhsCols; + size_t rhsRows; + size_t rhsCols; + size_t outRows; + size_t outCols; +}; + +static const TestCase kCases[] = { + {"f16_16x16x16", LaunchTMOV2LEFT_f16_16x16x16, 16, 16, 16, 16, 16, 16}, +}; +static constexpr size_t kNumCases = sizeof(kCases) / sizeof(kCases[0]); + +static int RunCase(const TestCase &tc, int deviceId, aclrtStream stream) { + (void)deviceId; + int rc = 0; + const size_t lhsElems = tc.lhsRows * tc.lhsCols; + const size_t rhsElems = tc.rhsRows * tc.rhsCols; + const size_t outElems = tc.outRows * tc.outCols; + const size_t lhsBytes = lhsElems * sizeof(uint16_t); + const size_t rhsBytes = rhsElems * sizeof(uint16_t); + const size_t outBytes = outElems * sizeof(float); + size_t lhsFileSize = lhsBytes; + size_t rhsFileSize = rhsBytes; + + std::printf( + "[INFO] === case: %s (lhs=%zux%zu, rhs=%zux%zu, out=%zux%zu) ===\n", + tc.name, + tc.lhsRows, + tc.lhsCols, + tc.rhsRows, + tc.rhsCols, + tc.outRows, + tc.outCols + ); + + std::string caseDir = std::string("./") + tc.name; + + void *lhsHost = nullptr; + void *rhsHost = nullptr; + void *outHost = nullptr; + void *lhsDevice = nullptr; + void *rhsDevice = nullptr; + void *outDevice = nullptr; + + aclrtMallocHost(&lhsHost, lhsBytes); + aclrtMallocHost(&rhsHost, rhsBytes); + aclrtMallocHost(&outHost, outBytes); + + aclrtMalloc(&lhsDevice, lhsBytes, ACL_MEM_MALLOC_HUGE_FIRST); + aclrtMalloc(&rhsDevice, rhsBytes, ACL_MEM_MALLOC_HUGE_FIRST); + aclrtMalloc(&outDevice, outBytes, ACL_MEM_MALLOC_HUGE_FIRST); + + if (!ReadFile((caseDir + "/input1.bin").c_str(), lhsFileSize, lhsHost, lhsBytes)) { + std::fprintf(stderr, "[ERROR] failed to read %s/input1.bin\n", caseDir.c_str()); + rc = 1; + } + if (rc == 0 && !ReadFile((caseDir + "/input2.bin").c_str(), rhsFileSize, rhsHost, rhsBytes)) { + std::fprintf(stderr, "[ERROR] failed to read %s/input2.bin\n", caseDir.c_str()); + rc = 1; + } + + if (rc == 0) { + aclrtMemcpy(lhsDevice, lhsBytes, lhsHost, lhsBytes, ACL_MEMCPY_HOST_TO_DEVICE); + aclrtMemcpy(rhsDevice, rhsBytes, rhsHost, rhsBytes, ACL_MEMCPY_HOST_TO_DEVICE); + + tc.launch( + static_cast(lhsDevice), + static_cast(rhsDevice), + static_cast(outDevice), + stream + ); + + aclrtSynchronizeStream(stream); + aclrtMemcpy(outHost, outBytes, outDevice, outBytes, ACL_MEMCPY_DEVICE_TO_HOST); + } + + if (rc == 0 && !WriteFile((caseDir + "/output.bin").c_str(), outHost, outBytes)) { + std::fprintf(stderr, "[ERROR] failed to write %s/output.bin\n", caseDir.c_str()); + rc = 1; + } + + if (lhsDevice != nullptr) + aclrtFree(lhsDevice); + if (rhsDevice != nullptr) + aclrtFree(rhsDevice); + if (outDevice != nullptr) + aclrtFree(outDevice); + if (lhsHost != nullptr) + aclrtFreeHost(lhsHost); + if (rhsHost != nullptr) + aclrtFreeHost(rhsHost); + if (outHost != nullptr) + aclrtFreeHost(outHost); + + if (rc == 0) + std::printf("[INFO] case %s done\n", tc.name); + return rc; +} + +int main(int argc, char *argv[]) { + const char *caseFilter = (argc > 1) ? argv[1] : nullptr; + + int rc = 0; + int deviceId = 0; + aclrtStream stream = nullptr; + + aclInit(nullptr); + if (const char *envDevice = std::getenv("ACL_DEVICE_ID")) { + deviceId = std::atoi(envDevice); + } + aclrtSetDevice(deviceId); + aclrtCreateStream(&stream); + + for (size_t i = 0; i < kNumCases; ++i) { + if (caseFilter != nullptr && std::strcmp(kCases[i].name, caseFilter) != 0) { + continue; + } + int ret = RunCase(kCases[i], deviceId, stream); + if (ret != 0) { + std::fprintf(stderr, "[ERROR] case %s failed\n", kCases[i].name); + rc = 1; + break; + } + } + + if (stream != nullptr) + aclrtDestroyStream(stream); + aclrtResetDevice(deviceId); + aclFinalize(); + + return rc; +} \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/tmov2left.pto b/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/tmov2left.pto new file mode 100644 index 0000000000..34514462f6 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2left/tmov2left.pto @@ -0,0 +1,106 @@ +// 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 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 ST kernel for cube tmov (Mat->Left): tests template expansion path. +// Uses pto.tmov OP which will be expanded via --enable-tile-op-expand. +// The template_tmov_m2l template should be selected when dst memory_space is "left". +// Compiled by ptoas --enable-insert-sync --enable-tile-op-expand --pto-backend=vpto + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @TMOV2LEFT_f16_16x16x16(%a_gm: !pto.ptr, %b_gm: !pto.ptr, %c_gm: !pto.ptr) attributes {pto.aicore} { + %c0_i64 = arith.constant 0 : i64 + %c1_i64 = arith.constant 1 : i64 + %c16_i64 = arith.constant 16 : i64 + %c32_i64 = arith.constant 32 : i64 + %c256_i64 = arith.constant 256 : i64 + %c512_i64 = arith.constant 512 : i64 + %false = arith.constant false + + // L1 Mat tiles (source for TMOV) + %l1_a_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + %l1_b_tile = pto.alloc_tile addr = %c512_i64 + : !pto.tile_buf + + // L0A/L0B tiles (destination for TMOV Mat->Left/Mat->Right) + %l0a_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + %l0b_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + + // Acc tile + %l0c_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + + // Get addresses for level3 operations + %l1_a = pto.tile_buf_addr %l1_a_tile + : !pto.tile_buf + -> !pto.ptr + %l1_b = pto.tile_buf_addr %l1_b_tile + : !pto.tile_buf + -> !pto.ptr + %l0c = pto.tile_buf_addr %l0c_tile + : !pto.tile_buf + -> !pto.ptr + + // TLOAD: GM -> L1 Mat (nd2nz conversion) + pto.mte_gm_l1_frac %a_gm, %l1_a, nd2nz, + shape(%c16_i64, %c16_i64), + src_layout(%c32_i64), + dst_group(%c1_i64, %c1_i64, %c16_i64, %c0_i64), + ctrl(%c0_i64, %false) + : !pto.ptr, !pto.ptr, nd2nz, + shape i64, i64, src_layout(i64), + dst_group i64, i64, i64, i64, ctrl i64, i1 + + pto.set_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID0"] + pto.wait_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID0"] + + // TMOV Mat->Left: L1 Mat -> L0A Left (using pto.tmov for template expansion) + pto.tmov ins(%l1_a_tile : !pto.tile_buf) + outs(%l0a_tile : !pto.tile_buf) + + // TLOAD B: GM -> L1 Mat + pto.mte_gm_l1_frac %b_gm, %l1_b, nd2nz, + shape(%c16_i64, %c16_i64), + src_layout(%c32_i64), + dst_group(%c1_i64, %c1_i64, %c16_i64, %c0_i64), + ctrl(%c0_i64, %false) + : !pto.ptr, !pto.ptr, nd2nz, + shape i64, i64, src_layout(i64), + dst_group i64, i64, i64, i64, ctrl i64, i1 + + pto.set_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID1"] + pto.wait_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID1"] + + // TMOV Mat->Right: L1 Mat -> L0B Right (using pto.tmov for template expansion) + pto.tmov ins(%l1_b_tile : !pto.tile_buf) + outs(%l0b_tile : !pto.tile_buf) + + // Sync before matmul + pto.set_flag["PIPE_MTE1", "PIPE_M", "EVENT_ID0"] + pto.wait_flag["PIPE_MTE1", "PIPE_M", "EVENT_ID0"] + + // TMATMUL: compute + pto.tmatmul ins(%l0a_tile, %l0b_tile : !pto.tile_buf, + !pto.tile_buf) + outs(%l0c_tile : !pto.tile_buf) + + // Sync before store + pto.set_flag["PIPE_M", "PIPE_FIX", "EVENT_ID1"] + pto.wait_flag["PIPE_M", "PIPE_FIX", "EVENT_ID1"] + + // TSTORE: Acc -> GM (nz2nd conversion) + pto.mte_l0c_gm %l0c, %c_gm, %c16_i64, %c16_i64, %c16_i64, %c16_i64, %c0_i64, %c0_i64, + nz2nd + : !pto.ptr, !pto.ptr, i64, i64, i64, i64, i64, i64 + + pto.barrier #pto.pipe + return + } +} \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/CMakeLists.txt b/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/CMakeLists.txt new file mode 100644 index 0000000000..d9eaabda2f --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/CMakeLists.txt @@ -0,0 +1,9 @@ +# 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. + +pto_tilelang_cube_st(tmov2right PTO_LEVEL level3) \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/cases.py b/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/cases.py new file mode 100644 index 0000000000..1fc3de0639 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/cases.py @@ -0,0 +1,36 @@ +# 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. + +# coding=utf-8 + +"""Single source of truth for tmov2right ST test cases. + +Tests the TMOV2RIGHT (Mat->Right) operation explicitly through a matmul flow: + - TLOAD: GM -> L1 Mat + - TMOV2LEFT: L1 Mat -> L0A Left (mte_l1_l0a) + - TMOV2RIGHT: L1 Mat -> L0B Right (mte_l1_l0b) - the operation being tested + - TMATMUL: compute Acc = Left x Right + - TSTORE: Acc -> GM + +The correctness of TMOV2RIGHT is verified by comparing the matmul output +with the expected golden result. +""" + +import numpy as np + + +CASES = [ + { + "name": "f16_16x16x16", + "dtype": np.float16, + "shape_a": (16, 16), + "shape_b": (16, 16), + "shape_c": (16, 16), + "eps": 1e-2, + }, +] \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/compare.py b/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/compare.py new file mode 100644 index 0000000000..11a83503a4 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/compare.py @@ -0,0 +1,45 @@ +# 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. + +# coding=utf-8 + +import os +import sys +import numpy as np + +from cases import CASES +from st_common import result_cmp, style_fail, style_pass + + +def main(): + case_filter = sys.argv[1] if len(sys.argv) > 1 else None + + all_passed = True + for case in CASES: + if case_filter is not None and case["name"] != case_filter: + continue + + case_dir = case["name"] + shape_c = case["shape_c"] + golden = np.fromfile(os.path.join(case_dir, "golden.bin"), dtype=np.float32).reshape(shape_c) + output = np.fromfile(os.path.join(case_dir, "output.bin"), dtype=np.float32).reshape(shape_c) + + ok = result_cmp(golden, output, case["eps"]) + if ok: + print(style_pass(f"[INFO] {case['name']}: compare passed")) + else: + print(style_fail(f"[ERROR] {case['name']}: compare failed")) + all_passed = False + + if not all_passed: + sys.exit(2) + print(style_pass("[INFO] all cases passed")) + + +if __name__ == "__main__": + main() \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/gen_data.py b/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/gen_data.py new file mode 100644 index 0000000000..3e79bc8707 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/gen_data.py @@ -0,0 +1,32 @@ +# 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. + +# coding=utf-8 + +import numpy as np + +from cases import CASES +from st_common import setup_case_rng, save_case_data + + +for case in CASES: + setup_case_rng(case) + + shape_a = case["shape_a"] + shape_b = case["shape_b"] + dtype = case["dtype"] + + lhs = np.random.uniform(-1.0, 1.0, size=shape_a).astype(dtype) + rhs = np.random.uniform(-1.0, 1.0, size=shape_b).astype(dtype) + golden = np.matmul(lhs.astype(np.float32), rhs.astype(np.float32)).astype(np.float32) + + save_case_data(case["name"], {"input1": lhs, "input2": rhs, "golden": golden}) + print( + f"[INFO] gen_data: {case['name']} " + f"lhs={shape_a} rhs={shape_b} out={case['shape_c']} dtype={dtype.__name__}" + ) \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/launch.cpp b/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/launch.cpp new file mode 100644 index 0000000000..13b3ac1b66 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/launch.cpp @@ -0,0 +1,19 @@ +// 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. + +#include + +#ifndef AICORE +#define AICORE [aicore] +#endif + +extern "C" __global__ AICORE void TMOV2RIGHT_f16_16x16x16(__gm__ uint16_t *a, __gm__ uint16_t *b, __gm__ float *c); + +void LaunchTMOV2RIGHT_f16_16x16x16(uint16_t *a, uint16_t *b, float *c, void *stream) { + TMOV2RIGHT_f16_16x16x16<<<1, nullptr, stream>>>((__gm__ uint16_t *)a, (__gm__ uint16_t *)b, (__gm__ float *)c); +} \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/main.cpp b/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/main.cpp new file mode 100644 index 0000000000..bfa16648a9 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/main.cpp @@ -0,0 +1,158 @@ +// 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. + +#include "acl/acl.h" +#include "test_common.h" +#include +#include +#include +#include +#include + +using namespace PtoTestCommon; + +void LaunchTMOV2RIGHT_f16_16x16x16(uint16_t *a, uint16_t *b, float *c, void *stream); + +using LaunchFn = void (*)(uint16_t *, uint16_t *, float *, void *); + +struct TestCase { + const char *name; + LaunchFn launch; + size_t lhsRows; + size_t lhsCols; + size_t rhsRows; + size_t rhsCols; + size_t outRows; + size_t outCols; +}; + +static const TestCase kCases[] = { + {"f16_16x16x16", LaunchTMOV2RIGHT_f16_16x16x16, 16, 16, 16, 16, 16, 16}, +}; +static constexpr size_t kNumCases = sizeof(kCases) / sizeof(kCases[0]); + +static int RunCase(const TestCase &tc, int deviceId, aclrtStream stream) { + (void)deviceId; + int rc = 0; + const size_t lhsElems = tc.lhsRows * tc.lhsCols; + const size_t rhsElems = tc.rhsRows * tc.rhsCols; + const size_t outElems = tc.outRows * tc.outCols; + const size_t lhsBytes = lhsElems * sizeof(uint16_t); + const size_t rhsBytes = rhsElems * sizeof(uint16_t); + const size_t outBytes = outElems * sizeof(float); + size_t lhsFileSize = lhsBytes; + size_t rhsFileSize = rhsBytes; + + std::printf( + "[INFO] === case: %s (lhs=%zux%zu, rhs=%zux%zu, out=%zux%zu) ===\n", + tc.name, + tc.lhsRows, + tc.lhsCols, + tc.rhsRows, + tc.rhsCols, + tc.outRows, + tc.outCols + ); + + std::string caseDir = std::string("./") + tc.name; + + void *lhsHost = nullptr; + void *rhsHost = nullptr; + void *outHost = nullptr; + void *lhsDevice = nullptr; + void *rhsDevice = nullptr; + void *outDevice = nullptr; + + aclrtMallocHost(&lhsHost, lhsBytes); + aclrtMallocHost(&rhsHost, rhsBytes); + aclrtMallocHost(&outHost, outBytes); + + aclrtMalloc(&lhsDevice, lhsBytes, ACL_MEM_MALLOC_HUGE_FIRST); + aclrtMalloc(&rhsDevice, rhsBytes, ACL_MEM_MALLOC_HUGE_FIRST); + aclrtMalloc(&outDevice, outBytes, ACL_MEM_MALLOC_HUGE_FIRST); + + if (!ReadFile((caseDir + "/input1.bin").c_str(), lhsFileSize, lhsHost, lhsBytes)) { + std::fprintf(stderr, "[ERROR] failed to read %s/input1.bin\n", caseDir.c_str()); + rc = 1; + } + if (rc == 0 && !ReadFile((caseDir + "/input2.bin").c_str(), rhsFileSize, rhsHost, rhsBytes)) { + std::fprintf(stderr, "[ERROR] failed to read %s/input2.bin\n", caseDir.c_str()); + rc = 1; + } + + if (rc == 0) { + aclrtMemcpy(lhsDevice, lhsBytes, lhsHost, lhsBytes, ACL_MEMCPY_HOST_TO_DEVICE); + aclrtMemcpy(rhsDevice, rhsBytes, rhsHost, rhsBytes, ACL_MEMCPY_HOST_TO_DEVICE); + + tc.launch( + static_cast(lhsDevice), + static_cast(rhsDevice), + static_cast(outDevice), + stream + ); + + aclrtSynchronizeStream(stream); + aclrtMemcpy(outHost, outBytes, outDevice, outBytes, ACL_MEMCPY_DEVICE_TO_HOST); + } + + if (rc == 0 && !WriteFile((caseDir + "/output.bin").c_str(), outHost, outBytes)) { + std::fprintf(stderr, "[ERROR] failed to write %s/output.bin\n", caseDir.c_str()); + rc = 1; + } + + if (lhsDevice != nullptr) + aclrtFree(lhsDevice); + if (rhsDevice != nullptr) + aclrtFree(rhsDevice); + if (outDevice != nullptr) + aclrtFree(outDevice); + if (lhsHost != nullptr) + aclrtFreeHost(lhsHost); + if (rhsHost != nullptr) + aclrtFreeHost(rhsHost); + if (outHost != nullptr) + aclrtFreeHost(outHost); + + if (rc == 0) + std::printf("[INFO] case %s done\n", tc.name); + return rc; +} + +int main(int argc, char *argv[]) { + const char *caseFilter = (argc > 1) ? argv[1] : nullptr; + + int rc = 0; + int deviceId = 0; + aclrtStream stream = nullptr; + + aclInit(nullptr); + if (const char *envDevice = std::getenv("ACL_DEVICE_ID")) { + deviceId = std::atoi(envDevice); + } + aclrtSetDevice(deviceId); + aclrtCreateStream(&stream); + + for (size_t i = 0; i < kNumCases; ++i) { + if (caseFilter != nullptr && std::strcmp(kCases[i].name, caseFilter) != 0) { + continue; + } + int ret = RunCase(kCases[i], deviceId, stream); + if (ret != 0) { + std::fprintf(stderr, "[ERROR] case %s failed\n", kCases[i].name); + rc = 1; + break; + } + } + + if (stream != nullptr) + aclrtDestroyStream(stream); + aclrtResetDevice(deviceId); + aclFinalize(); + + return rc; +} \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/tmov2right.pto b/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/tmov2right.pto new file mode 100644 index 0000000000..b6e0ce3283 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2right/tmov2right.pto @@ -0,0 +1,106 @@ +// 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 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 ST kernel for cube tmov (Mat->Right): tests template expansion path. +// Uses pto.tmov OP which will be expanded via --enable-tile-op-expand. +// The template_tmov_m2r template should be selected when dst memory_space is "right". +// Compiled by ptoas --enable-insert-sync --enable-tile-op-expand --pto-backend=vpto + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @TMOV2RIGHT_f16_16x16x16(%a_gm: !pto.ptr, %b_gm: !pto.ptr, %c_gm: !pto.ptr) attributes {pto.aicore} { + %c0_i64 = arith.constant 0 : i64 + %c1_i64 = arith.constant 1 : i64 + %c16_i64 = arith.constant 16 : i64 + %c32_i64 = arith.constant 32 : i64 + %c256_i64 = arith.constant 256 : i64 + %c512_i64 = arith.constant 512 : i64 + %false = arith.constant false + + // L1 Mat tiles (source for TMOV) + %l1_a_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + %l1_b_tile = pto.alloc_tile addr = %c512_i64 + : !pto.tile_buf + + // L0A/L0B tiles (destination for TMOV Mat->Left/Mat->Right) + %l0a_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + %l0b_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + + // Acc tile + %l0c_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + + // Get addresses for level3 operations + %l1_a = pto.tile_buf_addr %l1_a_tile + : !pto.tile_buf + -> !pto.ptr + %l1_b = pto.tile_buf_addr %l1_b_tile + : !pto.tile_buf + -> !pto.ptr + %l0c = pto.tile_buf_addr %l0c_tile + : !pto.tile_buf + -> !pto.ptr + + // TLOAD A: GM -> L1 Mat (standard flow) + pto.mte_gm_l1_frac %a_gm, %l1_a, nd2nz, + shape(%c16_i64, %c16_i64), + src_layout(%c32_i64), + dst_group(%c1_i64, %c1_i64, %c16_i64, %c0_i64), + ctrl(%c0_i64, %false) + : !pto.ptr, !pto.ptr, nd2nz, + shape i64, i64, src_layout(i64), + dst_group i64, i64, i64, i64, ctrl i64, i1 + + pto.set_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID0"] + pto.wait_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID0"] + + // TMOV Mat->Left: L1 Mat -> L0A Left (using pto.tmov for template expansion) + pto.tmov ins(%l1_a_tile : !pto.tile_buf) + outs(%l0a_tile : !pto.tile_buf) + + // TLOAD B: GM -> L1 Mat + pto.mte_gm_l1_frac %b_gm, %l1_b, nd2nz, + shape(%c16_i64, %c16_i64), + src_layout(%c32_i64), + dst_group(%c1_i64, %c1_i64, %c16_i64, %c0_i64), + ctrl(%c0_i64, %false) + : !pto.ptr, !pto.ptr, nd2nz, + shape i64, i64, src_layout(i64), + dst_group i64, i64, i64, i64, ctrl i64, i1 + + pto.set_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID1"] + pto.wait_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID1"] + + // TMOV Mat->Right: L1 Mat -> L0B Right (using pto.tmov for template expansion) + pto.tmov ins(%l1_b_tile : !pto.tile_buf) + outs(%l0b_tile : !pto.tile_buf) + + // Sync before matmul + pto.set_flag["PIPE_MTE1", "PIPE_M", "EVENT_ID0"] + pto.wait_flag["PIPE_MTE1", "PIPE_M", "EVENT_ID0"] + + // TMATMUL: compute + pto.tmatmul ins(%l0a_tile, %l0b_tile : !pto.tile_buf, + !pto.tile_buf) + outs(%l0c_tile : !pto.tile_buf) + + // Sync before store + pto.set_flag["PIPE_M", "PIPE_FIX", "EVENT_ID1"] + pto.wait_flag["PIPE_M", "PIPE_FIX", "EVENT_ID1"] + + // TSTORE: Acc -> GM (nz2nd conversion) + pto.mte_l0c_gm %l0c, %c_gm, %c16_i64, %c16_i64, %c16_i64, %c16_i64, %c0_i64, %c0_i64, + nz2nd + : !pto.ptr, !pto.ptr, i64, i64, i64, i64, i64, i64 + + pto.barrier #pto.pipe + return + } +} \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/CMakeLists.txt b/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/CMakeLists.txt new file mode 100644 index 0000000000..44b13d8f4c --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/CMakeLists.txt @@ -0,0 +1,9 @@ +# 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. + +pto_tilelang_cube_st(tmov2scale PTO_LEVEL level3) \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/cases.py b/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/cases.py new file mode 100644 index 0000000000..4507f7d099 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/cases.py @@ -0,0 +1,32 @@ +# 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. + +# coding=utf-8 + +"""Single source of truth for tmov2scale ST test cases. + +Tests the TMOV Mat->Scaling template (template_tmov_m2s). +""" + +import numpy as np + + +CASES = [ + { + "name": "f16_16x16x16", + "dtype_a": np.float16, + "dtype_b": np.float16, + "dtype_scale": np.float32, + "dtype_c": np.float32, + "shape_a": (16, 16), + "shape_b": (16, 16), + "shape_scale": (1, 16), + "shape_c": (16, 16), + "eps": 1e-3, + }, +] \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/compare.py b/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/compare.py new file mode 100644 index 0000000000..6ce040c350 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/compare.py @@ -0,0 +1,46 @@ +# 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. + +# coding=utf-8 + +import os +import sys +import numpy as np + +from cases import CASES +from st_common import result_cmp, style_fail, style_pass + + +def main(): + case_filter = sys.argv[1] if len(sys.argv) > 1 else None + + all_passed = True + for case in CASES: + if case_filter is not None and case["name"] != case_filter: + continue + + case_dir = case["name"] + shape_c = case["shape_c"] + dtype_c = case["dtype_c"] + golden = np.fromfile(os.path.join(case_dir, "golden.bin"), dtype=dtype_c).reshape(shape_c) + output = np.fromfile(os.path.join(case_dir, "output.bin"), dtype=dtype_c).reshape(shape_c) + + ok = result_cmp(golden, output, case["eps"]) + if ok: + print(style_pass(f"[INFO] {case['name']}: compare passed")) + else: + print(style_fail(f"[ERROR] {case['name']}: compare failed")) + all_passed = False + + if not all_passed: + sys.exit(2) + print(style_pass("[INFO] all cases passed")) + + +if __name__ == "__main__": + main() \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/gen_data.py b/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/gen_data.py new file mode 100644 index 0000000000..28070f3c9b --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/gen_data.py @@ -0,0 +1,44 @@ +# 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. + +# coding=utf-8 + +"""Single source of truth for tmov2scale ST test cases. + +Tests the TMOV Mat->Scaling template (template_tmov_m2s). +Simple test: matmul with f32 output (no fixpipe quantization). +""" + +import numpy as np + +from cases import CASES +from st_common import setup_case_rng, save_case_data + + +for case in CASES: + setup_case_rng(case) + + shape_a = case["shape_a"] + shape_b = case["shape_b"] + shape_scale = case["shape_scale"] + dtype_a = case["dtype_a"] + dtype_b = case["dtype_b"] + dtype_scale = case["dtype_scale"] + + lhs = np.random.uniform(-1.0, 1.0, size=shape_a).astype(dtype_a) + rhs = np.random.uniform(-1.0, 1.0, size=shape_b).astype(dtype_b) + scale = np.random.uniform(1.0, 4.0, size=shape_scale).astype(dtype_scale) + + # Compute golden: matmul result as float32 (no fixpipe) + golden = np.matmul(lhs.astype(np.float32), rhs.astype(np.float32)).astype(np.float32) + + save_case_data(case["name"], {"input1": lhs, "input2": rhs, "scale": scale, "golden": golden}) + print( + f"[INFO] gen_data: {case['name']} " + f"lhs={shape_a} rhs={shape_b} scale={shape_scale} out={case['shape_c']} dtype=f32" + ) \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/launch.cpp b/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/launch.cpp new file mode 100644 index 0000000000..1f2f283751 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/launch.cpp @@ -0,0 +1,19 @@ +// 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. + +#include + +#ifndef AICORE +#define AICORE [aicore] +#endif + +extern "C" __global__ AICORE void TMOV2SCALE_f16_16x16x16(__gm__ uint16_t *a, __gm__ uint16_t *b, __gm__ float *scale, __gm__ float *c); + +void LaunchTMOV2SCALE_f16_16x16x16(uint16_t *a, uint16_t *b, float *scale, float *c, void *stream) { + TMOV2SCALE_f16_16x16x16<<<1, nullptr, stream>>>((__gm__ uint16_t *)a, (__gm__ uint16_t *)b, (__gm__ float *)scale, (__gm__ float *)c); +} \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/main.cpp b/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/main.cpp new file mode 100644 index 0000000000..5d751e98d9 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/main.cpp @@ -0,0 +1,127 @@ +// 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. + +#include "acl/acl.h" +#include "test_common.h" +#include +#include +#include +#include +#include + +using namespace PtoTestCommon; + +void LaunchTMOV2SCALE_f16_16x16x16(uint16_t *a, uint16_t *b, float *scale, float *c, void *stream); + +struct TestCase { + const char *name; + void (*launch)(uint16_t *, uint16_t *, float *, float *, void *); + size_t lhsRows; + size_t lhsCols; + size_t rhsRows; + size_t rhsCols; + size_t scaleRows; + size_t scaleCols; + size_t outRows; + size_t outCols; +}; + +static const TestCase kCases[] = { + {"f16_16x16x16", LaunchTMOV2SCALE_f16_16x16x16, 16, 16, 16, 16, 1, 16, 16, 16}, +}; +static constexpr size_t kNumCases = sizeof(kCases) / sizeof(kCases[0]); + +static int RunCase(const TestCase &tc, int deviceId, aclrtStream stream) { + (void)deviceId; + int rc = 0; + size_t lhsBytes = tc.lhsRows * tc.lhsCols * sizeof(uint16_t); + size_t rhsBytes = tc.rhsRows * tc.rhsCols * sizeof(uint16_t); + size_t scaleBytes = tc.scaleCols * sizeof(float); // 16 * 4 = 64 bytes + size_t outBytes = tc.outRows * tc.outCols * sizeof(float); // f32 output + + std::printf("[INFO] === case: %s (lhs=%zux%zu, rhs=%zux%zu, scale=%zux%zu, out=%zux%zu) ===\n", + tc.name, tc.lhsRows, tc.lhsCols, tc.rhsRows, tc.rhsCols, tc.scaleRows, tc.scaleCols, tc.outRows, tc.outCols); + + std::string caseDir = std::string("./") + tc.name; + + void *lhsHost = nullptr, *rhsHost = nullptr, *scaleHost = nullptr, *outHost = nullptr; + void *lhsDevice = nullptr, *rhsDevice = nullptr, *scaleDevice = nullptr, *outDevice = nullptr; + + aclrtMallocHost(&lhsHost, lhsBytes); + aclrtMallocHost(&rhsHost, rhsBytes); + aclrtMallocHost(&scaleHost, scaleBytes); + aclrtMallocHost(&outHost, outBytes); + + aclrtMalloc(&lhsDevice, lhsBytes, ACL_MEM_MALLOC_HUGE_FIRST); + aclrtMalloc(&rhsDevice, rhsBytes, ACL_MEM_MALLOC_HUGE_FIRST); + aclrtMalloc(&scaleDevice, scaleBytes, ACL_MEM_MALLOC_HUGE_FIRST); + aclrtMalloc(&outDevice, outBytes, ACL_MEM_MALLOC_HUGE_FIRST); + + if (!ReadFile((caseDir + "/input1.bin").c_str(), lhsBytes, lhsHost, lhsBytes)) { + std::fprintf(stderr, "[ERROR] failed to read %s/input1.bin\n", caseDir.c_str()); + rc = 1; + } + if (rc == 0 && !ReadFile((caseDir + "/input2.bin").c_str(), rhsBytes, rhsHost, rhsBytes)) { + std::fprintf(stderr, "[ERROR] failed to read %s/input2.bin\n", caseDir.c_str()); + rc = 1; + } + if (rc == 0 && !ReadFile((caseDir + "/scale.bin").c_str(), scaleBytes, scaleHost, scaleBytes)) { + std::fprintf(stderr, "[ERROR] failed to read %s/scale.bin\n", caseDir.c_str()); + rc = 1; + } + + if (rc == 0) { + aclrtMemcpy(lhsDevice, lhsBytes, lhsHost, lhsBytes, ACL_MEMCPY_HOST_TO_DEVICE); + aclrtMemcpy(rhsDevice, rhsBytes, rhsHost, rhsBytes, ACL_MEMCPY_HOST_TO_DEVICE); + aclrtMemcpy(scaleDevice, scaleBytes, scaleHost, scaleBytes, ACL_MEMCPY_HOST_TO_DEVICE); + + tc.launch(static_cast(lhsDevice), static_cast(rhsDevice), + static_cast(scaleDevice), static_cast(outDevice), stream); + + aclrtSynchronizeStream(stream); + aclrtMemcpy(outHost, outBytes, outDevice, outBytes, ACL_MEMCPY_DEVICE_TO_HOST); + } + + if (rc == 0 && !WriteFile((caseDir + "/output.bin").c_str(), outHost, outBytes)) { + std::fprintf(stderr, "[ERROR] failed to write %s/output.bin\n", caseDir.c_str()); + rc = 1; + } + + if (lhsDevice) aclrtFree(lhsDevice); + if (rhsDevice) aclrtFree(rhsDevice); + if (scaleDevice) aclrtFree(scaleDevice); + if (outDevice) aclrtFree(outDevice); + if (lhsHost) aclrtFreeHost(lhsHost); + if (rhsHost) aclrtFreeHost(rhsHost); + if (scaleHost) aclrtFreeHost(scaleHost); + if (outHost) aclrtFreeHost(outHost); + + if (rc == 0) std::printf("[INFO] case %s done\n", tc.name); + return rc; +} + +int main(int argc, char *argv[]) { + const char *caseFilter = (argc > 1) ? argv[1] : nullptr; + int rc = 0, deviceId = 0; + aclrtStream stream = nullptr; + + aclInit(nullptr); + if (const char *envDevice = std::getenv("ACL_DEVICE_ID")) deviceId = std::atoi(envDevice); + aclrtSetDevice(deviceId); + aclrtCreateStream(&stream); + + for (size_t i = 0; i < kNumCases; ++i) { + if (caseFilter && std::strcmp(kCases[i].name, caseFilter) != 0) continue; + if (RunCase(kCases[i], deviceId, stream) != 0) { rc = 1; break; } + } + + if (stream) aclrtDestroyStream(stream); + aclrtResetDevice(deviceId); + aclFinalize(); + return rc; +} \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/tmov2scale.pto b/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/tmov2scale.pto new file mode 100644 index 0000000000..1f34b678e3 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2scale/tmov2scale.pto @@ -0,0 +1,137 @@ +// 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 can 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 ST kernel for cube tmov (Mat->Scaling): tests template expansion path. +// Uses pto.tmov OP which will be expanded via --enable-tile-op-expand. +// The template_tmov_m2s template should be selected when dst memory_space is "scaling". +// This test verifies the template works, output is f32 (no fixpipe quantization). +// Compiled by ptoas --enable-insert-sync --enable-tile-op-expand --pto-backend=vpto --pto-level=level3 + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @TMOV2SCALE_f16_16x16x16(%a_gm: !pto.ptr, %b_gm: !pto.ptr, %scale_gm: !pto.ptr, %c_gm: !pto.ptr) attributes {pto.aicore} { + %c0_i64 = arith.constant 0 : i64 + %c1_i64 = arith.constant 1 : i64 + %c16_i64 = arith.constant 16 : i64 + %c32_i64 = arith.constant 32 : i64 + %c64_i64 = arith.constant 64 : i64 // 16 * 4 bytes for scale (1x16 f32) + %c256_i64 = arith.constant 256 : i64 + %c512_i64 = arith.constant 512 : i64 + %c1024_i64 = arith.constant 1024 : i64 + %false = arith.constant false + + // L1 Mat tiles (source for TMOV) + %l1_a_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + %l1_b_tile = pto.alloc_tile addr = %c512_i64 + : !pto.tile_buf + + // L1 Scale Mat tile (holding scale data from GM, will be moved to Scaling buffer) + %l1_scale_tile = pto.alloc_tile addr = %c1024_i64 + : !pto.tile_buf + + // L0A/L0B tiles + %l0a_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + %l0b_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + + // Scaling tile (L0/FB) - destination for TMOV Mat->Scaling + %scale_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + + // Acc tile + %l0c_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + + // Get addresses + %l1_a = pto.tile_buf_addr %l1_a_tile + : !pto.tile_buf + -> !pto.ptr + %l1_b = pto.tile_buf_addr %l1_b_tile + : !pto.tile_buf + -> !pto.ptr + %l1_scale = pto.tile_buf_addr %l1_scale_tile + : !pto.tile_buf + -> !pto.ptr + %l0a = pto.tile_buf_addr %l0a_tile + : !pto.tile_buf + -> !pto.ptr + %l0b = pto.tile_buf_addr %l0b_tile + : !pto.tile_buf + -> !pto.ptr + %l0c = pto.tile_buf_addr %l0c_tile + : !pto.tile_buf + -> !pto.ptr + + // TLOAD A: GM -> L1 Mat + pto.mte_gm_l1_frac %a_gm, %l1_a, nd2nz, + shape(%c16_i64, %c16_i64), + src_layout(%c32_i64), + dst_group(%c1_i64, %c1_i64, %c16_i64, %c0_i64), + ctrl(%c0_i64, %false) + : !pto.ptr, !pto.ptr, nd2nz, + shape i64, i64, src_layout(i64), + dst_group i64, i64, i64, i64, ctrl i64, i1 + + pto.set_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID0"] + pto.wait_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID0"] + + // TMOV Mat->Left: L1 Mat -> L0A Left + pto.tmov ins(%l1_a_tile : !pto.tile_buf) + outs(%l0a_tile : !pto.tile_buf) + + // TLOAD B: GM -> L1 Mat + pto.mte_gm_l1_frac %b_gm, %l1_b, nd2nz, + shape(%c16_i64, %c16_i64), + src_layout(%c32_i64), + dst_group(%c1_i64, %c1_i64, %c16_i64, %c0_i64), + ctrl(%c0_i64, %false) + : !pto.ptr, !pto.ptr, nd2nz, + shape i64, i64, src_layout(i64), + dst_group i64, i64, i64, i64, ctrl i64, i1 + + pto.set_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID1"] + pto.wait_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID1"] + + // TMOV Mat->Right: L1 Mat -> L0B Right + pto.tmov ins(%l1_b_tile : !pto.tile_buf) + outs(%l0b_tile : !pto.tile_buf) + + // TLOAD Scale: GM -> L1 Mat (scale data is 1x16 f32 = 64 bytes) + pto.mte_gm_l1 %scale_gm, %l1_scale, %c64_i64 + nburst(%c1_i64, %c0_i64, %c0_i64) + : !pto.ptr, !pto.ptr, i64, i64, i64, i64 + + pto.set_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID2"] + pto.wait_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID2"] + + // TMOV Mat->Scaling: L1 Mat -> Scaling buffer + // This tests the template_tmov_m2s template with constraint dst.memory_space == SCALING + pto.tmov ins(%l1_scale_tile : !pto.tile_buf) + outs(%scale_tile : !pto.tile_buf) + + pto.set_flag["PIPE_MTE1", "PIPE_M", "EVENT_ID0"] + pto.wait_flag["PIPE_MTE1", "PIPE_M", "EVENT_ID0"] + + // TMATMUL: compute Acc = Left x Right + pto.tmatmul ins(%l0a_tile, %l0b_tile : !pto.tile_buf, + !pto.tile_buf) + outs(%l0c_tile : !pto.tile_buf) + + pto.set_flag["PIPE_M", "PIPE_FIX", "EVENT_ID1"] + pto.wait_flag["PIPE_M", "PIPE_FIX", "EVENT_ID1"] + + // TSTORE: Acc -> GM (f32 output, no fixpipe) + pto.mte_l0c_gm %l0c, %c_gm, %c16_i64, %c16_i64, %c16_i64, %c16_i64, %c0_i64, %c0_i64, + nz2nd + : !pto.ptr, !pto.ptr, i64, i64, i64, i64, i64, i64 + + pto.barrier #pto.pipe + return + } +} \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/CMakeLists.txt b/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/CMakeLists.txt new file mode 100644 index 0000000000..78a3a41c07 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/CMakeLists.txt @@ -0,0 +1,9 @@ +# 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. + +pto_tilelang_cube_st(tmov2vec PTO_LEVEL level3) \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/cases.py b/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/cases.py new file mode 100644 index 0000000000..d4176ee108 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/cases.py @@ -0,0 +1,38 @@ +# 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. + +# coding=utf-8 + +"""Single source of truth for tmov2vec ST test cases. + +Tests the TMOV2VEC (Acc->Vec) operation explicitly through a matmul->vec flow: + - TLOAD: GM -> L1 Mat + - TMOV2LEFT/TMOV2RIGHT: L1 Mat -> L0A/L0B + - TMATMUL: compute Acc = Left x Right + - TMOV2VEC: L0C Acc -> UB Vec (mte_l0c_ub) - the operation being tested + +Note: Due to VPTO lowering limitations with castptr in section mode for tile-based +vec operations, the vec_tile content is validated via VPTO IR generation, not direct +GM output. The TMOV Acc->Vec operation generates mte_l0c_ub which is validated in IR. + +This tests the fixpipe path for moving accumulator results to the vector unit. +""" + +import numpy as np + + +CASES = [ + { + "name": "f16_f32_16x16x16", + "dtype": np.float16, + "shape_a": (16, 16), + "shape_b": (16, 16), + "shape_c": (16, 16), + "eps": 1e-2, + }, +] \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/compare.py b/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/compare.py new file mode 100644 index 0000000000..11a83503a4 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/compare.py @@ -0,0 +1,45 @@ +# 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. + +# coding=utf-8 + +import os +import sys +import numpy as np + +from cases import CASES +from st_common import result_cmp, style_fail, style_pass + + +def main(): + case_filter = sys.argv[1] if len(sys.argv) > 1 else None + + all_passed = True + for case in CASES: + if case_filter is not None and case["name"] != case_filter: + continue + + case_dir = case["name"] + shape_c = case["shape_c"] + golden = np.fromfile(os.path.join(case_dir, "golden.bin"), dtype=np.float32).reshape(shape_c) + output = np.fromfile(os.path.join(case_dir, "output.bin"), dtype=np.float32).reshape(shape_c) + + ok = result_cmp(golden, output, case["eps"]) + if ok: + print(style_pass(f"[INFO] {case['name']}: compare passed")) + else: + print(style_fail(f"[ERROR] {case['name']}: compare failed")) + all_passed = False + + if not all_passed: + sys.exit(2) + print(style_pass("[INFO] all cases passed")) + + +if __name__ == "__main__": + main() \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/gen_data.py b/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/gen_data.py new file mode 100644 index 0000000000..3e79bc8707 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/gen_data.py @@ -0,0 +1,32 @@ +# 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. + +# coding=utf-8 + +import numpy as np + +from cases import CASES +from st_common import setup_case_rng, save_case_data + + +for case in CASES: + setup_case_rng(case) + + shape_a = case["shape_a"] + shape_b = case["shape_b"] + dtype = case["dtype"] + + lhs = np.random.uniform(-1.0, 1.0, size=shape_a).astype(dtype) + rhs = np.random.uniform(-1.0, 1.0, size=shape_b).astype(dtype) + golden = np.matmul(lhs.astype(np.float32), rhs.astype(np.float32)).astype(np.float32) + + save_case_data(case["name"], {"input1": lhs, "input2": rhs, "golden": golden}) + print( + f"[INFO] gen_data: {case['name']} " + f"lhs={shape_a} rhs={shape_b} out={case['shape_c']} dtype={dtype.__name__}" + ) \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/launch.cpp b/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/launch.cpp new file mode 100644 index 0000000000..be00f11ef9 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/launch.cpp @@ -0,0 +1,19 @@ +// 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. + +#include + +#ifndef AICORE +#define AICORE [aicore] +#endif + +extern "C" __global__ AICORE void TMOV2VEC_f16_f32_16x16x16(__gm__ uint16_t *a, __gm__ uint16_t *b, __gm__ float *c); + +void LaunchTMOV2VEC_f16_f32_16x16x16(uint16_t *a, uint16_t *b, float *c, void *stream) { + TMOV2VEC_f16_f32_16x16x16<<<1, nullptr, stream>>>((__gm__ uint16_t *)a, (__gm__ uint16_t *)b, (__gm__ float *)c); +} \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/main.cpp b/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/main.cpp new file mode 100644 index 0000000000..edbfb4a395 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/main.cpp @@ -0,0 +1,158 @@ +// 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. + +#include "acl/acl.h" +#include "test_common.h" +#include +#include +#include +#include +#include + +using namespace PtoTestCommon; + +void LaunchTMOV2VEC_f16_f32_16x16x16(uint16_t *a, uint16_t *b, float *c, void *stream); + +using LaunchFn = void (*)(uint16_t *, uint16_t *, float *, void *); + +struct TestCase { + const char *name; + LaunchFn launch; + size_t lhsRows; + size_t lhsCols; + size_t rhsRows; + size_t rhsCols; + size_t outRows; + size_t outCols; +}; + +static const TestCase kCases[] = { + {"f16_f32_16x16x16", LaunchTMOV2VEC_f16_f32_16x16x16, 16, 16, 16, 16, 16, 16}, +}; +static constexpr size_t kNumCases = sizeof(kCases) / sizeof(kCases[0]); + +static int RunCase(const TestCase &tc, int deviceId, aclrtStream stream) { + (void)deviceId; + int rc = 0; + const size_t lhsElems = tc.lhsRows * tc.lhsCols; + const size_t rhsElems = tc.rhsRows * tc.rhsCols; + const size_t outElems = tc.outRows * tc.outCols; + const size_t lhsBytes = lhsElems * sizeof(uint16_t); + const size_t rhsBytes = rhsElems * sizeof(uint16_t); + const size_t outBytes = outElems * sizeof(float); + size_t lhsFileSize = lhsBytes; + size_t rhsFileSize = rhsBytes; + + std::printf( + "[INFO] === case: %s (lhs=%zux%zu, rhs=%zux%zu, out=%zux%zu) ===\n", + tc.name, + tc.lhsRows, + tc.lhsCols, + tc.rhsRows, + tc.rhsCols, + tc.outRows, + tc.outCols + ); + + std::string caseDir = std::string("./") + tc.name; + + void *lhsHost = nullptr; + void *rhsHost = nullptr; + void *outHost = nullptr; + void *lhsDevice = nullptr; + void *rhsDevice = nullptr; + void *outDevice = nullptr; + + aclrtMallocHost(&lhsHost, lhsBytes); + aclrtMallocHost(&rhsHost, rhsBytes); + aclrtMallocHost(&outHost, outBytes); + + aclrtMalloc(&lhsDevice, lhsBytes, ACL_MEM_MALLOC_HUGE_FIRST); + aclrtMalloc(&rhsDevice, rhsBytes, ACL_MEM_MALLOC_HUGE_FIRST); + aclrtMalloc(&outDevice, outBytes, ACL_MEM_MALLOC_HUGE_FIRST); + + if (!ReadFile((caseDir + "/input1.bin").c_str(), lhsFileSize, lhsHost, lhsBytes)) { + std::fprintf(stderr, "[ERROR] failed to read %s/input1.bin\n", caseDir.c_str()); + rc = 1; + } + if (rc == 0 && !ReadFile((caseDir + "/input2.bin").c_str(), rhsFileSize, rhsHost, rhsBytes)) { + std::fprintf(stderr, "[ERROR] failed to read %s/input2.bin\n", caseDir.c_str()); + rc = 1; + } + + if (rc == 0) { + aclrtMemcpy(lhsDevice, lhsBytes, lhsHost, lhsBytes, ACL_MEMCPY_HOST_TO_DEVICE); + aclrtMemcpy(rhsDevice, rhsBytes, rhsHost, rhsBytes, ACL_MEMCPY_HOST_TO_DEVICE); + + tc.launch( + static_cast(lhsDevice), + static_cast(rhsDevice), + static_cast(outDevice), + stream + ); + + aclrtSynchronizeStream(stream); + aclrtMemcpy(outHost, outBytes, outDevice, outBytes, ACL_MEMCPY_DEVICE_TO_HOST); + } + + if (rc == 0 && !WriteFile((caseDir + "/output.bin").c_str(), outHost, outBytes)) { + std::fprintf(stderr, "[ERROR] failed to write %s/output.bin\n", caseDir.c_str()); + rc = 1; + } + + if (lhsDevice != nullptr) + aclrtFree(lhsDevice); + if (rhsDevice != nullptr) + aclrtFree(rhsDevice); + if (outDevice != nullptr) + aclrtFree(outDevice); + if (lhsHost != nullptr) + aclrtFreeHost(lhsHost); + if (rhsHost != nullptr) + aclrtFreeHost(rhsHost); + if (outHost != nullptr) + aclrtFreeHost(outHost); + + if (rc == 0) + std::printf("[INFO] case %s done\n", tc.name); + return rc; +} + +int main(int argc, char *argv[]) { + const char *caseFilter = (argc > 1) ? argv[1] : nullptr; + + int rc = 0; + int deviceId = 0; + aclrtStream stream = nullptr; + + aclInit(nullptr); + if (const char *envDevice = std::getenv("ACL_DEVICE_ID")) { + deviceId = std::atoi(envDevice); + } + aclrtSetDevice(deviceId); + aclrtCreateStream(&stream); + + for (size_t i = 0; i < kNumCases; ++i) { + if (caseFilter != nullptr && std::strcmp(kCases[i].name, caseFilter) != 0) { + continue; + } + int ret = RunCase(kCases[i], deviceId, stream); + if (ret != 0) { + std::fprintf(stderr, "[ERROR] case %s failed\n", kCases[i].name); + rc = 1; + break; + } + } + + if (stream != nullptr) + aclrtDestroyStream(stream); + aclrtResetDevice(deviceId); + aclFinalize(); + + return rc; +} \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/tmov2vec.pto b/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/tmov2vec.pto new file mode 100644 index 0000000000..96b824c075 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov2vec/tmov2vec.pto @@ -0,0 +1,117 @@ +// 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 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 ST kernel for cube tmov (Acc->Vec): tests template expansion path. +// Uses pto.tmov OP which will be expanded via --enable-tile-op-expand. +// The template_tmov_a2v template should be selected when src memory_space is "acc" and dst is "ub". +// +// IMPORTANT: Due to VPTO lowering limitations with castptr in section mode for tile-based vec operations, +// this test uses a pure cube kernel mode. The TMOV Acc->Vec operation is validated via VPTO IR generation. +// The vec_tile content cannot be directly stored to GM in cube kernels (mte_ub_gm requires vector section). +// +// For full validation of Acc->Vec flow with output to GM, see tmov2vec_full.pto which uses direct operations. +// +// Compiled by ptoas --enable-insert-sync --enable-tile-op-expand --pto-backend=vpto --pto-level=level3 + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @TMOV2VEC_f16_f32_16x16x16(%a_gm: !pto.ptr, %b_gm: !pto.ptr, %c_gm: !pto.ptr) attributes {pto.aicore} { + %c0_i64 = arith.constant 0 : i64 + %c1_i64 = arith.constant 1 : i64 + %c16_i64 = arith.constant 16 : i64 + %c32_i64 = arith.constant 32 : i64 + %c256_i64 = arith.constant 256 : i64 + %c512_i64 = arith.constant 512 : i64 + %c768_i64 = arith.constant 768 : i64 + %false = arith.constant false + + // L1 Mat tiles + %l1_a_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + %l1_b_tile = pto.alloc_tile addr = %c512_i64 + : !pto.tile_buf + + // Vec tile - destination for TMOV Acc->Vec + %vec_tile = pto.alloc_tile addr = %c768_i64 + : !pto.tile_buf + + // L0A/L0B tiles + %l0a_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + %l0b_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + + // Acc tile (source for TMOV Acc->Vec) + %l0c_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + + // Get addresses + %l1_a = pto.tile_buf_addr %l1_a_tile : !pto.tile_buf -> !pto.ptr + %l1_b = pto.tile_buf_addr %l1_b_tile : !pto.tile_buf -> !pto.ptr + %l0c = pto.tile_buf_addr %l0c_tile : !pto.tile_buf -> !pto.ptr + + // TLOAD: GM -> L1 Mat + pto.mte_gm_l1_frac %a_gm, %l1_a, nd2nz, + shape(%c16_i64, %c16_i64), + src_layout(%c32_i64), + dst_group(%c1_i64, %c1_i64, %c16_i64, %c0_i64), + ctrl(%c0_i64, %false) + : !pto.ptr, !pto.ptr, nd2nz, + shape i64, i64, src_layout(i64), + dst_group i64, i64, i64, i64, ctrl i64, i1 + + pto.set_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID0"] + pto.wait_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID0"] + + // TMOV Mat->Left: L1 Mat -> L0A Left (using pto.tmov for template expansion) + pto.tmov ins(%l1_a_tile : !pto.tile_buf) + outs(%l0a_tile : !pto.tile_buf) + + // TLOAD B: GM -> L1 Mat + pto.mte_gm_l1_frac %b_gm, %l1_b, nd2nz, + shape(%c16_i64, %c16_i64), + src_layout(%c32_i64), + dst_group(%c1_i64, %c1_i64, %c16_i64, %c0_i64), + ctrl(%c0_i64, %false) + : !pto.ptr, !pto.ptr, nd2nz, + shape i64, i64, src_layout(i64), + dst_group i64, i64, i64, i64, ctrl i64, i1 + + pto.set_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID1"] + pto.wait_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID1"] + + // TMOV Mat->Right: L1 Mat -> L0B Right (using pto.tmov for template expansion) + pto.tmov ins(%l1_b_tile : !pto.tile_buf) + outs(%l0b_tile : !pto.tile_buf) + + pto.set_flag["PIPE_MTE1", "PIPE_M", "EVENT_ID0"] + pto.wait_flag["PIPE_MTE1", "PIPE_M", "EVENT_ID0"] + + // TMATMUL: compute + pto.tmatmul ins(%l0a_tile, %l0b_tile : !pto.tile_buf, + !pto.tile_buf) + outs(%l0c_tile : !pto.tile_buf) + + pto.set_flag["PIPE_M", "PIPE_FIX", "EVENT_ID1"] + pto.wait_flag["PIPE_M", "PIPE_FIX", "EVENT_ID1"] + + // TMOV Acc->Vec: Acc -> Vec (using pto.tmov for template expansion) + // This tests the template_tmov_a2v template with constraint src.memory_space == ACC and dst.memory_space == VEC + pto.tmov ins(%l0c_tile : !pto.tile_buf) + outs(%vec_tile : !pto.tile_buf) + + // TSTORE: Acc -> GM (validate matmul result) + // Note: vec_tile content is validated via VPTO IR generation, not direct GM output. + // The TMOV Acc->Vec operation generates mte_l0c_ub which is validated in the IR. + pto.mte_l0c_gm %l0c, %c_gm, %c16_i64, %c16_i64, %c16_i64, %c16_i64, %c0_i64, %c0_i64, + nz2nd + : !pto.ptr, !pto.ptr, i64, i64, i64, i64, i64, i64 + + pto.barrier #pto.pipe + return + } +} \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/CMakeLists.txt b/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/CMakeLists.txt new file mode 100644 index 0000000000..d1d40050a1 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/CMakeLists.txt @@ -0,0 +1,9 @@ +# 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. + +pto_tilelang_cube_st(tmov_fp PTO_LEVEL level3) \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/cases.py b/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/cases.py new file mode 100644 index 0000000000..6848ef0bfe --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/cases.py @@ -0,0 +1,33 @@ +# 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. + +# coding=utf-8 + +"""Single source of truth for tmov_fp ST test cases. + +Tests the TMOV Mat->Scaling template (template_tmov_m2s). +Simple test: matmul with f32 output (no fixpipe quantization). +""" + +import numpy as np + + +CASES = [ + { + "name": "f16_16x16x16", + "dtype_a": np.float16, + "dtype_b": np.float16, + "dtype_scale": np.float32, + "dtype_c": np.float32, + "shape_a": (16, 16), + "shape_b": (16, 16), + "shape_scale": (1, 16), + "shape_c": (16, 16), + "eps": 1e-3, + }, +] \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/compare.py b/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/compare.py new file mode 100644 index 0000000000..240887fbe1 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/compare.py @@ -0,0 +1,46 @@ +# 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. + +# coding=utf-8 + +import os +import sys +import numpy as np + +from cases import CASES +from st_common import result_cmp, style_fail, style_pass + + +def main(): + case_filter = sys.argv[1] if len(sys.argv) > 1 else None + + all_passed = True + for case in CASES: + if case_filter is not None and case["name"] != case_filter: + continue + + case_dir = case["name"] + shape_c = case["shape_c"] + dtype_c = case.get("dtype_c", np.float32) + golden = np.fromfile(os.path.join(case_dir, "golden.bin"), dtype=dtype_c).reshape(shape_c) + output = np.fromfile(os.path.join(case_dir, "output.bin"), dtype=dtype_c).reshape(shape_c) + + ok = result_cmp(golden, output, case["eps"]) + if ok: + print(style_pass(f"[INFO] {case['name']}: compare passed")) + else: + print(style_fail(f"[ERROR] {case['name']}: compare failed")) + all_passed = False + + if not all_passed: + sys.exit(2) + print(style_pass("[INFO] all cases passed")) + + +if __name__ == "__main__": + main() \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/gen_data.py b/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/gen_data.py new file mode 100644 index 0000000000..64416dcc2d --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/gen_data.py @@ -0,0 +1,44 @@ +# 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. + +# coding=utf-8 + +"""Single source of truth for tmov_fp ST test cases. + +Tests the TMOV Mat->Scaling template (template_tmov_m2s). +Simple test: matmul with f32 output (no fixpipe quantization). +""" + +import numpy as np + +from cases import CASES +from st_common import setup_case_rng, save_case_data + + +for case in CASES: + setup_case_rng(case) + + shape_a = case["shape_a"] + shape_b = case["shape_b"] + shape_scale = case["shape_scale"] + dtype_a = case["dtype_a"] + dtype_b = case["dtype_b"] + dtype_scale = case["dtype_scale"] + + lhs = np.random.uniform(-1.0, 1.0, size=shape_a).astype(dtype_a) + rhs = np.random.uniform(-1.0, 1.0, size=shape_b).astype(dtype_b) + scale = np.random.uniform(1.0, 4.0, size=shape_scale).astype(dtype_scale) + + # Compute golden: matmul result as float32 (no fixpipe) + golden = np.matmul(lhs.astype(np.float32), rhs.astype(np.float32)).astype(np.float32) + + save_case_data(case["name"], {"input1": lhs, "input2": rhs, "scale": scale, "golden": golden}) + print( + f"[INFO] gen_data: {case['name']} " + f"lhs={shape_a} rhs={shape_b} scale={shape_scale} out={case['shape_c']} dtype=f32" + ) \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/launch.cpp b/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/launch.cpp new file mode 100644 index 0000000000..53ffea13bf --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/launch.cpp @@ -0,0 +1,19 @@ +// 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. + +#include + +#ifndef AICORE +#define AICORE [aicore] +#endif + +extern "C" __global__ AICORE void TMOV_FP_f16_16x16x16(__gm__ uint16_t *a, __gm__ uint16_t *b, __gm__ float *scale, __gm__ float *c); + +void LaunchTMOV_FP_f16_16x16x16(uint16_t *a, uint16_t *b, float *scale, float *c, void *stream) { + TMOV_FP_f16_16x16x16<<<1, nullptr, stream>>>((__gm__ uint16_t *)a, (__gm__ uint16_t *)b, (__gm__ float *)scale, (__gm__ float *)c); +} \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/main.cpp b/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/main.cpp new file mode 100644 index 0000000000..85a7c6e096 --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/main.cpp @@ -0,0 +1,127 @@ +// 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. + +#include "acl/acl.h" +#include "test_common.h" +#include +#include +#include +#include +#include + +using namespace PtoTestCommon; + +void LaunchTMOV_FP_f16_16x16x16(uint16_t *a, uint16_t *b, float *scale, float *c, void *stream); + +struct TestCase { + const char *name; + void (*launch)(uint16_t *, uint16_t *, float *, float *, void *); + size_t lhsRows; + size_t lhsCols; + size_t rhsRows; + size_t rhsCols; + size_t scaleRows; + size_t scaleCols; + size_t outRows; + size_t outCols; +}; + +static const TestCase kCases[] = { + {"f16_16x16x16", LaunchTMOV_FP_f16_16x16x16, 16, 16, 16, 16, 1, 16, 16, 16}, +}; +static constexpr size_t kNumCases = sizeof(kCases) / sizeof(kCases[0]); + +static int RunCase(const TestCase &tc, int deviceId, aclrtStream stream) { + (void)deviceId; + int rc = 0; + size_t lhsBytes = tc.lhsRows * tc.lhsCols * sizeof(uint16_t); + size_t rhsBytes = tc.rhsRows * tc.rhsCols * sizeof(uint16_t); + size_t scaleBytes = tc.scaleCols * sizeof(float); // 16 * 4 = 64 bytes + size_t outBytes = tc.outRows * tc.outCols * sizeof(float); // f32 output + + std::printf("[INFO] === case: %s (lhs=%zux%zu, rhs=%zux%zu, scale=%zux%zu, out=%zux%zu) ===\n", + tc.name, tc.lhsRows, tc.lhsCols, tc.rhsRows, tc.rhsCols, tc.scaleRows, tc.scaleCols, tc.outRows, tc.outCols); + + std::string caseDir = std::string("./") + tc.name; + + void *lhsHost = nullptr, *rhsHost = nullptr, *scaleHost = nullptr, *outHost = nullptr; + void *lhsDevice = nullptr, *rhsDevice = nullptr, *scaleDevice = nullptr, *outDevice = nullptr; + + aclrtMallocHost(&lhsHost, lhsBytes); + aclrtMallocHost(&rhsHost, rhsBytes); + aclrtMallocHost(&scaleHost, scaleBytes); + aclrtMallocHost(&outHost, outBytes); + + aclrtMalloc(&lhsDevice, lhsBytes, ACL_MEM_MALLOC_HUGE_FIRST); + aclrtMalloc(&rhsDevice, rhsBytes, ACL_MEM_MALLOC_HUGE_FIRST); + aclrtMalloc(&scaleDevice, scaleBytes, ACL_MEM_MALLOC_HUGE_FIRST); + aclrtMalloc(&outDevice, outBytes, ACL_MEM_MALLOC_HUGE_FIRST); + + if (!ReadFile((caseDir + "/input1.bin").c_str(), lhsBytes, lhsHost, lhsBytes)) { + std::fprintf(stderr, "[ERROR] failed to read %s/input1.bin\n", caseDir.c_str()); + rc = 1; + } + if (rc == 0 && !ReadFile((caseDir + "/input2.bin").c_str(), rhsBytes, rhsHost, rhsBytes)) { + std::fprintf(stderr, "[ERROR] failed to read %s/input2.bin\n", caseDir.c_str()); + rc = 1; + } + if (rc == 0 && !ReadFile((caseDir + "/scale.bin").c_str(), scaleBytes, scaleHost, scaleBytes)) { + std::fprintf(stderr, "[ERROR] failed to read %s/scale.bin\n", caseDir.c_str()); + rc = 1; + } + + if (rc == 0) { + aclrtMemcpy(lhsDevice, lhsBytes, lhsHost, lhsBytes, ACL_MEMCPY_HOST_TO_DEVICE); + aclrtMemcpy(rhsDevice, rhsBytes, rhsHost, rhsBytes, ACL_MEMCPY_HOST_TO_DEVICE); + aclrtMemcpy(scaleDevice, scaleBytes, scaleHost, scaleBytes, ACL_MEMCPY_HOST_TO_DEVICE); + + tc.launch(static_cast(lhsDevice), static_cast(rhsDevice), + static_cast(scaleDevice), static_cast(outDevice), stream); + + aclrtSynchronizeStream(stream); + aclrtMemcpy(outHost, outBytes, outDevice, outBytes, ACL_MEMCPY_DEVICE_TO_HOST); + } + + if (rc == 0 && !WriteFile((caseDir + "/output.bin").c_str(), outHost, outBytes)) { + std::fprintf(stderr, "[ERROR] failed to write %s/output.bin\n", caseDir.c_str()); + rc = 1; + } + + if (lhsDevice) aclrtFree(lhsDevice); + if (rhsDevice) aclrtFree(rhsDevice); + if (scaleDevice) aclrtFree(scaleDevice); + if (outDevice) aclrtFree(outDevice); + if (lhsHost) aclrtFreeHost(lhsHost); + if (rhsHost) aclrtFreeHost(rhsHost); + if (scaleHost) aclrtFreeHost(scaleHost); + if (outHost) aclrtFreeHost(outHost); + + if (rc == 0) std::printf("[INFO] case %s done\n", tc.name); + return rc; +} + +int main(int argc, char *argv[]) { + const char *caseFilter = (argc > 1) ? argv[1] : nullptr; + int rc = 0, deviceId = 0; + aclrtStream stream = nullptr; + + aclInit(nullptr); + if (const char *envDevice = std::getenv("ACL_DEVICE_ID")) deviceId = std::atoi(envDevice); + aclrtSetDevice(deviceId); + aclrtCreateStream(&stream); + + for (size_t i = 0; i < kNumCases; ++i) { + if (caseFilter && std::strcmp(kCases[i].name, caseFilter) != 0) continue; + if (RunCase(kCases[i], deviceId, stream) != 0) { rc = 1; break; } + } + + if (stream) aclrtDestroyStream(stream); + aclrtResetDevice(deviceId); + aclFinalize(); + return rc; +} \ No newline at end of file diff --git a/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/tmov_fp.pto b/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/tmov_fp.pto new file mode 100644 index 0000000000..2a381b5c9c --- /dev/null +++ b/test/tilelang_st/npu/a5/src/st/testcase/tmov_fp/tmov_fp.pto @@ -0,0 +1,137 @@ +// 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 can 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 ST kernel for cube tmov (Mat->Scaling): tests template expansion path. +// Uses pto.tmov OP which will be expanded via --enable-tile-op-expand. +// The template_tmov_m2s template should be selected when dst memory_space is "scaling". +// This test verifies the template works, output is f32 (no fixpipe quantization). +// Compiled by ptoas --enable-insert-sync --enable-tile-op-expand --pto-backend=vpto --pto-level=level3 + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @TMOV_FP_f16_16x16x16(%a_gm: !pto.ptr, %b_gm: !pto.ptr, %scale_gm: !pto.ptr, %c_gm: !pto.ptr) attributes {pto.aicore} { + %c0_i64 = arith.constant 0 : i64 + %c1_i64 = arith.constant 1 : i64 + %c16_i64 = arith.constant 16 : i64 + %c32_i64 = arith.constant 32 : i64 + %c64_i64 = arith.constant 64 : i64 // 16 * 4 bytes for scale (1x16 f32) + %c256_i64 = arith.constant 256 : i64 + %c512_i64 = arith.constant 512 : i64 + %c1024_i64 = arith.constant 1024 : i64 + %false = arith.constant false + + // L1 Mat tiles (source for TMOV) + %l1_a_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + %l1_b_tile = pto.alloc_tile addr = %c512_i64 + : !pto.tile_buf + + // L1 Scale Mat tile (holding scale data from GM, will be moved to Scaling buffer) + %l1_scale_tile = pto.alloc_tile addr = %c1024_i64 + : !pto.tile_buf + + // L0A/L0B tiles + %l0a_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + %l0b_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + + // Scaling tile (L0/FB) - destination for TMOV Mat->Scaling + %scale_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + + // Acc tile + %l0c_tile = pto.alloc_tile addr = %c0_i64 + : !pto.tile_buf + + // Get addresses + %l1_a = pto.tile_buf_addr %l1_a_tile + : !pto.tile_buf + -> !pto.ptr + %l1_b = pto.tile_buf_addr %l1_b_tile + : !pto.tile_buf + -> !pto.ptr + %l1_scale = pto.tile_buf_addr %l1_scale_tile + : !pto.tile_buf + -> !pto.ptr + %l0a = pto.tile_buf_addr %l0a_tile + : !pto.tile_buf + -> !pto.ptr + %l0b = pto.tile_buf_addr %l0b_tile + : !pto.tile_buf + -> !pto.ptr + %l0c = pto.tile_buf_addr %l0c_tile + : !pto.tile_buf + -> !pto.ptr + + // TLOAD A: GM -> L1 Mat + pto.mte_gm_l1_frac %a_gm, %l1_a, nd2nz, + shape(%c16_i64, %c16_i64), + src_layout(%c32_i64), + dst_group(%c1_i64, %c1_i64, %c16_i64, %c0_i64), + ctrl(%c0_i64, %false) + : !pto.ptr, !pto.ptr, nd2nz, + shape i64, i64, src_layout(i64), + dst_group i64, i64, i64, i64, ctrl i64, i1 + + pto.set_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID0"] + pto.wait_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID0"] + + // TMOV Mat->Left: L1 Mat -> L0A Left + pto.tmov ins(%l1_a_tile : !pto.tile_buf) + outs(%l0a_tile : !pto.tile_buf) + + // TLOAD B: GM -> L1 Mat + pto.mte_gm_l1_frac %b_gm, %l1_b, nd2nz, + shape(%c16_i64, %c16_i64), + src_layout(%c32_i64), + dst_group(%c1_i64, %c1_i64, %c16_i64, %c0_i64), + ctrl(%c0_i64, %false) + : !pto.ptr, !pto.ptr, nd2nz, + shape i64, i64, src_layout(i64), + dst_group i64, i64, i64, i64, ctrl i64, i1 + + pto.set_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID1"] + pto.wait_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID1"] + + // TMOV Mat->Right: L1 Mat -> L0B Right + pto.tmov ins(%l1_b_tile : !pto.tile_buf) + outs(%l0b_tile : !pto.tile_buf) + + // TLOAD Scale: GM -> L1 Mat (scale data is 1x16 f32 = 64 bytes) + pto.mte_gm_l1 %scale_gm, %l1_scale, %c64_i64 + nburst(%c1_i64, %c0_i64, %c0_i64) + : !pto.ptr, !pto.ptr, i64, i64, i64, i64 + + pto.set_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID2"] + pto.wait_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID2"] + + // TMOV Mat->Scaling: L1 Mat -> Scaling buffer + // This tests the template_tmov_m2s template with constraint dst.memory_space == SCALING + pto.tmov ins(%l1_scale_tile : !pto.tile_buf) + outs(%scale_tile : !pto.tile_buf) + + pto.set_flag["PIPE_MTE1", "PIPE_M", "EVENT_ID0"] + pto.wait_flag["PIPE_MTE1", "PIPE_M", "EVENT_ID0"] + + // TMATMUL: compute Acc = Left x Right + pto.tmatmul ins(%l0a_tile, %l0b_tile : !pto.tile_buf, + !pto.tile_buf) + outs(%l0c_tile : !pto.tile_buf) + + pto.set_flag["PIPE_M", "PIPE_FIX", "EVENT_ID1"] + pto.wait_flag["PIPE_M", "PIPE_FIX", "EVENT_ID1"] + + // TSTORE: Acc -> GM (f32 output, no fixpipe) + pto.mte_l0c_gm %l0c, %c_gm, %c16_i64, %c16_i64, %c16_i64, %c16_i64, %c0_i64, %c0_i64, + nz2nd + : !pto.ptr, !pto.ptr, i64, i64, i64, i64, i64, i64 + + pto.barrier #pto.pipe + return + } +} \ No newline at end of file diff --git a/tilelang-dsl/python/tilelang_dsl/expand_helper.py b/tilelang-dsl/python/tilelang_dsl/expand_helper.py index b7581469c8..6798a6f787 100644 --- a/tilelang-dsl/python/tilelang_dsl/expand_helper.py +++ b/tilelang-dsl/python/tilelang_dsl/expand_helper.py @@ -73,6 +73,7 @@ def _populate_dtype_map() -> None: _MEMSPACE_MAP = { "ub": MemorySpace.UB, + "vec": MemorySpace.UB, # MLIR uses "vec" for UB (Unified Buffer) "gm": MemorySpace.GM, "mat": MemorySpace.MAT, "left": MemorySpace.LEFT,