C/V 融合计算总参考Init 与 Process【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills本文档适用于 C/V 融合算子或设备侧存在多个Scope、多个计算阶段、AIC/AIV 协同的 AscendC 实现。 它不是用来替代vector或cube详细参考而是给出融合场景下的组合阅读顺序与协同组织方式。 概览与判断规则见references/dsl2Ascendc.md。第三章Kernel 入口C/V 融合总览1. 阅读顺序纯 Vector 算子只看references/dsl2Ascendc_compute_vector.md纯 Cube 算子只看references/dsl2Ascendc_compute_cube.mdC/V 融合算子先看本文再分别看references/dsl2Ascendc_compute_cube.md和references/dsl2Ascendc_compute_vector.md2. kernel 入口形态C/V 融合算子的入口通常同时接收输入、输出、workspace 和 tilingextern C __global__ __aicore__ void kernel_custom(GM_ADDR ...inputs..., GM_ADDR workspace, GM_ADDR tiling) { KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_1); AscendC::TPipe pipe; KernelClass kernel; kernel.Init(..., workspace, tiling, pipe); kernel.Process(); }3.vec_num与 block 组成DSLvec_numKERNEL_TYPE每个 block 组成GetSubBlockNum()1KERNEL_TYPE_MIX_AIC_1_11 AIC 1 AIV22KERNEL_TYPE_MIX_AIC_1_21 AIC 2 AIV3第四章主 Kernel 类C/V 融合参考archive_tasks/matmul_leakyrelu/kernel/matmul_leakyrelu.hC/V 融合主Kernel类建议按Init()和Process()两个大阶段组织。1.Init()接收 tiling 字段并初始化 GM / workspace / 子模块Init()主要负责读取并保存 tiling 字段绑定输入 / 输出 GM tensor初始化调度器与 workspace分别初始化 Cube 子模块和 Vector 子模块A. tiling 字段、GM 绑定与调度常见模式CopyTiling(tiling_, tilingGM)SetGlobalBuffer(...)绑定 A/B/C 等 GM tensor根据GetBlockIdx()、GetSubBlockNum()派生coreIdx初始化调度器如sched_.Init(...)B. workspace 与跨核协同如果 C/V 之间通过 workspace 传递中间结果通常在Init()中完成workspace 基址和每个 core 的偏移计算ring buffer /WorkspaceQueue初始化C/V 协同所需 flag 或队列的初始化若存在跨核同步或 producer / consumer 关系继续结合references/dsl2Ascendc_cross_core_sync.md。C. 子模块初始化融合场景下通常同时存在Cube 子模块如matmul.hVector 子模块如leakyrelu.h、scale.h推荐在Init()中按分支初始化ASCEND_IS_AIC分支初始化 Cube 子模块ASCEND_IS_AIV分支初始化 Vector 子模块2.Process()组织调度、AIC/AIV 分支与阶段调用Process()负责把工作负载循环、AIC/AIV 分支和模块调用串起来。A. 工作负载循环常见骨架__aicore__ inline void KernelClass::Process() { int mIdx, nIdx; while (sched_.HasNext()) { sched_.Next(mIdx, nIdx); if ASCEND_IS_AIC { // Cube 侧 } if ASCEND_IS_AIV { // Vector 侧 } } }B. AIC 分支AIC 分支通常负责从 GM 取当前 tile 的输入获取 workspace 生产者槽位调用 Cube 子模块如mm_.ComputeBlock(...)释放生产者槽位或发送完成信号C. AIV 分支AIV 分支通常负责获取 workspace 消费者槽位根据GetSubBlockIdx()计算当前子块偏移从 workspace 中取本子块负责的数据调用 Vector 子模块完成后处理并写回 GM释放消费者槽位或发送完成信号D. 何时拆单独子模块当满足以下任一条件时建议拆出单独计算子模块文件TileLang 设备侧有多个职责清晰的Scope同时存在 Cube 计算阶段和 Vector 后处理阶段需要在主Kernel类中复用某段计算逻辑建议让 TileLang 中一个主要Scope对应 AscendC 中一个子模块。【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
CANN/cannbot-skills C/V融合计算参考
发布时间：2026/6/18 8:05:39
C/V 融合计算总参考Init 与 Process【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills本文档适用于 C/V 融合算子或设备侧存在多个Scope、多个计算阶段、AIC/AIV 协同的 AscendC 实现。 它不是用来替代vector或cube详细参考而是给出融合场景下的组合阅读顺序与协同组织方式。 概览与判断规则见references/dsl2Ascendc.md。第三章Kernel 入口C/V 融合总览1. 阅读顺序纯 Vector 算子只看references/dsl2Ascendc_compute_vector.md纯 Cube 算子只看references/dsl2Ascendc_compute_cube.mdC/V 融合算子先看本文再分别看references/dsl2Ascendc_compute_cube.md和references/dsl2Ascendc_compute_vector.md2. kernel 入口形态C/V 融合算子的入口通常同时接收输入、输出、workspace 和 tilingextern C __global__ __aicore__ void kernel_custom(GM_ADDR ...inputs..., GM_ADDR workspace, GM_ADDR tiling) { KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_1); AscendC::TPipe pipe; KernelClass kernel; kernel.Init(..., workspace, tiling, pipe); kernel.Process(); }3.vec_num与 block 组成DSLvec_numKERNEL_TYPE每个 block 组成GetSubBlockNum()1KERNEL_TYPE_MIX_AIC_1_11 AIC 1 AIV22KERNEL_TYPE_MIX_AIC_1_21 AIC 2 AIV3第四章主 Kernel 类C/V 融合参考archive_tasks/matmul_leakyrelu/kernel/matmul_leakyrelu.hC/V 融合主Kernel类建议按Init()和Process()两个大阶段组织。1.Init()接收 tiling 字段并初始化 GM / workspace / 子模块Init()主要负责读取并保存 tiling 字段绑定输入 / 输出 GM tensor初始化调度器与 workspace分别初始化 Cube 子模块和 Vector 子模块A. tiling 字段、GM 绑定与调度常见模式CopyTiling(tiling_, tilingGM)SetGlobalBuffer(...)绑定 A/B/C 等 GM tensor根据GetBlockIdx()、GetSubBlockNum()派生coreIdx初始化调度器如sched_.Init(...)B. workspace 与跨核协同如果 C/V 之间通过 workspace 传递中间结果通常在Init()中完成workspace 基址和每个 core 的偏移计算ring buffer /WorkspaceQueue初始化C/V 协同所需 flag 或队列的初始化若存在跨核同步或 producer / consumer 关系继续结合references/dsl2Ascendc_cross_core_sync.md。C. 子模块初始化融合场景下通常同时存在Cube 子模块如matmul.hVector 子模块如leakyrelu.h、scale.h推荐在Init()中按分支初始化ASCEND_IS_AIC分支初始化 Cube 子模块ASCEND_IS_AIV分支初始化 Vector 子模块2.Process()组织调度、AIC/AIV 分支与阶段调用Process()负责把工作负载循环、AIC/AIV 分支和模块调用串起来。A. 工作负载循环常见骨架__aicore__ inline void KernelClass::Process() { int mIdx, nIdx; while (sched_.HasNext()) { sched_.Next(mIdx, nIdx); if ASCEND_IS_AIC { // Cube 侧 } if ASCEND_IS_AIV { // Vector 侧 } } }B. AIC 分支AIC 分支通常负责从 GM 取当前 tile 的输入获取 workspace 生产者槽位调用 Cube 子模块如mm_.ComputeBlock(...)释放生产者槽位或发送完成信号C. AIV 分支AIV 分支通常负责获取 workspace 消费者槽位根据GetSubBlockIdx()计算当前子块偏移从 workspace 中取本子块负责的数据调用 Vector 子模块完成后处理并写回 GM释放消费者槽位或发送完成信号D. 何时拆单独子模块当满足以下任一条件时建议拆出单独计算子模块文件TileLang 设备侧有多个职责清晰的Scope同时存在 Cube 计算阶段和 Vector 后处理阶段需要在主Kernel类中复用某段计算逻辑建议让 TileLang 中一个主要Scope对应 AscendC 中一个子模块。【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
