人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载本文围绕 asc-devkit 仓库中的Tiling Selector 特性样例02_tiling_selector展开以两个一维halfTensor 相加的矢量算子为例系统讲解如何在 Host 侧根据输入数据量与参与计算的 AI Core 数量计算 Tiling 参数并在主块均分、尾块均分、尾核切分、尾核尾块切分四种典型策略之间完成选择。读完本文你将掌握多核矢量算子 Tiling 切分的完整实现套路Tiling 结构体设计、Host 侧逐级切分算法、Device 侧按核偏移计算与分块循环处理以及样例在 NPU/CPU/仿真模式下的编译运行与精度验证方法。1. 为什么要做多核 Tiling背景与核心概念在昇腾 AI Core 上核函数Kernel是 Device 侧执行的并行函数。为了提升算子执行效率通常将输入数据切分后分配到多个逻辑核并行处理同时单个核的 Local Memory如 Unified Buffer即 UB容量有限无法一次容纳全部输入输出必须分批搬运数据、分块计算。这种对输入数据进行切分、分块计算的过程称为Tiling切分数据的算法即 Tiling 策略而根据 shape 等信息计算切分参数块大小、循环次数等的程序称为Tiling 实现Tiling 函数。相关背景可参阅 multi_core_tiling/overview.md。Tiling 切分遵循三条基本原则以 DataBlock 为最小切分单位受硬件限制UB 相关的数据搬运与矢量计算要求数据长度和起始地址 32 字节对齐。一个 DataBlock 为 32 字节输入数据不满足对齐要求时需按 DataBlock 向上对齐。充分利用 UB 空间减少从 Global Memory 到 UB 的搬运次数避免频繁搬运造成的性能瓶颈。均衡利用多核计算能力将计算尽量均衡地分配到各 AI Core充分发挥多核并行优势。围绕核间与核内两级切分本文涉及四个关键概念主块基于 Local Memory 可用空间确定的单个核一次处理的最大数据块长度用mainTileLength表示按 32 字节对齐。尾块单个核内最后一个长度小于mainTileLength的数据块长度用lastTileLength表示。整核核间数据无法完全均分时处理数据量较多的核。尾核核间数据无法完全均分时处理数据量较少的核。2. 样例总览与目录结构02_tiling_selector 目录展示的是 Ascend C SIMDTiling Selector特性样例根据输入数据量和 AI Core 数量计算 Tiling 参数并选择核间inter-core与核内intra-core切分策略。其下唯一的算子开发样例为 vector_tiling_strategy功能描述如下目录名称功能描述vector_tiling_strategy以两个一维 Tensor 相加为例展示矢量计算场景下多核 Tiling 切分策略的计算与选择流程。vector_tiling_strategy样例的目录结构如下详见 vector_tiling_strategy/README.md├── vector_tiling_strategy │ ├── scripts │ │ ├── gen_data.py // 输入数据和真值数据生成脚本 │ │ └── verify_result.py // 输出结果验证脚本 │ ├── CMakeLists.txt // 编译工程文件 │ ├── vector_tiling_strategy.asc // Ascend C样例实现与调用样例 │ ├── data_utils.h // 数据读入写出函数 │ └── README.md // 样例说明文档2.1 支持的产品及 CANN 软件版本产品CANN 软件版本Ascend 950PR 950DT 系列产品≥ CANN 9.1.0Atlas A3 系列产品≥ CANN 9.0.0Atlas A2 系列产品≥ CANN 9.0.02.2 样例规格总览本样例计算逻辑为z x y两个输入与一个输出均为形状[1, N]、数据类型half、格式 ND 的一维 Tensor核函数名为add_custom。样例固定使用8 个核主块数据量为3200 个half元素即 200 个 DataBlock通过 CMake 编译参数SCENARIO_NUM选择不同数据量场景共覆盖四种 Tiling 切分场景场景SCENARIO_NUM输入元素数每核处理元素数核内组成切分特点主块均分02560003200010 个主块各核数据量相同核内只有主块尾块均分12600963251210 个主块 512 元素尾块各核数据量相同尾块等长尾核切分2256064前 4 核 32016 / 后 4 核 32000整核10 主块 16 元素尾块尾核10 主块前 4 核为整核尾核无尾块尾核尾块切分3258112前 4 核 32272 / 后 4 核 32256整核10 主块 272 元素尾块尾核10 主块 256 元素尾块所有核都有尾块整核多处理 1 个 DataBlock主块mainTileLength是用户基于 UB 可用空间确定的一次可以处理的最大数据量且满足 32 字节对齐。3. 核心常量与 Tiling 结构体设计3.1 常量定义样例在 vector_tiling_strategy.asc 开头定义了全部切分基础常量constexpr uint32_t scenarioNum SCENARIO_NUM; // 场景编号由CMake编译参数注入 constexpr uint32_t numBlocks 8; // 参与计算的逻辑核数 constexpr uint32_t dataBlockSize 32; // 一个DataBlock的字节数 constexpr uint32_t halfSize 2; // 一个half元素的字节数 constexpr uint32_t alignNum dataBlockSize / halfSize; // 一个DataBlock包含的half元素数量即16 constexpr uint32_t mainTileLength 3200; // 一个主块包含的half元素数量其中alignNum 16是全部分支判断的核心粒度任何剩余数据都以一个 DataBlock 含 16 个half为单位进行整除与取余。3.2 四种 Tiling 结构体不同场景需要传递的切分信息不同样例定义了四种 Tiling 结构体见 vector_tiling_strategy.asc主块均分场景各核数据量相同仅需每核长度与主块个数。struct MainTileOnlyTiling { uint32_t blockLength; // 每个核处理的元素数量 uint32_t tileNum; // 主块数量 };尾块均分场景在上一结构体基础上新增尾块长度。struct MainTileWithTailBlockTiling { uint32_t blockLength; // 每个核处理的元素数量 uint32_t tileNum; // 主块数量 uint32_t lastTileLength; // 尾块元素数量 };尾核切分场景整核与尾核参数不同用former*描述整核、tail*描述尾核尾核只含主块无需tailLastTileLength。struct MainTileWithTailCoreTiling { uint32_t formerNum; // 整核数量 uint32_t formerLength; // 整核处理的元素数量 uint32_t formerTileNum; // 整核中主块的数量 uint32_t formerLastTileLength; // 整核中尾块的元素数量 uint32_t tailLength; // 尾核处理的元素数量 uint32_t tailTileNum; // 尾核中主块的数量 };尾核尾块切分场景尾核也存在尾块补充tailLastTileLength。struct TailBlockAndTailCoreTiling { uint32_t formerNum; // 整核数量 uint32_t formerLength; // 整核处理的元素数量 uint32_t formerTileNum; // 整核中主块的数量 uint32_t formerLastTileLength; // 整核中尾块的元素数量 uint32_t tailLength; // 尾核处理的元素数量 uint32_t tailTileNum; // 尾核中主块的数量 uint32_t tailLastTileLength; // 尾核中尾块的元素数量 };4. 四种切分场景详解4.1 场景 0主块均分Main Tile Equal Split输入x、y均为[1, 256000]个half元素输出z为[1, 256000]个half元素。Tiling 切分8 个核各处理 32000 个元素每个核包含 10 个主块每个主块 3200 个元素。说明所有核处理数据量相同核内只包含主块。数据量校验256000 / (8 × 3200) 10整除无余数因此每个核恰好分到 10 个主块。对应的MainTileOnlyTiling结构体取值为{32000, 10}。4.2 场景 1尾块均分Tail Block Equal Split输入x、y均为[1, 260096]个half元素输出z为[1, 260096]个half元素。Tiling 切分8 个核各处理 32512 个元素每个核包含 10 个主块和 512 个元素的尾块。说明所有核处理数据量相同每个核在主块后都有等长尾块。数据量校验260096 8 × 32512其中32512 10 × 3200 512512 32 × 16即每核多分 32 个 DataBlock 作为等长尾块。MainTileWithTailBlockTiling结构体取值为{32512, 10, 512}。4.3 场景 2尾核切分Tail Core Split输入x、y均为[1, 256064]个half元素输出z为[1, 256064]个half元素。Tiling 切分前 4 个整核各处理 32016 个元素后 4 个尾核各处理 32000 个元素整核包含 10 个主块和 16 个元素的尾块尾核包含 10 个主块。说明前 4 个整核比尾核多处理 1 个 DataBlock尾核无尾块。数据量校验256064 8 × 32000 64剩余 64 个元素即 4 个 DataBlock只能分给前 4 个核每核多 1 个 DataBlock即 16 个元素因此formerNum 4。MainTileWithTailCoreTiling结构体取值为{4, 32016, 10, 16, 32000, 10}。4.4 场景 3尾核尾块切分Tail Core Split with Tail Blocks输入x、y均为[1, 258112]个half元素输出z为[1, 258112]个half元素。Tiling 切分前 4 个整核各处理 32272 个元素后 4 个尾核各处理 32256 个元素整核包含 10 个主块和 272 个元素的尾块尾核包含 10 个主块和 256 个元素的尾块。说明所有核都有尾块前 4 个整核再多处理 1 个 DataBlock。数据量校验258112 8 × 32000 21122112 8 × 256 64即每核先分 16 个 DataBlock 基础尾块256 元素剩余 64 元素即 4 个 DataBlock 再分给前 4 核。TailBlockAndTailCoreTiling结构体取值为{4, 32272, 10, 272, 32256, 10, 256}。四种场景的输入输出规格表名称 / shape / 数据类型 / 格式汇总如下均满足输入输出同为 ND 格式的half核函数名add_custom场景输入 x / y shape输出 z shape数据类型格式核函数名0[1, 256000][1, 256000]halfNDadd_custom1[1, 260096][1, 260096]halfNDadd_custom2[1, 256064][1, 256064]halfNDadd_custom3[1, 258112][1, 258112]halfNDadd_custom5. Host 侧 Tiling 计算流程Host 侧输入为totalLength逐级计算主块数量、尾块数量、尾核数量最终输出 Tiling 切分结果。完整实现位于 vector_tiling_strategy.asc 的LaunchAddByTiling函数中。第 0 步按 DataBlock 对齐总数据量。uint32_t totalLengthAligned AlignUp(totalLength, alignNum);其中AlignUp实现为((value align - 1U) / align) * align。本样例四种场景的totalLength均为alignNum16的整数倍因此对齐前后数值不变。第 1 步在所有核上均分相同数量的主块。uint32_t mainTileNum totalLengthAligned / (numBlocks * mainTileLength); uint32_t mainTileRemainder totalLengthAligned % (numBlocks * mainTileLength);此时每个核至少处理mainTileNum * mainTileLength个元素mainTileRemainder是主块分配后的剩余数据量。当mainTileRemainder 0时进入主块均分场景并结束if (mainTileRemainder 0U) { MainTileOnlyTiling tiling{baseLength, baseTileNum}; printf(Tiling mode: main_tile_only\n); add_customMainTileOnlyTilingnumBlocks, 0, stream(xDevice, yDevice, zDevice, tiling); return; }第 2 步mainTileRemainder不为 0 时按 DataBlock 给所有核补齐等长尾块。uint32_t tailBlockNumEachCore mainTileRemainder / (numBlocks * alignNum); uint32_t remainingTailLength mainTileRemainder % (numBlocks * alignNum); uint32_t baseLength mainTileNum * mainTileLength tailBlockNumEachCore * alignNum;baseLength是每个核处理的基础数据量remainingTailLength是每核补齐等长 DataBlock 尾块后尚未分配的元素数量其值一定是alignNum的整数倍。第 3 步判断尾块均分场景。remainingTailLength 0表示所有核均分到等长尾块进入尾块均分场景if (remainingTailLength 0U) { MainTileWithTailBlockTiling tiling{baseLength, baseTileNum, baseLastTileLength}; printf(Tiling mode: main_tile_with_tail_block\n); add_customMainTileWithTailBlockTilingnumBlocks, 0, stream(xDevice, yDevice, zDevice, tiling); return; }第 4 步remainingTailLength不为 0 时将最后剩余的 DataBlock 分给前若干个整核并计算整核与尾核的数据长度。uint32_t formerNum remainingTailLength / alignNum; uint32_t formerLength baseLength alignNum; // 前formerNum个核都多处理一个DataBlock的数据 uint32_t tailLength baseLength; // 尾核处理的数据量第 5 步判断尾核切分 vs 尾核尾块切分。tailBlockNumEachCore 0时各核未分配到基础尾块尾核内只包含主块进入尾核切分场景否则进入尾核尾块切分场景if (tailBlockNumEachCore 0U) { uint32_t formerLastTileLength alignNum; MainTileWithTailCoreTiling tiling{formerNum, formerLength, formerTileNum, formerLastTileLength, tailLength, tailTileNum}; printf(Tiling mode: main_tile_with_tail_core\n); add_customMainTileWithTailCoreTilingnumBlocks, 0, stream(xDevice, yDevice, zDevice, tiling); } else { uint32_t formerLastTileLength baseLastTileLength alignNum; if (formerLastTileLength mainTileLength) { // 若整核尾块恰好凑满一个主块则合并为主块 formerTileNum 1U; formerLastTileLength 0U; } uint32_t tailLastTileLength baseLastTileLength; TailBlockAndTailCoreTiling tiling{formerNum, formerLength, formerTileNum, formerLastTileLength, tailLength, tailTileNum, tailLastTileLength}; printf(Tiling mode: tail_block_and_tail_core\n); add_customTailBlockAndTailCoreTilingnumBlocks, 0, stream(xDevice, yDevice, zDevice, tiling); }整核尾块长度的合并处理是一个容易被忽略的细节当baseLastTileLength alignNum恰好等于mainTileLength时整核的主块 尾块会退化为多一个完整主块此时应把formerLastTileLength置 0、formerTileNum加 1保证核内循环按主块粒度处理。以四个场景代入上述流程可完整还原结构体取值场景关键中间量命中分支结构体取值0mainTileRemainder 0main_tile_only{32000, 10}1mainTileRemainder 4096remainingTailLength 0main_tile_with_tail_block{32512, 10, 512}2remainingTailLength 64tailBlockNumEachCore 0main_tile_with_tail_core{4, 32016, 10, 16, 32000, 10}3remainingTailLength 64tailBlockNumEachCore 16tail_block_and_tail_core{4, 32272, 10, 272, 32256, 10, 256}6. Device 侧核函数实现6.1 核函数入口与模板化 Tiling 传递核函数add_custom是模板函数TilingData类型即第 3 节定义的四种结构体之一。Host 侧调用add_customTilingDatanumBlocks, 0, stream时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(); }6.2 Init按核索引计算 Global Memory 偏移每个核访问的 Global Memory 地址范围不同Init通过 GetBlockIdx() 获取当前逻辑 AI Core 索引并结合 Tiling 参数计算起始偏移。GetBlockIdx 的返回值范围为[0, numBlocks)numBlocks由调用配置。主块均分 / 尾块均分各核数据量相同偏移为blockIdx × blockLength尾块均分场景额外保存lastTileLength。__aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, MainTileWithTailBlockTiling tiling) { uint32_t offset AscendC::GetBlockIdx() * tiling.blockLength; InitGm(x, y, z, offset, tiling.blockLength); this-tileNum tiling.tileNum; this-lastTileLength tiling.lastTileLength; }尾核切分 / 尾核尾块切分整核与尾核数据量不同先判断当前核是否为整核blockIdx tiling.formerNum整核偏移为formerLength × blockIdx尾核偏移需跳过全部整核数据为formerLength × formerNum tailLength × (blockIdx - formerNum)__aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, TailBlockAndTailCoreTiling tiling) { uint32_t blockIdx AscendC::GetBlockIdx(); if (blockIdx tiling.formerNum) { // 整核 uint32_t offset tiling.formerLength * blockIdx; InitGm(x, y, z, offset, tiling.formerLength); this-tileNum tiling.formerTileNum; this-lastTileLength tiling.formerLastTileLength; } else { // 尾核 uint32_t offset tiling.formerLength * tiling.formerNum tiling.tailLength * (blockIdx - tiling.formerNum); InitGm(x, y, z, offset, tiling.tailLength); this-tileNum tiling.tailTileNum; this-lastTileLength tiling.tailLastTileLength; } }InitGm通过 SetGlobalBuffer 为输入xGm、yGm和输出zGm设置 Global Memory 地址(__gm__ half*)x offset为当前核数据的起始地址length为该核处理的元素数量。6.3 Process分块循环搬运、计算与同步Process采用静态 Tensor 编程方式实现参考 static_tensor_programming.md通过LocalTensor构造函数直接指定 VECCALC 位置上的固定地址避免 TPipe 框架初始化开销。xAddr、yAddr、zAddr依次线性排布在 UB 上static constexpr uint32_t xAddr 0; static constexpr uint32_t yAddr mainTileLength * halfSize; static constexpr uint32_t zAddr yAddr mainTileLength * halfSize;循环次数采用通用形式tileNum (lastTileLength 0 ? 1 : 0)计算前tileNum轮处理主块最后一轮若存在处理尾块curLen为当前轮元素数startElement loopIdx * mainTileLength为核内偏移。由于静态 Tensor 编程需要开发者自行管理流水同步Process内通过SetFlag/WaitFlag维护 MT E2 搬入、Vector 计算、MTE3 搬出三条流水之间的依赖__aicore__ inline void Process() { AscendC::LocalTensorhalf xLocal(AscendC::TPosition::VECCALC, xAddr, mainTileLength); AscendC::LocalTensorhalf yLocal(AscendC::TPosition::VECCALC, yAddr, mainTileLength); AscendC::LocalTensorhalf zLocal(AscendC::TPosition::VECCALC, zAddr, mainTileLength); 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; if (loopIdx ! 0U) { AscendC::WaitFlagAscendC::HardEvent::V_MTE2(EVENT_ID0); // 反向同步上一次计算完成后再搬入 } AscendC::DataCopy(xLocal, xGm[startElement], curLen); // 搬入输入x AscendC::DataCopy(yLocal, yGm[startElement], curLen); // 搬入输入y AscendC::SetFlagAscendC::HardEvent::MTE2_V(EVENT_ID0); AscendC::WaitFlagAscendC::HardEvent::MTE2_V(EVENT_ID0); // 正向同步搬入完成后再计算 if (loopIdx ! 0U) { AscendC::WaitFlagAscendC::HardEvent::MTE3_V(EVENT_ID0); // 反向同步上一次搬出完成后再计算 } AscendC::Add(zLocal, xLocal, yLocal, curLen); // 矢量加计算 if (loopIdx ! totalLoop - 1U) { AscendC::SetFlagAscendC::HardEvent::V_MTE2(EVENT_ID0); } AscendC::SetFlagAscendC::HardEvent::V_MTE3(EVENT_ID0); AscendC::WaitFlagAscendC::HardEvent::V_MTE3(EVENT_ID0); // 正向同步计算完成后再搬出 AscendC::DataCopy(zGm[startElement], zLocal, curLen); // 搬出输出z if (loopIdx ! totalLoop - 1U) { AscendC::SetFlagAscendC::HardEvent::MTE3_V(EVENT_ID0); } } }这里的同步模式与多核 Tiling 概述文档中的描述一致正向同步MTE2_V、V_MTE3保证先搬入后计算、先计算后搬出反向同步V_MTE2、MTE3_V防止上一轮数据被覆盖。四种场景共用同一个Process函数差异全部由tileNum与lastTileLength两个成员变量体现这正是 Tiling 结构体设计的价值所在。7. 编译运行与结果验证7.1 环境准备根据当前环境上 CANN 开发套件包的安装方式配置环境变量source ${install_path}/cann/set_env.sh说明${install_path}为 CANN 包安装目录未指定安装目录时默认安装至/usr/local/Ascend下。7.2 样例执行在样例根目录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。7.3 编译选项说明编译选项的默认值与取值校验定义在 CMakeLists.txt 中SCENARIO_NUM通过target_compile_definitions以宏定义形式注入源码constexpr uint32_t scenarioNum SCENARIO_NUM;且 CMake 侧做了取值校验SCENARIO_NUM必须匹配^[0-3]$否则FATAL_ERROR选项可选值说明CMAKE_ASC_RUN_MODEnpu默认、cpu、sim运行模式NPU 运行、CPU 调试、NPU 仿真CMAKE_ASC_ARCHITECTURESdav-2201默认、dav-3510NPU 架构dav-2201对应 Atlas A2 系列产品和 Atlas A3 系列产品dav-3510对应 Ascend 950PR 950DT 系列产品SCENARIO_NUM0默认、1、2、3场景编号0主块均分、1尾块均分、2尾核切分、3尾核尾块切分7.4 数据生成与精度验证数据生成脚本 gen_data.py 内置了四种场景的数据长度映射SCENARIO_LENGTHS {0: 256000, 1: 260096, 2: 256064, 3: 258112}以固定随机种子2026 scenario在[-8, 8]均匀分布中采样float16数据生成input_x.bin、input_y.bin输入与golden.bin真值。验证脚本 verify_result.py 读取output.bin与golden.bin使用np.isclose按相对容差1e-4、绝对容差1e-5逐元素比对统计错误率并允许ERROR_TOL 1e-4以内的误差half 浮点运算的舍入误差。执行结果如下即说明精度对比成功test pass!7.5 主函数执行流程vector_tiling_strategy.asc 中的main函数展示了完整的 ACL 运行时调用链aclInit→aclrtSetDevice→aclrtCreateStream→aclrtMallocHost/aclrtMalloc申请内存 → data_utils.h 中的ReadFile读取输入 →aclrtMemcpy拷贝至 Device →LaunchAddByTiling发起计算 →aclrtSynchronizeStream等待完成 → 拷回 Host 并WriteFile落盘 → 释放资源并aclFinalize。运行时会打印当前场景与数据量Usage: ./demo (compile-time SCENARIO_NUM0) Running scenario 0, totalLength256000 Tiling mode: main_tile_only Scenario 0 completed.8. 策略选择准则与扩展思考综合第 5 节的逐级判断四种 Tiling 切分策略的选择完全由两个取余结果驱动可归纳为如下决策树mainTileRemainder 0→主块均分MainTileOnlyTiling核内仅主块各核完全均衡remainingTailLength 0→尾块均分MainTileWithTailBlockTiling各核均衡核内含等长尾块tailBlockNumEachCore 0→尾核切分MainTileWithTailCoreTiling出现整核/尾核尾核无尾块其余情况 →尾核尾块切分TailBlockAndTailCoreTiling整核与尾核都含尾块。该决策逻辑可以直接移植到真实算子开发中Tiling 结构体按需扩展例如整核/尾核参数分组、Host 侧以 DataBlockalignNum为粒度做整除取余、Device 侧用GetBlockIdx()分支计算偏移这三步是所有多核矢量算子 Tiling 的通用范式。进一步扩展时可参考仓库中 multi_core_tiling 的四个分册主块均分、尾块均分、尾核切分、尾核尾块切分了解每个场景更细粒度的实现推导矢量计算 API如 Add与静态 Tensor 编程的同步管理细节可分别查阅对应文档。若要在自有工程中落地可将样例的LaunchAddByTiling泛化为输入 shape → 切分参数 → 核函数调用的工具函数并根据实际 UB 容量重新计算mainTileLength。赞分享人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载相关推荐Ascend 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人工智能深度学习算子库CANNAscendPyPTO Tiling 机制全解析多核切分、TilingKey/TilingData 与片上 Tile 数据流PyPTO Tiling 机制全解析多核切分、TilingKey/TilingData 与片上 Tile 数据流 本文是 CANN pypto 项目PyPT人工智能编译器模型编译深度学习高性能计算CANNAscend创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考