Ascend C 开发模式与入口点选择指南

开篇：你该选哪种开发模式？

如果你正在开始一个 Ascend C 算子项目，第一个问题不是"用什么 API"，而是**“我该选哪种开发模式”**：

• 直调模式：像写普通 C++ 函数一样，直接调用 kernel
• 算子框架模式：接入 CANN 算子生态，通过 aclnnXxx API 调用

选错了模式，后续的入口点选择、性能优化、部署方式都会走弯路。本文将从实际应用场景出发，帮你做出正确选择。

---

一、应用场景分析：你该选哪种模式？

1.1 场景一：适配现有算法库（如 PyTorch、vLLM）

典型需求：

• 将自定义算子接入 PyTorch / TensorFlow / vLLM 等框架
• 需要通过 torch.ops 或类似机制调用
• 需要支持图模式、自动微分等特性

推荐：算子框架模式

现有算法库
    ↓ 调用
CANN 算子库（.so）
    ↓ 内部
Ascend C Kernel + Tiling + Runtime

原因：

• CANN 算子生态与 PyTorch 等框架深度集成
• 自动支持图模式、算子融合、内存复用
• 可以被 vLLM、MindSpore 等上层框架直接调用
• tiling 策略由框架自动生成，减少手动调优

实际案例：

• ops-nn 中的所有算子（foreach、quant、matmul 等）都是算子框架模式
• vLLM-Ascend 的自定义算子也采用框架模式

1.2 场景二：研究原型 / 性能验证

典型需求：

• 快速验证一个新算法的可行性
• 测试某个 kernel 的性能上限
• 不需要部署到生产环境

推荐：直调模式

原因：

• 开发周期短，可以快速迭代
• 不需要处理复杂的 tiling 和算子注册
• 可以直接在可执行文件中测试，调试方便
• 适合论文实验、性能分析

实际案例：

• 性能对比实验（如 GEMV Vector vs Cube）

1.3 场景三：独立算子 / 性能关键路径

典型需求：

• 一个独立的算子，不需要与其他算子融合
• 性能极其关键，需要精细控制
• 不依赖图模式

推荐：直调模式

原因：

• 可以完全控制 kernel launch 参数
• 减少框架开销
• 可以手动优化 tiling 策略

注意：这种场景较少见，大多数生产环境还是需要框架模式。

1.4 场景四：需要 Cube + Vector 并行

典型需求：

• 算子需要同时使用 Cube（矩阵乘）和 Vector（后处理）
• 希望两者并行执行以提高性能

推荐：算子框架模式（MIX 模式）

原因：

• 直调模式不支持 MIX 模式（会 hang）
• 框架模式的 KFC（Kernel Flow Control）可以自动调度 AIC 和 AIV

1.5 选择决策树

你的需求是什么？
│
├─ 适配现有算法库（PyTorch/vLLM/...）
│   └─ ✅ 算子框架模式
│
├─ 研究原型 / 性能验证
│   └─ ✅ 直调模式
│
├─ 需要图模式 / 算子融合
│   └─ ✅ 算子框架模式
│
├─ 需要 Cube + Vector 并行（MIX）
│   └─ ✅ 算子框架模式（直调不支持）
│
└─ 独立算子 / 不依赖框架
    └─ ⚠️ 直调模式（少数场景）

---

二、两种模式的核心差异

2.1 核心差异对比

特性	直调模式	算子框架模式
代码量	少（kernel + host）	多（kernel + tiling + proto）
编译产物	单个可执行文件 .out	算子库 .so
调用方式	kernel<<<>>>(args)	aclnnXxx(args)
Tiling	手动管理	框架自动生成
Workspace	手动管理	框架自动计算
KFC 框架	❌ 不可用	✅ 可用
MIX 模式	❌ 不支持	✅ 支持

2.2 代码对比

直调模式：

// ===== Kernel 端 (.asc) =====
extern "C" __global__ __aicore__ void my_kernel(GM_ADDR in, GM_ADDR out)
{
    // 直接写 kernel 逻辑
    AscendC::DataCopy(...);
    AscendC::Add(...);
}

// ===== Host 端 (.cpp) =====
// 声明 kernel 函数（普通 C++ 函数签名）
void my_kernel(uint32_t blockDim, void* l2ctrl, void* stream,
               uint8_t* in, uint8_t* out);

int main() {
    aclInit(nullptr);
    aclrtSetDevice(0);
    
    // 分配内存
    void *d_in, *d_out;
    aclrtMalloc(&d_in, size, ...);
    aclrtMalloc(&d_out, size, ...);
    
    // 直接调用 kernel！就像调用普通函数
    my_kernel(1, nullptr, nullptr, (uint8_t*)d_in, (uint8_t*)d_out);
    
    aclrtSynchronizeStream(nullptr);
    aclFinalize();
}

算子框架模式：

// ===== Kernel 端 (.cpp) =====
extern "C" __global__ __aicore__ void my_kernel(
    GM_ADDR in, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling)
{
    KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);  // 告诉框架调度到 AIV
    GET_TILING_DATA(tilingData, tiling);
    // ... kernel 逻辑
}

// ===== Host 端 =====
// 需要实现完整的算子注册流程（通常由 msopgen 工具生成）：
// - op_kernel/*.cpp（kernel 实现）
// - op_host/*.cpp（tiling 策略 + aclnn API）
// - op_proto/*.cpp（算子原型定义）

// 用户调用方式（两阶段 API）：
size_t workspaceSize;
aclnnMyOpGetWorkspaceSize(..., &workspaceSize);
aclrtMalloc(&workspace, workspaceSize, ...);
aclnnMyOp(workspace, stream, ...);

---

三、NPU 硬件架构与 Vector/Cube 选择

3.1 AI Core 的内部结构

在讨论入口点之前，必须先理解 NPU 的硬件架构。

┌─────────────────────────────────────────────────────────┐
│                    AI Core (AIC)                        │
│  ┌─────────────────────────────────────────────────────┐│
│  │ Cube Unit (矩阵计算单元)                             ││
│  │ • MAC 阵列：高吞吐矩阵乘法                           ││
│  │ • 最优场景：M, N, K 都较大 (如 1024×1024×1024)       ││
│  │ • 典型 API：Matmul, Mmad                             ││
│  └─────────────────────────────────────────────────────┘│
│  ┌─────────────────────────────────────────────────────┐│
│  │ Vector Unit (向量计算单元)                           ││
│  │ • SIMD：逐元素运算 (Add, Mul, Cast...)               ││
│  │ • Reduce：归约操作 (ReduceSum, ReduceMax...)         ││
│  │ • DMA：数据搬运 (DataCopy, DataCopyPad)              ││
│  └─────────────────────────────────────────────────────┘│
│  ┌─────────────────────────────────────────────────────┐│
│  │ Storage (存储层次)                                   ││
│  │ • UB (Unified Buffer): Vector 的工作空间             ││
│  │ • L1: Cube 的工作空间                                ││
│  │ • L2: 片上共享缓存                                   ││
│  └─────────────────────────────────────────────────────┘│
└─────────────────────────────────────────────────────────┘

3.2 分离架构（Atlas A2）

在 Atlas A2 (dav-2201) 上，架构进一步分离：

┌─────────────────────────────┐
│     AI Core (AIC)           │  ← Cube + Vector（但分离调度）
└─────────────────────────────┘
              ↓ 独立调度
┌─────────────────────────────┐
│    Vector Core (AIV)        │  ← 独立的 Vector Unit + UB
│    数量：AIC:AIV = 1:2      │
└─────────────────────────────┘

关键点：在分离架构下，AIC 和 AIV 可以并行执行，但也带来了协调问题。

3.3 Vector vs Cube 的性能特征

场景	Cube 方案	Vector 方案	推荐
GEMM (大 N)	✅ Cube 利用率高	❌ 效率低	Cube
GEMV (N=1)	❌ MTE2 96%, Cube < 1%	✅ ReduceSum 高效	Vector
逐元素运算	❌ 不适合	✅ SIMD 高效	Vector
归约操作	❌ 不适合	✅ ReduceSum/ReduceMin	Vector
量化 MatMul	✅ Cube Matmul	—	双 Kernel

3.4 GEMV 的典型案例

问题：GEMV (mat[M,K] @ vec[K], N=1) 用 Cube Matmul 性能极差

原因：

• MTE2 占比 96-99%（几乎全部时间在等数据）
• Cube MAC ratio < 0.5%（计算单元几乎空闲）
• GM→L1 带宽利用率仅 0.21-0.48%

Vector 方案：逐行 MulAdd + ReduceSum

// Vector kernel：逐行点积
for (int32_t row = 0; row < rowsThisCore; row++) {
    Duplicate(rowSumLocal, (T)0, 1);  // 清零累加器
    
    for (int32_t k = 0; k < totalK; k += TILE_K) {
        DataCopy(matLocal, matGm[row * K + k], tileK);
        DataCopy(vecLocal, vecGm[k], tileK);
        Mul(tmpLocal, matLocal, vecLocal, tileK);
        ReduceSum(rowSumLocal, tmpLocal, rowSumLocal, tileK);
    }
    DataCopy(outGm[row], rowSumLocal, 1);
}

3.5 Vector/Cube 选择决策

你的算子需要什么计算？
│
├─ 矩阵乘法 (GEMM)
│   │
│   ├─ N 较大 (N > 16)?
│   │   └─ Cube Matmul（高吞吐）
│   │
│   └─ N = 1 (GEMV)?
│       └─ Vector MulAdd + ReduceSum（避免 Cube 空转）
│
├─ 逐元素运算
│   └─ Vector（Cast, Add, Mul, Gelu...）
│
├─ 归约
│   └─ Vector（单核即可，避免多核开销）
│
└─ 混合计算
    │
    ├─ 算子框架模式?
    │   └─ MIX 模式（框架调度）
    │
    └─ 直调模式?
        └─ 双 Kernel：先 Vector，后 Cube

---

四、入口点选择：基于模式决定

确定了开发模式后，才需要考虑入口点选择。

4.1 入口点修饰符设计

修饰符	含义	硬件单元	使用场景
__aicore__	AI Core 入口	AIC (Cube + Vector)	Cube/Matmul Kernel、算子框架模式
__vector__	Vector Core 入口	AIV (纯 Vector)	纯 Vector Kernel（直调模式）
❌ __cube__	不存在	-	Cube 逻辑通过 __aicore__ + ASCENDC_CUBE_ONLY 实现

设计理念：

• __aicore__ = 通用入口，通过宏和运行时调度区分模式
• __vector__ = 专用入口，用于直调模式下隔离 Vector Core

4.2 入口点选择规则

模式	Kernel 类型	入口点写法
直调	纯 Vector	__vector__
直调	纯 Cube/Matmul	__aicore__ + ASCENDC_CUBE_ONLY
直调	混合	双 Kernel（Vector + Cube 分离）
框架	纯 Vector	__aicore__ + KERNEL_TYPE_AIV_ONLY
框架	纯 Cube	__aicore__ + KERNEL_TYPE_AIC_ONLY
框架	混合	__aicore__ + MIX 模式

4.3 直调模式的关键陷阱

问题场景：直调模式下，Vector Kernel 使用 __aicore__ 入口，会干扰后续 Cube Matmul。

实验数据：

Shape (M×K×N)	__vector__	__aicore__
16×16×16	✅ PASS	✅ PASS
128×256×128	✅ PASS	❌ FAIL
256×512×256	✅ PASS	❌ FAIL
512×1024×512	✅ PASS	✅ PASS

结论：直调模式的纯 Vector Kernel 必须使用 __vector__ 入口。

4.4 算子框架模式的优势

算子框架模式下，所有 kernel 都使用 __aicore__ 入口，通过宏告诉框架调度：

extern "C" __global__ __aicore__ void my_kernel(...)
{
    // 框架根据这个宏调度到正确的硬件单元
    KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
    // ...
}

优势：

• 不存在"干扰后续 Kernel"的问题
• KFC 框架正确管理资源调度
• 支持 MIX 模式（AIC+AIV 并行）

---

五、实战案例：量化 MatMul

5.1 场景描述

实现量化矩阵乘：out = dequant(INT8_weight) @ FP16_x

需要：

1. Vector Kernel：INT8 → FP16 反量化
2. Cube Kernel：FP16 矩阵乘

5.2 直调模式实现

// ===== dequant_kernel.asc =====
extern "C" __global__ __vector__ void dequant_kernel(  // 注意：用 __vector__
    GM_ADDR int8_weight, GM_ADDR fp16_weight, GM_ADDR tiling)
{
    // Vector 操作：Cast + Muls
}

// ===== matmul_kernel.asc =====
#define ASCENDC_CUBE_ONLY
extern "C" __global__ __aicore__ void matmul_kernel(
    GM_ADDR x1, GM_ADDR fp16_weight, GM_ADDR out, GM_ADDR tiling)
{
    // Cube 操作：Matmul
}

// ===== host.cpp =====
int main() {
    // 先执行 Vector Kernel
    dequant_kernel(1, nullptr, nullptr, d_int8, d_fp16, d_tiling);
    
    // 再执行 Cube Kernel
    matmul_kernel(1, nullptr, nullptr, d_x1, d_fp16, d_out, d_tiling);
    
    aclrtSynchronizeStream(nullptr);
}

5.3 算子框架模式实现

// ===== quant_matmul_kernel.cpp =====
extern "C" __global__ __aicore__ void quant_matmul_kernel(
    GM_ADDR x1, GM_ADDR int8_weight, GM_ADDR out, 
    GM_ADDR workspace, GM_ADDR tiling)
{
    // 使用 MIX 模式：AIC 和 AIV 并行
    if (g_coreType == AIV) {
        // Vector 侧：反量化
    } else {
        // Cube 侧：Matmul
    }
}

对比：

• 直调模式：需要两个独立 kernel，顺序执行
• 框架模式：一个 kernel，MIX 模式并行执行

---

六、常见问题

Q1：__cube__ 修饰符存在吗？

不存在。Cube-only 模式通过 __aicore__ + ASCENDC_CUBE_ONLY 宏实现。

Q2：GEMV (N=1) 应该用 Cube 还是 Vector？

Vector。GEMV 用 Cube 时，MTE2 占比 96%，Cube 利用率 < 1%。用 Vector 的 ReduceSum 效率高得多。

Q3：生产部署必须用框架模式吗？

推荐用框架模式。原因：

• 与 PyTorch 等框架集成
• 支持图模式和算子融合
• 自动 tiling 和内存管理
• 社区支持和文档完善

Q4：直调模式什么时候用？

• 研究原型验证
• 性能基准测试
• 独立小工具
• 学习 Ascend C

---

七、总结

模式选择（第一决策）

场景	推荐模式
适配算法库（PyTorch/vLLM）	算子框架
研究原型 / 性能验证	直调
需要图模式 / 算子融合	算子框架
需要 MIX 并行	算子框架（直调不支持）

Vector/Cube 选择（第二决策）

场景	推荐
GEMM (大 N)	Cube
GEMV (N=1)	Vector
逐元素运算	Vector
归约操作	Vector

入口点选择（第三决策）

模式	Vector Kernel	Cube Kernel
直调	__vector__	__aicore__ + ASCENDC_CUBE_ONLY
框架	__aicore__ + KERNEL_TYPE_AIV_ONLY	__aicore__ + KERNEL_TYPE_AIC_ONLY

核心原则

1. 先定模式，再定入口点
2. 生产部署用框架，研究原型用直调
3. 直调模式下纯 Vector Kernel 必须用 __vector__
4. N=1 用 Vector，N 大用 Cube

---

本文基于 CANN 8.5.0 和 Atlas A2 (dav-2201) 验证，不同硬件和CANN版本结论可能存在差异。