人工智能指令集算子库CANNAscend【免费下载链接】pto-isaParallel Tile Operation (PTO) is a virtual instruction set architecture designed by Ascend CANN, focusing on tile-level operations. This repository offers high-performance, cross-platform tile operations across Ascend platforms.项目地址https://gitcode.com/cann/pto-isa点击查看免费下载TCOLEXPANDADD 是 CANN PTOParallel Tile Operation虚拟指令集 PTOISA.md 中 TCOLEXPAND 族的一员用于将src1提供的每列一个标量广播到整列并与src0对应元素执行加法结果写入dst。本文从数学语义、两级汇编形式、C 内建接口、数据类型与布局约束出发结合 A2/A3/A5 后端与 CPU 模拟器源码及 NPU 测试用例完整还原该指令从声明、发射到向量单元执行的全链路帮助算子开发者在自动模式与手动模式下正确使用并理解其底层行为。指令示意图指令语义与数学解释TCOLEXPANDADD 属于按列广播的二元向量运算它不是对整个 Tile 执行标量广播而是把src1中的第j个标量值广播到第j列的所有行再逐元素执行加法。设R dst.GetValidRow()、C dst.GetValidCol()s_j为从src1中取出的第j列标量则对于0 i R且0 j C$$ \mathrm{dst}{i,j} \mathrm{src0}{i,j} s_j $$直观理解src0与dst的有效形状相同R行C列逐元素参与运算src1的有效形状只需覆盖C个值通常为1 x C或R_src1 x C其中R_src1可为 1其第j个元素沿列方向扩展这与 TCOLEXPAND仅做纯广播、无运算形成对照TCOLEXPANDADD 将列扩展与二元加法融合为单条指令省去先广播再相加的两步操作。从指令族的角度看TCOLEXPAND 族在 include/pto/common/event.hpp 中按二元算子成族排列TCOLEXPANDDIV、TCOLEXPANDMUL、TCOLEXPANDADD、TCOLEXPANDMAX、TCOLEXPANDMIN、TCOLEXPANDSUB、TCOLEXPANDEXPDIF。其中TCOLEXPANDADD在 include/pto/common/event.hpp 被映射到PIPE_V向量计算单元即该指令由 AICVector 侧执行。汇编语法同步形式%dst tcolexpandadd %src0, %src1 : !pto.tile..., !pto.tile... - !pto.tile...AS Level 1SSA%dst pto.tcolexpandadd %src0, %src1 : !pto.tile..., !pto.tile... - !pto.tile...AS Level 2DPSpto.tcolexpandadd ins(%src0, %src1 : !pto.tile_buf..., !pto.tile_buf...) outs(%dst : !pto.tile_buf...)两级汇编的区别在于AS Level 1SSA 形式以!pto.tile...逻辑 Tile 为操作数由编译器负责资源放置与调度主要用于自动模式AS Level 2DPS 形式操作数变为显式绑定存储地址的!pto.tile_buf...ins(...)列出输入、outs(...)列出输出是更贴近硬件资源分配的下沉形式。C 内建接口TCOLEXPANDADD 的 C 接口声明于 include/pto/common/pto_instr.hpp公共包含头为pto/pto-inst.hpptemplate typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents PTO_INST RecordEvent TCOLEXPANDADD(TileDataDst dst, TileDataSrc0 src0, TileDataSrc1 src1, WaitEvents ... events);接口特点模板参数TileDataDst / TileDataSrc0 / TileDataSrc1决定了 Tile 的形状、数据类型与布局RowMajor变参WaitEvents ... events允许传入依赖事件调用会先执行detail::PtoWaitEvents(events...)等待前置指令完成再经MAP_INSTR_IMPL(TCOLEXPANDADD, dst, src0, src1)映射到后端实现返回值RecordEvent记录了本次发射可继续传给后续指令构成事件链从而实现手动模式下跨管道MTE2 → V → MTE3的显式同步。约束条件数据类型TileDataDst::DType、TileDataSrc0::DType、TileDataSrc1::DType必须一致且属于以下集合half、float、int16、int32适用于 Atlas A2 训练系列产品 / Atlas A2 推理系列产品、Atlas A3 训练系列产品 / Atlas A3 推理系列产品以及 Ascend 950PR / Ascend 950DTuint16、uint32、bfloat16_t、int8、uint8、int64、uint64仅适用于 Ascend 950PR / Ascend 950DT即 A5。该约束在源码层面对应双重检查后端实现中的static_assert如 include/pto/npu/a2a3/TColExpandBinOp.hpp编译期校验 A2/A3 侧的数据类型集合CPU 模拟器在 include/pto/cpu/TColExpandOp.hpp 中通过IsColExpandAllowedType与CheckColExtendTiles约束类型一致性dst与src0、src1类型必须相同并规定OP_ADD允许float、half及int64/uint64/int32/int16/uint32/uint16等整型OP_EXPDIF除外。Tile 形状与布局编译期约束TileDataDst::isRowMajor必须为true即只支持RowMajor行主序布局CPU 侧同样要求TileDst/TileSrc0/TileSrc1均为 RowMajor见 include/pto/cpu/TColExpandOp.hpp 的static_assert。src1 的形状约束src1预期提供每列一个标量即其有效形状必须覆盖C个值通常1 x C而不是与src0相同的R x C从 include/pto/npu/a2a3/TColExpandBinOp.hpp 的实现看当src1有效形状与dst相同时src1eqdst其行步长退化为与dst一致仍只使用每列首个标量参与广播而TCOLEXPANDOP_IMPL还会自动比较src0与dst的有效形状src0eqdst决定以src0还是src1作为主操作数进行发射兼顾了调用方传入顺序的灵活性。布局与分形约束确切的布局/分形约束是目标特定的需参见include/pto/npu/*/TColExpand*.hpp下的后端头文件如 include/pto/npu/a2a3/TColExpandBinOp.hpp、include/pto/npu/a5/TColExpandBinOp.hpp。64 位元素类型A5 特有int64/uint64仅在 Ascend 950PR / Ascend 950DT 上受支持且存在两条关键规则寄存器对模拟A5 没有原生的 64 位向量 ALU指令通过一对 32 位寄存器分别保存每个元素的低 32 位与高 32 位模拟实现。对应实现见 include/pto/npu/a5/TColExpandAdd.hpp在PTO_NPU_ARCH_A5 / PTO_NPU_ARCH_A6下调用Int64BinaryCalcRegsInt64Op::Add, T(dstLow, dstHigh, src0Low, src0High, src1Low, src1High, preg)完成高低字分别参与运算的 64 位加法解交织布局每列标量操作数src1与全尺寸操作数使用相同的解交织de-interleaved布局读取保证高低字配对正确。计算结果为精确的 64 位二进制补码值。Tile 对齐遵循 64 位元素的通用规则RowMajor 的 Tile 要求Cols % 4 0。底层实现解析A2 / A3 后端基于重复步长repeat stride的 vadd在 include/pto/npu/a2a3/TColExpandAdd.hpp 中ColExpandAddOp将运算落到硬件vadd指令上vadd(dst, src0, src1, repeats, 1, 1, 1, 8, 8, 8); // 广播模式src1 的 repeat 步长为 0不前进 vadd(dst, src0, src1, repeats, 1, 1, 1, dstRepeatStride, src0RepeatStride, 0); // 通用步长版本其关键思想是src1的 repeat 步长传0使向量硬件在连续 repeat 中反复读取同一份标量数据从而实现列广播而不需要显式拷贝展开。调度逻辑位于 include/pto/npu/a2a3/TColExpandBinOp.hppNormMode常规模式当 Tile 的Cols ValidCol列无 padding或Rows 1时以ElementsPerRepeat每个 repeat 可承载的元素数为粒度拆分循环分整段与余数两阶段设置连续掩码SetContMaskByDTypeT后发射CountMode计数模式其余情况存在列 padding退化为逐行循环每行通过SetVectorCount(validCol)限定有效列数后单独发射避免 padding 区被误写。发射前还会依据blockSizeElem BLOCK_BYTE_SIZE / sizeof(DType)与elementsPerRepeat REPEAT_BYTE / sizeof(DType)见 include/pto/npu/a2a3/TColExpandBinOp.hpp换算出行步长与 repeat 步长保证不同数据类型下都能对齐硬件最小粒度。A5 后端掩码寄存器与 64 位模拟A5 后端include/pto/npu/a5/TColExpandAdd.hpp使用寄存器张量RegTensor与掩码寄存器vadd(reg_dst, reg_src0, reg_src1, preg, MODE_ZEROING);preg掩码配合MODE_ZEROING处理列尾部无效元素64 位类型则走上述高低字寄存器对模拟路径。CPU 模拟器并行按列广播CPU 侧实现位于 include/pto/cpu/TColExpandOp.hppTColExpand_Op先做类型/布局编译期检查随后通过cpu::parallel_for_1d按列并行每列取出src1的标量值src1Val再对整列逐行执行ElementOpCalT, OP_ADD::apply完成广播加法。该实现同时是单元测试与功能仿真的参考实现保证指令语义在 CPU 上可验证、可调试。使用示例自动模式与手动模式自动模式Auto Mode自动模式下由编译器/运行时负责 Tile 的资源放置与调度用户只表达数据流# Auto mode: compiler/runtime-managed placement and scheduling. %dst pto.tcolexpandadd %src0, %src1 : !pto.tile..., !pto.tile... - !pto.tile...手动模式Manual Mode手动模式下资源必须显式绑定后再发射指令可通过pto.tassign将虚拟操作数绑定到具体存储地址# Manual mode: resources must be bound explicitly before issuing the instruction. # Optional for tile operands: # pto.tassign %arg0, tile(0x1000) # pto.tassign %arg1, tile(0x2000) %dst pto.tcolexpandadd %src0, %src1 : !pto.tile..., !pto.tile... - !pto.tile...PTO 汇编形式%dst tcolexpandadd %src0, %src1 : !pto.tile..., !pto.tile... - !pto.tile... # AS Level 2 (DPS) pto.tcolexpandadd ins(%src0, %src1 : !pto.tile_buf..., !pto.tile_buf...) outs(%dst : !pto.tile_buf...)内核实测示例NPU 测试用例仓库在 A2/A3、A5、kirin9030、kirinDev0000、kirinX90 等多个平台均提供了 TCOLEXPANDADD 的系统测试用例例如 tests/npu/a5/src/st/testcase/tcolexpandadd/tcolexpandadd_kernel.cpptemplate typename T, uint32_t dstRow, uint32_t dstCol, uint32_t src1Row, uint32_t src1Col __global__ AICORE void runCOLEXPANDADD(__gm__ T __out__* out, __gm__ T __in__* src0, __gm__ T __in__* src1) { using DynShapeDim5 Shape1, 1, 1, src1Row, src1Col; using DynStridDim5 pto::Stride1, 1, 1, src1Col, 1; using GlobalData GlobalTensorT, DynShapeDim5, DynStridDim5; using TileData TileTileType::Vec, T, src1Row, src1Col, BLayout::RowMajor, -1, -1; using DstDynShapeDim5 Shape1, 1, 1, dstRow, dstCol; using DstDynStridDim5 pto::Stride1, 1, 1, dstCol, 1; using DstGlobalData GlobalTensorT, DstDynShapeDim5, DstDynStridDim5; using DstTileData TileTileType::Vec, T, dstRow, dstCol, BLayout::RowMajor, -1, -1; DstTileData src0Tile(dstRow, dstCol); TileData src1Tile(src1Row, src1Col); DstTileData dstTile(dstRow, dstCol); TASSIGN(src0Tile, 0x0); TASSIGN(src1Tile, 0x10000); TASSIGN(dstTile, 0x20000); int offset 0; DstGlobalData src0Global(src0 offset); GlobalData src1Global(src1 offset); DstGlobalData dstGlobal(out offset); TLOAD(dstTile, dstGlobal); TLOAD(src0Tile, src0Global); TLOAD(src1Tile, src1Global); #ifndef __PTO_AUTO__ set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); #endif TCOLEXPANDADD(dstTile, src0Tile, src1Tile); #ifndef __PTO_AUTO__ set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); #endif TSTORE(dstGlobal, dstTile); #ifndef __PTO_AUTO__ set_flag(PIPE_MTE3, PIPE_S, EVENT_ID0); wait_flag(PIPE_MTE3, PIPE_S, EVENT_ID0); pipe_barrier(PIPE_ALL); #endif out dstGlobal.data(); }该用例展示了完整的装载TLOAD→ 运算TCOLEXPANDADD→ 存储TSTORE流水TASSIGN显式绑定三个 Tile 的存储基址0x0/0x10000/0x20000src1Tile的形状为src1Row x src1Col测试中取1 x C即每列一个标量dstTile与src0Tile形状为dstRow x dstCol。__PTO_AUTO__宏区分自动模式与手动模式非自动模式下需用set_flag / wait_flag显式同步 MTE2 → V → MTE3 → S 管道事件。测试用例覆盖的形状与数据类型组合包括float16x128、32x32、1x128aclFloat16half4x256、10x64后者验证非对齐列数下的掩码/计数路径int32_t16x32int16_t16x64int64_t、uint64_t16x32A5 的 64 位寄存器对模拟路径。数据生成脚本见 tests/npu/a5/src/st/testcase/tcolexpandadd/gen_data.pyA2/A3 对应用例见 tests/npu/a2a3/src/st/testcase/tcolexpandadd/CPU 仿真测试见 tests/cpu/st/testcase/tcolexpandop/。进一步阅读同族指令参考TCOLEXPAND、TCOLEXPANDDIV、TCOLEXPANDMUL、TCOLEXPANDSUB、TCOLEXPANDMAX、TCOLEXPANDMIN、TCOLEXPANDEXPDIF以及与行广播对应的 TROWEXPANDADD虚拟指令集整体说明PTO-Virtual-ISA-Manual.md 与 PTOISA.md后端实现include/pto/npu/a2a3/TColExpandBinOp.hpp、include/pto/npu/a2a3/TColExpandAdd.hpp、include/pto/npu/a5/TColExpandAdd.hppCPU 参考实现include/pto/cpu/TColExpandOp.hpp编程入门教程docs/coding/tutorials/README.md 与 docs/coding/ProgrammingModel.md。赞分享人工智能指令集算子库CANNAscend【免费下载链接】pto-isaParallel Tile Operation (PTO) is a virtual instruction set architecture designed by Ascend CANN, focusing on tile-level operations. This repository offers high-performance, cross-platform tile operations across Ascend platforms.项目地址https://gitcode.com/cann/pto-isa点击查看免费下载相关推荐PTO ISA 指令详解TCOLEXPANDMUL 列广播乘法Column-wise Broadcast MultiplyPTO ISA 指令详解TCOLEXPANDMUL 列广播乘法Column wise Broadcast Multiply 导读 TCOLEXPANDMU人工智能指令集算子库CANNAscendPTO-ISA TCOLEXPANDDIV 指令详解列广播除法Column-wise Broadcast Divide的语义、汇编与 C 编程指南PTO ISA TCOLEXPANDDIV 指令详解列广播除法Column wise Broadcast Divide的语义、汇编与 C 编程指南 T人工智能指令集算子库CANNAscendPTO-ISA TCOLEXPANDDIV 指令详解列广播除法Column-wise Broadcast Divide的语义、编程接口与多平台实现PTO ISA TCOLEXPANDDIV 指令详解列广播除法Column wise Broadcast Divide的语义、编程接口与多平台实现 导读人工智能指令集算子库CANNAscend创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考