PTO ISA TSQRT 指令详解：Tile 级逐元素平方根的语义、精度算法与源码实现

PTO ISA TSQRT 指令详解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本文围绕 PTOParallel Tile Operation指令集中的TSQRT指令展开完整覆盖其数学语义、汇编语法同步形式、AS Level 1 SSA、AS Level 2 DPS、C 内建接口与 NPU 实现约束并结合当前仓库中的 NPU/CPU 实现源码与标准测试用例深入讲解SqrtAlgorithm::HIGH_PRECISION高精度开方算法的构造原理与事件化流水调用方式帮助开发者在 Ascend 平台上正确、高效地编写 Tile 级开方算子。一、指令概述与数学语义TSQRT是 PTO 虚拟 ISA 中的一条逐元素elementwise平方根指令以输出 Tile 的有效区域为迭代域对源 Tile 的每个元素执行开方运算。原文档定义见 docs/isa/TSQRT.md中文版见 docs/isa/TSQRT_zh.md给出的数学语义为$$ \mathrm{dst}{i,j} \sqrt{\mathrm{src}{i,j}} \quad \text{对所有有效区域内的元素 (i, j)} $$作为单源、单目标dst/src的向量一元运算TSQRT与TRELU、TNEG、TLOG等指令同属 PTO 向量一元算子家族与之相关的还有倒数平方根指令TRSQRT两者在 include/pto/common/pto_instr.hpp 中相邻声明可对照阅读。二、汇编语法三种描述层级TSQRT的汇编描述遵循 PTO 的多层级语法体系原文档给出了三种形式开发者在阅读 IR 或手写汇编时需加以区分1. 同步形式基础助记符%dst tsqrt %src : !pto.tile...2. AS Level 1SSA 形式%dst pto.tsqrt %src : !pto.tile... - !pto.tile...3. AS Level 2DPS 形式pto.tsqrt ins(%src : !pto.tile_buf...) outs(%dst : !pto.tile_buf...)三种形式中!pto.tile.../!pto.tile_buf...为 Tile 类型占位符行/列维度、数据类型等具体由上下文推导DPSDifferential Programming Syntax层级则显式区分输入ins与输出outs操作数与自动放置/手动绑定的资源模型相对应。自动模式与手动模式的区别原文档的 ASM 示例部分进一步说明了两种资源管理模式的差异自动模式Auto Mode由编译器/运行时负责 Tile 资源的放置与调度开发者只需写出指令本身# Auto mode: compiler/runtime-managed placement and scheduling. %dst pto.tsqrt %src : !pto.tile... - !pto.tile...手动模式Manual Mode在发射指令之前必须显式将 Tile 操作数绑定到向量缓冲区的指定地址tile(addr)即通过pto.tassign完成资源绑定# 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.tsqrt %src : !pto.tile... - !pto.tile...两种模式与 C 层面对应自动模式直接使用Tile对象手动模式则需先调用TASSIGN显式指定 Tile 地址见第五节示例。三、C 内建接口与精度控制TSQRT的 C 内建函数声明于 include/pto/common/pto_instr.hpp公共包含头为pto/pto-inst.hpp。原文档给出的签名如下template typename TileDataDst, typename TileDataSrc, typename... WaitEvents PTO_INST RecordEvent TSQRT(TileDataDst dst, TileDataSrc src, WaitEvents ... events);结合仓库源码可以进一步确认实际实现中TSQRT在三个类型参数之前还有一个编译期精度模板参数默认值SqrtAlgorithm::DEFAULT完整签名见 include/pto/common/pto_instr.hpptemplate auto PrecisionType SqrtAlgorithm::DEFAULT, typename TileDataDst, typename TileDataSrc, typename... WaitEvents PTO_INST RecordEvent TSQRT(TileDataDst dst, TileDataSrc src, WaitEvents... events) { detail::PtoWaitEvents(events...); TSQRT_IMPLPrecisionType(dst, src); return {}; }由此可以得出三个关键事实精度是一等公民SqrtAlgorithm定义于 include/pto/common/type.hpp为enum class SqrtAlgorithm : uint8_t { DEFAULT, HIGH_PRECISION }。调用方可以通过TSQRTSqrtAlgorithm::HIGH_PRECISION(dst, src, event)的方式要求更高精度的实现不指定时走DEFAULT路径。事件驱动的同步模型模板参数包WaitEvents...允许传入若干Event函数体内通过detail::PtoWaitEvents(events...)等待前序指令如TLOAD完成后才发射本指令返回值RecordEvent记录本次指令可作为后续TSTORE等指令的前置事件构成完整的指令级依赖链。分发机制TSQRT是统一的入口PTO_INST宏标注实际按目标平台分发到各后端的TSQRT_IMPL——NPU A5 位于 include/pto/npu/a5/TUnaryOp.hppCPU 模拟后端位于 include/pto/cpu/TSqrt.hppA2/A3 后端位于 include/pto/npu/a2a3/TUnaryOp.hpp。四、约束条件与有效区域原文档“约束”一节列出的是NPU 实现侧的检查项从源码可以逐条印证见 include/pto/npu/a5/TUnaryOp.hpp 中的TSQRT_IMPL及其调用的TUnaryCheck约束项要求检查阶段数据类型TileData::DType必须是float或half编译期静态断言Tile 位置必须是向量位置TileData::Loc TileType::Vec编译期静态有效边界TileData::ValidRow TileData::Rows且TileData::ValidCol TileData::Cols编译期运行时形状一致src.GetValidRow() dst.GetValidRow()且src.GetValidCol() dst.GetValidCol()运行时PTO_ASSERTTile 布局必须行主序TileData::isRowMajor编译期其中运行时形状一致性检查可直接在源码中看到对应实现include/pto/npu/a5/TUnaryOp.hppunsigned dstValidRow dst.GetValidRow(); unsigned dstValidCol dst.GetValidCol(); PTO_ASSERT(dstValidCol src.GetValidCol(), TSQRT: Number of columns of src and dst must be the same.); PTO_ASSERT(dstValidRow src.GetValidRow(), TSQRT: Number of rows of src and dst must be the same.);有效区域Valid Region该指令以dst.GetValidRow()/dst.GetValidCol()作为迭代域即只计算 dst 声明的有效行×有效列范围内的元素超出部分保持不动。这一点在 tests/npu/a5/src/st/testcase/tsqrt/tsqrt_kernel.cpp 的测试用例中得到体现源/目标 Tile 可以具有不同的物理尺寸如 128×128 vs 64×64但运行时有效区域必须一致。定义域 / NaN 行为原文档明确指出对负数输入等非法定义域情况的处理行为是target-defined由目标平台定义的文档层面不做跨平台保证——这意味着在移植代码时应以目标平台的实测行为为准。CPU 模拟后端的行为差异值得注意的是CPU 后端include/pto/cpu/TSqrt.hpp的类型约束比 NPU 更宽静态断言接受bfloat16_t、half、float三种类型要求 src/dst 类型相同且以double精度调用std::sqrt再回转const auto x static_castdouble(src.data()[GetTileElementOffsetTileDataSrc(r, c)]); dst.data()[GetTileElementOffsetTileDataDst(r, c)] static_casttypename TileDataDst::DType(std::sqrt(x));因此使用 PTO CPU 模拟器cpu_sim验证TSQRT时bfloat16 用例也能通过编译而在 NPU 上会被静态断言拦截。做平台间一致性验证时应注意这一差异。五、NPU 实现纵深从模板分发到向量寄存器以 A5 后端为例TSQRT_IMPL的调用链为include/pto/npu/a5/TUnaryOp.hppTSQRT (pto_instr.hpp, 公共入口) └─ TSQRT_IMPLPrecisionType // 检查 取 validRow/validCol └─ TSqrtPrecisionType, Dst, Src // __tf__ PTO_INTERNAL, OP_TYPE(element_wise) └─ TUnaryOpDst, Src, SqrtOpPrecisionType, T └─ SqrtOp::UnaryInstr // 逐行作用于 RegTensor几个源码层面的要点TSqrt标注为__tf__tile function并带OP_TYPE(element_wise)属性表明它面向 Tile 级数据布局做向量化展开TUnaryOp是 PTO 中一元向量运算的通用骨架负责按有效区域逐行取RegTensor向量寄存器组并调用SqrtOp::UnaryInstr同时接收MaskReg参与运算MODE_ZEROING掩码模式下被屏蔽的 lane 结果清零默认路径PrecisionType DEFAULT直接落到硬件向量指令vsqrt} else { vsqrt(dstReg, srcReg, pReg, MODE_ZEROING); }HIGH_PRECISION 算法亚正规数缩放 牛顿迭代精修当选择SqrtAlgorithm::HIGH_PRECISION时SqrtOp::UnaryInstr会根据元素类型分派到两个定制实现include/pto/npu/a5/custom/TSqrtHp.hppfloat 路径SqrtFloatImpl核心思路是对vsqrt的一次近似结果做牛顿法精修并妥善处理亚正规数与特殊值亚正规数保护对src 1.0的 lane先乘16777216.0f即 2²⁴把操作数放大出亚正规区间再开方开方完成后乘回0.000244140625f2⁻¹⁹ 的缩放因子恢复量级——通过vcmps_lt/vsel对两条路径按 lane 选择一次牛顿迭代由x的平方根估计s vsqrt(x)计算倒数估计r 1/s误差项e x - r·x按牛顿迭代式更新估计值源码中以vmul/vmula/vmadd组合实现特殊值旁路单独用vcmps_eq检测Inf0x7f800000u并通过位或技巧与0x80000000u按位或后等于-0.0同时识别0.0/-0.0对这些 lane 直接vsel回原始输入√(∞)∞、√(±0)±0避免迭代污染特殊值语义。half 路径SqrtPrecisionImpl逻辑相同但用 half 位模式常量表达以0x03FF为亚正规阈值做比较先乘0x6C002¹²放大、开方后乘0x24002⁻⁶还原。这些实现细节解释了为什么HIGH_PRECISION需要更多向量寄存器临时空间也印证了原文档“行为由目标定义”的表述——精度算法本身就是目标平台A5 后端的具体实现选择。六、完整使用示例6.1 最小示例Auto / Manual原文档给出两个最小可编译示例。自动模式下 Tile 地址由编译器/运行时管理#include pto/pto-inst.hpp using namespace pto; void example_auto() { using TileT TileTileType::Vec, float, 16, 16; TileT src, dst; TSQRT(dst, src); }手动模式下需先TASSIGN绑定地址再发射指令TASSIGN对应汇编中的pto.tassign#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); TSQRT(dst, src); }6.2 完整 Kernel事件化流水与精度选择更贴近实战的写法可以参考仓库中的标准测试 kernel tests/npu/a5/src/st/testcase/tsqrt/tsqrt_kernel.cpp它展示了TSQRT在真实算子中的完整用法全局张量加载、事件同步、精度选择、in-place 处理template typename T, int dstRow, int dstCol, int srcRow, int srcCol, int validRow, int validCol, bool highPrecision, bool isInPlace __global__ AICORE void runTsqrt(__gm__ T __out__* out, __gm__ T __in__* src) { using DynShapeDim5 Shape1, 1, 1, -1, -1; using SrcGlobalData GlobalTensorT, DynShapeDim5, pto::Stride1, 1, srcRow, srcCol, 1; using DstGlobalData GlobalTensorT, DynShapeDim5, pto::Stride1, 1, dstRow, dstCol, 1; DstGlobalData dstGlobal(out, DynShapeDim5(validRow, validCol)); SrcGlobalData srcGlobal(src, DynShapeDim5(validRow, validCol)); using SrcTileData TileTileType::Vec, T, srcRow, srcCol, BLayout::RowMajor, -1, -1; using DstTileData TileTileType::Vec, T, dstRow, dstCol, BLayout::RowMajor, -1, -1; SrcTileData srcTile(validRow, validCol); DstTileData dstTile(validRow, validCol); TASSIGN(dstTile, isInPlace ? 0x0 : srcRow * srcCol * sizeof(T)); TASSIGN(srcTile, 0x0); EventOp::TLOAD, Op::TSQRT event0; EventOp::TSQRT, Op::TSTORE_VEC event1; constexpr auto precisionType highPrecision ? SqrtAlgorithm::HIGH_PRECISION : SqrtAlgorithm::DEFAULT; event0 TLOAD(srcTile, srcGlobal); event1 TSQRTprecisionType(dstTile, srcTile, event0); // 等待 TLOAD 完成 TSTORE(dstGlobal, dstTile, event1); }该 kernel 的要点类型参数化highPrecision编译期折叠为constexpr auto precisionType作为模板实参传给TSQRT...两种精度生成独立实例无运行时分支开销事件链TLOAD返回EventOp::TLOAD, Op::TSQRTTSQRT消费该事件并返回EventOp::TSQRT, Op::TSTORE_VEC供TSTORE使用——这正是第三节WaitEvents参数包在事件系统中的具体形态in-place 模式当输出与输入地址重叠时TASSIGN将 dst 绑定到0x0与 src 同址此时必须保证TSQRT的读-写模式支持原地更新逐元素一元运算在语义上满足显式实例化覆盖多种形状组合如LaunchTSqrtfloat, 128, 128, 64, 64, 64, 64, false, falsedst 为 128×128、src 为 64×64、有效区域 64×64与第四节的“物理尺寸可以不同、有效区域必须一致”的约束相互印证。同目录下还有数据生成脚本 tests/npu/a5/src/st/testcase/tsqrt/gen_data.py 与驱动 tests/npu/a5/src/st/testcase/tsqrt/main.cppA2/A3、Kirin 系 NPU 与 CPU 后端均有对应的tsqrt测试目录如 tests/npu/a2a3/src/st/testcase/tsqrt、tests/cpu/st/testcase/tsqrt可作为多平台一致性验证的入口。七、小结与延伸TSQRT的语义是严格逐元素的dst[i,j] sqrt(src[i,j])迭代域由 dst 的有效行/列决定三种汇编层级同步 / AS1 SSA / AS2 DPS与 C 内建的 Auto/Manual 两种资源模式一一对应NPU 后端要求float/half类型、向量位置、行主序布局src/dst 有效形状在运行时一致CPU 模拟后端额外支持bfloat16SqrtAlgorithm::HIGH_PRECISION通过亚正规数缩放加牛顿迭代精修float 与 half 各有一套位级实现并保留±0、Inf等特殊值语义适用于对精度敏感的数值算法负数输入等定义域外行为是 target-defined 的跨平台移植时应以实测为准。延伸方向同族的倒数平方根指令TRSQRTinclude/pto/common/pto_instr.hpp注意它额外提供了带tmp临时 Tile 的重载、以及TSQRT与TROWEXPAND的融合形态仓库提供了 tests/npu/a5/src/st/testcase/trowexpand_tsqrt 融合测试都可在本仓库中按相同的路径结构继续追踪。【免费下载链接】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），仅供参考