已合并
dynamic_mx_quant fix oom #7291
季骏创建于 7月9日
dynamic_mx_quant fix oom #7291
已合并
共 2 个文件变更+12-9
| @@ -1,5 +1,5 @@ | |||
| 1 | /** | 1 | /** |
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | 2 | + * Copyright (c) 2026-2026 Huawei Technologies Co., Ltd. |
| 3 | * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | 3 | * 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. |
| @@ -157,31 +157,36 @@ ASCENDC_TPL_SEL( | |||
| 157 | ASCENDC_TPL_UINT_SEL(scaleAlg, ASCENDC_TPL_UI_LIST, TPL_SCALE_ALG_0, TPL_SCALE_ALG_1, | 157 | ASCENDC_TPL_UINT_SEL(scaleAlg, ASCENDC_TPL_UI_LIST, TPL_SCALE_ALG_0, TPL_SCALE_ALG_1, |
| 158 | TPL_SCALE_ALG_2_SPEC_DST_MAX_VALUE), | 158 | TPL_SCALE_ALG_2_SPEC_DST_MAX_VALUE), |
| 159 | ASCENDC_TPL_UINT_SEL(roundMode, ASCENDC_TPL_UI_LIST, TPL_NOT_CARE_ROUND_MODE), | 159 | ASCENDC_TPL_UINT_SEL(roundMode, ASCENDC_TPL_UI_LIST, TPL_NOT_CARE_ROUND_MODE), |
| 160 | - ASCENDC_TPL_UINT_SEL(isOddScale, ASCENDC_TPL_UI_LIST, TPL_NOT_CARE_SCALE)), | 160 | + ASCENDC_TPL_UINT_SEL(isOddScale, ASCENDC_TPL_UI_LIST, TPL_NOT_CARE_SCALE), |
| 161 | + ASCENDC_TPL_TILING_STRUCT_SEL(DynamicMxQuantTailAxisTilingData)), | ||
| 161 | // 尾轴量化，blocksize!=32量化 | 162 | // 尾轴量化，blocksize!=32量化 |
| 162 | ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(optiMode, ASCENDC_TPL_UI_LIST, TPL_TAIL_AXIS_QUANT_NORMAL), | 163 | ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(optiMode, ASCENDC_TPL_UI_LIST, TPL_TAIL_AXIS_QUANT_NORMAL), |
| 163 | ASCENDC_TPL_UINT_SEL(scaleAlg, ASCENDC_TPL_UI_LIST, TPL_NOT_CARE_SCALE_ALG), | 164 | ASCENDC_TPL_UINT_SEL(scaleAlg, ASCENDC_TPL_UI_LIST, TPL_NOT_CARE_SCALE_ALG), |
| 164 | ASCENDC_TPL_UINT_SEL(roundMode, ASCENDC_TPL_UI_LIST, TPL_NOT_CARE_ROUND_MODE), | 165 | ASCENDC_TPL_UINT_SEL(roundMode, ASCENDC_TPL_UI_LIST, TPL_NOT_CARE_ROUND_MODE), |
| 165 | - ASCENDC_TPL_UINT_SEL(isOddScale, ASCENDC_TPL_UI_LIST, TPL_ODD_SCALE, TPL_EVEN_SCALE)), | 166 | + ASCENDC_TPL_UINT_SEL(isOddScale, ASCENDC_TPL_UI_LIST, TPL_ODD_SCALE, TPL_EVEN_SCALE), |
| 167 | + ASCENDC_TPL_TILING_STRUCT_SEL(DynamicMxQuantTilingData)), | ||
| 166 | // 非尾轴量化，blocksize=32，小尾轴优化模板 | 168 | // 非尾轴量化，blocksize=32，小尾轴优化模板 |
| 167 | ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(optiMode, ASCENDC_TPL_UI_LIST, TPL_NOT_TAIL_AXIS_QUANT_SMALL_OPTI), | 169 | ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(optiMode, ASCENDC_TPL_UI_LIST, TPL_NOT_TAIL_AXIS_QUANT_SMALL_OPTI), |
| 168 | ASCENDC_TPL_UINT_SEL(scaleAlg, ASCENDC_TPL_UI_LIST, TPL_SCALE_ALG_0, TPL_SCALE_ALG_1, | 170 | ASCENDC_TPL_UINT_SEL(scaleAlg, ASCENDC_TPL_UI_LIST, TPL_SCALE_ALG_0, TPL_SCALE_ALG_1, |
| 169 | TPL_SCALE_ALG_2_SPEC_DST_MAX_VALUE, TPL_SCALE_ALG_2_OTHERS), | 171 | TPL_SCALE_ALG_2_SPEC_DST_MAX_VALUE, TPL_SCALE_ALG_2_OTHERS), |
| 170 | ASCENDC_TPL_UINT_SEL(roundMode, ASCENDC_TPL_UI_LIST, TPL_RINT, TPL_ROUND, TPL_FLOOR), | 172 | ASCENDC_TPL_UINT_SEL(roundMode, ASCENDC_TPL_UI_LIST, TPL_RINT, TPL_ROUND, TPL_FLOOR), |
| 171 | - ASCENDC_TPL_UINT_SEL(isOddScale, ASCENDC_TPL_UI_LIST, TPL_NOT_CARE_SCALE)), | 173 | + ASCENDC_TPL_UINT_SEL(isOddScale, ASCENDC_TPL_UI_LIST, TPL_NOT_CARE_SCALE), |
| 174 | + ASCENDC_TPL_TILING_STRUCT_SEL(DynamicMxQuant4OptimizeTilingData)), | ||
| 172 | // 非尾轴量化，大尾轴优化模板 | 175 | // 非尾轴量化，大尾轴优化模板 |
| 173 | ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(optiMode, ASCENDC_TPL_UI_LIST, TPL_NOT_TAIL_AXIS_QUANT_LARGE_OPTI), | 176 | ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(optiMode, ASCENDC_TPL_UI_LIST, TPL_NOT_TAIL_AXIS_QUANT_LARGE_OPTI), |
| 174 | ASCENDC_TPL_UINT_SEL(scaleAlg, ASCENDC_TPL_UI_LIST, TPL_SCALE_ALG_0, TPL_SCALE_ALG_1, | 177 | ASCENDC_TPL_UINT_SEL(scaleAlg, ASCENDC_TPL_UI_LIST, TPL_SCALE_ALG_0, TPL_SCALE_ALG_1, |
| 175 | TPL_SCALE_ALG_2_SPEC_DST_MAX_VALUE, TPL_SCALE_ALG_2_OTHERS), | 178 | TPL_SCALE_ALG_2_SPEC_DST_MAX_VALUE, TPL_SCALE_ALG_2_OTHERS), |
| 176 | ASCENDC_TPL_UINT_SEL(roundMode, ASCENDC_TPL_UI_LIST, TPL_RINT, TPL_ROUND, TPL_FLOOR), | 179 | ASCENDC_TPL_UINT_SEL(roundMode, ASCENDC_TPL_UI_LIST, TPL_RINT, TPL_ROUND, TPL_FLOOR), |
| 177 | - ASCENDC_TPL_UINT_SEL(isOddScale, ASCENDC_TPL_UI_LIST, TPL_NOT_CARE_SCALE)), | 180 | + ASCENDC_TPL_UINT_SEL(isOddScale, ASCENDC_TPL_UI_LIST, TPL_NOT_CARE_SCALE), |
| 181 | + ASCENDC_TPL_TILING_STRUCT_SEL(DynamicMxQuant4OptimizeTilingData)), | ||
| 178 | // 非尾轴，普通模板和多行处理模板 | 182 | // 非尾轴，普通模板和多行处理模板 |
| 179 | ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(optiMode, ASCENDC_TPL_UI_LIST, TPL_NOT_TAIL_AXIS_QUANT_NORMAL, | 183 | ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(optiMode, ASCENDC_TPL_UI_LIST, TPL_NOT_TAIL_AXIS_QUANT_NORMAL, |
| 180 | TPL_NOT_TAIL_AXIS_QUANT_OPTI), | 184 | TPL_NOT_TAIL_AXIS_QUANT_OPTI), |
| 181 | ASCENDC_TPL_UINT_SEL(scaleAlg, ASCENDC_TPL_UI_LIST, TPL_NOT_CARE_SCALE_ALG), | 185 | ASCENDC_TPL_UINT_SEL(scaleAlg, ASCENDC_TPL_UI_LIST, TPL_NOT_CARE_SCALE_ALG), |
| 182 | ASCENDC_TPL_UINT_SEL(roundMode, ASCENDC_TPL_UI_LIST, TPL_NOT_CARE_ROUND_MODE), | 186 | ASCENDC_TPL_UINT_SEL(roundMode, ASCENDC_TPL_UI_LIST, TPL_NOT_CARE_ROUND_MODE), |
| 183 | - ASCENDC_TPL_UINT_SEL(isOddScale, ASCENDC_TPL_UI_LIST, TPL_ODD_SCALE, TPL_EVEN_SCALE)), ); | 187 | + ASCENDC_TPL_UINT_SEL(isOddScale, ASCENDC_TPL_UI_LIST, TPL_ODD_SCALE, TPL_EVEN_SCALE), |
| 188 | + ASCENDC_TPL_TILING_STRUCT_SEL(DynamicMxQuantTilingData))); | ||
| 184 | 189 | ||
| 185 | } // namespace DynamicMxQuant | 190 | } // namespace DynamicMxQuant |
| 186 | 191 | ||
| 187 | -#endif // _DYNAMIC_MX_QUANT_STRUCT_H_ | 192 | +#endif // _DYNAMIC_MX_QUANT_STRUCT_H_ |
| @@ -139,8 +139,6 @@ template <uint64_t optiMode, uint64_t scaleAlg, uint64_t roundMode, uint64_t isO | |||
| 139 | __global__ __aicore__ void dynamic_mx_quant(GM_ADDR x, GM_ADDR y, GM_ADDR mxScale, GM_ADDR workspace, GM_ADDR tiling) | 139 | __global__ __aicore__ void dynamic_mx_quant(GM_ADDR x, GM_ADDR y, GM_ADDR mxScale, GM_ADDR workspace, GM_ADDR tiling) |
| 140 | { | 140 | { |
| 141 | REGISTER_TILING_DEFAULT(DynamicMxQuant4OptimizeTilingData); | 141 | REGISTER_TILING_DEFAULT(DynamicMxQuant4OptimizeTilingData); |
| 142 | - REGISTER_TILING_FOR_TILINGKEY("NOT TAIL AXIS TILING DATA", DynamicMxQuantTilingData); | ||
| 143 | - REGISTER_TILING_FOR_TILINGKEY("TAIL AXIS WITH BLOCKSIZE 32 TILING DATA", DynamicMxQuantTailAxisTilingData); | ||
| 144 | KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIV_1_0); | 142 | KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIV_1_0); |
| 145 | 143 | ||
| 146 | #if (__NPU_ARCH__ == 3510) | 144 | #if (__NPU_ARCH__ == 3510) |