PTO-ISA TPARTADD 指令详解部分有效区域逐元素加法tpartadd的语义、接口与平台实现【免费下载链接】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导读本文以 CANN pto-isa 开源仓库的指令文档 TPARTADD_zh.md 为骨架系统讲解TPARTADDtpartadd这条 Tile 级向量指令它面向目标有效区域valid region执行逐元素加法支持两个源 Tile 的有效区域与目标不一致的部分有效场景。读完本文你将掌握 TPARTADD 的数学语义、SSA/DPS 两级汇编形式、C 内建接口签名、A2A3 与 Ascend 950A5平台的实现差异与类型约束并能结合 TPartAdd.hpp、TPartOp.hpp、TPartBinOps.hpp 等源码与 tpartadd 测试用例 写出可运行、可验证的算子代码。指令概述与数学语义TPARTADD 是 PTO 指令集中用于部分有效区域逐元素加法的向量指令全称 Tile Partial Add。它不属于普通的TADD要求各 Tile 有效区域完全一致而是允许两个源 Tilesrc0、src1与目标 Tiledst的有效区域存在差异加法只在 dst 的有效区域内、且两个输入同时有效的元素上执行。语义定义在目标有效区域内的每个元素(i, j)结果按以下规则确定若src0与src1在该位置均有效dst[i][j] src0[i][j] src1[i][j]若仅src0有效dst[i][j] src0[i][j]直接拷贝若仅src1有效dst[i][j] src1[i][j]直接拷贝。其余有效区域不匹配的情形例如某个输入在目标有效区域内既不是完全有效也不是完全无效由具体实现定义编程时应当避免依赖这类行为。这一语义可以形式化为dst[i][j] src0[i][j] src1[i][j] 若 src0、src1 在 (i,j) 均有效 src0[i][j] 若仅 src0 在 (i,j) 有效 src1[i][j] 若仅 src1 在 (i,j) 有效指令示意图如下直观展示了两个输入有效区域叠加到目标区域的过程与 TADD 的定位差异普通TADD要求参与运算的 Tile 有效区域完全对齐而 TPARTADD 解决的是目标区域更大、某个输入只有部分区域有数据的场景。从 a2a3 实现 的断言可以看出其核心约束src0与src1中至少有一个源 Tile 的有效区域与dst完全相等另一个源的有效区域在两个维度上都不能超过dst。这是该指令唯一被正式支持的部分有效模式。汇编语法与两级指令形式TPARTADD 与其他 PTO 指令一样在文档中给出同步形式与两个汇编抽象层级AS Level的表达同步形式PTO 汇编%dst tpartadd %src0, %src1 : !pto.tile... - !pto.tile...AS Level 1SSA 形式SSA 形式使用显式的pto.前缀与函数调用风格两个源操作数均为!pto.tile...%dst pto.tpartadd %src0, %src1 : (!pto.tile..., !pto.tile...) - !pto.tile...AS Level 2DPS 形式DPSData Parallel Style形式将输入与输出分开声明操作数类型为!pto.tile_buf...pto.tpartadd ins(%src0, %src1 : !pto.tile_buf..., !pto.tile_buf...) outs(%dst : !pto.tile_buf...)在自动模式下指令的 Tile 资源放置与调度由编译器/运行时完成在手动模式下需要在发射指令前先通过pto.tassign将 Tile 显式绑定到地址例如# 手动模式先显式绑定资源再发射指令 # pto.tassign %arg0, tile(0x1000) # pto.tassign %arg1, tile(0x2000) %dst pto.tpartadd %src0, %src1 : (!pto.tile..., !pto.tile...) - !pto.tile...C 内建接口TPARTADD 的 C 内建接口声明于 include/pto/common/pto_instr.hpp公共包含头为pto/pto-inst.hpptemplate typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents PTO_INST RecordEvent TPARTADD(TileDataDst dst, TileDataSrc0 src0, TileDataSrc1 src1, WaitEvents ... events);接口要点模板参数TileDataDst、TileDataSrc0、TileDataSrc1为三个 Tile 的编译期类型由TileTileType::Vec, T, R, C, BLayout, RowStride, ColStride描述WaitEvents...为可选的事件依赖参数用于异步流水场景下声明该指令依赖的先前事件。返回值PTO_INST RecordEvent即该指令完成后发出的事件可继续作为下游指令如TSTORE的WaitEvents传入。运行时语义实际运算逻辑通过平台相关的TPARTADD_IMPL展开详见下文平台实现一节接口本身由宏PTO_INST标记为内建指令。在 A5 平台的 tpartadd 测试用例 中可以看到事件依赖的典型用法EventOp::TLOAD, Op::TPARTADD event0; EventOp::TPARTADD, Op::TSTORE_VEC event1; TLOAD(src0Tile, src0Global); event0 TLOAD(src1Tile, src1Global); event1 TPARTADDTileDataDst, TileDataSrc0, TileDataSrc1(dstTile, src0Tile, src1Tile, event0); TSTORE(dstGlobal, dstTile, event1);这样TPARTADD会等待TLOAD完成而TSTORE会等待TPARTADD完成从而在无同步屏障的手动模式下保证数据依赖正确。有效区域语义与约束详解通用约束根据 TPARTADD_zh.md 与实现代码TPARTADD 的通用约束如下dst、src0、src1的元素类型必须一致否则编译期报错。A2A3 实现中的static_assert直接给出错误信息TPARTADD src and dst data type is different!见 TPartAdd.hpp。目标有效区域定义结果的计算范围。对目标有效区域内的每个元素两个输入都有效则执行加法仅一个有效则拷贝该输入。若dst的有效区域为零行或列任一为 0指令直接返回不产生任何计算。A2A3 实现中对应if (dstValidRow 0 || dstValidCol 0) return;TPartAdd.hppA5 的 TPartOp 同样有该早退检查。支持的部分有效区域模式要求至少一个源 Tile 的有效区域与dst完全一致另一个源 Tile 的有效区域在两个维度上均不能超过dst。超出该范围的有效区域组合行为由具体实现定义。部分有效模式在源码中的体现A2A3 实现的 TPartInstr 对三种受支持的情形分派src1的行数小于dstsrc1ValidRow dstValidRow且列相等先对src1ValidRow行执行逐元素加法再通过TPartCopyInstr把src0的剩余行拷贝到dstsrc1的列数小于dst列小于且行不超过先整体拷贝src0到dst再对src1ValidCol列在 dst 上原地累加此时src1与dst共同作为 TPartOps 的两个输入src0 src1 dst三者有效区域完全一致直接执行全区域加法。当两个输入有效区域均小于 dst即都不与 dst 相等时PTO_ASSERT会触发编译期/运行期检查At most one entry in the valid-rows and valid-cols of src0 and src1 is smaller than dst见 TPartOp.hpp。A2A3 实现检查Atlas A2/A3 训练与推理系列支持的元素类型int32_t、int16_t、half、float。dst、src0、src1必须全部为行主序isRowMajor否则static_assert报错 TPARTADD not supported BLayout type.见 TPartAdd.hpp。底层映射到向量加法指令vadd(dst, src0, src1, repeats, 1, 1, 1, dstRepeatStride, src0RepeatStride, src1RepeatStride)其中repeats由有效区域按elementsPerRepeat单次 repeat 可处理的元素数与REPEAT_MAX单条指令最大 repeat 数切分而来当行跨度超过REPEAT_STRIDE_MAX或小区域处理更划算时会退化为PartCountModecounter 计数模式见 TPartOp.hpp 与 TPartAdd.hpp。Ascend 950PR / Ascend 950DTA5实现检查支持的元素类型更广uint8_t、int8_t、uint16_t、int16_t、uint32_t、int32_t、int64_t、uint64_t、half、float、bfloat16_t。A5 实现位于 include/pto/npu/a5/TPartAdd.hpp其static_assert逐项列出了上述 11 种类型TPartAdd.hpp。64 位整数特殊路径当元素类型为int64_t/uint64_t时A5 无法用单条向量加法指令直接处理会走Int64PartInt64Op::Add, ...专用路径将 64 位数据拆分为低 32 位与高 32 位两部分分别加载、运算、交织写回见 TPartBinOps.hpp 的Int64PartCalcRegs/Int64PartSameStrideRepeat与 TPartAdd.hpp。该路径针对三者 stride 一致Int64PartSameStride与一般性有效区域Int64PartGeneral分别优化其中一般性路径通过谓词寄存器逐元素判断src0/src1的有效性并在重叠区用加法结果覆盖、非重叠区用有效输入填充完整实现了文档定义的三分支语义TPartBinOps.hpp。非 64 位类型走TPARTOP_IMPLPartAddOpT, ...PartAddOp将二元运算定义为vadd(dst, src0, src1, preg, MODE_ZEROING)TPartAdd.hpp上层 TPartOp 按重叠行/列区间先做加法、超出部分从较大源拷贝的流程组织TPartProcRow负责单行内先计算重叠区再搬运剩余区TPartCopySrc负责整体拷贝。从源码结构看A5 的非 64 位路径还支持VFImplKind向量融合实现版本参数用于选择不同的底层实现变体。CPU 模拟实现CPU 侧用于调试与 golden 比对的实现位于 include/pto/cpu/TPartAdd.hpp其核心运算退化为最直观的标量表达dst[DstOffset] src0[Src0Offset] src1[Src1Offset];CPU 实现不做类型与布局的编译期检查直接按三者的有效区域GetValidRow()/GetValidCol()驱动TPartInstr逐元素执行是理解指令语义最直接的可读参考。编程示例以下示例均来自 TPARTADD_zh.md可直接编译运行。自动模式Auto自动模式下Tile 的缓冲区地址与调度由编译器/运行时负责代码只需声明 Tile 并调用指令#include pto/pto-inst.hpp using namespace pto; void example_auto() { using TileT TileTileType::Vec, float, 16, 16; TileT src0, src1, dst; TPARTADD(dst, src0, src1); }手动模式Manual手动模式下需先用TASSIGN显式为 Tile 绑定地址再发射指令#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); TPARTADD(dst, src0, src1); }汇编形式对照自动模式编译器/运行时负责资源放置与调度%dst pto.tpartadd %src0, %src1 : (!pto.tile..., !pto.tile...) - !pto.tile...手动模式先显式绑定资源再发射指令# 可选当该指令包含 tile 操作数时 # pto.tassign %arg0, tile(0x1000) # pto.tassign %arg1, tile(0x2000) %dst pto.tpartadd %src0, %src1 : (!pto.tile..., !pto.tile...) - !pto.tile...PTO 汇编同步形式 AS Level 2 DPS 形式%dst tpartadd %src0, %src1 : !pto.tile... - !pto.tile... pto.tpartadd ins(%src0, %src1 : !pto.tile_buf..., !pto.tile_buf...) outs(%dst : !pto.tile_buf...)端到端使用流程与测试验证在真实算子中TPARTADD 通常与TLOAD/TSTORE/TASSIGN组合使用形成全局内存加载 → 部分加法 → 全局内存回写的完整链路。A2A3 的 tpartadd_kernel.cpp 给出了完整模板__global__ AICORE void runTPartAdd(__gm__ T __out__* out, __gm__ T __in__* src0, __gm__ T __in__* src1) { using GlobalDataDst GlobalTensorT, Shape1, 1, 1, dstVR, dstVC, Stride1, 1, 1, dstVC, 1; using TileDataDst TileTileType::Vec, T, dstVR, dstVC, BLayout::RowMajor, -1, -1; TileDataDst dstTile(dstVR, dstVC); TASSIGN(src0Tile, 0x0); TASSIGN(src1Tile, 0x10000); TASSIGN(dstTile, 0x20000); TLOAD(src0Tile, src0Global); TLOAD(src1Tile, src1Global); // 手动模式下需显式插入 MTE2 - V 的流水屏障 TPARTADDTileDataDst, TileDataSrc0, TileDataSrc1(dstTile, src0Tile, src1Tile); TSTORE(dstGlobal, dstTile); }该测试通过模板实例化覆盖了多种有效区域组合例如LaunchTPartAddfloat, 64, 64, 64, 64, 64, 64三者有效区域完全一致64×64LaunchTPartAddfloat, 64, 64, 8, 64, 64, 64src0仅 8 行有效LaunchTPartAddfloat, 64, 64, 64, 8, 64, 64、LaunchTPartAddfloat, 64, 64, 64, 64, 8, 64src0/src1仅 8 列有效LaunchTPartAddint16_t, 8, 48, 8, 48, 8, 16src1列数小于dst48 vs 16对应列不匹配分支LaunchTPartAddaclFloat16, 8, 768, 8, 512, 8, 768src0列数小于dst768 vs 512。这些实例化覆盖了文档定义的全部受支持部分有效模式行小于、列小于、完全一致。CPU 侧测试 tests/cpu/st/testcase/tpartadd/main.cpp 通过 gtest 框架对float、int32_t、int64_t、uint64_t、int16_t、uint16_t、uint32_t以及开启CPU_SIM_BFLOAT_ENABLED时的bfloat16_t逐一执行 64×64 目标区域、src1有效区域 32×32 的部分加法并与 golden 数据做逐元素比对kEpsilon 0.0f即要求完全一致。A5 侧还额外提供了runTPartAddInplace原地in-place测试用例验证dst与src0共用同一缓冲区时结果仍然正确tpartadd_kernel.cpp。使用建议与注意事项遵循有效区域约束尽量保证src0与dst有效区域完全一致仅让src1或反之在单一维度上小于目标这是所有平台都支持的最安全用法两个源都小于目标的组合会触发断言。类型与布局对齐A2A3 上仅支持int32_t/int16_t/half/float且必须行主序A5 支持类型更广含bfloat16_t与 64 位整型但 64 位整型走独立实现路径。跨平台移植时以目标平台的类型表为准。零有效区域dst有效区域为 0 时指令直接返回可作为条件分支的替代写法。异步流水在手动模式下务必通过事件依赖或显式流水屏障如 A2A3 测试中的set_flag(PIPE_MTE2, PIPE_V, ...)/wait_flag(...)或 A5 测试中的EventOp::TLOAD, Op::TPARTADD保证TLOAD结果对向量指令可见再让TSTORE等待向量指令完成。延伸阅读指令总览与操作数约定PTO-Virtual-ISA-Manual_zh.md 与 PTOISA_zh.md相邻逐元素运算指令TADD_zh.md全区域加法、TMUL_zh.md、TMAX_zh.mdTile 编程模型与有效区域概念Tile_zh.md、ProgrammingModel_zh.md指令内建接口声明include/pto/common/pto_instr.hpp公共头文件 include/pto/pto-inst.hpp平台实现与测试include/pto/npu/a2a3/TPartAdd.hpp、include/pto/npu/a5/TPartAdd.hpp、include/pto/cpu/TPartAdd.hpp以及 tests/npu/a2a3/src/st/testcase/tpartadd/、tests/npu/a5/src/st/testcase/tpartadd/、tests/cpu/st/testcase/tpartadd/ 三套测试用例。【免费下载链接】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），仅供参考 PTO-ISA TPARTADD 指令详解：部分有效区域逐元素加法（tpartadd）的语义、接口与平台实现
PTO-ISA TPARTADD 指令详解部分有效区域逐元素加法tpartadd的语义、接口与平台实现【免费下载链接】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导读本文以 CANN pto-isa 开源仓库的指令文档 TPARTADD_zh.md 为骨架系统讲解TPARTADDtpartadd这条 Tile 级向量指令它面向目标有效区域valid region执行逐元素加法支持两个源 Tile 的有效区域与目标不一致的部分有效场景。读完本文你将掌握 TPARTADD 的数学语义、SSA/DPS 两级汇编形式、C 内建接口签名、A2A3 与 Ascend 950A5平台的实现差异与类型约束并能结合 TPartAdd.hpp、TPartOp.hpp、TPartBinOps.hpp 等源码与 tpartadd 测试用例 写出可运行、可验证的算子代码。指令概述与数学语义TPARTADD 是 PTO 指令集中用于部分有效区域逐元素加法的向量指令全称 Tile Partial Add。它不属于普通的TADD要求各 Tile 有效区域完全一致而是允许两个源 Tilesrc0、src1与目标 Tiledst的有效区域存在差异加法只在 dst 的有效区域内、且两个输入同时有效的元素上执行。语义定义在目标有效区域内的每个元素(i, j)结果按以下规则确定若src0与src1在该位置均有效dst[i][j] src0[i][j] src1[i][j]若仅src0有效dst[i][j] src0[i][j]直接拷贝若仅src1有效dst[i][j] src1[i][j]直接拷贝。其余有效区域不匹配的情形例如某个输入在目标有效区域内既不是完全有效也不是完全无效由具体实现定义编程时应当避免依赖这类行为。这一语义可以形式化为dst[i][j] src0[i][j] src1[i][j] 若 src0、src1 在 (i,j) 均有效 src0[i][j] 若仅 src0 在 (i,j) 有效 src1[i][j] 若仅 src1 在 (i,j) 有效指令示意图如下直观展示了两个输入有效区域叠加到目标区域的过程与 TADD 的定位差异普通TADD要求参与运算的 Tile 有效区域完全对齐而 TPARTADD 解决的是目标区域更大、某个输入只有部分区域有数据的场景。从 a2a3 实现 的断言可以看出其核心约束src0与src1中至少有一个源 Tile 的有效区域与dst完全相等另一个源的有效区域在两个维度上都不能超过dst。这是该指令唯一被正式支持的部分有效模式。汇编语法与两级指令形式TPARTADD 与其他 PTO 指令一样在文档中给出同步形式与两个汇编抽象层级AS Level的表达同步形式PTO 汇编%dst tpartadd %src0, %src1 : !pto.tile... - !pto.tile...AS Level 1SSA 形式SSA 形式使用显式的pto.前缀与函数调用风格两个源操作数均为!pto.tile...%dst pto.tpartadd %src0, %src1 : (!pto.tile..., !pto.tile...) - !pto.tile...AS Level 2DPS 形式DPSData Parallel Style形式将输入与输出分开声明操作数类型为!pto.tile_buf...pto.tpartadd ins(%src0, %src1 : !pto.tile_buf..., !pto.tile_buf...) outs(%dst : !pto.tile_buf...)在自动模式下指令的 Tile 资源放置与调度由编译器/运行时完成在手动模式下需要在发射指令前先通过pto.tassign将 Tile 显式绑定到地址例如# 手动模式先显式绑定资源再发射指令 # pto.tassign %arg0, tile(0x1000) # pto.tassign %arg1, tile(0x2000) %dst pto.tpartadd %src0, %src1 : (!pto.tile..., !pto.tile...) - !pto.tile...C 内建接口TPARTADD 的 C 内建接口声明于 include/pto/common/pto_instr.hpp公共包含头为pto/pto-inst.hpptemplate typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents PTO_INST RecordEvent TPARTADD(TileDataDst dst, TileDataSrc0 src0, TileDataSrc1 src1, WaitEvents ... events);接口要点模板参数TileDataDst、TileDataSrc0、TileDataSrc1为三个 Tile 的编译期类型由TileTileType::Vec, T, R, C, BLayout, RowStride, ColStride描述WaitEvents...为可选的事件依赖参数用于异步流水场景下声明该指令依赖的先前事件。返回值PTO_INST RecordEvent即该指令完成后发出的事件可继续作为下游指令如TSTORE的WaitEvents传入。运行时语义实际运算逻辑通过平台相关的TPARTADD_IMPL展开详见下文平台实现一节接口本身由宏PTO_INST标记为内建指令。在 A5 平台的 tpartadd 测试用例 中可以看到事件依赖的典型用法EventOp::TLOAD, Op::TPARTADD event0; EventOp::TPARTADD, Op::TSTORE_VEC event1; TLOAD(src0Tile, src0Global); event0 TLOAD(src1Tile, src1Global); event1 TPARTADDTileDataDst, TileDataSrc0, TileDataSrc1(dstTile, src0Tile, src1Tile, event0); TSTORE(dstGlobal, dstTile, event1);这样TPARTADD会等待TLOAD完成而TSTORE会等待TPARTADD完成从而在无同步屏障的手动模式下保证数据依赖正确。有效区域语义与约束详解通用约束根据 TPARTADD_zh.md 与实现代码TPARTADD 的通用约束如下dst、src0、src1的元素类型必须一致否则编译期报错。A2A3 实现中的static_assert直接给出错误信息TPARTADD src and dst data type is different!见 TPartAdd.hpp。目标有效区域定义结果的计算范围。对目标有效区域内的每个元素两个输入都有效则执行加法仅一个有效则拷贝该输入。若dst的有效区域为零行或列任一为 0指令直接返回不产生任何计算。A2A3 实现中对应if (dstValidRow 0 || dstValidCol 0) return;TPartAdd.hppA5 的 TPartOp 同样有该早退检查。支持的部分有效区域模式要求至少一个源 Tile 的有效区域与dst完全一致另一个源 Tile 的有效区域在两个维度上均不能超过dst。超出该范围的有效区域组合行为由具体实现定义。部分有效模式在源码中的体现A2A3 实现的 TPartInstr 对三种受支持的情形分派src1的行数小于dstsrc1ValidRow dstValidRow且列相等先对src1ValidRow行执行逐元素加法再通过TPartCopyInstr把src0的剩余行拷贝到dstsrc1的列数小于dst列小于且行不超过先整体拷贝src0到dst再对src1ValidCol列在 dst 上原地累加此时src1与dst共同作为 TPartOps 的两个输入src0 src1 dst三者有效区域完全一致直接执行全区域加法。当两个输入有效区域均小于 dst即都不与 dst 相等时PTO_ASSERT会触发编译期/运行期检查At most one entry in the valid-rows and valid-cols of src0 and src1 is smaller than dst见 TPartOp.hpp。A2A3 实现检查Atlas A2/A3 训练与推理系列支持的元素类型int32_t、int16_t、half、float。dst、src0、src1必须全部为行主序isRowMajor否则static_assert报错 TPARTADD not supported BLayout type.见 TPartAdd.hpp。底层映射到向量加法指令vadd(dst, src0, src1, repeats, 1, 1, 1, dstRepeatStride, src0RepeatStride, src1RepeatStride)其中repeats由有效区域按elementsPerRepeat单次 repeat 可处理的元素数与REPEAT_MAX单条指令最大 repeat 数切分而来当行跨度超过REPEAT_STRIDE_MAX或小区域处理更划算时会退化为PartCountModecounter 计数模式见 TPartOp.hpp 与 TPartAdd.hpp。Ascend 950PR / Ascend 950DTA5实现检查支持的元素类型更广uint8_t、int8_t、uint16_t、int16_t、uint32_t、int32_t、int64_t、uint64_t、half、float、bfloat16_t。A5 实现位于 include/pto/npu/a5/TPartAdd.hpp其static_assert逐项列出了上述 11 种类型TPartAdd.hpp。64 位整数特殊路径当元素类型为int64_t/uint64_t时A5 无法用单条向量加法指令直接处理会走Int64PartInt64Op::Add, ...专用路径将 64 位数据拆分为低 32 位与高 32 位两部分分别加载、运算、交织写回见 TPartBinOps.hpp 的Int64PartCalcRegs/Int64PartSameStrideRepeat与 TPartAdd.hpp。该路径针对三者 stride 一致Int64PartSameStride与一般性有效区域Int64PartGeneral分别优化其中一般性路径通过谓词寄存器逐元素判断src0/src1的有效性并在重叠区用加法结果覆盖、非重叠区用有效输入填充完整实现了文档定义的三分支语义TPartBinOps.hpp。非 64 位类型走TPARTOP_IMPLPartAddOpT, ...PartAddOp将二元运算定义为vadd(dst, src0, src1, preg, MODE_ZEROING)TPartAdd.hpp上层 TPartOp 按重叠行/列区间先做加法、超出部分从较大源拷贝的流程组织TPartProcRow负责单行内先计算重叠区再搬运剩余区TPartCopySrc负责整体拷贝。从源码结构看A5 的非 64 位路径还支持VFImplKind向量融合实现版本参数用于选择不同的底层实现变体。CPU 模拟实现CPU 侧用于调试与 golden 比对的实现位于 include/pto/cpu/TPartAdd.hpp其核心运算退化为最直观的标量表达dst[DstOffset] src0[Src0Offset] src1[Src1Offset];CPU 实现不做类型与布局的编译期检查直接按三者的有效区域GetValidRow()/GetValidCol()驱动TPartInstr逐元素执行是理解指令语义最直接的可读参考。编程示例以下示例均来自 TPARTADD_zh.md可直接编译运行。自动模式Auto自动模式下Tile 的缓冲区地址与调度由编译器/运行时负责代码只需声明 Tile 并调用指令#include pto/pto-inst.hpp using namespace pto; void example_auto() { using TileT TileTileType::Vec, float, 16, 16; TileT src0, src1, dst; TPARTADD(dst, src0, src1); }手动模式Manual手动模式下需先用TASSIGN显式为 Tile 绑定地址再发射指令#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); TPARTADD(dst, src0, src1); }汇编形式对照自动模式编译器/运行时负责资源放置与调度%dst pto.tpartadd %src0, %src1 : (!pto.tile..., !pto.tile...) - !pto.tile...手动模式先显式绑定资源再发射指令# 可选当该指令包含 tile 操作数时 # pto.tassign %arg0, tile(0x1000) # pto.tassign %arg1, tile(0x2000) %dst pto.tpartadd %src0, %src1 : (!pto.tile..., !pto.tile...) - !pto.tile...PTO 汇编同步形式 AS Level 2 DPS 形式%dst tpartadd %src0, %src1 : !pto.tile... - !pto.tile... pto.tpartadd ins(%src0, %src1 : !pto.tile_buf..., !pto.tile_buf...) outs(%dst : !pto.tile_buf...)端到端使用流程与测试验证在真实算子中TPARTADD 通常与TLOAD/TSTORE/TASSIGN组合使用形成全局内存加载 → 部分加法 → 全局内存回写的完整链路。A2A3 的 tpartadd_kernel.cpp 给出了完整模板__global__ AICORE void runTPartAdd(__gm__ T __out__* out, __gm__ T __in__* src0, __gm__ T __in__* src1) { using GlobalDataDst GlobalTensorT, Shape1, 1, 1, dstVR, dstVC, Stride1, 1, 1, dstVC, 1; using TileDataDst TileTileType::Vec, T, dstVR, dstVC, BLayout::RowMajor, -1, -1; TileDataDst dstTile(dstVR, dstVC); TASSIGN(src0Tile, 0x0); TASSIGN(src1Tile, 0x10000); TASSIGN(dstTile, 0x20000); TLOAD(src0Tile, src0Global); TLOAD(src1Tile, src1Global); // 手动模式下需显式插入 MTE2 - V 的流水屏障 TPARTADDTileDataDst, TileDataSrc0, TileDataSrc1(dstTile, src0Tile, src1Tile); TSTORE(dstGlobal, dstTile); }该测试通过模板实例化覆盖了多种有效区域组合例如LaunchTPartAddfloat, 64, 64, 64, 64, 64, 64三者有效区域完全一致64×64LaunchTPartAddfloat, 64, 64, 8, 64, 64, 64src0仅 8 行有效LaunchTPartAddfloat, 64, 64, 64, 8, 64, 64、LaunchTPartAddfloat, 64, 64, 64, 64, 8, 64src0/src1仅 8 列有效LaunchTPartAddint16_t, 8, 48, 8, 48, 8, 16src1列数小于dst48 vs 16对应列不匹配分支LaunchTPartAddaclFloat16, 8, 768, 8, 512, 8, 768src0列数小于dst768 vs 512。这些实例化覆盖了文档定义的全部受支持部分有效模式行小于、列小于、完全一致。CPU 侧测试 tests/cpu/st/testcase/tpartadd/main.cpp 通过 gtest 框架对float、int32_t、int64_t、uint64_t、int16_t、uint16_t、uint32_t以及开启CPU_SIM_BFLOAT_ENABLED时的bfloat16_t逐一执行 64×64 目标区域、src1有效区域 32×32 的部分加法并与 golden 数据做逐元素比对kEpsilon 0.0f即要求完全一致。A5 侧还额外提供了runTPartAddInplace原地in-place测试用例验证dst与src0共用同一缓冲区时结果仍然正确tpartadd_kernel.cpp。使用建议与注意事项遵循有效区域约束尽量保证src0与dst有效区域完全一致仅让src1或反之在单一维度上小于目标这是所有平台都支持的最安全用法两个源都小于目标的组合会触发断言。类型与布局对齐A2A3 上仅支持int32_t/int16_t/half/float且必须行主序A5 支持类型更广含bfloat16_t与 64 位整型但 64 位整型走独立实现路径。跨平台移植时以目标平台的类型表为准。零有效区域dst有效区域为 0 时指令直接返回可作为条件分支的替代写法。异步流水在手动模式下务必通过事件依赖或显式流水屏障如 A2A3 测试中的set_flag(PIPE_MTE2, PIPE_V, ...)/wait_flag(...)或 A5 测试中的EventOp::TLOAD, Op::TPARTADD保证TLOAD结果对向量指令可见再让TSTORE等待向量指令完成。延伸阅读指令总览与操作数约定PTO-Virtual-ISA-Manual_zh.md 与 PTOISA_zh.md相邻逐元素运算指令TADD_zh.md全区域加法、TMUL_zh.md、TMAX_zh.mdTile 编程模型与有效区域概念Tile_zh.md、ProgrammingModel_zh.md指令内建接口声明include/pto/common/pto_instr.hpp公共头文件 include/pto/pto-inst.hpp平台实现与测试include/pto/npu/a2a3/TPartAdd.hpp、include/pto/npu/a5/TPartAdd.hpp、include/pto/cpu/TPartAdd.hpp以及 tests/npu/a2a3/src/st/testcase/tpartadd/、tests/npu/a5/src/st/testcase/tpartadd/、tests/cpu/st/testcase/tpartadd/ 三套测试用例。【免费下载链接】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），仅供参考