PTO Tile Intrinsics 编程模型-平芜编程栈

PTO Tile Intrinsics 编程模型

【免费下载链接】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 Tile Lib 提供以Tile 粒度为核心的 C++ 内建接口（intrinsics），并可映射到 PTO ISA。该模型的设计目标是：

跨设备代际的可移植性：硬件细节可能变化（指令细节、存储布局、调度约束等），但编程模型保持稳定。
接近硬件的性能表达能力：Tile 与 GlobalTensor 足够底层，可表达高效的数据搬运与计算。
覆盖两类开发者：偏“编译器做重活”的高效编程方式，以及偏“显式控制放置与同步”的专家调优方式。

抽象执行模型（core/device/host）参见：docs/machine/abstract-machine_zh.md。

核心概念

Tile：固定容量的二维片上缓冲区（概念上类似 tile 寄存器 / SRAM 块），也是大多数 PTO 指令的主要计算单元。参见：docs/coding/Tile_zh.md。
GlobalTensor：全局内存（GM）的轻量级视图，带 5 维 shape/stride/layout 元数据；被TLOAD、TSTORE等内存类指令消费。参见：docs/coding/GlobalTensor_zh.md。
Scalar：用于参数化指令的立即数与枚举（舍入模式、比较模式、原子模式等）。参见：docs/coding/Scalar_zh.md。
Event：显式的依赖 token，用于在不引入全局屏障的情况下表达流水线类之间的顺序约束。参见：docs/coding/Event_zh.md。

两种开发风格

PTO-Auto

PTO-Auto 面向希望获得简单、可移植体验的开发者：

编译器/运行时选择内存放置与地址绑定策略。
编译器插入必需的同步。
编译器调度操作并在可能时做融合。

该模式适合作为正确性与可移植性的起点。

PTO-Manual

PTO-Manual 面向需要完全控制以进行性能调优的开发者：

开发者控制内存放置与绑定（例如通过TASSIGN）。
开发者显式表达顺序（events 和/或TSYNC）。
开发者控制操作调度与流水线结构。

该模式使关键内核能够进行专家级优化，同时仍复用同一套 Tile/GlobalTensor 抽象。

执行模型：SPMD 与 MPMD

PTO 支持SPMD与MPMD两种执行模型。

这两种模型描述的是工作如何映射到核心；它们与Auto vs Manual开发风格是正交的（可以写 SPMD-Auto、SPMD-Manual、MPMD-Auto 或 MPMD-Manual）。

SPMD（Single Program, Multiple Data）

在 SPMD 中，所有参与的核心运行同一入口函数，每个核心使用自身的运行时身份（例如block_idx）选择其数据区域。

当存在 sub-block 分解时，可以构造稳定的“虚拟 id”：

auto cid = get_block_idx(); auto vid = get_block_idx() * get_subblockdim() + get_subblockid();

SPMD 适合规则的张量 tiling（GEMM、按行 softmax、逐元素算子等）。

MPMD（Multiple Program, Multiple Data）

在 MPMD 中，不同核心（或核心组）可以在同一 tile 图中执行不同的 tile 程序。概念上由Device Machine 调度器决定某个核心运行哪段“程序”。

一种可移植写法是：调度器提供一个task id给内核入口函数，在内核中按 task 分派：

__global__ __aicore__ void KernelMPMD(__gm__ float* out, __gm__ const float* in, uint32_t task_id) { switch (task_id) { case 0: return ProducerStage(out, in); case 1: return ConsumerStage(out, in); default: return; } }

说明：