Ascend C 矢量计算多核 Tiling 切分策略实战：以 vector_tiling_strategy 样例解析核间与核内数据分配

人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载导读本文以 asc-devkit 仓库中的 vector_tiling_strategy 样例 为主线系统讲解昇腾 AI 处理器上矢量计算的多核 Tiling数据切分策略。该样例以两个一维 Tensor 相加z x y为载体演示如何根据输入数据总量与 AI Core 数量的关系逐级计算核间不同核之间的数据分配与核内单个核内主块/尾块的循环粒度切分参数使数据在多个 AI Core 之间尽量均衡分配。读完本文你将掌握主块均分、尾块均分、尾核切分、尾核尾块切分四种典型 Tiling 场景的判定条件与计算方法并能独立编译、运行、验证该样例将其思路迁移到自己的算子开发中。样例概述从数据量推导切分参数本样例固定使用8 个 AI Core参与计算核内主块数据量为3200 个half元素即 200 个 DataBlock。所谓主块mainTileLength是用户基于 Unified BufferUB可用空间确定的一次可处理的最大数据量并满足 32 字节对齐样例中一个 DataBlock 为 32 字节一个half元素占 2 字节因此一个 DataBlock 包含 16 个half元素对应源码常量alignNum dataBlockSize / halfSize 16。样例通过 CMake 编译参数SCENARIO_NUM选择不同数据量场景覆盖四种典型切分形态场景编号场景名称数据总量half元素切分特征0主块均分2560008 核数据量完全相同核内只有主块1尾块均分2600968 核数据量完全相同每核主块后追加等长尾块2尾核切分256064前 4 个整核多处理 1 个 DataBlock尾核无尾块3尾核尾块切分258112所有核都有尾块前 4 个整核再多处理 1 个 DataBlock这些场景编号与数据量之间的对应关系在数据生成脚本 scripts/gen_data.py 中通过字典SCENARIO_LENGTHS {0: 256000, 1: 260096, 2: 256064, 3: 258112}显式维护并在 vector_tiling_strategy.asc 的main()中按scenarioNum三元表达式选择对应dataLen。支持的产品与 CANN 软件版本产品CANN 软件版本Ascend 950PR 950DT 系列产品 CANN 9.1.0Atlas A3 系列产品 CANN 9.0.0Atlas A2 系列产品 CANN 9.0.0从编译选项CMAKE_ASC_ARCHITECTURES的角度看dav-2201架构对应 Atlas A2 / A3 系列产品CANN 9.0.0dav-3510架构对应 Ascend 950PR 950DT 系列产品CANN 9.1.0。目录结构介绍vector_tiling_strategy ├── scripts │ ├── gen_data.py // 输入数据和真值数据生成脚本 │ └── verify_result.py // 输出结果验证脚本 ├── CMakeLists.txt // 编译工程文件 ├── vector_tiling_strategy.asc // Ascend C 样例实现与调用样例 ├── data_utils.h // 数据读入写出函数 └── README.md // 样例说明文档各文件的职责如下vector_tiling_strategy.asc核心实现包含 Host 侧 Tiling 参数计算LaunchAddByTiling、四个 Tiling 数据结构、设备侧核函数add_custom与算子逻辑类KernelAdd以及基于 ACLAscendCL的主机端数据搬运与 kernel 启动代码CMakeLists.txt编译工程文件定义CMAKE_ASC_RUN_MODE、CMAKE_ASC_ARCHITECTURES、SCENARIO_NUM三个可配置项并通过target_compile_definitions将SCENARIO_NUM注入源码编译data_utils.h提供ReadFile/WriteFile二进制文件读写函数供主程序读取输入input_x.bin、input_y.bin并写出计算结果output.binscripts/gen_data.py按场景生成随机输入数据与 golden 真值数据scripts/verify_result.py对输出与真值做精度比对。四种 Tiling 切分场景详解所有场景的数据格式均为 ND输入输出均为half类型核函数名为add_custom计算逻辑统一为z x y。8 个核被分为两类整核former core处理更多数据尾核tail core处理基础数据量。场景0主块均分输入x、y均为 [1, 256000] 个half元素输出z为 [1, 256000] 个half元素Tiling 切分8 个核各处理 32000 个元素每个核包含 10 个主块每个主块 3200 个元素说明所有核处理数据量相同核内只包含主块推导验证256000 8 × 10 × 3200总数据量恰好整除numBlocks × mainTileLength因此无需任何尾块或尾核对应源码中的MainTileOnlyTiling分支。场景1尾块均分输入x、y均为 [1, 260096] 个half元素输出z为 [1, 260096] 个half元素Tiling 切分8 个核各处理 32512 个元素每个核包含 10 个主块和 512 个元素的尾块说明所有核处理数据量相同每个核在主块后都有等长尾块推导验证260096 - 256000 40964096 / (8 × 16) 32即每个核可均分到 32 个 DataBlock 尾块32 × 16 512个元素无剩余对应MainTileWithTailBlockTiling分支。场景2尾核切分输入x、y均为 [1, 256064] 个half元素输出z为 [1, 256064] 个half元素Tiling 切分前 4 个整核各处理 32016 个元素后 4 个尾核各处理 32000 个元素整核包含 10 个主块和 16 个元素的尾块尾核包含 10 个主块说明前 4 个整核比尾核多处理 1 个 DataBlock尾核无尾块推导验证256064 - 256000 64个剩余元素64 / 16 4个 DataBlock恰好分给前 4 个核formerNum 4每个整核多处理 1 个 DataBlock16 个元素。此时tailBlockNumEachCore 0对应MainTileWithTailCoreTiling分支。场景3尾核尾块切分输入x、y均为 [1, 258112] 个half元素输出z为 [1, 258112] 个half元素Tiling 切分前 4 个整核各处理 32272 个元素后 4 个尾核各处理 32256 个元素整核包含 10 个主块和 272 个元素的尾块尾核包含 10 个主块和 256 个元素的尾块说明所有核都有尾块前 4 个整核再多处理 1 个 DataBlock推导验证258112 - 256000 21122112 / (8 × 16) 16余64。即每核先均分 16 个 DataBlock 尾块256 个元素baseLastTileLength 256余下 64 个元素对应 4 个 DataBlock 再分给前 4 个整核整核尾块变为256 16 272个元素。此时tailBlockNumEachCore ! 0对应TailBlockAndTailCoreTiling分支。样例输入输出规格场景0类型nameshapedata typeformat样例输入x[1, 256000]halfND样例输入y[1, 256000]halfND样例输出z[1, 256000]halfND核函数名add_custom---场景1类型nameshapedata typeformat样例输入x[1, 260096]halfND样例输入y[1, 260096]halfND样例输出z[1, 260096]halfND核函数名add_custom---场景2类型nameshapedata typeformat样例输入x[1, 256064]halfND样例输入y[1, 256064]halfND样例输出z[1, 256064]halfND核函数名add_custom---场景3类型nameshapedata typeformat样例输入x[1, 258112]halfND样例输入y[1, 258112]halfND样例输出z[1, 258112]halfND核函数名add_custom---Host 侧 Tiling 计算过程本样例的 Tiling 计算在 Host 侧完成输入为总数据量totalLength逐级计算主块数量、尾块数量、尾核数量最终输出 Tiling 切分结果。以下过程与 vector_tiling_strategy.asc 中LaunchAddByTiling的实现一一对应。第 1 步按 DataBlock 对齐总数据量totalLengthAligned AlignUp(totalLength, alignNum);AlignUp的定义为((value align - 1U) / align) * align。由于设备侧DataCopy、矢量计算等操作要求数据满足 32 字节DataBlock对齐因此先对总长度向上取整对齐。样例中四个场景的totalLength本身已是 16 的整数倍对齐后数值不变。第 2 步所有核均分主块mainTileNum totalLengthAligned / (numBlocks * mainTileLength); mainTileRemainder totalLengthAligned % (numBlocks * mainTileLength);此时每个核至少处理mainTileNum * mainTileLength个元素mainTileRemainder表示完成主块分配后的剩余数据量。当mainTileRemainder 0时所有核数据量一致且只含主块进入主块均分场景并结束 Tiling 计算if (mainTileRemainder 0) { // 使用 MainTileOnlyTiling。 return; }第 3 步按 DataBlock 给所有核补齐等长尾块tailBlockNumEachCore mainTileRemainder / (numBlocks * alignNum); remainingTailLength mainTileRemainder % (numBlocks * alignNum); baseLength mainTileNum * mainTileLength tailBlockNumEachCore * alignNum;baseLength表示每个核处理的基础数据量主块 等长 DataBlock 尾块。remainingTailLength表示每核补齐等长尾块后尚未分配的元素数量其值恒为alignNum的整数倍——因为剩余部分无法再让所有 8 核各分到 1 个 DataBlock只能分配给部分核。第 4 步判断尾块均分场景if (remainingTailLength 0) { // 使用 MainTileWithTailBlockTiling。 return; }当remainingTailLength 0时所有核均分到等长尾块进入尾块均分场景。第 5 步剩余 DataBlock 分给前若干个整核formerNum remainingTailLength / alignNum; formerLength baseLength alignNum; tailLength baseLength;将最后剩余的 DataBlock 分给前formerNum个整核整核比尾核多处理 1 个 DataBlock因此formerLength baseLength alignNum尾核仍为baseLength。第 6 步区分尾核切分与尾核尾块切分if (tailBlockNumEachCore 0) { // 使用 MainTileWithTailCoreTiling。 } else { // 使用 TailBlockAndTailCoreTiling。 }tailBlockNumEachCore 0时尾核内只包含主块、没有基础尾块进入尾核切分场景否则所有核都有尾块进入尾核尾块切分场景。一个重要的边界处理在尾核尾块切分场景的源码中还存在一个细节优化见 vector_tiling_strategy.ascformerLastTileLength baseLastTileLength alignNum; if (formerLastTileLength mainTileLength) { formerTileNum 1U; formerLastTileLength 0U; }当整核的尾块长度恰好等于一个主块长度3200时不额外创建一个独立尾块而是将整核的主块数量formerTileNum加 1、尾块长度置 0从而让整核完全按主块粒度处理数据。这种尾块升级为主块的合并策略可以避免在主块循环之外产生一个几乎等长的尾块循环简化核内流水调度。四种 Tiling 数据结构与核内执行流程四个 Tiling 数据结构Host 侧计算出切分参数后会按场景实例化对应结构体并通过add_customTilingData模板核函数传入设备侧见 vector_tiling_strategy.asc结构体对应场景关键字段MainTileOnlyTiling主块均分blockLength、tileNumMainTileWithTailBlockTiling尾块均分上述字段 lastTileLengthMainTileWithTailCoreTiling尾核切分formerNum、formerLength、formerTileNum、formerLastTileLength、tailLength、tailTileNumTailBlockAndTailCoreTiling尾核尾块切分上述字段 tailLastTileLength核函数启动采用模板参数传递 Tiling 数据的方式template typename TilingData __vector__ __global__ void add_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, TilingData tiling) { AscendC::InitSocState(); KernelAdd op; op.Init(x, y, z, tiling); op.Process(); AscendC::PipeBarrierPIPE_ALL(); }每个核通过AscendC::GetBlockIdx()该接口声明见 include/basic_api/kernel_operator_sys_var_intf.h获取自身的逻辑核号进而定位自己负责的 GMGlobal Memory数据段主块均分 / 尾块均分场景offset GetBlockIdx() * tiling.blockLength各核数据段连续且等长尾核切分 / 尾核尾块切分场景先判断blockIdx tiling.formerNum整核段为formerLength * blockIdx尾核段起始位置为formerLength * formerNum tailLength * (blockIdx - formerNum)保证前后核数据不重叠、不遗漏。核内主循环的流水与同步KernelAdd::Process()见 vector_tiling_strategy.asc是核内计算的核心按主块 可选尾块组织循环uint32_t totalLoop this-tileNum (this-lastTileLength 0U ? 1U : 0U); for (uint32_t loopIdx 0; loopIdx totalLoop; loopIdx) { uint32_t curLen (loopIdx this-tileNum) ? mainTileLength : this-lastTileLength; uint32_t startElement loopIdx * mainTileLength; // ... AscendC::DataCopy(xLocal, xGm[startElement], curLen); // GM - UBMTE2 AscendC::DataCopy(yLocal, yGm[startElement], curLen); AscendC::Add(zLocal, xLocal, yLocal, curLen); // UB 矢量加法Vector AscendC::DataCopy(zGm[startElement], zLocal, curLen); // UB - GMMTE3 }每轮循环完成搬入MTE2→ 计算Vector→ 搬出MTE3三个阶段并通过WaitFlag/SetFlag成对的事件同步HardEvent::V_MTE2、MTE2_V、V_MTE3、MTE3_V保证相邻循环之间的数据依赖例如第loopIdx轮搬入前先等待上一轮矢量计算结束WaitFlagV_MTE2避免覆盖上一轮尚未消费完的 UB 缓冲区。UB 侧在样例中为三个主块预留了连续缓冲区xAddr 0、yAddr mainTileLength * halfSize、zAddr yAddr mainTileLength * halfSize即各占 3200 × 2 6400 字节见 vector_tiling_strategy.asc。这解释了为什么主块长度必须基于 UB 可用空间确定——三个缓冲区之和必须能同时装入 UB。值得说明的是样例中的AscendC::Add(zLocal, xLocal, yLocal, curLen)与AscendC::DataCopy均来自 asc-devkit 提供的标准 API矢量加法 Level 2 接口dst src0 src1声明见 include/basic_api/kernel_operator_vec_binary_intf.h数据搬移接口DataCopy的重载定义见 include/basic_api/kernel_operator_data_copy_intf.h。其中DataCopy支持从 GlobalTensor 到 LocalTensor 的搬运计数元素个数这正是样例按元素长度curLen搬运x、y两路输入的依据。Host 侧整体流程主函数main()见 vector_tiling_strategy.asc执行完整的 ACL 流程aclInit初始化 →aclrtSetDevice指定设备 → 创建 stream →aclrtMallocHost/aclrtMalloc分配 Host 与 Device 内存 →ReadFile读入输入 →aclrtMemcpy拷贝至 Device →LaunchAddByTiling启动核函数 →aclrtSynchronizeStream同步 → 结果回拷 →WriteFile写出output.bin→ 释放资源。其中输入输出文件均为原始二进制格式dataByteSize dataLen * halfSize。编译与运行配置环境变量请根据当前环境上 CANN 开发套件包的安装方式配置环境变量source ${install_path}/cann/set_env.sh说明${install_path}为 CANN 包安装目录未指定安装目录时默认安装至/usr/local/Ascend下。样例执行在样例根目录 vector_tiling_strategy 下执行如下命令SCENARIO_NUM0 # 设置场景编号 mkdir -p build cd build; # 创建并进入build目录 cmake .. -DCMAKE_ASC_ARCHITECTURESdav-2201 -DSCENARIO_NUM$SCENARIO_NUM;make -j; # 编译工程 python3 ../scripts/gen_data.py -scenarioNum$SCENARIO_NUM # 生成测试输入数据 ./demo # 执行编译生成的可执行程序执行样例 python3 ../scripts/verify_result.py ./output/output.bin ./output/golden.bin # 验证输出结果是否正确使用 CPU 调试或 NPU 仿真模式时添加-DCMAKE_ASC_RUN_MODEcpu或-DCMAKE_ASC_RUN_MODEsim参数即可示例如下cmake .. -DCMAKE_ASC_RUN_MODEcpu -DCMAKE_ASC_ARCHITECTURESdav-2201 -DSCENARIO_NUM$SCENARIO_NUM;make -j; # CPU调试模式 cmake .. -DCMAKE_ASC_RUN_MODEsim -DCMAKE_ASC_ARCHITECTURESdav-2201 -DSCENARIO_NUM$SCENARIO_NUM;make -j; # NPU仿真模式注意切换编译模式前需清理 cmake 缓存可在 build 目录下执行rm CMakeCache.txt后重新 cmake。编译选项说明选项可选值说明CMAKE_ASC_RUN_MODEnpu默认、cpu、sim运行模式NPU 运行、CPU 调试、NPU 仿真CMAKE_ASC_ARCHITECTURESdav-2201默认、dav-3510NPU 架构dav-2201对应 Atlas A2 系列产品和 Atlas A3 系列产品dav-3510对应 Ascend 950PR950DT 系列产品SCENARIO_NUM0默认、1、2、3场景编号0主块均分、1尾块均分、2尾核切分、3尾核尾块切分其中SCENARIO_NUM在 CMakeLists.txt 中做了取值校验——非0~3的取值会触发FATAL_ERROR避免无效场景导致数据长度与切分逻辑不匹配。执行结果执行成功且精度对比通过时输出如下test pass!若编译模式切换后未清理缓存导致参数未生效可参考上文注意项清理CMakeCache.txt后重试。数据生成与结果验证输入与真值数据生成scripts/gen_data.py 使用 NumPy 生成随机数据input_x、input_y在[-8, 8]区间内均匀采样并转换为float16即halfgolden 真值为两者相加后再转回float16。随机数生成器以2026 scenario为种子保证同一场景多次生成的数据可复现。生成结果写入./input/input_x.bin、./input/input_y.bin与./output/golden.bin。输出结果验证scripts/verify_result.py 将output.bin与golden.bin读入后逐元素比较采用np.isclose判断相对容差1e-4、绝对容差1e-5同时计算整体误差率只要误差率不超过1e-4即判定通过并打印test pass!。当存在差异时脚本会输出最多 100 个差异点的索引、期望值、实际值与相对误差便于定位问题size mismatch会被视为失败并退出。样例设计要点总结从源码结构可以归纳出本样例体现的几个 Tiling 设计要点可直接复用到其他矢量算子对齐优先任何切分都以 DataBlock32 字节为最小分配单位先整体对齐、再逐级分配保证搬运与计算接口的数据对齐要求主块固定、尾块兜底主块长度由 UB 容量决定且保持不变剩余数据一律以尾块形态按 DataBlock 粒度继续分配避免破坏主块的流水效率核间均衡优先核内合并兜底优先让所有核数据量相等场景 0/1无法完全相等时只让少数整核多出 1 个 DataBlock场景 2/3并在整核尾块恰好等于主块长度时将其合并为主块简化循环结构Tiling 参数结构体按场景区分四种结构体精确承载各自场景所需的字段通过模板核函数传递避免冗余字段与运行时类型判断。该样例所属的 Tiling Selector 特性目录中还有更整体的介绍见 02_tiling_selector 目录说明即根据输入数据量和 AI Core 数量计算 Tiling 参数并选择核间和核内切分策略的典型流程。开发者可以在此基础上将本例的切分思路推广到更多核数、更大数据量以及多维度 Tensor 的算子场景。赞分享人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载相关推荐Ascend C 多核 Tiling 切分策略全解析以 Tiling Selector 矢量算子样例解读核间与核内数据切分Ascend C 多核 Tiling 切分策略全解析以 Tiling Selector 矢量算子样例解读核间与核内数据切分 本文围绕 asc devkit 仓人工智能深度学习算子库CANNAscendAscend C SIMD Tiling Selector多核Tiling切分策略的计算、选择与实战Ascend C SIMD Tiling Selector多核Tiling切分策略的计算、选择与实战 Ascend C 提供了 SIMD Tiling Sel人工智能深度学习算子库CANNAscend昇腾 Ascend C AddBroadcast 算子核函数直调实战多核并行、Tiling 切分与 Broadcast 广播语义昇腾 Ascend C AddBroadcast 算子核函数直调实战多核并行、Tiling 切分与 Broadcast 广播语义 导读 本篇文章围绕 asc人工智能深度学习算子库CANNAscend上一篇nlohmann_json is_array() 深度解析JSON 数组类型判断的语义、实现与测试验证下一篇NEP 31 全解析NumPy API 的全局与上下文局部覆盖机制uarray/unumpy 方案创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考