【昇腾/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 Cube1.3 场景三独立算子 / 性能关键路径典型需求一个独立的算子不需要与其他算子融合性能极其关键需要精细控制不依赖图模式推荐直调模式原因可以完全控制 kernel launch 参数减少框架开销可以手动优化 tiling 策略注意这种场景较少见大多数生产环境还是需要框架模式。1.4 场景四需要 Cube Vector 并行典型需求算子需要同时使用 Cube矩阵乘和 Vector后处理希望两者并行执行以提高性能推荐算子框架模式MIX 模式原因直调模式不支持 MIX 模式会 hang框架模式的 KFCKernel Flow Control可以自动调度 AIC 和 AIV1.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) externC__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) externC__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/*.cppkernel 实现// - op_host/*.cpptiling 策略 aclnn API// - op_proto/*.cpp算子原型定义// 用户调用方式两阶段 APIsize_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) ││ │ │ • 典型 APIMatmul, 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 利用率高❌ 效率低CubeGEMV (N1)❌ MTE2 96%, Cube 1%✅ ReduceSum 高效Vector逐元素运算❌ 不适合✅ SIMD 高效Vector归约操作❌ 不适合✅ ReduceSum/ReduceMinVector量化 MatMul✅ Cube Matmul—双 Kernel3.4 GEMV 的典型案例问题GEMV (mat[M,K] vec[K], N1) 用 Cube Matmul 性能极差原因MTE2 占比 96-99%几乎全部时间在等数据Cube MAC ratio 0.5%计算单元几乎空闲GM→L1 带宽利用率仅 0.21-0.48%Vector 方案逐行 MulAdd ReduceSum// Vector kernel逐行点积for(int32_trow0;rowrowsThisCore;row){Duplicate(rowSumLocal,(T)0,1);// 清零累加器for(int32_tk0;ktotalK;kTILE_K){DataCopy(matLocal,matGm[row*Kk],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 空转 │ ├─ 逐元素运算 │ └─ VectorCast, 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 Core4.2 入口点选择规则模式Kernel 类型入口点写法直调纯 Vector__vector__直调纯 Cube/Matmul__aicore__ASCENDC_CUBE_ONLY直调混合双 KernelVector 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✅ PASS128×256×128✅ PASS❌ FAIL256×512×256✅ PASS❌ FAIL512×1024×512✅ PASS✅ PASS结论直调模式的纯 Vector Kernel必须使用__vector__入口。4.4 算子框架模式的优势算子框架模式下所有 kernel 都使用__aicore__入口通过宏告诉框架调度externC__global__ __aicore__voidmy_kernel(...){// 框架根据这个宏调度到正确的硬件单元KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);// ...}优势不存在干扰后续 Kernel的问题KFC 框架正确管理资源调度支持 MIX 模式AICAIV 并行五、实战案例量化 MatMul5.1 场景描述实现量化矩阵乘out dequant(INT8_weight) FP16_x需要Vector KernelINT8 → FP16 反量化Cube KernelFP16 矩阵乘5.2 直调模式实现// dequant_kernel.asc externC__global__ __vector__voiddequant_kernel(// 注意用 __vector__GM_ADDR int8_weight,GM_ADDR fp16_weight,GM_ADDR tiling){// Vector 操作Cast Muls}// matmul_kernel.asc #defineASCENDC_CUBE_ONLYexternC__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 externC__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_coreTypeAIV){// Vector 侧反量化}else{// Cube 侧Matmul}}对比直调模式需要两个独立 kernel顺序执行框架模式一个 kernelMIX 模式并行执行六、常见问题Q1__cube__修饰符存在吗不存在。Cube-only 模式通过__aicore__ASCENDC_CUBE_ONLY宏实现。Q2GEMV (N1) 应该用 Cube 还是 VectorVector。GEMV 用 Cube 时MTE2 占比 96%Cube 利用率 1%。用 Vector 的 ReduceSum 效率高得多。Q3生产部署必须用框架模式吗推荐用框架模式。原因与 PyTorch 等框架集成支持图模式和算子融合自动 tiling 和内存管理社区支持和文档完善Q4直调模式什么时候用研究原型验证性能基准测试独立小工具学习 Ascend C七、总结模式选择第一决策场景推荐模式适配算法库PyTorch/vLLM算子框架研究原型 / 性能验证直调需要图模式 / 算子融合算子框架需要 MIX 并行算子框架直调不支持Vector/Cube 选择第二决策场景推荐GEMM (大 N)CubeGEMV (N1)Vector逐元素运算Vector归约操作Vector入口点选择第三决策模式Vector KernelCube Kernel直调__vector____aicore__ASCENDC_CUBE_ONLY框架__aicore__KERNEL_TYPE_AIV_ONLY__aicore__KERNEL_TYPE_AIC_ONLY核心原则先定模式再定入口点生产部署用框架研究原型用直调直调模式下纯 Vector Kernel 必须用__vector__N1 用 VectorN 大用 Cube本文基于 CANN 8.5.0 和 Atlas A2 (dav-2201) 验证不同硬件和CANN版本结论可能存在差异。