PTO TMINS 指令详解：Ascend Tile 与标量逐元素最小值运算的数学语义、约束与实战用法

PTO TMINS 指令详解：Ascend Tile 与标量逐元素最小值运算的数学语义、约束与实战用法

【免费下载链接】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

导读

TMINS 是 CANN pto-isa（Parallel Tile Operation，面向昇腾平台的 Tile 级虚拟指令集）中用于执行Tile 与标量逐元素最小值运算的核心向量指令：它以整个 Tile 的每个元素为操作对象，与一个标量值逐一比较并取较小者写入目标 Tile。本文以 TMINS_zh.md 为骨架，结合仓库内 NPU（A2/A3、Ascend 950PR/950DT）与 CPU 模拟器上的源码实现及完整测试用例，系统讲解 TMINS 的数学语义、三级汇编形式、C++ 内建接口、平台约束、自动/手动两种编程模式，并给出可复制运行的实战示例。读完本文，你将能够在 PTO 编程框架中正确声明与调用 TMINS，理解其有效区域（valid region）迭代规则与底层vmins指令的映射关系，并能参照测试用例编写自己的标量归约类算子。

指令一览

TMINS 属于 PTO 指令集中 Tile 与标量的二元运算族（S 后缀表示 scalar，即标量操作数），与 TMIN（Tile 与 Tile 逐元素最小值）互为补充，也与 TADDS、TSUBS、TMULS、TDIVS、TMAXS 等标量广播类指令共享同一套编程模型与硬件流水线。

数学语义

TMINS 对 Tile 的**有效区域（valid region）**内每个元素执行一次min运算：对有效区域中的每个元素(i, j)，其目标值为源 Tile 对应元素与标量值两者中的较小者：

$$ \mathrm{dst}{i,j} = \min(\mathrm{src}{i,j}, \mathrm{scalar}) $$

这里scalar是标量操作数，会被隐式广播到每个元素位置参与比较。需要注意的是，TMINS 与 TMIN 的区别在于第二操作数：TMIN 比较的是两个 Tile 的逐元素值（dst[i][j] = min(src0[i][j], src1[i][j])，见 TMIN_zh.md），而 TMINS 中所有元素共享同一个标量阈值。

典型应用场景包括：数据裁剪（clamp 上界）、激活函数实现中的上界约束（如min(x, alpha)）、量化/归一化中的阈值截断等需要将整个 Tile 与一个固定常数比较的运算。

汇编语法：三级抽象层次

TMINS 在 PTO 指令集中同时存在三种汇编形式，分别对应不同的编译/调度阶段。

同步形式（PTO 汇编）

最直接的指令形式，%scalar直接给出标量值，类型以f32等具体 dtype 标注：

%dst = tmins %src, %scalar : !pto.tile<...>, f32

AS Level 1（SSA 形式）

SSA（静态单赋值）中间表示，操作数与结果均为!pto.tile<...>类型的值（value），由编译器后续进行资源分配：

%dst = pto.tmins %src, %scalar : (!pto.tile<...>, dtype) -> !pto.tile<...>

AS Level 2（DPS 形式）

DPS（Destructive/显式 place-and-schedule）形式，此时 Tile 已被绑定到具体缓冲区!pto.tile_buf<...>，输入与输出分离声明：

pto.tmins ins(%src, %scalar : !pto.tile_buf<...>, dtype) outs(%dst : !pto.tile_buf<...>)

从源码结构看，pto.tmins指令名称中的 "s" 后缀贯穿三个抽象层次保持一致，便于后端在编译流水线中逐级降级映射。

C++ 内建接口

TMINS 的 C++ 内建接口模板声明于 include/pto/common/pto_instr.hpp，对开发者开放的公共包含头为<pto/pto-inst.hpp>：

template <typename TileDataDst, typename TileDataSrc, typename... WaitEvents> PTO_INST RecordEvent TMINS(TileDataDst &dst, TileDataSrc &src, typename TileDataSrc::DType scalar, WaitEvents &... events);

接口要点：

• 参数顺序：目标 Tiledst、源 Tilesrc、标量scalar（其类型必须等于TileDataSrc::DType）、可选的事件列表events（用于跨流水线同步等待）；
• 返回类型：PTO_INST RecordEvent，即返回一个记录事件，可与后续指令构成事件依赖链；
• 宏分发：TMINS通过MAP_INSTR_IMPL(TMINS, dst, src, scalar)宏映射到平台相关的TMINS_IMPL实现（见 include/pto/common/pto_instr.hpp），同一份用户代码可无缝编译到不同昇腾平台或 CPU 模拟器。

底层实现调用链

在 NPU 侧，TMINS_IMPL最终落到硬件向量指令vmins：

• Atlas A2/A3 平台实现在 include/pto/npu/a2a3/TMins.hpp：MinSOp::BinSInstr直接调用vmins(dst, src0, src1, repeats, 1, 1, 8, 8)，并支持通过dstRepeatStride/srcRepeatStride自定义 repeat 间步长；
• Ascend 950 平台实现在 include/pto/npu/a5/TMins.hpp：MinSOp::BinSInstr调用vmins(reg_dst, reg_src0, src1, preg, MODE_ZEROING)，即带掩码寄存器与清零模式（zeroing mode）的向量比较形式；对于int64_t/uint64_t这类宽类型，A5 实现走Int64Scalar<Int64Op::Min, ...>专用路径，将 64 位标量比较拆分为内部多次 32 位运算组合完成（见 include/pto/npu/a5/TMins.hpp）。

在 CPU 模拟器侧，TMINS 与 TADDS/TSUBS/TMULS/TMAXS 等标量指令共用同一套UnaryTileScalarOpImpl逐元素循环框架，仅以ElementOp::OP_MINS区分运算类型，实现在 include/pto/cpu/TBinSOps.hpp。这意味着在无昇腾硬件的开发环境中，CPU 模拟可以给出与硬件一致的数值行为，便于算子逻辑先行验证。

约束与有效区域

使用 TMINS 时必须满足以下平台与通用约束，违反约束会在编译期（static_assert）或运行期（PTO_ASSERT）报错。

Atlas A2/A3 训练/推理系列产品（实现检查）

• TileData::DType必须是以下类型之一：int32_t、int、int16_t、half、float16_t、float、float32_t；
• 运行时：src.GetValidRow() == dst.GetValidRow()且src.GetValidCol() == dst.GetValidCol()。

对应实现中，A2/A3 的TMINS_IMPL通过static_assert校验数据类型集合，并通过PTO_ASSERT(src.GetValidCol() == dst.GetValidCol(), ...)与PTO_ASSERT(src.GetValidRow() == dst.GetValidRow(), ...)在运行期强制行列有效边界一致（见 include/pto/npu/a2a3/TMins.hpp）。

Ascend 950PR / Ascend 950DT（实现检查）

• TileData::DType必须是以下类型之一：uint8_t、int8_t、uint16_t、int16_t、uint32_t、int32_t、int64_t、uint64_t、half、float、bfloat16_t——相比 A2/A3 大幅扩展了整数与低精度类型覆盖，并额外支持bfloat16_t；
• 运行时：src.GetValidCol() == dst.GetValidCol()。

对应实现中，A5 的TMINS_IMPL同样以static_assert校验类型集合与TileType::Vec位置，以PTO_ASSERT(src0.GetValidCol() == dst.GetValidCol(), ...)校验列数一致（见 include/pto/npu/a5/TMins.hpp）。

通用约束（所有平台）

• dst与src必须使用相同的元素类型（A2/A3 实现在类型层通过static_assert(std::is_same_v<T, typename TileDataDst::DType>, ...)直接强制）；
• 标量类型必须与 Tile 数据类型一致（接口签名中scalar类型即TileDataSrc::DType）；
• Tile 位置必须是向量：TileData::Loc == TileType::Vec。

有效区域

TMINS 以dst.GetValidRow()/dst.GetValidCol()作为迭代域：运算只覆盖目标 Tile 声明的有效行列范围内，超出部分不参与计算。因此调用方应保证src的有效区域至少覆盖dst的有效区域，避免读取越界数据。

实战示例：自动模式与手动模式

PTO 提供两种 Tile 资源管理方式。自动模式由编译器/运行时负责 Tile 的内存放置与调度；手动模式通过TASSIGN显式将 Tile 绑定到指定地址，再由指令发射。

自动（Auto）模式

#include <pto/pto-inst.hpp> using namespace pto; void example_auto() { using TileT = Tile<TileType::Vec, float, 16, 16>; TileT src, dst; TMINS(dst, src, 0.0f); }

声明一个 16×16 的float向量 Tile，调用TMINS将src中每个元素与0.0f比较取小写入dst。0.0f的类型与TileT::DType（float）一致，满足通用约束。

手动（Manual）模式

#include <pto/pto-inst.hpp> using namespace pto; void example_manual() { using TileT = Tile<TileType::Vec, float, 16, 16>; TileT src, dst; TASSIGN(src, 0x1000); TASSIGN(dst, 0x2000); TMINS(dst, src, 0.0f); }

先通过TASSIGN将src绑定到地址0x1000、dst绑定到0x2000，再发射TMINS。这种模式适合开发者需要精确控制缓冲区布局（如复用固定的片上 buffer 池）的场景。

汇编示例（ASM）

自动模式

自动模式下资源放置与调度由编译器/运行时完成，指令以 SSA 值形式出现：

# 自动模式：由编译器/运行时负责资源放置与调度。 %dst = pto.tmins %src, %scalar : (!pto.tile<...>, dtype) -> !pto.tile<...>

手动模式

手动模式先显式绑定资源（pto.tassign将 tile 操作数绑定到物理地址），再发射指令：

# 手动模式：先显式绑定资源，再发射指令。 # 可选（当该指令包含 tile 操作数时）： # pto.tassign %arg0, @tile(0x1000) # pto.tassign %arg1, @tile(0x2000) %dst = pto.tmins %src, %scalar : (!pto.tile<...>, dtype) -> !pto.tile<...>

PTO 汇编形式汇总

%dst = tmins %src, %scalar : !pto.tile<...>, f32 # AS Level 2 (DPS) pto.tmins ins(%src, %scalar : !pto.tile_buf<...>, dtype) outs(%dst : !pto.tile_buf<...>)

端到端算子示例：TLOAD → TMINS → TSTORE

仓库中 tests/cpu/st/testcase/tmins/tmins_kernel.cpp 给出了一个完整的端到端 kernel：将全局内存数据加载到 Tile，执行 TMINS 标量裁剪，再存回全局内存，并用set_flag/wait_flag在 MTE2（加载）、V（向量运算）、MTE3（存储）三条流水线之间建立事件同步：

template <typename T, int kGRows_, int kGCols_, int kTRows_, int kTCols_> __global__ AICORE void runTMins(__gm__ T __out__* out, __gm__ T __in__* src0, __gm__ T __in__* src1) { using DynShapeDim5 = Shape<1, 1, 1, kGRows_, kGCols_>; using DynStridDim5 = Stride<1, 1, 1, kGCols_, 1>; using GlobalData = GlobalTensor<T, DynShapeDim5, DynStridDim5>; using TileData = Tile<TileType::Vec, T, kTRows_, kTCols_, BLayout::RowMajor, -1, -1>; TileData src0Tile(kTRows_, kTCols_); TileData dstTile(kTRows_, kTCols_); TASSIGN(src0Tile, 0x0 + 0x400); TASSIGN(dstTile, 0x8000 + 0x400); GlobalData src0Global(src0); GlobalData dstGlobal(out); TLOAD(src0Tile, src0Global); set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); TMINS(dstTile, src0Tile, src1[0]); // 标量来自全局内存首元素 set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); TSTORE(dstGlobal, dstTile); out = dstGlobal.data(); }

该示例同时展示了一个实用技巧：标量操作数不一定是编译期常量，也可以是src1[0]这类从全局内存读取的运行时值，使 TMINS 可以用于动态阈值的裁剪逻辑。

测试验证与 golden 机制

TMINS 在仓库中拥有完整的多平台测试覆盖，包括 CPU 模拟（tests/cpu/st/testcase/tmins/）、A2/A3（tests/npu/a2a3/src/st/testcase/tmins/）、A5（tests/npu/a5/src/st/testcase/tmins/）以及 costmodel 相关测试（tests/costmodel/st/testcase/tmins/main.cpp）。

以 CPU 侧测试 tests/cpu/st/testcase/tmins/main.cpp 为例，其验证流程为：

1. 通过gen_data.py生成随机输入input1.bin、标量输入input_scalar.bin与参考 golden 文件golden.bin；
2. 使用 ACL 运行时 API（aclrtMalloc/aclrtMemcpy）分配主机与设备内存，将输入拷贝到设备；
3. 启动LaunchTMinskernel，同步流后将结果拷回主机写为output.bin；
4. 将output.bin与golden.bin逐元素比较（容差 0.0001f），通过EXPECT_TRUE(ret)判定测试通过。

覆盖的数据类型与形状包括float/int32_t/int64_t/uint64_t/int16_t/half（以及开启CPU_SIM_BFLOAT_ENABLED时的bfloat16_t），形状覆盖 64×64 与 16×256 两种 Tile 配置（见 tests/cpu/st/testcase/tmins/main.cpp），可直接作为自行扩展 TMINS 用例的模板。

TMINS 与 TMIN 的对比

维度	TMINS（Tile × 标量）	TMIN（Tile × Tile）
数学语义	dst[i][j] = min(src[i][j], scalar)	dst[i][j] = min(src0[i][j], src1[i][j])
第二操作数	标量值（TileDataSrc::DType）	另一个 Tile
A2/A3 数据类型	int32_t/int/int16_t/half/float16_t/float/float32_t	int32_t/int16_t/half/float（并额外要求行主序布局与静态有效边界检查）
运行时校验	行列有效边界一致	src0/src1/dst的validRow/validCol相同
典型用途	阈值截断、上界裁剪	逐元素融合比较（如 ReLU6、逐点 clamp 下限）

两者都要求TileType::Vec向量位置、以dst有效区域为迭代域，且均声明于 include/pto/common/pto_instr.hpp，接口风格完全一致，可在同一 kernel 中混合使用。

总结

TMINS 是 PTO 指令集中一个简洁但高频使用的 Tile-标量二元指令：数学上等价于对有效区域内每个元素执行min(x, scalar)，编程上通过TMINS(dst, src, scalar)一行即可完成整个 Tile 的标量裁剪。使用时需要特别关注两点：一是平台相关的数据类型支持范围（A2/A3 与 950 系列差异明显，950 额外支持 8/16/64 位整数与bfloat16_t）；二是有效区域一致性约束（src与dst的 valid 行列必须匹配）。其底层实现无论是 A2/A3 的vminsrepeat 循环、A5 的带掩码vmins与 64 位专用路径，还是 CPU 模拟器的UnaryTileScalarOpImpl逐元素循环，都保证了跨平台一致的语义，配合仓库内覆盖多平台、多数据类型、多 Tile 形状的完整测试用例，开发者可以放心在昇腾 NPU 上以自动或手动模式使用 TMINS 实现高效的逐元素最小值运算。

【免费下载链接】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