人工智能指令集算子库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点击查看免费下载TMULATile Multiply-Add是 CANN PTOParallel Tile Operation虚拟指令集中面向向量 Tile 的三元逐元素运算指令将乘与累加融合为单条指令执行数学语义为dst src0 * src1 dst。本文以 docs/isa/TMULA.md及其中文版 docs/isa/TMULA_zh.md为骨架结合仓库中的 NPU 后端实现a2a3/a5与 CPU/NPU 测试用例完整讲解 TMULA 的数学语义、两级汇编语法、C 内建接口、使用约束、底层硬件指令映射与端到端测试方法。读完本文你将能够在 PTO 编程模型中正确、高效地使用 TMULA 完成乘加融合运算并理解它在不同 Ascend 平台上的实现差异与精度语义。一、指令概述为什么需要一条乘加融合指令TMULA 是一条三元逐元素elementwise运算指令一次指令执行完成两个动作将src0与src1逐元素相乘将乘积与dst的当前值累加并写回dst自身。即dst src0 * src1 dstdst同时承担累加器角色读改写语义。这在归一化、加权求和、矩阵乘的累加链、LayerNorm/Softmax 后处理等场景中非常常见。相比先TMUL再TADD的两条指令方案TMULA 将乘法与加法融合为一条指令减少了指令发射次数与中间结果的搬运是典型的计算融合优化手段。数学语义对有效区域valid region内的每个元素(i, j)$$ \mathrm{dst}{i,j} \mathrm{src0}{i,j} \cdot \mathrm{src1}{i,j} \mathrm{dst}{i,j} $$其中dst、src0、src1均为同尺寸的向量 Tile运算完全逐元素进行不涉及跨行跨列的归约。二、汇编语法同步形式与两级抽象PTO ISA 为 TMULA 提供了多种表达层次从面向人类的同步伪汇编到面向编译器/后端的 SSA 与 DPS 形式。2.1 同步汇编形式%dst tmula %src0, %src1 : !pto.tile...这是文档中的同步形式Synchronous form语义直观src0 * src1 dst结果写回dst。2.2 AS Level 1SSA 形式%dst pto.tmula %src0, %src1 : (!pto.tile..., !pto.tile...) - !pto.tile...AS Level 1 采用纯 SSA 风格显式给出输入操作数类型与返回值类型两个!pto.tile...输入一个!pto.tile...输出。2.3 AS Level 2DPS 形式pto.tmula ins(%src0, %src1 : !pto.tile_buf..., !pto.tile_buf...) outs(%dst : !pto.tile_buf...)AS Level 2 采用 DPSDestination-Passing Style描述输入通过ins()声明输出通过outs()声明操作数类型为物理缓冲类型!pto.tile_buf...更贴近硬件资源视图。2.4 自动模式与手动模式PTO 同时支持两种资源管理与调度方式# 自动模式由编译器/运行时负责资源放置与调度。 %dst pto.tmula %src0, %src1 : (!pto.tile..., !pto.tile...) - !pto.tile...# 手动模式先显式绑定资源再发射指令。 # 可选当该指令包含 tile 操作数时 # pto.tassign %arg0, tile(0x1000) # pto.tassign %arg1, tile(0x2000) %dst pto.tmula %src0, %src1 : (!pto.tile..., !pto.tile...) - !pto.tile...手动模式下开发者需要先用pto.tassign将逻辑 tile 操作数显式绑定到物理 tile 地址如tile(0x1000)再发射指令自动模式下则完全交给编译器/运行时完成资源放置与调度。这一模式切换在测试代码中体现为__PTO_AUTO__宏详见下文测试用例章节。三、C 内建接口与头文件TMULA 的 C 内建接口声明于公共指令头文件 include/pto/common/pto_instr.hpp第 1712-1718 行template typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents PTO_INST RecordEvent TMULA(TileDataDst dst, TileDataSrc0 src0, TileDataSrc1 src1, WaitEvents ...events);接口要点泛型模板参数TileDataDst、TileDataSrc0、TileDataSrc1分别对应输出与两个输入的 Tile 数据类型WaitEvents...为可变参数用于向指令传递依赖事件等待前序指令完成。返回RecordEvent指令发射后会返回一个记录事件可继续传递给后续指令建立依赖链。实现分发从源码结构看TMULA在等待事件detail::PtoWaitEvents(events...)之后通过MAP_INSTR_IMPL(TMULA, dst, src0, src1)宏按平台CPU_SIM / NPU a2a3 / NPU a5 等分发到对应的TMULA_IMPL实现。公共包含头实际编程时包含pto/pto-inst.hpp即可内部声明位于pto/common/pto_instr.hpp中文文档明确标注了这一头文件层级关系。四、使用约束与合法性检查TMULA 的使用有一组严格的编译期与运行期约束违反时将触发static_assert或PTO_ASSERT报错。4.1 数据类型编译期检查TileData::DType必须是half、float中文文档或源码实现中扩展的float16_t、float32_t之一从 a2a3/a5 后端的TMulaCheck实现看编译期断言为static_assert( std::is_same_vT, half || std::is_same_vT, float16_t || std::is_same_vT, float || std::is_same_vT, float32_t, Fix: TMULA has invalid data type.);dst、src0、src1三者数据类型必须一致static_assert( std::is_same_vT, typename TileDataSrc0::DType std::is_same_vT, typename TileDataSrc1::DType, Fix: TMULA the data type of dst must be consistent with of src0 and src1.);4.2 布局与位置约束行主序TileData::isRowMajor必须为真三个操作数均要求行主序布局static_assert( TileDataDst::isRowMajor TileDataSrc0::isRowMajor TileDataSrc1::isRowMajor, Fix: TMULA only support row major layout.);向量 Tile 位置中文文档补充的通用约束Tile 位置必须是向量即TileData::Loc TileType::Vec。4.3 有效区域valid region约束静态有效边界TileData::ValidRow TileData::Rows且TileData::ValidCol TileData::Cols运行时形状一致性dst、src0、src1的有效行列数必须完全相同否则运行期断言报错PTO_ASSERT( src0.GetValidRow() validRows src0.GetValidCol() validCols src1.GetValidRow() validRows src1.GetValidCol() validCols, Fix: TMULA input tile src0 valid shape mismatch with output tile dst shape.);运算迭代范围以dst.GetValidRow()/dst.GetValidCol()为准。4.4half精度语义CPU_SIMCPU_SIM 模拟器中half路径的舍入规则为乘积先舍入为half再执行累加累加结果再次舍入为half。这一中间舍入语义对精度敏感场景如低精度训练推理至关重要在 tests/cpu/st/testcase/tmula/gen_data.py 等测试脚本生成 golden 数据时也需要与之保持一致。五、底层实现从内建接口到硬件指令TMULA 在 NPU 端按 Ascend 平台分别落地于 include/pto/npu/a2a3/TMula.hpp 与 include/pto/npu/a5/TMula.hpp。两者共享相同的对外语义但底层指令映射不同。5.1 a2a3 平台基于vmla的双操作数指令形态a2a3 后端的MulaOp将 TMULA 映射为vmla向量指令template typename T struct MulaOp { PTO_INTERNAL static void BinInstr(__ubuf__ T* dst, __ubuf__ T* src0, __ubuf__ T* src1, uint8_t repeats) { vmla(dst, src0, src1, repeats, 1, 1, 1, 8, 8, 8); } PTO_INTERNAL static void BinInstr( __ubuf__ T* dst, __ubuf__ T* src0, __ubuf__ T* src1, uint8_t repeats, uint8_t dstRepeatStride, uint8_t src0RepeatStride, uint8_t src1RepeatStride) { vmla(dst, src0, src1, repeats, 1, 1, 1, dstRepeatStride, src0RepeatStride, src1RepeatStride); } };可以看到默认的 repeat 间步长为1, 1, 1连续 repeatblock 间步长默认8, 8, 8重载版本允许显式指定三个操作数各自的 repeat 步长用于非连续布局。TMULA_IMPL会依据 Tile 的RowStride编译期常量推导dstRowStride / src0RowStride / src1RowStride三者相同时走统一的BinaryInstr路径不同时走带分别步长的BinaryInstr重载。块大小与单次 repeat 处理元素数由字节常量推导constexpr unsigned blockSizeElem BLOCK_BYTE_SIZE / sizeof(T); // 每个 block 的元素数 constexpr unsigned elementsPerRepeat REPEAT_BYTE / sizeof(T); // 每次 repeat 的元素数5.2 a5 平台half走vmul vadd序列其余走vmulaa5 后端采用三元指令形态TernInstr且对half与其它类型做了不同的指令映射template typename T struct MulaOp { PTO_INTERNAL static void TernInstr( RegTensorT reg_dst, RegTensorT reg_src0, RegTensorT reg_src1, MaskReg preg) { if constexpr (std::is_same_vT, half) { vmul(reg_src0, reg_src0, reg_src1, preg, MODE_ZEROING); vadd(reg_dst, reg_dst, reg_src0, preg, MODE_ZEROING); } else { vmula(reg_dst, reg_src0, reg_src1, preg, MODE_ZEROING); } } };要点非half类型如float直接使用vmula单条指令完成乘加half类型则拆为vmul将src0 * src1写入src0寄存器后再vadd累加到dst中间结果始终保存在寄存器中且均带MODE_ZEROING掩码模式——从源码结构看这与 CPU_SIM 中乘积先舍入为 half 再累加的语义相呼应即 a5 的half路径在两条指令间自然产生一次 half 精度中间值a5 的elementsPerRepeat CCE_VL / sizeof(T)即每次 repeat 处理一个向量寄存器CCE_VL 字节的元素量。这种同一 ISA 语义、不同平台不同指令序列的设计正是 PTO 虚拟指令集屏蔽平台差异、提供统一编程视图的典型体现。六、编程示例与端到端测试实战6.1 最小使用示例文档给出的最小示例16×16 的 float 向量 Tile#include pto/pto-inst.hpp using namespace pto; void example() { using TileT TileTileType::Vec, float, 16, 16; TileT a, b, out; TMULA(out, a, b); }其中TileTileType::Vec, float, 16, 16声明了一个 16 行 16 列、位于向量Vec位置的 float TileTMULA(out, a, b)执行out a * b out。6.2 完整 CPU 测试内核TMULA 的典型使用链路仓库 CPU 侧测试 tests/cpu/st/testcase/tmula/tmula_kernel.cpp 展示了 TMULA 在完整 kernel 中的标准使用链路TASSIGN 绑定资源 → TLOAD 加载数据 → TMULA 计算 → TSTORE 写回。template typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows, int vCols __global__ AICORE void runTMULA(__gm__ T __out__* out, __gm__ T __in__* src0, __gm__ T __in__* src1) { using DynShape pto::Shape-1, -1, -1, -1, -1; using DynStride pto::Stride-1, -1, -1, -1, -1; using GlobalData GlobalTensorT, DynShape, DynStride; GlobalData dstGlobal(out, pto::Shape(1, 1, 1, vRows, vCols), pto::Stride(dstTileH * dstTileW, dstTileH * dstTileW, dstTileH * dstTileW, dstTileW, 1)); // src0Global / src1Global 结构相同略 using TileDataDst TileTileType::Vec, T, dstTileH, dstTileW, BLayout::RowMajor, -1, -1; // ... TASSIGN(src0Tile, 0x0); TASSIGN(src1Tile, src0TileH * src0TileW * sizeof(T)); TASSIGN(dstTile, (src0TileH * src0TileW src1TileH * src1TileW) * sizeof(T)); TLOAD(src0Tile, src0Global); TLOAD(src1Tile, src1Global); TLOAD(dstTile, dstGlobal); #ifndef __PTO_AUTO__ set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); #endif TMULATileDataDst, TileDataSrc0, TileDataSrc1(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); out dstGlobal.data(); }该测试内核同时展示了 PTO 的两套调度方式自动模式定义了__PTO_AUTO__直接发射TMULA由编译器/运行时自动处理 MTE2→V→MTE3 流水线依赖手动模式未定义__PTO_AUTO__通过set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0)/wait_flag(...)显式等待数据搬入MTE2完成后再计算计算完成后再次用事件同步保证TSTOREMTE3读到的数据是计算后的结果。注意这里dst也必须先TLOAD加载——因为 TMULA 的语义是累加到dst的原值上dst 的初值参与运算。测试还覆盖了多种形状组合模板实例化template void LaunchTMULAfloat, 64, 64, 64, 64, 64, 64, 64, 64(...); template void LaunchTMULAfloat, 32, 128, 32, 192, 32, 256, 32, 127(...); template void LaunchTMULAaclFloat16, 64, 64, 64, 64, 64, 64, 64, 64(...); template void LaunchTMULAaclFloat16, 32, 128, 32, 192, 32, 256, 32, 127(...); template void LaunchTMULAaclFloat16, 1, 16384, 1, 16384, 1, 16384, 1, 16384(...);覆盖了 float/half、正方形与非正方形 Tile、以及1 行 16384 列的极窄长条形状用于验证 TMULA 在不同有效形状下的正确性。6.3 NPU 侧 ST 测试golden 数据对比验证NPU 侧的系统测试用例位于 tests/npu/a2a3/src/st/testcase/tmula/main.cpp其验证流程为通过aclInit/aclrtSetDevice/aclrtCreateStream初始化 Ascend 运行时aclrtMallocHost/aclrtMalloc分配主机与设备内存读入input_dst.bindst 初值、input0.bin、input1.bin三个输入拷贝到设备后LaunchTMULA启动 kernelaclrtSynchronizeStream同步写回output.bin与golden.bin逐元素对比ResultCmpT(golden, devFinal, 0.001f)误差阈值 0.001EXPECT_TRUE(ret)判定用例通过。对应测试用例TEST_F(TMULATest, case_float_64x64_64x64_64x64_64x64) { test_TMULAfloat, 64, 64, 64, 64, 64, 64, 64, 64(); } TEST_F(TMULATest, case_float_32x128_32x192_32x256_32x127) { test_TMULAfloat, 32, 128, 32, 192, 32, 256, 32, 127(); } TEST_F(TMULATest, case_half_64x64_64x64_64x64_64x64) { test_TMULAaclFloat16, 64, 64, 64, 64, 64, 64, 64, 64(); } TEST_F(TMULATest, case_half_32x128_32x192_32x256_32x127) { ... }其中用例名case_half_32x128_32x192_32x256_32x127依次编码了 dst/src0/src1 的 Tile 尺寸与有效区域vRows32, vCols127例如32x128表示 dst Tile 为 32 行 128 列、有效列 127——用于验证有效区域小于物理 Tile 尺寸的合法场景。这些用例通过 tests/run_st.sh 脚本统一驱动运行。七、使用要点小结语义是融合累加dst同时是输出与累加器使用前必须保证dst中已载入需要累加的初值如归约结果或上一次累加结果。类型与布局约束严格仅支持half/float及float16_t/float32_t必须行主序、TileType::Vec位置三个操作数类型一致违反约束会触发编译期static_assert。有效区域必须一致dst/src0/src1的有效行列数必须完全相同运算范围以dst.GetValidRow()/GetValidCol()为准。half精度注意CPU_SIM 与 a5 平台的half路径均在乘积后先舍入为half再累加属于语义的一部分做精度对比如 golden 生成时需对齐该舍入行为。平台指令映射不同a2a3 用vmla双操作数形态a5 非 half 用vmula、half 用vmul vadd序列——上层代码无需感知这正是 PTO 虚拟指令集的价值所在。流水线依赖按模式处理自动模式__PTO_AUTO__由编译器管理依赖手动模式下需用set_flag/wait_flag在TLOADMTE2与TMULAV、TMULA与TSTOREMTE3之间显式建立事件同步。若需深入了解 TMULA 所在算术类指令族或 Tile 编程模型可进一步阅读 docs/isa/README.md、docs/menu/arithmetic_zh.md 以及 docs/coding/ProgrammingModel.md 与 docs/coding/ProgrammingModel_zh.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点击查看免费下载相关推荐CANN pto-isa TMADD 指令详解src0*dstsrc1 乘加融合的 Tile 级向量运算CANN pto isa TMADD 指令详解src0 dstsrc1 乘加融合的 Tile 级向量运算 TMADD 是 CANN pto isaPara人工智能指令集算子库CANNAscendPyPTO 向量函数 vf.mul_dst_add 详解dst×src0src1 乘加融合FMA寄存器运算PyPTO 向量函数 vf.mul_dst_add 详解dst×src0src1 乘加融合FMA寄存器运算 本文围绕 PyPTOParallel Te人工智能编译器模型编译高性能计算深度学习CANNPTO-ISA 指令详解TCOLEXPANDEXPDIF 列指数差运算exp(src0 - src1)PTO ISA 指令详解TCOLEXPANDEXPDIF 列指数差运算exp src0 src1 TCOLEXPANDEXPDIF 是 CANN Par人工智能指令集算子库CANNAscend创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考