CANN pto-isa：90+ Tile 级虚拟指令速查手册

个人主页：ujainu

前言

pto-isa 是昇腾CANN生态里的 PTO 虚拟指令集架构，它定义了 90+ 条标准 Tile 级操作，让算子开发者用一套虚拟指令就能在昇腾NPU上跨平台开发。打个比方：如果算子开发是做菜，那 pto-isa 就是一本标准化菜谱——不管你厨房里是煤气灶还是电磁炉，菜谱写的是"中火翻炒30秒"，而不是"把旋钮转到第3档"。虚拟指令就是这种跟具体硬件解耦的"中火翻炒"。

仓库定位

为什么需要一套虚拟指令？

pto-isa 仓库的核心使命很明确：把算子逻辑从硬件细节里拔出来。这背后有一个现实的痛点——昇腾芯片经历了多代迭代（从310到910系列），每一代的计算单元数量、片上内存大小、DMA通道数都不一样。如果算子开发者直接对着某一代芯片写代码，等下一代芯片出来，就得逐行适配，工程量跟重写差不多。

更麻烦的是，上层框架（PyTorch、MindSpore）的算子成百上千，每一个都去适配底层硬件细节，根本不现实。所以昇腾CANN团队的设计思路是：在中间插一层"虚拟指令"，让算子开发者只关心"对一块数据做什么操作"，至于这块数据怎么切分、怎么搬运、怎么映射到物理核上——交给编译器去操心。

# 传统方式：直接操作硬件细节 # 换芯片 = 全部重写 Reg_Write(0x3F, data_ptr); // 写寄存器 DMACopy(src, dst, 256); // 硬件DMA通道，通道数因芯片而异 CU_Execute(pipeline_0); // 指定计算单元，新一代可能没有这个pipeline # pto-isa 方式：Tile 级虚拟指令 # 换芯片？指令不变，编译器适配 TILE_COPY(src_tile, dst_tile) // 搬一块数据 TILE_MAC(a_tile, b_tile, c_tile) // 乘加一块数据

这就像你跟快递员说"把这箱书送到3楼"，而不是说"左转走20米、右转上楼梯、第15个台阶左拐……"。哪天小区改了门禁系统、换了楼梯口位置，你那句话依然管用，变的是快递员自己的路线规划。

仓库里到底有什么？

pto-isa 仓库的核心产出是一套虚拟指令定义文件，包含每条指令的操作码、操作数格式、语义说明。开发者不直接"运行"这些指令，而是把它们当作编译器的中间表示（IR）来使用。

# 克隆仓库看看结构gitclone https://atomgit.com/cann/pto-isa.gitcdpto-isa&&find.-typef|head-20

核心能力

90+ 条标准 Tile 级操作

pto-isa 定义了覆盖计算、搬运、规约、控制等场景的 90+ 条虚拟指令。按功能大致分几类：

数据搬运类——把数据搬到该去的地方

TILE_COPY — Tile 间数据拷贝，就像把一箱货物从A仓库搬到B仓库 TILE_GATHER — 按索引收集数据，像快递员按单取件 TILE_SCATTER — 按索引散布数据，像快递员按地址派件 TILE_BROADCAST — 把一个值复制到整块Tile，像大喇叭广播通知

实际写算子时，搬运往往是第一步。比如你想让两个矩阵相乘，得先把数据从全局内存搬到AI Core的本地缓冲区：

# Python 伪代码：典型的数据搬运流程defmatmul_tile(gmem_A,gmem_B,local_buf):# 从全局内存搬运到本地tile_a=TILE_LOAD(gmem_A,offset=0,shape=[16,16])tile_b=TILE_LOAD(gmem_B,offset=0,shape=[16,16])# 计算在本地完成tile_c=TILE_MATMUL(tile_a,tile_b)# 结果写回全局内存TILE_STORE(tile_c,gmem_C,offset=0)

算术计算类——真正的"干活"指令

TILE_ADD — 逐元素加法，A+B=C TILE_MUL — 逐元素乘法，A*B=C TILE_MAC — 乘累加，C += A*B（矩阵运算的核心原语） TILE_MATMUL — Tile 级矩阵乘，这是大模型算子的心脏 TILE_TRANSPOSE — 矩阵转置，行变列、列变行

其中 TILE_MAC（乘累加）是最常用的计算指令。大语言模型的推理和训练，底层几乎全是矩阵乘，而矩阵乘拆开来看就是无数次"取两个数相乘、再加到累加器上"。

// 两个 Tile 相加——最朴素的用法 TILE_ADD(tile_a, tile_b, tile_out);

规约类——把一堆数浓缩成一个

TILE_REDUCE_SUM — 求和规约，把整块Tile里的数全加起来 TILE_REDUCE_MAX — 求最大值规约 TILE_REDUCE_MEAN — 求平均值规约

规约在深度学习里到处都是。比如 Softmax 函数需要先求指数和（一个 SUM 规约），再拿每个元素除以这个和；Layer Norm 需要算均值和方差——都离不开规约。

// Softmax 的规约步骤（简化版） TILE_EXP(input_tile, exp_tile); // e^x TILE_REDUCE_SUM(exp_tile, sum_scalar); // 求和，得到一个标量 TILE_BROADCAST(sum_scalar, sum_tile); // 广播回 Tile 形状 TILE_DIV(exp_tile, sum_tile, prob_tile); // 每个元素除以总和

元素变换类——逐个元素做数学变换

TILE_EXP — 指数运算 e^x TILE_LOG — 对数运算 ln(x) TILE_RECIP — 取倒数 1/x TILE_SQRT — 开平方 TILE_RELU — ReLU 激活，负数变零、正数不变 TILE_GELU — GELU 激活，大模型里常用的激活函数

位操作与类型转换类

TILE_CAST — 数据类型转换，比如 FP16→FP32 做累加防溢出 TILE_AND / TILE_OR / TILE_XOR — 位运算，某些量化算子会用到 TILE_SHIFT — 移位操作，也是量化场景的常客

类型转换在高精度累加场景特别常见。FP16 的表示范围小，做长序列的乘累加容易溢出，所以常见做法是：平时用 FP16 搬运和存储，累加时临时转成 FP32，算完再转回来。

// FP16 计算，FP32 累加的经典模式 TILE_CAST(tile_a_fp16, tile_a_fp32, FP16_TO_FP32); TILE_CAST(tile_b_fp16, tile_b_fp32, FP16_TO_FP32); TILE_MAC(tile_a_fp32, tile_b_fp32, acc_fp32); // FP32 精度累加 TILE_CAST(acc_fp32, tile_out_fp16, FP32_TO_FP16); // 转回 FP16 输出

递进示例：从最简单到实用

1️⃣ 最简单：两个 Tile 相加

// 就像两叠纸逐页对齐相加 TILE_ADD(tile_a, tile_b, tile_out);

2️⃣ 基础：矩阵乘 + 搬运

// 从全局内存搬到本地，做完矩阵乘再搬回去 // 就像：取食材 → 下锅炒 → 装盘上菜 TILE_COPY(gmem_tile_a, local_tile_a); TILE_COPY(gmem_tile_b, local_tile_b); TILE_MATMUL(local_tile_a, local_tile_b, local_tile_c); TILE_COPY(local_tile_c, gmem_tile_c);

3️⃣ 实用：FlashAttention 核心片段

// FlashAttention 的关键：分块算注意力，算完一块扔一块 // 就像流水线上的工人，处理完一个零件立刻传给下一位 for (int i = 0; i < num_tiles; i++) { TILE_COPY(kv_tile[i], local_kv); // 取一块KV TILE_MATMUL(q_tile, local_kv, score); // Q × K^T TILE_SOFTMAX(score, prob); // Softmax归一化 TILE_MAC(prob, local_kv, out_tile); // × V 累加 }

4️⃣ 进阶：Layer Norm 的完整实现

// Layer Norm：先算均值和方差，再归一化 // 就像给全班考试成绩做"标准化"——让平均分变0、标准差变1 TILE_REDUCE_MEAN(input_tile, mean_tile); TILE_SUB(input_tile, mean_tile, diff_tile); TILE_MUL(diff_tile, diff_tile, sq_tile); TILE_REDUCE_MEAN(sq_tile, var_tile); TILE_ADD(var_tile, eps_tile, safe_var); // 加一个小常数防除零 TILE_SQRT(safe_var, std_tile); TILE_DIV(diff_tile, std_tile, norm_tile); // 归一化 TILE_MUL(norm_tile, gamma_tile, out_tile); // 乘以可学习缩放参数

在架构中的位置

pto-isa 位于 CANN 五层架构的第3层编译层，跟 Graph Compiler 同层。它不是跑在硬件上的指令，而是编译器的输入——编译器把虚拟指令翻译成昇腾达芬奇架构真正能执行的机器指令。

第1层 AscendCL ← 应用层接口 第2层 AOL/AOE ← 算子库 + 调优 第3层 Graph Compiler / pto-isa ← 编译层：虚拟指令在这里变成真指令 第4层 Runtime ← 执行层 第5层 驱动 ← 基础层

把这个层级关系翻译成大白话：第1-2层是用户直接接触的"前台"，你调一个 PyTorch 的torch.matmul，经过 AscendCL 送到第2层；第3层是"后厨"，pto-isa 定义的操作到了这里才被翻译成达芬奇核能懂的二进制指令；第4-5层是"配送"，Runtime 负责把编译好的指令调度到具体的 AI Core 上执行，驱动负责跟硬件打交道。

# 一条虚拟指令的"旅行"过程 用户代码: tile_c = TILE_MATMUL(tile_a, tile_b) ↓ 编译器前端: 语法检查、类型推断、Tile形状推导 ↓ 编译器后端: 指令选择 → 将 TILE_MATMUL 拆成 Cube 指令 ↓ 寄存器分配: 分配物理寄存器、DMA通道 ↓ 二进制生成: 达芬奇架构机器码 ↓ 运行时: 调度到 AI Core 执行

与其他仓库的关系

pto-isa 不是孤立存在的，它在 CANN 工具链里扮演"中间语言"的角色：

pypto——编程框架，pto-isa 的"语法糖"

pypto 提供了 Tile 编程模型和多层次 IR 系统，pto-isa 的虚拟指令就是 pypto 编程模型的底层语义。打个比方：pypto 像是 Python 语言本身，pto-isa 像是 Python 字节码——你用 pypto 写高层代码，编译器帮你"翻译"成 pto-isa 虚拟指令。pypto 的优势是提供循环、分支、张量抽象等高级结构，让开发者不用手拼一条条虚拟指令。

# pypto 编程风格（伪代码）——更像在写正常程序@tile_kerneldeffused_add_relu(a:Tile[FP16],b:Tile[FP16]):c=a+b# pypto 编译成 TILE_ADDreturnrelu(c)# pypto 编译成 TILE_RELU + TILE_MUL

asc-devkit——另一条算子开发路径

Ascend C 是昇腾官方的算子开发高级语言，语法类似 C++，封装了数据管理、流水线并行等能力。pypto/pto-isa 提供了另一条 Tile 级编程路径。两条路径各有适用场景：Ascend C 适合想快速开发单个算子的工程师，pto-isa 更适合需要精细控制 Tile 级操作、或者做编译器基础设施开发的场景。

// Ascend C 风格对比（伪代码）——更"高级"，不用自己管Tile__global__ __aicore__voidadd_kernel(GM_ADDR a,GM_ADDR b,GM_ADDR c){LocalTensor<half>x,y,z;DataCopy(x,g_a);DataCopy(y,g_b);Add(z,x,y);DataCopy(g_c,z);}

graph-autofusion——算子融合的 codegen 引擎

算子自动融合框架的 SuperKernel codegen JIT 会将多个算子融合后生成 pto-isa 虚拟指令。比如模型里有连续的"加法→ReLU→归一化"三个算子，graph-autofusion 会把它们合并成一个 SuperKernel，中间数据不用写回内存，直接用 pto-isa 的虚拟指令串联起来。这种融合对性能的影响非常大——减少了片外内存访问次数，省下来的时间远比多算几条指令花的时间多。

# autofusion 配置片段（示意）fusion_patterns:-name:add_relu_normsequence:[Add,ReLU,LayerNorm]backend:pto-isa# 融合后的 SuperKernel 输出为 pto-isa 指令tile_policy:same_shape# 要求相邻算子的 Tile 形状一致才能融合

ascend-boost-comm——算子公共平台中间件

ascend-boost-comm 南向对接算子库、北向支撑加速库，pto-isa 定义的 Tile 操作是这条链路里的基础语义单元。简单说，boost-comm 是个"翻译官"，把上层框架的需求翻译成底层算子能理解的调用。在这个翻译过程中，Tile 级操作的语义定义就是 pto-isa 提供的——它是两边都能听懂的"普通话"。

pypto（编程范式） ↓ 编译到 pto-isa（虚拟指令） ↓ 翻译为 Graph Compiler → 达芬奇机器指令 graph-autofusion（融合算子 JIT） ↓ 生成 pto-isa（虚拟指令） ↓ 同样走 Graph Compiler 下沉

实战补充：Profiling Tile 操作

开发完算子，怎么知道 Tile 指令的执行效率？昇腾CANN 提供了 Profiling 工具，可以查看每条虚拟指令的耗时：

# 用 msprof 采集 Tile 操作的性能数据（伪代码）importtorch_npu.profilerasprofilerwithprofiler.profile(activities=[profiler.ProfilerActivity.CPU,profiler.ProfilerActivity.NPU],record_shapes=True,with_stack=True)asp:model(input_tensor)# 跑你的模型# 查看 Tile 级指令的耗时分布p.table(sort_by="self_npu_time_total",row_limit=20)

你也可以在代码里手动插入标记，精确测量某一段虚拟指令序列花了多少时间：

// 插桩计时（伪代码） TILE_PROFILE_BEGIN("matmul_block"); TILE_COPY(gmem_a, local_a); TILE_COPY(gmem_b, local_b); TILE_MATMUL(local_a, local_b, local_c); TILE_PROFILE_END("matmul_block"); // 运行后会在 Profiling 报告里看到 "matmul_block" 的详细耗时

记住这些就够了

pto-isa 是 PTO 虚拟指令集架构，定义了 90+ 条 Tile 级操作，让算子逻辑跟硬件解耦。虚拟指令写一次，编译器适配多代芯片。它跟 pypto 是一对——pypto 是编程框架，pto-isa 是指令定义；graph-autofusion 的 SuperKernel 融合也依赖它做 codegen。想深入了解？直接去仓库翻指令定义：

https://atomgit.com/cann/pto-isa