Star912
1827
代码
代码
Issues224
Pull Requests309
Actions
项目讨论
Wiki
分析
Star912
1827
已合并
编译优化：减小mm算子的头文件包含，提高编译速度 #848
chen-shuai创建于 1月21日
编译优化：减小mm算子的头文件包含，提高编译速度 #848
已合并
chen-shuai创建于 1月21日
从chen-shuai/ops-nn_7392replace_matmul_head合入到cann/ops-nnmaster
共 93 个文件变更+183-169
Mmatmul/batch_mat_mul_v3/op_kernel/arch35/batch_mat_mul_v3_asw_al1_full_load_kernel_advanced.h+2-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -17,7 +17,7 @@
17 17 
18#include "batch_mat_mul_v3_asw_block_advanced.h"18#include "batch_mat_mul_v3_asw_block_advanced.h"
19#include "../../mat_mul_v3/arch35/mat_mul_v3_full_load_kernel_helper.h"19#include "../../mat_mul_v3/arch35/mat_mul_v3_full_load_kernel_helper.h"
20-#include "kernel_operator.h"20+#include "kernel_basic_intf.h"
21#include "lib/matmul_intf.h"21#include "lib/matmul_intf.h"
22 22 
23namespace BatchMatMulV3Advanced {23namespace BatchMatMulV3Advanced {
Mmatmul/batch_mat_mul_v3/op_kernel/arch35/batch_mat_mul_v3_asw_bl1_full_load_kernel_advanced.h+2-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -17,7 +17,7 @@
17 17 
18#include "batch_mat_mul_v3_asw_block_advanced.h"18#include "batch_mat_mul_v3_asw_block_advanced.h"
19#include "../../mat_mul_v3/arch35/mat_mul_v3_full_load_kernel_helper.h"19#include "../../mat_mul_v3/arch35/mat_mul_v3_full_load_kernel_helper.h"
20-#include "kernel_operator.h"20+#include "kernel_basic_intf.h"
21#include "lib/matmul_intf.h"21#include "lib/matmul_intf.h"
22 22 
23namespace BatchMatMulV3Advanced {23namespace BatchMatMulV3Advanced {
Mmatmul/batch_mat_mul_v3/op_kernel/arch35/batch_mat_mul_v3_asw_block_advanced.h+2-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -15,7 +15,7 @@
15#ifndef BATCH_MATMUL_V3_ASW_BLOCK_ADVANCED_H15#ifndef BATCH_MATMUL_V3_ASW_BLOCK_ADVANCED_H
16#define BATCH_MATMUL_V3_ASW_BLOCK_ADVANCED_H16#define BATCH_MATMUL_V3_ASW_BLOCK_ADVANCED_H
17 17 
18-#include "kernel_operator.h"18+#include "kernel_basic_intf.h"
19#include "lib/matmul_intf.h"19#include "lib/matmul_intf.h"
20#include "../../mat_mul_v3/mat_mul_v3_common.h"20#include "../../mat_mul_v3/mat_mul_v3_common.h"
21#include "../../mat_mul_v3/arch35/mat_mul_tiling_data.h"21#include "../../mat_mul_v3/arch35/mat_mul_tiling_data.h"
Mmatmul/batch_mat_mul_v3/op_kernel/arch35/batch_mat_mul_v3_asw_kernel_advanced.h+2-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,7 +16,7 @@
16#define BATCH_MAT_MUL_V3_ASW_KERNEL_ADVANCED_H16#define BATCH_MAT_MUL_V3_ASW_KERNEL_ADVANCED_H
17 17 
18#include "batch_mat_mul_v3_asw_block_advanced.h"18#include "batch_mat_mul_v3_asw_block_advanced.h"
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20#include "lib/matmul_intf.h"20#include "lib/matmul_intf.h"
21 21 
22namespace BatchMatMulV3Advanced {22namespace BatchMatMulV3Advanced {
Mmatmul/batch_mat_mul_v3/op_kernel/arch35/batch_mat_mul_v3_iterbatch_block_advanced.h+2-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -15,7 +15,7 @@
15#ifndef BATCH_MATMUL_V3_ITERBATCH_BLOCK_ADVANCED_H15#ifndef BATCH_MATMUL_V3_ITERBATCH_BLOCK_ADVANCED_H
16#define BATCH_MATMUL_V3_ITERBATCH_BLOCK_ADVANCED_H16#define BATCH_MATMUL_V3_ITERBATCH_BLOCK_ADVANCED_H
17 17 
18-#include "kernel_operator.h"18+#include "kernel_basic_intf.h"
19#include "lib/matmul_intf.h"19#include "lib/matmul_intf.h"
20 20 
21namespace BatchMatMulV3Advanced {21namespace BatchMatMulV3Advanced {
Mmatmul/batch_mat_mul_v3/op_kernel/arch35/batch_mat_mul_v3_iterbatch_kernel_advanced.h+2-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,7 +16,7 @@
16#define BATCH_MAT_MUL_V3_ITERBATCH_KERNEL_ADVANCED_H16#define BATCH_MAT_MUL_V3_ITERBATCH_KERNEL_ADVANCED_H
17 17 
18#include "batch_mat_mul_v3_iterbatch_block_advanced.h"18#include "batch_mat_mul_v3_iterbatch_block_advanced.h"
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20#include "lib/matmul_intf.h"20#include "lib/matmul_intf.h"
21 21 
22namespace BatchMatMulV3Advanced {22namespace BatchMatMulV3Advanced {
Mmatmul/batch_mat_mul_v3/tests/ut/op_kernel/test_batch_mat_mul_v3.cpp+1-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -31,7 +31,6 @@
31#include <cstdint>31#include <cstdint>
32 32 
33#include "kernel_tiling/kernel_tiling.h"33#include "kernel_tiling/kernel_tiling.h"
34-// #include "kernel_operator.h"
35using namespace std;34using namespace std;
36 35 
37class batch_mat_mul_v3_test : public testing::Test {36class batch_mat_mul_v3_test : public testing::Test {
Mmatmul/common/cmct/block/block_mmad_a_prefetch_b_prologue.h+1-1
@@ -21,7 +21,7 @@
21#include "../utils/tensor_utils.h"21#include "../utils/tensor_utils.h"
22#include "../policy/dispatch_policy.h"22#include "../policy/dispatch_policy.h"
23#include "../utils/tuple_utils.h"23#include "../utils/tuple_utils.h"
24-#include "kernel_operator_intf.h"24+#include "kernel_basic_intf.h"
25#include "../utils/underscore.h"25#include "../utils/underscore.h"
26#include "../utils/tensor_traits.h"26#include "../utils/tensor_traits.h"
27#include "../utils/math_utils.h"27#include "../utils/math_utils.h"
Mmatmul/common/cmct/block/block_mmad_b_prologue_mx.h+1-1
@@ -22,7 +22,7 @@
22#include "../utils/tuple_utils.h"22#include "../utils/tuple_utils.h"
23#include "../policy/dispatch_policy.h"23#include "../policy/dispatch_policy.h"
24#include "block_mmad.h"24#include "block_mmad.h"
25-#include "kernel_operator.h"25+#include "kernel_basic_intf.h"
26#include "lib/matmul_intf.h"26#include "lib/matmul_intf.h"
27 27 
28namespace Cmct::Gemm::Block {28namespace Cmct::Gemm::Block {
Mmatmul/common/cmct/block/block_mmad_builder.h+1-1
@@ -20,7 +20,7 @@
20 20 
21#define ASCENDC_CUBE_ONLY21#define ASCENDC_CUBE_ONLY
22#include "../utils/host_utils.h"22#include "../utils/host_utils.h"
23-#include "kernel_operator.h"23+#include "kernel_basic_intf.h"
24 24 
25#include "../utils/common_utils.h"25#include "../utils/common_utils.h"
26#include "../utils/layout_utils.h"26#include "../utils/layout_utils.h"
Mmatmul/common/cmct/block/block_mmad_pertile_param.h+1-1
@@ -16,7 +16,7 @@
16#ifndef MATMUL_BLOCK_BLOCK_MMAD_PERTILE_PARAM_H16#ifndef MATMUL_BLOCK_BLOCK_MMAD_PERTILE_PARAM_H
17#define MATMUL_BLOCK_BLOCK_MMAD_PERTILE_PARAM_H17#define MATMUL_BLOCK_BLOCK_MMAD_PERTILE_PARAM_H
18 18 
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20#include "../utils/common_utils.h"20#include "../utils/common_utils.h"
21#include "../utils/layout_utils.h"21#include "../utils/layout_utils.h"
22#include "../utils/tuple_utils.h"22#include "../utils/tuple_utils.h"
Mmatmul/common/cmct/epilogue/block_epilogue.h+1-1
@@ -16,7 +16,7 @@
16#ifndef EPILOGUE_BLOCK_EPILOGUE_H16#ifndef EPILOGUE_BLOCK_EPILOGUE_H
17#define EPILOGUE_BLOCK_EPILOGUE_H17#define EPILOGUE_BLOCK_EPILOGUE_H
18#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3101)18#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3101)
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20#include "../utils/common_utils.h"20#include "../utils/common_utils.h"
21#include "../utils/device_utils.h"21#include "../utils/device_utils.h"
22#include "fusion/default_fusion_op.h"22#include "fusion/default_fusion_op.h"
Mmatmul/common/cmct/epilogue/block_epilogue_empty.h+1-1
@@ -15,7 +15,7 @@
15 15 
16#ifndef EPILOGUE_BLOCK_EPILOGUE_EMPTY_H16#ifndef EPILOGUE_BLOCK_EPILOGUE_EMPTY_H
17#define EPILOGUE_BLOCK_EPILOGUE_EMPTY_H17#define EPILOGUE_BLOCK_EPILOGUE_EMPTY_H
18-#include "kernel_operator.h"18+#include "kernel_basic_intf.h"
19#include "fusion/default_fusion_op.h"19#include "fusion/default_fusion_op.h"
20#include "../utils/common_utils.h"20#include "../utils/common_utils.h"
21#include "../utils/device_utils.h"21#include "../utils/device_utils.h"
Mmatmul/common/cmct/epilogue/block_epilogue_fixpipe.h+1-1
@@ -16,7 +16,7 @@
16#ifndef CMCT_INCLUDE_EPILOGUE_BLOCK_EPILOGUE_FIXPIPE_H16#ifndef CMCT_INCLUDE_EPILOGUE_BLOCK_EPILOGUE_FIXPIPE_H
17#define CMCT_INCLUDE_EPILOGUE_BLOCK_EPILOGUE_FIXPIPE_H17#define CMCT_INCLUDE_EPILOGUE_BLOCK_EPILOGUE_FIXPIPE_H
18#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3101)18#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3101)
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20#include "fusion/default_fusion_op.h"20#include "fusion/default_fusion_op.h"
21 21 
22namespace Cmct {22namespace Cmct {
Mmatmul/common/cmct/epilogue/block_epilogue_iterbatch.h+1-1
@@ -15,7 +15,7 @@
15 15 
16#ifndef CMCT_INCLUDE_EPILOGUE_BLOCK_EPILOGUE_ITERBATCH_H16#ifndef CMCT_INCLUDE_EPILOGUE_BLOCK_EPILOGUE_ITERBATCH_H
17#define CMCT_INCLUDE_EPILOGUE_BLOCK_EPILOGUE_ITERBATCH_H17#define CMCT_INCLUDE_EPILOGUE_BLOCK_EPILOGUE_ITERBATCH_H
18-#include "kernel_operator.h"18+#include "kernel_basic_intf.h"
19#include "fusion/default_fusion_op.h"19#include "fusion/default_fusion_op.h"
20#include "../utils/common_utils.h"20#include "../utils/common_utils.h"
21#include "../utils/device_utils.h"21#include "../utils/device_utils.h"
Mmatmul/common/cmct/epilogue/block_epilogue_pertile.h+1-1
@@ -15,7 +15,7 @@
15 15 
16#ifndef EPILOGUE_BLOCK_EPILOGUE_PERTILE_H16#ifndef EPILOGUE_BLOCK_EPILOGUE_PERTILE_H
17#define EPILOGUE_BLOCK_EPILOGUE_PERTILE_H17#define EPILOGUE_BLOCK_EPILOGUE_PERTILE_H
18-#include "kernel_operator.h"18+#include "kernel_basic_intf.h"
19#include "../utils/common_utils.h"19#include "../utils/common_utils.h"
20#include "../utils/grouped_matmul_constant.h"20#include "../utils/grouped_matmul_constant.h"
21#include "../utils/layout_utils.h"21#include "../utils/layout_utils.h"
Mmatmul/common/cmct/epilogue/block_epilogue_streamk.h+1-1
@@ -15,7 +15,7 @@
15 15 
16#ifndef CMCT_INCLUDE_EPILOGUE_BLOCK_EPILOGUE_STREAMK_H16#ifndef CMCT_INCLUDE_EPILOGUE_BLOCK_EPILOGUE_STREAMK_H
17#define CMCT_INCLUDE_EPILOGUE_BLOCK_EPILOGUE_STREAMK_H17#define CMCT_INCLUDE_EPILOGUE_BLOCK_EPILOGUE_STREAMK_H
18-#include "kernel_operator.h"18+#include "kernel_basic_intf.h"
19#include "fusion/default_fusion_op.h"19#include "fusion/default_fusion_op.h"
20#include "../utils/common_utils.h"20#include "../utils/common_utils.h"
21#include "../utils/device_utils.h"21#include "../utils/device_utils.h"
Mmatmul/common/cmct/epilogue/fusion/default_fusion_op.h+1-1
@@ -15,7 +15,7 @@
15 15 
16#ifndef EPILOGUE_FUSION_DEFAULT_FUSION_OP_H16#ifndef EPILOGUE_FUSION_DEFAULT_FUSION_OP_H
17#define EPILOGUE_FUSION_DEFAULT_FUSION_OP_H17#define EPILOGUE_FUSION_DEFAULT_FUSION_OP_H
18-#include "kernel_operator.h"18+#include "kernel_basic_intf.h"
19#include "../../utils/common_utils.h"19#include "../../utils/common_utils.h"
20#include "../../utils/device_utils.h"20#include "../../utils/device_utils.h"
21 21 
Mmatmul/common/cmct/epilogue/fusion/fusion_add.h+1-1
@@ -15,7 +15,7 @@
15 15 
16#ifndef EPILOGUE_FUSION_EPILOGUE_FUSION_ADD_H16#ifndef EPILOGUE_FUSION_EPILOGUE_FUSION_ADD_H
17#define EPILOGUE_FUSION_EPILOGUE_FUSION_ADD_H17#define EPILOGUE_FUSION_EPILOGUE_FUSION_ADD_H
18-#include "kernel_operator.h"18+#include "kernel_basic_intf.h"
19#include "../../utils/common_utils.h"19#include "../../utils/common_utils.h"
20#include "../../utils/device_utils.h"20#include "../../utils/device_utils.h"
21 21 
Mmatmul/common/cmct/epilogue/fusion/fusion_gelu.h+2-1
@@ -15,7 +15,8 @@
15 15 
16#ifndef EPILOGUE_FUSION_FUSION_GELU_H16#ifndef EPILOGUE_FUSION_FUSION_GELU_H
17#define EPILOGUE_FUSION_FUSION_GELU_H17#define EPILOGUE_FUSION_FUSION_GELU_H
18-#include "kernel_operator.h"18+#include "kernel_basic_intf.h"
19+#include "math/erf.h"
19#include "../../utils/common_utils.h"20#include "../../utils/common_utils.h"
20#include "../../utils/device_utils.h"21#include "../../utils/device_utils.h"
21 22 
Mmatmul/common/cmct/epilogue/fusion/fusion_mul.h+1-1
@@ -15,7 +15,7 @@
15 15 
16#ifndef EPILOGUE_FUSION_FUSION_MUL_H16#ifndef EPILOGUE_FUSION_FUSION_MUL_H
17#define EPILOGUE_FUSION_FUSION_MUL_H17#define EPILOGUE_FUSION_FUSION_MUL_H
18-#include "kernel_operator.h"18+#include "kernel_basic_intf.h"
19#include "../../utils/common_utils.h"19#include "../../utils/common_utils.h"
20#include "../../utils/device_utils.h"20#include "../../utils/device_utils.h"
21 21 
Mmatmul/common/cmct/kernel/kernel_batch_matmul_to_mul.h+1-1
@@ -16,7 +16,7 @@
16#ifndef MATMUL_KERNEL_KERNEL_BATCH_MATMUL_TO_MUL_H16#ifndef MATMUL_KERNEL_KERNEL_BATCH_MATMUL_TO_MUL_H
17#define MATMUL_KERNEL_KERNEL_BATCH_MATMUL_TO_MUL_H17#define MATMUL_KERNEL_KERNEL_BATCH_MATMUL_TO_MUL_H
18 18 
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20#include "lib/matmul_intf.h"20#include "lib/matmul_intf.h"
21 21 
22#include "../utils/common_utils.h"22#include "../utils/common_utils.h"
Mmatmul/common/cmct/kernel/kernel_matmul.h+1-1
@@ -17,7 +17,7 @@
17#define MATMUL_KERNEL_KERNEL_MATMUL_H17#define MATMUL_KERNEL_KERNEL_MATMUL_H
18 18 
19#define ASCENDC_CUBE_ONLY19#define ASCENDC_CUBE_ONLY
20-#include "kernel_operator.h"20+#include "kernel_basic_intf.h"
21#include "lib/matmul_intf.h"21#include "lib/matmul_intf.h"
22 22 
23#include "../utils/common_utils.h"23#include "../utils/common_utils.h"
Mmatmul/common/cmct/kernel/kernel_matmul_a_prefetch_b_antiquant.h+1-1
@@ -17,7 +17,7 @@
17 17 
18#include "../utils/integral_constant.h"18#include "../utils/integral_constant.h"
19#include "../utils/tuple_utils.h"19#include "../utils/tuple_utils.h"
20-#include "kernel_operator.h"20+#include "kernel_basic_intf.h"
21#include "kernel_operator_list_tensor_intf.h"21#include "kernel_operator_list_tensor_intf.h"
22#include "lib/matmul_intf.h"22#include "lib/matmul_intf.h"
23#include "../prologue/constant.h"23#include "../prologue/constant.h"
Mmatmul/common/cmct/kernel/kernel_matmul_iterbatch.h+1-1
@@ -17,7 +17,7 @@
17#define MATMUL_KERNEL_KERNEL_MATMUL_ITERBATCH_H17#define MATMUL_KERNEL_KERNEL_MATMUL_ITERBATCH_H
18 18 
19#define ASCENDC_CUBE_ONLY19#define ASCENDC_CUBE_ONLY
20-#include "kernel_operator.h"20+#include "kernel_basic_intf.h"
21#include "lib/matmul_intf.h"21#include "lib/matmul_intf.h"
22 22 
23#include "../utils/common_utils.h"23#include "../utils/common_utils.h"
Mmatmul/common/cmct/kernel/kernel_matmul_merge_batch.h+1-1
@@ -17,7 +17,7 @@
17#define MATMUL_KERNEL_KERNEL_MATMUL_MERGE_BATCH_H17#define MATMUL_KERNEL_KERNEL_MATMUL_MERGE_BATCH_H
18 18 
19#define ASCENDC_CUBE_ONLY19#define ASCENDC_CUBE_ONLY
20-#include "kernel_operator.h"20+#include "kernel_basic_intf.h"
21#include "lib/matmul_intf.h"21#include "lib/matmul_intf.h"
22 22 
23#include "../utils/common_utils.h"23#include "../utils/common_utils.h"
Mmatmul/common/cmct/kernel/kernel_matmul_mix.h+1-1
@@ -17,7 +17,7 @@
17#define MATMUL_KERNEL_KERNEL_MATMUL_MIX_H17#define MATMUL_KERNEL_KERNEL_MATMUL_MIX_H
18 18 
19#define ASCENDC_CUBE_ONLY19#define ASCENDC_CUBE_ONLY
20-#include "kernel_operator.h"20+#include "kernel_basic_intf.h"
21#include "lib/matmul_intf.h"21#include "lib/matmul_intf.h"
22 22 
23#include "../utils/common_utils.h"23#include "../utils/common_utils.h"
Mmatmul/common/cmct/kernel/kernel_matmul_mix_fixpipe_opti.h+1-1
@@ -17,7 +17,7 @@
17#define CMCT_INCLUDE_MATMUL_KERNEL_KERNEL_MATMUL_MIX_FIXPIPE_OPTI_H17#define CMCT_INCLUDE_MATMUL_KERNEL_KERNEL_MATMUL_MIX_FIXPIPE_OPTI_H
18 18 
19#define ASCENDC_CUBE_ONLY19#define ASCENDC_CUBE_ONLY
20-#include "kernel_operator.h"20+#include "kernel_basic_intf.h"
21#include "lib/matmul_intf.h"21#include "lib/matmul_intf.h"
22 22 
23#include "../utils/common_utils.h"23#include "../utils/common_utils.h"
Mmatmul/common/cmct/kernel/kernel_matmul_mix_with_weight_prologue.h+1-1
@@ -23,7 +23,7 @@
23#include "../utils/tuple_utils.h"23#include "../utils/tuple_utils.h"
24#include "../utils/underscore.h"24#include "../utils/underscore.h"
25#include "../block/block_mmad_b_prologue_mx.h"25#include "../block/block_mmad_b_prologue_mx.h"
26-#include "kernel_operator.h"26+#include "kernel_basic_intf.h"
27#include "kernel_operator_list_tensor_intf.h"27#include "kernel_operator_list_tensor_intf.h"
28#include "lib/matmul_intf.h"28#include "lib/matmul_intf.h"
29 29 
Mmatmul/common/cmct/kernel/kernel_matmul_streamk.h+1-1
@@ -17,7 +17,7 @@
17#define CMCT_INCLUDE_MATMUL_KERNEL_KERNEL_MATMUL_STREAMK_H17#define CMCT_INCLUDE_MATMUL_KERNEL_KERNEL_MATMUL_STREAMK_H
18 18 
19#define ASCENDC_CUBE_ONLY19#define ASCENDC_CUBE_ONLY
20-#include "kernel_operator.h"20+#include "kernel_basic_intf.h"
21#include "lib/matmul_intf.h"21#include "lib/matmul_intf.h"
22 22 
23#include "../utils/common_utils.h"23#include "../utils/common_utils.h"
Mmatmul/common/cmct/kernel/kernel_matmul_without_que.h+1-1
@@ -17,7 +17,7 @@
17#define MATMUL_KERNEL_KERNEL_MATMUL_WITHOUT_QUE_H17#define MATMUL_KERNEL_KERNEL_MATMUL_WITHOUT_QUE_H
18 18 
19#define ASCENDC_CUBE_ONLY19#define ASCENDC_CUBE_ONLY
20-#include "kernel_operator.h"20+#include "kernel_basic_intf.h"
21#include "lib/matmul_intf.h"21#include "lib/matmul_intf.h"
22 22 
23#include "../utils/common_utils.h"23#include "../utils/common_utils.h"
Mmatmul/common/cmct/kernel/kernel_qbmm_mx.h+2-3
@@ -15,8 +15,7 @@
15 15 
16#ifndef MATMUL_KERNEL_KERNEL_QBMM_MX_H16#ifndef MATMUL_KERNEL_KERNEL_QBMM_MX_H
17#define MATMUL_KERNEL_KERNEL_QBMM_MX_H17#define MATMUL_KERNEL_KERNEL_QBMM_MX_H
18-#include "kernel_operator.h"18+#include "kernel_basic_intf.h"
19-#include "kernel_operator_intf.h"
20#include "../utils/common_utils.h"19#include "../utils/common_utils.h"
21#include "../utils/fill_utils.h"20#include "../utils/fill_utils.h"
22#include "../utils/grouped_matmul_constant.h"21#include "../utils/grouped_matmul_constant.h"
@@ -280,4 +279,4 @@ __aicore__ inline void QuantMmBatchMX<QBMM_MX_KERNEL_FUN_TEM_PARAMS>::ProcessSin
280} // namespace Gemm279} // namespace Gemm
281} // namespace Cmct280} // namespace Cmct
282 281 
283-#endif282+#endif
Mmatmul/common/cmct/kernel/kernel_qbmm_pertile.h+1-2
@@ -16,8 +16,7 @@
16#ifndef MATMUL_KERNEL_KERNEL_QBMM_PERTILE_H16#ifndef MATMUL_KERNEL_KERNEL_QBMM_PERTILE_H
17#define MATMUL_KERNEL_KERNEL_QBMM_PERTILE_H17#define MATMUL_KERNEL_KERNEL_QBMM_PERTILE_H
18 18 
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20-#include "kernel_operator_intf.h"
21 20 
22#include "../utils/common_utils.h"21#include "../utils/common_utils.h"
23#include "../utils/fill_utils.h"22#include "../utils/fill_utils.h"
Mmatmul/common/cmct/prologue/constant.h+1-1
@@ -10,7 +10,7 @@
10#ifndef PROLOGUE_CONSTANT_H10#ifndef PROLOGUE_CONSTANT_H
11#define PROLOGUE_CONSTANT_H11#define PROLOGUE_CONSTANT_H
12 12 
13-#include "kernel_operator_intf.h"13+#include "kernel_basic_intf.h"
14#include "../utils/tensor_traits.h"14#include "../utils/tensor_traits.h"
15 15 
16namespace Cmct::Prologue {16namespace Cmct::Prologue {
Mmatmul/common/cmct/prologue/tile/antiquant_nd_kn.h+1-1
@@ -10,7 +10,7 @@
10#ifndef PROLOGUE_TILE_ANTIQUANT_ND_KN_H10#ifndef PROLOGUE_TILE_ANTIQUANT_ND_KN_H
11#define PROLOGUE_TILE_ANTIQUANT_ND_KN_H11#define PROLOGUE_TILE_ANTIQUANT_ND_KN_H
12#include "../../utils/underscore.h"12#include "../../utils/underscore.h"
13-#include "kernel_operator_intf.h"13+#include "kernel_basic_intf.h"
14namespace Cmct::Prologue::Tile {14namespace Cmct::Prologue::Tile {
15 15 
16using AscendC::MicroAPI::RegTensor;16using AscendC::MicroAPI::RegTensor;
Mmatmul/common/cmct/prologue/tile/antiquant_nd_nk.h+1-1
@@ -9,7 +9,7 @@
9 */9 */
10#ifndef PROLOGUE_TILE_ANTIQUANT_ND_NK_H10#ifndef PROLOGUE_TILE_ANTIQUANT_ND_NK_H
11#define PROLOGUE_TILE_ANTIQUANT_ND_NK_H11#define PROLOGUE_TILE_ANTIQUANT_ND_NK_H
12-#include "kernel_operator_intf.h"12+#include "kernel_basic_intf.h"
13#include "../../utils/underscore.h"13#include "../../utils/underscore.h"
14#include "../../utils/constant.h"14#include "../../utils/constant.h"
15#include "kernel_micro_datacopy_intf_impl.h"15#include "kernel_micro_datacopy_intf_impl.h"
Mmatmul/common/cmct/prologue/tile/antiquant_zn.h+1-1
@@ -11,7 +11,7 @@
11#define PROLOGUE_TILE_ANTIQUANT_ZN_H11#define PROLOGUE_TILE_ANTIQUANT_ZN_H
12#include "../../utils/constant.h"12#include "../../utils/constant.h"
13#include "../../utils/underscore.h"13#include "../../utils/underscore.h"
14-#include "kernel_operator_intf.h"14+#include "kernel_basic_intf.h"
15namespace Cmct::Prologue::Tile {15namespace Cmct::Prologue::Tile {
16using AscendC::MicroAPI::RegTensor;16using AscendC::MicroAPI::RegTensor;
17using Gemm::BLK_ELEM;17using Gemm::BLK_ELEM;
Mmatmul/common/cmct/prologue/tile/tile_cast_b8_to_b16.h+1-1
@@ -10,7 +10,7 @@
10 */10 */
11#ifndef PROLOGUE_TILE_CAST_MX_H11#ifndef PROLOGUE_TILE_CAST_MX_H
12#define PROLOGUE_TILE_CAST_MX_H12#define PROLOGUE_TILE_CAST_MX_H
13-#include "kernel_operator_intf.h"13+#include "kernel_basic_intf.h"
14#include "../../utils/math_utils.h"14#include "../../utils/math_utils.h"
15#include "../../utils/underscore.h"15#include "../../utils/underscore.h"
16 16 
Mmatmul/common/cmct/tile/tile_copy_if.h+1-1
@@ -10,7 +10,7 @@
10#ifndef TILE_TILE_COPY_IF_H10#ifndef TILE_TILE_COPY_IF_H
11#define TILE_TILE_COPY_IF_H11#define TILE_TILE_COPY_IF_H
12 12 
13-#include "kernel_operator_intf.h"13+#include "kernel_basic_intf.h"
14#include "../utils/arch.h"14#include "../utils/arch.h"
15 15 
16namespace Cmct::Gemm::Tile {16namespace Cmct::Gemm::Tile {
Mmatmul/common/cmct/utils/common_utils.h+2-1
@@ -17,7 +17,8 @@
17#define UTILS_COMMON_UTILS_H17#define UTILS_COMMON_UTILS_H
18 18 
19#include "integral_constant.h"19#include "integral_constant.h"
20-#include "kernel_operator.h"20+#include "kernel_basic_intf.h"
21+#include "std/algorithm.h"
21#include "lib/matmul_intf.h"22#include "lib/matmul_intf.h"
22namespace Cmct {23namespace Cmct {
23namespace Gemm {24namespace Gemm {
Mmatmul/common/cmct/utils/device_utils.h+1-1
@@ -15,7 +15,7 @@
15 15 
16#ifndef UTILS_DEVICE_UTILS_H16#ifndef UTILS_DEVICE_UTILS_H
17#define UTILS_DEVICE_UTILS_H17#define UTILS_DEVICE_UTILS_H
18-#include "kernel_operator.h"18+#include "kernel_basic_intf.h"
19namespace Cmct {19namespace Cmct {
20namespace Gemm {20namespace Gemm {
21 21 
Mmatmul/common/cmct/utils/fill_utils.h+1-1
@@ -16,7 +16,7 @@
16#ifndef UTILS_FILL_UTILS_H16#ifndef UTILS_FILL_UTILS_H
17#define UTILS_FILL_UTILS_H17#define UTILS_FILL_UTILS_H
18#include "common_utils.h"18#include "common_utils.h"
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20 20 
21namespace Cmct {21namespace Cmct {
22namespace Gemm {22namespace Gemm {
Mmatmul/common/cmct/utils/integral_constant.h+1-1
@@ -14,7 +14,7 @@
14 */14 */
15#ifndef UTILS_INTEGRAL_CONSTANT_H15#ifndef UTILS_INTEGRAL_CONSTANT_H
16#define UTILS_INTEGRAL_CONSTANT_H16#define UTILS_INTEGRAL_CONSTANT_H
17-#include "kernel_operator.h"17+#include "kernel_basic_intf.h"
18 18 
19namespace AscendC {19namespace AscendC {
20namespace Std {20namespace Std {
Mmatmul/common/cmct/utils/math_utils.h+2-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -11,7 +11,7 @@
11#define UTILS_MATH_UTILS_H11#define UTILS_MATH_UTILS_H
12 12 
13#include "device_utils.h"13#include "device_utils.h"
14-#include "kernel_operator_intf.h"14+#include "kernel_basic_intf.h"
15 15 
16namespace Cmct {16namespace Cmct {
17template <typename T>17template <typename T>
Mmatmul/common/cmct/utils/tensor_traits.h+2-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -9,7 +9,7 @@
9 */9 */
10#ifndef UTILS_TENSOR_TRAITS_H10#ifndef UTILS_TENSOR_TRAITS_H
11#define UTILS_TENSOR_TRAITS_H11#define UTILS_TENSOR_TRAITS_H
12-#include "kernel_operator.h"12+#include "kernel_basic_intf.h"
13#include "integral_constant.h"13#include "integral_constant.h"
14 14 
15namespace Cmct::Gemm {15namespace Cmct::Gemm {
Mmatmul/common/op_host/CMakeLists.txt+1-1
@@ -7,4 +7,4 @@
7# See LICENSE in the root of the software repository for the full text of the License.7# See LICENSE in the root of the software repository for the full text of the License.
8#/8#/
9 9 
10-add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE common ACLNNTYPE aclnn_exclude)10+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE common ACLNNTYPE aclnn_exclude DEPENDENCIES batch_mat_mul_v3)
Mmatmul/fused_mat_mul/op_kernel/arch35/fused_mat_mul_tilingkey.h+2-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -19,7 +19,7 @@
19#include "ascendc/host_api/tiling/template_argument.h"19#include "ascendc/host_api/tiling/template_argument.h"
20#include "fused_mat_mul_tiling_key_public.h"20#include "fused_mat_mul_tiling_key_public.h"
21#if defined(__CCE_AICORE__)21#if defined(__CCE_AICORE__)
22-#include "kernel_operator.h"22+#include "kernel_basic_intf.h"
23#endif23#endif
24 24 
25#if defined(ORIG_DTYPE_X1) && ((defined(DT_BF16) && ORIG_DTYPE_X1 == DT_BF16)) || \25#if defined(ORIG_DTYPE_X1) && ((defined(DT_BF16) && ORIG_DTYPE_X1 == DT_BF16)) || \
Mmatmul/gemm_v3/op_kernel/gemm_v3.cpp+2-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -19,6 +19,7 @@
19#endif19#endif
20 20 
21#include "../mat_mul_v3/mat_mul_v3_common.h"21#include "../mat_mul_v3/mat_mul_v3_common.h"
22+#include "kernel_operator.h"
22#include "../mat_mul_v3/mat_mul_base_kernel.h"23#include "../mat_mul_v3/mat_mul_base_kernel.h"
23#include "../mat_mul_v3/mat_mul_deterministic_splitk_kernel.h"24#include "../mat_mul_v3/mat_mul_deterministic_splitk_kernel.h"
24#include "../mat_mul_v3/mat_mul_sc_splitk_kernel.h"25#include "../mat_mul_v3/mat_mul_sc_splitk_kernel.h"
Mmatmul/mat_mul_v3/op_kernel/arch35/mat_mul_input_k_eq_zero_clear_output.h+3-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -23,7 +23,8 @@
23#endif23#endif
24 24 
25#include "kernel_tiling/kernel_tiling.h"25#include "kernel_tiling/kernel_tiling.h"
26-#include "kernel_operator.h"26+#include "kernel_basic_intf.h"
27+#include "std/algorithm.h"
27#include "lib/matmul_intf.h"28#include "lib/matmul_intf.h"
28#include "../mat_mul_v3_common.h"29#include "../mat_mul_v3_common.h"
29 30 
Mmatmul/mat_mul_v3/op_kernel/mat_mul_base_block.h+2-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,6 +16,7 @@
16#define __OP_KERNEL_MATMUL_V3_BASE_BLOCK_H__16#define __OP_KERNEL_MATMUL_V3_BASE_BLOCK_H__
17 17 
18#include "mat_mul_v3_common.h"18#include "mat_mul_v3_common.h"
19+#include "kernel_operator.h"
19 20 
20using namespace AscendC;21using namespace AscendC;
21using namespace matmul;22using namespace matmul;
Mmatmul/mat_mul_v3/op_kernel/mat_mul_deterministic_splitk_kernel.h+2-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,6 +16,7 @@
16#define __OP_KERNEL_MATMUL_V3_DETERMINISTIC_SPLITK_KERNEL_H__16#define __OP_KERNEL_MATMUL_V3_DETERMINISTIC_SPLITK_KERNEL_H__
17 17 
18#include "mat_mul_v3_common.h"18#include "mat_mul_v3_common.h"
19+#include "kernel_operator.h"
19 20 
20const uint64_t ND2NZ_AIV_SYNC_AIC_FLAG = 8;21const uint64_t ND2NZ_AIV_SYNC_AIC_FLAG = 8;
21const uint64_t MAX_NUM = 12 * 1024; //ub最大容纳的元素个数22const uint64_t MAX_NUM = 12 * 1024; //ub最大容纳的元素个数
Mmatmul/mat_mul_v3/op_kernel/mat_mul_nd2nz_util.h+2-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,6 +16,7 @@
16#define __OP_KERNEL_MATMUL_V3_ND2NZ_UTIL__16#define __OP_KERNEL_MATMUL_V3_ND2NZ_UTIL__
17 17 
18#include "mat_mul_v3_common.h"18#include "mat_mul_v3_common.h"
19+#include "kernel_operator.h"
19 20 
20using namespace AscendC;21using namespace AscendC;
21using namespace matmul;22using namespace matmul;
Mmatmul/mat_mul_v3/op_kernel/mat_mul_sc_splitk_block.h+2-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,6 +16,7 @@
16#define __OP_KERNEL_MATMUL_V3_SC_SPLITK_BLOCK_H__16#define __OP_KERNEL_MATMUL_V3_SC_SPLITK_BLOCK_H__
17 17 
18#include "mat_mul_v3_common.h"18#include "mat_mul_v3_common.h"
19+#include "kernel_operator.h"
19 20 
20using namespace AscendC;21using namespace AscendC;
21using namespace matmul;22using namespace matmul;
Mmatmul/mat_mul_v3/op_kernel/mat_mul_unaligned_sc_splitk_kernel.h+2-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,6 +16,7 @@
16#define __OP_KERNEL_MATMUL_V3_UNALIGNED_SC_SPLITK_KERNEL_H__16#define __OP_KERNEL_MATMUL_V3_UNALIGNED_SC_SPLITK_KERNEL_H__
17 17 
18#include "mat_mul_v3_common.h"18#include "mat_mul_v3_common.h"
19+#include "kernel_operator.h"
19#include "mat_mul_sc_splitk_block.h"20#include "mat_mul_sc_splitk_block.h"
20#include "mat_mul_nd2nz.h"21#include "mat_mul_nd2nz.h"
21 22 
Mmatmul/mat_mul_v3/op_kernel/mat_mul_unaligned_sc_splitk_kernel_gm_to_l1.h+2-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,6 +16,7 @@
16#define __OP_KERNEL_MATMUL_V3_UNALIGNED_SC_SPLITK_KERNEL_GM_TO_L1_H__16#define __OP_KERNEL_MATMUL_V3_UNALIGNED_SC_SPLITK_KERNEL_GM_TO_L1_H__
17 17 
18#include "mat_mul_v3_common.h"18#include "mat_mul_v3_common.h"
19+#include "kernel_operator.h"
19#include "mat_mul_sc_splitk_block.h"20#include "mat_mul_sc_splitk_block.h"
20#include "mat_mul_nd2nz.h"21#include "mat_mul_nd2nz.h"
21 22 
Mmatmul/mat_mul_v3/op_kernel/mat_mul_v3.cpp+2-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -13,6 +13,7 @@
13 * \brief13 * \brief
14 */14 */
15 15 
16+#include "kernel_operator.h"
16#include "mat_mul_v3_common.h"17#include "mat_mul_v3_common.h"
17#include "mat_mul_base_kernel.h"18#include "mat_mul_base_kernel.h"
18#include "mat_mul_deterministic_splitk_kernel.h"19#include "mat_mul_deterministic_splitk_kernel.h"
Mmatmul/mat_mul_v3/op_kernel/mat_mul_v3_common.h+3-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -14,7 +14,8 @@
14 */14 */
15#ifndef __OP_KERNEL_MATMUL_V3_COMMON_H__15#ifndef __OP_KERNEL_MATMUL_V3_COMMON_H__
16#define __OP_KERNEL_MATMUL_V3_COMMON_H__16#define __OP_KERNEL_MATMUL_V3_COMMON_H__
17-#include "kernel_operator.h"17+#include "kernel_basic_intf.h"
18+#include "std/algorithm.h"
18#include "lib/matmul_intf.h"19#include "lib/matmul_intf.h"
19#include "mat_mul_v3_tiling_data.h"20#include "mat_mul_v3_tiling_data.h"
20 21 
Mmatmul/quant_batch_matmul_v3/op_kernel/arch35/qbmm_mix_online_dynamic.h+1-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,7 +16,6 @@
16#define QBMM_MIX_ONLINE_DYNAMIC_H16#define QBMM_MIX_ONLINE_DYNAMIC_H
17 17 
18#include "../../inc/platform.h"18#include "../../inc/platform.h"
19-#include "kernel_operator.h"
20#include "mm_extension_interface/qbmm_custom_mm_policy.h"19#include "mm_extension_interface/qbmm_custom_mm_policy.h"
21#include "qbmm_asw_block.h"20#include "qbmm_asw_block.h"
22#include "../quant_batch_matmul_v3_base.h"21#include "../quant_batch_matmul_v3_base.h"
Mmatmul/quant_batch_matmul_v3/op_kernel/quant_batch_matmul_v3.h+3-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -15,6 +15,8 @@
15#ifndef QUANT_BATCH_MATMUL_V3_H15#ifndef QUANT_BATCH_MATMUL_V3_H
16#define QUANT_BATCH_MATMUL_V3_H16#define QUANT_BATCH_MATMUL_V3_H
17 17 
18+#include "kernel_operator.h"
19+#include "kernel_operator_intf.h"
18#include "quant_batch_matmul_v3_base.h"20#include "quant_batch_matmul_v3_base.h"
19 21 
20namespace AscendC {22namespace AscendC {
Mmatmul/quant_batch_matmul_v3/op_kernel/quant_batch_matmul_v3_apt.cpp+2-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -30,7 +30,7 @@
30#if (ORIG_DTYPE_SCALE == DT_FLOAT8_E8M0)30#if (ORIG_DTYPE_SCALE == DT_FLOAT8_E8M0)
31#include "arch35/qbmm_mx_basic_api_cmct.h"31#include "arch35/qbmm_mx_basic_api_cmct.h"
32#endif32#endif
33-#include "kernel_operator.h"33+#include "kernel_basic_intf.h"
34 34 
35#ifdef __CCE_KT_TEST__35#ifdef __CCE_KT_TEST__
36#define UT_STATIC static36#define UT_STATIC static
Mmatmul/quant_batch_matmul_v3/op_kernel/quant_batch_matmul_v3_base.h+5-3
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,8 +16,10 @@
16#define QUANT_BATCH_MATMUL_V3_BASE_H16#define QUANT_BATCH_MATMUL_V3_BASE_H
17 17 
18#include <cstdint>18#include <cstdint>
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20-#include "kernel_operator_intf.h"20+#include "quantization/ascend_dequant_utils.h"
21+#include "pad/broadcast.h"
22+#include "quantization/ascend_dequant.h"
21#include "kernel_type.h"23#include "kernel_type.h"
22#include "lib/matmul_intf.h"24#include "lib/matmul_intf.h"
23#include "quant_batch_matmul_v3_kernel_tiling_data.h"25#include "quant_batch_matmul_v3_kernel_tiling_data.h"
Mmatmul/quant_batch_matmul_v3/op_kernel/quant_batch_matmul_v3_basic_epilogue.h+3-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,6 +16,8 @@
16#define QUANT_BATCH_MATMUL_V3_BASIC_EPILOGUE_H16#define QUANT_BATCH_MATMUL_V3_BASIC_EPILOGUE_H
17 17 
18#include "quant_batch_matmul_v3_base.h"18#include "quant_batch_matmul_v3_base.h"
19+#include "kernel_operator.h"
20+#include "kernel_operator_intf.h"
19 21 
20namespace AscendC {22namespace AscendC {
21 23 
Mmatmul/quant_batch_matmul_v3/op_kernel/quant_batch_matmul_v3_bf16.h+2-0
@@ -16,6 +16,8 @@
16#define QUANT_BATCH_MATMUL_V3_BF16_H16#define QUANT_BATCH_MATMUL_V3_BF16_H
17 17 
18#include "quant_batch_matmul_v3_base.h"18#include "quant_batch_matmul_v3_base.h"
19+#include "kernel_operator.h"
20+#include "kernel_operator_intf.h"
19 21 
20namespace AscendC {22namespace AscendC {
21 23 
Mmatmul/quant_batch_matmul_v3/op_kernel/quant_batch_matmul_v3_bf16_opt.h+3-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,6 +16,8 @@
16#define QUANT_BATCH_MATMUL_V3_BF16_OPT_H16#define QUANT_BATCH_MATMUL_V3_BF16_OPT_H
17 17 
18#include "quant_batch_matmul_v3_base.h"18#include "quant_batch_matmul_v3_base.h"
19+#include "kernel_operator.h"
20+#include "kernel_operator_intf.h"
19 21 
20namespace AscendC {22namespace AscendC {
21 23 
Mmatmul/quant_batch_matmul_v3/op_kernel/quant_batch_matmul_v3_block.h+3-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,6 +16,8 @@
16#define QUANT_BATCH_MATMUL_V3_BLOCK_H16#define QUANT_BATCH_MATMUL_V3_BLOCK_H
17 17 
18#include "quant_batch_matmul_v3_base.h"18#include "quant_batch_matmul_v3_base.h"
19+#include "kernel_operator.h"
20+#include "kernel_operator_intf.h"
19 21 
20namespace AscendC {22namespace AscendC {
21class QuantBatchMatmulV3BaseBlock {23class QuantBatchMatmulV3BaseBlock {
Mmatmul/quant_batch_matmul_v3/op_kernel/quant_batch_matmul_v3_init_output.h+3-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,6 +16,8 @@
16#define QUANT_BATCH_MATMUL_V3_INIT_OUTPUT_H16#define QUANT_BATCH_MATMUL_V3_INIT_OUTPUT_H
17 17 
18#include "quant_batch_matmul_v3_base.h"18#include "quant_batch_matmul_v3_base.h"
19+#include "kernel_operator.h"
20+#include "kernel_operator_intf.h"
19 21 
20#if (defined(__NPU_ARCH__) && __NPU_ARCH__ == 3003)22#if (defined(__NPU_ARCH__) && __NPU_ARCH__ == 3003)
21#define SYNC_AIC_FLAG 123#define SYNC_AIC_FLAG 1
Mmatmul/quant_batch_matmul_v3/op_kernel/quant_batch_matmul_v3_pertoken.h+2-0
@@ -16,6 +16,8 @@
16#define QUANT_BATCH_MATMUL_V3_PERTOKEN_H16#define QUANT_BATCH_MATMUL_V3_PERTOKEN_H
17 17 
18#include "quant_batch_matmul_v3_base.h"18#include "quant_batch_matmul_v3_base.h"
19+#include "kernel_operator.h"
20+#include "kernel_operator_intf.h"
19 21 
20namespace AscendC {22namespace AscendC {
21 23 
Mmatmul/quant_batch_matmul_v3/op_kernel/quant_batch_matmul_v3_pertoken_opt.h+3-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,6 +16,8 @@
16#define QUANT_BATCH_MATMUL_V3_PERTOKEN_OPT_H16#define QUANT_BATCH_MATMUL_V3_PERTOKEN_OPT_H
17 17 
18#include "quant_batch_matmul_v3_base.h"18#include "quant_batch_matmul_v3_base.h"
19+#include "kernel_operator.h"
20+#include "kernel_operator_intf.h"
19 21 
20namespace AscendC {22namespace AscendC {
21 23 
Mmatmul/quant_batch_matmul_v3/op_kernel/quant_batch_matmul_v3_update.h+3-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,6 +16,8 @@
16#define QUANT_BATCH_MATMUL_V3_UPDATE_H16#define QUANT_BATCH_MATMUL_V3_UPDATE_H
17 17 
18#include "quant_batch_matmul_v3_base.h"18#include "quant_batch_matmul_v3_base.h"
19+#include "kernel_operator.h"
20+#include "kernel_operator_intf.h"
19 21 
20namespace AscendC {22namespace AscendC {
21 23 
Mmatmul/quant_batch_matmul_v4/op_kernel/arch35/cmct_convertor.h+1-1
@@ -22,7 +22,7 @@
22#include "cmct/policy/dispatch_policy.h"22#include "cmct/policy/dispatch_policy.h"
23#include "cmct/utils/gemm_type.h"23#include "cmct/utils/gemm_type.h"
24#include "cmct/utils/integral_constant.h"24#include "cmct/utils/integral_constant.h"
25-#include "kernel_operator.h"25+#include "kernel_basic_intf.h"
26#include "lib/std/type_traits.h"26#include "lib/std/type_traits.h"
27#include "quant_batch_matmul_v4_tiling_data.h"27#include "quant_batch_matmul_v4_tiling_data.h"
28 28 
Mmatmul/quant_batch_matmul_v4/op_kernel/arch35/quant_batch_matmul_v4_vf.h+2-3
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,8 +16,7 @@
16 #ifndef QUANT_BATCH_MATMUL_V4_VF_H16 #ifndef QUANT_BATCH_MATMUL_V4_VF_H
17 #define QUANT_BATCH_MATMUL_V4_VF_H17 #define QUANT_BATCH_MATMUL_V4_VF_H
18 18 
19- #include "kernel_operator.h"19+ #include "kernel_basic_intf.h"
20- #include "kernel_operator_intf.h"
21 20 
22 namespace MicroAPI = AscendC::MicroAPI;21 namespace MicroAPI = AscendC::MicroAPI;
23 22 
Mmatmul/quant_batch_matmul_v4/op_kernel/arch35/tool.h+2-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,7 +16,7 @@
16#define TOOL_H16#define TOOL_H
17 17 
18#include "kernel_log.h"18#include "kernel_log.h"
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20#include "kernel_utils.h"20#include "kernel_utils.h"
21 21 
22using AscendC::AIC;22using AscendC::AIC;
Mmatmul/quant_batch_matmul_v4/op_kernel/quant_batch_matmul_v4_apt.cpp+2-3
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -14,8 +14,7 @@
14 */14 */
15 15 
16#define K_MAX_SHAPE_DIM 016#define K_MAX_SHAPE_DIM 0
17-#include "kernel_operator.h"17+#include "kernel_basic_intf.h"
18-#include "kernel_operator_intf.h"
19// if run with ttk without bias, can't get DTYPE_BIAS macro18// if run with ttk without bias, can't get DTYPE_BIAS macro
20#ifndef DTYPE_BIAS19#ifndef DTYPE_BIAS
21#define DTYPE_BIAS DTYPE_Y20#define DTYPE_BIAS DTYPE_Y
Mmatmul/transpose_batch_mat_mul/op_kernel/arch35/transpose_batch_mat_mul_asw_block_advanced.h+2-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -15,7 +15,7 @@
15#ifndef TRANSPOSE_BATCH_MAT_MUL_ASW_BLOCK_ADVANCED_H15#ifndef TRANSPOSE_BATCH_MAT_MUL_ASW_BLOCK_ADVANCED_H
16#define TRANSPOSE_BATCH_MAT_MUL_ASW_BLOCK_ADVANCED_H16#define TRANSPOSE_BATCH_MAT_MUL_ASW_BLOCK_ADVANCED_H
17 17 
18-#include "kernel_operator.h"18+#include "kernel_basic_intf.h"
19#include "lib/matmul_intf.h"19#include "lib/matmul_intf.h"
20#include "../../mat_mul_v3/mat_mul_v3_common.h"20#include "../../mat_mul_v3/mat_mul_v3_common.h"
21#include "../../mat_mul_v3/arch35/mat_mul_tiling_data.h"21#include "../../mat_mul_v3/arch35/mat_mul_tiling_data.h"
Mmatmul/transpose_batch_mat_mul/op_kernel/arch35/transpose_batch_mat_mul_asw_kernel_advanced.h+2-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,7 +16,7 @@
16#define TRANSPOSE_BATCH_MAT_MUL_ASW_KERNEL_ADVANCED_H16#define TRANSPOSE_BATCH_MAT_MUL_ASW_KERNEL_ADVANCED_H
17 17 
18#include "transpose_batch_mat_mul_asw_block_advanced.h"18#include "transpose_batch_mat_mul_asw_block_advanced.h"
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20#include "lib/matmul_intf.h"20#include "lib/matmul_intf.h"
21 21 
22namespace TransposeBatchMatMulAdvanced {22namespace TransposeBatchMatMulAdvanced {
Mmatmul/weight_quant_batch_matmul_v2/op_kernel/arch35/n_first/anti_quant_y_vf.h+2-3
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,8 +16,7 @@
16#define WEIGHT_QUANT_BATCHMATMUL_V2_ANTI_QUANT_Y_VF_H16#define WEIGHT_QUANT_BATCHMATMUL_V2_ANTI_QUANT_Y_VF_H
17 17 
18#include "basic_block_config.h"18#include "basic_block_config.h"
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20-#include "kernel_operator_intf.h"
21 20 
22namespace MicroAPI = AscendC::MicroAPI;21namespace MicroAPI = AscendC::MicroAPI;
23using AscendC::MicroAPI::RegTensor;22using AscendC::MicroAPI::RegTensor;
Mmatmul/weight_quant_batch_matmul_v2/op_kernel/arch35/n_first/basic_block_config.h+2-3
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -17,8 +17,7 @@
17 17 
18#include "../tool_arch35.h"18#include "../tool_arch35.h"
19#include "../../weight_quant_batch_matmul_v2_constant.h"19#include "../../weight_quant_batch_matmul_v2_constant.h"
20-#include "kernel_operator.h"20+#include "kernel_basic_intf.h"
21-#include "kernel_operator_intf.h"
22#include "lib/matmul_intf.h"21#include "lib/matmul_intf.h"
23 22 
24namespace WeightQuantBatchMatmulV2::Arch35 {23namespace WeightQuantBatchMatmulV2::Arch35 {
Mmatmul/weight_quant_batch_matmul_v2/op_kernel/arch35/n_first/basic_block_vf_mx.h+2-3
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,8 +16,7 @@
16#define WEIGHT_QUANT_BATCHMATMUL_V2_BASIC_BLOCK_VF_MX_H16#define WEIGHT_QUANT_BATCHMATMUL_V2_BASIC_BLOCK_VF_MX_H
17 17 
18#include "basic_block_config.h"18#include "basic_block_config.h"
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20-#include "kernel_operator_intf.h"
21 20 
22namespace MicroAPI = AscendC::MicroAPI;21namespace MicroAPI = AscendC::MicroAPI;
23using AscendC::BLOCK_CUBE;22using AscendC::BLOCK_CUBE;
Mmatmul/weight_quant_batch_matmul_v2/op_kernel/arch35/n_first/basic_block_vf_nd.h+2-3
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,8 +16,7 @@
16#define WEIGHT_QUANT_BATCHMATMUL_V2_BASIC_BLOCK_VF_ND_H16#define WEIGHT_QUANT_BATCHMATMUL_V2_BASIC_BLOCK_VF_ND_H
17 17 
18#include "basic_block_config.h"18#include "basic_block_config.h"
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20-#include "kernel_operator_intf.h"
21 20 
22namespace MicroAPI = AscendC::MicroAPI;21namespace MicroAPI = AscendC::MicroAPI;
23using AscendC::VECTOR_REG_WIDTH;22using AscendC::VECTOR_REG_WIDTH;
Mmatmul/weight_quant_batch_matmul_v2/op_kernel/arch35/n_first/basic_block_vf_nz.h+2-3
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,8 +16,7 @@
16#define WEIGHT_QUANT_BATCHMATMUL_V2_BASIC_BLOCK_VF_NZ_H16#define WEIGHT_QUANT_BATCHMATMUL_V2_BASIC_BLOCK_VF_NZ_H
17 17 
18#include "basic_block_config.h"18#include "basic_block_config.h"
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20-#include "kernel_operator_intf.h"
21 20 
22namespace MicroAPI = AscendC::MicroAPI;21namespace MicroAPI = AscendC::MicroAPI;
23using AscendC::BLOCK_CUBE;22using AscendC::BLOCK_CUBE;
Mmatmul/weight_quant_batch_matmul_v2/op_kernel/arch35/n_first/weight_quant_batch_matmul_v2_basic_block.h+2-3
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,8 +16,7 @@
16#define WEIGHT_QUANT_BATCHMATMUL_V2_BASIC_BLOCK_H16#define WEIGHT_QUANT_BATCHMATMUL_V2_BASIC_BLOCK_H
17 17 
18#include "basic_block_config.h"18#include "basic_block_config.h"
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20-#include "kernel_operator_intf.h"
21#include "lib/matmul_intf.h"20#include "lib/matmul_intf.h"
22#include "weight_quant_batch_matmul_v2_cube_compute.h"21#include "weight_quant_batch_matmul_v2_cube_compute.h"
23#include "weight_quant_batch_matmul_v2_vec_compute.h"22#include "weight_quant_batch_matmul_v2_vec_compute.h"
Mmatmul/weight_quant_batch_matmul_v2/op_kernel/arch35/n_first/weight_quant_batch_matmul_v2_cube_compute.h+2-3
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,8 +16,7 @@
16#define WEIGHT_QUANT_BATCHMATMUL_V2_CUBE_COMPUTE_H16#define WEIGHT_QUANT_BATCHMATMUL_V2_CUBE_COMPUTE_H
17 17 
18#include "basic_block_config.h"18#include "basic_block_config.h"
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20-#include "kernel_operator_intf.h"
21#include "lib/matmul_intf.h"20#include "lib/matmul_intf.h"
22 21 
23using AscendC::BLOCK_CUBE;22using AscendC::BLOCK_CUBE;
Mmatmul/weight_quant_batch_matmul_v2/op_kernel/arch35/n_first/weight_quant_batch_matmul_v2_vcv_basic_block.h+2-3
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,8 +16,7 @@
16#define WEIGHT_QUANT_BATCHMATMUL_V2_VCV_BASIC_BLOCK_H16#define WEIGHT_QUANT_BATCHMATMUL_V2_VCV_BASIC_BLOCK_H
17 17 
18#include "basic_block_config.h"18#include "basic_block_config.h"
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20-#include "kernel_operator_intf.h"
21#include "lib/matmul_intf.h"20#include "lib/matmul_intf.h"
22#include "weight_quant_batch_matmul_v2_cube_compute.h"21#include "weight_quant_batch_matmul_v2_cube_compute.h"
23#include "weight_quant_batch_matmul_v2_vec_compute.h"22#include "weight_quant_batch_matmul_v2_vec_compute.h"
Mmatmul/weight_quant_batch_matmul_v2/op_kernel/arch35/n_first/weight_quant_batch_matmul_v2_vec_compute.h+2-3
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -20,8 +20,7 @@
20#include "basic_block_vf_mx.h"20#include "basic_block_vf_mx.h"
21#include "basic_block_vf_nd.h"21#include "basic_block_vf_nd.h"
22#include "basic_block_vf_nz.h"22#include "basic_block_vf_nz.h"
23-#include "kernel_operator.h"23+#include "kernel_basic_intf.h"
24-#include "kernel_operator_intf.h"
25 24 
26using AscendC::BLOCK_CUBE;25using AscendC::BLOCK_CUBE;
27using AscendC::CacheMode;26using AscendC::CacheMode;
Mmatmul/weight_quant_batch_matmul_v2/op_kernel/arch35/tool_arch35.h+2-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -18,7 +18,7 @@
18#include <limits>18#include <limits>
19 19 
20#include "kernel_log.h"20#include "kernel_log.h"
21-#include "kernel_operator.h"21+#include "kernel_basic_intf.h"
22#include "kernel_utils.h"22#include "kernel_utils.h"
23#include "lib/matmul_intf.h"23#include "lib/matmul_intf.h"
24 24 
Mmatmul/weight_quant_batch_matmul_v2/op_kernel/arch35/weight_quant_batch_matmul_v2_adaptive_sliding_window.h+2-3
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -19,8 +19,7 @@
19#include "weight_quant_batch_matmul_v2_asw_block.h"19#include "weight_quant_batch_matmul_v2_asw_block.h"
20#include "../weight_quant_batch_matmul_v2_constant.h"20#include "../weight_quant_batch_matmul_v2_constant.h"
21#include "weight_quant_batch_matmul_v2_arch35_tiling_data.h"21#include "weight_quant_batch_matmul_v2_arch35_tiling_data.h"
22-#include "kernel_operator.h"22+#include "kernel_basic_intf.h"
23-#include "kernel_operator_intf.h"
24#include "lib/matmul_intf.h"23#include "lib/matmul_intf.h"
25#include "../tool.h"24#include "../tool.h"
26 25 
Mmatmul/weight_quant_batch_matmul_v2/op_kernel/arch35/weight_quant_batch_matmul_v2_arch35_tiling_key.h+19-19
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -25,7 +25,7 @@
25#endif25#endif
26#if defined(ORIG_DTYPE_WEIGHT) && \26#if defined(ORIG_DTYPE_WEIGHT) && \
27 ((defined(DT_INT4) && ORIG_DTYPE_WEIGHT == DT_INT4) || (defined(DT_INT32) && ORIG_DTYPE_WEIGHT == DT_INT32))27 ((defined(DT_INT4) && ORIG_DTYPE_WEIGHT == DT_INT4) || (defined(DT_INT32) && ORIG_DTYPE_WEIGHT == DT_INT32))
28-#define S428+#define WQBMMV2_S4
29#undef DTYPE_WEIGHT29#undef DTYPE_WEIGHT
30#define DTYPE_WEIGHT AscendC::int4b_t30#define DTYPE_WEIGHT AscendC::int4b_t
31#undef ORIG_DTYPE_WEIGHT31#undef ORIG_DTYPE_WEIGHT
@@ -45,7 +45,7 @@
45#endif45#endif
46#endif46#endif
47#if defined(ORIG_DTYPE_WEIGHT) && (defined(DT_INT8) && ORIG_DTYPE_WEIGHT == DT_INT8)47#if defined(ORIG_DTYPE_WEIGHT) && (defined(DT_INT8) && ORIG_DTYPE_WEIGHT == DT_INT8)
48-#define S848+#define WQBMMV2_S8
49#endif49#endif
50#if defined(ORIG_DTYPE_X) && defined(DT_BF16) && ORIG_DTYPE_X == DT_BF1650#if defined(ORIG_DTYPE_X) && defined(DT_BF16) && ORIG_DTYPE_X == DT_BF16
51#define X_BF1651#define X_BF16
@@ -244,7 +244,7 @@ ASCENDC_TPL_SEL(
244 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)244 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)
245 ),245 ),
246 #endif246 #endif
247- #if !defined(__CCE_AICORE__) || ((defined(S8) || defined(S4)) && defined(WEIGHT_ND))247+ #if !defined(__CCE_AICORE__) || ((defined(WQBMMV2_S8) || defined(WQBMMV2_S4)) && defined(WEIGHT_ND))
248 ASCENDC_TPL_ARGS_SEL(248 ASCENDC_TPL_ARGS_SEL(
249 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),249 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
250 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),250 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),
@@ -265,7 +265,7 @@ ASCENDC_TPL_SEL(
265 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2RegBaseTilingDataParams)265 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2RegBaseTilingDataParams)
266 ),266 ),
267 #endif267 #endif
268- #if !defined(__CCE_AICORE__) || ((defined(S4) || defined(F4)) && !defined(WEIGHT_ND))268+ #if !defined(__CCE_AICORE__) || ((defined(WQBMMV2_S4) || defined(F4)) && !defined(WEIGHT_ND))
269 ASCENDC_TPL_ARGS_SEL(269 ASCENDC_TPL_ARGS_SEL(
270 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),270 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
271 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),271 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),
@@ -286,7 +286,7 @@ ASCENDC_TPL_SEL(
286 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2RegBaseTilingDataParams)286 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2RegBaseTilingDataParams)
287 ),287 ),
288 #endif288 #endif
289- #if !defined(__CCE_AICORE__) || (defined(S4) && !defined(WEIGHT_ND))289+ #if !defined(__CCE_AICORE__) || (defined(WQBMMV2_S4) && !defined(WEIGHT_ND))
290 ASCENDC_TPL_ARGS_SEL(290 ASCENDC_TPL_ARGS_SEL(
291 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),291 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
292 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),292 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),
@@ -307,7 +307,7 @@ ASCENDC_TPL_SEL(
307 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)307 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)
308 ),308 ),
309 #endif309 #endif
310- #if !defined(__CCE_AICORE__) || (defined(S4) && !defined(WEIGHT_ND))310+ #if !defined(__CCE_AICORE__) || (defined(WQBMMV2_S4) && !defined(WEIGHT_ND))
311 ASCENDC_TPL_ARGS_SEL(311 ASCENDC_TPL_ARGS_SEL(
312 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),312 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
313 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),313 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),
@@ -328,7 +328,7 @@ ASCENDC_TPL_SEL(
328 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)328 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)
329 ),329 ),
330 #endif330 #endif
331- #if !defined(__CCE_AICORE__) || ((defined(S8) || defined(S4)) && defined(WEIGHT_ND))331+ #if !defined(__CCE_AICORE__) || ((defined(WQBMMV2_S8) || defined(WQBMMV2_S4)) && defined(WEIGHT_ND))
332 ASCENDC_TPL_ARGS_SEL(332 ASCENDC_TPL_ARGS_SEL(
333 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),333 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
334 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),334 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),
@@ -349,7 +349,7 @@ ASCENDC_TPL_SEL(
349 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)349 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)
350 ),350 ),
351 #endif351 #endif
352- #if !defined(__CCE_AICORE__) || ((defined(S8) || defined(WEIGHT_F8_INPUT) || defined(S4)) && defined(WEIGHT_ND))352+ #if !defined(__CCE_AICORE__) || ((defined(WQBMMV2_S8) || defined(WEIGHT_F8_INPUT) || defined(WQBMMV2_S4)) && defined(WEIGHT_ND))
353 ASCENDC_TPL_ARGS_SEL(353 ASCENDC_TPL_ARGS_SEL(
354 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),354 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
355 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),355 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),
@@ -370,7 +370,7 @@ ASCENDC_TPL_SEL(
370 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)370 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)
371 ),371 ),
372 #endif372 #endif
373- #if !defined(__CCE_AICORE__) || ((defined(S8) || defined(S4)) && defined(WEIGHT_ND))373+ #if !defined(__CCE_AICORE__) || ((defined(WQBMMV2_S8) || defined(WQBMMV2_S4)) && defined(WEIGHT_ND))
374 ASCENDC_TPL_ARGS_SEL(374 ASCENDC_TPL_ARGS_SEL(
375 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),375 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
376 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),376 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),
@@ -391,7 +391,7 @@ ASCENDC_TPL_SEL(
391 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)391 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)
392 ),392 ),
393 #endif393 #endif
394- #if !defined(__CCE_AICORE__) || ((defined(S8) || defined(WEIGHT_F8_INPUT) || defined(S4)) && defined(WEIGHT_ND))394+ #if !defined(__CCE_AICORE__) || ((defined(WQBMMV2_S8) || defined(WEIGHT_F8_INPUT) || defined(WQBMMV2_S4)) && defined(WEIGHT_ND))
395 ASCENDC_TPL_ARGS_SEL(395 ASCENDC_TPL_ARGS_SEL(
396 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),396 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
397 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),397 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),
@@ -474,7 +474,7 @@ ASCENDC_TPL_SEL(
474 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)474 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)
475 ),475 ),
476 #endif476 #endif
477- #if !defined(__CCE_AICORE__) || ((defined(S8) || defined(S4)) && defined(WEIGHT_ND))477+ #if !defined(__CCE_AICORE__) || ((defined(WQBMMV2_S8) || defined(WQBMMV2_S4)) && defined(WEIGHT_ND))
478 ASCENDC_TPL_ARGS_SEL(478 ASCENDC_TPL_ARGS_SEL(
479 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),479 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
480 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),480 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),
@@ -495,7 +495,7 @@ ASCENDC_TPL_SEL(
495 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2RegBaseTilingDataParams)495 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2RegBaseTilingDataParams)
496 ),496 ),
497 #endif497 #endif
498- #if !defined(__CCE_AICORE__) || ((defined(S4) || defined(F4)) && !defined(WEIGHT_ND))498+ #if !defined(__CCE_AICORE__) || ((defined(WQBMMV2_S4) || defined(F4)) && !defined(WEIGHT_ND))
499 ASCENDC_TPL_ARGS_SEL(499 ASCENDC_TPL_ARGS_SEL(
500 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),500 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
501 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),501 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),
@@ -516,7 +516,7 @@ ASCENDC_TPL_SEL(
516 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2RegBaseTilingDataParams)516 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2RegBaseTilingDataParams)
517 ),517 ),
518 #endif518 #endif
519- #if !defined(__CCE_AICORE__) || (defined(S4) && !defined(WEIGHT_ND))519+ #if !defined(__CCE_AICORE__) || (defined(WQBMMV2_S4) && !defined(WEIGHT_ND))
520 ASCENDC_TPL_ARGS_SEL(520 ASCENDC_TPL_ARGS_SEL(
521 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),521 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
522 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),522 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),
@@ -537,7 +537,7 @@ ASCENDC_TPL_SEL(
537 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)537 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)
538 ),538 ),
539 #endif539 #endif
540- #if !defined(__CCE_AICORE__) || (defined(S4) && !defined(WEIGHT_ND))540+ #if !defined(__CCE_AICORE__) || (defined(WQBMMV2_S4) && !defined(WEIGHT_ND))
541 ASCENDC_TPL_ARGS_SEL(541 ASCENDC_TPL_ARGS_SEL(
542 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),542 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
543 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),543 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),
@@ -558,7 +558,7 @@ ASCENDC_TPL_SEL(
558 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)558 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)
559 ),559 ),
560 #endif560 #endif
561- #if !defined(__CCE_AICORE__) || ((defined(S8) || defined(S4)) && defined(WEIGHT_ND))561+ #if !defined(__CCE_AICORE__) || ((defined(WQBMMV2_S8) || defined(WQBMMV2_S4)) && defined(WEIGHT_ND))
562 ASCENDC_TPL_ARGS_SEL(562 ASCENDC_TPL_ARGS_SEL(
563 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),563 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
564 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),564 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),
@@ -579,7 +579,7 @@ ASCENDC_TPL_SEL(
579 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)579 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)
580 ),580 ),
581 #endif581 #endif
582- #if !defined(__CCE_AICORE__) || ((defined(S8) || defined(WEIGHT_F8_INPUT) || defined(S4)) && defined(WEIGHT_ND))582+ #if !defined(__CCE_AICORE__) || ((defined(WQBMMV2_S8) || defined(WEIGHT_F8_INPUT) || defined(WQBMMV2_S4)) && defined(WEIGHT_ND))
583 ASCENDC_TPL_ARGS_SEL(583 ASCENDC_TPL_ARGS_SEL(
584 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),584 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
585 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),585 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),
@@ -601,7 +601,7 @@ ASCENDC_TPL_SEL(
601 ),601 ),
602 #endif602 #endif
603 603 
604- #if !defined(__CCE_AICORE__) || ((defined(S8) || defined(S4)) && defined(WEIGHT_ND))604+ #if !defined(__CCE_AICORE__) || ((defined(WQBMMV2_S8) || defined(WQBMMV2_S4)) && defined(WEIGHT_ND))
605 ASCENDC_TPL_ARGS_SEL(605 ASCENDC_TPL_ARGS_SEL(
606 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),606 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
607 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),607 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),
@@ -622,7 +622,7 @@ ASCENDC_TPL_SEL(
622 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)622 ASCENDC_TPL_TILING_STRUCT_SEL(wqbmmv2_tiling::WeightQuantBatchMatmulV2ASTilingDataParams)
623 ),623 ),
624 #endif624 #endif
625- #if !defined(__CCE_AICORE__) || ((defined(S8) || defined(WEIGHT_F8_INPUT) || defined(S4)) && defined(WEIGHT_ND))625+ #if !defined(__CCE_AICORE__) || ((defined(WQBMMV2_S8) || defined(WEIGHT_F8_INPUT) || defined(WQBMMV2_S4)) && defined(WEIGHT_ND))
626 ASCENDC_TPL_ARGS_SEL(626 ASCENDC_TPL_ARGS_SEL(
627 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),627 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
628 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),628 ASCENDC_TPL_UINT_SEL(SOC_VERSION_TYPE, ASCENDC_TPL_UI_LIST, WQBMMV2_SOC_SUPPORT_L1_TO_BT_BF16),
Mmatmul/weight_quant_batch_matmul_v2/op_kernel/arch35/weight_quant_batch_matmul_v2_asw_block.h+2-3
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -15,8 +15,7 @@
15#ifndef WEIGHT_QUANT_BMMV2_ASW_BLOCK_H15#ifndef WEIGHT_QUANT_BMMV2_ASW_BLOCK_H
16#define WEIGHT_QUANT_BMMV2_ASW_BLOCK_H16#define WEIGHT_QUANT_BMMV2_ASW_BLOCK_H
17 17 
18-#include "kernel_operator.h"18+#include "kernel_basic_intf.h"
19-#include "kernel_operator_intf.h"
20#include "lib/matmul_intf.h"19#include "lib/matmul_intf.h"
21#include "../tool.h"20#include "../tool.h"
22#include "weight_quant_batch_matmul_v2_arch35_tiling_data.h"21#include "weight_quant_batch_matmul_v2_arch35_tiling_data.h"
Mmatmul/weight_quant_batch_matmul_v2/op_kernel/arch35/weight_quant_batch_matmul_v2_reg_base_common.h+2-3
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,8 +16,7 @@
16#ifndef WEIGHT_QUANT_BATCHMATMUL_V2_REG_BASE_COMMON_H16#ifndef WEIGHT_QUANT_BATCHMATMUL_V2_REG_BASE_COMMON_H
17#define WEIGHT_QUANT_BATCHMATMUL_V2_REG_BASE_COMMON_H17#define WEIGHT_QUANT_BATCHMATMUL_V2_REG_BASE_COMMON_H
18 18 
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20-#include "kernel_operator_intf.h"
21#include "lib/matmul_intf.h"20#include "lib/matmul_intf.h"
22#include "tool_arch35.h"21#include "tool_arch35.h"
23#include "../weight_quant_batch_matmul_v2_constant.h"22#include "../weight_quant_batch_matmul_v2_constant.h"
Mmatmul/weight_quant_batch_matmul_v2/op_kernel/arch35/weight_quant_batch_matmul_v2_vf.h+2-3
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,8 +16,7 @@
16#ifndef WEIGHT_QUANT_BATCHMATMUL_V2_VF_H16#ifndef WEIGHT_QUANT_BATCHMATMUL_V2_VF_H
17#define WEIGHT_QUANT_BATCHMATMUL_V2_VF_H17#define WEIGHT_QUANT_BATCHMATMUL_V2_VF_H
18 18 
19-#include "kernel_operator.h"19+#include "kernel_basic_intf.h"
20-#include "kernel_operator_intf.h"
21 20 
22namespace MicroAPI = AscendC::MicroAPI;21namespace MicroAPI = AscendC::MicroAPI;
23 22 
Mmatmul/weight_quant_batch_matmul_v2/op_kernel/tool.h+2-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -17,7 +17,7 @@
17 17 
18#include <limits>18#include <limits>
19#include "kernel_log.h"19#include "kernel_log.h"
20-#include "kernel_operator.h"20+#include "kernel_basic_intf.h"
21#include "kernel_utils.h"21#include "kernel_utils.h"
22#include "lib/matmul_intf.h"22#include "lib/matmul_intf.h"
23 23 
Mmatmul/weight_quant_batch_matmul_v2/op_kernel/weight_quant_batch_matmul_v2.cpp+3-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -18,6 +18,8 @@
18#if defined(__CCE_AICORE__) && __CCE_AICORE__ == 220 || (defined(__NPU_ARCH__) && __NPU_ARCH__ == 3003)18#if defined(__CCE_AICORE__) && __CCE_AICORE__ == 220 || (defined(__NPU_ARCH__) && __NPU_ARCH__ == 3003)
19#include "weight_quant_batch_matmul_v2_constant.h"19#include "weight_quant_batch_matmul_v2_constant.h"
20#include "tool.h"20#include "tool.h"
21+#include "kernel_operator.h"
22+#include "kernel_operator_intf.h"
21#if ( \23#if ( \
22 defined(ORIG_DTYPE_ANTIQUANT_SCALE) && \24 defined(ORIG_DTYPE_ANTIQUANT_SCALE) && \
23 ((ORIG_DTYPE_ANTIQUANT_SCALE == DT_UINT64) || (ORIG_DTYPE_ANTIQUANT_SCALE == DT_INT64)))25 ((ORIG_DTYPE_ANTIQUANT_SCALE == DT_UINT64) || (ORIG_DTYPE_ANTIQUANT_SCALE == DT_INT64)))
Mmatmul/weight_quant_batch_matmul_v2/op_kernel/weight_quant_batch_matmul_v2_apt.cpp+6-7
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -15,20 +15,19 @@
15 15 
16#define ENABLE_L2_CACHE16#define ENABLE_L2_CACHE
17#include "weight_quant_batch_matmul_v2_constant.h"17#include "weight_quant_batch_matmul_v2_constant.h"
18-#include "kernel_operator.h"18+#include "kernel_basic_intf.h"
19-#include "kernel_operator_intf.h"
20#include "lib/matmul_intf.h"19#include "lib/matmul_intf.h"
21#include "arch35/weight_quant_batch_matmul_v2_arch35_tiling_key.h"20#include "arch35/weight_quant_batch_matmul_v2_arch35_tiling_key.h"
22#include "arch35/weight_quant_batch_matmul_v2_arch35_tiling_data.h"21#include "arch35/weight_quant_batch_matmul_v2_arch35_tiling_data.h"
23 22 
24#if !defined(__DAV_310R6__) && !(defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102))23#if !defined(__DAV_310R6__) && !(defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102))
25-#if defined(A16) && (defined(S4) || defined(F4) || defined(S8))24+#if defined(A16) && (defined(WQBMMV2_S4) || defined(F4) || defined(WQBMMV2_S8))
26#include "arch35/weight_quant_batch_matmul_v2_reg_base.h"25#include "arch35/weight_quant_batch_matmul_v2_reg_base.h"
27#endif26#endif
28 27 
29#include "arch35/n_first/weight_quant_batch_matmul_v2_basic_block_controller.h"28#include "arch35/n_first/weight_quant_batch_matmul_v2_basic_block_controller.h"
30 29 
31-#if defined(A16) && (defined(S4) || defined(S8) || defined(WEIGHT_F8_INPUT) || defined(MICROSCALING))30+#if defined(A16) && (defined(WQBMMV2_S4) || defined(WQBMMV2_S8) || defined(WEIGHT_F8_INPUT) || defined(MICROSCALING))
32#include "arch35/cmct_convertor.h"31#include "arch35/cmct_convertor.h"
33using WeightQuantBatchMatmulV2::InvokeKernel;32using WeightQuantBatchMatmulV2::InvokeKernel;
34#endif33#endif
@@ -116,13 +115,13 @@ __global__ __aicore__ void weight_quant_batch_matmul_v2(
116 using DtypeBias = AscendC::Std::conditional_t<IS_BIAS_FP32, float, DTYPE_X>;115 using DtypeBias = AscendC::Std::conditional_t<IS_BIAS_FP32, float, DTYPE_X>;
117 static constexpr QuantType antiquantType = static_cast<QuantType>(ANTIQUANT_TYPE);116 static constexpr QuantType antiquantType = static_cast<QuantType>(ANTIQUANT_TYPE);
118 if constexpr (SUB_ALGORITHM == WQBMMV2_SUB_ALGO_ANTI_REG_SCSC) {117 if constexpr (SUB_ALGORITHM == WQBMMV2_SUB_ALGO_ANTI_REG_SCSC) {
119-#if defined(A16) && (defined(S4) || defined(F4) || defined(S8))118+#if defined(A16) && (defined(WQBMMV2_S4) || defined(F4) || defined(WQBMMV2_S8))
120 INVOKE_WEIGHT_QUANT_BMM_OP_REG_BASE_IMPL(119 INVOKE_WEIGHT_QUANT_BMM_OP_REG_BASE_IMPL(
121 DtypeBias, WeightQuantBatchMatmulV2RegBaseKernel, TRANS_A, TRANS_B, HAS_ANTIQUANT_OFFSET, antiquantType,120 DtypeBias, WeightQuantBatchMatmulV2RegBaseKernel, TRANS_A, TRANS_B, HAS_ANTIQUANT_OFFSET, antiquantType,
122 IS_WEIGHT_NZ);121 IS_WEIGHT_NZ);
123#endif122#endif
124 } else {123 } else {
125-#if defined(A16) && (defined(S4) || defined(S8) || defined(WEIGHT_F8_INPUT) || defined(MICROSCALING))124+#if defined(A16) && (defined(WQBMMV2_S4) || defined(WQBMMV2_S8) || defined(WEIGHT_F8_INPUT) || defined(MICROSCALING))
126 InvokeKernel<125 InvokeKernel<
127 TEMPLATE_CUSTOM, TRANS_A, TRANS_B, ANTIQUANT_TYPE, HAS_ANTIQUANT_OFFSET, IS_BIAS_FP32, IS_WEIGHT_NZ>(126 TEMPLATE_CUSTOM, TRANS_A, TRANS_B, ANTIQUANT_TYPE, HAS_ANTIQUANT_OFFSET, IS_BIAS_FP32, IS_WEIGHT_NZ>(
128 KERNEL_PARAMS);127 KERNEL_PARAMS);