已合并
add add_high_performance #1262
cellur_z创建于 3月28日
add add_high_performance #1262
已合并
共 8 个文件变更+1365-0
| @@ -0,0 +1,44 @@ | |||
| 1 | +# ---------------------------------------------------------------------------------------------------------- | ||
C | |||
| 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 | ||
| 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. | ||
| 6 | +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | +# See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | +# ---------------------------------------------------------------------------------------------------------- | ||
| 10 | + | ||
| 11 | + | ||
| 12 | +cmake_minimum_required(VERSION 3.16) | ||
| 13 | + | ||
| 14 | +find_package(ASC REQUIRED HINTS $ENV{ASCEND_INSTALL_PATH}/compiler/tikcpp/ascendc_kernel_cmake) | ||
| 15 | + | ||
| 16 | +project(kernel_samples LANGUAGES ASC CXX) | ||
| 17 | + | ||
| 18 | +set(CASE_TYPE 0 CACHE STRING "Case type to compile (0-7)") | ||
| 19 | + | ||
| 20 | +add_executable(demo | ||
| 21 | + add.asc | ||
| 22 | +) | ||
| 23 | + | ||
| 24 | +target_compile_definitions(demo PRIVATE | ||
| 25 | + CASE_TYPE=${CASE_TYPE} | ||
| 26 | +) | ||
| 27 | + | ||
| 28 | +target_link_libraries(demo PRIVATE | ||
| 29 | + tiling_api | ||
| 30 | + register | ||
| 31 | + platform | ||
| 32 | + m | ||
| 33 | + dl | ||
| 34 | +) | ||
| 35 | + | ||
| 36 | +# ====================================================================================== | ||
| 37 | +# NPU 编译选项配置 | ||
| 38 | +# | ||
| 39 | +# 说明： | ||
| 40 | +# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。 | ||
| 41 | +# ====================================================================================== | ||
| 42 | +target_compile_options(demo PRIVATE | ||
| 43 | + $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201> | ||
| 44 | +) | ||
| @@ -0,0 +1,570 @@ | |||
| 1 | +# Add性能调优样例 | ||
| 2 | + | ||
| 3 | +## 概述 | ||
| 4 | + | ||
| 5 | +本样例以加法为例，介绍基于静态Tensor方式编程的性能调优方法。整个调优过程分为七个步骤（case 0-6），逐步展示从标量运算到向量运算、从单核到多核、从基础实现到深度优化的完整调优路径。 | ||
| 6 | + | ||
| 7 | +**优化路径**： | ||
| 8 | +- Case 0: 单核标量版本（基准） | ||
| 9 | +- Case 1: 单核向量版本 | ||
| 10 | +- Case 2: 多核均匀切分 + 小块搬运 | ||
| 11 | +- Case 3: 多核均匀切分 + 大块搬运 | ||
| 12 | +- Case 4: 多核均匀切分 + 双缓冲优化 | ||
| 13 | +- Case 5: 多核均匀切分 + 双缓冲 + L2Cache bypass | ||
| 14 | +- Case 6: 多核均匀切分 + 双缓冲 + L2Cache bypass + 避免Bank Conflict | ||
| 15 | + | ||
| 16 | +## 支持的产品 | ||
| 17 | + | ||
| 18 | +- Atlas A3 训练系列产品/Atlas A3 推理系列产品 | ||
| 19 | +- Atlas A2 训练系列产品/Atlas A2 推理系列产品 | ||
B cmakelist工程里有# <COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510>，这里是不是需要加上A5支持？    cellur_z 3月30日 评论： 3月30日 评论： | |||
| 20 | + | ||
| 21 | +## 目录结构介绍 | ||
| 22 | + | ||
| 23 | +``` | ||
| 24 | +├── 00_add_high_performance | ||
| 25 | +│ ├── scripts | ||
| 26 | +│ │ ├── gen_data.py // 输入数据和真值数据生成脚本文件 | ||
| 27 | +│ │ └── verify_result.py // 真值对比文件 | ||
| 28 | +│ ├── CMakeLists.txt // 编译工程文件 | ||
| 29 | +│ ├── data_utils.h // 数据读入写出函数 | ||
| 30 | +│ ├── add.asc // Ascend C样例实现（包含7个优化case） | ||
| 31 | +``` | ||
| 32 | + | ||
| 33 | +## 样例描述 | ||
| 34 | + | ||
| 35 | +- **样例功能**： | ||
| 36 | + | ||
| 37 | + 样例实现的是固定shape为8192×8192的两个矩阵相加。 | ||
| 38 | + | ||
| 39 | + Add的计算公式为： | ||
| 40 | + | ||
| 41 | +$$ | ||
| 42 | + z = x + y | ||
| 43 | +$$ | ||
| 44 | + | ||
| 45 | + - x：输入，形状为[8192, 8192]，数据类型为half； | ||
| 46 | + - y：输入，形状为[8192, 8192]，数据类型为half； | ||
| 47 | + - z：输出，形状为[8192, 8192]，数据类型为half； | ||
| 48 | + | ||
| 49 | +**表1 AI Core 性能指标字段说明表** | ||
| 50 | +| 字段名 | 字段含义 | | ||
| 51 | +|:---:|:---| | ||
| 52 | +|Task Duration(μs)|Task整体耗时，包含调度到加速器的时间、加速器上的执行时间以及响应结束时间。| | ||
 单位有的是us，有的是μs | |||
| 53 | +|aiv_time|Task在AI Vector Core上的理论执行时间，单位为μs。| | ||
| 54 | +| aiv_vec_time(μs) | vec类型指令（向量类运算指令）耗时，单位μs。 | | ||
| 55 | +| aiv_vec_ratio | vec类型指令（向量类运算指令）的cycle数在total cycle数中的占用比。 | | ||
| 56 | +| aiv_scalar_time(μs) | scalar类型指令（标量类运算指令）耗时，单位μs。 | | ||
| 57 | +| aiv_scalar_ratio | scalar类型指令（标量类运算指令）的cycle数在total cycle数中的占用比。 | | ||
| 58 | +| aiv_mte2_time(μs) | mte2类型指令（GM->UB搬运类指令）耗时，单位μs。 | | ||
| 59 | +| aiv_mte2_ratio | mte2类型指令（GM->UB搬运类指令）的cycle数在total cycle数中的占用比。 | | ||
| 60 | +| aiv_mte3_time(μs) | mte3类型指令（UB->GM搬运类指令）耗时，单位μs。 | | ||
| 61 | +| aiv_mte3_ratio | mte3类型指令（UB->GM搬运类指令）的cycle数在total cycle数中的占用比。 | | ||
| 62 | + | ||
| 63 | +## 样例实现 | ||
| 64 | + | ||
| 65 | +### Case 0: 单核标量版本（基准程序） | ||
| 66 | + | ||
| 67 | +**实现方式**：参考 `KernelAdd::ProcessScalar()` 函数实现 | ||
| 68 | + | ||
| 69 | +基准程序实现了`half`类型的两组输入数据的加法，采用`for`循环`scalar`运算的方式进行计算。 | ||
| 70 | + | ||
| 71 | +**关键代码**： | ||
| 72 | +```cpp | ||
| 73 | +for (uint32_t i = 0; i < curLen; i++) { | ||
| 74 | + float xVal = (float)xLocal.GetValue(i); | ||
| 75 | + float yVal = (float)yLocal.GetValue(i); | ||
| 76 | + zLocal.SetValue(i, (half)(xVal + yVal)); | ||
| 77 | +} | ||
| 78 | +``` | ||
| 79 | + | ||
| 80 | +**配置**： | ||
| 81 | +- 单核标量运算 | ||
| 82 | +- `dataCopyLen = 4096` 为每次搬运的数据量元素个数 | ||
| 83 | +- 单次搬运的数据量为 8192 Byte，单次scalar处理的数据量为4 Byte | ||
| 84 | + | ||
| 85 | +**性能数据**： | ||
| 86 | + | ||
| 87 | +| Task Duration(μs) | aiv_time(μs) | aiv_vec_time(μs) | aiv_vec_ratio | aiv_scalar_time(μs) | aiv_scalar_ratio | aiv_mte2_time(μs) | aiv_mte2_ratio | aiv_mte3_time(μs) | aiv_mte3_ratio | | ||
| 88 | +|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:| | ||
| 89 | +| 1239734.12 | 1239733.62 | 0.015 | 0% | 1233756.661 | 99.5% | 5889.469 | 0.005 | 2426.024 | 0.002 | | ||
| 90 | + | ||
| 91 | +**优化效果分析**： | ||
| 92 | +- 端到端耗时：**1239734.12μs**（约1.24秒） | ||
| 93 | +- 标量指令耗时：1233756.661μs，占比 **99.5%** | ||
| 94 | +- 向量指令耗时：0.015μs，占比接近0% | ||
| 95 | +- 性能瓶颈：标量运算串行执行，无法利用硬件并行能力。该场景仅作为Add运算性能对比样例，在实际业务场景中不建议用户使用Scalar运算 | ||
| 96 | + | ||
| 97 | +**原理说明**： | ||
| 98 | +- 标量运算每次只能处理1个数据元素，需要逐元素循环 | ||
| 99 | +- AI Core的硬件优势在于向量/矩阵并行计算，标量运算无法发挥硬件能力 | ||
| 100 | + | ||
| 101 | +**性能优化建议**： | ||
| 102 | +> ⚠️ **避免标量循环，使用向量指令** | ||
| 103 | +> | ||
| 104 | +> 在Ascend C编程中，应避免使用`for`循环配合`GetValue/SetValue`的标量运算。使用`AscendC::Add`等向量指令可带来数量级的性能提升。 | ||
| 105 | + | ||
| 106 | +--- | ||
| 107 | + | ||
| 108 | +### Case 1: 单核向量版本 | ||
| 109 | + | ||
| 110 | +**实现方式**：参考 `KernelAdd::ProcessSingle()` 函数实现 | ||
| 111 | + | ||
| 112 | +将标量运算转换为向量运算，使用`AscendC::Add`向量指令替代标量循环，大幅提升计算效率。。 | ||
| 113 | + | ||
| 114 | +**关键代码**： | ||
| 115 | +```cpp | ||
| 116 | +AscendC::Add(zLocal, xLocal, yLocal, curLen); | ||
| 117 | +``` | ||
| 118 | + | ||
| 119 | +**配置**： | ||
| 120 | +- 单核运算 | ||
| 121 | +- `dataCopyLen = 4096` 为每次搬运的数据量元素个数 | ||
| 122 | +- 单次搬运操作`DataCopy`的数据量为 8192 Byte | ||
| 123 | +- 单次`Add`处理两个输入`Tensor`，处理的总数据量为 16384 Byte | ||
| 124 | + | ||
| 125 | + | ||
| 126 | +**性能数据**： | ||
| 127 | + | ||
| 128 | +| Task Duration(μs) | aiv_time(μs) | aiv_vec_time(μs) | aiv_vec_ratio | aiv_scalar_time(μs) | aiv_scalar_ratio | aiv_mte2_time(μs) | aiv_mte2_ratio | aiv_mte3_time(μs) | aiv_mte3_ratio | | ||
| 129 | +|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:| | ||
| 130 | +| 6906.26 | 6905.68 | 761.649 | 0.11 | 337.964 | 0.049 | 6007.112 | 0.87 | 2705.84 | 0.392 | | ||
| 131 | + | ||
| 132 | +**优化效果分析**： | ||
| 133 | +- 端到端性能：6906.26μs，相比Case 0提升 **179.6倍** | ||
| 134 | +- 标量指令耗时：从1233756.66μs降至337.964μs，减少 **99.97%** | ||
| 135 | +- 向量指令耗时：761.649μs，占比11.1% | ||
| 136 | +- 数据搬运耗时：6007.112μs，占比87%，成为新瓶颈 | ||
| 137 | + | ||
| 138 | +**原理说明**： | ||
| 139 | +- 向量指令单次可处理多个数据元素（本例中单次处理4096个half元素） | ||
| 140 | +- 向量单元的并行计算能力远超标量单元 | ||
| 141 | +- 但数据搬运成为瓶颈，说明计算速度已经快于数据供给，单核时数据搬运的请求量不足，导致带宽未用满 | ||
| 142 | + | ||
| 143 | +**性能优化建议**： | ||
| 144 | +> 💡 **使用向量指令替代标量循环** | ||
| 145 | +> | ||
| 146 | +> 使用`AscendC::Add`、`AscendC::Mul`等向量API替代逐元素的标量循环，可充分利用AI Core的向量计算单元，获得百倍以上的性能提升。 | ||
| 147 | + | ||
| 148 | +> 💡 **不建议用户仅使用单核** | ||
| 149 | + | ||
| 150 | +**下一步优化方向**： | ||
| 151 | +- 数据搬运（MTE2）占比87%，成为主要瓶颈 | ||
| 152 | +- 需要通过多核并行和增大搬运粒度来提升带宽利用率 | ||
| 153 | + | ||
| 154 | +--- | ||
| 155 | + | ||
| 156 | +### Case 2: 多核均匀切分 + 小块搬运 | ||
| 157 | + | ||
| 158 | +**实现方式**：参考 `KernelAdd::Process()` 函数实现 | ||
| 159 | + | ||
| 160 | +开启多核并行计算，将8192×8192的矩阵切分到多个AIV Core上并行处理，采用均匀切分策略。 | ||
| 161 | + | ||
| 162 | +**配置**： | ||
| 163 | +- 行方向切分6份，列方向切分8份，将数据均匀切分至48个核运算 | ||
| 164 | +- `dataCopyLen = 4096` 为每次切分的数据量元素个数 | ||
| 165 | +- 单次搬运操作`DataCopy`的数据量为 8192 Byte | ||
| 166 | +- 单次`Add`处理两个输入`Tensor`，处理的总数据量为 16384 Byte | ||
| 167 | + | ||
| 168 | +**关键代码**： | ||
| 169 | +```cpp | ||
| 170 | +// 均匀切分计算每个核处理的行列数 | ||
| 171 | +uint32_t baseCoreN = totalN / splitN; | ||
| 172 | +uint32_t remainderN = totalN % splitN; | ||
| 173 | +if (blockIdxN < remainderN) { | ||
| 174 | + actualCoreN = baseCoreN + 1; // 均匀分配余数 | ||
| 175 | +} | ||
| 176 | +``` | ||
| 177 | + | ||
| 178 | +**性能数据**： | ||
| 179 | + | ||
| 180 | +| Task Duration(μs) | aiv_time(μs) | aiv_vec_time(μs) | aiv_vec_ratio | aiv_scalar_time(μs) | aiv_scalar_ratio | aiv_mte2_time(μs) | aiv_mte2_ratio | aiv_mte3_time(μs) | aiv_mte3_ratio | | ||
| 181 | +|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:| | ||
| 182 | +| 312.22 | 306.29 | 15.897 | 0.052 | 8.446 | 0.028 | 223.329 | 0.729 | 54.891 | 0.179 | | ||
| 183 | + | ||
| 184 | +**优化效果分析**： | ||
| 185 | +- 端到端任务耗时为312.22μs，相比Case 1提升 **22.1倍** | ||
| 186 | +- 数据搬运MTE2耗时为223.329μs，占比72.9% | ||
| 187 | + | ||
| 188 | +**原理说明**： | ||
| 189 | +- 48个AI Core并行处理，理论上可获得48倍加速 | ||
| 190 | +- 实际加速比22.1倍，低于理论值的原因： | ||
| 191 | + - 数据搬运仍是瓶颈（mte2占比72.9%） | ||
| 192 | +- 均匀切分确保各核负载均衡，避免长尾效应 | ||
| 193 | + | ||
| 194 | +**性能优化建议**： | ||
| 195 | +> 💡 **充分利用多核并行，采用均匀切分策略** | ||
| 196 | +> | ||
| 197 | +> 1. 将数据均匀切分到多个AI Core，实现并行计算 | ||
| 198 | +> 2. 使用均匀切分策略（余数分配到前几个核），确保负载均衡 | ||
| 199 | +> 3. 切分粒度需考虑：核数、数据量、UB空间大小 | ||
| 200 | + | ||
| 201 | +**下一步优化方向**： | ||
| 202 | +- MTE2占比72.9%，MTE3占比17.9%，搬运仍是瓶颈 | ||
| 203 | +- 计算仅占5.2%，说明"计算快、搬运慢" | ||
| 204 | +- 可通过增大单次搬运数据量提升带宽利用率 | ||
| 205 | +--- | ||
| 206 | + | ||
| 207 | +### Case 3: 多核均匀切分 + 大块数据搬运 | ||
| 208 | + | ||
| 209 | +**实现方式**：参考 `KernelAdd::Process()` 函数实现 | ||
| 210 | + | ||
| 211 | +为了充分利用带宽资源，增大搬运指令的数据量。 | ||
| 212 | + | ||
| 213 | +**配置**： | ||
| 214 | +- 行方向切分6份，列方向切分8份，将数据均匀切分至48个核运算 | ||
| 215 | +- `dataCopyLen = 16384` 为每次切分的数据量元素个数（4倍于Case2） | ||
| 216 | +- 单次搬运操作`DataCopy`的数据量为 32678 Byte | ||
| 217 | +- 单次`Add`处理两个输入`Tensor`，处理的总数据量为 65536 Byte | ||
| 218 | + | ||
| 219 | + | ||
| 220 | +**性能数据**： | ||
| 221 | + | ||
| 222 | +| Task Duration(μs) | aiv_time(μs) | aiv_vec_time(μs) | aiv_vec_ratio | aiv_scalar_time(μs) | aiv_scalar_ratio | aiv_mte2_time(μs) | aiv_mte2_ratio | aiv_mte3_time(μs) | aiv_mte3_ratio | | ||
| 223 | +|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:| | ||
| 224 | +| 270.76 |266.22 | 12.853 | 0.048 | 4.949 | 0.019 | 188.331 | 0.707 | 54.936 | 0.206 | | ||
| 225 | + | ||
| 226 | +**优化效果分析**： | ||
| 227 | +- 端到端性能：270.76μs，相比Case 2提升 **15.3%** | ||
| 228 | +- 通过增大单次搬运的数据量，MTE2耗时从223.329μs降至188.331μs，减少 **15.7%** | ||
| 229 | + | ||
| 230 | +**原理说明**： | ||
| 231 | +- 增大单次搬运数据量可减少搬运次数 | ||
| 232 | +- 每次搬运的启动开销固定，搬运更多数据可摊薄开销 | ||
| 233 | +- 更大的连续数据块可更好地利用内存带宽 | ||
| 234 | +- 但数据量受限于UB空间大小（本例UB需容纳x、y、z三份数据） | ||
| 235 | + | ||
| 236 | +**性能优化建议**： | ||
| 237 | +> 💡 **增大单次数据搬运量，减少搬运次数** | ||
| 238 | +> | ||
| 239 | +> 1. 在UB空间允许的范围内，尽量增大`dataCopyLen` | ||
| 240 | +> 2. 使用连续的大块数据搬运，避免频繁的小数据块搬运 | ||
| 241 | +> 3. 需平衡UB空间使用和搬运效率 | ||
| 242 | + | ||
| 243 | + | ||
| 244 | +> ⚠️ **注意：dataCopyLen并非越大越好** | ||
| 245 | +> | ||
| 246 | +> 在Case 3的基础上，如果进一步增大dataCopyLen（如从16384增至16512），端到端性能反而略有下降（271.5μs vs 270.76μs，+0.27%）。建议结合数据总量、UB空间、对齐要求综合考虑，确定最优的dataCopyLen值。 | ||
| 247 | + | ||
| 248 | +--- | ||
| 249 | + | ||
| 250 | +### Case 4: 双缓冲优化 | ||
| 251 | + | ||
| 252 | +**实现方式**：参考 `KernelAdd::ProcessDoubleBuffer()` 函数实现 | ||
| 253 | + | ||
| 254 | +采用双缓冲（Double Buffer）技术，实现数据搬运与计算的流水线并行，隐藏内存访问延迟。 | ||
| 255 | + | ||
| 256 | +**关键代码**： | ||
| 257 | +```cpp | ||
| 258 | +// Ping-Pong双缓冲地址 | ||
| 259 | +static constexpr uint32_t xAddrPing = 0; | ||
| 260 | +static constexpr uint32_t yAddrPing = MAX_DATA_COPY_LEN * sizeof(half); | ||
| 261 | +static constexpr uint32_t zAddrPing = yAddrPing + MAX_DATA_COPY_LEN * sizeof(half); | ||
| 262 | +static constexpr uint32_t xAddrPong = zAddrPing + MAX_DATA_COPY_LEN * sizeof(half); | ||
| 263 | +static constexpr uint32_t yAddrPong = xAddrPong + MAX_DATA_COPY_LEN * sizeof(half); | ||
| 264 | +static constexpr uint32_t zAddrPong = yAddrPong + MAX_DATA_COPY_LEN * sizeof(half); | ||
| 265 | + | ||
| 266 | +// 双缓冲流水线：交替使用两个事件ID和两组缓冲区 | ||
| 267 | +for (uint32_t loopIdx = 0; loopIdx < totalBlocks; loopIdx++) { | ||
| 268 | + int32_t eventID = (loopIdx % 2 == 0 ? EVENT_ID0 : EVENT_ID1); | ||
| 269 | + AscendC::LocalTensor<half> &xLocal = (loopIdx % 2 == 0 ? xPing : xPong); | ||
| 270 | + // ... 数据搬运和计算，使用对应的eventID同步 | ||
| 271 | + AscendC::Add(zLocal, xLocal, yLocal, curLen); | ||
| 272 | +} | ||
| 273 | +``` | ||
| 274 | + | ||
| 275 | +**配置**： | ||
| 276 | +- 行方向切分6份，列方向切分8份，将数据均匀切分至48个核运算 | ||
| 277 | +- `dataCopyLen = 16384` 为每次切分的数据量元素个数 | ||
| 278 | +- 单次搬运操作`DataCopy`的数据量为 32678 Byte | ||
| 279 | +- 单次`Add`处理两个输入`Tensor`，处理的总数据量为 65536 Byte | ||
| 280 | +- 将待处理的数据一分为二，由此数据的进出搬运和Vector计算实现并行执行 | ||
| 281 | + | ||
| 282 | +**内存布局**： | ||
| 283 | + | ||
| 284 | +TODO： 待配图 | ||
| 285 | +``` | ||
| 286 | +UB内存分配（双缓冲）： | ||
| 287 | +┌──────────────┐ | ||
| 288 | +│ xPing │ 0x00000 | ||
| 289 | +│ 16384*2B │ | ||
| 290 | +├──────────────┤ | ||
| 291 | +│ yPing │ 0x08000 (32768) | ||
| 292 | +│ 16384*2B │ | ||
| 293 | +├──────────────┤ | ||
| 294 | +│ zPing │ 0x10000 (65536) | ||
| 295 | +│ 16384*2B │ | ||
| 296 | +├──────────────┤ | ||
| 297 | +│ xPong │ 0x18000 (98304) | ||
| 298 | +│ 16384*2B │ | ||
| 299 | +├──────────────┤ | ||
| 300 | +│ yPong │ 0x20000 (131072) | ||
| 301 | +│ 16384*2B │ | ||
| 302 | +├──────────────┤ | ||
| 303 | +│ zPong │ 0x28000 (163840) | ||
| 304 | +│ 16384*2B │ | ||
| 305 | +└──────────────┘ | ||
| 306 | +``` | ||
| 307 | + | ||
| 308 | +**性能数据**： | ||
| 309 | + | ||
| 310 | +| Task Duration(μs) | aiv_time(μs) | aiv_vec_time(μs) | aiv_vec_ratio | aiv_scalar_time(μs) | aiv_scalar_ratio | aiv_mte2_time(μs) | aiv_mte2_ratio | aiv_mte3_time(μs) | aiv_mte3_ratio | | ||
| 311 | +|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:| | ||
| 312 | +| 265.96 | 261.15 | 12.853 | 0.049 | 4.332 | 0.017 | 251.088 | 0.961 | 95.478 | 0.336 | | ||
| 313 | + | ||
| 314 | +**优化效果分析**： | ||
| 315 | +- 端到端性能：265.96μs，相比Case 3提升 **1.8%** | ||
| 316 | +- MTE2耗时从188.331μs增至251.088μs（+33.3%），MTE3耗时从54.936μs增至95.478μs（+73.8%），此时从串行的纯读带宽变成混合读写带宽，因此耗时增加，用户应需更多地关注端到端耗时的减少 | ||
| 317 | +- 由于开启了双缓冲，在流水中搬运和计算并行执行，隐藏了部分延迟 | ||
| 318 | + | ||
| 319 | +**原理说明**： | ||
| 320 | +- **Ping-Pong机制**： | ||
| 321 | + - Ping缓冲区进行计算时，Pong缓冲区进行数据搬运 | ||
| 322 | + - 交替执行，实现计算与搬运的流水线并行 | ||
| 323 | +- **事件同步**： | ||
| 324 | + - 使用两个Event ID分别管理Ping和Pong的同步 | ||
| 325 | + - 确保"上一块写完"才能"下一块读入" | ||
| 326 | +- **空间翻倍**：需要2倍空间存储Ping和Pong数据 | ||
| 327 | + | ||
| 328 | +**性能优化建议**： | ||
| 329 | +> 💡 **使用双缓冲实现搬运与计算并行** | ||
| 330 | +> | ||
| 331 | +> 1. 当计算与搬运时间相近时，双缓冲收益最大 | ||
| 332 | +> 2. 需要足够的UB空间（约2倍单缓冲空间） | ||
| 333 | +> 3. 使用独立的Event ID管理两组缓冲区的同步 | ||
| 334 | + | ||
| 335 | +**下一步优化方向**： | ||
| 336 | +- 双缓冲的收益有限，说明瓶颈在搬运速度本身 | ||
| 337 | +- 可尝试L2 Cache优化来提升搬运效率 | ||
| 338 | +--- | ||
| 339 | + | ||
| 340 | +### Case 5: 双缓冲 + L2 Cache bypass | ||
| 341 | + | ||
| 342 | +**实现方式**：参考 `KernelAdd::ProcessDoubleBufferL2Bypass()` 函数实现（内部调用 `ProcessDoubleBuffer()`，区别在于 `Init` 时设置 `enableL2Bypass=true`） | ||
| 343 | + | ||
| 344 | +在双缓冲基础上，对于只需要载入一次的数据量可以设置L2 Cache bypass，直接从HBM载入到AICORE内部。 | ||
| 345 | + | ||
| 346 | +**关键代码**： | ||
| 347 | +```cpp | ||
| 348 | +// 在Init时设置L2 Cache bypass | ||
| 349 | +if (enableL2Bypass) { | ||
| 350 | + xGm.SetL2CacheHint(AscendC::CacheMode::CACHE_MODE_DISABLE); | ||
| 351 | + yGm.SetL2CacheHint(AscendC::CacheMode::CACHE_MODE_DISABLE); | ||
| 352 | +} | ||
| 353 | +// ProcessDoubleBufferL2Bypass内部调用ProcessDoubleBuffer | ||
| 354 | +``` | ||
| 355 | + | ||
| 356 | +**配置**： | ||
| 357 | +- 行方向切分6份，列方向切分8份，将数据均匀切分至48个核运算 | ||
| 358 | +- `dataCopyLen = 16384` 为每次切分的数据量元素个数 | ||
| 359 | +- 单次搬运操作`DataCopy`的数据量为 32678 Byte | ||
| 360 | +- 单次`Add`处理两个输入`Tensor`，处理的总数据量为 65536 Byte | ||
| 361 | +- 将待处理的数据一分为二，由此数据的进出搬运和Vector计算实现并行执行 | ||
| 362 | + | ||
| 363 | +**L2 Cache策略**： | ||
| 364 | +- xGm：禁用L2 Cache（一次性读取） | ||
| 365 | +- yGm：禁用L2 Cache（一次性读取） | ||
| 366 | + | ||
| 367 | +**性能数据**： | ||
| 368 | + | ||
| 369 | +| Task Duration(μs) | aiv_time(μs) | aiv_vec_time(μs) | aiv_vec_ratio | aiv_scalar_time(μs) | aiv_scalar_ratio | aiv_mte2_time(μs) | aiv_mte2_ratio | aiv_mte3_time(μs) | aiv_mte3_ratio | | ||
| 370 | +|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:| | ||
| 371 | +| 187.1| 183.62 | 12.853 | 0.07 | 5.416 | 0.029 | 171.163 | 0.932 | 81.061 | 0.441 | | ||
| 372 | + | ||
| 373 | +**优化效果分析**： | ||
| 374 | +- 端到端性能：187.1μs，相比Case 4提升 **29.7%** | ||
| 375 | +- MTE2耗时：从251.088μs降至171.163μs，减少 **31.8%** | ||
| 376 | +- MTE3耗时：从95.478μs降至81.061μs，减少 **15.1%** | ||
| 377 | +- 向量指令耗时：12.853μs，保持不变 | ||
| 378 | + | ||
| 379 | +**原理说明**： | ||
| 380 | +- **L2 Cache的作用**： | ||
| 381 | + - L2 Cache是AI Core和HBM之间的缓存层 | ||
| 382 | + - 重复访问的数据可从L2 Cache读取，速度更快 | ||
| 383 | +- **流式访问特点**： | ||
| 384 | + - Add算子的输入数据只读取一次，不存在数据复用 | ||
| 385 | + - 由于本样例场景数据量较大，超出L2 Cache大小，导致数据写回操作，从而引起额外耗时 | ||
| 386 | + | ||
| 387 | +**性能优化建议**： | ||
| 388 | +> 💡 **合理采用L2 Cache bypass** | ||
| 389 | +> | ||
| 390 | +> 1. 对于只读取一次的输入数据（如本例的x、y），设置`SetL2CacheHint(CACHE_MODE_DISABLE)` | ||
| 391 | +> 2. 对于需要重复访问的数据（如卷积的权重），保留L2 Cache | ||
| 392 | +> 3. 建议用户按照实测数据进行配置优化，在实际的模型和训练场景中，需要结合上下游算子进行合理配置 | ||
| 393 | + | ||
| 394 | +**下一步优化方向**： | ||
| 395 | +- 搬运效率已提升，但向量指令效率仍有优化空间 | ||
| 396 | +- 可尝试优化UB内存布局避免Bank Conflict | ||
| 397 | +--- | ||
| 398 | + | ||
| 399 | +### Case 6: 双缓冲 + L2 Cache bypass + 避免Bank Conflict | ||
| 400 | + | ||
| 401 | +**实现方式**：参考 `KernelAdd::ProcessDoubleBufferBankConflict()` 函数实现 | ||
| 402 | + | ||
| 403 | +在双缓冲+L2 Cache bypass基础上，优化内存地址布局，避免UB（Unified Buffer）的Bank Conflict，实现最优性能。 | ||
| 404 | + | ||
| 405 | +**关键代码**： | ||
| 406 | +```cpp | ||
| 407 | +// 设置L2 Cache bypass | ||
| 408 | +xGm.SetL2CacheHint(AscendC::CacheMode::CACHE_MODE_DISABLE); | ||
| 409 | +yGm.SetL2CacheHint(AscendC::CacheMode::CACHE_MODE_DISABLE); | ||
| 410 | + | ||
| 411 | +// 优化后的地址布局（避免Bank Conflict） | ||
| 412 | +static constexpr uint32_t xAddrPingBC = 0; | ||
| 413 | +static constexpr uint32_t yAddrPingBC = BANK_CONFLICT_DATA_COPY_LEN * sizeof(half); | ||
| 414 | +static constexpr uint32_t xAddrPongBC = MAX_DATA_COPY_LEN * sizeof(half) * 2; | ||
| 415 | +static constexpr uint32_t yAddrPongBC = xAddrPongBC + BANK_CONFLICT_DATA_COPY_LEN * sizeof(half); | ||
| 416 | +static constexpr uint32_t zAddrPingBC = MAX_DATA_COPY_LEN * sizeof(half) * 4; | ||
| 417 | +static constexpr uint32_t zAddrPongBC = zAddrPingBC + BANK_CONFLICT_DATA_COPY_LEN * sizeof(half); | ||
| 418 | +``` | ||
| 419 | + | ||
| 420 | +**配置**： | ||
| 421 | +- 行方向切分6份，列方向切分8份，将数据均匀切分至48个核运算 | ||
| 422 | +- `dataCopyLen = 16256` 为每次切分的数据量元素个数 | ||
| 423 | +- 单次搬运操作`DataCopy`的数据量为 32512 Byte | ||
| 424 | +- 单次`Add`处理两个输入`Tensor`，处理的总数据量为 65204 Byte | ||
| 425 | +- 将待处理的数据一分为二，由此数据的进出搬运和Vector计算实现并行执行 | ||
| 426 | + | ||
| 427 | +**内存布局优化**： | ||
| 428 | + | ||
| 429 | +未优化前的UB Bank内存布局（即case5） | ||
| 430 | +<img src="figure/UBBankConflict.png" width="100%"> | ||
| 431 | + | ||
| 432 | +可以看到这样同时存在一个bank内的读写冲突，一个bankgroup内的读读冲突以及写写冲突。 | ||
| 433 | + | ||
| 434 | +优化后的UB Bank内存布局 | ||
| 435 | +<img src="figure/UBBankConflictResolution.png" width="90%"> | ||
| 436 | + | ||
| 437 | +由于vec指令一拍读取512B的数据（即同时读取8个block的数据），如上图xping、yping的起始地址正好错开了512B，有效消解了ub bank冲突。 | ||
| 438 | + | ||
| 439 | +**Bank Conflict详解**： | ||
| 440 | +- UB分为多个Bank Group，同时读写同一Bank Group会导致冲突 | ||
| 441 | +- 通过调整dataCopyLen（16384→16256）使数据起始地址偏移 | ||
| 442 | +- 确保vec指令一拍访存的数据分布在不同的Bank Group | ||
| 443 | + | ||
| 444 | +**性能数据**： | ||
| 445 | + | ||
| 446 | +| Task Duration(μs) | aiv_time(μs) | aiv_vec_time(μs) | aiv_vec_ratio | aiv_scalar_time(μs) | aiv_scalar_ratio | aiv_mte2_time(μs) | aiv_mte2_ratio | aiv_mte3_time(μs) | aiv_mte3_ratio | | ||
| 447 | +|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:| | ||
| 448 | +| 183 | 178.32 | 7.012 | 0.039 | 11.332 | 0.064 | 169.776 | 0.952 | 43.292 | 0.243 | | ||
| 449 | + | ||
| 450 | +**优化效果分析**： | ||
| 451 | +- 端到端性能：183μs，相比Case 5提升 **2.2%** | ||
| 452 | +- 向量指令耗时：从12.853μs降至7.012μs，减少 **45.4%** | ||
| 453 | +- MTE3耗时：从81.061μs降至43.292μs，减少 **46.6%** | ||
| 454 | + | ||
| 455 | +**原理说明**： | ||
| 456 | +- **Bank Conflict问题**： | ||
| 457 | + - UB（Unified Buffer）分为多个Bank Group | ||
| 458 | + - 向量指令一次读写的数据如果落在同一个Bank，会产生读写冲突 | ||
| 459 | + - 向量指令一次读/写的数据如果落在同一个Bank Group，会产生读读冲突或写写冲突 | ||
| 460 | + - Bank Conflict会导致访存串行化，降低向量指令效率 | ||
| 461 | +- **解决方法**： | ||
| 462 | + - 减小dataCopyLen（16384→16256），使数据起始地址产生偏移 | ||
| 463 | + - 重新设计内存布局，确保同一时刻访问的数据分布在不同Bank | ||
| 464 | + | ||
| 465 | +**性能优化建议**： | ||
| 466 | +> 💡 **优化UB内存布局，避免Bank Conflict** | ||
| 467 | +> | ||
| 468 | +> 1. 当`aiv_vec_time`异常高时，可能存在Bank Conflict | ||
| 469 | +> 2. 通过调整dataCopyLen或内存布局偏移，使数据分布在不同Bank Group | ||
| 470 | +> 4. Bank Conflict优化对vector-bound场景收益明显 | ||
| 471 | + | ||
| 472 | +**最终性能总结**： | ||
| 473 | +- 相比基准Case 0：性能提升 **6774倍**（1239734.12μs → 183μs） | ||
| 474 | +- 相比单核向量Case 1：性能提升 **37.7倍**（6906.26μs → 183μs） | ||
| 475 | + | ||
| 476 | + | ||
| 477 | +--- | ||
| 478 | + | ||
| 479 | +## 性能对比总结 | ||
| 480 | + | ||
| 481 | +| Case | 优化策略 | 核数 | dataCopyLen | Task Duration(μs) | 相对Case 0 | | ||
| 482 | +|:---:|:---|:---:|:---:|:---:|:---:| | ||
| 483 | +| 0 | 单核标量（基准） | 1 | 4096 | 1239734.12 | 1x | | ||
| 484 | +| 1 | 单核向量 | 1 | 4096 | 6906.26 | 179.6x | | ||
| 485 | +| 2 | 多核均匀切分 | 48 | 4096 | 312.22 | 3971x | | ||
| 486 | +| 3 | 增大搬运粒度 | 48 | 16384 | 270.76 | 4581x | | ||
| 487 | +| 4 | 双缓冲 | 48 | 16384 | 265.96 | 4661x | | ||
| 488 | +| 5 | L2 Cache bypass | 48 | 16384 | 187.1 | 6626x | | ||
| 489 | +| 6 | Bank Conflict优化 | 48 | 16256 | 183 | 6774x | | ||
| 490 | + | ||
| 491 | +**优化要点总结**： | ||
| 492 | + | ||
| 493 | +| 优化手段 | 核心原理 | 适用场景 | | ||
| 494 | +|:---|:---:|:---| | ||
| 495 | +| 标量→向量 | 向量指令并行处理多个元素 | 所有计算密集型算子 | | ||
| 496 | +| 单核→多核 | 多核并行，负载均衡 | 大数据量场景 | | ||
| 497 | +| 增大搬运粒度 | 减少搬运次数，摊薄启动开销 | 搬运密集型场景 | | ||
| 498 | +| 双缓冲 | 搬运与计算流水线并行 | 计算与搬运时间相近 | | ||
| 499 | +| L2 Cache bypass | 避免Cache污染，减少开销 | 流式访问（只读一次） | | ||
| 500 | +| Bank Conflict优化 | 优化内存布局，避免访存冲突 | Vector-bound场景 | | ||
| 501 | + | ||
| 502 | +--- | ||
| 503 | + | ||
| 504 | +## 编译运行 | ||
| 505 | + | ||
| 506 | +### 切换Case | ||
| 507 | + | ||
| 508 | +在 cmake 编译时通过 `-DCASE_TYPE=N` 指定要编译的 case： | ||
| 509 | + | ||
| 510 | +```bash | ||
| 511 | +cmake -DCASE_TYPE=6 .. # 编译 case 6（可替换为0-6） | ||
| 512 | +``` | ||
| 513 | + | ||
| 514 | +各 case 说明： | ||
| 515 | +- `0`: 单核标量版本 | ||
| 516 | +- `1`: 单核向量版本 | ||
| 517 | +- `2`: 多核均匀切分 (dataCopyLen=4096) | ||
| 518 | +- `3`: 多核均匀切分 (dataCopyLen=16384) | ||
| 519 | +- `4`: 双缓冲优化 | ||
| 520 | +- `5`: 双缓冲+L2Cache bypass | ||
| 521 | +- `6`: 双缓冲+L2Cache bypass+避免Bank Conflict | ||
| 522 | + | ||
| 523 | +### 编译执行 | ||
| 524 | + | ||
| 525 | +在本样例根目录下执行如下步骤，编译并执行算子： | ||
| 526 | + | ||
| 527 | +- **配置环境变量** | ||
| 528 | + 请根据当前环境上CANN开发套件包的[安装方式](../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。 | ||
| 529 | + - 默认路径，root用户安装CANN软件包 | ||
| 530 | + ```bash | ||
| 531 | + source /μsr/local/Ascend/cann/set_env.sh | ||
| 532 | + ``` | ||
| 533 | + | ||
| 534 | + - 默认路径，非root用户安装CANN软件包 | ||
| 535 | + ```bash | ||
| 536 | + source $HOME/Ascend/cann/set_env.sh | ||
| 537 | + ``` | ||
| 538 | + | ||
| 539 | + - 指定路径install_path，安装CANN软件包 | ||
| 540 | + ```bash | ||
| 541 | + source ${install_path}/cann/set_env.sh | ||
| 542 | + ``` | ||
| 543 | + | ||
| 544 | +- **样例执行** | ||
| 545 | + ```bash | ||
| 546 | + mkdir -p build && cd build; # 创建并进入build目录 | ||
| 547 | + cmake -DCASE_TYPE=6 ..;make -j; # 编译指定case（默认为0，可替换为0-6） | ||
| 548 | + python3 ../scripts/gen_data.py # 生成测试输入数据 | ||
| 549 | + ./demo # 执行（使用编译时指定的case） | ||
| 550 | + python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 验证输出结果是否正确，确认算法逻辑正确 | ||
| 551 | + ``` | ||
| 552 | + | ||
| 553 | + 执行结果如下，说明精度对比成功。 | ||
| 554 | + ```bash | ||
| 555 | + test pass! | ||
| 556 | + ``` | ||
| 557 | + | ||
| 558 | +### 性能分析 | ||
| 559 | + | ||
| 560 | +使用 `msprof` 工具获取详细性能数据： | ||
| 561 | + | ||
| 562 | +```bash | ||
| 563 | +msprof ./demo # 分析性能 | ||
| 564 | +``` | ||
| 565 | + | ||
| 566 | +查看性能分析结果： | ||
| 567 | +```bash | ||
| 568 | +# 查看Task Duration 以及各项数据 | ||
| 569 | +cat ./prof_*/mindstudio_profiler_output/op_summary_*.csv | ||
| 570 | +``` | ||
Aexamples/01_simd_cpp_api/04_best_practices/00_vector_compute_practices/add_high_performance/add.asc+564-0
| @@ -0,0 +1,564 @@ | |||
| 1 | +/** | ||
| 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 | ||
| 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. | ||
| 6 | +* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | +* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | +* See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | +*/ | ||
| 10 | + | ||
| 11 | + | ||
| 12 | +/* ! | ||
| 13 | + * \file add.asc | ||
| 14 | + * \brief Add operator performance tuning sample | ||
| 15 | + */ | ||
| 16 | + | ||
| 17 | +#include "data_utils.h" | ||
| 18 | +#include "kernel_tiling/kernel_tiling.h" | ||
| 19 | +#include "tiling/platform/platform_ascendc.h" | ||
| 20 | +#include "acl/acl.h" | ||
| 21 | +#include "kernel_operator.h" | ||
| 22 | +#include "graph/tensor.h" | ||
| 23 | +#include "tiling/tiling_api.h" | ||
| 24 | + | ||
| 25 | +struct AddCustomTilingData { | ||
| 26 | + uint32_t caseType; | ||
| 27 | + uint32_t totalM; | ||
| 28 | + uint32_t totalN; | ||
| 29 | + uint32_t splitM; | ||
| 30 | + uint32_t splitN; | ||
| 31 | + uint32_t singleCoreM; | ||
| 32 | + uint32_t singleCoreN; | ||
| 33 | + uint32_t dataCopyLen; | ||
| 34 | + uint32_t enableL2Bypass; | ||
| 35 | +}; | ||
| 36 | + | ||
| 37 | +using AscendC::TPosition; | ||
| 38 | +namespace { | ||
| 39 | +constexpr uint32_t MAX_DATA_COPY_LEN = 16384; | ||
| 40 | +constexpr uint32_t BANK_CONFLICT_DATA_COPY_LEN = 16256; | ||
| 41 | +} | ||
| 42 | + | ||
| 43 | +#ifndef CASE_TYPE | ||
| 44 | +#define CASE_TYPE 6 | ||
| 45 | +#endif | ||
| 46 | + | ||
| 47 | +class KernelAdd { | ||
| 48 | +public: | ||
| 49 | + __aicore__ inline KernelAdd() = default; | ||
| 50 | + __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, | ||
| 51 | + uint32_t totalM, uint32_t totalN, | ||
| 52 | + uint32_t splitM, uint32_t splitN, | ||
| 53 | + uint32_t singleCoreM, uint32_t singleCoreN, | ||
| 54 | + uint32_t dataCopyLen, | ||
| 55 | + bool enableL2Bypass = false) | ||
| 56 | + { | ||
| 57 | + uint32_t blockIdx = AscendC::GetBlockIdx(); | ||
| 58 | + uint32_t blockIdxM = blockIdx / splitN; | ||
| 59 | + uint32_t blockIdxN = blockIdx % splitN; | ||
| 60 | + | ||
| 61 | + this->totalM = totalM; | ||
| 62 | + this->totalN = totalN; | ||
| 63 | + this->dataCopyLen = dataCopyLen; | ||
| 64 | + | ||
| 65 | + uint32_t baseCoreM = totalM / splitM; | ||
| 66 | + uint32_t remainderM = totalM % splitM; | ||
| 67 | + uint32_t actualCoreM; | ||
| 68 | + uint32_t startM; | ||
| 69 | + if (blockIdxM < remainderM) { | ||
| 70 | + actualCoreM = baseCoreM + 1; | ||
| 71 | + startM = blockIdxM * actualCoreM; | ||
| 72 | + } else { | ||
| 73 | + actualCoreM = baseCoreM; | ||
| 74 | + startM = remainderM * (baseCoreM + 1) + (blockIdxM - remainderM) * baseCoreM; | ||
| 75 | + } | ||
| 76 | + this->singleCoreM = actualCoreM; | ||
| 77 | + | ||
| 78 | + uint32_t baseCoreN = totalN / splitN; | ||
| 79 | + uint32_t remainderN = totalN % splitN; | ||
| 80 | + uint32_t actualCoreN; | ||
| 81 | + uint32_t startN; | ||
| 82 | + if (blockIdxN < remainderN) { | ||
| 83 | + actualCoreN = baseCoreN + 1; | ||
| 84 | + startN = blockIdxN * actualCoreN; | ||
| 85 | + } else { | ||
| 86 | + actualCoreN = baseCoreN; | ||
| 87 | + startN = remainderN * (baseCoreN + 1) + (blockIdxN - remainderN) * baseCoreN; | ||
| 88 | + } | ||
| 89 | + this->singleCoreN = actualCoreN; | ||
| 90 | + | ||
| 91 | + uint32_t bufferSize = singleCoreM * totalN; | ||
| 92 | + xGm.SetGlobalBuffer((__gm__ half *)x + startM * totalN + startN, bufferSize); | ||
| 93 | + yGm.SetGlobalBuffer((__gm__ half *)y + startM * totalN + startN, bufferSize); | ||
| 94 | + zGm.SetGlobalBuffer((__gm__ half *)z + startM * totalN + startN, bufferSize); | ||
| 95 | + | ||
| 96 | + if (enableL2Bypass) { | ||
| 97 | + xGm.SetL2CacheHint(AscendC::CacheMode::CACHE_MODE_DISABLE); | ||
| 98 | + yGm.SetL2CacheHint(AscendC::CacheMode::CACHE_MODE_DISABLE); | ||
| 99 | + } | ||
| 100 | + } | ||
| 101 | + | ||
| 102 | + __aicore__ inline void Process() | ||
| 103 | + { | ||
| 104 | + AscendC::LocalTensor<half> xLocal(AscendC::TPosition::VECCALC, xAddr, MAX_DATA_COPY_LEN); | ||
| 105 | + AscendC::LocalTensor<half> yLocal(AscendC::TPosition::VECCALC, yAddr, MAX_DATA_COPY_LEN); | ||
| 106 | + AscendC::LocalTensor<half> zLocal(AscendC::TPosition::VECCALC, zAddr, MAX_DATA_COPY_LEN); | ||
| 107 | + | ||
| 108 | + uint32_t rowsPerBlock = dataCopyLen / singleCoreN; | ||
| 109 | + uint32_t fullBlocks = singleCoreM / rowsPerBlock; | ||
| 110 | + uint32_t tailRows = singleCoreM % rowsPerBlock; | ||
| 111 | + uint32_t totalBlocks = fullBlocks + (tailRows > 0 ? 1 : 0); | ||
| 112 | + | ||
| 113 | + AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0); | ||
| 114 | + | ||
| 115 | + for (uint32_t loopIdx = 0; loopIdx < totalBlocks; loopIdx++) { | ||
| 116 | + uint32_t startM = loopIdx * rowsPerBlock; | ||
| 117 | + uint32_t curRows = (loopIdx < fullBlocks) ? rowsPerBlock : tailRows; | ||
| 118 | + uint32_t curLen = curRows * singleCoreN; | ||
| 119 | + | ||
| 120 | + AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0); | ||
| 121 | + | ||
| 122 | + uint32_t blockLen = static_cast<uint32_t>(singleCoreN * sizeof(half)); | ||
| 123 | + uint32_t srcStride = static_cast<uint32_t>((totalN - singleCoreN) * sizeof(half)); | ||
| 124 | + uint32_t dstStride = 0; | ||
| 125 | + AscendC::DataCopyExtParams copyParams = {static_cast<uint16_t>(curRows), blockLen, srcStride, dstStride, 0}; | ||
| 126 | + AscendC::DataCopyPadExtParams<half> padParams = {false, 0, 0, 0}; | ||
| 127 | + AscendC::DataCopyPad<half>(xLocal, xGm[startM * totalN], copyParams, padParams); | ||
| 128 | + AscendC::DataCopyPad<half>(yLocal, yGm[startM * totalN], copyParams, padParams); | ||
| 129 | + | ||
| 130 | + AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0); | ||
| 131 | + AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0); | ||
| 132 | + | ||
| 133 | + AscendC::Add(zLocal, xLocal, yLocal, curLen); | ||
| 134 | + | ||
| 135 | + AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0); | ||
| 136 | + AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0); | ||
| 137 | + | ||
| 138 | + copyParams.srcStride = 0; | ||
| 139 | + copyParams.dstStride = srcStride; | ||
| 140 | + AscendC::DataCopyPad<half>(zGm[startM * totalN], zLocal, copyParams); | ||
| 141 | + | ||
| 142 | + AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0); | ||
| 143 | + } | ||
| 144 | + | ||
| 145 | + AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0); | ||
| 146 | + } | ||
| 147 | + | ||
| 148 | + __aicore__ inline void ProcessScalar() | ||
| 149 | + { | ||
| 150 | + AscendC::LocalTensor<half> xLocal(AscendC::TPosition::VECCALC, xAddr, MAX_DATA_COPY_LEN); | ||
| 151 | + AscendC::LocalTensor<half> yLocal(AscendC::TPosition::VECCALC, yAddr, MAX_DATA_COPY_LEN); | ||
| 152 | + AscendC::LocalTensor<half> zLocal(AscendC::TPosition::VECCALC, zAddr, MAX_DATA_COPY_LEN); | ||
| 153 | + | ||
| 154 | + uint32_t totalLoop = singleCoreM * ((singleCoreN + dataCopyLen - 1) / dataCopyLen); | ||
| 155 | + uint32_t loopIdx = 0; | ||
| 156 | + | ||
| 157 | + for (uint32_t m = 0; m < singleCoreM; m++) { | ||
| 158 | + for (uint32_t n = 0; n < singleCoreN; n += dataCopyLen) { | ||
| 159 | + uint32_t curLen = (singleCoreN - n) > dataCopyLen ? dataCopyLen : (singleCoreN - n); | ||
| 160 | + | ||
| 161 | + if (loopIdx != 0) { | ||
| 162 | + AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0); | ||
| 163 | + } | ||
| 164 | + | ||
| 165 | + AscendC::DataCopy(xLocal, xGm[m * totalN + n], curLen); | ||
| 166 | + AscendC::DataCopy(yLocal, yGm[m * totalN + n], curLen); | ||
| 167 | + AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0); | ||
| 168 | + AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0); | ||
| 169 | + | ||
| 170 | + if (loopIdx != 0) { | ||
| 171 | + AscendC::WaitFlag<AscendC::HardEvent::MTE3_V>(EVENT_ID0); | ||
| 172 | + } | ||
| 173 | + | ||
| 174 | + for (uint32_t i = 0; i < curLen; i++) { | ||
| 175 | + float xVal = (float)xLocal.GetValue(i); | ||
| 176 | + float yVal = (float)yLocal.GetValue(i); | ||
| 177 | + zLocal.SetValue(i, (half)(xVal + yVal)); | ||
| 178 | + } | ||
| 179 | + | ||
| 180 | + if (loopIdx != (totalLoop - 1)) { | ||
| 181 | + AscendC::SetFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0); | ||
| 182 | + } | ||
| 183 | + AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0); | ||
| 184 | + AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0); | ||
| 185 | + | ||
| 186 | + AscendC::DataCopy(zGm[m * totalN + n], zLocal, curLen); | ||
| 187 | + if (loopIdx != (totalLoop - 1)) { | ||
| 188 | + AscendC::SetFlag<AscendC::HardEvent::MTE3_V>(EVENT_ID0); | ||
| 189 | + } | ||
| 190 | + | ||
| 191 | + loopIdx++; | ||
| 192 | + } | ||
| 193 | + } | ||
| 194 | + } | ||
| 195 | + | ||
| 196 | + __aicore__ inline void ProcessSingle() | ||
| 197 | + { | ||
| 198 | + AscendC::LocalTensor<half> xLocal(AscendC::TPosition::VECCALC, xAddr, MAX_DATA_COPY_LEN); | ||
| 199 | + AscendC::LocalTensor<half> yLocal(AscendC::TPosition::VECCALC, yAddr, MAX_DATA_COPY_LEN); | ||
| 200 | + AscendC::LocalTensor<half> zLocal(AscendC::TPosition::VECCALC, zAddr, MAX_DATA_COPY_LEN); | ||
| 201 | + | ||
| 202 | + uint32_t totalLoop = singleCoreM * ((singleCoreN + dataCopyLen - 1) / dataCopyLen); | ||
| 203 | + uint32_t loopIdx = 0; | ||
| 204 | + | ||
| 205 | + for (uint32_t m = 0; m < singleCoreM; m++) { | ||
| 206 | + for (uint32_t n = 0; n < singleCoreN; n += dataCopyLen) { | ||
| 207 | + uint32_t curLen = (singleCoreN - n) > dataCopyLen ? dataCopyLen : (singleCoreN - n); | ||
| 208 | + | ||
| 209 | + if (loopIdx != 0) { | ||
| 210 | + AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0); | ||
| 211 | + } | ||
| 212 | + | ||
| 213 | + AscendC::DataCopy(xLocal, xGm[m * totalN + n], curLen); | ||
| 214 | + AscendC::DataCopy(yLocal, yGm[m * totalN + n], curLen); | ||
| 215 | + AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0); | ||
| 216 | + AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0); | ||
| 217 | + | ||
| 218 | + if (loopIdx != 0) { | ||
| 219 | + AscendC::WaitFlag<AscendC::HardEvent::MTE3_V>(EVENT_ID0); | ||
| 220 | + } | ||
| 221 | + | ||
| 222 | + AscendC::Add(zLocal, xLocal, yLocal, curLen); | ||
| 223 | + | ||
| 224 | + if (loopIdx != (totalLoop - 1)) { | ||
| 225 | + AscendC::SetFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0); | ||
| 226 | + } | ||
| 227 | + AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0); | ||
| 228 | + AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0); | ||
| 229 | + | ||
| 230 | + AscendC::DataCopy(zGm[m * totalN + n], zLocal, curLen); | ||
| 231 | + if (loopIdx != (totalLoop - 1)) { | ||
| 232 | + AscendC::SetFlag<AscendC::HardEvent::MTE3_V>(EVENT_ID0); | ||
| 233 | + } | ||
| 234 | + | ||
| 235 | + loopIdx++; | ||
| 236 | + } | ||
| 237 | + } | ||
| 238 | + } | ||
| 239 | + | ||
| 240 | + __aicore__ inline void ProcessDoubleBuffer() | ||
| 241 | + { | ||
| 242 | + AscendC::LocalTensor<half> xPing(AscendC::TPosition::VECCALC, xAddrPing, MAX_DATA_COPY_LEN); | ||
| 243 | + AscendC::LocalTensor<half> yPing(AscendC::TPosition::VECCALC, yAddrPing, MAX_DATA_COPY_LEN); | ||
| 244 | + AscendC::LocalTensor<half> zPing(AscendC::TPosition::VECCALC, zAddrPing, MAX_DATA_COPY_LEN); | ||
| 245 | + AscendC::LocalTensor<half> xPong(AscendC::TPosition::VECCALC, xAddrPong, MAX_DATA_COPY_LEN); | ||
| 246 | + AscendC::LocalTensor<half> yPong(AscendC::TPosition::VECCALC, yAddrPong, MAX_DATA_COPY_LEN); | ||
| 247 | + AscendC::LocalTensor<half> zPong(AscendC::TPosition::VECCALC, zAddrPong, MAX_DATA_COPY_LEN); | ||
| 248 | + | ||
| 249 | + uint32_t rowsPerBlock = dataCopyLen / singleCoreN; | ||
| 250 | + uint32_t fullBlocks = singleCoreM / rowsPerBlock; | ||
| 251 | + uint32_t tailRows = singleCoreM % rowsPerBlock; | ||
| 252 | + uint32_t totalBlocks = fullBlocks + (tailRows > 0 ? 1 : 0); | ||
| 253 | + | ||
| 254 | + AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0); | ||
| 255 | + AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID1); | ||
| 256 | + | ||
| 257 | + for (uint32_t loopIdx = 0; loopIdx < totalBlocks; loopIdx++) { | ||
| 258 | + uint32_t startM = loopIdx * rowsPerBlock; | ||
| 259 | + uint32_t curRows = (loopIdx < fullBlocks) ? rowsPerBlock : tailRows; | ||
| 260 | + uint32_t curLen = curRows * singleCoreN; | ||
| 261 | + | ||
| 262 | + int32_t eventID = (loopIdx % 2 == 0 ? EVENT_ID0 : EVENT_ID1); | ||
| 263 | + AscendC::LocalTensor<half> &xLocal = (loopIdx % 2 == 0 ? xPing : xPong); | ||
| 264 | + AscendC::LocalTensor<half> &yLocal = (loopIdx % 2 == 0 ? yPing : yPong); | ||
| 265 | + AscendC::LocalTensor<half> &zLocal = (loopIdx % 2 == 0 ? zPing : zPong); | ||
| 266 | + | ||
| 267 | + AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(eventID); | ||
| 268 | + | ||
| 269 | + uint32_t blockLen = static_cast<uint32_t>(singleCoreN * sizeof(half)); | ||
| 270 | + uint32_t srcStride = static_cast<uint32_t>((totalN - singleCoreN) * sizeof(half)); | ||
| 271 | + uint32_t dstStride = 0; | ||
| 272 | + AscendC::DataCopyExtParams copyParams = {static_cast<uint16_t>(curRows), blockLen, srcStride, dstStride, 0}; | ||
| 273 | + AscendC::DataCopyPadExtParams<half> padParams = {false, 0, 0, 0}; | ||
| 274 | + AscendC::DataCopyPad<half>(xLocal, xGm[startM * totalN], copyParams, padParams); | ||
| 275 | + AscendC::DataCopyPad<half>(yLocal, yGm[startM * totalN], copyParams, padParams); | ||
| 276 | + | ||
| 277 | + AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(eventID); | ||
| 278 | + AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(eventID); | ||
| 279 | + | ||
| 280 | + AscendC::Add(zLocal, xLocal, yLocal, curLen); | ||
| 281 | + | ||
| 282 | + AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(eventID); | ||
| 283 | + AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(eventID); | ||
| 284 | + | ||
| 285 | + copyParams.srcStride = 0; | ||
| 286 | + copyParams.dstStride = srcStride; | ||
| 287 | + AscendC::DataCopyPad<half>(zGm[startM * totalN], zLocal, copyParams); | ||
| 288 | + | ||
| 289 | + AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(eventID); | ||
| 290 | + } | ||
| 291 | + | ||
| 292 | + AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0); | ||
| 293 | + AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID1); | ||
| 294 | + } | ||
| 295 | + | ||
| 296 | + __aicore__ inline void ProcessDoubleBufferL2Bypass() | ||
| 297 | + { | ||
| 298 | + ProcessDoubleBuffer(); | ||
| 299 | + } | ||
| 300 | + | ||
| 301 | + __aicore__ inline void ProcessDoubleBufferBankConflict() | ||
| 302 | + { | ||
| 303 | + xGm.SetL2CacheHint(AscendC::CacheMode::CACHE_MODE_DISABLE); | ||
| 304 | + yGm.SetL2CacheHint(AscendC::CacheMode::CACHE_MODE_DISABLE); | ||
| 305 | + | ||
| 306 | + AscendC::LocalTensor<half> xPing(AscendC::TPosition::VECCALC, xAddrPingBC, MAX_DATA_COPY_LEN); | ||
| 307 | + AscendC::LocalTensor<half> yPing(AscendC::TPosition::VECCALC, yAddrPingBC, MAX_DATA_COPY_LEN); | ||
| 308 | + AscendC::LocalTensor<half> zPing(AscendC::TPosition::VECCALC, zAddrPingBC, MAX_DATA_COPY_LEN); | ||
| 309 | + AscendC::LocalTensor<half> xPong(AscendC::TPosition::VECCALC, xAddrPongBC, MAX_DATA_COPY_LEN); | ||
| 310 | + AscendC::LocalTensor<half> yPong(AscendC::TPosition::VECCALC, yAddrPongBC, MAX_DATA_COPY_LEN); | ||
| 311 | + AscendC::LocalTensor<half> zPong(AscendC::TPosition::VECCALC, zAddrPongBC, MAX_DATA_COPY_LEN); | ||
| 312 | + | ||
| 313 | + uint32_t rowsPerBlock = dataCopyLen / singleCoreN; | ||
| 314 | + uint32_t fullBlocks = singleCoreM / rowsPerBlock; | ||
| 315 | + uint32_t tailRows = singleCoreM % rowsPerBlock; | ||
| 316 | + uint32_t totalBlocks = fullBlocks + (tailRows > 0 ? 1 : 0); | ||
| 317 | + | ||
| 318 | + AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0); | ||
| 319 | + AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID1); | ||
| 320 | + | ||
| 321 | + for (uint32_t loopIdx = 0; loopIdx < totalBlocks; loopIdx++) { | ||
| 322 | + uint32_t startM = loopIdx * rowsPerBlock; | ||
| 323 | + uint32_t curRows = (loopIdx < fullBlocks) ? rowsPerBlock : tailRows; | ||
| 324 | + uint32_t curLen = curRows * singleCoreN; | ||
| 325 | + | ||
| 326 | + int32_t eventID = (loopIdx % 2 == 0 ? EVENT_ID0 : EVENT_ID1); | ||
| 327 | + AscendC::LocalTensor<half> &xLocal = (loopIdx % 2 == 0 ? xPing : xPong); | ||
| 328 | + AscendC::LocalTensor<half> &yLocal = (loopIdx % 2 == 0 ? yPing : yPong); | ||
| 329 | + AscendC::LocalTensor<half> &zLocal = (loopIdx % 2 == 0 ? zPing : zPong); | ||
| 330 | + | ||
| 331 | + AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(eventID); | ||
| 332 | + | ||
| 333 | + if (curRows == rowsPerBlock) { | ||
| 334 | + uint32_t blockLen = static_cast<uint32_t>(singleCoreN * sizeof(half)); | ||
| 335 | + uint32_t srcStride = static_cast<uint32_t>((totalN - singleCoreN) * sizeof(half)); | ||
| 336 | + uint32_t dstStride = 0; | ||
| 337 | + AscendC::DataCopyExtParams copyParams = {static_cast<uint16_t>(curRows), blockLen, srcStride, dstStride, 0}; | ||
| 338 | + AscendC::DataCopyPadExtParams<half> padParams = {false, 0, 0, 0}; | ||
| 339 | + AscendC::DataCopyPad<half>(xLocal, xGm[startM * totalN], copyParams, padParams); | ||
| 340 | + AscendC::DataCopyPad<half>(yLocal, yGm[startM * totalN], copyParams, padParams); | ||
| 341 | + } else { | ||
| 342 | + for (uint32_t m = 0; m < curRows; m++) { | ||
| 343 | + AscendC::DataCopy(xLocal[m * singleCoreN], xGm[(startM + m) * totalN], singleCoreN); | ||
| 344 | + AscendC::DataCopy(yLocal[m * singleCoreN], yGm[(startM + m) * totalN], singleCoreN); | ||
| 345 | + } | ||
| 346 | + } | ||
| 347 | + | ||
| 348 | + AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(eventID); | ||
| 349 | + AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(eventID); | ||
| 350 | + | ||
| 351 | + AscendC::Add(zLocal, xLocal, yLocal, curLen); | ||
| 352 | + | ||
| 353 | + AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(eventID); | ||
| 354 | + AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(eventID); | ||
| 355 | + | ||
| 356 | + if (curRows == rowsPerBlock) { | ||
| 357 | + uint32_t blockLen = static_cast<uint32_t>(singleCoreN * sizeof(half)); | ||
| 358 | + uint32_t srcStride = 0; | ||
| 359 | + uint32_t dstStride = static_cast<uint32_t>((totalN - singleCoreN) * sizeof(half)); | ||
| 360 | + AscendC::DataCopyExtParams copyParams = {static_cast<uint16_t>(curRows), blockLen, srcStride, dstStride, 0}; | ||
| 361 | + AscendC::DataCopyPad<half>(zGm[startM * totalN], zLocal, copyParams); | ||
| 362 | + } else { | ||
| 363 | + for (uint32_t m = 0; m < curRows; m++) { | ||
| 364 | + AscendC::DataCopy(zGm[(startM + m) * totalN], zLocal[m * singleCoreN], singleCoreN); | ||
| 365 | + } | ||
| 366 | + } | ||
| 367 | + | ||
| 368 | + AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(eventID); | ||
| 369 | + } | ||
| 370 | + | ||
| 371 | + AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0); | ||
| 372 | + AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID1); | ||
| 373 | + } | ||
| 374 | + | ||
| 375 | +private: | ||
| 376 | + static constexpr uint32_t xAddr = 0; | ||
| 377 | + static constexpr uint32_t yAddr = MAX_DATA_COPY_LEN * sizeof(half); | ||
| 378 | + static constexpr uint32_t zAddr = yAddr + MAX_DATA_COPY_LEN * sizeof(half); | ||
| 379 | + | ||
| 380 | + static constexpr uint32_t xAddrPing = 0; | ||
| 381 | + static constexpr uint32_t yAddrPing = MAX_DATA_COPY_LEN * sizeof(half); | ||
| 382 | + static constexpr uint32_t zAddrPing = yAddrPing + MAX_DATA_COPY_LEN * sizeof(half); | ||
| 383 | + static constexpr uint32_t xAddrPong = zAddrPing + MAX_DATA_COPY_LEN * sizeof(half); | ||
| 384 | + static constexpr uint32_t yAddrPong = xAddrPong + MAX_DATA_COPY_LEN * sizeof(half); | ||
| 385 | + static constexpr uint32_t zAddrPong = yAddrPong + MAX_DATA_COPY_LEN * sizeof(half); | ||
| 386 | + | ||
| 387 | + static constexpr uint32_t xAddrPingBC = 0; | ||
| 388 | + static constexpr uint32_t yAddrPingBC = BANK_CONFLICT_DATA_COPY_LEN * sizeof(half); | ||
| 389 | + static constexpr uint32_t xAddrPongBC = MAX_DATA_COPY_LEN * sizeof(half) * 2; | ||
| 390 | + static constexpr uint32_t yAddrPongBC = xAddrPongBC + BANK_CONFLICT_DATA_COPY_LEN * sizeof(half); | ||
| 391 | + static constexpr uint32_t zAddrPingBC = MAX_DATA_COPY_LEN * sizeof(half) * 4; | ||
| 392 | + static constexpr uint32_t zAddrPongBC = zAddrPingBC + BANK_CONFLICT_DATA_COPY_LEN * sizeof(half); | ||
| 393 | + | ||
| 394 | + AscendC::GlobalTensor<half> xGm; | ||
| 395 | + AscendC::GlobalTensor<half> yGm; | ||
| 396 | + AscendC::GlobalTensor<half> zGm; | ||
| 397 | + uint32_t totalM; | ||
| 398 | + uint32_t totalN; | ||
| 399 | + uint32_t singleCoreM; | ||
| 400 | + uint32_t singleCoreN; | ||
| 401 | + uint32_t dataCopyLen; | ||
| 402 | +}; | ||
| 403 | + | ||
| 404 | +extern "C" __vector__ __global__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR tiling) | ||
| 405 | +{ | ||
| 406 | + AscendC::InitSocState(); | ||
| 407 | + KernelAdd op; | ||
| 408 | + auto tilingData = (__gm__ AddCustomTilingData *)tiling; | ||
| 409 | + | ||
| 410 | + if constexpr (CASE_TYPE == 5) { | ||
| 411 | + op.Init(x, y, z, | ||
| 412 | + tilingData->totalM, tilingData->totalN, | ||
| 413 | + tilingData->splitM, tilingData->splitN, | ||
| 414 | + tilingData->singleCoreM, tilingData->singleCoreN, | ||
| 415 | + tilingData->dataCopyLen, | ||
| 416 | + true); | ||
| 417 | + op.ProcessDoubleBufferL2Bypass(); | ||
| 418 | + } else if constexpr (CASE_TYPE == 6) { | ||
| 419 | + op.Init(x, y, z, | ||
| 420 | + tilingData->totalM, tilingData->totalN, | ||
| 421 | + tilingData->splitM, tilingData->splitN, | ||
| 422 | + tilingData->singleCoreM, tilingData->singleCoreN, | ||
| 423 | + tilingData->dataCopyLen, | ||
| 424 | + false); | ||
| 425 | + op.ProcessDoubleBufferBankConflict(); | ||
| 426 | + } else { | ||
| 427 | + op.Init(x, y, z, | ||
| 428 | + tilingData->totalM, tilingData->totalN, | ||
| 429 | + tilingData->splitM, tilingData->splitN, | ||
| 430 | + tilingData->singleCoreM, tilingData->singleCoreN, | ||
| 431 | + tilingData->dataCopyLen, | ||
| 432 | + false); | ||
| 433 | + | ||
| 434 | + if constexpr (CASE_TYPE == 0) { | ||
| 435 | + op.ProcessScalar(); | ||
| 436 | + } else if constexpr (CASE_TYPE == 1) { | ||
| 437 | + op.ProcessSingle(); | ||
| 438 | + } else if constexpr (CASE_TYPE == 4) { | ||
| 439 | + op.ProcessDoubleBuffer(); | ||
| 440 | + } else { | ||
| 441 | + op.Process(); | ||
| 442 | + } | ||
| 443 | + } | ||
| 444 | +} | ||
| 445 | + | ||
| 446 | +struct ArgInfo { | ||
| 447 | + std::string fileName; | ||
| 448 | + size_t length; | ||
| 449 | +}; | ||
| 450 | + | ||
| 451 | +int32_t main(int32_t argc, char* argv[]) | ||
| 452 | +{ | ||
| 453 | + printf("Usage: %s (compile-time CASE_TYPE=%u)\n", argv[0], (uint32_t)CASE_TYPE); | ||
| 454 | + printf("Running case %u...\n", (uint32_t)CASE_TYPE); | ||
| 455 | + | ||
| 456 | + uint32_t totalM = 8192; | ||
| 457 | + uint32_t totalN = 8192; | ||
| 458 | + uint32_t splitM; | ||
| 459 | + uint32_t splitN; | ||
| 460 | + uint32_t dataCopyLen; | ||
| 461 | + | ||
| 462 | + if constexpr (CASE_TYPE == 0) { | ||
| 463 | + splitM = 1; | ||
| 464 | + splitN = 1; | ||
| 465 | + dataCopyLen = 4096; | ||
| 466 | + } else if constexpr (CASE_TYPE == 1) { | ||
| 467 | + splitM = 1; | ||
| 468 | + splitN = 1; | ||
| 469 | + dataCopyLen = 4096; | ||
| 470 | + } else if constexpr (CASE_TYPE == 2) { | ||
| 471 | + splitM = 6; | ||
| 472 | + splitN = 8; | ||
| 473 | + dataCopyLen = 4096; | ||
| 474 | + } else if constexpr (CASE_TYPE == 3) { | ||
| 475 | + splitM = 6; | ||
| 476 | + splitN = 8; | ||
| 477 | + dataCopyLen = 16384; | ||
| 478 | + } else if constexpr (CASE_TYPE == 4) { | ||
| 479 | + splitM = 6; | ||
| 480 | + splitN = 8; | ||
| 481 | + dataCopyLen = 16384; | ||
| 482 | + } else if constexpr (CASE_TYPE == 5) { | ||
| 483 | + splitM = 6; | ||
| 484 | + splitN = 8; | ||
| 485 | + dataCopyLen = 16384; | ||
| 486 | + } else if constexpr (CASE_TYPE == 6) { | ||
| 487 | + splitM = 6; | ||
| 488 | + splitN = 8; | ||
| 489 | + dataCopyLen = 16256; | ||
| 490 | + } else { | ||
| 491 | + splitM = 1; splitN = 1; dataCopyLen = 4096; | ||
| 492 | + } | ||
| 493 | + | ||
| 494 | + uint32_t numBlocks = splitM * splitN; | ||
| 495 | + | ||
| 496 | + uint32_t singleCoreM = totalM / splitM; | ||
| 497 | + uint32_t singleCoreN = totalN / splitN; | ||
| 498 | + uint32_t dataLen = totalM * totalN; | ||
| 499 | + | ||
| 500 | + size_t inputByteSize = dataLen * sizeof(uint16_t); | ||
| 501 | + size_t outputByteSize = dataLen * sizeof(uint16_t); | ||
| 502 | + AddCustomTilingData tiling; | ||
| 503 | + tiling.caseType = CASE_TYPE; | ||
| 504 | + tiling.totalM = totalM; | ||
| 505 | + tiling.totalN = totalN; | ||
| 506 | + tiling.splitM = splitM; | ||
| 507 | + tiling.splitN = splitN; | ||
| 508 | + tiling.singleCoreM = singleCoreM; | ||
| 509 | + tiling.singleCoreN = singleCoreN; | ||
| 510 | + tiling.dataCopyLen = dataCopyLen; | ||
| 511 | + | ||
| 512 | + std::vector<ArgInfo> inputsInfo = {{"./input/input_x.bin", inputByteSize}, {"./input/input_y.bin", inputByteSize}}; | ||
| 513 | + std::vector<ArgInfo> outputsInfo = {{"./output/output.bin", outputByteSize}}; | ||
| 514 | + | ||
| 515 | + aclInit(nullptr); | ||
| 516 | + int32_t deviceId = 0; | ||
| 517 | + aclrtSetDevice(deviceId); | ||
| 518 | + aclrtStream stream = nullptr; | ||
| 519 | + aclrtCreateStream(&stream); | ||
| 520 | + | ||
| 521 | + std::vector<uint8_t *> inputHost(inputsInfo.size()); | ||
| 522 | + std::vector<uint8_t *> inputDevice(inputsInfo.size()); | ||
| 523 | + std::vector<uint8_t *> outputHost(outputsInfo.size()); | ||
| 524 | + std::vector<uint8_t *> outputDevice(outputsInfo.size()); | ||
| 525 | + uint8_t *tilingDevice; | ||
| 526 | + | ||
| 527 | + aclrtMalloc((void **)(&tilingDevice), sizeof(AddCustomTilingData), ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 528 | + aclrtMemcpy(tilingDevice, sizeof(AddCustomTilingData), (uint8_t *)&tiling, sizeof(AddCustomTilingData), ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 529 | + | ||
| 530 | + for (uint32_t i = 0; i < inputsInfo.size(); i++) { | ||
| 531 | + aclrtMallocHost((void **)(&inputHost[i]), inputsInfo[i].length); | ||
| 532 | + aclrtMalloc((void **)(&inputDevice[i]), inputsInfo[i].length, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 533 | + ReadFile(inputsInfo[i].fileName, inputsInfo[i].length, inputHost[i], inputsInfo[i].length); | ||
| 534 | + aclrtMemcpy(inputDevice[i], inputsInfo[i].length, inputHost[i], inputsInfo[i].length, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 535 | + } | ||
| 536 | + | ||
| 537 | + for (uint32_t i = 0; i < outputsInfo.size(); i++) { | ||
| 538 | + aclrtMallocHost((void **)(&outputHost[i]), outputsInfo[i].length); | ||
| 539 | + aclrtMalloc((void **)(&outputDevice[i]), outputsInfo[i].length, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 540 | + } | ||
| 541 | + | ||
| 542 | + add_custom<<<numBlocks, nullptr, stream>>>(inputDevice[0], inputDevice[1], outputDevice[0], tilingDevice); | ||
| 543 | + aclrtSynchronizeStream(stream); | ||
| 544 | + | ||
| 545 | + aclrtFree(tilingDevice); | ||
| 546 | + for (uint32_t i = 0; i < outputsInfo.size(); i++) { | ||
| 547 | + aclrtMemcpy(outputHost[i], outputsInfo[i].length, outputDevice[i], outputsInfo[i].length, ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 548 | + WriteFile(outputsInfo[i].fileName, outputHost[i], outputsInfo[i].length); | ||
| 549 | + aclrtFree(outputDevice[i]); | ||
| 550 | + aclrtFreeHost(outputHost[i]); | ||
| 551 | + } | ||
| 552 | + | ||
| 553 | + for (uint32_t i = 0; i < inputsInfo.size(); i++) { | ||
| 554 | + aclrtFree(inputDevice[i]); | ||
| 555 | + aclrtFreeHost(inputHost[i]); | ||
| 556 | + } | ||
| 557 | + | ||
| 558 | + aclrtDestroyStream(stream); | ||
| 559 | + aclrtResetDevice(deviceId); | ||
| 560 | + aclFinalize(); | ||
| 561 | + | ||
| 562 | + printf("Case %u completed.\n", (uint32_t)CASE_TYPE); | ||
| 563 | + return 0; | ||
| 564 | +} | ||
| @@ -0,0 +1,93 @@ | |||
| 1 | +/** | ||
| 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 | ||
| 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. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + */ | ||
| 10 | + | ||
| 11 | +/* ! | ||
| 12 | + * \file data_utils.h | ||
| 13 | + * \brief Utility functions for reading and writing binary files | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | + | ||
| 25 | +bool ReadFile(const std::string &filePath, size_t &fileSize, void *buffer, size_t bufferSize) | ||
| 26 | +{ | ||
| 27 | + struct stat sBuf; | ||
| 28 | + int fileStatus = stat(filePath.data(), &sBuf); | ||
| 29 | + if (fileStatus == -1) { | ||
| 30 | + ERROR_LOG("failed to get file"); | ||
| 31 | + return false; | ||
| 32 | + } | ||
| 33 | + if (S_ISREG(sBuf.st_mode) == 0) { | ||
| 34 | + ERROR_LOG("%s is not a file, please enter a file", filePath.c_str()); | ||
| 35 | + return false; | ||
| 36 | + } | ||
| 37 | + | ||
| 38 | + std::ifstream file; | ||
| 39 | + file.open(filePath, std::ios::binary); | ||
| 40 | + if (!file.is_open()) { | ||
| 41 | + ERROR_LOG("Open file failed. path = %s", filePath.c_str()); | ||
| 42 | + return false; | ||
| 43 | + } | ||
| 44 | + | ||
| 45 | + std::filebuf *buf = file.rdbuf(); | ||
| 46 | + size_t size = buf->pubseekoff(0, std::ios::end, std::ios::in); | ||
| 47 | + if (size == 0) { | ||
| 48 | + ERROR_LOG("file size is 0"); | ||
| 49 | + file.close(); | ||
| 50 | + return false; | ||
| 51 | + } | ||
| 52 | + if (size > bufferSize) { | ||
| 53 | + ERROR_LOG("file size is larger than buffer size"); | ||
| 54 | + file.close(); | ||
| 55 | + return false; | ||
| 56 | + } | ||
| 57 | + buf->pubseekpos(0, std::ios::in); | ||
| 58 | + buf->sgetn(static_cast<char *>(buffer), size); | ||
| 59 | + fileSize = size; | ||
| 60 | + file.close(); | ||
| 61 | + return true; | ||
| 62 | +} | ||
| 63 | + | ||
| 64 | +/** | ||
| 65 | + * @brief Write data to file | ||
| 66 | + * @param [in] filePath: file path | ||
| 67 | + * @param [in] buffer: data to write to file | ||
| 68 | + * @param [in] size: size to write | ||
| 69 | + * @return write result | ||
| 70 | + */ | ||
| 71 | +bool WriteFile(const std::string &filePath, const void *buffer, size_t size) | ||
| 72 | +{ | ||
| 73 | + if (buffer == nullptr) { | ||
| 74 | + ERROR_LOG("Write file failed. buffer is nullptr"); | ||
| 75 | + return false; | ||
| 76 | + } | ||
| 77 | + | ||
| 78 | + int fd = open(filePath.c_str(), O_RDWR | O_CREAT | O_TRUNC, S_IRUSR | S_IWRITE); | ||
| 79 | + if (fd < 0) { | ||
| 80 | + ERROR_LOG("Open file failed. path = %s", filePath.c_str()); | ||
| 81 | + return false; | ||
| 82 | + } | ||
| 83 | + | ||
| 84 | + size_t writeSize = write(fd, buffer, size); | ||
| 85 | + (void)close(fd); | ||
| 86 | + if (writeSize != size) { | ||
| 87 | + ERROR_LOG("Write file Failed."); | ||
| 88 | + return false; | ||
| 89 | + } | ||
| 90 | + | ||
| 91 | + return true; | ||
| 92 | +} | ||
| 93 | + | ||
| @@ -0,0 +1,3 @@ | |||
| 1 | +version https://git-lfs.github.com/spec/v1 | ||
| 2 | +oid sha256:2b11e480a482a3f7e8eabd62a7963d56412031be98a2c9216943e8ea1d5f9692 | ||
| 3 | +size 1936805 | ||
| @@ -0,0 +1,3 @@ | |||
| 1 | +version https://git-lfs.github.com/spec/v1 | ||
| 2 | +oid sha256:e26e8551c39d0724019b6d3a42ae1da538af3360da715607cfc9e55885f22ce1 | ||
| 3 | +size 2000067 | ||
| @@ -0,0 +1,31 @@ | |||
| 1 | +#!/usr/bin/python3 | ||
| 2 | +# coding=utf-8 | ||
| 3 | + | ||
| 4 | +# ---------------------------------------------------------------------------------------------------------- | ||
| 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 | ||
| 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. | ||
| 9 | +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 10 | +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 11 | +# See LICENSE in the root of the software repository for the full text of the License. | ||
| 12 | +# ---------------------------------------------------------------------------------------------------------- | ||
| 13 | + | ||
| 14 | + | ||
| 15 | +import os | ||
| 16 | +import numpy as np | ||
| 17 | + | ||
| 18 | + | ||
| 19 | +def gen_golden_data_simple(): | ||
| 20 | + input_x = np.random.uniform(1, 100, [8192, 8192]).astype(np.float16) | ||
| 21 | + input_y = np.random.uniform(1, 100, [8192, 8192]).astype(np.float16) | ||
| 22 | + golden = (input_x + input_y).astype(np.float16) | ||
| 23 | + os.makedirs("input", exist_ok=True) | ||
| 24 | + os.makedirs("output", exist_ok=True) | ||
| 25 | + input_x.tofile("./input/input_x.bin") | ||
| 26 | + input_y.tofile("./input/input_y.bin") | ||
| 27 | + golden.tofile("./output/golden.bin") | ||
| 28 | + | ||
| 29 | + | ||
| 30 | +if __name__ == "__main__": | ||
| 31 | + gen_golden_data_simple() | ||
| @@ -0,0 +1,57 @@ | |||
| 1 | +#!/usr/bin/python3 | ||
| 2 | +# coding=utf-8 | ||
| 3 | + | ||
| 4 | +# ---------------------------------------------------------------------------------------------------------- | ||
| 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 | ||
| 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. | ||
| 9 | +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 10 | +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 11 | +# See LICENSE in the root of the software repository for the full text of the License. | ||
| 12 | +# ---------------------------------------------------------------------------------------------------------- | ||
| 13 | + | ||
| 14 | + | ||
| 15 | +import sys | ||
| 16 | +import numpy as np | ||
| 17 | + | ||
| 18 | + | ||
| 19 | +RELATIVE_TOL = 1e-4 | ||
| 20 | +ABSOLUTE_TOL = 1e-5 | ||
| 21 | +ERROR_TOL = 1e-4 | ||
| 22 | + | ||
| 23 | + | ||
| 24 | +def verify_result(output, golden): | ||
| 25 | + output = np.fromfile(output, dtype=np.float16).reshape(-1) | ||
| 26 | + golden = np.fromfile(golden, dtype=np.float16).reshape(-1) | ||
| 27 | + different_element_results = np.isclose(output, | ||
| 28 | + golden, | ||
| 29 | + rtol=RELATIVE_TOL, | ||
| 30 | + atol=ABSOLUTE_TOL, | ||
| 31 | + equal_nan=True) | ||
| 32 | + different_element_indexes = np.where(different_element_results == False)[0] | ||
| 33 | + for index in range(len(different_element_indexes)): | ||
| 34 | + real_index = different_element_indexes[index] | ||
| 35 | + golden_data = golden[real_index] | ||
| 36 | + output_data = output[real_index] | ||
| 37 | + print( | ||
| 38 | + "data index: %06d, expected: %-.9f, actual: %-.9f, rdiff: %-.6f" % | ||
| 39 | + (real_index, golden_data, output_data, | ||
| 40 | + abs(output_data - golden_data) / golden_data)) | ||
| 41 | + if index == 100: | ||
| 42 | + break | ||
| 43 | + error_ratio = float(different_element_indexes.size) / golden.size | ||
| 44 | + print("error ratio: %.4f, tolerance: %.4f" % (error_ratio, ERROR_TOL)) | ||
| 45 | + return error_ratio <= ERROR_TOL | ||
| 46 | + | ||
| 47 | + | ||
| 48 | +if __name__ == '__main__': | ||
| 49 | + try: | ||
| 50 | + res = verify_result(sys.argv[1], sys.argv[2]) | ||
| 51 | + if not res: | ||
| 52 | + raise ValueError("[ERROR] result error") | ||
| 53 | + else: | ||
| 54 | + print("test pass!") | ||
| 55 | + except Exception as e: | ||
| 56 | + print(e) | ||
| 57 | + sys.exit(1) | ||
样例目录名去掉“00”编号