# add_op

**Repository Path**: arepse/add_op

## Basic Information

- **Project Name**: add_op
- **Description**: Ascend C 向量加法算子 (z = x + y)，aclnn 两阶段 API + PyTorch binding，含 ascend_ops fast kernel launch 性能对比
- **Primary Language**: Unknown
- **License**: Apache-2.0
- **Default Branch**: master
- **Homepage**: None
- **GVP Project**: No

## Statistics

- **Stars**: 0
- **Forks**: 0
- **Created**: 2026-05-26
- **Last Updated**: 2026-05-27

## Categories & Tags

**Categories**: Uncategorized

**Tags**: None

## README

# Add 算子开发笔记

## 1. 项目概述

基于 Ascend C 实现一个简单的向量加法算子 `z = x + y`，完成 CPU 验证和 NPU 注册流程。

**环境**: CANN 8.5.0, Ascend910_9392 (48 AIV cores), `dav-c220-vec` / `dav-2201` 架构, PyTorch 2.9.0 + torch_npu 2.9.0

## 2. 文件结构

```
add_op/
  add_custom.cpp      — Ascend C kernel（TPipe/TQue 三段式流水）
  main.cpp            — CPU 模式测试主程序（tikicpulib）
  main_npu.cpp        — NPU 模式测试主程序（ccec 混合编译）
  CMakeLists.txt      — 构建配置（CPU + NPU target）
  run.sh              — 编译运行脚本
  run_profiling.sh    — msprof simulator profiling 脚本
  ascendc/
    src/add_custom/
      CMakeLists.txt
      op_host/                 # Schema、InferShape、Tiling
        add_custom_def.cpp
        add_custom_proto.cpp
        add_custom_tiling.h
        add_custom_tiling.cpp
      op_kernel/               # Ascend C kernel (三段式流水)
        add_custom.cpp
      op_api/                  # aclnn API 层 (两阶段调用)
        add_custom.h / .cpp
        aclnn_add_custom.h / .cpp
    torch_binding/
      npu_add_custom.cpp       # torch.ops.custom.add_custom binding
    build.sh                   # 构建 .run 包 (借用 ASVG)
    install.sh                 # 安装到 CANN OPP
    test.py                    # torch 端到端验证
    benchmark.py               # 1D/2D/3D NPU 性能对比
    README.md
```

## 3. Kernel 实现

多核三段式流水：最多 48 核 × 8 tiles/core

```
CopyIn (GM→UB) → Compute (Add) → CopyOut (UB→GM)
```

- 每个 block 处理 `coreDataNum` 个元素（tiling 在 host 侧计算）
- 使用 `TPipe`, `TQue<VECIN>`, `TQue<VECOUT>`, `GlobalTensor`, `LocalTensor`
- kernel binary 仅注册 1D simplifiedKey；多维输入在 torch binding 层 flatten→1D→reshape 透传

## 4. 编译命令

### CPU 模式（g++ + tikicpulib）

```bash
g++ -std=c++17 -g \
  -D__CCE_AICORE__=220 -D__DAV_C220__ -D__CCE_KT_TEST__=1 \
  -DASCENDC_CPU_DEBUG=1 -DASCENDC_DEBUG=1 -D_GLIBCXX_USE_CXX11_ABI=0 \
  -I${ASCEND_HOME}/compiler/tikcpp/tikcfw \
  -I${ASCEND_HOME}/compiler/tikcpp/tikcfw/impl \
  -I${ASCEND_HOME}/compiler/tikcpp/tikcfw/interface \
  -I${ASCEND_HOME}/tools/tikicpulib/lib/include \
  -I${ASCEND_HOME}/acllib/include \
  -L${ASCEND_HOME}/tools/tikicpulib/lib/Ascend910_9392 \
  -L${ASCEND_HOME}/tools/simulator/Ascend910_9392/lib \
  -L${ASCEND_HOME}/atc/lib64 \
  add_custom.cpp main.cpp \
  -ltikicpulib_stubreg -ltikcpp_debug -lascendcl -lc_sec -lruntime_cmodel \
  -Wl,-rpath,... -o add_custom_cpu
```

### NPU kernel 编译（ccec → .o）

```bash
ccec -x cce add_custom.cpp \
  --cce-aicore-arch=dav-c220-vec --cce-aicore-only --cce-auto-sync \
  -mllvm -cce-aicore-stack-size=0x8000 \
  -mllvm -cce-aicore-function-stack-size=0x8000 \
  -I${ASCEND_HOME}/compiler/tikcpp/tikcfw \
  -I${ASCEND_HOME}/compiler/tikcpp/tikcfw/impl \
  -I${ASCEND_HOME}/compiler/tikcpp/tikcfw/interface \
  -I${ASCEND_HOME}/tools/tikicpulib/lib/include \
  -I${ASCEND_HOME}/acllib/include \
  -O2 -std=c++17 -c -o add_custom.o
```

### NPU 可执行文件（ccec 混合编译）

```bash
ccec -x cce add_custom.cpp -x c++ main_npu.cpp \
  --cce-aicore-arch=dav-c220-vec --cce-auto-sync \
  -mllvm -cce-aicore-stack-size=0x8000 \
  -mllvm -cce-aicore-function-stack-size=0x8000 \
  [includes] [lib paths] \
  -lruntime -lascendcl -lstdc++ -o add_custom_npu
```

## 5. 验证结果

```
CPU 模式: TEST PASSED (16384/16384 correct), 8 cores all succeeded
```

## 6. 开发流程（aclnn 路径）

### 6.1 构建

```bash
# 默认: ASVG_DIR=/data/wc/ASVG, CANN_PATH=/usr/local/Ascend/cann-8.5.0
cd /data/wc/add_op/ascendc
bash build.sh

# 自定义路径
ASVG_DIR=/path/to/ASVG CANN_PATH=/path/to/cann bash build.sh
```

build.sh 流程：
1. 将 `src/add_custom/` 软链接到 ASVG 的 `ops/ascendc/src/`
2. 在 ASVG 的 `func.cmake` 中插入 add_custom 注册行
3. 调用 ASVG 的 `build.sh -n "add_custom"` 编译
4. 拷贝 `.run` 包到本项目 `output/`
5. **trap EXIT** 自动恢复 ASVG 原状

### 6.2 安装

```bash
bash install.sh
source /usr/local/Ascend/cann-8.5.0/opp/vendors/customize/vendors/customize/bin/set_env.bash
```

### 6.3 Torch Binding

`torch_binding/npu_add_custom.cpp` 需编入 ASVG 的 `custom_ops` wheel：

```bash
# 1. 复制 binding
cp torch_binding/npu_add_custom.cpp \
   /data/wc/ASVG/ops/ascendc/torch_ops_extension/custom_ops/csrc/

# 2. 在 ops_def_registration.cpp 的 TORCH_LIBRARY(custom, m) 中添加:
#    m.def("add_custom(Tensor x, Tensor y) -> Tensor");

# 3. 重新安装 wheel
cd /data/wc/ASVG/ops/ascendc/torch_ops_extension
pip install --no-build-isolation --force-reinstall --no-deps .
```

### 6.4 测试

```bash
source /usr/local/Ascend/ascend-toolkit/set_env.sh
source /usr/local/Ascend/cann-8.5.0/opp/vendors/customize/vendors/customize/bin/set_env.bash
python3 /data/wc/add_op/ascendc/test.py
```

### 6.5 正确性测试代码 (test.py)

```python
import torch
import torch_npu
import custom_ops

for N in [64, 256, 1024, 8192, 16384, 100000]:
    x = torch.randn(N, dtype=torch.float32).npu()
    y = torch.randn(N, dtype=torch.float32).npu()
    z = torch.ops.custom.add_custom(x, y)
    expected = x.cpu() + y.cpu()
    max_err = (z.cpu() - expected).abs().max().item()
    print(f"  N={N:>7d}: max_err={max_err:.6e} [{'PASS' if max_err < 1e-5 else 'FAIL'}]")

# 精确值验证
x = torch.tensor([1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0], dtype=torch.float32).npu()
y = torch.tensor([10.0, 20.0, 30.0, 40.0, 50.0, 60.0, 70.0, 80.0], dtype=torch.float32).npu()
z = torch.ops.custom.add_custom(x, y)
print(f"  exact: {z.cpu().tolist()}")  # [11.0, 22.0, 33.0, 44.0, 55.0, 66.0, 77.0, 88.0]
```

### 6.6 两种算子对比测试代码

对比 AddCustom (aclnn 路径) 与 ascend_ops.add (fast kernel launch) 的正确性和性能：

```python
import torch
import torch_npu
import custom_ops   # add_custom (aclnn)
import ascend_ops   # fast kernel launch
import time

device = torch.device("npu:0")
sizes = [1024, 8192, 65536, 262144, 1048576]
warmup, repeats = 10, 100

print("=" * 70)
print("正确性对比")
print("=" * 70)

for N in sizes:
    x = torch.randn(N, dtype=torch.float32, device=device)
    y = torch.randn(N, dtype=torch.float32, device=device)
    z_custom = torch.ops.custom.add_custom(x, y)
    z_ascend = torch.ops.ascend_ops.add(x, y)
    max_err = (z_custom - z_ascend).abs().max().item()
    print(f"  N={N:>10d}: max_err={max_err:.6e} [{'MATCH' if max_err == 0 else 'DIFF'}]")

print()
print("=" * 70)
print(f"性能对比 (warmup={warmup}, repeats={repeats})")
print("=" * 70)
print(f"{'N':>10s} | {'add_custom (ms)':>16s} | {'ascend_ops (ms)':>16s} | {'ratio':>6s}")
print("-" * 60)

for N in sizes:
    x = torch.randn(N, dtype=torch.float32, device=device)
    y = torch.randn(N, dtype=torch.float32, device=device)

    # warmup
    for _ in range(warmup):
        torch.ops.custom.add_custom(x, y)
        torch.ops.ascend_ops.add(x, y)
    torch.npu.synchronize()

    # benchmark add_custom (aclnn)
    torch.npu.synchronize()
    t0 = time.perf_counter()
    for _ in range(repeats):
        torch.ops.custom.add_custom(x, y)
    torch.npu.synchronize()
    t_custom = (time.perf_counter() - t0) / repeats * 1000

    # benchmark ascend_ops (fast kernel launch)
    torch.npu.synchronize()
    t0 = time.perf_counter()
    for _ in range(repeats):
        torch.ops.ascend_ops.add(x, y)
    torch.npu.synchronize()
    t_ascend = (time.perf_counter() - t0) / repeats * 1000

    ratio = t_custom / t_ascend
    print(f"{N:>10d} | {t_custom:>16.4f} | {t_ascend:>16.4f} | {ratio:>5.2f}x")
```

运行方式：

```bash
source /usr/local/Ascend/ascend-toolkit/set_env.sh
source /usr/local/Ascend/cann-8.5.0/opp/vendors/customize/vendors/customize/bin/set_env.bash
python3 compare_test.py
```

## 7. PyTorch Dispatcher 机制

### 7.1 Dispatcher 分发流程

PyTorch 的所有算子调用都经过 Dispatcher 统一分发。以 `torch.ops.custom.add_custom(x, y)` 为例：

```
Python 调用
  → torch._ops.OpOverload.__call__
    → C++ Dispatcher::call(schema, args)
      → 1. 从输入 tensor 推断 dispatch key set (如 PrivateUse1)
      → 2. 按优先级遍历 dispatch key，查找注册的 kernel
      → 3. 调用匹配的后端实现 (如 add_custom_npu)
```

Dispatch key 的推断规则：
- **设备 key**: 由输入 tensor 的 device 决定（CPU/CUDA/PrivateUse1/...）
- **功能 key**: 叠加在设备 key 之上（Autograd、Functionalize、FuncTorchBatched 等）
- **优先级**: 功能 key > 设备 key。若注册了 `AutogradPrivateUse1`，它在 forward 时先拦截，record autograd graph，再 redispatch 到 `PrivateUse1` 执行实际计算

### 7.2 TORCH_LIBRARY_IMPL 宏注册机制

`TORCH_LIBRARY_IMPL` 是三级注册结构：**namespace × dispatch key × 算子名**

```cpp
// 第一步: 定义 schema (与后端无关)
TORCH_LIBRARY(custom, m) {
    m.def("add_custom(Tensor x, Tensor y) -> Tensor");
}

// 第二步: 为特定 dispatch key 注册实现
TORCH_LIBRARY_IMPL(custom, PrivateUse1, m) {
    m.impl("add_custom", &npu_add_custom_npu);   // NPU 实现
}
TORCH_LIBRARY_IMPL(custom, Meta, m) {
    m.impl("add_custom", &npu_add_custom_meta);  // shape/dtype 推导
}
```

宏展开后生成一个静态初始化器，在 `.so` 加载时自动向全局 Dispatcher 表注册。同一算子可在多个 `.so` 中为不同 key 注册实现，Dispatcher 在运行时按 key 查表调用。

`TORCH_LIBRARY` 与 `TORCH_LIBRARY_FRAGMENT` 的区别：
- `TORCH_LIBRARY(ns, m)`: 创建 namespace，每个 namespace 只能出现一次（ODR），适合集中注册
- `TORCH_LIBRARY_FRAGMENT(ns, m)`: 向已有 namespace 追加定义，可出现多次，适合每个算子自包含注册

### 7.3 PrivateUse1 Device Key 注册原理

PyTorch 预留了 `PrivateUse1`（以及 PrivateUse2/3）作为第三方硬件的扩展点。torch_npu 的注册流程：

```python
# torch_npu 初始化时执行
torch.utils.rename_privateuse1_backend("npu")
```

此调用完成三件事：
1. **设备别名**: `torch.device("npu")` 等价于 `torch.device("privateuseone")`
2. **Module 注入**: 向 `torch` 注入 `torch.npu` 模块（`torch.npu.synchronize()` 等）
3. **Dispatch key 映射**: `PrivateUse1` key 对应所有 `.npu()` tensor 的算子分发

在 C++ 侧，torch_npu 还注册了：
- 设备 guard (`c10::impl::DeviceGuardImpl`)
- 内存分配器 (`c10::Allocator`)
- Generator（随机数）
- Stream/Event

### 7.4 CUDA vs PrivateUse1 (NPU) 对比

| 维度 | CUDA | NPU (torch_npu) |
|------|------|------------------|
| **Dispatch Key** | `CUDA` (内置) | `PrivateUse1` (扩展槽位) |
| **设备类型** | `torch.device("cuda")` | `torch.device("npu")` (别名) |
| **注册宏** | `TORCH_LIBRARY_IMPL(..., CUDA, m)` | `TORCH_LIBRARY_IMPL(..., PrivateUse1, m)` |
| **自动微分 Key** | `AutogradCUDA` (内置) | `AutogradPrivateUse1` (需显式注册) |
| **torch.compile** | 原生 Triton 后端 | 需 torchair 桥接到 GE 图引擎 |
| **集成方式** | PyTorch 源码内编译 | 独立 `.so` 插件，`import torch_npu` 时动态注册 |
| **扩展槽位** | 无限制 | 全局仅 3 个 PrivateUse 槽位 |
| **算子覆盖** | 自动 fallback 到 CPU | 未注册的算子抛 NotImplementedError |

核心差异：CUDA 是 PyTorch 的一等公民，dispatch key、autograd、JIT、compile 全链路内置支持。NPU 通过 PrivateUse1 插件机制接入，在运行时动态注册后端，需要额外适配 torchair (compile)、custom autograd (反向) 等。两者共享同一个 Dispatcher 框架，算子注册和分发的 API 完全一致。

### 7.5 Schema 注册方式对比

**AddCustom (aclnn 路径)**：Schema 在 ASVG 的 `ops_def_registration.cpp` 中集中注册，所有 custom 算子共享一个 `TORCH_LIBRARY(custom, m)` block。新增算子需修改公共文件。

**ascend_ops (fast kernel launch)**：Schema 使用 `TORCH_LIBRARY_FRAGMENT` 在 `.asc` 文件中就地注册，每个算子自包含，无需修改公共文件：

```cpp
TORCH_LIBRARY_FRAGMENT(EXTENSION_MODULE_NAME, m) {
    m.def("add(Tensor x, Tensor y) -> Tensor");  // schema 定义
}
TORCH_LIBRARY_IMPL(EXTENSION_MODULE_NAME, PrivateUse1, m) {
    m.impl("add", &add_npu);                     // NPU 实现
}
TORCH_LIBRARY_IMPL(EXTENSION_MODULE_NAME, Meta, m) {
    m.impl("add", &add_meta);                    // shape 推导
}
```

这种模式的优势：每个 `.asc` 文件包含 schema + kernel + tiling + 注册，新增算子不触及任何公共文件，编译和维护完全解耦。

## 8. 性能对比: AddCustom (aclnn) vs ascend_ops (fast kernel launch)

测试环境: Ascend910_9392, CANN 8.5.0, PyTorch 2.9.0, float32, NPU, warmup=10, repeats=100

### 正确性

1D/2D/3D 全部 9 组 shape 均 bit-exact match（max error = 0）。

### 性能 (100 次取平均)

| 维度 | Shape | 元素数 | add_custom (ms) | ascend_ops (ms) | 比值 |
|------|-------|--------|----------------|-----------------|------|
| 1D | (1024,) | 1K | 0.0358 | 0.0135 | 2.64x |
| 1D | (65536,) | 66K | 0.0362 | 0.0131 | 2.77x |
| 1D | (1048576,) | 1.0M | 0.0355 | 0.0137 | 2.58x |
| 2D | (32, 256) | 8K | 0.0366 | 0.0134 | 2.74x |
| 2D | (128, 1024) | 131K | 0.0365 | 0.0136 | 2.68x |
| 2D | (256, 4096) | 1.0M | 0.0364 | 0.0131 | 2.77x |
| 3D | (4, 128, 256) | 131K | 0.0363 | 0.0133 | 2.73x |
| 3D | (8, 64, 1024) | 524K | 0.0375 | 0.0138 | 2.71x |
| 3D | (16, 32, 2048) | 1.0M | 0.0361 | 0.0131 | 2.75x |

ascend_ops 的 fast kernel launch 在所有维度/数据量下均快约 **2.6-2.8x**，且性能与 shape 维度无关（add_custom 通过 flatten→1D→reshape 透传，实际 kernel 执行路径相同）。性能差异完全来自调用路径开销：aclnn 两阶段（GetWorkspaceSize + Execute）vs fast launch 直接 `<<<blocks, nullptr, stream>>>`。

复现命令：

```bash
source /usr/local/Ascend/ascend-toolkit/set_env.sh
source /usr/local/Ascend/cann-8.5.0/opp/vendors/customize/vendors/customize/bin/set_env.bash
ASCEND_RT_VISIBLE_DEVICES=1 python3 /data/wc/add_op/ascendc/benchmark.py
```

### 架构差异总结

| 维度 | AddCustom (本项目) | ascend_ops.add |
|------|-------------------|----------------|
| 调用路径 | aclnn 两阶段: GetWorkspaceSize → Execute | `<<<blocks, nullptr, stream>>>` 直接启动 |
| 文件数量 | 10 个文件 (op_host / op_kernel / op_api) | 1 个 `.asc` 文件包含全部 |
| Tiling | 独立 `_tiling.cpp`，host 侧计算 | 内联 `calc_tiling_params()`，device 侧 |
| dtype | 仅 float32 | float32 + float16 + int32 |
| 核数 | 固定 48 核 | 动态按数据量分配 |
| 构建 | 借用 ASVG build.sh → .run 包 | 独立 CMake → pip wheel |
| torch 注册 | 需修改 ASVG 的 ops_def_registration.cpp | TORCH_LIBRARY_FRAGMENT 自注册 |
| shape 限制 | 任意 shape (torch binding 层 flatten→1D→reshape) | 任意 shape |

## 9. 参考: ops-math 仓库架构

### 9.1 ascend_ops 构建

[ops-math](https://gitcode.com/cann/ops-math) 的 `examples/fast_kernel_launch_example/` 提供 `ascend_ops` 包。

```bash
source /usr/local/Ascend/ascend-toolkit/set_env.sh
cd /data/wc/ops-math/examples/fast_kernel_launch_example
python3 setup.py clean
NPU_ARCH=dav-2201 python3 -m build --wheel --no-isolation
pip install dist/*.whl --force-reinstall --no-deps
pytest tests/add/test_add.py -v    # 58 tests passed
```

### 9.2 算子组织结构

华为官方 [ops-math](https://gitcode.com/cann/ops-math) 的算子组织（以 `add` 为例）：

```
math/add/
  op_host/     — Schema (_def.cpp)、InferShape (_infershape.cpp)、Tiling
  op_graph/    — GE 图层 proto (_proto.h)、图推导 (_graph_infer.cpp)
  op_kernel/   — Ascend C kernel (_apt.cpp)、DAG 定义
  op_api/      — aclnn API: aclnn_*.{h,cpp} (L2) + *.{h,cpp} (L0)
  CMakeLists.txt
```

**Schema 注册**: `OP_ADD(Add)` + `REG_OP(Add)`
**Meta 函数**: `IMPL_OP_INFERSHAPE(Add)` + `IMPL_OP(Add).InferDataType`
**NPU 实现**: 生产模式用 DAG (`Bind<Vec::Add<T>>`)，简单模式用手写 TPipe/TQue（本项目采用此模式）

### 9.3 Schema 注册

**主注册（运行时）** — `op_host/add_def.cpp`:

```cpp
class Add : public OpDef {
public:
    explicit Add(const char* name) : OpDef(name) {
        this->Input("x1")
            .ParamType(REQUIRED)
            .DataType({ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ...})
            .Format({ge::FORMAT_ND, ...});
        this->Input("x2") ...
        this->Output("y") ...

        OpAICoreConfig aicoreConfig;
        aicoreConfig.DynamicCompileStaticFlag(true)
            .ExtendCfgInfo("opFile.value", "add_apt");  // 链接到 kernel 文件名
        this->AICore().AddConfig("ascend950", aicoreConfig);
    }
};
OP_ADD(Add);
```

- `OP_ADD(OpName)` 宏注册到全局 registry
- `ExtendCfgInfo("opFile.value", "add_apt")` 关联 kernel 源文件 `add_apt.cpp`
- DataType 列表是位置对应的：第 N 个 dtype 组合

**图层注册** — `op_graph/add_proto.h`:

```cpp
REG_OP(Add)
    .INPUT(x1, TensorType({DT_BOOL, DT_FLOAT, DT_INT32, ...}))
    .INPUT(x2, TensorType({...}))
    .OUTPUT(y, TensorType({...}))
    .OP_END_FACTORY_REG(Add)
```

### 9.4 Meta 函数（形状/类型推导）

**Shape 推导** — `op_host/add_infershape.cpp`:

```cpp
static ge::graphStatus InferShapeForAdd(gert::InferShapeContext* context) {
    return Ops::Base::InferShape4Broadcast(context);
}
IMPL_OP_INFERSHAPE(Add).InferShape(InferShapeForAdd);
```

通用工具函数：
- `InferShape4Elewise` — 逐元素 ops: 输出 shape = 输入 shape
- `InferShape4Broadcast` — 广播 ops: 输出 shape = broadcast(input1, input2)
- `InferShape4Reduce` — 归约 ops

**Dtype 推导** — `op_graph/add_graph_infer.cpp`:

```cpp
static ge::graphStatus InferDataTypeForAdd(gert::InferDataTypeContext* context) {
    auto x1Dtype = context->GetInputDataType(0);
    auto x2Dtype = context->GetInputDataType(1);
    context->SetOutputDataType(0, PromoteType(x1Dtype, x2Dtype));
    return ge::GRAPH_SUCCESS;
}
IMPL_OP(Add).InferDataType(InferDataTypeForAdd);
```

### 9.5 NPU 实现（两种模式）

**生产模式 — DAG 框架** (`op_kernel/arch35/add_dag.h`):

```cpp
template <typename T>
struct AddWithoutCastCompute {
    using OpInputX1 = Bind<Vec::CopyInBrc<T>, Placeholder::In0<T>>;
    using OpInputX2 = Bind<Vec::CopyInBrc<T>, Placeholder::In1<T>>;
    using OpAddRes  = Bind<Vec::Add<T>, OpInputX1, OpInputX2>;
    using OpCopyOut = Bind<Vec::CopyOut<T>, Placeholder::Out0<T>, OpAddRes>;
    using Outputs = Elems<OpCopyOut>;
    using OpDag = DAGSch<Outputs, void, MemCfg>;
};
```

Kernel 入口 (`op_kernel/add_apt.cpp`):
```cpp
template <uint64_t schMode>
__global__ __aicore__ void add(GM_ADDR x1, GM_ADDR x2, GM_ADDR y,
                                GM_ADDR workspace, GM_ADDR tiling) {
    using OpDag = AddOp::AddWithoutCastCompute<DTYPE_X1>::OpDag;
    BroadcastSch<schMode, OpDag> sch(tiling);
    sch.Process(x1, x2, y);
}
```

**简单模式 — 手写流水** (`examples/add_example/`):

```cpp
class AddExample {
    TPipe pipe;
    TQue<QuePosition::VECIN, BUFFER_NUM> inputQueueX, inputQueueY;
    TQue<QuePosition::VECOUT, BUFFER_NUM> outputQueueZ;

    void Process() {
        for (int64_t i = 0; i < loopCount; i++) {
            CopyIn(i, currentNum);    // GM → UB
            Compute(currentNum);      // AscendC::Add()
            CopyOut(i, currentNum);   // UB → GM
        }
    }
};
```

### 9.6 Tiling 机制

注册:
```cpp
IMPL_OP_OPTILING(Add).Tiling(TilingForAdd).TilingParse<AddCompileInfo>(TilingPrepareForAdd);
REGISTER_OPS_TILING_TEMPLATE(Add, AddTiling, ADD_COMMON_TILING_PRIORITY);
```

Tiling 基类执行流程:
```
DoTiling() {
    GetShapeAttrsInfo()  → 读输入 shape/attr
    GetPlatformInfo()    → 读 coreNum, UB size
    IsCapable()          → 不支持则 fallthrough 到下一个 tiling class
    DoOpTiling()         → 计算 tiling 数据
    GetWorkspaceSize()   → 计算 workspace
    PostTiling()         → 写入 context
    SetTilingKey()       → 选择 kernel 模板
}
```

### 9.7 API 层

两阶段调用模式:
```
Phase 1: aclnnAddGetWorkspaceSize(x1, x2, y, &wsSize, &executor)  — 构建执行计划
Phase 2: aclnnAdd(workspace, wsSize, executor, stream)              — 执行
```

### 9.8 构建系统

每个算子的 CMakeLists.txt 只有一行宏:
```cmake
set(SUPPORT_COMPUTE_UNIT "ascend950")
set(SUPPORT_TILING_DIR "arch35")
add_all_modules_sources(OPTYPE add ...)
```

`add_all_modules_sources` 自动 glob 各子目录的 .cpp 文件并注册到对应的编译 target。

### 9.9 关键注册宏汇总

| 宏 | 文件位置 | 作用 |
|---|---|---|
| `OP_ADD(OpName)` | `op_host/*_def.cpp` | 注册 OpDef schema |
| `REG_OP(OpName)...OP_END_FACTORY_REG` | `op_graph/*_proto.h` | 注册 GE 图层 proto |
| `IMPL_OP_INFERSHAPE(OpName).InferShape(fn)` | `op_host/*_infershape.cpp` | 注册形状推导 |
| `IMPL_OP(OpName).InferDataType(fn)` | `op_graph/*_graph_infer.cpp` | 注册类型推导 |
| `IMPL_OP_OPTILING(OpName).Tiling(fn)` | `op_host/arch35/*_tiling.cpp` | 注册 tiling |
| `OP_TYPE_REGISTER(OpName)` | `op_api/*.cpp` | 注册 L0 op 类型 |