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-isaParallel Tile Operation面向昇腾平台的 Tile 级虚拟指令集中用于执行Tile 与标量逐元素最小值运算的核心向量指令它以整个 Tile 的每个元素为操作对象与一个标量值逐一比较并取较小者写入目标 Tile。本文以 TMINS_zh.md 为骨架结合仓库内 NPUA2/A3、Ascend 950PR/950DT与 CPU 模拟器上的源码实现及完整测试用例系统讲解 TMINS 的数学语义、三级汇编形式、C 内建接口、平台约束、自动/手动两种编程模式并给出可复制运行的实战示例。读完本文你将能够在 PTO 编程框架中正确声明与调用 TMINS理解其有效区域valid region迭代规则与底层vmins指令的映射关系并能参照测试用例编写自己的标量归约类算子。指令一览TMINS 属于 PTO 指令集中 Tile 与标量的二元运算族S 后缀表示 scalar即标量操作数与 TMINTile 与 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..., f32AS Level 1SSA 形式SSA静态单赋值中间表示操作数与结果均为!pto.tile...类型的值value由编译器后续进行资源分配%dst pto.tmins %src, %scalar : (!pto.tile..., dtype) - !pto.tile...AS Level 2DPS 形式DPSDestructive/显式 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.hpptemplate 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最终落到硬件向量指令vminsAtlas A2/A3 平台实现在 include/pto/npu/a2a3/TMins.hppMinSOp::BinSInstr直接调用vmins(dst, src0, src1, repeats, 1, 1, 8, 8)并支持通过dstRepeatStride/srcRepeatStride自定义 repeat 间步长Ascend 950 平台实现在 include/pto/npu/a5/TMins.hppMinSOp::BinSInstr调用vmins(reg_dst, reg_src0, src1, preg, MODE_ZEROING)即带掩码寄存器与清零模式zeroing mode的向量比较形式对于int64_t/uint64_t这类宽类型A5 实现走Int64ScalarInt64Op::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_vT, typename TileDataDst::DType, ...)直接强制标量类型必须与 Tile 数据类型一致接口签名中scalar类型即TileDataSrc::DTypeTile 位置必须是向量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 TileTileType::Vec, float, 16, 16; TileT src, dst; TMINS(dst, src, 0.0f); }声明一个 16×16 的float向量 Tile调用TMINS将src中每个元素与0.0f比较取小写入dst。0.0f的类型与TileT::DTypefloat一致满足通用约束。手动Manual模式#include pto/pto-inst.hpp using namespace pto; void example_manual() { using TileT TileTileType::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 Shape1, 1, 1, kGRows_, kGCols_; using DynStridDim5 Stride1, 1, 1, kGCols_, 1; using GlobalData GlobalTensorT, DynShapeDim5, DynStridDim5; using TileData TileTileType::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/A3tests/npu/a2a3/src/st/testcase/tmins/、A5tests/npu/a5/src/st/testcase/tmins/以及 costmodel 相关测试tests/costmodel/st/testcase/tmins/main.cpp。以 CPU 侧测试 tests/cpu/st/testcase/tmins/main.cpp 为例其验证流程为通过gen_data.py生成随机输入input1.bin、标量输入input_scalar.bin与参考 golden 文件golden.bin使用 ACL 运行时 APIaclrtMalloc/aclrtMemcpy分配主机与设备内存将输入拷贝到设备启动LaunchTMinskernel同步流后将结果拷回主机写为output.bin将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 的对比维度TMINSTile × 标量TMINTile × Tile数学语义dst[i][j] min(src[i][j], scalar)dst[i][j] min(src0[i][j], src1[i][j])第二操作数标量值TileDataSrc::DType另一个 TileA2/A3 数据类型int32_t/int/int16_t/half/float16_t/float/float32_tint32_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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 
长河君 · 2026/9/19 18:06:32 · 阅读 1.2k 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-isaParallel Tile Operation面向昇腾平台的 Tile 级虚拟指令集中用于执行Tile 与标量逐元素最小值运算的核心向量指令它以整个 Tile 的每个元素为操作对象与一个标量值逐一比较并取较小者写入目标 Tile。本文以 TMINS_zh.md 为骨架结合仓库内 NPUA2/A3、Ascend 950PR/950DT与 CPU 模拟器上的源码实现及完整测试用例系统讲解 TMINS 的数学语义、三级汇编形式、C 内建接口、平台约束、自动/手动两种编程模式并给出可复制运行的实战示例。读完本文你将能够在 PTO 编程框架中正确声明与调用 TMINS理解其有效区域valid region迭代规则与底层vmins指令的映射关系并能参照测试用例编写自己的标量归约类算子。指令一览TMINS 属于 PTO 指令集中 Tile 与标量的二元运算族S 后缀表示 scalar即标量操作数与 TMINTile 与 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..., f32AS Level 1SSA 形式SSA静态单赋值中间表示操作数与结果均为!pto.tile...类型的值value由编译器后续进行资源分配%dst pto.tmins %src, %scalar : (!pto.tile..., dtype) - !pto.tile...AS Level 2DPS 形式DPSDestructive/显式 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.hpptemplate 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最终落到硬件向量指令vminsAtlas A2/A3 平台实现在 include/pto/npu/a2a3/TMins.hppMinSOp::BinSInstr直接调用vmins(dst, src0, src1, repeats, 1, 1, 8, 8)并支持通过dstRepeatStride/srcRepeatStride自定义 repeat 间步长Ascend 950 平台实现在 include/pto/npu/a5/TMins.hppMinSOp::BinSInstr调用vmins(reg_dst, reg_src0, src1, preg, MODE_ZEROING)即带掩码寄存器与清零模式zeroing mode的向量比较形式对于int64_t/uint64_t这类宽类型A5 实现走Int64ScalarInt64Op::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_vT, typename TileDataDst::DType, ...)直接强制标量类型必须与 Tile 数据类型一致接口签名中scalar类型即TileDataSrc::DTypeTile 位置必须是向量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 TileTileType::Vec, float, 16, 16; TileT src, dst; TMINS(dst, src, 0.0f); }声明一个 16×16 的float向量 Tile调用TMINS将src中每个元素与0.0f比较取小写入dst。0.0f的类型与TileT::DTypefloat一致满足通用约束。手动Manual模式#include pto/pto-inst.hpp using namespace pto; void example_manual() { using TileT TileTileType::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 Shape1, 1, 1, kGRows_, kGCols_; using DynStridDim5 Stride1, 1, 1, kGCols_, 1; using GlobalData GlobalTensorT, DynShapeDim5, DynStridDim5; using TileData TileTileType::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/A3tests/npu/a2a3/src/st/testcase/tmins/、A5tests/npu/a5/src/st/testcase/tmins/以及 costmodel 相关测试tests/costmodel/st/testcase/tmins/main.cpp。以 CPU 侧测试 tests/cpu/st/testcase/tmins/main.cpp 为例其验证流程为通过gen_data.py生成随机输入input1.bin、标量输入input_scalar.bin与参考 golden 文件golden.bin使用 ACL 运行时 APIaclrtMalloc/aclrtMemcpy分配主机与设备内存将输入拷贝到设备启动LaunchTMinskernel同步流后将结果拷回主机写为output.bin将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 的对比维度TMINSTile × 标量TMINTile × Tile数学语义dst[i][j] min(src[i][j], scalar)dst[i][j] min(src0[i][j], src1[i][j])第二操作数标量值TileDataSrc::DType另一个 TileA2/A3 数据类型int32_t/int/int16_t/half/float16_t/float/float32_tint32_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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考