人工智能指令集算子库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点击查看免费下载TPUT_ASYNC_NOTIFY是 CANN pto-isa 通信指令集中远程写 信号通知二合一的关键原语它把本地 GM 的一段非空负载异步搬运到远端 GM随后更新远端一个 32 位信号一次调用同时完成数据投递与就绪通知。本文以 docs/isa/comm/TPUT_ASYNC_NOTIFY.md 为骨架结合仓库头文件实现与多平台测试用例完整讲解其数据流、三种 DMA 后端SDMA/URMA/RDMA的模板参数与会话构造、操作语义、约束条件、完成语义及并发规则并给出可直接落地的 Set / AtomicAdd / 收发双方完整代码示例。读完本文你将掌握如何在 A2/A3、A5 与 Ascend950 平台上正确使用该指令实现先到数据、后到通知的跨 NPU 同步通信。指令定位一条指令同时完成远程写与信号通知TPUT_ASYNC_NOTIFY的语义是一条带通知的异步远程写asynchronous remote write with notification。它做两件事且严格按顺序完成将一段非空负载从本地 GM 传输到远端 GM负载传输完成后再更新远端一个 32 位信号。其数据流可表示为srcGlobalData (local GM) → DMA engine → dstGlobalData (remote GM) → update dstSignalData (remote GM)负载的源srcGlobalData位于本地GM目的dstGlobalData与信号dstSignalData的地址都由调用方显式提供指向远端GM。返回的AsyncEvent同时覆盖负载传输完成和信号更新完成两个事件Wait成功后既可安全读取远端负载也可依赖远端信号已置位。如果只需要发信号、不需要搬运任何负载应当改用纯通知原语TNOTIFY见 docs/isa/comm/TNOTIFY.md。TNOTIFY只对远端int32_t信号执行Set或AtomicAdd不产生任何数据搬运。与同族的TPUT_ASYNC纯异步远程写见 docs/isa/comm/TPUT_ASYNC.md相比TPUT_ASYNC_NOTIFY在负载后追加了一步信号更新因此特别适合生产者写数据 → 消费者等信号的流水线场景例如跨 NPU 的直写式 AllGather 尾部同步、标志位驱动的生产者-消费者队列等。模板参数与 DMA 后端编译期选择引擎engine模板参数在编译期选择 DMA 后端默认值为DmaEngine::SDMA。仓库中DmaEngine枚举定义于 include/pto/comm/comm_types.hppenum class DmaEngine : uint8_t。三个后端的能力与平台限制如下表引擎平台限制支持的NotifyOpDmaEngine::SDMA默认A2/A3 与 A5Set、AtomicAddDmaEngine::URMA仅 Ascend950NPU_ARCH 3510要求 CANN Toolkit 9.1.0Set、AtomicAddDmaEngine::RDMA仅 Ascend950NPU_ARCH 3510当前支持 HNS1825 RoCE 网卡仅Set选择建议A2/A3 与 A5 上的默认跨核通信使用DmaEngine::SDMA无需额外硬件与版本前置条件peer参数被忽略。Ascend950 上的低延迟 RDMA 路径使用DmaEngine::URMA用户态 RDMA 内存访问要求 CANN Toolkit 9.1.0或DmaEngine::RDMA标准跨节点组网当前仅 HNS1825 RoCE。模板化的engine设计使代码对未来的新后端如 CCU保持前向兼容——这是 pto-isa 通信层一贯的编译期抽象策略。C 内置函数签名与参数详解TPUT_ASYNC_NOTIFY声明于 include/pto/comm/pto_comm_inst.hpp通过#if defined(PTO_NPU_ARCH_A2A3) || defined(PTO_NPU_ARCH_A5) || defined(__CPU_SIM)宏在 A2/A3、A5 与 CPU 仿真平台上可用template DmaEngine engine DmaEngine::SDMA, typename GlobalDstData, typename GlobalSrcData, typename GlobalSignalData, typename... WaitEvents PTO_INST AsyncEvent TPUT_ASYNC_NOTIFY(GlobalDstData dstGlobalData, GlobalSrcData srcGlobalData, GlobalSignalData dstSignalData, int32_t signalValue, NotifyOp notifyOp, const AsyncSession session, uint32_t peer, WaitEvents ... events);注意该签名比TPUT_ASYNC多了信号相关参数并且带显式peer参数。参数说明如下参数说明dstGlobalData负载目的全局张量位于远端GM。srcGlobalData负载源全局张量位于本地GM。dstSignalData位于远端 GM 的一个 32 位信号数据类型必须是int32_t。signalValueSet模式下赋给信号的值AtomicAdd模式下作为加法的增量。notifyOp信号更新操作NotifyOp::Set或NotifyOp::AtomicAdd。session为engine模板参数构建的AsyncSession用BuildAsyncSession构建一次之后复用于所有异步调用与事件等待。peerURMA/RDMA 的目标 rank用于选择通信队列与远端内存信息SDMA 不使用该参数。events零个或多个前置 PTO 流水线事件每个事件必须提供零参Wait()方法。events的处理逻辑可在 include/pto/comm/pto_comm_inst.hpp 中看到函数体首先调用WaitAllEvents(events...)等待所有前置事件再分发到架构相关的TPUT_ASYNC_NOTIFY_IMPLengine。有两点需要特别注意AsyncEvent不能作为events参数传入AsyncEvent提供的是带 session 的Wait(session)而非零参Wait()因此必须显式等待并携带其 sessionpeer的语义随引擎不同URMA/RDMA 用peer选择目标 rank 的队列与已注册内存此时dstGlobalData与dstSignalData都必须属于该 peerSDMA 的远端地址直接来自全局张量本身peer被忽略。信号类型comm::Signalcomm::Signal是一个int32_t信号的全局张量别名定义于comm命名空间using Signal GlobalTensorint32_t, Shape1, 1, 1, 1, 1, Stride1, 1, 1, 1, 1, Layout::ND;构造Signal只包装调用方提供的 GM 地址不分配也不初始化底层内存。调用方必须自行分配信号内存并按通信协议完成初始化。两种NotifyOp对signalValue的用法不同Set把signalValue赋值给信号直接覆盖AtomicAdd把signalValue作为有符号增量累加到信号上原子加。NotifyOp枚举同样定义于 include/pto/comm/comm_types.hppenum class NotifyOp : uint8_t。操作语义严格先负载、后信号一次调用按以下顺序执行等待所有events前置流水线事件把完整的负载从srcGlobalData传输到dstGlobalData负载传输完成后按notifyOp更新dstSignalData。两种通知操作的数学语义NotifyOp::Set$$ \mathrm{signal}^{\mathrm{remote}} \mathrm{signalValue} $$NotifyOp::AtomicAdd$$ \mathrm{signal}^{\mathrm{remote}} \mathrel{} \mathrm{signalValue} \quad (\text{atomic}) $$关于原子性与一致性的关键结论AtomicAdd对信号的更新是原子的。多个生产者可以并发对同一个信号做AtomicAdd最终增量等于各signalValue之和——这是实现多生产者完成计数器的惯用手段。原子性不适用于负载写入多个生产者并发搬运负载时目的地址区间不得重叠否则数据竞态由应用负责规避。当多个生产者并发对同一信号执行Set时指令不保证最终值在没有应用级同步的前提下不要在同一信号上混用Set、AtomicAdd与普通 store。DmaEngine::RDMA不支持NotifyOp::AtomicAdd。上述负载先于信号的顺序只约束单次调用内部不定义不同 session 或独立执行流之间的顺序。AsyncSession 构造三种后端的差异所有异步原语都要求先构建一个AsyncSession。构造接口BuildAsyncSession位于 include/pto/comm/async_common/async_event_impl.hpp按引擎提供不同重载。构造失败时返回false只有构建成功的 session 才能用于异步原语与事件等待。SDMA 构造默认后端template DmaEngine engine DmaEngine::SDMA, typename ScratchTile PTO_INTERNAL bool BuildAsyncSession( ScratchTile scratchTile, __gm__ uint8_t *workspace, AsyncSession session, uint32_t syncId 0, const sdma::SdmaBaseConfig baseConfig { sdma::kDefaultSdmaBlockBytes, 0, 1}, uint32_t channelGroupIdx sdma::kAutoChannelGroupIdx);参数默认值说明scratchTile—SDMA 控制元数据用的 UB 临时 Tile。workspace—由宿主机侧SdmaWorkspaceManager分配的 GM 指针。session—输出的AsyncSession。syncId0MTE3/MTE2 流水线同步事件 ID取值范围 0-7若内核在该 ID 上使用了其他流水线屏障需要覆盖。baseConfig{kDefaultSdmaBlockBytes, 0, 1}SDMA 块字节数、通信块偏移、队列数。channelGroupIdxkAutoChannelGroupIdxSDMA 通道组索引默认取get_block_idx()映射到当前 AI 核。SDMA 并发要点A2/A3并发 AIV 必须各自构建独立 session并使用不同的 Channel Group。默认的kAutoChannelGroupIdx从get_block_idx()选组显式传组的调用方也必须保证每个并发 AIV 独占一个组queue_numbaseConfig第三项可以大于 1但一次TPUT_ASYNC_NOTIFY的负载与信号只通过该组的 queue 0 提交以保证负载先于信号的顺序增大queue_num不会把该通知负载分条到多个队列。URMA 构造仅 NPU_ARCH 3510#ifdef PTO_URMA_SUPPORTED template DmaEngine engine PTO_INTERNAL bool BuildAsyncSession(__gm__ uint8_t *workspace, AsyncSession session); #endif参数说明workspace由宿主机侧UrmaWorkspaceManager分配的 GM 指针。session输出的AsyncSession。URMA不需要scratchTile轮询直接使用ld_dev/st_dev硬件内建函数要求 CANN Toolkit 9.1.0。URMA session不绑定目标 rank目标 rank 由内置函数调用时的peer参数选择——这也是它与TPUT_ASYNC文档中session 绑定destRankId差异之所在使用时务必注意。RDMA 构造仅 NPU_ARCH 3510#ifdef PTO_RDMA_SUPPORTED template DmaEngine engine, typename ScratchTile PTO_INTERNAL bool BuildAsyncSession(ScratchTile scratchTile, __gm__ uint8_t *workspace, uint32_t myPe, AsyncSession session, uint32_t syncId 0); #endif参数说明scratchTileRDMA 使用的 UB/Vec 临时 Tile至少 64 字节用于 WQE/CQE 控制数据。workspace宿主机侧 RDMA 初始化流程返回的 GM 指针。myPe本地 rank ID用于选择已注册的本地内存区域。session输出的AsyncSession。syncIdMTE/标量同步事件 ID范围 0-7。RDMA session 同样不绑定目标 peer目标 rank 由内置函数的peer参数选择comm::AsyncSession session; if (comm::BuildAsyncSessioncomm::DmaEngine::RDMA( scratchTile, rdmaWorkspace, myPe, session)) { auto event comm::TPUT_ASYNC_NOTIFYcomm::DmaEngine::RDMA( dstGlobalData, srcGlobalData, remoteSignal, 1, comm::NotifyOp::Set, session, peer); (void)event.Wait(session); }scratchTile 的角色与尺寸要求scratchTile是 SDMA/RDMA session 使用的临时 UB 工作区不包含用户负载。它被转换为TmpBuffer用于读写 SDMA 控制字flag、sq_tail、channel_info、轮询事件完成标志、提交队列尾指针等控制与同步元数据。负载数据在 GM 缓冲区之间直接搬运不经由scratchTile。约束必须是pto::Tile类型且位于 UB/Vec 内存ScratchTile::Loc TileType::Vec必须保持有效直到关联事件全部完成SDMA 至少需要 8 个可用字节推荐类型为TileTileType::Vec, uint8_t, 1, comm::sdma::UB_ALIGN_SIZE256 字节RDMA 至少需要 64 个可用字节。约束条件编译期断言与运行时校验除文档明示的约束外仓库实现 include/pto/comm/async_common/TPutAsyncCommonDetail.hpp 用static_assert与PTO_ASSERT在编译期/运行期做了同样严格的校验二者一一对应类型与布局GlobalSrcData::RawDType必须等于GlobalDstData::RawDTypestatic_assertTPUT_ASYNC: src/dst element type mismatchGlobalSrcData::layout必须等于GlobalDstData::layoutTPUT_ASYNC: src/dst layout mismatch。扁平连续 1D源与目的的负载张量必须是扁平、连续的逻辑 1D 张量。实现中的TPutAsyncIsFlatContiguous1D检查打包布局pitch41且各维 stride 满足pitch_{n} dim_{n1} * pitch_{n1}且只有一行dim0..dim3 1。不满足 1D 连续要求时当前实现会返回无效的异步事件handle 0。容量目的元素容量必须不小于源元素个数TPutAsyncValidatePayload中PTO_ASSERT(dstElems srcElems, ...)实际传输字节数为srcElems * sizeof(RawDType)。负载非空负载大小必须大于零纯信号操作请用TNOTIFY。信号dstSignalData必须恰好包含远端 GM 中的一个int32_tstatic_assertsignal type must be int32_t地址必须非空signal pointer must not be null且 4 字节对齐signal address must be 4-byte aligned即reinterpret_castuint64_t(data) (alignof(int32_t)-1) 0notifyOp必须是Set或AtomicAddnotifyOp must be Set or AtomicAdd。信号内存由调用方负责分配与初始化。地址区间负载目的区间不得与dstSignalData重叠。URMA/RDMA 的 peer 与注册内存负载目的与信号必须属于同一目标 peer本地负载、远端负载、远端信号三段区间都必须落在宿主机初始化时注册的内存区域内。URMA 负载大小单次调用无大小上限——超过 256 MB单个 WQE 的传输上限的负载会被实现自动拆分为多个 WQE对调用方透明且信号保证在所有负载分片完成后才更新。RDMA 单次负载不得超过0x7fffffff字节。RDMA 仅支持Set不要对 RDMA 使用NotifyOp::AtomicAdd。workspace 来源SDMA workspace 必须由宿主机侧SdmaWorkspaceManager初始化URMA workspace 必须由宿主机侧UrmaWorkspaceManager初始化。传给UrmaWorkspaceManager::Init()的对称数据缓冲区必须是 HCCL 可注册的设备内存其分配方式需满足当前 CANN/HCCL 运行时的要求可参考 include/pto/comm/async/urma/urma_workspace_manager.hpp。生命周期session、workspace、scratchTile 必须保持存活直到所有关联事件完成。完成语义Wait / Test 与 quiet 语义event.Wait(session)阻塞直到事件完成event.Test(session)非阻塞地测试是否完成两者都必须使用发起操作时的同一个 session成功完成覆盖完整负载传输 随后信号更新。与TPUT_ASYNC一致的 quiet 语义同一 session、同一后端队列中的异步操作按提交顺序完成。连续提交多个操作后只需等待最后一个返回的AsyncEvent即可覆盖前面所有未决操作。URMA/RDMA 面向不同 peer的队列相互独立完成因此需要对每个 peer 分别等待最后一个事件Wait/Test无需再次传入peer。接收方视角的关键保证当接收方观察到信号被更新时对应负载一定已经传输到远端 GM。但指令不保证已有的接收侧负载缓存项被刷新读取负载前调用方必须按目标平台与运行时的内存一致性规则保证可见性示例中接收方在TWAIT之后、读负载之前需要做相应的同步处理。并发与 session 所有权规则A2/A3 SDMA使用独立 session 与独立 Channel Group 的 AIV 可以并发访问同一 rank 或不同 rank。URMA不同 AIV 可并发访问不同 peer对同一 peer的访问必须串行化。并发的负载区间不得重叠共享信号用AtomicAddSet应使用各自独立的信号。禁止多个执行流并发使用同一个 session。A2/A3 SDMA 专属每个并发 AIV 必须使用独立 session 和独立 Channel Group多个 AIV 不得并发向同一组提交。当queue_num N时最多只有kSdmaMaxChannelGroups / N个组可用。并发Set生产者应使用独立的远端信号多 AIV 共享一个信号做完成计数器时应使用AtomicAdd两种模式下负载目的区间都不得重叠。URMA/RDMA即使调用方基于同一 workspace 构建了不同 session对同一 peer/QP 的提交也必须串行化不同 peer 使用独立队列。重建 session 或复用其后端队列之前必须先完成此前所有事件。使用 URMA 时释放通信资源前必须完成每个活跃 peer/QP 的最后一个事件并同步所有使用CommContext的宿主流若通过existingComm提供了外部 HCCL 通信器还需先停止其其他用户并销毁通信器以释放其 channel/MR然后再调用DestroyComm、Reset、用BuildComm重建或销毁CommContext。完整示例六种典型用法以下示例假定宿主通信运行时已翻译好远端地址并为所选引擎初始化好 workspace。1. SDMA Set完整内核#include pto/comm/pto_comm_inst.hpp #include pto/common/pto_tile.hpp using namespace pto; template typename T __global__ AICORE void PutAndNotifySdma(__gm__ T *remoteDst, __gm__ T *localSrc, __gm__ int32_t *remoteSignalPtr, __gm__ uint8_t *sdmaWorkspace, uint32_t peer) { using ShapeDyn ShapeDYNAMIC, DYNAMIC, DYNAMIC, DYNAMIC, DYNAMIC; using StrideDyn StrideDYNAMIC, DYNAMIC, DYNAMIC, DYNAMIC, DYNAMIC; using GT GlobalTensorT, ShapeDyn, StrideDyn, Layout::ND; using ScratchTile TileTileType::Vec, uint8_t, 1, comm::sdma::UB_ALIGN_SIZE; ShapeDyn shape(1, 1, 1, 1, 1024); StrideDyn stride(1024, 1024, 1024, 1024, 1); GT dstGlobalData(remoteDst, shape, stride); GT srcGlobalData(localSrc, shape, stride); comm::Signal remoteSignal(remoteSignalPtr); ScratchTile scratchTile; TASSIGN(scratchTile, 0x0); comm::AsyncSession session; if (!comm::BuildAsyncSessioncomm::DmaEngine::SDMA( scratchTile, sdmaWorkspace, session)) { return; } auto event comm::TPUT_ASYNC_NOTIFYcomm::DmaEngine::SDMA( dstGlobalData, srcGlobalData, remoteSignal, 1, comm::NotifyOp::Set, session, peer); (void)event.Wait(session); }要点负载形状(1,1,1,1,1024)满足扁平连续 1D 要求scratchTile在栈上分配并TASSIGN(scratchTile, 0x0)清零后传给BuildAsyncSessionSDMA 路径的peer实际被忽略但签名要求提供。2. SDMA AtomicAdd与上一例使用相同的 session 与张量构造仅将notifyOp改为AtomicAddcomm::AsyncSession session; if (comm::BuildAsyncSessioncomm::DmaEngine::SDMA( scratchTile, sdmaWorkspace, session)) { auto event comm::TPUT_ASYNC_NOTIFYcomm::DmaEngine::SDMA( dstGlobalData, srcGlobalData, remoteSignal, 1, comm::NotifyOp::AtomicAdd, session, peer); (void)event.Wait(session); }3. URMA SetURMA session 不绑定目标 rank目标由peer选择comm::AsyncSession session; if (comm::BuildAsyncSessioncomm::DmaEngine::URMA( urmaWorkspace, session)) { auto event comm::TPUT_ASYNC_NOTIFYcomm::DmaEngine::URMA( dstGlobalData, srcGlobalData, remoteSignal, 1, comm::NotifyOp::Set, session, peer); (void)event.Wait(session); }4. URMA AtomicAddcomm::AsyncSession session; if (comm::BuildAsyncSessioncomm::DmaEngine::URMA( urmaWorkspace, session)) { auto event comm::TPUT_ASYNC_NOTIFYcomm::DmaEngine::URMA( dstGlobalData, srcGlobalData, remoteSignal, 1, comm::NotifyOp::AtomicAdd, session, peer); (void)event.Wait(session); }5. RDMA SetRDMA 仅支持Set用peer选择目标 rankcomm::AsyncSession session; if (comm::BuildAsyncSessioncomm::DmaEngine::RDMA( scratchTile, rdmaWorkspace, myPe, session)) { auto event comm::TPUT_ASYNC_NOTIFYcomm::DmaEngine::RDMA( dstGlobalData, srcGlobalData, remoteSignal, 1, comm::NotifyOp::Set, session, peer); (void)event.Wait(session); }6. 接收方消费者comm::Signal ready(localSignalPtr); comm::TWAIT(ready, 1, comm::WaitCmp::EQ); // 读取负载前按目标平台与运行时规则确保负载可见性。接收方用TWAIT等待信号变为 1EQ比较此后按平台内存一致性规则同步后即可安全读取负载。完整的TWAIT语义可参考 docs/isa/comm/TWAIT.md纯信号发送侧的原语参考 docs/isa/comm/TNOTIFY.md。测试佐证跨平台用例覆盖仓库为TPUT_ASYNC_NOTIFY提供了 CPU 仿真、A2/A3、A5 三套用例验证了Set/AtomicAdd的负载搬运与信号更新行为tests/cpu/st/testcase/tput_async_notify/main.cpp宿主侧用例把信号初始化为INIT_SIGNAL 5调用内核后断言信号值——AtomicAdd模式下期望5 7 12Set模式下期望7同时把负载结果写回output.bin与 golden 比对直接验证了本文第 4 节的数学语义tests/cpu/st/testcase/tput_async_notify/tput_async_notify_kernel.cppCPU 仿真内核实现tests/npu/a2a3/comm/st/testcase/tput_async_notify/tput_async_notify_kernel.cppA2/A3 平台 SDMA 路径用例tests/npu/a5/comm/st/testcase/tput_async_notify/tput_async_notify_kernel.cppA5 平台用例tests/npu/a5/comm/st/testcase/tput_async_notify_urma/tput_async_notify_urma_kernel.cpp 与tests/npu/a5/comm/st/testcase/tput_async_notify_rdma/Ascend950 的 URMA/RDMA 路径用例及配套 README。从源码实现看不同架构的TPUT_ASYNC_NOTIFY_IMPL策略也不同include/pto/comm/pto_comm_inst.hpp 中的注释明确指出A2/A3 的 SDMA 把负载与信号提交到同一条 SQA5 的 SDMA 命名路径采用同步 MTE 后接标量SET/AtomicAdd返回handle 0的已完成事件A5 的 URMA/RDMA 用peer选择目标队列与注册内存CPU stub 则把负载委托给TPUT_ASYNC、信号更新委托给TNOTIFY。这套分层实现保证了上层调用语义的一致性与平台可移植性。小结TPUT_ASYNC_NOTIFY在 pto-isa 通信体系中提供了一条负载 信号原子序的异步原语其核心使用要点可概括为按平台选引擎A2/A3、A5 用 SDMAAscend950 用 URMA 或 RDMA、按引擎构建 session注意 URMA/RDMA 的peer动态选目标、SDMA 的 Channel Group 隔离、保证负载扁平连续 1D 且目的容量充足、信号为远端 4 字节对齐的单个int32_t、共享信号计数用AtomicAdd、并发写负载区间不重叠、最后用同一 session 的Wait收尾。遵循这些规则即可在多核/多机场景下安全地实现先数据后通知的跨 NPU 同步通信。赞分享人工智能指令集算子库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 完全指南点对点传输、信号同步与集合通信编程实战CANN pto-isaPTO 通信 ISA 完全指南点对点传输、信号同步与集合通信编程实战CANN pto isa 本文是 PTOParallel Tile Operatio人工智能指令集算子库CANNAscendPTO TTEST 指令详解基于轮询的非阻塞信号同步检测CANN pto-isaPTO TTEST 指令详解基于轮询的非阻塞信号同步检测CANN pto isa 本篇技术指南聚焦 CANN pto isa 仓库中 PTO 通信指令集的人工智能指令集算子库CANNAscendPTO-ISA 异步远程写原语 TPUT_ASYNC 详解SDMA/URMA/RDMA 三引擎会话构建与 Quiet 完成语义PTO ISA 异步远程写原语 TPUT_ASYNC 详解SDMA/URMA/RDMA 三引擎会话构建与 Quiet 完成语义 TPUT_ASYNC 是 CA人工智能指令集算子库CANNAscend上一篇用一句中文驱动浏览器干活Midscene Chrome扩展十分钟上手下一篇突破Documenso免费版限制从技术原理到合规解决方案创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考