CANN PTO-ISA 逐元素按位或指令 TOR 完全指南语法、C 内建接口与实现原理【免费下载链接】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-isaTORTile OR是 CANN PTO-ISA 虚拟指令集中用于对两个 Tile 执行逐元素按位或bitwise OR的二元运算指令在图像处理、位掩码计算、特征图逻辑合并等整数向量场景中有着直接应用。本文以 docs/isa/TOR_zh.md 为骨架结合仓库内 NPUA2/A3 与 A5实现头文件与 CPU/NPU 测试用例系统讲解 TOR 的数学语义、三级汇编语法、C 内建接口调用方式、平台约束以及底层vor指令发射与类型转换的实现细节。读完本文你将能够在自动模式与手动模式下正确编写并验证使用 TOR 指令的 PTO 算子。指令概述TOR 是一条逐元素elementwise的二元按位或指令它对两个 Tile 中位于有效区域valid region内的每一对元素执行按位或运算并将结果写入目标 Tile。该指令属于 PTO 指令集中的向量二元逻辑运算族与之配套的还有标量版本的TORSTile OR with Scalar见 pto_instr.hpp。从仓库目录看TOR 与其他指令一样拥有完整的文档 — 声明 — 实现 — 测试链路指令规范docs/isa/TOR_zh.md公共头文件入口include/pto/pto-inst.hppA2/A3 平台实现include/pto/npu/a2a3/TOr.hppA5 平台实现include/pto/npu/a5/TOr.hppNPU 测试用例tests/npu/a5/src/st/testcase/tor/tor_kernel.cppCPU 仿真测试用例tests/cpu/st/testcase/tor/tor_kernel.cpp数学语义对于有效区域内的每个元素(i, j)TOR 计算$$ \mathrm{dst}{i,j} \mathrm{src0}{i,j} ;|; \mathrm{src1}_{i,j} $$其中|为按位或运算符src0、src1为两个输入 Tiledst为输出 Tile。运算逐位进行因此结果的位宽与输入保持一致且不涉及进位或跨元素依赖天然适合 SIMD 向量化执行。汇编语法从同步形式到三级 ASTOR 在 PTO 汇编中有三种表达层次分别对应不同的资源管理粒度。同步形式最简表达类型标注为 tile%dst tor %src0, %src1 : !pto.tile...AS Level 1SSA 形式显式写出操作符pto.tor与完整的类型签名输入输出均为!pto.tile...%dst pto.tor %src0, %src1 : (!pto.tile..., !pto.tile...) - !pto.tile...AS Level 2DPS 形式使用ins(...) outs(...)显式区分输入与输出缓冲区操作数类型为物理缓冲区类型!pto.tile_buf...pto.tor ins(%src0, %src1 : !pto.tile_buf..., !pto.tile_buf...) outs(%dst : !pto.tile_buf...)三种形式描述的是同一运算在不同抽象层级上的表示同步形式是给程序员看的最简写法SSA 形式面向编译器中间表示DPS 形式则更贴近物理寄存与缓冲区的分配。C 内建接口在 C 算子代码中推荐直接调用内建函数TOR。其声明位于 include/pto/common/pto_instr.hpptemplate typename TileData, typename... WaitEvents PTO_INST RecordEvent TOR(TileData dst, TileData src0, TileData src1, WaitEvents ... events);接口说明dst、src0、src1均为 Tile 类型的引用分别表示输出与两个输入可变参数events用于传入等待事件实现指令间的依赖同步返回值RecordEvent可被后续指令如TSTORE作为依赖事件继续传递从而构建流水线公共包含头为pto/pto-inst.hpp内部声明位于pto/common/pto_instr.hpp从源码看该接口通过MAP_INSTR_IMPL(TOR, dst, src0, src1)宏分发到平台相关的TOR_IMPL实现见 pto_instr.hpp。平台约束与有效区域TOR 在不同硬件平台上支持的元素类型范围不同使用时必须注意。下表汇总自 TOR_zh.md 的约束章节约束项Atlas A2/A3 训练/推理系列产品Ascend 950PR / Ascend 950DT支持的元素类型uint8_t、int8_t、uint16_t、int16_t、uint32_t、int32_t在左侧基础上额外支持int64_t、uint64_tdst/src0/src1元素类型必须相同必须相同内存布局必须行主序row-major必须行主序row-major有效形状src0/src1的GetValidRow()/GetValidCol()必须与dst一致同左有效区域指令以dst.GetValidRow()/dst.GetValidCol()作为迭代域即只对目标 Tile 的有效行、列范围进行运算越界区域不参与计算。这些约束在源码中以static_assert与运行时断言的形式固化。以 a2a3/TOr.hpp 的TOrCheck为例static_assert( std::is_sameT, typename TileDataSrc0::DType::value std::is_sameT, typename TileDataSrc1::DType::value, Fix: TOR the data type of dst must be consistent with of src0 and src1.); static_assert( TileDataDst::isRowMajor TileDataSrc0::isRowMajor TileDataSrc1::isRowMajor, Fix: TOR only support row major layout.); PTO_ASSERT( src0.GetValidRow() validRows src0.GetValidCol() validCols, Fix: TOR input tile src0 valid shape mismatch with output tile dst shape.);其中数据类型与行主序约束在编译期检查有效形状一致性在运行时检查违反任一约束都会得到带 Fix: 前缀的明确报错信息便于快速定位问题。源码级实现剖析底层指令发射与类型转换TOR 的通用封装背后是平台相关的向量按位或指令。理解实现细节有助于预测性能特征与排查边界问题。A2/A3 平台vor指令与 32 字节块语义在 a2a3/TOr.hpp 中OrOp封装了底层向量指令template typename T struct OrOp { PTO_INTERNAL static void BinInstr(__ubuf__ T* dst, __ubuf__ T* src0, __ubuf__ T* src1, uint8_t repeats) { vor(dst, src0, src1, repeats, 1, 1, 1, 8, 8, 8); } ... };实现要点数据位于__ubuf__统一缓冲区即片上向量存储中输入输出均为 Tile 指针当三个 Tile 的行跨步一致且无需类型变换时走BinaryInstrOrOpT, T, TileDataDst, elementsPerRepeat, blockSizeElem, dstRowStride快速路径否则走带独立行跨步的通用路径见 a2a3/TOr.hpp关键常量定义于 include/pto/common/constants.hppREPEAT_BYTE 256、BLOCK_BYTE_SIZE 32。每个 repeat 处理 256 字节块大小为 32 字节由此可推导出每个 repeat 处理的元素数为256 / sizeof(T)每块元素数为32 / sizeof(T)对 4 字节类型使用B322B16Trait、对其他类型使用B82B16Trait做数据类型转换TransType、TransSize、TransStride以保证与向量指令的位宽对齐要求匹配。A5 平台寄存器向量指令与 64 位整数扩展在 a5/TOr.hpp 中OrOp使用寄存器张量RegTensor与掩码寄存器MaskReg发射vor指令64 位整数int64_t/uint64_t走独立的Int64BinaryInt64Op::Or, ...路径其余类型走BinaryInstrOrOpT, ...路径。TOr函数还带VFImplKind version参数支持向量函数VF实现的选择默认值为VFIMPL_DEFAULT。由此可见A5 平台对 TOR 的 64 位整数支持并非简单地复用 32 位路径而是有专门的 64 位二元实现来保证uint64_t/int64_t的逐位语义正确。完整示例从最小内核到事件流水线最小示例文档版文档 TOR_zh.md 给出的最小用法如下可直接复制编译#include pto/pto-inst.hpp using namespace pto; void example() { using TileT TileTileType::Vec, int32_t, 16, 16; TileT a, b, out; TOR(out, a, b); }这里TileTileType::Vec, int32_t, 16, 16声明了一个 16×16 的int32_t向量 TileTOR(out, a, b)完成out a | b。带事件依赖的完整内核测试用例版真实算子中TOR 通常与 TLOAD/TSTORE 组成加载 — 计算 — 存储流水线。仓库 NPU 测试用例 tests/npu/a5/src/st/testcase/tor/tor_kernel.cpp 展示了标准写法template typename T, int kTRows_, int kTCols_, int vRows, int vCols __global__ AICORE void runTOr(__gm__ T __out__* out, __gm__ T __in__* src0, __gm__ T __in__* src1) { using DynShapeDim5 Shape1, 1, 1, vRows, vCols; using DynStridDim5 pto::Stride1, 1, 1, vCols, 1; using GlobalData GlobalTensorT, DynShapeDim5, DynStridDim5; using TileData TileTileType::Vec, T, kTRows_, kTCols_, BLayout::RowMajor, -1, -1; TileData src0Tile(vRows, vCols); TileData src1Tile(vRows, vCols); TileData dstTile(vRows, vCols); TASSIGN(src0Tile, 0x0); TASSIGN(src1Tile, 0x10000); TASSIGN(dstTile, 0x20000); GlobalData src0Global(src0); GlobalData src1Global(src1); GlobalData dstGlobal(out); EventOp::TLOAD, Op::TOR event0; EventOp::TOR, Op::TSTORE_VEC event1; TLOAD(src0Tile, src0Global); event0 TLOAD(src1Tile, src1Global); event1 TOR(dstTile, src0Tile, src1Tile, event0); TSTORE(dstGlobal, dstTile, event1); out dstGlobal.data(); }该用例的要点通过TASSIGN显式绑定 Tile 到缓冲区地址0x0、0x10000、0x20000这是手动模式下的资源管理方式使用EventOp::TLOAD, Op::TOR/EventOp::TOR, Op::TSTORE_VEC声明事件让第二次TLOAD的完成事件驱动TORTOR的完成事件再驱动TSTORE实现三阶段流水重叠同一文件底部还实例化了多组形状/类型组合包括int64_t/uint64_t的非方阵与窄有效区域如4×15用例覆盖了 A5 平台的 64 位路径与有效区域边界。CPU 仿真版用例 tests/cpu/st/testcase/tor/tor_kernel.cpp 结构类似不含事件链直接顺序执行TLOAD→TOR→TSTORE适合在没有 NPU 硬件时验证功能正确性。汇编形式示例ASM自动模式自动模式下Tile 的放置与调度由编译器/运行时负责程序员只需给出运算本身# 自动模式由编译器/运行时负责资源放置与调度。 %dst pto.tor %src0, %src1 : (!pto.tile..., !pto.tile...) - !pto.tile...手动模式手动模式要求先显式绑定资源再发射指令。对含 tile 操作数的指令可选地通过pto.tassign将 SSA 值绑定到指定地址的 tile 资源# 手动模式先显式绑定资源再发射指令。 # 可选当该指令包含 tile 操作数时 # pto.tassign %arg0, tile(0x1000) # pto.tassign %arg1, tile(0x2000) %dst pto.tor %src0, %src1 : (!pto.tile..., !pto.tile...) - !pto.tile...PTO 汇编形式%dst tor %src0, %src1 : !pto.tile... # AS Level 2 (DPS) pto.tor ins(%src0, %src1 : !pto.tile_buf..., !pto.tile_buf...) outs(%dst : !pto.tile_buf...)测试与验证路径仓库为 TOR 提供了跨平台的验证覆盖可作为编写与调试自身算子的参考NPU A5 平台tests/npu/a5/src/st/testcase/tor/tor_kernel.cpp覆盖uint8_t、int8_t、uint16_t、int16_t、uint32_t、int32_t、int64_t、uint64_t及多种 Tile 形状含 1×16384 长条与 2048×16 宽幅形状NPU A2/A3 平台tests/npu/a2a3/src/st/testcase/tor/tor_kernel.cpp覆盖 32 位及以下整数类型CPU 仿真tests/cpu/st/testcase/tor/tor_kernel.cpp可在无 NPU 环境下先行验证逻辑正确性其构建入口见 tests/cpu/st/testcase/tor/CMakeLists.txtpto_cpu_sim_st(tor)此外 tests/npu/kirin9030/src/st/testcase/tor/tor_kernel.cpp 与 tests/npu/kirinDev0000/src/st/testcase/tor/tor_kernel.cpp 提供了更多平台的覆盖。小结TOR 是 PTO-ISA 中语义最简单、用途最广的向量二元逻辑指令之一数学上是逐元素按位或汇编上支持同步形式与 AS Level 1/2 三级表达C 侧通过TOR(dst, src0, src1, ...)内建接口即可调用。使用时需重点把握两点约束三操作数元素类型必须一致且为行主序以及src0/src1有效形状必须与dst一致。在实现层面A2/A3 平台通过vor向量指令按 256 字节 repeat、32 字节块组织运算并做类型转换适配A5 平台则额外提供了独立的 64 位整数实现路径。对于需要标量按位或的场景可参考配套的TORS指令声明于 pto_instr.hpp。建议新算子开发时优先参考上述测试用例中的事件流水线写法以充分发挥片上并行能力。【免费下载链接】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），仅供参考