---
tags: [CANN]
doc_type: knowledge-note
返回目录: "[[计算机系/昇腾与CANN/大模型与算子专题学习导航|大模型与算子专题学习导航]]"
上一篇: "[[计算机系/昇腾与CANN/Ascend C接入PyTorch单算子|Ascend C接入PyTorch单算子]]"
下一篇: "[[计算机系/昇腾与CANN/Ascend C调试与性能调优|Ascend C调试与性能调优]]"
前置知识: "[[计算机系/昇腾与CANN/Ascend C接入PyTorch单算子|Ascend C接入PyTorch单算子]]"
---

# Ascend C算子入图与图执行

单算子接入解决“框架如何调用这个函数”；入图进一步解决“如何描述它与其他算子的依赖，以及如何编译或重放整个执行过程”。前置的 Schema 和设备注册见 [PyTorch 单算子](Ascend%20C接入PyTorch单算子.md)。

## 1. 三条路线的作用不同

| 路线 | 处理对象 | 主要收益 | 需要提供什么 |
| --- | --- | --- | --- |
| PyTorch + TorchAir + GE | 从模型捕获的计算图 | 图级优化与编译执行 | Schema、Meta、转换、GE 实现 |
| AclGraph | 已下发到 stream 的任务序列 | 减少重复 Host 调度开销 | 可捕获的操作与稳定有效的资源 |
| 直接构建 GE 图 | Ascend IR 中的 Tensor、Node、Graph | 图编译、资源规划和执行 | 算子描述、图连接与执行实现 |

TorchAir 与 GE 往往处在同一链路的上下游，而不是互斥方案。AclGraph 重放任务不等于自动重写 Kernel；图编译有机会进行融合等优化，但自定义不透明节点也可能限制融合范围。

## 2. torch.compile 如何获取计算图

Eager 模式按程序调用逐步提交操作，便于调试；NPU 操作仍可异步执行，并不意味着每条 Python 语句都等待硬件完成。

`torch.compile` 的前端 TorchDynamo 通过 Python 帧求值机制分析字节码，捕获可编译的算子子图，并用 guards 约束此次编译成立的输入和状态。某些形状、类型或控制流变化可能触发重新编译。

默认后端常涉及 AOTAutograd 和 TorchInductor，但采用 TorchAir 时，下游是昇腾的转换与 GE 执行路径，不能一概说最后都会生成 Triton Kernel。

![PyTorch、TorchAir 与 GE 的关系](.assets/AscendC进阶/graph-11.png)

不支持的 Python 操作可能造成 graph break；默认配置可能分段捕获和回退，`fullgraph=True` 则要求完整捕获，遇到图中断通常报错。动态控制流并非一律能无代价处理。

## 3. 自定义算子进入 TorchAir 的四层契约

### 3.1 Schema：让图知道这是哪一个算子

例如 `ascendc_add(Tensor x, Tensor y) -> Tensor`。Schema 定义已经在单算子篇展开，同一命名空间/算子不能在多个库中重复定义而不做注册组织。

### 3.2 Meta/Fake：推导输出描述

编译阶段需要知道输出 shape、dtype 等信息，却不应执行真实 NPU 数据计算。对于只支持同形 float32 Add 的实现：

```cpp
at::Tensor AscendcAddMeta(const at::Tensor& x, const at::Tensor& y)
{
    TORCH_CHECK(x.sizes() == y.sizes(), "shape mismatch");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                y.scalar_type() == at::kFloat, "float32 required");
    return at::empty_like(x);
}

TORCH_LIBRARY_IMPL(ascendc_ops, Meta, m)
{
    m.impl("ascendc_add", TORCH_FN(AscendcAddMeta));
}
```

如果实际 Kernel 只支持 `[8,2048]`，Meta 和运行时都要体现这个约束。不能让 Meta 宣称任意形状都能输出，实际 Kernel 却仍按固定尺寸访问。

### 3.3 Converter：把 FX 节点转换为 GE 节点

```python
@torchair.register_fx_node_ge_converter(
    torch.ops.ascendc_ops.ascendc_add.default
)
def convert_ascendc_add(x, y, meta_outputs=None):
    return torchair.ge.custom_op(
        "AddCustom", inputs={"x": x, "y": y}, outputs=["z"]
    )
```

转换器表达图节点映射，不在 Python 中计算 `x+y`。输入名、输出名和 GE 原型必须一致。

### 3.4 GE 实现：在图运行时执行 Kernel

对应的 GE 扩展方案通过 `EagerExecuteOp` 与 `ShapeInferOp`，实现 `Execute`、`InferShape`、`InferDataType`，并注册 `AddCustom` 原型与执行映射。

`Execute` 取得输入、校验约束、申请输出，在 `ctx->GetStream()` 上启动 Ascend C Kernel；另外两个接口供 GE 推导 Tensor 描述。PyTorch Meta 与 GE 推导是不同阶段的契约，都应与实际实现保持一致。

这些扩展类和注册方式依赖配套 CANN/GE 版本；使用具有相应接口的开发环境，不能把新接口名称直接套进旧工具包。

## 4. 编译与验证图模式

加载注册库、GE 扩展及转换器后，可用以下方式编译模型：

```python
class AddModel(torch.nn.Module):
    def forward(self, x, y):
        return torch.ops.ascendc_ops.ascendc_add(x, y)

config = torchair.CompilerConfig()
backend = torchair.get_npu_backend(compiler_config=config)
model = torch.compile(
    AddModel().npu(), backend=backend, fullgraph=True, dynamic=False
)
```

`dynamic=False` 不会让不支持动态形状的 Kernel 自动获得动态能力。应先验证单算子，再验证整图结果；性能对比要分开首次编译与后续稳定执行。

遇到错误时按层定位：算子找不到查加载/Schema；Fake 推导失败查 Meta；转换失败查 converter；GE 不识别节点查原型和映射；执行错误查资源、Kernel 和输入约束。

## 5. AclGraph：捕获一次，多次重放

![AclGraph 捕获和重放](.assets/AscendC进阶/graph-17.png)

对于许多短 Kernel，Host 逐次提交的开销可能占较大比例。AclGraph 将一段支持捕获的 stream 操作记录为模型运行实例，随后一次提交重放。

```cpp
aclmdlRICaptureBegin(mainStream, ACL_MODEL_RI_CAPTURE_MODE_GLOBAL);
aclrtMemcpyAsync(xDevice, bytes, xHost, bytes,
                 ACL_MEMCPY_HOST_TO_DEVICE, mainStream);
aclrtMemcpyAsync(yDevice, bytes, yHost, bytes,
                 ACL_MEMCPY_HOST_TO_DEVICE, mainStream);
add_custom<blockLength><<<numBlocks, 0, mainStream>>>(
    xDevice, yDevice, zDevice);
aclrtMemcpyAsync(zHost, bytes, zDevice, bytes,
                 ACL_MEMCPY_DEVICE_TO_HOST, mainStream);
aclmdlRICaptureEnd(mainStream, &modelRI);

for (int i = 0; i < 5; ++i) {
    aclmdlRIExecuteAsync(modelRI, mainStream);
    aclrtSynchronizeStream(mainStream);
}
aclmdlRIDestroy(modelRI);
```

这里的设备/Host 缓冲、stream 和 `modelRI` 需要事先创建，异步拷贝使用符合运行时要求的 Host 内存，并检查接口返回值。

重放依赖记录下来的地址和执行配置。输入数值可通过有效缓冲更新，但不能在上次异步执行仍读取它时覆盖。shape、地址或任务结构变化时，要按支持的更新机制处理或重新捕获。

销毁运行实例和释放缓冲前要确保相关任务已完成。循环内同步便于逐轮读取结果；吞吐测试则应根据真正的数据依赖安排同步，避免把过多 Host 等待计入调度收益。

## 6. 直接构建 GE 图

Ascend IR 的三个核心对象是：

- Tensor：数据及 name、shape、dtype、format 等描述。
- Node/Operator：算子类型、名字、输入输出与属性。
- Graph：节点集合、连边、图输入与输出。

图可以通过三种方式构建：解析 ONNX/PB 等模型、按算子原型手工连接、通过 ES 函数式接口构图。模型解析仍需完成算子映射，文件扩展名受支持不等于其中每个自定义算子都有实现。

例如两个 Add 节点：

```text
x ──┐
    Add0 ──┐
y ──┘      Add1 → output
y ─────────┘
```

最终输出为 $x+2y$。构图时创建 x/y 的 Data 节点，连接 Add0 和 Add1，再设置 Graph 的输入为 x/y、输出为 Add1。运行前还要初始化 GE、创建 Session、加入图及准备 Tensor；`RunGraph` 完成后检查状态和输出，最后移除图并释放资源。

### ES 的“Eager Style”指语法风格

![ES 构图组成](.assets/AscendC进阶/graph-24.png)

`EsGraphBuilder`、`EsTensorHolder` 和由原型生成的构图 API，让用户通过函数返回值表达连边。它更接近写普通函数组合，但这里仍是在构建图，并非每调用一次函数就立刻运行对应 NPU Kernel。

### 图编译在做什么

![GE 图编译与优化](.assets/AscendC进阶/graph-25.png)

主要过程包括输出描述推导、常量折叠和冗余消除，按执行引擎拆分子图，进行支持的图优化，规划内存和 stream，再形成可执行产物。离线部署路径可生成 `.om`；在线图执行不要求用户每次先手工生成 `.om` 文件。

判断优化收益时，区分算子计算时间、Host 下发时间、编译时间和数据拷贝时间。Kernel 本身很慢时，仅使用捕获重放并不会解决其访存或计算瓶颈。

样例入口：[TorchAir + GE](https://gitcode.com/cann/asc-devkit/tree/master/examples/01_simd_cpp_api/02_features/00_framework/00_pytorch/ge_torchair)、[GE](https://gitcode.com/cann/asc-devkit/tree/master/examples/01_simd_cpp_api/02_features/00_framework/03_ge)、[AclGraph](https://gitcode.com/cann/asc-devkit/tree/master/examples/01_simd_cpp_api/02_features/00_framework/04_aclgraph)。