【昇腾/AscendC开发】直调模式 VS 算子框架模式? Ascend C 开发模式与入口点选择指南

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

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

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

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

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

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

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 核心差异对比

2.2 代码对比

直调模式：

// ===== Kernel 端 (.asc) =====extern"C"__global__ __aicore__voidmy_kernel(GM_ADDR in,GM_ADDR out){// 直接写 kernel 逻辑AscendC::DataCopy(...);AscendC::Add(...);}// ===== Host 端 (.cpp) =====// 声明 kernel 函数（普通 C++ 函数签名）voidmy_kernel(uint32_tblockDim,void*l2ctrl,void*stream,uint8_t*in,uint8_t*out);intmain(){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__voidmy_kernel(GM_ADDR in,GM_ADDR out,GM_ADDR workspace,GM_ADDR tiling){KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);// 告诉框架调度到 AIVGET_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 的性能特征

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_trow=0;row<rowsThisCore;row++){Duplicate(rowSumLocal,(T)0,1);// 清零累加器for(int32_tk=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__= 通用入口，通过宏和运行时调度区分模式
• __vector__= 专用入口，用于直调模式下隔离 Vector Core

4.2 入口点选择规则

4.3 直调模式的关键陷阱

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

实验数据：

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

4.4 算子框架模式的优势

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

extern"C"__global__ __aicore__voidmy_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__voiddequant_kernel(// 注意：用 __vector__GM_ADDR int8_weight,GM_ADDR fp16_weight,GM_ADDR tiling){// Vector 操作：Cast + Muls}// ===== matmul_kernel.asc =====#defineASCENDC_CUBE_ONLYextern"C"__global__ __aicore__voidmatmul_kernel(GM_ADDR x1,GM_ADDR fp16_weight,GM_ADDR out,GM_ADDR tiling){// Cube 操作：Matmul}// ===== host.cpp =====intmain(){// 先执行 Vector Kerneldequant_kernel(1,nullptr,nullptr,d_int8,d_fp16,d_tiling);// 再执行 Cube Kernelmatmul_kernel(1,nullptr,nullptr,d_x1,d_fp16,d_out,d_tiling);aclrtSynchronizeStream(nullptr);}

5.3 算子框架模式实现

// ===== quant_matmul_kernel.cpp =====extern"C"__global__ __aicore__voidquant_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

七、总结

模式选择（第一决策）

Vector/Cube 选择（第二决策）

入口点选择（第三决策）

核心原则

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

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