已合并
mmad样例优化重构 #1454
wx369创建于 4月7日
mmad样例优化重构 #1454
已合并
共 11 个文件变更+686-205
| @@ -22,7 +22,7 @@ | |||
| 22 | | [load_data_with_transpose_b16](./load_data_with_transpose_b16) | 本样例介绍基础API LoadDataWithTranspose b16数据类型下的使用 | | 22 | | [load_data_with_transpose_b16](./load_data_with_transpose_b16) | 本样例介绍基础API LoadDataWithTranspose b16数据类型下的使用 | |
| 23 | | [load_data_with_transpose_b32](./load_data_with_transpose_b32) | 本样例介绍基础API LoadDataWithTranspose b32数据类型下的使用 | | 23 | | [load_data_with_transpose_b32](./load_data_with_transpose_b32) | 本样例介绍基础API LoadDataWithTranspose b32数据类型下的使用 | |
| 24 | | [batch_mmad](./batch_mmad) | 本样例介绍在输入为float数据类型并且左、右矩阵均不转置的场景下，带batch的矩阵乘法，其中从GM-->L1、L0C-->GM、L0C-->L1这三条通路分别采用了DataCopy ND2NZ和Fixpipe批量搬运数据，从L1-->L0A/L0B以及Mmad执行矩阵乘这两个步骤则是循环batch次，每次循环内只处理一对左、右矩阵 | | 24 | | [batch_mmad](./batch_mmad) | 本样例介绍在输入为float数据类型并且左、右矩阵均不转置的场景下，带batch的矩阵乘法，其中从GM-->L1、L0C-->GM、L0C-->L1这三条通路分别采用了DataCopy ND2NZ和Fixpipe批量搬运数据，从L1-->L0A/L0B以及Mmad执行矩阵乘这两个步骤则是循环batch次，每次循环内只处理一对左、右矩阵 | |
| 25 | -| [mmad](./mmad) | 本样例介绍基于基础API Mmad实现矩阵乘 | | 25 | +| [mmad](./mmad) | 本样例介绍了输入为ND格式，B4 / B8 / B16 / B32输入数据类型（具体以int4_t / int8_t / half / float为例），四种输入数据类型下的矩阵乘法，说明如何通过Mmad指令实现矩阵乘法计算（C = A x B + Bias） | |
| 26 | | [mmad_load3dv2](./mmad_load3dv2) | 本样例介绍LoadData3DV2指令将A、B矩阵从L1搬运到L0A/L0B的过程，其中 A 和 B 分别表示矩阵乘法的左右输入矩阵。LoadData3DV2指令参数配置及执行指令前后各个矩阵数据排布变化，均配合示意图进行了说明 | | 26 | | [mmad_load3dv2](./mmad_load3dv2) | 本样例介绍LoadData3DV2指令将A、B矩阵从L1搬运到L0A/L0B的过程，其中 A 和 B 分别表示矩阵乘法的左右输入矩阵。LoadData3DV2指令参数配置及执行指令前后各个矩阵数据排布变化，均配合示意图进行了说明 | |
| 27 | | [mmad_s8_f16_f32_with_A_B_transpose_option](./mmad_s8_f16_f32_with_A_B_transpose_option) | 本样例介绍了在 int8_t / half / float 三种数据类型下，以及左、右矩阵均不转置 / 左矩阵不转置、右矩阵转置 / 左矩阵转置、右矩阵不转置 / 左、右矩阵均转置 共 12 种矩阵乘法场景中，相关指令的使用方法，其中 A 和 B 分别表示矩阵乘法的左右输入矩阵。| | 27 | | [mmad_s8_f16_f32_with_A_B_transpose_option](./mmad_s8_f16_f32_with_A_B_transpose_option) | 本样例介绍了在 int8_t / half / float 三种数据类型下，以及左、右矩阵均不转置 / 左矩阵不转置、右矩阵转置 / 左矩阵转置、右矩阵不转置 / 左、右矩阵均转置 共 12 种矩阵乘法场景中，相关指令的使用方法，其中 A 和 B 分别表示矩阵乘法的左右输入矩阵。| |
| 28 | | [mmad_unitflag](./mmad_unitflag) | 本样例介绍是否使能unitFlag对于Mmad指令执行矩阵乘法性能的影响。 | | 28 | | [mmad_unitflag](./mmad_unitflag) | 本样例介绍是否使能unitFlag对于Mmad指令执行矩阵乘法性能的影响。 | |
| @@ -1,5 +1,5 @@ | |||
| 1 | # ---------------------------------------------------------------------------------------------------------- | 1 | # ---------------------------------------------------------------------------------------------------------- |
| 2 | -# Copyright (c) 2025 Huawei Technologies Co., Ltd. | 2 | +# Copyright (c) 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. |
| @@ -8,9 +8,11 @@ | |||
| 8 | # See LICENSE in the root of the software repository for the full text of the License. | 8 | # See LICENSE in the root of the software repository for the full text of the License. |
| 9 | # ---------------------------------------------------------------------------------------------------------- | 9 | # ---------------------------------------------------------------------------------------------------------- |
| 10 | 10 | ||
| 11 | - | ||
| 12 | cmake_minimum_required(VERSION 3.16) | 11 | cmake_minimum_required(VERSION 3.16) |
| 13 | 12 | ||
| 13 | +set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim") | ||
| 14 | +set(CMAKE_ASC_ARCHITECTURES "dav-2201" CACHE STRING "NPU architecture: dav-2201, dav-3510") | ||
| 15 | + | ||
| 14 | find_package(ASC REQUIRED) | 16 | find_package(ASC REQUIRED) |
| 15 | 17 | ||
| 16 | project(kernel_samples LANGUAGES ASC CXX) | 18 | project(kernel_samples LANGUAGES ASC CXX) |
| @@ -19,20 +21,16 @@ add_executable(demo | |||
| 19 | mmad.asc | 21 | mmad.asc |
| 20 | ) | 22 | ) |
| 21 | 23 | ||
| 22 | -target_link_libraries(demo PRIVATE | 24 | +set(SCENARIO_NUM "1" CACHE STRING "SCENARIO_NUM, e.g. 1 2 3 4") |
| 23 | - tiling_api | 25 | +if(SCENARIO_NUM EQUAL 1 AND CMAKE_ASC_ARCHITECTURES STREQUAL "dav-3510") |
| 24 | - register | 26 | + message(FATAL_ERROR "When SCENARIO_NUM=1, only CMAKE_ASC_ARCHITECTURES=dav-2201 is supported.") |
| 25 | - platform | 27 | +endif() |
| 26 | - m | 28 | + |
| 29 | +message(STATUS "[INFO] SCENARIO_NUM is ${SCENARIO_NUM}") | ||
| 30 | +target_compile_definitions(demo PRIVATE | ||
| 31 | + SCENARIO_NUM=${SCENARIO_NUM} | ||
| 27 | ) | 32 | ) |
| 28 | 33 | ||
| 29 | -# ====================================================================================== | ||
| 30 | -# NPU 编译选项配置 | ||
| 31 | -# | ||
| 32 | -# 说明： | ||
| 33 | -# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。 | ||
| 34 | -# ====================================================================================== | ||
| 35 | target_compile_options(demo PRIVATE | 34 | target_compile_options(demo PRIVATE |
| 36 | - $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201> | 35 | + $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}> |
| 37 | - # $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510> | ||
| 38 | ) | 36 | ) |
| @@ -1,6 +1,6 @@ | |||
| 1 | # Mmad样例 | 1 | # Mmad样例 |
| 2 | ## 概述 | 2 | ## 概述 |
| 3 | -本样例介绍基于基础API实现矩阵乘Mmad样例。 | 3 | +本样例介绍了输入为ND格式，B4 / B8 / B16 / B32输入数据类型（具体以int4_t / int8_t / half / float为例），四种输入数据类型下的矩阵乘法，说明如何通过Mmad指令实现矩阵乘法计算（C = A x B + Bias）。 |
| 4 | 4 | ||
| 5 | ## 支持的产品 | 5 | ## 支持的产品 |
| 6 | - Ascend 950PR/Ascend 950DT | 6 | - Ascend 950PR/Ascend 950DT |
| @@ -10,6 +10,7 @@ | |||
| 10 | ## 目录结构介绍 | 10 | ## 目录结构介绍 |
| 11 | ``` | 11 | ``` |
| 12 | ├── mmad | 12 | ├── mmad |
| 13 | +│ ├── img // 本文中的插图文件 | ||
| 13 | │ ├── scripts | 14 | │ ├── scripts |
| 14 | │ │ ├── gen_data.py // 输入数据和真值数据生成脚本 | 15 | │ │ ├── gen_data.py // 输入数据和真值数据生成脚本 |
| 15 | │ │ └── verify_result.py // 验证输出数据和真值数据是否一致的验证脚本 | 16 | │ │ └── verify_result.py // 验证输出数据和真值数据是否一致的验证脚本 |
| @@ -18,35 +19,174 @@ | |||
| 18 | │ └── mmad.asc // Ascend C算子实现 & 调用样例 | 19 | │ └── mmad.asc // Ascend C算子实现 & 调用样例 |
| 19 | ``` | 20 | ``` |
| 20 | 21 | ||
| 21 | -## 算子描述 | 22 | +## 样例描述 |
| 22 | -- 算子功能： | 23 | +一次完整的矩阵乘法涉及的数据搬运过程包括：GM-->L1、L1-->L0A/L0B、L1-->BT（BiasTable Buffer）、L0C-->GM，其中不同存储单元的数据排布格式，如下表1所示： |
| 23 | - 本样例中实现的是[m, n, k]固定为[32, 32, 32]的Matmul算子，并使用Ascend C基础Api实现。对应的数学表达式为： | 24 | + |
| 24 | - ``` | 25 | +<table border="2" align="center"> |
| 25 | - C = A * B | 26 | +<caption>表1：不同存储单元的数据排布格式</caption> |
| 26 | - ``` | 27 | + <tr> |
| 27 | -- 算子规格： | 28 | + <td >存储单元</td> |
| 28 | - <table> | 29 | + <td>数据排布格式</td> |
| 29 | - <tr><td rowspan="1" align="center">算子类型(OpType)</td><td colspan="4" align="center">Add</td></tr> | 30 | + </tr> |
| 30 | - </tr> | 31 | + <tr> |
| 31 | - <tr><td rowspan="3" align="center">算子输入</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr> | 32 | + <td>GM</td> |
| 32 | - <tr><td align="center">x</td><td align="center">32 * 32</td><td align="center">half</td><td align="center">ND</td></tr> | 33 | + <td>输入A、B矩阵和输出C矩阵为ND排列。</td> |
| 33 | - <tr><td align="center">y</td><td align="center">32 * 32</td><td align="center">half</td><td align="center">ND</td></tr> | 34 | + </tr> |
| 34 | - </tr> | 35 | + <tr> |
| 35 | - </tr> | 36 | + <td>L1</td> |
| 36 | - <tr><td rowspan="1" align="center">算子输出</td><td align="center">z</td><td align="center">32 * 32</td><td align="center">float</td><td align="center">ND</td></tr> | 37 | + <td>A、B矩阵为Nz排列。</td> |
| 37 | - </tr> | 38 | + </tr> |
| 38 | - <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">mmad_custom</td></tr> | 39 | + <tr> |
| 39 | - </table> | 40 | + <td>L0A</td> |
| 40 | -- 算子实现： | 41 | + <td>Ascend 950PR/Ascend 950DT产品，A矩阵为Nz排列；<br>Atlas A3 训练系列产品/Atlas A3 推理系列产品和Atlas A2 训练系列产品/Atlas A2 推理系列产品，A矩阵为Zz排列。</td> |
| 41 | - Mmad算子的数学表达式为： | 42 | + </tr> |
| 42 | - ``` | 43 | + <tr> |
| 43 | - C = A * B | 44 | + <td>L0B</td> |
| 44 | - ``` | 45 | + <td>B矩阵为Zn排列。</td> |
| 45 | - 计算逻辑是：Ascend C提供的矩阵乘计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储并进行分形转换，然后使用计算接口完成两个输入参数矩阵乘运算，得到最终结果，再搬出到外部存储上。 | 46 | + </tr> |
| 47 | + <tr> | ||
| 48 | + <td>BT（BiasTable Buffer）</td> | ||
| 49 | + <td>Bias是shape为[N]的一维Tensor。</td> | ||
| 50 | + </tr> | ||
| 51 | + <tr> | ||
| 52 | + <td>L0C</td> | ||
| 53 | + <td>C矩阵为Nz排列。</td> | ||
| 54 | + </tr> | ||
| 55 | +</table> | ||
| 56 | + | ||
| 57 | +通常的矩阵乘法计算公式：C = A × B + Bias，其中A、B、Bias、C矩阵的需要满足的shape分别为[M,K]、[K,N]、[N]和[M,N]。Bias的数据类型与C矩阵数据类型的对应关系，如表2所示： | ||
| 58 | +<table border="2" align="center"> | ||
| 59 | +<caption>表 2：L0C 与输入 Bias 的数据类型对应关系</caption> | ||
| 60 | + <tr> | ||
| 61 | + <td>Bias在GM/L1上的数据类型</td> | ||
| 62 | + <td>Bias在BT（BiasTable Buffer）上的数据类型</td> | ||
| 63 | + <td>矩阵计算输出到L0C上的数据类型</td> | ||
| 64 | + </tr> | ||
| 65 | + <tr> | ||
| 66 | + <td>int32_t</td> | ||
| 67 | + <td>int32_t</td> | ||
| 68 | + <td>int32_t</td> | ||
| 69 | + </tr> | ||
| 70 | + <tr> | ||
| 71 | + <td>bfloat16</td> | ||
| 72 | + <td rowspan="3">float</td> | ||
| 73 | + <td rowspan="3">float</td> | ||
| 74 | + </tr> | ||
| 75 | + <tr> | ||
| 76 | + <td>half</td> | ||
| 77 | + </tr> | ||
| 78 | + <tr> | ||
| 79 | + <td>float</td> | ||
| 80 | + </tr> | ||
| 81 | +</table> | ||
| 82 | + | ||
| 83 | +程序中scenarioNum参数不同取值对应的场景，如下表3所示： | ||
| 84 | + | ||
| 85 | +<table border="2" align="center"> | ||
| 86 | +<caption>表3：scenarioNum不同取值的含义</caption> | ||
| 87 | + <tr> | ||
| 88 | + <td >scenarioNum</td> | ||
| 89 | + <td>输入数据类型</td> | ||
| 90 | + <td>输出数据类型</td> | ||
| 91 | + <td>A矩阵</td> | ||
| 92 | + <td>B矩阵</td> | ||
| 93 | + <td>Bias</td> | ||
| 94 | + </tr> | ||
| 95 | + <tr> | ||
| 96 | + <td>1</td> | ||
| 97 | + <td>int4b_t</td> | ||
| 98 | + <td>int32_t</td> | ||
| 99 | + <td>不转置</td> | ||
| 100 | + <td>转置</td> | ||
| 101 | + <td>不启用Bias，C矩阵的初始值为0</td> | ||
| 102 | + </tr> | ||
| 103 | + <tr> | ||
| 104 | + <td>2</td> | ||
| 105 | + <td>int8_t</td> | ||
| 106 | + <td>int32_t</td> | ||
| 107 | + <td>不转置</td> | ||
| 108 | + <td>不转置</td> | ||
| 109 | + <td>带Bias且不传入biasTensor，C矩阵的初始值来源于C2</td> | ||
| 110 | + </tr> | ||
| 111 | + <tr> | ||
| 112 | + <td>3</td> | ||
| 113 | + <td>bfloat16</td> | ||
| 114 | + <td>bfloat16</td> | ||
| 115 | + <td>不转置</td> | ||
| 116 | + <td>转置</td> | ||
| 117 | + <td>不带Bias，C矩阵累加来源于CO1初始值</td> | ||
| 118 | + </tr> | ||
| 119 | + <tr> | ||
| 120 | + <td>4</td> | ||
| 121 | + <td>float</td> | ||
| 122 | + <td>float</td> | ||
| 123 | + <td>转置</td> | ||
| 124 | + <td>转置</td> | ||
| 125 | + <td>带Bias且传入biasTensor的场景。</td> | ||
| 126 | + </tr> | ||
| 127 | +</table> | ||
| 128 | + | ||
| 129 | +### 场景详细说明 | ||
| 130 | + 本样例通过编译参数 `SCENARIO_NUM` 选择不同的输出场景，所有场景基于相同的矩阵乘规格：[M, N, K] = [30, 40, 70]，核函数名为 `mmad_custom`。 | ||
| 131 | + | ||
| 132 | +**场景1 int4b_t输入，int32_t输出，C矩阵初始值为0** | ||
| 133 | + | ||
| 134 | +- 输入：A不转置 [30, 70] int4b_t类型，ND格式；B转置 [40, 70] int4b_t类型，ND格式；不带Bias | ||
| 135 | +- 输出：C [30, 40] int32_t类型，ND格式 | ||
| 136 | +- 实现：使用`Mmad`实现矩阵乘法运算，通过参数：`mmadParams.cmatrixInitVal = true`，设置C矩阵初始值为0 | ||
| 137 | +- 说明：该场景仅支持Atlas A3 训练系列产品/Atlas A3 推理系列产品/Atlas A2 训练系列产品/Atlas A2 推理系列产品，且不支持使用传入biasTensor的方式加Bias（场景4）。 | ||
| 138 | + | ||
| 139 | +**场景2 int8_t输入，int32_t输出，C矩阵初始值来源于C2** | ||
| 140 | +- 输入：A不转置 [30, 70] int8_t类型，ND格式；B不转置 [70, 40] int8_t类型，ND格式；Bias [40] int32_t类型 | ||
| 141 | +- 输出：C [30, 40] int32_t类型，ND格式 | ||
| 142 | +- 实现：使用`Mmad`实现矩阵乘法运算，不传入biasTensor通过参数：`mmadParams.cmatrixInitVal = false、mmadParams.cmatrixSource = true`，设置C矩阵初始值来源于C2 | ||
| 143 | +- 说明：int8_t类型输入，B矩阵不转置场景下，N轴向2 * 16对齐，填充了全部是无效数据的32 * 16的分形。如下图1所示，如果设置`mmadParams.n = N`，就会导致读入编号为3、7的分形，同时又没能将包含有效数据的编号为9、10的分形读入。因此需要设置：`mmadParams.n = CeilAlign(N, BLOCK_CUBE * fractalNum)`，此时会读入全部分型，虽然矩阵计算结果中包含了无效数据参与计算的结果，但是在Fixpipe指令搬出数据时通过设置`fixpipeParams.nSize = N`来保证无效数据参与计算的结果不会被搬出。 | ||
| 144 | +<p align="center"> | ||
| 145 | + <img src="img/mmad_s8_L0B_转置.png" width="700"> | ||
W | |||
| 146 | +</p> | ||
| 147 | +<p align="center"> | ||
| 148 | +图1：int8_t类型，B不转置，N轴实际对齐要求与Mmad指令默认不一致 | ||
| 149 | +</p> | ||
| 150 | + | ||
| 151 | +**场景3 bfloat16输入，float输出，A不转置，B转置，C矩阵初始值来源于CO1** | ||
| 152 | +- 输入：A不转置 [30, 70] bfloat16类型，ND格式；B转置 [40, 70] bfloat16类型，ND格式；不带Bias，C矩阵初始值来源于CO1 | ||
| 153 | +- 输出：C [30, 40] float类型，ND格式 | ||
| 154 | +- 实现：使用`Mmad`实现矩阵乘法运算，通过参数：`mmadParams.cmatrixInitVal = false、mmadParams.cmatrixSource = false`，设置C矩阵初始值来源于CO1 | ||
| 155 | +- 说明：该场景进行两次Mmad计算，第一次计算结果存储在CO1，作为下一次计算的C矩阵初始值，最终累加两次Mmad计算结果。 | ||
| 156 | + | ||
| 157 | +**场景4 float输入，float输出，A转置，B转置，传入biasTensor，kDirectionAlign值设为true** | ||
| 158 | +- 输入：A转置 [70, 30] float类型，ND格式；B转置 [40, 70] float类型，ND格式；Bias [40] float类型 | ||
| 159 | +- 输出：C [30, 40] float类型，ND格式 | ||
| 160 | +- 实现：使用`Mmad`实现矩阵乘法运算，传入biasTensor，该场景下`mmadParams.cmatrixSource`参数无效 | ||
| 161 | +- 说明：float类型输入，A矩阵转置场景下，需要使用`mmadParams.kDirectionAlign`来解决K轴实际向`CeilAlign(K, 8*2)`对齐，与Mmad指令默认要求对齐到`CeilAlign(K, 8)`不一的问题。该场景下该参数设置为真，K轴对齐到`CeilAlign(K, 16)`，矩阵计算单元从L0A读取数据会跳过填充的无效数据，其余场景下该参数默认为flase，K轴仍向`CeilAlign(K, 8)`对齐，如下图2所示： | ||
| 162 | +<p align="center"> | ||
| 163 | + <img src="img/mmad_f32_L0A_转置.png" width="1100"> | ||
| 164 | +</p> | ||
W 这个图最下面的字“图：float类型，A转置....”不要写在图里，单独用文字写。不然字体会大小跟其他不一致。   | |||
| 165 | +<p align="center"> | ||
| 166 | +图2：float类型，A转置，K轴实际对齐与Mmad指令默认要求不一致 | ||
| 167 | +</p> | ||
| 168 | + | ||
| 169 | +### 矩阵乘法（Mmad） | ||
| 170 | + | ||
| 171 | +下面将介绍如何配置 | ||
| 172 | +[Mmad](https://www.hiascend.com/document/detail/zh/canncommercial/850/API/ascendcopapi/atlasascendc_api_07_0249.html)指令的MmadParams结构体的成员，各个成员变量的具体含义这里不再赘述。 | ||
| 173 | + | ||
| 174 | +需要注意的是当Mmad指令执行时，矩阵计算单元会从L0A/L0B连续读入多个分形参与矩阵乘计算，读入分形的数量根据MmadParams结构体的成员变量m、n、k的取值以及Mmad指令对L0A/L0B上A矩阵和B矩阵各个轴的对齐要求来计算的。以输入为b16类型为例：Mmad指令是按照A矩阵分形为[16,16]、B矩阵分形为[16,16]来连续读入分形的，也就是说矩阵计算单元从L0A/L0B连续读入的分形总数目分别为：2x5=10、5x3=15，写入到L0C的分形总数为2x3=6。如下图所示，图3表示Atlas A3 训练系列产品/Atlas A3 推理系列产品和Atlas A2 训练系列产品/Atlas A2 推理系列产品，图4表示Ascend 950PR/Ascend 950DT，两者在L0A上的数据排布不一致，前者为Zz，后者则是Nz。 | ||
| 175 | +<p align="center"> | ||
| 176 | + <img src="img/mmad_f16_A3.png" width="900"> | ||
| 177 | +</p> | ||
| 178 | +<p align="center"> | ||
| 179 | +图3：bfloat16类型，L0A上Zz排布，Mmad数据排布示意图 | ||
| 180 | +</p> | ||
| 181 | +<p align="center"> | ||
| 182 | + <img src="img/mmad_f16_A5.png" width="900"> | ||
| 183 | +</p> | ||
| 184 | +<p align="center"> | ||
| 185 | +图4：bfloat16类型，L0A上Nz排布，Mmad数据排布示意图 | ||
| 186 | +</p> | ||
| 187 | + | ||
| 188 | +Mmad计算中包含了补齐的无效数据，需要配合Fixpipe指令在L0C搬出到GM过程中，排除参数Mmad计算过程中填充的无效数据。 | ||
| 46 | 189 | ||
| 47 | - Mmad算子的实现流程分为基本任务：CopyIn，SplitA，SplitB，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入inputGM搬运到Local Memory A1/B1中，搬运过程中进行ND至NZ分形转换。SplitA/SplitB分别将数据进一步搬运至接口所要求Local Memory A2/B2，Compute任务负责对数据进行矩阵乘运算，计算结果存储在Local Memory CO1中。CopyOut任务负责将输出数据从CO1搬运至Global Memory上的输出outputGm中，同时完成NZ到ND的分形转换。 | ||
| 48 | - - 调用实现 | ||
| 49 | - 使用内核调用符<<<>>>调用核函数。 | ||
| 50 | 190 | ||
| 51 | ## 编译运行 | 191 | ## 编译运行 |
| 52 | 在本样例根目录下执行如下步骤，编译并执行算子。 | 192 | 在本样例根目录下执行如下步骤，编译并执行算子。 |
| @@ -69,12 +209,33 @@ | |||
| 69 | 209 | ||
| 70 | - 样例执行 | 210 | - 样例执行 |
| 71 | ```bash | 211 | ```bash |
| 212 | + SCENARIO=1 | ||
| 72 | mkdir -p build && cd build; # 创建并进入build目录 | 213 | mkdir -p build && cd build; # 创建并进入build目录 |
| 73 | - cmake ..;make -j; # 编译工程 | 214 | + cmake .. -DCMAKE_ASC_ARCHITECTURES=dav-2201 -DSCENARIO_NUM=$SCENARIO;make -j; # 编译工程，默认npu模式 |
| 74 | - python3 ../scripts/gen_data.py # 生成测试输入数据 | 215 | + python3 ../scripts/gen_data.py -scenarioNum=$SCENARIO # 生成测试输入数据 |
| 75 | ./demo # 执行编译生成的可执行程序，执行样例 | 216 | ./demo # 执行编译生成的可执行程序，执行样例 |
| 76 | - python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 验证输出结果是否正确，确认算法逻辑正确 | 217 | + python3 ../scripts/verify_result.py -scenarioNum=$SCENARIO output/output.bin output/golden.bin # 验证输出结果是否正确，确认算法逻辑正确 |
| 77 | ``` | 218 | ``` |
| 219 | + | ||
| 220 | + 使用NPU仿真 模式时，添加`-DCMAKE_ASC_RUN_MODE=sim` 参数即可。 | ||
| 221 | + | ||
| 222 | + 示例如： | ||
| 223 | + ```bash | ||
| 224 | + cmake .. -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201 -DSCENARIO_NUM=$SCENARIO;make -j; # NPU仿真模式 | ||
| 225 | + ``` | ||
| 226 | + | ||
| 227 | + > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。 | ||
| 228 | + | ||
| 229 | +- 编译选项说明 | ||
| 230 | + | ||
| 231 | + | 参数 | 说明 | 可选值 | 默认值 | | ||
| 232 | + |------|------|--------|--------| | ||
| 233 | + | CMAKE_ASC_RUN_MODE | 运行模式 | npu、sim | npu | | ||
| 234 | + | CMAKE_ASC_ARCHITECTURES | NPU硬件架构 | dav-2201、dav-3510 | dav-2201 | | ||
| 235 | + | SCENARIO_NUM | 场景编号 | 1-4 | 1 | | ||
| 236 | + | ||
| 237 | +- 执行结果 | ||
| 238 | + | ||
| 78 | 执行结果如下，说明精度对比成功。 | 239 | 执行结果如下，说明精度对比成功。 |
| 79 | ```bash | 240 | ```bash |
| 80 | test pass! | 241 | test pass! |
| @@ -1,5 +1,5 @@ | |||
| 1 | /** | 1 | /** |
| 2 | -* Copyright (c) 2025 Huawei Technologies Co., Ltd. | 2 | +* Copyright (c) 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. |
| @@ -0,0 +1,3 @@ | |||
| 1 | +version https://git-lfs.github.com/spec/v1 | ||
| 2 | +oid sha256:864ff9243c6c95ce0e61d15eb73eb8e30c9512f9def78ff98546d168d196c4bc | ||
| 3 | +size 59044 | ||
| @@ -0,0 +1,3 @@ | |||
| 1 | +version https://git-lfs.github.com/spec/v1 | ||
| 2 | +oid sha256:72d96363ac7a36005aee2de877caa3568c4a0495b06dbf09f96974c86a6f500c | ||
| 3 | +size 57652 | ||
Aexamples/01_simd_cpp_api/02_features/03_basic_api/01_matrix_compute/mmad/img/mmad_f32_L0A_转置.png+3-0
| @@ -0,0 +1,3 @@ | |||
| 1 | +version https://git-lfs.github.com/spec/v1 | ||
| 2 | +oid sha256:333826ed2bbc6a9f6d98df9dc03f2ae6e819e09c873e13123d7587336b1c2ff7 | ||
| 3 | +size 55443 | ||
Aexamples/01_simd_cpp_api/02_features/03_basic_api/01_matrix_compute/mmad/img/mmad_s8_L0B_转置.png+3-0
| @@ -0,0 +1,3 @@ | |||
| 1 | +version https://git-lfs.github.com/spec/v1 | ||
| 2 | +oid sha256:d11706405c15e9a1943e925df11d2b3d242f12315211fd532321af3cd4ed897b | ||
| 3 | +size 79422 | ||
| @@ -1,5 +1,5 @@ | |||
| 1 | /** | 1 | /** |
| 2 | -* Copyright (c) 2025 Huawei Technologies Co., Ltd. | 2 | +* Copyright (c) 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. |
| @@ -11,193 +11,443 @@ | |||
| 11 | 11 | ||
| 12 | /* ! | 12 | /* ! |
| 13 | * \file mmad.asc | 13 | * \file mmad.asc |
| 14 | - * \brief | 14 | + * \brief 本文介绍了int4b_t / int8_t / half / float，四种数据类型下的矩阵乘法，包含3种启用bias的方式和kDirectionAlign参数的使用。 |
| 15 | */ | 15 | */ |
| 16 | - | 16 | +#ifdef ASCENDC_CPU_DEBUG |
| 17 | +#include "cpu_debug_launch.h" | ||
| 18 | +#endif | ||
| 17 | #include "acl/acl.h" | 19 | #include "acl/acl.h" |
| 18 | #include "kernel_operator.h" | 20 | #include "kernel_operator.h" |
| 19 | #include "data_utils.h" | 21 | #include "data_utils.h" |
| 20 | 22 | ||
| 21 | -// half type, cube block: [16, 16] | 23 | +constexpr uint32_t scenarioNum = SCENARIO_NUM; |
| 22 | -constexpr uint32_t CUBE_BLOCK = 16; | 24 | +constexpr bool HAS_BIAS_TENSOR = (SCENARIO_NUM == 2 || SCENARIO_NUM == 4); |
| 23 | -constexpr uint32_t CUBE_BLOCK_SIZE = 16 * 16; | 25 | +constexpr uint32_t BLOCK_CUBE = 16; |
| 26 | +constexpr uint64_t ADDR_0 = 0; | ||
| 24 | 27 | ||
| 28 | +template <class T, class U, uint32_t M, uint32_t N, uint32_t K, bool isATranspose, bool isBTranspose> | ||
| 25 | class KernelMmad { | 29 | class KernelMmad { |
| 26 | public: | 30 | public: |
| 27 | __aicore__ inline KernelMmad() | 31 | __aicore__ inline KernelMmad() |
| 28 | { | 32 | { |
| 29 | - aSize = m * k; | 33 | + // 根据不同输入类型设置c0Size |
| 30 | - bSize = k * n; | 34 | + if constexpr (SCENARIO_NUM == 1) { |
| 31 | - cSize = m * n; | 35 | + c0Size = 64; |
| 36 | + fractalNum = 4; | ||
| 37 | + // A不转，B转 | ||
| 38 | + aSizeAlignL1 = CeilAlign(M, BLOCK_CUBE) * CeilAlign(K, c0Size); | ||
| 39 | + aSizeAlignL0 = CeilAlign(M, BLOCK_CUBE) * CeilAlign(K, c0Size); | ||
| 40 | + bSizeAlignL1 = CeilAlign(N, BLOCK_CUBE) * CeilAlign(K, c0Size); | ||
| 41 | + bSizeAlignL0 = CeilAlign(K, c0Size) * CeilAlign(N, BLOCK_CUBE); | ||
| 42 | + } else if constexpr (SCENARIO_NUM == 2) { | ||
| 43 | + c0Size = 32; | ||
| 44 | + fractalNum = 2; | ||
| 45 | + // A、B 都不转 | ||
| 46 | + aSizeAlignL1 = CeilAlign(M, BLOCK_CUBE) * CeilAlign(K, c0Size); | ||
| 47 | + aSizeAlignL0 = CeilAlign(M, BLOCK_CUBE) * CeilAlign(K, c0Size); | ||
| 48 | + bSizeAlignL1 = CeilAlign(K, BLOCK_CUBE * fractalNum) * CeilAlign(N, c0Size); | ||
| 49 | + bSizeAlignL0 = CeilAlign(K, c0Size) * CeilAlign(N, BLOCK_CUBE * fractalNum); | ||
| 50 | + } else if constexpr (SCENARIO_NUM == 3) { | ||
| 51 | + c0Size = 16; | ||
| 52 | + fractalNum = 1; | ||
| 53 | + // A不转，B转 | ||
| 54 | + aSizeAlignL1 = CeilAlign(M, BLOCK_CUBE) * CeilAlign(K, c0Size); | ||
| 55 | + aSizeAlignL0 = CeilAlign(M, BLOCK_CUBE) * CeilAlign(K, c0Size); | ||
| 56 | + bSizeAlignL1 = CeilAlign(N, BLOCK_CUBE) * CeilAlign(K, c0Size); | ||
| 57 | + bSizeAlignL0 = CeilAlign(K, c0Size) * CeilAlign(N, BLOCK_CUBE); | ||
| 58 | + } else { | ||
| 59 | + c0Size = 8; | ||
| 60 | + fractalNum = 2; | ||
| 61 | + // A、B都转 | ||
| 62 | + aSizeAlignL1 = CeilAlign(K, BLOCK_CUBE) * CeilAlign(M, c0Size * fractalNum); | ||
| 63 | + aSizeAlignL0 = CeilAlign(M, BLOCK_CUBE) * CeilAlign(K, c0Size * fractalNum); | ||
| 64 | + bSizeAlignL1 = CeilAlign(N, BLOCK_CUBE) * CeilAlign(K, c0Size); | ||
| 65 | + bSizeAlignL0 = CeilAlign(K, c0Size) * CeilAlign(N, BLOCK_CUBE); | ||
| 66 | + } | ||
| 67 | + fractalSize = BLOCK_CUBE * c0Size; | ||
| 68 | + cSizeAlignL0 = CeilAlign(M, BLOCK_CUBE) * CeilAlign(N, BLOCK_CUBE); | ||
| 69 | + | ||
| 70 | +#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510) | ||
| 71 | + // 对于Ascend 950PR/Ascend 950DT，C1->C2，要求待搬运数据的连续传输数据块长度（blockLen）的单位是32B，且数据类型为b32时值需为偶数 | ||
| 72 | + biasSizeAlign = CeilAlign(CeilDiv(N * sizeof(T), 32), 2) * 32; | ||
| 73 | +#else | ||
| 74 | + // Atlas A3 训练系列产品/Atlas A3 推理系列产品、Atlas A2 训练系列产品/Atlas A2 推理系列产品，C1->C2，要求blockLen的单位是64B | ||
| 75 | + biasSizeAlign = CeilAlign(N * sizeof(T), 64); | ||
| 76 | +#endif | ||
| 77 | + a1Addr = ADDR_0; | ||
| 78 | + b1Addr = a1Addr + aSizeAlignL1 * sizeof(U); | ||
| 79 | + c1Addr = a1Addr + b1Addr + bSizeAlignL1 * sizeof(U); | ||
| 32 | } | 80 | } |
| 33 | - __aicore__ inline void Init(GM_ADDR a, GM_ADDR b, GM_ADDR c) | 81 | + __aicore__ inline void Init(GM_ADDR a, GM_ADDR b, GM_ADDR bias, GM_ADDR c) |
| 34 | { | 82 | { |
| 35 | - aGM.SetGlobalBuffer((__gm__ half *)a); | 83 | + aGM.SetGlobalBuffer((__gm__ U *)a); |
| 36 | - bGM.SetGlobalBuffer((__gm__ half *)b); | 84 | + bGM.SetGlobalBuffer((__gm__ U *)b); |
| 37 | - cGM.SetGlobalBuffer((__gm__ float *)c); | 85 | + biasGM.SetGlobalBuffer((__gm__ T *)bias); |
| 38 | - pipe.InitBuffer(inQueueA1, 1, aSize * sizeof(half)); | 86 | + cGM.SetGlobalBuffer((__gm__ T *)c); |
| 39 | - pipe.InitBuffer(inQueueA2, 1, aSize * sizeof(half)); | ||
| 40 | - pipe.InitBuffer(inQueueB1, 1, bSize * sizeof(half)); | ||
| 41 | - pipe.InitBuffer(inQueueB2, 1, bSize * sizeof(half)); | ||
| 42 | - pipe.InitBuffer(outQueueCO1, 1, cSize * sizeof(float)); | ||
| 43 | } | 87 | } |
| 44 | __aicore__ inline void Process() | 88 | __aicore__ inline void Process() |
| 45 | { | 89 | { |
| 46 | - CopyIn(); | 90 | + AscendC::LocalTensor<U> a1Local(AscendC::TPosition::A1, a1Addr, aSizeAlignL1); |
| 47 | - SplitA(); | 91 | + AscendC::LocalTensor<U> b1Local(AscendC::TPosition::B1, b1Addr, bSizeAlignL1); |
| 48 | - SplitB(); | 92 | + AscendC::LocalTensor<T> bias1Local(AscendC::TPosition::C1, c1Addr, biasSizeAlign); |
| 49 | - Compute(); | 93 | + |
| 50 | - CopyOut(); | 94 | + // GM -> L1 |
| 95 | + CopyIn(a1Local, b1Local, bias1Local); | ||
| 96 | + // L1 -> L0A | ||
| 97 | + AscendC::LocalTensor<U> a2Local(AscendC::TPosition::A2, ADDR_0, aSizeAlignL0); | ||
| 98 | + if constexpr (isATranspose) { | ||
| 99 | + SplitATranspose(a1Local, a2Local); | ||
| 100 | + } else { | ||
| 101 | + SplitA(a1Local, a2Local); | ||
| 102 | + } | ||
| 103 | + // L1 -> L0B | ||
| 104 | + AscendC::LocalTensor<U> b2Local(AscendC::TPosition::B2, ADDR_0, bSizeAlignL0); | ||
| 105 | + if constexpr (!isBTranspose) { | ||
| 106 | + SplitBTranspose(b1Local, b2Local); | ||
| 107 | + } else { | ||
| 108 | + SplitB(b1Local, b2Local); | ||
| 109 | + } | ||
| 110 | + // L1 -> BT | ||
| 111 | + AscendC::LocalTensor<T> bias2Local(AscendC::TPosition::C2, ADDR_0, biasSizeAlign); | ||
| 112 | + if constexpr (HAS_BIAS_TENSOR) { | ||
| 113 | + SplitBias(bias1Local, bias2Local); | ||
| 114 | + } | ||
| 115 | + // Mmad | ||
| 116 | + AscendC::LocalTensor<T> co1Local(AscendC::TPosition::CO1, ADDR_0, cSizeAlignL0); | ||
| 117 | + Compute(a2Local, b2Local, bias2Local, co1Local); | ||
| 118 | + // L0C -> GM | ||
| 119 | + CopyOut(co1Local); | ||
| 51 | } | 120 | } |
| 52 | 121 | ||
| 53 | private: | 122 | private: |
| 54 | - __aicore__ inline uint32_t CeilCubeBlock(uint32_t len) { | 123 | + // 向上整除 |
| 55 | - return (len + CUBE_BLOCK - 1) / CUBE_BLOCK; | 124 | + __aicore__ inline uint32_t CeilDiv(uint32_t numerator, uint32_t denominator) |
| 125 | + { | ||
| 126 | + return (numerator + denominator - 1) / denominator; | ||
| 127 | + } | ||
| 128 | + // 向上对齐 | ||
| 129 | + __aicore__ inline uint32_t CeilAlign(uint32_t numerator, uint32_t denominator) | ||
| 130 | + { | ||
| 131 | + return (numerator + denominator - 1) / denominator * denominator; | ||
| 56 | } | 132 | } |
| 57 | 133 | ||
| 58 | - __aicore__ inline void CopyIn() | 134 | + __aicore__ inline void CopyIn(AscendC::LocalTensor<U>& a1Local, AscendC::LocalTensor<U>& b1Local, |
| 135 | + AscendC::LocalTensor<T>& bias1Local) | ||
| 59 | { | 136 | { |
| 60 | - AscendC::LocalTensor<half> a1Local = inQueueA1.AllocTensor<half>(); | ||
| 61 | - AscendC::LocalTensor<half> b1Local = inQueueB1.AllocTensor<half>(); | ||
| 62 | - | ||
| 63 | AscendC::Nd2NzParams nd2nzA1Params; | 137 | AscendC::Nd2NzParams nd2nzA1Params; |
| 64 | - nd2nzA1Params.ndNum = 1; | 138 | + if constexpr (isATranspose) { |
| 65 | - nd2nzA1Params.nValue = m; | 139 | +#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510) |
| 66 | - nd2nzA1Params.dValue = k; | 140 | + // ND -> Nz |
| 67 | - nd2nzA1Params.srcNdMatrixStride = 0; | 141 | + nd2nzA1Params.ndNum = 1; |
| 68 | - nd2nzA1Params.srcDValue = k; | 142 | + nd2nzA1Params.nValue = K; |
| 69 | - nd2nzA1Params.dstNzC0Stride = CeilCubeBlock(m) * CUBE_BLOCK; | 143 | + nd2nzA1Params.dValue = M; |
| 70 | - nd2nzA1Params.dstNzNStride = 1; | 144 | + nd2nzA1Params.srcNdMatrixStride = 0; |
| 71 | - nd2nzA1Params.dstNzMatrixStride = 0; | 145 | + nd2nzA1Params.srcDValue = M; |
| 146 | + nd2nzA1Params.dstNzC0Stride = CeilAlign(K, BLOCK_CUBE); | ||
| 147 | + nd2nzA1Params.dstNzNStride = 1; | ||
| 148 | + nd2nzA1Params.dstNzMatrixStride = 0; | ||
| 149 | +#else | ||
| 150 | + // ND -> Zz | ||
| 151 | + nd2nzA1Params.ndNum = CeilDiv(K, BLOCK_CUBE); | ||
| 152 | + nd2nzA1Params.nValue = BLOCK_CUBE; | ||
| 153 | + nd2nzA1Params.dValue = M; | ||
| 154 | + nd2nzA1Params.srcNdMatrixStride = BLOCK_CUBE * M; | ||
| 155 | + nd2nzA1Params.srcDValue = M; | ||
| 156 | + nd2nzA1Params.dstNzC0Stride = BLOCK_CUBE; | ||
| 157 | + nd2nzA1Params.dstNzNStride = 1; | ||
| 158 | + nd2nzA1Params.dstNzMatrixStride = BLOCK_CUBE * CeilAlign(M, BLOCK_CUBE * fractalNum); | ||
| 159 | +#endif | ||
| 160 | + } else { | ||
| 161 | + // ND -> Nz | ||
| 162 | + nd2nzA1Params.ndNum = 1; | ||
| 163 | + nd2nzA1Params.nValue = M; | ||
| 164 | + nd2nzA1Params.dValue = K; | ||
| 165 | + nd2nzA1Params.srcNdMatrixStride = 0; | ||
| 166 | + nd2nzA1Params.srcDValue = K; | ||
| 167 | + nd2nzA1Params.dstNzC0Stride = CeilAlign(M, BLOCK_CUBE); | ||
| 168 | + nd2nzA1Params.dstNzNStride = 1; | ||
| 169 | + nd2nzA1Params.dstNzMatrixStride = 0; | ||
| 170 | + if constexpr (AscendC::IsSameType<U, AscendC::int4b_t>::value) { | ||
| 171 | + // 两个int4b_t合并成int8_t, 进行GM-->L1, nd2nz的搬运 | ||
| 172 | + nd2nzA1Params.dValue = K >> 1; | ||
| 173 | + nd2nzA1Params.srcDValue = K >> 1; | ||
| 174 | + } | ||
| 175 | + } | ||
| 72 | AscendC::DataCopy(a1Local, aGM, nd2nzA1Params); | 176 | AscendC::DataCopy(a1Local, aGM, nd2nzA1Params); |
| 177 | + AscendC::SetFlag<AscendC::HardEvent::MTE2_MTE1>(EVENT_ID0); | ||
| 73 | 178 | ||
| 74 | AscendC::Nd2NzParams nd2nzB1Params; | 179 | AscendC::Nd2NzParams nd2nzB1Params; |
| 75 | - nd2nzB1Params.ndNum = 1; | 180 | + if constexpr (isBTranspose) { |
| 76 | - nd2nzB1Params.nValue = k; | 181 | + nd2nzB1Params.ndNum = 1; |
| 77 | - nd2nzB1Params.dValue = n; | 182 | + nd2nzB1Params.nValue = N; |
| 78 | - nd2nzB1Params.srcNdMatrixStride = 0; | 183 | + nd2nzB1Params.dValue = K; |
| 79 | - nd2nzB1Params.srcDValue = n; | 184 | + nd2nzB1Params.srcNdMatrixStride = 0; |
| 80 | - nd2nzB1Params.dstNzC0Stride = CeilCubeBlock(k) * CUBE_BLOCK; | 185 | + nd2nzB1Params.srcDValue = K; |
| 81 | - nd2nzB1Params.dstNzNStride = 1; | 186 | + nd2nzB1Params.dstNzC0Stride = CeilAlign(N, BLOCK_CUBE); |
| 82 | - nd2nzB1Params.dstNzMatrixStride = 0; | 187 | + nd2nzB1Params.dstNzNStride = 1; |
| 188 | + nd2nzB1Params.dstNzMatrixStride = 0; | ||
| 189 | + if constexpr (AscendC::IsSameType<U, AscendC::int4b_t>::value) { | ||
| 190 | + // 两个int4b_t合并成int8_t, 进行GM-->L1, nd2nz的搬运 | ||
| 191 | + nd2nzB1Params.dValue = K >> 1; | ||
| 192 | + nd2nzB1Params.srcDValue = K >> 1; | ||
| 193 | + } | ||
| 194 | + } else { | ||
| 195 | + nd2nzB1Params.ndNum = 1; | ||
| 196 | + nd2nzB1Params.nValue = K; | ||
| 197 | + nd2nzB1Params.dValue = N; | ||
| 198 | + nd2nzB1Params.srcNdMatrixStride = 0; | ||
| 199 | + nd2nzB1Params.srcDValue = N; | ||
| 200 | + nd2nzB1Params.dstNzC0Stride = CeilAlign(K, BLOCK_CUBE * fractalNum); | ||
| 201 | + nd2nzB1Params.dstNzNStride = 1; | ||
| 202 | + nd2nzB1Params.dstNzMatrixStride = 0; | ||
| 203 | + } | ||
| 83 | AscendC::DataCopy(b1Local, bGM, nd2nzB1Params); | 204 | AscendC::DataCopy(b1Local, bGM, nd2nzB1Params); |
| 205 | + AscendC::SetFlag<AscendC::HardEvent::MTE2_MTE1>(EVENT_ID1); | ||
| 84 | 206 | ||
| 85 | - inQueueA1.EnQue(a1Local); | 207 | + if constexpr (HAS_BIAS_TENSOR) { |
| 86 | - inQueueB1.EnQue(b1Local); | 208 | + // 由于N可能非32B对齐，“基础数据搬运”无法处理非对齐数据，且“DataCopyPad”不支持GM到L1的搬运， |
| 209 | + // 因此，使用“随路转换ND2Nz搬运”实现Bias的GM到L1搬运，确保搬运到L1数据32B对齐。 | ||
| 210 | + AscendC::Nd2NzParams nd2nzC1Params = {1, 1, N, 0, N, 1, 1, 0}; | ||
| 211 | + AscendC::DataCopy(bias1Local, biasGM, nd2nzC1Params); | ||
| 212 | + AscendC::SetFlag<AscendC::HardEvent::MTE2_MTE1>(EVENT_ID2); | ||
| 213 | + } | ||
| 87 | } | 214 | } |
| 88 | 215 | ||
| 89 | - __aicore__ inline void SplitA() | 216 | + __aicore__ inline void SplitA(AscendC::LocalTensor<U>& a1Local, AscendC::LocalTensor<U>& a2Local) |
| 90 | { | 217 | { |
| 91 | - AscendC::LocalTensor<half> a1Local = inQueueA1.DeQue<half>(); | 218 | + AscendC::WaitFlag<AscendC::HardEvent::MTE2_MTE1>(EVENT_ID0); |
| 92 | - AscendC::LocalTensor<half> a2Local = inQueueA2.AllocTensor<half>(); | 219 | +#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510) |
| 93 | - | 220 | + // Nz -> Nz, NPU_ARCH为3510时，要求A2上的数据排布为: Nz |
| 94 | -#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 2201 || __NPU_ARCH__ == 2202) | ||
| 95 | - uint32_t dstOffset = CeilCubeBlock(k) * CUBE_BLOCK_SIZE; | ||
| 96 | - uint32_t srcOffset = CUBE_BLOCK_SIZE; | ||
| 97 | - // Nz -> Zz | ||
| 98 | - AscendC::LoadData2DParams loadDataParams; | ||
| 99 | - loadDataParams.repeatTimes = CeilCubeBlock(k); | ||
| 100 | - loadDataParams.srcStride = CeilCubeBlock(m); | ||
| 101 | - loadDataParams.dstGap = 0; | ||
| 102 | - loadDataParams.ifTranspose = false; | ||
| 103 | - for (int i = 0; i < CeilCubeBlock(m); ++i) { | ||
| 104 | - AscendC::LoadData(a2Local[i * dstOffset], a1Local[i * srcOffset], loadDataParams); | ||
| 105 | - } | ||
| 106 | -#elif defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510) | ||
| 107 | AscendC::LoadData2DParamsV2 loadDataParams; | 221 | AscendC::LoadData2DParamsV2 loadDataParams; |
| 108 | loadDataParams.mStartPosition = 0; | 222 | loadDataParams.mStartPosition = 0; |
| 109 | loadDataParams.kStartPosition = 0; | 223 | loadDataParams.kStartPosition = 0; |
| 110 | - loadDataParams.mStep = AscendC::DivCeil(m, 16); | 224 | + loadDataParams.mStep = CeilDiv(M, BLOCK_CUBE); |
| 111 | - loadDataParams.kStep = AscendC::DivCeil(k * sizeof(half), 32); | 225 | + loadDataParams.kStep = CeilDiv(K, c0Size); |
| 112 | - loadDataParams.srcStride = AscendC::DivCeil(m, 16); | 226 | + loadDataParams.srcStride = CeilDiv(M, BLOCK_CUBE); |
| 113 | - loadDataParams.dstStride = AscendC::DivCeil(m, 16); | 227 | + loadDataParams.dstStride = CeilDiv(M, BLOCK_CUBE); |
| 114 | - loadDataParams.sid = 0; | ||
| 115 | loadDataParams.ifTranspose = false; | 228 | loadDataParams.ifTranspose = false; |
| 116 | AscendC::LoadData(a2Local, a1Local, loadDataParams); | 229 | AscendC::LoadData(a2Local, a1Local, loadDataParams); |
| 117 | -#endif | 230 | +#else |
| 118 | - inQueueA2.EnQue<half>(a2Local); | 231 | + // Nz -> Zz, NPU_ARCH为2201时，要求A2上的数据排布为: Zz |
| 119 | - inQueueA1.FreeTensor(a1Local); | 232 | + uint32_t dstOffset = CeilDiv(K, c0Size) * fractalSize; |
| 120 | - } | 233 | + uint32_t srcOffset = fractalSize; |
| 121 | - __aicore__ inline void SplitB() | ||
| 122 | - { | ||
| 123 | - AscendC::LocalTensor<half> b1Local = inQueueB1.DeQue<half>(); | ||
| 124 | - AscendC::LocalTensor<half> b2Local = inQueueB2.AllocTensor<half>(); | ||
| 125 | - | ||
| 126 | - uint32_t dstOffset = CeilCubeBlock(n) * CUBE_BLOCK_SIZE; | ||
| 127 | - uint32_t srcOffset = CUBE_BLOCK_SIZE; | ||
| 128 | - // Nz -> Zn | ||
| 129 | AscendC::LoadData2DParams loadDataParams; | 234 | AscendC::LoadData2DParams loadDataParams; |
| 130 | - loadDataParams.repeatTimes = CeilCubeBlock(n); | 235 | + loadDataParams.startIndex = 0; |
| 131 | - loadDataParams.srcStride = CeilCubeBlock(k); | 236 | + loadDataParams.repeatTimes = CeilDiv(K, c0Size); |
| 237 | + loadDataParams.srcStride = CeilDiv(M, BLOCK_CUBE); | ||
| 132 | loadDataParams.dstGap = 0; | 238 | loadDataParams.dstGap = 0; |
| 239 | + loadDataParams.ifTranspose = false; | ||
| 240 | + for (int i = 0; i < CeilDiv(M, BLOCK_CUBE); ++i) { | ||
| 241 | + AscendC::LoadData(a2Local[i * dstOffset], a1Local[i * srcOffset], loadDataParams); | ||
| 242 | + } | ||
| 243 | +#endif | ||
| 244 | + AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(EVENT_ID0); | ||
| 245 | + } | ||
| 246 | + | ||
| 247 | + __aicore__ inline void SplitATranspose(AscendC::LocalTensor<U>& a1Local, AscendC::LocalTensor<U>& a2Local) | ||
| 248 | + { | ||
| 249 | + AscendC::WaitFlag<AscendC::HardEvent::MTE2_MTE1>(EVENT_ID0); | ||
| 250 | +#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510) | ||
| 251 | + // Nz -> Nz, NPU_ARCH为3510时，要求A2上的数据排布为: Nz | ||
| 252 | + AscendC::LoadData2DParamsV2 loadDataParams; | ||
| 253 | + loadDataParams.mStartPosition = 0; | ||
| 254 | + loadDataParams.kStartPosition = 0; | ||
| 255 | + // 开启转置时，b4 mStep必须是4的倍数，b8 mStep必须是2的倍数，b32 kStep必须是2的倍数 | ||
| 256 | + loadDataParams.mStep = CeilDiv(K, BLOCK_CUBE); | ||
| 257 | + loadDataParams.kStep = CeilAlign(CeilDiv(M, c0Size), 2); | ||
| 258 | + loadDataParams.srcStride = CeilDiv(K, BLOCK_CUBE); | ||
| 259 | + loadDataParams.dstStride = CeilDiv(M, c0Size * fractalNum); | ||
| 133 | loadDataParams.ifTranspose = true; | 260 | loadDataParams.ifTranspose = true; |
| 134 | - for (int i = 0; i < CeilCubeBlock(k); ++i) { | 261 | + AscendC::LoadData(a2Local, a1Local, loadDataParams); |
| 135 | - AscendC::LoadData(b2Local[i * dstOffset], b1Local[i * srcOffset], loadDataParams); | 262 | +#else |
| 263 | + // Nz -> Zz, NPU_ARCH为2201时，要求A2上的数据排布为: Zz | ||
| 264 | + uint32_t dstOffset = CeilDiv(K, c0Size * fractalNum) * fractalSize * fractalNum; | ||
| 265 | + uint32_t srcOffset = fractalSize * fractalNum; | ||
| 266 | + | ||
| 267 | + AscendC::LoadData2dTransposeParams loadDataParams; | ||
| 268 | + loadDataParams.startIndex = 0; | ||
| 269 | + loadDataParams.repeatTimes = CeilDiv(K, c0Size * fractalNum); | ||
| 270 | + loadDataParams.srcStride = CeilDiv(M, c0Size * fractalNum); | ||
| 271 | + loadDataParams.dstGap = 1; | ||
| 272 | + loadDataParams.dstFracGap = 0; | ||
| 273 | + for (int i = 0; i < CeilDiv(M, c0Size * fractalNum); ++i) { | ||
| 274 | + AscendC::LoadDataWithTranspose(a2Local[i * dstOffset], a1Local[i * srcOffset], loadDataParams); | ||
| 275 | + } | ||
| 276 | +#endif | ||
| 277 | + AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(EVENT_ID0); | ||
| 278 | + } | ||
| 279 | + | ||
| 280 | + __aicore__ inline void SplitB(AscendC::LocalTensor<U>& b1Local, AscendC::LocalTensor<U>& b2Local) | ||
| 281 | + { | ||
| 282 | + AscendC::WaitFlag<AscendC::HardEvent::MTE2_MTE1>(EVENT_ID1); | ||
| 283 | + // B2上的数据排布为: Zn | ||
| 284 | + AscendC::LoadData2DParams loadDataParams; | ||
| 285 | + loadDataParams.repeatTimes = CeilDiv(N, BLOCK_CUBE) * CeilDiv(K, c0Size); | ||
| 286 | + loadDataParams.srcStride = 1; | ||
| 287 | + loadDataParams.dstGap = 0; | ||
| 288 | + loadDataParams.ifTranspose = false; | ||
| 289 | + AscendC::LoadData(b2Local, b1Local, loadDataParams); | ||
| 290 | + AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(EVENT_ID1); | ||
| 291 | + } | ||
| 292 | + | ||
| 293 | + __aicore__ inline void SplitBTranspose(AscendC::LocalTensor<U>& b1Local, AscendC::LocalTensor<U>& b2Local) | ||
| 294 | + { | ||
| 295 | + AscendC::WaitFlag<AscendC::HardEvent::MTE2_MTE1>(EVENT_ID1); | ||
| 296 | + uint32_t dstOffset = CeilDiv(N, BLOCK_CUBE * fractalNum) * fractalSize * fractalNum; | ||
| 297 | + uint32_t srcOffset = fractalSize * fractalNum; | ||
| 298 | + AscendC::LoadData2dTransposeParams loadDataParams; | ||
| 299 | + loadDataParams.startIndex = 0; | ||
| 300 | + loadDataParams.repeatTimes = CeilDiv(N, c0Size); | ||
| 301 | + loadDataParams.srcStride = CeilDiv(K, BLOCK_CUBE * fractalNum); | ||
| 302 | + loadDataParams.dstGap = 1; | ||
| 303 | + loadDataParams.dstFracGap = 0; | ||
| 304 | + for (int i = 0; i < CeilDiv(K, BLOCK_CUBE * fractalNum); ++i) { | ||
| 305 | + AscendC::LoadDataWithTranspose(b2Local[i * dstOffset], b1Local[i * srcOffset], loadDataParams); | ||
| 306 | + } | ||
| 307 | + AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(EVENT_ID1); | ||
| 308 | + } | ||
| 309 | + | ||
| 310 | + __aicore__ inline void SplitBias(AscendC::LocalTensor<T>& bias1Local, AscendC::LocalTensor<T>& bias2Local) | ||
| 311 | + { | ||
| 312 | + AscendC::WaitFlag<AscendC::HardEvent::MTE2_MTE1>(EVENT_ID2); | ||
| 313 | +#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510) | ||
| 314 | + AscendC::DataCopyParams c12c2Params = {1, static_cast<uint16_t>(biasSizeAlign / 32), 0, 0}; | ||
| 315 | +#else | ||
| 316 | + AscendC::DataCopyParams c12c2Params = {1, static_cast<uint16_t>(biasSizeAlign / 64), 0, 0}; | ||
| 317 | +#endif | ||
| 318 | + AscendC::DataCopy(bias2Local, bias1Local, c12c2Params); | ||
| 319 | + AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(EVENT_ID2); | ||
| 320 | + } | ||
| 321 | + | ||
| 322 | + __aicore__ inline void Compute(AscendC::LocalTensor<U>& a2Local, AscendC::LocalTensor<U>& b2Local, | ||
| 323 | + AscendC::LocalTensor<T>& bias2Local, AscendC::LocalTensor<T>& co1Local) | ||
| 324 | + { | ||
| 325 | + AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(EVENT_ID0); | ||
| 326 | + AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(EVENT_ID1); | ||
| 327 | + | ||
| 328 | + AscendC::MmadParams mmadParams; | ||
| 329 | + mmadParams.m = M; | ||
| 330 | + mmadParams.n = N; | ||
| 331 | + mmadParams.k = K; | ||
| 332 | + | ||
| 333 | + if constexpr (AscendC::IsSameType<U, int8_t>::value && !isBTranspose) { | ||
| 334 | + mmadParams.n = CeilAlign(N, BLOCK_CUBE * fractalNum); | ||
| 335 | + } | ||
| 336 | + if constexpr (AscendC::IsSameType<U, float>::value && isATranspose) { | ||
| 337 | + mmadParams.kDirectionAlign = true; | ||
| 136 | } | 338 | } |
| 137 | 339 | ||
| 138 | - inQueueB1.FreeTensor(b1Local); | 340 | + if constexpr (SCENARIO_NUM == 1) { |
| 139 | - inQueueB2.EnQue<half>(b2Local); | 341 | + // 不带Bias场景 |
| 342 | + mmadParams.cmatrixInitVal = true; // C矩阵初始值为0 | ||
| 343 | + AscendC::Mmad(co1Local, a2Local, b2Local, mmadParams); | ||
| 344 | + } else if constexpr (SCENARIO_NUM == 2) { | ||
| 345 | + // 带Bias且不传入biasTensor，C矩阵的初始值来源于C2 | ||
| 346 | + mmadParams.cmatrixInitVal = false; | ||
| 347 | + mmadParams.cmatrixSource = true; | ||
| 348 | + AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(EVENT_ID2); | ||
| 349 | + AscendC::Mmad(co1Local, a2Local, b2Local, mmadParams); | ||
| 350 | + } else if constexpr (SCENARIO_NUM == 3) { | ||
| 351 | + // 不带Bias，C矩阵累加来源于CO1初始值 | ||
| 352 | + // 第一次mmad计算，CO1的计算结果作为下一次mmad计算C矩阵的初始值 | ||
| 353 | + AscendC::Mmad(co1Local, a2Local, b2Local, mmadParams); | ||
| 354 | + mmadParams.cmatrixInitVal = false; | ||
| 355 | + mmadParams.cmatrixSource = false; // C矩阵的初始值来源于CO1 | ||
| 356 | + AscendC::Mmad(co1Local, a2Local, b2Local, mmadParams); | ||
| 357 | + } else { | ||
| 358 | + // 带Bias且传入biasTensor的场景，该场景下cmatrixSource参数无效 | ||
| 359 | + mmadParams.cmatrixInitVal = false; | ||
| 360 | + AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(EVENT_ID2); | ||
| 361 | + AscendC::Mmad(co1Local, a2Local, b2Local, bias2Local, mmadParams); | ||
| 362 | + } | ||
| 363 | + AscendC::SetFlag<AscendC::HardEvent::M_FIX>(EVENT_ID0); | ||
| 140 | } | 364 | } |
| 141 | - __aicore__ inline void Compute() | 365 | + __aicore__ inline void CopyOut(AscendC::LocalTensor<T>& co1Local) |
| 142 | { | 366 | { |
| 143 | - AscendC::LocalTensor<half> a2Local = inQueueA2.DeQue<half>(); | 367 | + AscendC::WaitFlag<AscendC::HardEvent::M_FIX>(EVENT_ID0); |
| 144 | - AscendC::LocalTensor<half> b2Local = inQueueB2.DeQue<half>(); | 368 | +#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510) |
| 145 | - AscendC::LocalTensor<float> c1Local = outQueueCO1.AllocTensor<float>(); | 369 | + AscendC::FixpipeParamsArch3510<AscendC::CO2Layout::ROW_MAJOR> fixpipeParams; |
| 146 | - AscendC::MmadParams mmadParams; | 370 | +#else |
| 147 | - mmadParams.m = m; | ||
| 148 | - mmadParams.n = n; | ||
| 149 | - mmadParams.k = k; | ||
| 150 | - AscendC::Mmad(c1Local, a2Local, b2Local, mmadParams); | ||
| 151 | - outQueueCO1.EnQue<float>(c1Local); | ||
| 152 | - inQueueA2.FreeTensor(a2Local); | ||
| 153 | - inQueueB2.FreeTensor(b2Local); | ||
| 154 | - } | ||
| 155 | - __aicore__ inline void CopyOut() | ||
| 156 | - { | ||
| 157 | - AscendC::LocalTensor<float> c1Local = outQueueCO1.DeQue<float>(); | ||
| 158 | AscendC::FixpipeParamsV220 fixpipeParams; | 371 | AscendC::FixpipeParamsV220 fixpipeParams; |
| 159 | - fixpipeParams.nSize = n; | ||
| 160 | - fixpipeParams.mSize = m; | ||
| 161 | - fixpipeParams.srcStride = m; | ||
| 162 | - fixpipeParams.dstStride = n; | ||
| 163 | - | ||
| 164 | fixpipeParams.ndNum = 1; | 372 | fixpipeParams.ndNum = 1; |
| 165 | fixpipeParams.srcNdStride = 0; | 373 | fixpipeParams.srcNdStride = 0; |
| 166 | fixpipeParams.dstNdStride = 0; | 374 | fixpipeParams.dstNdStride = 0; |
| 167 | - AscendC::Fixpipe(cGM, c1Local, fixpipeParams); | 375 | +#endif |
| 168 | - outQueueCO1.FreeTensor(c1Local); | 376 | + fixpipeParams.mSize = M; |
| 377 | + fixpipeParams.nSize = N; | ||
| 378 | + fixpipeParams.srcStride = CeilAlign(M, BLOCK_CUBE); | ||
| 379 | + fixpipeParams.dstStride = N; | ||
| 380 | + | ||
| 381 | + AscendC::Fixpipe(cGM, co1Local, fixpipeParams); | ||
| 169 | } | 382 | } |
| 170 | 383 | ||
| 171 | private: | 384 | private: |
| 172 | - AscendC::TPipe pipe; | 385 | + AscendC::GlobalTensor<U> aGM; |
| 173 | - AscendC::TQue<AscendC::TPosition::A1, 1> inQueueA1; | 386 | + AscendC::GlobalTensor<U> bGM; |
| 174 | - AscendC::TQue<AscendC::TPosition::A2, 1> inQueueA2; | 387 | + AscendC::GlobalTensor<T> biasGM; |
| 175 | - AscendC::TQue<AscendC::TPosition::B1, 1> inQueueB1; | 388 | + AscendC::GlobalTensor<T> cGM; |
| 176 | - AscendC::TQue<AscendC::TPosition::B2, 1> inQueueB2; | ||
| 177 | - AscendC::TQue<AscendC::TPosition::CO1, 1> outQueueCO1; | ||
| 178 | 389 | ||
| 179 | - AscendC::GlobalTensor<half> aGM; | 390 | + uint32_t c0Size, fractalSize, fractalNum; |
| 180 | - AscendC::GlobalTensor<half> bGM; | 391 | + uint32_t aSizeAlignL1, bSizeAlignL1, biasSizeAlign, aSizeAlignL0, bSizeAlignL0, cSizeAlignL0; |
| 181 | - AscendC::GlobalTensor<float> cGM; | 392 | + |
| 182 | - uint16_t m = 32, k = 32, n = 32; | 393 | + uint64_t a1Addr, b1Addr, c1Addr; |
| 183 | - uint16_t aSize, bSize, cSize; | ||
| 184 | }; | 394 | }; |
| 185 | 395 | ||
| 186 | -extern "C" __global__ __cube__ void mmad_custom(GM_ADDR a, GM_ADDR b, GM_ADDR c) | 396 | +template <uint32_t M, uint32_t N, uint32_t K> |
| 397 | +__global__ __cube__ void mmad_custom(GM_ADDR a, GM_ADDR b, GM_ADDR bias, GM_ADDR c) | ||
| 187 | { | 398 | { |
| 188 | - KernelMmad op; | 399 | + AscendC::InitSocState(); |
| 189 | - op.Init(a, b, c); | 400 | + if constexpr (scenarioNum == 1) { |
| 190 | - op.Process(); | 401 | + KernelMmad<int32_t, AscendC::int4b_t, M, N, K, false, true> op; |
| 402 | + op.Init(a, b, bias, c); | ||
| 403 | + op.Process(); | ||
| 404 | + } else if constexpr (scenarioNum == 2) { | ||
| 405 | + KernelMmad<int32_t, int8_t, M, N, K, false, false> op; | ||
| 406 | + op.Init(a, b, bias, c); | ||
| 407 | + op.Process(); | ||
| 408 | + } else if constexpr (scenarioNum == 3) { | ||
| 409 | + KernelMmad<float, bfloat16_t, M, N, K, false, true> op; | ||
| 410 | + op.Init(a, b, bias, c); | ||
| 411 | + op.Process(); | ||
| 412 | + } else { | ||
| 413 | + KernelMmad<float, float, M, N, K, true, true> op; | ||
| 414 | + op.Init(a, b, bias, c); | ||
| 415 | + op.Process(); | ||
| 416 | + } | ||
| 417 | + AscendC::PipeBarrier<PIPE_ALL>(); | ||
| 191 | } | 418 | } |
| 192 | 419 | ||
| 193 | int32_t main(int32_t argc, char *argv[]) | 420 | int32_t main(int32_t argc, char *argv[]) |
| 194 | { | 421 | { |
| 195 | - uint32_t M = 32; | 422 | + constexpr uint32_t M = 30; |
| 196 | - uint32_t N = 32; | 423 | + constexpr uint32_t N = 40; |
| 197 | - uint32_t K = 32; | 424 | + constexpr uint32_t K = 70; |
| 198 | - size_t aFileSize = M * K * sizeof(int16_t); // uint16_t represent half | 425 | + size_t aFileSize = 0; |
| 199 | - size_t bFileSize = K * N * sizeof(int16_t); // uint16_t represent half | 426 | + size_t bFileSize = 0; |
| 200 | - size_t cFileSize = M * N * sizeof(float); | 427 | + size_t biasFileSize = 0; |
| 428 | + size_t cFileSize = 0; | ||
| 429 | + | ||
| 430 | + if constexpr (scenarioNum == 1) { | ||
| 431 | + aFileSize = M * K * sizeof(int8_t) / 2; | ||
| 432 | + bFileSize = K * N * sizeof(int8_t) / 2; | ||
| 433 | + biasFileSize = N * sizeof(int32_t); | ||
| 434 | + cFileSize = M * N * sizeof(int32_t); | ||
| 435 | + } else if constexpr (scenarioNum == 2) { | ||
| 436 | + aFileSize = M * K * sizeof(int8_t); | ||
| 437 | + bFileSize = K * N * sizeof(int8_t); | ||
| 438 | + biasFileSize = N * sizeof(int32_t); | ||
| 439 | + cFileSize = M * N * sizeof(int32_t); | ||
| 440 | + } else if constexpr (scenarioNum == 3) { | ||
| 441 | + aFileSize = M * K * sizeof(bfloat16_t); | ||
| 442 | + bFileSize = K * N * sizeof(bfloat16_t); | ||
| 443 | + biasFileSize = N * sizeof(float); | ||
| 444 | + cFileSize = M * N * sizeof(float); | ||
| 445 | + } else { | ||
| 446 | + aFileSize = M * K * sizeof(float); | ||
| 447 | + bFileSize = K * N * sizeof(float); | ||
| 448 | + biasFileSize = N * sizeof(float); | ||
| 449 | + cFileSize = M * N * sizeof(float); | ||
| 450 | + } | ||
| 201 | uint32_t numBlocks = 1; | 451 | uint32_t numBlocks = 1; |
| 202 | 452 | ||
| 203 | aclInit(nullptr); | 453 | aclInit(nullptr); |
| @@ -220,12 +470,21 @@ int32_t main(int32_t argc, char *argv[]) | |||
| 220 | ReadFile("./input/x2_gm.bin", bFileSize, bHost, bFileSize); | 470 | ReadFile("./input/x2_gm.bin", bFileSize, bHost, bFileSize); |
| 221 | aclrtMemcpy(bDevice, bFileSize, bHost, bFileSize, ACL_MEMCPY_HOST_TO_DEVICE); | 471 | aclrtMemcpy(bDevice, bFileSize, bHost, bFileSize, ACL_MEMCPY_HOST_TO_DEVICE); |
| 222 | 472 | ||
| 473 | + uint8_t *biasHost; | ||
| 474 | + uint8_t *biasDevice; | ||
| 475 | + if constexpr (HAS_BIAS_TENSOR) { | ||
| 476 | + aclrtMallocHost((void **)(&biasHost), biasFileSize); | ||
| 477 | + aclrtMalloc((void **)&biasDevice, biasFileSize, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 478 | + ReadFile("./input/bias_gm.bin", biasFileSize, biasHost, biasFileSize); | ||
| 479 | + aclrtMemcpy(biasDevice, biasFileSize, biasHost, biasFileSize, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 480 | + } | ||
| 481 | + | ||
| 223 | uint8_t *cHost; | 482 | uint8_t *cHost; |
| 224 | uint8_t *cDevice; | 483 | uint8_t *cDevice; |
| 225 | aclrtMallocHost((void **)(&cHost), cFileSize); | 484 | aclrtMallocHost((void **)(&cHost), cFileSize); |
| 226 | aclrtMalloc((void **)&cDevice, cFileSize, ACL_MEM_MALLOC_HUGE_FIRST); | 485 | aclrtMalloc((void **)&cDevice, cFileSize, ACL_MEM_MALLOC_HUGE_FIRST); |
| 227 | 486 | ||
| 228 | - mmad_custom<<<numBlocks, nullptr, stream>>>(aDevice, bDevice, cDevice); | 487 | + mmad_custom<M, N, K><<<numBlocks, nullptr, stream>>>(aDevice, bDevice, biasDevice, cDevice); |
| 229 | aclrtSynchronizeStream(stream); | 488 | aclrtSynchronizeStream(stream); |
| 230 | 489 | ||
| 231 | aclrtMemcpy(cHost, cFileSize, cDevice, cFileSize, ACL_MEMCPY_DEVICE_TO_HOST); | 490 | aclrtMemcpy(cHost, cFileSize, cDevice, cFileSize, ACL_MEMCPY_DEVICE_TO_HOST); |
| @@ -235,6 +494,10 @@ int32_t main(int32_t argc, char *argv[]) | |||
| 235 | aclrtFreeHost(aHost); | 494 | aclrtFreeHost(aHost); |
| 236 | aclrtFree(bDevice); | 495 | aclrtFree(bDevice); |
| 237 | aclrtFreeHost(bHost); | 496 | aclrtFreeHost(bHost); |
| 497 | + if constexpr (HAS_BIAS_TENSOR) { | ||
| 498 | + aclrtFree(biasDevice); | ||
| 499 | + aclrtFreeHost(biasHost); | ||
| 500 | + } | ||
| 238 | aclrtFree(cDevice); | 501 | aclrtFree(cDevice); |
| 239 | aclrtFreeHost(cHost); | 502 | aclrtFreeHost(cHost); |
| 240 | 503 | ||
| @@ -2,7 +2,7 @@ | |||
| 2 | # coding=utf-8 | 2 | # coding=utf-8 |
| 3 | 3 | ||
| 4 | # ---------------------------------------------------------------------------------------------------------- | 4 | # ---------------------------------------------------------------------------------------------------------- |
| 5 | -# Copyright (c) 2025 Huawei Technologies Co., Ltd. | 5 | +# Copyright (c) 2026 Huawei Technologies Co., Ltd. |
| 6 | # This program is free software, you can redistribute it and/or modify it under the terms and conditions of | 6 | # This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 7 | # CANN Open Software License Agreement Version 2.0 (the "License"). | 7 | # CANN Open Software License Agreement Version 2.0 (the "License"). |
| 8 | # Please refer to the License for details. You may not use this file except in compliance with the License. | 8 | # Please refer to the License for details. You may not use this file except in compliance with the License. |
| @@ -11,26 +11,63 @@ | |||
| 11 | # See LICENSE in the root of the software repository for the full text of the License. | 11 | # See LICENSE in the root of the software repository for the full text of the License. |
| 12 | # ---------------------------------------------------------------------------------------------------------- | 12 | # ---------------------------------------------------------------------------------------------------------- |
| 13 | 13 | ||
| 14 | - | 14 | +import argparse |
| 15 | import os | 15 | import os |
| 16 | import numpy as np | 16 | import numpy as np |
| 17 | +from ml_dtypes import bfloat16 | ||
| 17 | 18 | ||
| 18 | 19 | ||
| 19 | -def gen_golden_data(): | 20 | +def gen_golden_data(scenario_num=1): |
| 20 | - M = 32 | 21 | + m=30 |
| 21 | - N = 32 | 22 | + k=70 |
| 22 | - K = 32 | 23 | + n=40 |
| 24 | + if scenario_num == 1: | ||
| 25 | + x1_gm_int8 = np.random.randint(1, 7, [m, k]).astype(np.int8) | ||
| 26 | + x2_gm_int8 = np.random.randint(1, 7, [k, n]).astype(np.int8) | ||
| 27 | + | ||
| 28 | + golden = np.matmul(x1_gm_int8.astype(np.int32), x2_gm_int8.astype(np.int32)).astype(np.int32) | ||
| 29 | + x2_gm_int8 = x2_gm_int8.transpose() | ||
| 30 | + | ||
| 31 | + x1_gm = np.zeros(shape=[m, k // 2]).astype(np.int8) | ||
| 32 | + for i in range(m): | ||
| 33 | + for j in range(k): | ||
| 34 | + if j % 2 == 0: | ||
| 35 | + x1_gm[i][j // 2] = (x1_gm_int8[i][j + 1] << 4) + (x1_gm_int8[i][j] & 0x0f) | ||
| 36 | + x2_gm = np.zeros(shape=[n, k // 2]).astype(np.int8) | ||
| 37 | + for i in range(n): | ||
| 38 | + for j in range(k): | ||
| 39 | + if j % 2 == 0: | ||
| 40 | + x2_gm[i][j // 2] = (x2_gm_int8[i][j + 1] << 4) + (x2_gm_int8[i][j] & 0x0f) | ||
| 41 | + elif scenario_num == 2: | ||
| 42 | + x1_gm = np.random.randint(1, 10, [m, k]).astype(np.int8) | ||
| 43 | + x2_gm = np.random.randint(1, 10, [k, n]).astype(np.int8) | ||
| 44 | + bias_gm = np.random.randint(1, 10, [n]).astype(np.int32) | ||
| 45 | + golden = np.matmul(x1_gm.astype(np.int32), x2_gm.astype(np.int32)).astype(np.int32) + bias_gm | ||
| 46 | + elif scenario_num == 3: | ||
| 47 | + x1_gm = np.random.uniform(1, 10, [m, k]).astype(bfloat16) | ||
| 48 | + x2_gm = np.random.uniform(1, 10, [k, n]).astype(bfloat16) | ||
| 49 | + golden = np.matmul(x1_gm.astype(np.float32), x2_gm.astype(np.float32)).astype(np.float32) * 2 | ||
| 50 | + x2_gm = x2_gm.transpose() | ||
| 51 | + else: | ||
| 52 | + x1_gm = np.random.uniform(1, 10, [m, k]).astype(np.float32) | ||
| 53 | + x2_gm = np.random.uniform(1, 10, [k, n]).astype(np.float32) | ||
| 54 | + bias_gm = np.random.uniform(1, 10, [n]).astype(np.float32) | ||
| 55 | + golden = np.matmul(x1_gm.astype(np.float32), x2_gm.astype(np.float32)).astype(np.float32) + bias_gm | ||
| 56 | + x1_gm = x1_gm.transpose() | ||
| 57 | + x2_gm = x2_gm.transpose() | ||
| 58 | + | ||
| 59 | + os.makedirs("input", exist_ok=True) | ||
| 60 | + os.makedirs("output", exist_ok=True) | ||
| 23 | 61 | ||
| 24 | - x1_gm = np.random.uniform(1, 10, [M, K]).astype(np.float16) | ||
| 25 | - x2_gm = np.random.uniform(1, 10, [K, N]).astype(np.float16) | ||
| 26 | - golden = np.matmul(x1_gm.astype(np.float32), x2_gm.astype(np.float32)) | ||
| 27 | - golden = golden.astype(np.float32) | ||
| 28 | - os.system("mkdir -p input") | ||
| 29 | - os.system("mkdir -p output") | ||
| 30 | x1_gm.tofile("./input/x1_gm.bin") | 62 | x1_gm.tofile("./input/x1_gm.bin") |
| 31 | x2_gm.tofile("./input/x2_gm.bin") | 63 | x2_gm.tofile("./input/x2_gm.bin") |
| 64 | + if scenario_num in [2, 4]: | ||
| 65 | + bias_gm.tofile("./input/bias_gm.bin") | ||
| 32 | golden.tofile("./output/golden.bin") | 66 | golden.tofile("./output/golden.bin") |
| 33 | 67 | ||
| 34 | 68 | ||
| 35 | if __name__ == "__main__": | 69 | if __name__ == "__main__": |
| 36 | - gen_golden_data() | 70 | + parser = argparse.ArgumentParser() |
| 71 | + parser.add_argument('-scenarioNum', type=int, default=1, choices=range(1, 5)) | ||
| 72 | + args = parser.parse_args() | ||
| 73 | + gen_golden_data(args.scenarioNum) | ||
Mexamples/01_simd_cpp_api/02_features/03_basic_api/01_matrix_compute/mmad/scripts/verify_result.py+21-11
| @@ -2,7 +2,7 @@ | |||
| 2 | # coding=utf-8 | 2 | # coding=utf-8 |
| 3 | 3 | ||
| 4 | # ---------------------------------------------------------------------------------------------------------- | 4 | # ---------------------------------------------------------------------------------------------------------- |
| 5 | -# Copyright (c) 2025 Huawei Technologies Co., Ltd. | 5 | +# Copyright (c) 2026 Huawei Technologies Co., Ltd. |
| 6 | # This program is free software, you can redistribute it and/or modify it under the terms and conditions of | 6 | # This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 7 | # CANN Open Software License Agreement Version 2.0 (the "License"). | 7 | # CANN Open Software License Agreement Version 2.0 (the "License"). |
| 8 | # Please refer to the License for details. You may not use this file except in compliance with the License. | 8 | # Please refer to the License for details. You may not use this file except in compliance with the License. |
| @@ -11,10 +11,9 @@ | |||
| 11 | # See LICENSE in the root of the software repository for the full text of the License. | 11 | # See LICENSE in the root of the software repository for the full text of the License. |
| 12 | # ---------------------------------------------------------------------------------------------------------- | 12 | # ---------------------------------------------------------------------------------------------------------- |
| 13 | 13 | ||
| 14 | - | 14 | +import argparse |
| 15 | import sys | 15 | import sys |
| 16 | import numpy as np | 16 | import numpy as np |
| 17 | -import tensorflow as tf | ||
| 18 | 17 | ||
| 19 | 18 | ||
| 20 | # for float32 | 19 | # for float32 |
| @@ -23,9 +22,13 @@ absolute_tol = 1e-9 | |||
| 23 | error_tol = 1e-4 | 22 | error_tol = 1e-4 |
| 24 | 23 | ||
| 25 | 24 | ||
| 26 | -def verify_result(output, golden): | 25 | +def verify_result(scenario_num, output, golden): |
| 27 | - output = np.fromfile(output, dtype=np.float32).reshape(-1) | 26 | + if scenario_num > 2: |
| 28 | - golden = np.fromfile(golden, dtype=np.float32).reshape(-1) | 27 | + output_type = np.float32 |
| 28 | + else: | ||
| 29 | + output_type = np.int32 | ||
| 30 | + output = np.fromfile(output, dtype=output_type).reshape(-1) | ||
| 31 | + golden = np.fromfile(golden, dtype=output_type).reshape(-1) | ||
| 29 | different_element_results = np.isclose(output, | 32 | different_element_results = np.isclose(output, |
| 30 | golden, | 33 | golden, |
| 31 | rtol=relative_tol, | 34 | rtol=relative_tol, |
| @@ -36,10 +39,12 @@ def verify_result(output, golden): | |||
| 36 | real_index = different_element_indexes[index] | 39 | real_index = different_element_indexes[index] |
| 37 | golden_data = golden[real_index] | 40 | golden_data = golden[real_index] |
| 38 | output_data = output[real_index] | 41 | output_data = output[real_index] |
| 39 | - print( | 42 | + if scenario_num > 2: |
| 40 | - "data index: %06d, expected: %-.9f, actual: %-.9f, rdiff: %-.6f" % | 43 | + print("data index: %06d, expected: %-.9f, actual: %-.9f, rdiff: %-.6f" % |
| 41 | - (real_index, golden_data, output_data, | 44 | + (real_index, golden_data, output_data, abs(output_data - golden_data) / golden_data)) |
| 42 | - abs(output_data - golden_data) / golden_data)) | 45 | + else: |
| 46 | + print("data index: %06d, expected: %d, actual: %d, rdiff: %-.6f" % | ||
| 47 | + (real_index, golden_data, output_data, abs(output_data - golden_data) / golden_data)) | ||
| 43 | if index == 100: | 48 | if index == 100: |
| 44 | break | 49 | break |
| 45 | error_ratio = float(different_element_indexes.size) / golden.size | 50 | error_ratio = float(different_element_indexes.size) / golden.size |
| @@ -48,8 +53,13 @@ def verify_result(output, golden): | |||
| 48 | 53 | ||
| 49 | 54 | ||
| 50 | if __name__ == '__main__': | 55 | if __name__ == '__main__': |
| 56 | + parser = argparse.ArgumentParser() | ||
| 57 | + parser.add_argument('-scenarioNum', type=int, default=1, choices=range(1, 5)) | ||
| 58 | + parser.add_argument('output', type=str) | ||
| 59 | + parser.add_argument('golden', type=str) | ||
| 60 | + args = parser.parse_args() | ||
| 51 | try: | 61 | try: |
| 52 | - res = verify_result(sys.argv[1], sys.argv[2]) | 62 | + res = verify_result(args.scenarioNum, args.output, args.golden) |
| 53 | if not res: | 63 | if not res: |
| 54 | raise ValueError("[ERROR] result error") | 64 | raise ValueError("[ERROR] result error") |
| 55 | else: | 65 | else: |
图加白色背景