PTO TPARTMIN 指令全解析CANN pto-isa 中基于有效区域valid region的逐元素最小值选择【免费下载链接】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本篇技术指南以 TPARTMIN 指令规范及对应的 中文版为核心主体结合 CANN pto-isa 开源仓库中的 NPUA2A3/A5实现、CPU 参考实现与 ST 测试用例系统讲解 PTO 虚拟指令集Parallel Tile Operation ISA中TPARTMIN指令的语义、汇编语法、C 内建接口、约束条件与底层实现原理。读完本文你将掌握如何在 Auto/Manual 两种模式下正确使用TPARTMIN完成两个 Tile 之间带部分有效区域partial validity的逐元素取最小值运算并理解其与普通TMIN指令、以及 A2A3/A5 两个平台实现差异的关键点。指令概述什么是 TPARTMINTPARTMINTile Partial Minimum是 CANN pto-isa 提供的一类部分有效区域二元运算指令族成员同族指令还包括TPARTADD、TPARTMUL、TPARTARGMAX、TPARTARGMIN等见 docs/isa 目录。它与普通逐元素最小值指令TMIN的核心区别在于TPARTMIN的结果域由目标 Tiledst的有效区域valid region决定而不是要求三个 Tile 的形状与有效区域完全一致。它允许两个输入源 Tile 的有效区域与目标 Tile 存在部分匹配关系并在有效区域重叠与不重叠的位置采取不同的取值策略在目标有效区域内的某个元素位置(i, j)上若src0与src1都有效则结果取min(src0, src1)若只有src0在该位置有效则结果直接复制src0的值若只有src1在该位置有效则结果直接复制src1的值。也就是说TPARTMIN本质上是以dst的有效区域为模板在两个输入之间做逐元素最小值选择并用有效输入的值补齐另一输入无效的区域。这使其非常适合处理形状不对齐如 padding、mask、边界裁剪的 Tile 数据合并场景。上图来自指令规范文档展示了TPARTMIN的 Tile 级运算示意三个 Tile 的有效区域关系与结果生成过程。数学语义对目标有效区域内的每个元素(i, j)指令的计算语义可形式化描述为$$ \mathrm{dst}{i,j} \begin{cases} \min(\mathrm{src0}{i,j}, \mathrm{src1}{i,j}) \text{if both inputs are defined at } (i,j) \ \mathrm{src0}{i,j} \text{if only src0 is defined at } (i,j) \ \mathrm{src1}_{i,j} \text{if only src1 is defined at } (i,j) \end{cases} $$这里defined有效指该位置落在对应 Tile 的有效区域内。注意语义描述的是目标有效区域内的行为对于目标有效区域之外的位置其结果不被定义也不在本指令的职责范围内。其余未显式列出的有效区域组合例如两个输入都无效、或输入有效但目标无效等其行为由具体实现定义implementation-defined。汇编语法从同步形式到两级 AS 层级同步形式PTO Assembly FormPTO 汇编层面的指令写法为%dst tpartmin %src0, %src1 : !pto.tile... - !pto.tile...它对应同步语义即该指令发射后需要等待其完成才能继续后续依赖它的操作。AS Level 1SSA 形式在虚拟指令集的第一级抽象SSA、!pto.tile类型中操作数以 SSA 值的形式出现由编译器/运行时负责资源分配与调度%dst pto.tpartmin %src0, %src1 : (!pto.tile..., !pto.tile...) - !pto.tile...AS Level 2DPS 形式在第二级抽象DPS即destructive/指名目的地形式、!pto.tile_buf类型中输入与输出显式分离使用ins(...)与outs(...)子句声明pto.tpartmin ins(%src0, %src1 : !pto.tile_buf..., !pto.tile_buf...) outs(%dst : !pto.tile_buf...)这三级形式之间是递进的降级关系SSA 形式经资源绑定后降级为 DPS 形式再最终映射到硬件向量指令详见下文底层实现一节。实际书写中...位置会被具体的 Tile 形状、布局参数如!pto.tilef32, 16, 16替代。C 内建接口IntrinsicTPARTMIN的 C 内建接口声明于 include/pto/common/pto_instr.hpp公共包含头为pto/pto-inst.hpp内部声明位于pto/common/pto_instr.hpptemplate typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents PTO_INST RecordEvent TPARTMIN(TileDataDst dst, TileDataSrc0 src0, TileDataSrc1 src1, WaitEvents ... events);参数说明参数含义dst目标 Tile其有效区域定义结果的计算范围src0、src1两个输入源 Tileevents可变参数可选的WaitEvents用于表达跨流水如 MTE2→V 向量流水的同步依赖返回值为RecordEvent可用于事件链式同步。在 pto_instr.hpp 中该接口通过MAP_INSTR_IMPL(TPARTMIN, dst, src0, src1)宏按平台NPU A2A3 / A5 / CPU分发到对应实现。约束条件通用约束各平台通用dst、src0、src1的元素类型必须一致目标有效区域定义结果域指令只在该区域内产生结果对目标有效区域内的每个元素两个输入均有效则取逐元素最小值仅一个输入有效则复制该输入的值若dst的有效区域为零行数或列数为 0指令直接返回不做任何计算部分有效区域支持模式要求至少有一个源 Tile 的有效区域与dst完全一致另一个源 Tile 的有效区域在两个维度上都不能超过dst上述列举之外的有效区域组合其行为均由具体实现定义。A2A3 实现检查Atlas A2/A3 训练与推理系列产品支持的元素类型int32_t、int、int16_t、half、float16_t、float、float32_tdst、src0、src1必须全部为行主序isRowMajor即BLayout::RowMajor不支持 BLayout 等其他布局。A5 实现检查Ascend 950PR / Ascend 950DT支持的元素类型int8_t、uint8_t、int16_t、uint16_t、int32_t、uint32_t、int64_t、uint64_t、half、bfloat16_t、float。对比可见A5 平台的类型覆盖面明显更广额外支持 8/16/64 位整型与bfloat16_t而 A2A3 平台对布局有硬性要求必须行主序A5 则未在规范中列出布局限制。底层实现原理从虚拟指令到硬件向量指令A2A3 实现基于vmin向量指令的分段执行A2A3 平台的实现位于 include/pto/npu/a2a3/TPartMin.hpp其核心逻辑分两层运算算子PartMinOp直接下发昇腾向量指令vmin(dst, src0, src1, repeats, ...)其中repeats1后面三个参数1,1,1为重复块的 mask 步长再后面的三个参数是 dst/src0/src1 的 repeat 步长include/pto/npu/a2a3/TPartMin.hppPTO_INTERNAL static void PartInstr(__ubuf__ T* dst, __ubuf__ T* src0, __ubuf__ T* src1, uint8_t repeats) { vmin(dst, src0, src1, repeats, 1, 1, 1, 8, 8, 8); }有效区域匹配判定TPartMin检查src0/src1的有效区域是否与dst完全一致condSrc0EqDst/condSrc1EqDst二者必须至少有一个成立否则触发编译期断言PTO_ASSERT(Fix: TPARTMIN At most one entry in the valid-rows and valid-cols of src0 and src1 is smaller than dst.)。若src0与dst一致则以src0为基准调用TPartInstr若只有src1与dst一致则交换 src0/src1 后以相同方式处理include/pto/npu/a2a3/TPartMin.hpp。执行前检查TPARTMIN_IMPL通过static_assert检查三者的DType一致、类型在 A2A3 支持列表中、且全部为isRowMajor随后通过GetValidRow()/GetValidCol()读取有效区域若dst有效区域为零则提前返回include/pto/npu/a2a3/TPartMin.hpp。真正处理部分有效区域的分发逻辑在通用工具 include/pto/npu/a2a3/TPartOp.hpp 的TPartInstr中它按src1相对dst的有效区域关系分三种路径src1行数小于dst列相等先在重叠的行区域执行TPartOps向量vmin运算再用TPartCopyInstr把src0剩余行复制到dst——对应仅一个输入有效则复制该输入的语义src1列数小于dst行不超过先把src0全区域复制到dst再插入pipe_barrier(PIPE_V)流水屏障在重叠列区域以dst自身为第三个操作数执行vmin覆盖式更新src1与dst完全一致直接对全有效区域执行TPartOps向量运算。TPartOps内部还会根据行/列步长是否超过REPEAT_STRIDE_MAX、行数是否大于REPEAT_MAX等条件自动在count 计数模式逐行循环与norm 归一模式按 repeat 批量发射之间选择更优的执行路径兼顾正确性与性能。A5 实现掩码向量指令与 64 位整数特殊路径A5 平台的实现位于 include/pto/npu/a5/TPartMin.hpp风格与 A2A3 不同运算算子TPartMinOp::BinInstr使用带谓词寄存器MaskReg preg的向量指令vmin(dst, src0, src1, preg, MODE_ZEROING)通过掩码直接表达有效区域include/pto/npu/a5/TPartMin.hpp对64 位整型int64_t/uint64_t走专门的Int64PartInt64Op::Min, ...路径位于 include/pto/npu/a5/TPartBinOps.hpp该路径将 64 位元素拆分为高/低 32 位寄存器对如Int64PartCalcRegs、Int64PartSameStrideRepeat、Int64PartCopyRow等函数所示进行运算并对重叠行/尾行分别处理其余类型统一走TPARTOP_IMPLTPartMinOpT, ...通用部分二元运算框架include/pto/npu/a5/TPartMin.hpp。CPU 参考实现CPU 参考实现位于 include/pto/cpu/TPartMin.hpp语义最直观对有效区域内的每个元素逐一调用std::min(src0[Src0Offset], src1[Src1Offset])并写入dstinclude/pto/cpu/TPartMin.hpp可作为理解指令语义、以及跨平台结果对齐的黄金参照。编程示例Auto 与 Manual 两种模式Auto 模式自动资源管理#include pto/pto-inst.hpp using namespace pto; void example_auto() { using TileT TileTileType::Vec, float, 16, 16; TileT src0, src1, dst; TPARTMIN(dst, src0, src1); }Auto 模式下编译器/运行时负责 Tile 的放置placement与调度scheduling开发者只需声明 16×16 的float向量 Tile 并调用指令即可。对应的汇编形态为# Auto mode: compiler/runtime-managed placement and scheduling. %dst pto.tpartmin %src0, %src1 : (!pto.tile..., !pto.tile...) - !pto.tile...Manual 模式显式资源绑定#include pto/pto-inst.hpp using namespace pto; void example_manual() { using TileT TileTileType::Vec, float, 16, 16; TileT src0, src1, dst; TASSIGN(src0, 0x1000); TASSIGN(src1, 0x2000); TASSIGN(dst, 0x3000); TPARTMIN(dst, src0, src1); }Manual 模式下开发者需先用TASSIGN将三个 Tile 显式绑定到具体的 tile buffer 地址如0x1000/0x2000/0x3000再发射TPARTMIN。对应汇编# 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.tpartmin %src0, %src1 : (!pto.tile..., !pto.tile...) - !pto.tile...完整 Kernel 级用法结合 TLOAD/TSTORE 与流水同步仓库 ST 测试给出了在真实 kernel 中组合使用TPARTMIN的完整范式见 tests/npu/a2a3/src/st/testcase/tpartmin/tpartmin_kernel.cpp#include pto/pto-inst.hpp using namespace pto; template typename T, int kGRowsD_, int kGColsD_, ... __global__ AICORE void runTPartMin(__gm__ T __out__* out, __gm__ T __in__* src0, __gm__ T __in__* src1) { using GlobalData GlobalTensorT, pto::Shape-1, -1, -1, -1, -1, pto::Stride-1, -1, -1, -1, -1; using dstTileData TileTileType::Vec, T, kTRowsD_, kTColsD_, BLayout::RowMajor, -1, -1; using src0TileData TileTileType::Vec, T, kTRowsS0_, kTColsS0_, BLayout::RowMajor, -1, -1; using src1TileData TileTileType::Vec, T, kTRowsS1_, kTColsS1_, BLayout::RowMajor, -1, -1; dstTileData dstTile(kGRowsD_, kGColsD_); src0TileData src0Tile(kGRowsS0_, kGColsS0_); src1TileData src1Tile(kGRowsS1_, kGColsS1_); // ... TASSIGN 绑定 tile buffer构造 GlobalData 全局视图 ... TLOAD(src0Tile, src0Global); TLOAD(src1Tile, src1Global); #ifndef __PTO_AUTO__ set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); // MTE2(搬运完成) - V(向量) 事件同步 wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); #endif TPARTMINdstTileData, src0TileData, src1TileData(dstTile, src0Tile, src1Tile); #ifndef __PTO_AUTO__ set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); // V(向量完成) - MTE3(搬出) 事件同步 wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); #endif TSTORE(dstGlobal, dstTile); }该示例展示了四个关键点使用TileTileType::Vec, T, Rows, Cols, BLayout::RowMajor, -1, -1声明行主序向量 Tile并用-1表示动态形状维度运行时以全局形状构造用TLOAD把全局内存__gm__数据搬入 Tile buffer 后再执行TPARTMIN最后用TSTORE将结果搬回全局内存跨流水同步手动模式下需用set_flag/wait_flag在 MTE2搬入→ V向量计算→ MTE3搬出之间插入事件同步保证数据就绪与结果可见Auto 模式定义了__PTO_AUTO__下则由编译器自动插入同步同一份 kernel 通过模板参数控制三个 Tile 的全局/ Tile 形状测试框架据此枚举不同的有效区域组合如 src1 行列数小于 dst、等于 dst 等。指令调用与分发入口从调用角度看TPARTMIN(dst, src0, src1)的调用链为pto_instr.hpp中的公开模板 →MAP_INSTR_IMPL平台分发 → 各平台的TPARTMIN_IMPLA2A3 见 include/pto/npu/a2a3/TPartMin.hppA5 见 include/pto/npu/a5/TPartMin.hppCPU 见 include/pto/cpu/TPartMin.hpp→ 硬件向量指令vmin。测试与验证仓库为TPARTMIN提供了跨平台、跨形状组合的 STSoftware Test测试用例覆盖 CPU 与多个 NPU 平台CPU tests/cpu/st/testcase/tpartmin含main.cpp、tpartmin_kernel.cpp、gen_data.pyNPU A2A3 tests/npu/a2a3/src/st/testcase/tpartminNPU A5 tests/npu/a5/src/st/testcase/tpartmin其他平台kirin9030 / kirinDev0000 / kirinX90同样包含同名用例目录。各测试目录结构一致gen_data.py负责生成输入/期望输出数据涵盖多种有效区域组合与数据类型tpartmin_kernel.cpp为被测 kernelmain.cpp负责 host 侧内存申请、数据搬移、kernel 启动与结果比对并统一登记在平台级 CMakeLists.txt 中。运行整个测试集可参考 tests/run_st.sh。这种数据生成器 kernel host 校验的结构使得TPARTMIN的部分有效区域语义可以在 CPU 参考实现与 NPU 实现之间做交叉验证确保各平台行为一致。总结TPARTMIN是 PTO 指令集中部分有效区域二元运算的代表性指令语义上它以dst有效区域为结果域重叠区取min单输入有效区做复制规则清晰且贴近实际算子中的 padding/mask 合并需求形态上从 PTO 汇编同步形式、AS Level 1SSA到 AS Level 2DPS再到 C intrinsic提供了一致而完备的编程接口实现上A2A3 通过TPartOp框架把部分匹配拆解为重叠区vmin 非重叠区复制A5 借助掩码向量指令与 64 位整型专用路径实现更宽的类型覆盖CPU 参考实现则用std::min给出最直观的语义基准使用上Auto 模式全自动管理资源Manual 模式需结合TASSIGN与set_flag/wait_flag完成资源绑定与流水同步仓库 ST 测试为两种模式及各类有效区域组合提供了可直接参考的完整代码。如需进一步了解同族指令如TPARTADD、TPARTMUL、TPARTARGMAX可继续查阅 docs/isa 目录下的对应规范文档指令总览与索引可参考 docs/isa/README.md。【免费下载链接】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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考