人工智能指令集算子库CANNAscend【免费下载链接】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点击查看免费下载TCOLEXPANDSUB 是 CANN PTOParallel Tile Operation虚拟指令集中用于列广播减法的核心向量指令它将每行中的每个元素减去一个每列一个标量的向量常与 TCOLEXPAND 系列指令配合用于逐列统计后的去均值center等 tile 级运算。本文以 docs/isa/TCOLEXPANDSUB.md 为骨架结合include/pto/下的 A2A3/A5 后端实现与tests/中的 ST 测试用例完整讲解其数学语义、汇编语法PTO Assembly / AS Level 1 SSA / AS Level 2 DPS、C 内建接口、类型与布局约束、64 位元素模拟机制以及底层向量指令调用链读者可在掌握语法后直接编写可运行的 PTO kernel。指令概述TCOLEXPANDSUB 的语义是列广播减法column-wise broadcast subtractsrc0是一个R × C的 Tilesrc1是一个长度为C每列一个标量的向量指令把src1中第j列的值s_j广播到每一行然后计算dst[i][j] src0[i][j] - s_j。其指令示意图如下来源docs/figures/isa/TCOLEXPANDSUB.svg它与同族指令的关系PTO 提供了一整组TCOLEXPAND系列二元运算指令ADD/DIV/MAX/MIN/MUL/SUB/EXPDIF 等底层共享同一套列广播二元运算模板见下文实现原理小节其中减法即 TCOLEXPANDSUB。TCOLEXPAND单操作数负责把源列首元素广播到整个目标列而TCOLEXPANDSUB等二元指令则在广播的同时完成逐元素运算常被组合用于归一化类算子。数学语义设R dst.GetValidRow()有效行数、C dst.GetValidCol()有效列数s_j为从src1中取出的第j列的标量每列一个值。指令的计算公式为$$ \mathrm{dst}{i,j} \mathrm{src0}{i,j} - s_j \qquad (0 \le i R,\ 0 \le j C) $$要点广播发生在列维度上s_j不随行号i变化同一列的所有行共享同一个减数src1只提供C个有效值一行、C列实现层面会为每一行重复读取这一行向量结果dst的形状与src0一致为R × C。汇编语法TCOLEXPANDSUB 的汇编书写分为三个层级PTO 汇编同步形式、AS Level 1SSA 形式与 AS Level 2DPS 形式。PTO 汇编同步形式%dst tcolexpandsub %src0, %src1 : !pto.tile..., !pto.tile... - !pto.tile...AS Level 1SSA 形式%dst pto.tcolexpandsub %src0, %src1 : !pto.tile..., !pto.tile... - !pto.tile...AS Level 2DPS 形式pto.tcolexpandsub ins(%src0, %src1 : !pto.tile_buf..., !pto.tile_buf...) outs(%dst : !pto.tile_buf...)三个层级的差异体现在操作数描述粒度上PTO 汇编与 AS Level 1 使用 SSA 值%dst/%src0/%src1与!pto.tile...类型描述AS Level 2 则显式区分输入ins与输出outs并落到!pto.tile_buf...的 buffer 类型对应更接近硬件资源视图的表示。同族指令TCOLEXPANDADD 等的语法结构完全一致仅指令名不同。C 内建接口指令通过 C 内建函数暴露声明位于 include/pto/common/pto_instr.hpp。公共包含头为pto/pto-inst.hpp内部声明位于pto/common/pto_instr.hpp。template typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents PTO_INST RecordEvent TCOLEXPANDSUB(TileDataDst dst, TileDataSrc0 src0, TileDataSrc1 src1, WaitEvents ... events);接口要点可从源码 pto_instr.hpp 中 TCOLEXPANDSUB 的实现 印证返回RecordEvent可用于指令间事件同步变参模板WaitEvents...支持传入事件对象实现先等待、再执行的依赖语义内部实现先调用detail::PtoWaitEvents(events...)等待传入事件再通过MAP_INSTR_IMPL(TCOLEXPANDSUB, dst, src0, src1)宏分派到平台相关实现三个 Tile 操作数模板参数分别对应dst、src0、src1与汇编形式中的操作数顺序一致。自动模式与手动模式自动模式Auto Mode资源放置与调度由编译器/运行时托管直接书写 SSA 指令即可# Auto mode: compiler/runtime-managed placement and scheduling. %dst pto.tcolexpandsub %src0, %src1 : !pto.tile..., !pto.tile... - !pto.tile...手动模式Manual Mode发射指令前需先显式绑定资源可通过pto.tassign将参数绑定到指定 tile 地址可选当指令包含 tile 操作数时# 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.tcolexpandsub %src0, %src1 : !pto.tile..., !pto.tile... - !pto.tile...两种模式仅影响资源绑定与调度方式指令语义不变。约束Constraints使用 TCOLEXPANDSUB 必须满足以下约束数据类型TileDataDst::DType与TileDataSrc1::DType必须是以下类型之一通用类型适用于 A2、A3 与 A5 平台half、float、int16、int32A5 专属扩展类型uint16、uint32、bfloat16_t、int8、uint8、int64、uint64。布局约束编译期TileDataDst::isRowMajor必须为真即目标 Tile 必须是行主序布局。src1形状src1预期提供每列一个标量即其有效形状必须覆盖C个值通常为1 × C的 Vec Tile。平台特定约束确切的布局/分形fractal约束是目标平台相关的参见 include/pto/npu/a2a3/TColExpand*.hpp 与 include/pto/npu/a5/TColExpand*.hpp 下的后端头文件。这些约束在源码中有对应的编译期检查。例如 a2a3/TColExpandBinOp.hpp 中的TCOLEXPANDOP_IMPL通过static_assert限定int32_t、int、int16_t、half、float16_t、float、float32_t等类型并断言TileData::isRowMajor否则直接编译失败错误提示Fix: TCOLEXPANDOP Invalid data type.与Fix: TCOLEXPANDOP not supported Layout typea5/TColExpandBinOp.hpp 则额外允许int64_t、uint64_t、uint32_t、uint16_t、bfloat16_t、int8_t、uint8_t与文档中的 A5 扩展类型一致。64 位元素类型A5 专属int64/uint64仅在 A5Ascend 950PR/Ascend 950DT上受支持其实现有两个关键特性模拟执行A5 没有原生的 64 位向量运算单元指令通过一对 32 位寄存器分别保存每个元素的低 32 位与高 32 位模拟实现每列标量操作数与全尺寸操作数使用相同的解交织de-interleaved布局读取。精度与对齐计算结果为精确的 64 位补码值Tile 对齐遵循 64 位元素的通用规则——RowMajor 的 Tile 要求Cols % 4 0物理列数需为 4 的倍数有效列数不要求对齐。源码级实现原理TCOLEXPANDSUB 在不同平台上有不同的底层向量指令路径均通过模板统一分派理解实现有助于把握性能特征与边界条件。A2/A3 平台基于 vsub 的重复计算A2/A3 后端实现在 include/pto/npu/a2a3/TColExpandSub.hpptemplate typename T struct ColExpandSubOp { PTO_INTERNAL static void ColExpandBinInstr(__ubuf__ T* dst, __ubuf__ T* src0, __ubuf__ T* src1, uint8_t repeats) { vsub(dst, src0, src1, repeats, 1, 1, 1, 8, 8, 8); } ... };其核心是向量减指令vsub并定义了两组算子ColExpandSubOp与ColExpandSubOp2后者交换src0/src1顺序用于处理操作数形状互换的情形。真正的广播循环在 a2a3/TColExpandBinOp.hpp 的TCOLEXPANDOP_IMPL中完成首先比较src0/src1与dst的有效形状是否一致src0eqdst/src1eqdst从而决定使用哪个算子、哪个操作数作为每列标量TColExpandBinOp.hpp#L107-L120若目标 tile 连续Cols ValidCol或Rows 1走TColExpandBinaryNormMode以一次向量指令配合 repeat 步长repeat stride完成多行广播否则走TColExpandBinaryCountMode逐行循环并配合SetVectorCount设置向量计数TColExpandBinOp.hpp#L76-L82。A5 平台寄存器张量 谓词掩码含 64 位模拟A5 后端实现在 include/pto/npu/a5/TColExpandSub.hpp其标量路径为template typename T struct ColExpandSubOp { PTO_INTERNAL static void ColExpandBinaryInstr( RegTensorT reg_dst, RegTensorT reg_src0, RegTensorT reg_src1, MaskReg preg) { vsub(reg_dst, reg_src0, reg_src1, preg, MODE_ZEROING); } ... };差异点在于 A5 使用RegTensor寄存器张量MaskReg谓词掩码MODE_ZEROING的编程模型并在 a5/TColExpandBinOp.hpp 中提供1D/2D与PostUpdate/NoPostUpdate多种实现变体VFImplKind可选默认走VFIMPL_DEFAULT连续 tile 使用 PostUpdate 版本以复用地址自增特性。64 位路径则由Int64ColExpandBinary承担a5/TColExpandBinOp.hpp#L155-L219每个 64 位元素拆为(low, high)两个 32 位寄存器通过Int64LoadBounded按列偏移读取Int64BinaryCalcRegsInt64Op::Sub, T完成低/高字对的减法运算最后用vintlv交织回写并配合pintlv_b32生成的低/高掩码分别存储。对于非 A5/A6 目标该函数仅有声明式 stub编译期不可用。使用示例C kernel 示例A5 平台测试用例以下示例取自 tests/npu/a5/src/st/testcase/tcolexpandsub/tcolexpandsub_kernel.cpp展示了完整的数据搬运—计算—存回流程#include pto/pto-inst.hpp using namespace pto; template typename T, uint32_t dstRow, uint32_t dstCol, uint32_t src1Row, uint32_t src1Col __global__ AICORE void runCOLEXPANDSUB(__gm__ T __out__* out, __gm__ T __in__* src0, __gm__ T __in__* src1) { using TileData TileTileType::Vec, T, src1Row, src1Col, BLayout::RowMajor, -1, -1; using DstTileData TileTileType::Vec, T, dstRow, dstCol, BLayout::RowMajor, -1, -1; DstTileData src0Tile(dstRow, dstCol); TileData src1Tile(src1Row, src1Col); DstTileData dstTile(dstRow, dstCol); TASSIGN(src0Tile, 0x0); TASSIGN(src1Tile, 0x10000); TASSIGN(dstTile, 0x20000); TLOAD(src0Tile, src0Global); TLOAD(src1Tile, src1Global); // 手动模式下的流水线同步等待 MTE2 搬运完成后向量单元再计算 set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); TCOLEXPANDSUB(dstTile, src0Tile, src1Tile); // 计算完成后等待再触发 MTE3 存回 set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); TSTORE(dstGlobal, dstTile); }该用例验证了src1形状为1 × Csrc1Row 1的典型用法并覆盖多种数据类型与形状组合例如tcolexpandsub_kernel.cpp#L74-L83float6×128、18×32halfaclFloat1610×256、12×64int32_t16×32int16_t16×6464 位类型int64_t、uint64_t16×32对应 A5 的 64 位模拟路径。CPU 参考实现CPU 侧的参考测试位于 tests/cpu/st/testcase/tcolexpandop/tcolexpandop_kernel.cpp其中LaunchTCOLEXPANDSUB以 lambda 方式调用TCOLEXPANDSUB(dst, src0, src1)src1声明为TileTileType::Vec, T, 1, iCol, BLayout::RowMajor, -1, -1即单行向量可用于在没有 NPU 的环境下做语义正确性对照。PTO 汇编完整形式%dst tcolexpandsub %src0, %src1 : !pto.tile..., !pto.tile... - !pto.tile... # AS Level 2 (DPS) pto.tcolexpandsub ins(%src0, %src1 : !pto.tile_buf..., !pto.tile_buf...) outs(%dst : !pto.tile_buf...)相关资源指令文档docs/isa/TCOLEXPANDSUB.md含中文版 TCOLEXPANDSUB_zh.md、示意图 docs/figures/isa/TCOLEXPANDSUB.svg同类单操作数广播指令docs/isa/TCOLEXPAND.mdC 接口声明include/pto/common/pto_instr.hpp公共头pto/pto-inst.hppA2/A3 后端include/pto/npu/a2a3/TColExpandSub.hpp 与 include/pto/npu/a2a3/TColExpandBinOp.hppA5 后端include/pto/npu/a5/TColExpandSub.hpp 与 include/pto/npu/a5/TColExpandBinOp.hpp测试用例A5 ST 用例 tests/npu/a5/src/st/testcase/tcolexpandsub/、CPU 参考用例 tests/cpu/st/testcase/tcolexpandop/。赞分享人工智能指令集算子库CANNAscend【免费下载链接】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点击查看免费下载相关推荐PTO-ISA TCOLEXPANDDIV 指令详解列广播除法Column-wise Broadcast Divide的语义、汇编与 C 编程指南PTO ISA TCOLEXPANDDIV 指令详解列广播除法Column wise Broadcast Divide的语义、汇编与 C 编程指南 T人工智能指令集算子库CANNAscendPTO-ISA TCOLEXPANDDIV 指令详解列广播除法Column-wise Broadcast Divide的语义、编程接口与多平台实现PTO ISA TCOLEXPANDDIV 指令详解列广播除法Column wise Broadcast Divide的语义、编程接口与多平台实现 导读人工智能指令集算子库CANNAscend基于 agno 的音频转写数据标注实战纯文本转写、说话人分离与时间戳分段基于 agno 的音频转写数据标注实战纯文本转写、说话人分离与时间戳分段 音频转写Speech to text是 agno 数据标注体系中音频模态最基础的人工智能指令集算子库CANNAscend上一篇【免费下载】 NVIDIA nvbandwidth 工具使用指南下一篇leak-check深度解析构建隐私保护的BFS图遍历与脱敏聚合系统创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
PTO TCOLEXPANDSUB 指令详解：列广播减法（Column-wise Broadcast Subtract）的语义、三级汇编语法与双平台源码实现
人工智能指令集算子库CANNAscend【免费下载链接】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点击查看免费下载TCOLEXPANDSUB 是 CANN PTOParallel Tile Operation虚拟指令集中用于列广播减法的核心向量指令它将每行中的每个元素减去一个每列一个标量的向量常与 TCOLEXPAND 系列指令配合用于逐列统计后的去均值center等 tile 级运算。本文以 docs/isa/TCOLEXPANDSUB.md 为骨架结合include/pto/下的 A2A3/A5 后端实现与tests/中的 ST 测试用例完整讲解其数学语义、汇编语法PTO Assembly / AS Level 1 SSA / AS Level 2 DPS、C 内建接口、类型与布局约束、64 位元素模拟机制以及底层向量指令调用链读者可在掌握语法后直接编写可运行的 PTO kernel。指令概述TCOLEXPANDSUB 的语义是列广播减法column-wise broadcast subtractsrc0是一个R × C的 Tilesrc1是一个长度为C每列一个标量的向量指令把src1中第j列的值s_j广播到每一行然后计算dst[i][j] src0[i][j] - s_j。其指令示意图如下来源docs/figures/isa/TCOLEXPANDSUB.svg它与同族指令的关系PTO 提供了一整组TCOLEXPAND系列二元运算指令ADD/DIV/MAX/MIN/MUL/SUB/EXPDIF 等底层共享同一套列广播二元运算模板见下文实现原理小节其中减法即 TCOLEXPANDSUB。TCOLEXPAND单操作数负责把源列首元素广播到整个目标列而TCOLEXPANDSUB等二元指令则在广播的同时完成逐元素运算常被组合用于归一化类算子。数学语义设R dst.GetValidRow()有效行数、C dst.GetValidCol()有效列数s_j为从src1中取出的第j列的标量每列一个值。指令的计算公式为$$ \mathrm{dst}{i,j} \mathrm{src0}{i,j} - s_j \qquad (0 \le i R,\ 0 \le j C) $$要点广播发生在列维度上s_j不随行号i变化同一列的所有行共享同一个减数src1只提供C个有效值一行、C列实现层面会为每一行重复读取这一行向量结果dst的形状与src0一致为R × C。汇编语法TCOLEXPANDSUB 的汇编书写分为三个层级PTO 汇编同步形式、AS Level 1SSA 形式与 AS Level 2DPS 形式。PTO 汇编同步形式%dst tcolexpandsub %src0, %src1 : !pto.tile..., !pto.tile... - !pto.tile...AS Level 1SSA 形式%dst pto.tcolexpandsub %src0, %src1 : !pto.tile..., !pto.tile... - !pto.tile...AS Level 2DPS 形式pto.tcolexpandsub ins(%src0, %src1 : !pto.tile_buf..., !pto.tile_buf...) outs(%dst : !pto.tile_buf...)三个层级的差异体现在操作数描述粒度上PTO 汇编与 AS Level 1 使用 SSA 值%dst/%src0/%src1与!pto.tile...类型描述AS Level 2 则显式区分输入ins与输出outs并落到!pto.tile_buf...的 buffer 类型对应更接近硬件资源视图的表示。同族指令TCOLEXPANDADD 等的语法结构完全一致仅指令名不同。C 内建接口指令通过 C 内建函数暴露声明位于 include/pto/common/pto_instr.hpp。公共包含头为pto/pto-inst.hpp内部声明位于pto/common/pto_instr.hpp。template typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents PTO_INST RecordEvent TCOLEXPANDSUB(TileDataDst dst, TileDataSrc0 src0, TileDataSrc1 src1, WaitEvents ... events);接口要点可从源码 pto_instr.hpp 中 TCOLEXPANDSUB 的实现 印证返回RecordEvent可用于指令间事件同步变参模板WaitEvents...支持传入事件对象实现先等待、再执行的依赖语义内部实现先调用detail::PtoWaitEvents(events...)等待传入事件再通过MAP_INSTR_IMPL(TCOLEXPANDSUB, dst, src0, src1)宏分派到平台相关实现三个 Tile 操作数模板参数分别对应dst、src0、src1与汇编形式中的操作数顺序一致。自动模式与手动模式自动模式Auto Mode资源放置与调度由编译器/运行时托管直接书写 SSA 指令即可# Auto mode: compiler/runtime-managed placement and scheduling. %dst pto.tcolexpandsub %src0, %src1 : !pto.tile..., !pto.tile... - !pto.tile...手动模式Manual Mode发射指令前需先显式绑定资源可通过pto.tassign将参数绑定到指定 tile 地址可选当指令包含 tile 操作数时# 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.tcolexpandsub %src0, %src1 : !pto.tile..., !pto.tile... - !pto.tile...两种模式仅影响资源绑定与调度方式指令语义不变。约束Constraints使用 TCOLEXPANDSUB 必须满足以下约束数据类型TileDataDst::DType与TileDataSrc1::DType必须是以下类型之一通用类型适用于 A2、A3 与 A5 平台half、float、int16、int32A5 专属扩展类型uint16、uint32、bfloat16_t、int8、uint8、int64、uint64。布局约束编译期TileDataDst::isRowMajor必须为真即目标 Tile 必须是行主序布局。src1形状src1预期提供每列一个标量即其有效形状必须覆盖C个值通常为1 × C的 Vec Tile。平台特定约束确切的布局/分形fractal约束是目标平台相关的参见 include/pto/npu/a2a3/TColExpand*.hpp 与 include/pto/npu/a5/TColExpand*.hpp 下的后端头文件。这些约束在源码中有对应的编译期检查。例如 a2a3/TColExpandBinOp.hpp 中的TCOLEXPANDOP_IMPL通过static_assert限定int32_t、int、int16_t、half、float16_t、float、float32_t等类型并断言TileData::isRowMajor否则直接编译失败错误提示Fix: TCOLEXPANDOP Invalid data type.与Fix: TCOLEXPANDOP not supported Layout typea5/TColExpandBinOp.hpp 则额外允许int64_t、uint64_t、uint32_t、uint16_t、bfloat16_t、int8_t、uint8_t与文档中的 A5 扩展类型一致。64 位元素类型A5 专属int64/uint64仅在 A5Ascend 950PR/Ascend 950DT上受支持其实现有两个关键特性模拟执行A5 没有原生的 64 位向量运算单元指令通过一对 32 位寄存器分别保存每个元素的低 32 位与高 32 位模拟实现每列标量操作数与全尺寸操作数使用相同的解交织de-interleaved布局读取。精度与对齐计算结果为精确的 64 位补码值Tile 对齐遵循 64 位元素的通用规则——RowMajor 的 Tile 要求Cols % 4 0物理列数需为 4 的倍数有效列数不要求对齐。源码级实现原理TCOLEXPANDSUB 在不同平台上有不同的底层向量指令路径均通过模板统一分派理解实现有助于把握性能特征与边界条件。A2/A3 平台基于 vsub 的重复计算A2/A3 后端实现在 include/pto/npu/a2a3/TColExpandSub.hpptemplate typename T struct ColExpandSubOp { PTO_INTERNAL static void ColExpandBinInstr(__ubuf__ T* dst, __ubuf__ T* src0, __ubuf__ T* src1, uint8_t repeats) { vsub(dst, src0, src1, repeats, 1, 1, 1, 8, 8, 8); } ... };其核心是向量减指令vsub并定义了两组算子ColExpandSubOp与ColExpandSubOp2后者交换src0/src1顺序用于处理操作数形状互换的情形。真正的广播循环在 a2a3/TColExpandBinOp.hpp 的TCOLEXPANDOP_IMPL中完成首先比较src0/src1与dst的有效形状是否一致src0eqdst/src1eqdst从而决定使用哪个算子、哪个操作数作为每列标量TColExpandBinOp.hpp#L107-L120若目标 tile 连续Cols ValidCol或Rows 1走TColExpandBinaryNormMode以一次向量指令配合 repeat 步长repeat stride完成多行广播否则走TColExpandBinaryCountMode逐行循环并配合SetVectorCount设置向量计数TColExpandBinOp.hpp#L76-L82。A5 平台寄存器张量 谓词掩码含 64 位模拟A5 后端实现在 include/pto/npu/a5/TColExpandSub.hpp其标量路径为template typename T struct ColExpandSubOp { PTO_INTERNAL static void ColExpandBinaryInstr( RegTensorT reg_dst, RegTensorT reg_src0, RegTensorT reg_src1, MaskReg preg) { vsub(reg_dst, reg_src0, reg_src1, preg, MODE_ZEROING); } ... };差异点在于 A5 使用RegTensor寄存器张量MaskReg谓词掩码MODE_ZEROING的编程模型并在 a5/TColExpandBinOp.hpp 中提供1D/2D与PostUpdate/NoPostUpdate多种实现变体VFImplKind可选默认走VFIMPL_DEFAULT连续 tile 使用 PostUpdate 版本以复用地址自增特性。64 位路径则由Int64ColExpandBinary承担a5/TColExpandBinOp.hpp#L155-L219每个 64 位元素拆为(low, high)两个 32 位寄存器通过Int64LoadBounded按列偏移读取Int64BinaryCalcRegsInt64Op::Sub, T完成低/高字对的减法运算最后用vintlv交织回写并配合pintlv_b32生成的低/高掩码分别存储。对于非 A5/A6 目标该函数仅有声明式 stub编译期不可用。使用示例C kernel 示例A5 平台测试用例以下示例取自 tests/npu/a5/src/st/testcase/tcolexpandsub/tcolexpandsub_kernel.cpp展示了完整的数据搬运—计算—存回流程#include pto/pto-inst.hpp using namespace pto; template typename T, uint32_t dstRow, uint32_t dstCol, uint32_t src1Row, uint32_t src1Col __global__ AICORE void runCOLEXPANDSUB(__gm__ T __out__* out, __gm__ T __in__* src0, __gm__ T __in__* src1) { using TileData TileTileType::Vec, T, src1Row, src1Col, BLayout::RowMajor, -1, -1; using DstTileData TileTileType::Vec, T, dstRow, dstCol, BLayout::RowMajor, -1, -1; DstTileData src0Tile(dstRow, dstCol); TileData src1Tile(src1Row, src1Col); DstTileData dstTile(dstRow, dstCol); TASSIGN(src0Tile, 0x0); TASSIGN(src1Tile, 0x10000); TASSIGN(dstTile, 0x20000); TLOAD(src0Tile, src0Global); TLOAD(src1Tile, src1Global); // 手动模式下的流水线同步等待 MTE2 搬运完成后向量单元再计算 set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); TCOLEXPANDSUB(dstTile, src0Tile, src1Tile); // 计算完成后等待再触发 MTE3 存回 set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); TSTORE(dstGlobal, dstTile); }该用例验证了src1形状为1 × Csrc1Row 1的典型用法并覆盖多种数据类型与形状组合例如tcolexpandsub_kernel.cpp#L74-L83float6×128、18×32halfaclFloat1610×256、12×64int32_t16×32int16_t16×6464 位类型int64_t、uint64_t16×32对应 A5 的 64 位模拟路径。CPU 参考实现CPU 侧的参考测试位于 tests/cpu/st/testcase/tcolexpandop/tcolexpandop_kernel.cpp其中LaunchTCOLEXPANDSUB以 lambda 方式调用TCOLEXPANDSUB(dst, src0, src1)src1声明为TileTileType::Vec, T, 1, iCol, BLayout::RowMajor, -1, -1即单行向量可用于在没有 NPU 的环境下做语义正确性对照。PTO 汇编完整形式%dst tcolexpandsub %src0, %src1 : !pto.tile..., !pto.tile... - !pto.tile... # AS Level 2 (DPS) pto.tcolexpandsub ins(%src0, %src1 : !pto.tile_buf..., !pto.tile_buf...) outs(%dst : !pto.tile_buf...)相关资源指令文档docs/isa/TCOLEXPANDSUB.md含中文版 TCOLEXPANDSUB_zh.md、示意图 docs/figures/isa/TCOLEXPANDSUB.svg同类单操作数广播指令docs/isa/TCOLEXPAND.mdC 接口声明include/pto/common/pto_instr.hpp公共头pto/pto-inst.hppA2/A3 后端include/pto/npu/a2a3/TColExpandSub.hpp 与 include/pto/npu/a2a3/TColExpandBinOp.hppA5 后端include/pto/npu/a5/TColExpandSub.hpp 与 include/pto/npu/a5/TColExpandBinOp.hpp测试用例A5 ST 用例 tests/npu/a5/src/st/testcase/tcolexpandsub/、CPU 参考用例 tests/cpu/st/testcase/tcolexpandop/。赞分享人工智能指令集算子库CANNAscend【免费下载链接】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点击查看免费下载相关推荐PTO-ISA TCOLEXPANDDIV 指令详解列广播除法Column-wise Broadcast Divide的语义、汇编与 C 编程指南PTO ISA TCOLEXPANDDIV 指令详解列广播除法Column wise Broadcast Divide的语义、汇编与 C 编程指南 T人工智能指令集算子库CANNAscendPTO-ISA TCOLEXPANDDIV 指令详解列广播除法Column-wise Broadcast Divide的语义、编程接口与多平台实现PTO ISA TCOLEXPANDDIV 指令详解列广播除法Column wise Broadcast Divide的语义、编程接口与多平台实现 导读人工智能指令集算子库CANNAscend基于 agno 的音频转写数据标注实战纯文本转写、说话人分离与时间戳分段基于 agno 的音频转写数据标注实战纯文本转写、说话人分离与时间戳分段 音频转写Speech to text是 agno 数据标注体系中音频模态最基础的人工智能指令集算子库CANNAscend上一篇【免费下载】 NVIDIA nvbandwidth 工具使用指南下一篇leak-check深度解析构建隐私保护的BFS图遍历与脱敏聚合系统创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考