企业官网建设流程全解析

• 人工智能
• 指令集
• 算子库
• CANN
• Ascend

导读

本文全面解析 CANN pto-isa 提供的Auto Mode（自动模式）编译能力：编译器自动为Tile分配片上缓冲区地址、自动在昇腾硬件不同 pipeline 之间插入同步指令，让 kernel 开发者免于手写TASSIGN与Event同步，同时保持与专家手调 manual 模式相当的性能。读完本文，你将掌握 auto mode 的抽象边界（TF 层以上）、三大核心特性（liveness 分析、自动同步、Tile 内存分配）、完整的 Ascend CANN 编译命令，以及 kernel 开发者在 auto mode 下必须遵守的控制流与内存规则。

一、什么是 PTO Auto Mode

PTO（Parallel Tile Operation）是昇腾 CANN 定义的一种面向 tile 级运算的虚拟指令集架构，而Auto Mode 是 PTO 的一种编译模式：它把"为 Tile 分配内存"和"跨 pipeline 插入同步指令"这两件事全部交给编译器完成。

与 manual 模式相比，auto mode 下编程方式几乎不变，但有以下两点关键差异：

• 不需要调用TASSIGN来为 Tile 手动绑定硬件缓冲区地址；
• 不需要显式的Event同步（在 auto mode 下，这些指令实际是 no-op，写了也不起作用）。

这一表述在 auto 模式总览 中亦有明确说明，并得到仓库演示工程的印证：在 demos/auto_mode/baseline/add 的 README 中明确指出，与 manual 模式不同，kernel 中无需手动调用TASSIGN和同步指令，编译器会代为处理。

PTO AUTO 模式的定位是提供两大收益（见 auto 模式 README）：

1. 简化开发：在开发高效 PTO 代码的同时，仍为 kernel 开发者保留实现优化所必需的机制；
2. 跨架构兼容：同一份 PTO 源码可针对不同代际的昇腾架构编译，无需修改源码即可保持性能——尤其不需要关心不同架构之间 Cube 与 Vector 计算协同方式的差异。

二、为什么使用 Auto Mode

auto mode 的目标非常明确：为使用者提供更高层的接口抽象以提升开发效率，同时确保与专家手调代码相比仍具备有竞争力的性能。围绕这一目标，它提供了两大抽象能力：

• 自动同步指令插入：在昇腾硬件不同 pipeline（如PIPE_MTE2、PIPE_V、PIPE_MTE3）之间自动插入同步指令；
• Tile 内存分配与管理：对Tile抽象对象自动完成片上 buffer 的分配。

这两大抽象正是本文后续第三、四节要展开的核心。

三、抽象层级：编译器工作的边界在 TF 层以上

理解 auto mode 之前，必须先理解 PTO 指令实现的分层抽象。按照 auto 模式总览 的说明，每一个 PTO 指令的实现通常包含以下层级（从高到低）：

PTO 编译器工作在 "Tile" 这一抽象层级上。这意味着：上文所述的所有 auto 特性只在 TF 层以上生效，因为 TF 层接口是 Tile 抽象的最后一层；一旦进入 tile function 内部，就不再存在 tile 级抽象，只剩裸指针和裸 CCE intrinsics——那是 CCE 编译器的领域。因此，tile function 对 PTO 编译器而言是一个完全的黑盒子，PTO 编译器的功能不会进入 tile function 内部运作。

这个边界在库开发者规则中有更具体的体现：库开发者规则文档 第 3 条指出，auto-sync 不会穿透 tile function，整个 auto 模式编译器都工作在 tile function 层级，tile function 内部对编译器完全不可见，因此 tile function 内部仍需要库开发者手动加入同步。

源码层面，这一"向量化 Tile"的设计可以从 include/pto/common/memory.hpp 得到佐证：MemoryQualifier<TileType::Vec, DType>在__PTO_AUTO__宏下定义为__ubuf__ DType（向量类型），而在 manual 模式下定义为__ubuf__ DType*（指针类型）；TileType::Mat、TileType::Left、TileType::Right、TileType::Acc等也有同样的区分。这正是 库开发者规则文档 第 1 条所说 ".data()在 auto mode 下返回向量类型而非指针类型" 的根源。

四、Auto Mode 三大核心特性

4.1 Tile 自动 liveness 分析（核心基础）

在 auto mode 下，PTO 编译会跟踪每一个 Tile 及其live-range（活跃区间）。这是 auto mode 的核心分析组件，为后面两项特性（自动同步、Tile 内存分配）提供基础支撑：只有知道每个 Tile 何时被写、何时最后一次被读，编译器才能决定缓冲区的复用时机和同步的插入位置。

这一点在 库开发者规则文档 第 4 条中也有印证：auto mode 下整个内存分配完全基于每个 Tile 的 liveness 分析，不依赖其他任何上下文——这正是TPUSH、TPOP当前无法在 auto mode 下工作的原因。

4.2 自动同步（Automatic Synchronization）

在 manual 模式下，程序员必须时刻关注硬件异步执行的特点，借助 PTO 的 事件模型（Event Model） 在代码的精确位置插入同步，以同时保证功能正确与高性能。这既繁琐又极易出错。

auto mode 编译则替程序员省去了这一负担：编译器会在底层自动确定需要插入同步的位置，保证功能正确的同时维持有竞争力的性能。

为了直观感受两者的差距，我们看 auto 模式示例文档 中的TADD（向量加法）对比。

Auto mode 版本：

#include <pto/pto-inst.hpp> #include <pto/common/constants.hpp> using namespace pto; AICORE void runTAdd(__gm__ float __out__ *out, __gm__ float __in__ *src0, __gm__ float __in__ *src1) { using DynShapeDim5 = Shape<1, 1, 1, 64, 64>; using DynStridDim5 = Stride<1, 1, 1, 64, 1>; using GlobalData = GlobalTensor<float, DynShapeDim5, DynStridDim5>; using TileData = Tile<TileType::Vec, float, 64, 64, BLayout::RowMajor, 64, 64>; TileData src0Tile(64, 64); TileData src1Tile(64, 64); TileData dstTile(64, 64); GlobalData src0Global(src0); GlobalData src1Global(src1); GlobalData dstGlobal(out); TLOAD(src0Tile, src0Global); TLOAD(src1Tile, src1Global); TADD(dstTile, src0Tile, src1Tile); TSTORE(dstGlobal, dstTile); }

Manual mode 版本：

#include <pto/pto-inst.hpp> #include <pto/common/constants.hpp> using namespace pto; AICORE void runTAdd(__gm__ float __out__ *out, __gm__ float __in__ *src0, __gm__ float __in__ *src1) { using DynShapeDim5 = Shape<1, 1, 1, 64, 64>; using DynStridDim5 = Stride<1, 1, 1, 64, 1>; using GlobalData = GlobalTensor<float, DynShapeDim5, DynStridDim5>; using TileData = Tile<TileType::Vec, float, 64, 64, BLayout::RowMajor, 64, 64>; TileData src0Tile(64, 64); TileData src1Tile(64, 64); TileData dstTile(64, 64); /* TAssign only in manual mode */ TASSIGN(src0Tile, 0x0); TASSIGN(src1Tile, 0x10000); TASSIGN(dstTile, 0x20000); GlobalData src0Global(src0); GlobalData src1Global(src1); GlobalData dstGlobal(out); /* event model only in manual mode */ Event<Op::TLOAD, Op::TADD> event0; Event<Op::TADD, Op::TSTORE_VEC> event1; TLOAD(src0Tile, src0Global); event0 = TLOAD(src1Tile, src1Global); event1 = TADD(dstTile, src0Tile, src1Tile, event0); TSTORE(dstGlobal, dstTile, event1); }

对比可见：manual 版本需要手写 3 条TASSIGN分配地址、声明 2 个Event对象，并手工把事件 token 从TLOAD一路传递到TSTORE；而 auto 版本只有纯粹的 load → compute → store 数据流，其余全部交给编译器。

4.3 Tile 内存分配（Tile Memory Allocation）

在 PTO 默认（manual）编译模式下，实例化Tile变量之后，还需要用TASSIGN指令手动为它绑定一个专用的 buffer 地址。而 auto mode 下这一步被省略：只要实例化Tile变量，编译器就会自动在底层为它分配 buffer 地址。

同样的差异也体现在矩阵乘TMATMUL的对比中（完整代码见 auto 模式示例文档）。下面分别给出 auto 与 manual 两个版本的核心片段：

Auto mode 版本（TMATMUL）：

template <typename cType, typename aType, typename bType, typename fbType, typename l0cType, int M, int K, int N, int ValidM, int ValidK, int ValidN> __global__ AICORE void runTMatMul(__gm__ cType *out, __gm__ aType *src0, __gm__ bType *src1, __gm__ fbType *src2) { using GlobalDataSrc0 = GlobalTensor<aType, pto::Shape<1, 1, 1, ValidM, ValidK>, pto::Stride<ValidM * ValidK, ValidM * ValidK, ValidM * ValidK, ValidK, 1>>; using GlobalDataSrc1 = GlobalTensor<bType, pto::Shape<1, 1, 1, ValidK, ValidN>, pto::Stride<ValidK * ValidN, ValidK * ValidN, ValidK * ValidN, ValidN, 1>>; using GlobalDataSrc2 = GlobalTensor<fbType, pto::Shape<1, 1, 1, 1, ValidN>, pto::Stride<ValidN, ValidN, ValidN, ValidN, 1>>; using GlobalDataOut = GlobalTensor<cType, pto::Shape<1, 1, 1, ValidM, ValidN>, pto::Stride<ValidM * ValidN, ValidM * ValidN, ValidM * ValidN, ValidN, 1>>; GlobalDataSrc0 src0Global(src0); GlobalDataSrc1 src1Global(src1); GlobalDataSrc2 src2Global(src2); GlobalDataOut dstGlobal(out); using TileMatAData = Tile<TileType::Mat, aType, M, K, BLayout::ColMajor, ValidM, ValidK, SLayout::RowMajor, 512>; using TileMatBData = Tile<TileType::Mat, bType, K, N, BLayout::ColMajor, ValidK, ValidN, SLayout::RowMajor, 512>; using TileMatFbData = Tile<TileType::Mat, fbType, 1, N, BLayout::RowMajor, 1, ValidN, SLayout::NoneBox>; using LeftTile = TileLeft<aType, M, K, ValidM, ValidK>; using RightTile = TileRight<bType, K, N, ValidK, ValidN>; using AccTile = TileAcc<l0cType, M, N, ValidM, ValidN>; using FbTile = Tile<TileType::Scaling, fbType, 1, N, BLayout::RowMajor, 1, ValidN, SLayout::NoneBox>; TileMatAData aMatTile; TileMatBData bMatTile; TileMatFbData fbMatTile; LeftTile aTile; RightTile bTile; AccTile cTile; FbTile fbTile; TLOAD(aMatTile, src0Global); TLOAD(bMatTile, src1Global); TLOAD(fbMatTile, src2Global); /**************************TMOV & TMATMUL**************************/ TMOV(aTile, aMatTile); TMOV(bTile, bMatTile); TMATMUL(cTile, aTile, bTile); TMOV(fbTile, fbMatTile); /********************************TSTORE****************************/ TSTORE_FP<AccTile, GlobalDataOut, FbTile>(dstGlobal, cTile, fbTile); }

Manual mode 版本（TMATMUL 差异部分）：

/* TAssign only in manual mode */ TASSIGN(aMatTile, 0x0); TASSIGN(bMatTile, 0x10000); TASSIGN(fbMatTile, 0x20000); LeftTile aTile; RightTile bTile; AccTile cTile; FbTile fbTile; /* TAssign only in manual mode */ TASSIGN(aTile, 0x0); TASSIGN(bTile, 0x0); TASSIGN(cTile, 0x0); TASSIGN(fbTile, 0x0); /* event model only in manual mode */ Event<Op::TLOAD, Op::TMOV_M2L> evtLoad_Mov; Event<Op::TMOV_M2B, Op::TMATMUL> evtMov_Matmul; Event<Op::TMATMUL, Op::TMOV_M2S> evtMatmul_MovM2s; TLOAD(aMatTile, src0Global); TLOAD(bMatTile, src1Global); evtLoad_Mov = TLOAD(fbMatTile, src2Global); /**************************TMOV & TMATMUL**************************/ TMOV(aTile, aMatTile, evtLoad_Mov); evtMov_Matmul = TMOV(bTile, bMatTile); evtMatmul_MovM2s = TMATMUL(cTile, aTile, bTile, evtMov_Matmul); TMOV(fbTile, fbMatTile, evtMatmul_MovM2s); /********************************TSTORE****************************/ TSTORE_FP<AccTile, GlobalDataOut, FbTile>(dstGlobal, cTile, fbTile);

可以看到，矩阵乘这类横跨 L1（TileMatA/B/Fb）、L0（LeftTile/RightTile/AccTile/FbTile）多个缓冲区的复杂 kernel，manual 模式需要多达 7 条TASSIGN和 3 个跨 pipeline 的Event；auto 模式则全部消解。

五、使用 Ascend CANN 编译 Auto Mode 代码

5.1 编译选项

要用 auto mode 编译 kernel，只需在 Bisheng CCE 工具链中使能 PTO auto mode 编译 pass，即追加两个编译选项：

--cce-pto-enable --cce-pto-auto-enable

此外有两个需要牢记的限制（见 auto 模式 README 与 demos 说明）：

• auto mode 目前只支持-O2优化选项；
• 需要根据目标 SoC 指定--cce-aicore-arch（本仓库示例中使用的取值包括dav-c220-vec、dav-c310-vec等）。

5.2 Device 侧编译示例

编译单个 CCE kernel 源文件为 object 文件：

source /usr/local/Ascend/ascend-toolkit/latest/bin/setenv.bash bisheng -c -x cce -O2 --cce-aicore-only \ --cce-aicore-arch=dav-c310-vec \ -std=c++17 \ --cce-pto-enable \ --cce-pto-auto-enable \ kernel.cpp -o kernel.o

参数说明：

• -c -x cce：将输入视为 CCE 源码并只编译不链接；
• -O2：auto mode 当前支持的唯一优化级别（必须使用）；
• --cce-aicore-only：仅生成 AI Core 侧代码；
• --cce-aicore-arch=dav-c310-vec：指定目标 SoC 的 AI Core 架构；
• -std=c++17：按 C++17 标准编译；
• --cce-pto-enable：使能 PTO 编译通道；
• --cce-pto-auto-enable：在 PTO 通道上叠加 auto mode 编译。

5.3 在工程构建（CMake）中使能

除了直接调用 bisheng，auto mode 也可以集成进ascendc_library工程。仓库中的 demos/auto_mode/baseline/add/CMakeLists.txt 展示了标准做法：

ascendc_library(no_workspace_kernel STATIC csrc/kernel/add_custom.cpp ) ascendc_compile_options(no_workspace_kernel PRIVATE --cce-pto-enable --cce-pto-auto-enable -O2)

同样需要注意：-O2必须显式带上，否则不满足 auto mode 的编译前提。

5.4 一个可运行的最小例子：TMUL

auto 模式 README 给出了元素级乘法TMUL的最小对比。Auto mode 版本如下，去掉全部手动同步与地址分配：

template <typename T, int kGRows_, int kGCols_, int kTRows_, int kTCols_> __global__ AICORE void runTMul(__gm__ T __out__ *out, __gm__ T __in__ *src0, __gm__ T __in__ *src1) { using DynShapeDim5 = Shape<1, 1, 1, kGRows_, kGCols_>; using DynStridDim5 = Stride<1, 1, 1, kGCols_, 1>; using GlobalData = GlobalTensor<T, DynShapeDim5, DynStridDim5>; using TileData = Tile<TileType::Vec, T, kTRows_, kTCols_, BLayout::RowMajor, -1, -1>; TileData src0Tile(kGRows_, kGCols_); TileData src1Tile(kGRows_, kGCols_); TileData dstTile(kGRows_, kGCols_); int offset = (block_idx / 4) * (64 * 16) + (block_idx % 4) * 16; GlobalData src0Global(src0 + offset); GlobalData src1Global(src1 + offset); GlobalData dstGlobal(out + offset); TLOAD(src0Tile, src0Global); TLOAD(src1Tile, src1Global); TMUL(dstTile, src0Tile, src1Tile); TSTORE(dstGlobal, dstTile); out = dstGlobal.data(); }

而manual 版本则需要显式的TASSIGN和set_flag/wait_flag同步：

template <typename T, int kGRows_, int kGCols_, int kTRows_, int kTCols_> __global__ AICORE void runTMul(__gm__ T __out__ *out, __gm__ T __in__ *src0, __gm__ T __in__ *src1) { using DynShapeDim5 = Shape<1, 1, 1, kGRows_, kGCols_>; using DynStridDim5 = Stride<1, 1, 1, kGCols_, 1>; using GlobalData = GlobalTensor<T, DynShapeDim5, DynStridDim5>; using TileData = Tile<TileType::Vec, T, kTRows_, kTCols_, BLayout::RowMajor, -1, -1>; TileData src0Tile(kGRows_, kGCols_); TileData src1Tile(kGRows_, kGCols_); TileData dstTile(kGRows_, kGCols_); TASSIGN(src0Tile, 0x0 + 0x400 * block_idx); TASSIGN(src1Tile, 0x4000 + 0x400 * block_idx); TASSIGN(dstTile, 0x8000 + 0x400 * block_idx); int offset = (block_idx / 4) * (64 * 16) + (block_idx % 4) * 16; GlobalData src0Global(src0 + offset); GlobalData src1Global(src1 + offset); GlobalData dstGlobal(out + offset); TLOAD(src0Tile, src0Global); TLOAD(src1Tile, src1Global); set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); TMUL(dstTile, src0Tile, src1Tile); set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); TSTORE(dstGlobal, dstTile); out = dstGlobal.data(); }

manual 版本中程序员必须手工计算并硬编码 buffer 地址（0x0 + 0x400 * block_idx等），还需显式插入两组set_flag/wait_flag来保证TLOAD → TMUL → TSTORE的 pipeline 顺序——这些在 auto 模式下全部由编译器接管。

六、Auto Mode 下的开发规则与限制（实战要点）

auto mode 能自动完成同步与分配的前提是：代码的可分析性。违反以下规则可能导致编译失败、功能错误（如精度问题）或性能退化。这些规则完整收录于 kernel 开发者规则与限制 与 库开发者规则与限制。

6.1 控制流规则

复杂的控制流（尤其是循环内部）会让编译器难以进行精确的跨 pipeline 并行与双缓冲优化。由于编译器必须保证程序正确性，它可能生成更保守的同步操作，从而带来性能损失。具体建议：

（1）首末迭代的守卫条件应可静态求值。把循环首、末迭代的守卫写成能被编译器静态评估的形式，编译器就能自动对首末迭代做 peel，从而大幅简化自动同步。例如：

for (int tile_id = 0; tile_id < total_tiles; tile_id++) { if (tile_id == 0) { TLOAD(srcTile, globalSrc); } ... if (tile_id == total_tiles-1) { TSTORE(globalDst, dstTile); } }

（2）循环嵌套中的循环不变量控制流应留在内层。不依赖内层循环归纳变量的 if 语句，应保留在内层循环内部，而非提到外层：

// 推荐：把 if 保留在内层循环里 for (int tile_id = 0; tile_id < total_tiles; tile_id++) { int next_tile = tile_id < total_tiles-1 ? tile_id + 1 : -1; ... for (int subtile_id = 0; subtile_id < total_subtiles; subtile_id++) { if (next_tile != -1) { ... // computation here } } }

（3）if 语句中的复杂逻辑表达式应预先求值。守卫 PTO 指令的复杂逻辑表达式，强烈建议先求值成 bool 变量再使用：

bool cond = (srcTile.GetValidRow() > 16 || srcTile.GetValidCol() > 16) && srcTile.GetKAligned(); if (cond) { TLOAD(srcTile, globalSrc1); } else { TLOAD(srcTile, globalSrc0); }

（4）当前阶段强烈建议不要使用双/多缓冲（double/multi buffering）。一旦 kernel 变复杂，双缓冲几乎必然引入复杂控制流，给编译器的自动同步带来巨大挑战。auto mode 团队正在设计带约束的专用抽象/接口来支持双缓冲，但在那之前应避免使用。

6.2 内存分配规则

（1）用TRESHAPE表达两个 Tile 同基地址（别名）。auto mode 下直接对两个 Tile 写TASSIGN(tileA, 0x0); TASSIGN(tileB, 0x0);是无效的；正确做法是：

// Invalid in auto mode TASSIGN(tileA, 0x0); TASSIGN(tileB, 0x0); // Correct in auto mode TRESHAPE(tileB, tileA);

（2）用 sub-tile aliasing 表达"Tile B 是 Tile A 的子视图"。用于表达addr(tileB) = addr(tileA) + 偏移，且偏移可以是运行时变量：

uint16_t rowOffset, colOffset; // can be runtime variable // addr(tileB) = addr(tileA) + offsets TileData tileA(...); TileData tileB(...); // Invalid in auto mode TASSIGN(tileA, 0x0); TASSIGN(tileB, 0x0 + rowOffset * TileData::Col + colOffset * 1 + sizeof(T)); // Correct for auto mode sub-tile aliasing(tileB, tileA, rowOffset, colOffset);

（3）auto mode 下 Tile 的内存地址在运行时不可改变。编译器会给每个声明的 Tile 变量分配恒定地址，无法像 manual 模式那样在循环里通过TASSIGN(tile, 0x100 * i)动态改址。auto mode 下每个 Tile 只会被分配一次内存。这里有一个关键心智模型：把 Tile 想象成 C++ 引用——一旦声明，其内存地址就已确定，不能再改变。

（4）正确理解TRESHAPE与 sub-tile aliasing 在 auto mode 下的语义。manual 模式下它们都是真正重新赋址的指令，而 auto mode 下它们只是向编译器提供两个 Tile 如何别名（alias）的提示：

• TRESHAPE：auto mode 下把源、目标 Tile 绑定到同一地址，绑定在整个作用域内有效（manual 模式则是在指令执行点赋址）；
• sub-tile aliasing：编译器计算子 Tile 的相对偏移，并加到基 Tile 的自动分配地址上。

由于 auto mode 下 Tile 地址在其作用域内不可变，一个 Tile 不能作为多条TRESHAPE/sub-tile aliasing 的目标（行为未定义）。好的实践是：把TRESHAPE/sub-tile aliasing 紧跟在目标 Tile 的声明之后。

6.3 通用规则

（1）不要在目标（输出）Tile 上调用TLOAD。如果一个 Tile 仅作输出、不需要从 GM 拷贝数据，就绝不应对它TLOAD。原因：auto mode 下编译器基于 liveness 复用内存——若dstTile与srcTile的活跃区间不重叠，编译器可能给它们分配同一地址，此时两个并发TLOAD会互相覆盖，产生数据竞争。

（2）不要直接调用 CCE intrinsics。原因有二：一是 CCE intrinsics 接收裸指针参数，而 auto mode 下 Tile 以向量类型（而非指针类型）表示，无法编译；二是自动内存分配与同步只建立在 PTO 指令分析之上，编译器无法识别其他指令。因此，不要在 kernel 里直接调用 Tile 的.data()成员函数——它本质上是给库开发者用的接口。若确实没有对应的 PTO 指令可用，应向 pto-isa 提交新增 PTO 指令的请求。

（3）优先使用PtoSetWaitFlag或 Event 同步，而不是裸的set_flag/wait_flag。PtoSetWaitFlag与 Event 同步在内部已对 manual/auto 两种模式做了防护：auto mode 编译时该接口是 no-op，不会与编译器的自动同步冲突。直接调用set_flag/wait_flag则必须手动用__PTO_AUTO__宏包起来，繁琐且易错。

6.4 库开发者规则要点

对于 pto-isa 库的维护者（实现 PTO 指令/库的人），库开发者规则与限制 还提出了额外的实现约束：

• .data()的返回类型在 auto mode 下是向量类型而非指针：pto::(Conv)Tile::data()返回TileDType，在 auto mode 下定义为向量类型，不能当作裸指针使用；
• 避免对结构体/类成员做默认初始化：默认初始化会让编译器的 SROA pass 无法消除AllocaInst及关联的 load/store，建议用#ifdef __PTO_AUTO__区分写法，并尽量采用兼容 C 语言的 POD 聚合编程；
• tile function 及其调用链内部仍需显式同步：使用set_flag、wait_flag或pipe_barrier；tile function 之外使用PtoSetWaitFlag（在 auto mode 下为 no-op）；
• 实现中避免使用TASSIGN：某些指令实现直接用了TASSIGN_IMPL，它在 auto mode 下是 no-op；若只是表达别名应改用TRESHAPE；
• *_IMPL函数约束：函数签名需带PTO_INTERNAL宏；实现内直接调用 tile function（除非内联，否则不调用非 tile 函数）；通过.data()传参或对.data()的返回值按引用返回（auto &src = srcTile.data();正确，auto dst = dstTile.data();错误）；
• tile function 参数规则：参数类型用typename <...>::TileDType而非DType *；按值传递；正确附加__in__/__out__属性；用__cce_get_tile_ptr获取底层 buffer 指针；返回值必须为void（其余返回值改为按值传出参数）；
• 避免在 tile function 前出现运行时控制流：如TROWEXPANDDIV_IMPL、TMULS_IMPL所示，运行时条件会严重干扰自动同步，应尽量移除或移入 tile function 内部。

七、动手实践：完整的 Auto Mode 工程示例

仓库的 demos/auto_mode/baseline/add 提供了一个开箱即用的 auto mode + PyTorch 自定义算子（torch_npu+KERNEL_LAUNCH）示例。其 kernel 源码 add_custom.cpp 正是 auto 模式写法的完整示范——定义Tile、TLOAD、TADD、TSTORE，全程无TASSIGN、无显式同步：

using TileData = Tile<TileType::Vec, T, bTileRows, bTileCols, BLayout::RowMajor, -1, -1>; TileData xTile(bTileRows, bTileCols), yTile(bTileRows, bTileCols), zTile(bTileRows, bTileCols); TLOAD(xTile, xGlobal); TLOAD(yTile, yGlobal); TADD(zTile, xTile, yTile); TSTORE(zGlobal, zTile);

该 kernel 通过extern "C" __global__ AICORE void add_custom(...)作为算子入口，host 侧在 csrc/host/my_add.cpp 中用TORCH_LIBRARY_FRAGMENT(npu, m)声明my_add算子、通过ACLRT_LAUNCH_KERNEL（示例封装为EXEC_KERNEL_CMD）下发执行。

构建与运行流程（详见 示例 README）：

# 1. 设置环境与 PTO 库路径 export ASCEND_HOME_PATH=/usr/local/Ascend/ source /usr/local/Ascend/ascend-toolkit/set_env.sh export PTO_LIB_PATH=[YOUR_PATH]/pto-isa # 2. 构建 wheel rm -rf build op_extension.egg-info python3 setup.py bdist_wheel # 3. 安装 wheel cd dist pip uninstall *.whl pip install *.whl # 4. 运行测试 cd test python3 test.py

注意：需要在 CMakeLists.txt 中把SOC_VERSION设置为目标 SoC（示例注释提到 A2A3 对应Ascend910B1），可用npu_smi info查询芯片名后以Ascend<芯片名>形式填写；编译选项务必包含--cce-pto-enable --cce-pto-auto-enable -O2。该示例当前不使用双缓冲，也再次印证了"auto mode 下暂不建议使用双缓冲"的约束。

八、总结

PTO Auto Mode 是 CANN pto-isa 面向生产力的一次关键抽象：它以 Tile liveness 分析为核心，在 TF 抽象层之上自动完成跨 pipeline 同步与 Tile 片上内存分配，使 kernel 开发者得以用"声明 Tile → 数据流计算 → 存储"的纯数据流方式编写 kernel，同时保持与 manual 模式接近的性能，并天然获得跨昇腾架构代际的源码级兼容。

使用上只需记住三个要点：编译时追加--cce-pto-enable --cce-pto-auto-enable、必须使用-O2、按目标 SoC 指定--cce-aicore-arch；编码时遵循本文第六节的规则——可静态分析的控制流、用TRESHAPE/sub-tile aliasing 表达别名、不在输出 Tile 上调TLOAD、不直接调用 CCE intrinsics、tile function 内部仍手动同步。

如果你想进一步深入，推荐按顺序阅读本仓库的以下材料：auto 模式与 manual 模式的完整代码对照见 Examples.md；kernel 开发者约束见 Kernel_Developer_Rules_And_Limitations.md；库开发者约束见 Library_Developer_Rules_And_Limitations.md；manual 模式下的事件同步模型见 Event 文档；可直接运行的最小示例见 demos/auto_mode/baseline/add。

• 人工智能
• 指令集
• 算子库
• CANN
• Ascend