KERNEL_TASK_TYPE_DEFAULT 昇腾CANN算子开发

zhangfeng1133

118人浏览 · 2026-08-30 17:19:33

zhangfeng1133 · 2026-08-30 17:19:33 发布

KERNEL_TASK_TYPE_DEFAULT 昇腾CANN算子开发

宏定义：KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE)，属于Ascend‑C 自定义算子（Kernel）任务类型宏，用于注册算子内核，告诉CANN调度器这个算子是什么类型、怎么执行。

1、它具体做什么

Ascend‑C算子开发，算子内核代码写完成后，必须用这个宏向框架注册算子任务属性，生成Kernel描述信息，给NPU调度器：

标记算子属于哪一类任务（AI Core / Vector Core / CPU）
声明算子内核函数入口
设置任务默认属性：栈大小、优先级别、阻塞模式等默认参数
编译阶段生成算子元数据，供GE/ATC编译时识别调用

KERNEL_TASK_TYPE_DEFAULT = 使用框架默认任务配置，绝大多数普通自定义算子优先选这个。

宏原型简化

KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE)

KERNEL_TYPE：你自己定义的算子任务类型名（宏，自定义标识符，比如MyAddKernel）

等价含义：使用CANN内置默认task配置，不手动改写task的栈、优先级、阻塞策略。

2、什么时候选 KERNEL_TASK_TYPE_DEFAULT ✅（90%场景选它）

适用场景：

AI Core / Vector Core 上执行的普通Ascend‑C算子
不需要特殊调参：不需要自定义栈大小、不需要高优先级、不需要非阻塞特殊配置
单算子、普通计算逻辑，没有超大局部栈占用，不需要特殊调度行为

示例完整注册片段（标准写法）

// 1.定义任务类型
KERNEL_TASK_TYPE_DEFAULT(MyAddKernel);

// 2.注册kernel入口，绑定上面的任务类型
KERNEL_REGISTER(MyAddKernel)
{
    // 算子内核逻辑，AI Core上运行
}

3、不选DEFAULT的情况（什么时候不用这个宏）

如果你的算子有下面需求，不能用DEFAULT，要用KERNEL_TASK_TYPE手动配置：

// 手动自定义task属性，不用DEFAULT
KERNEL_TASK_TYPE(MySpecialKernel)
    .stack_size = 81920,    // 自定义栈大小，局部变量特别大时
    .task_priority = 2,     // 自定义任务优先级
    .block_type = TASK_NON_BLOCKING // 非阻塞任务

触发条件：

算子内部局部数组/变量巨大，默认栈溢出，需要调大stack_size
需要修改任务优先级
需要非阻塞任务模式
CPU侧Kernel（Host CPU算子）部分场景需要定制参数

⚠️绝大多数AI Core算子直接用KERNEL_TASK_TYPE_DEFAULT，不要手动乱改栈，栈设置过大会浪费片上内存。

4、DEFAULT到底默认了哪些参数？

CANN内置默认值（CANN 8.x/9.x）

stack_size：默认 32768（32KB）AI Core栈
task_priority：默认优先级0
block_type：阻塞任务 TASK_BLOCKING
其余属性使用框架默认

5、容易踩坑点

KERNEL_TYPE只是标识符，不能和其他算子重名，重名编译不会报错，运行时算子错乱。
KERNEL_TASK_TYPE_DEFAULT必须写在全局作用域，不能写在函数内部。
如果算子运行报栈溢出 stack overflow，就不能继续用DEFAULT，改用KERNEL_TASK_TYPE调大stack_size。
不要混淆：
- KERNEL_TASK_TYPE_DEFAULT：定义任务类型
- KERNEL_REGISTER(xxx)：注册内核函数入口，参数必须和前面任务类型名字完全一致。

最简完整示例模板

#include "kernel_operator.h"

// 使用默认任务类型，绝大多数算子直接这一行
KERNEL_TASK_TYPE_DEFAULT(KernelAddDemo);

// 算子内核实现
KERNEL_REGISTER(KernelAddDemo)
{
    // AI Core计算逻辑
}

如果你手上有报错日志（栈溢出、算子调度失败），可以贴出来，我帮你判断是继续用DEFAULT还是需要手动改task参数。

Error 95 根因分析与 KERNEL_TYPE 完整配置方案

核心结论：你的 error 95 最可能的根因是“纯 Cube 逻辑的算子声明了 `KERNEL_TYPE_MIX_AIC_1_2`，违反了 CANN 的 kernel type 与算子实现必须匹配的约束”。你有两条路可走：保持 `KERNEL_TYPE_AIC_ONLY`（推荐，零改动、性能无损），或按真正的 MIX 架构改造 kernel（需要动 kernel 分支逻辑 + workspace + blockDim 三处）。下面给出完整的技术原理、排查流程和两套可直接使用的配置模板。

一、Error 95 为什么会发生

1.1 已验证的约束条件（有官方文档依据）

CANN 对 KERNEL_TASK_TYPE_DEFAULT 的使用有几条硬性约束，任何一条违反都可能导致编译或运行报错：
约束 1：算子实现必须与 kernel type 匹配。 官方约束原文：“当设置具体的 kernel task type 时，用户的算子实现需要与 kernel type 相匹配。比如用户设置 kernel type 为 KERNEL_TYPE_MIX_AIC_1_2，则算子内部实现应与核配比 AIC:AIV 为 1:2 相对应；若用户设置 kernel type 为 KERNEL_TYPE_AIC_ONLY，则算子内部实现应该为纯 cube 逻辑，不应该存在 vector 部分的逻辑。”
约束 2：纯 Cube/Vector 算子强制设 MIX 时 workspace 不能为 0。 官方约束原文：“当纯 cube 或者纯 vec 算子强制设定 kernel type 为 MIX 类型时，workspace 的大小不能设置为 0，需要设置一个大于 0 的值（比如 16、32 等）。”
约束 3：blockDim 不能超过物理核数。 官方约束原文：“针对 Vector/Cube 融合计算的算子，启动时按照 AIV 和 AIC 组合启动，blockDim 用于设置启动多少个组合执行……注意：该场景下，设置的 blockDim 逻辑核的核数不能超过物理核（2 个 Vector 核和 1 个 Cube 核组合为 1 个物理核）的核数。”
约束 4：多核函数须逐个显式设置类型。 如果同一编译单元有多个 __global__ __aicore__ 核函数，必须为每一个都设置 Kernel 类型。

1.2 你的场景为什么踩中这些约束

对照你的情况逐条检查：

约束	你的情况	是否命中
类型-实现匹配	kernel 只有 Mmad 等 Cube 逻辑，声明了 MIX_AIC_1_2（要求 1:2 的 AIC:AIV 分支逻辑）	❌ 大概率命中
workspace 非零	纯 Cube 算子的 tiling 通常把 workspace 算成 0，MIX 强制要求 > 0	❌ 大概率命中
blockDim 语义	MIX_AIC_1_2 的 blockDim 是“组合数”（每组 = 1 Cube + 2 Vector），如果按 Cube 习惯设置可能超限	⚠️ 需检查
最可能的因果链：你的 kernel 里只有 Cube 计算（没有 `ASCEND_IS_AIC` / `ASCEND_IS_AIV` 分支），声明 MIX_AIC_1_2 后，框架会尝试启动 Vector 核去执行 Cube 指令流（或反之），触发校验失败，报出 error 95。

1.3 关于 error 95 本身的诚实说明

需要向你明确一点：“error 95” 这个具体错误码的官方定义，我在 CANN 文档中没有找到直接对应条目。 CANN 的错误码体系是 6 位字符格式（如 EZ9999、E50000-E89999 属于 AICORE 模块），单独一个数字"95"可能是编译日志里的行号、内部校验编号，或被截断的错误码尾部。
因此上述根因分析是基于“改一行 Kernel type 就消除报错”这一现象 + 官方约束条件做的高置信度推断，而非从 error 95 的官方释义直接推导。如果你能贴出完整报错日志（含 error 95 前后各 10 行），可以进一步精确定位。

二、KERNEL_TYPE_* 枚举值完整对照表

枚举值	启动的核（blockDim=10 时）	适用算子特征	950PR 支持
`KERNEL_TYPE_AIC_ONLY`	10 个 Cube 核	纯 Cube 逻辑（Mmad/Matmul）	✅
`KERNEL_TYPE_AIV_ONLY`	10 个 Vector 核	纯 Vector 逻辑（elementwise/softmax）	✅
`KERNEL_TYPE_MIX_AIC_1_0`	10 个 Cube 核（带硬同步）	纯 Cube + 多核控制指令	✅
`KERNEL_TYPE_MIX_AIV_1_0`	10 个 Vector 核（带硬同步）	纯 Vector + 多核控制指令	✅
`KERNEL_TYPE_MIX_AIC_1_1`	10 Cube + 10 Vector	AIC:AIV = 1:1 混合	✅
`KERNEL_TYPE_MIX_AIC_1_2`	10 Cube + 20 Vector	AIC:AIV = 1:2 混合	✅
`KERNEL_TYPE_AICORE`	10 个 AI Core（不区分）	耦合架构通用	❌
`KERNEL_TYPE_VECTORCORE`	—	预留参数，暂不支持	❌
`KERNEL_TYPE_MIX_AICORE`	—	预留参数，暂不支持	❌
数据来源：CANN 官方文档的产品支持情况表和枚举值说明。
一个关键的硬件事实：Atlas 950PR 配备最多 32 个 Cube 核 + 64 个 Vector 核，物理配比恰好是 1:2。这就是 `MIX_AIC_1_2` 存在的意义——它正是为这类芯片的物理核配比设计的。所以 950PR 上用 MIX_AIC_1_2 本身是合理的，前提是你的算子真的需要两类核并行。

三、Error 95 排查流程（如改配置后仍报错）

按以下顺序逐项检查，每步都给出具体操作：

第 1 步：打开详细日志，拿到完整报错上下文

export ASCEND_GLOBAL_LOG_LEVEL=0          # debug 级别，信息最全
export ASCEND_SLOG_PRINT_TO_STDOUT=1      # 日志直接打屏

重新运行后，在日志中搜索 error 95 前后的完整报错串。CANN 报错通常带有 EZ/E5/EI 段前缀和描述文字，错误码前后的描述文字比"95"这个数字本身信息量大得多。

第 2 步：核对 kernel 内部是否有 Vector API 调用

在 kernel 源码中搜索以下 Vector API：

AscendC::Add、AscendC::Mul（双目 Vector 运算）
AscendC::Exp、AscendC::Softmax（激活类）
AscendC::ReduceMax、AscendC::ReduceSum（规约）
如果声明了 AIC_ONLY 却有这些调用 → 类型不匹配，同样会报错。

第 3 步：核对 workspace 设置

检查 Host 侧 tiling 函数中 workspace 的计算逻辑。MIX 类型要求非零 workspace：

// Host 侧 tiling 函数中
size_t usrSize = 256;  // 用户 workspace
auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();  // 系统 workspace
size_t *currentWorkspace = context->GetWorkspaceSizes(1);
currentWorkspace[0] = usrSize + sysWorkspaceSize;  // 总大小 = 用户 + 系统

参考 SetSysWorkSpace 官方文档。

第 4 步：核对 blockDim 与物理核数

auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
uint32_t aicNum = ascendcPlatform.GetCoreNumAic();   // Cube 核数
uint32_t aivNum = ascendcPlatform.GetCoreNumAiv();   // Vector 核数
// AIC_ONLY: blockDim = aicNum
// AIV_ONLY: blockDim = aivNum
// MIX_AIC_1_2: blockDim = 组合数 = min(aicNum, aivNum / 2)

第 5 步：确认所有核函数都设置了类型

如果 `.cpp` 里有多个 `global aicore` 函数，每一个都必须设置 Kernel 类型，否则编译器无法自动推导。

四、方案 A：保持 AIC_ONLY（推荐，零改动方案）

适用条件：kernel 只有 Cube 计算，无 Vector 侧逻辑。你目前的情况就是这种。

Kernel 侧完整模板

#include "kernel_operator.h"
// 算子处理类
class MyCubeOnlyOp {
public:
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z,
                                const MyTilingData& tiling)
    {
        // 初始化 Global Tensor
        xGm.SetGlobalBuffer((__gm__ half*)x);
        yGm.SetGlobalBuffer((__gm__ half*)y);
        zGm.SetGlobalBuffer((__gm__ float*)z);
        // ... tiling 参数初始化
    }
    __aicore__ inline void Process()
    {
        // 纯 Cube 逻辑：只有 Mmad / Matmul 等
        // 不要出现 Add/Exp/Softmax 等 Vector API
    }
private:
    GlobalTensor<half> xGm, yGm;
    GlobalTensor<float> zGm;
};
// 核函数
extern "C" __global__ __aicore__ void my_cube_op_custom(
    GM_ADDR x, GM_ADDR y, GM_ADDR z,
    GM_ADDR workspace, GM_ADDR tiling)
{
    GET_TILING_DATA(tilingData, tiling);
    if (workspace != nullptr) {
        SetSysWorkspace(workspace);  // 如使用 Matmul 高阶 API 需要系统 workspace
    }
    MyCubeOnlyOp op;
    op.Init(x, y, z, tilingData);
    // 关键：声明为纯 Cube 类型
    KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIC_ONLY);
    op.Process();  // 只有 Cube 逻辑
}

Host 侧 Tiling 完整模板

static ge::graphStatus TilingFunc(gert::TilingContext* context)
{
    // 1. shape 推导
    auto shape = context->GetInputShape(0)->GetStorageShape();
    // ...
    // 2. 多核切分：blockDim = Cube 核数
    auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
    uint32_t coreNum = ascendcPlatform.GetCoreNumAic();  // AIC_ONLY: 取 Cube 核数
    context->SetBlockDim(coreNum);
    // 3. workspace 设置（纯 Cube 类型允许为 0，除非用了 Matmul 高阶 API）
    size_t usrSize = 0;  // 无用户 workspace 需求时可为 0
    uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();  // 系统需求
    size_t *currentWorkspace = context->GetWorkspaceSizes(1);
    currentWorkspace[0] = usrSize + sysWorkspaceSize;
    // 4. tiling 数据填充
    MyTilingData tiling;
    // ... 填充 tiling 字段
    tiling.SaveToBuffer(context->GetRawTilingData()->GetData(),
                        context->GetRawTilingData()->GetCapacity());
    return ge::GRAPH_SUCCESS;
}

五、方案 B：真正改造为 MIX_AIC_1_2（含完整代码）

适用条件：算子确实需要 Cube（矩阵乘）+ Vector（激活/softmax/量化等）两类核并行。例如 Matmul+LeakyRelu、Matmul+GELU 这类融合算子。

5.1 改造核心：三件事必须同时做

改造项	具体内容	不做的后果
① kernel 内部分支	用 `ASCEND_IS_AIC` / `ASCEND_IS_AIV` 隔离两类核的代码	Vector 核拿到 Cube 指令流，无法执行
② workspace 设非零	`usrSize + sysWorkspaceSize`，且 `usrSize > 0`	编译校验失败
③ blockDim 语义调整	设为“组合数”而非 Cube 核数	核数超物理上限

5.2 Kernel 侧完整模板

#include "kernel_operator.h"
// 混合算子处理类：Cube 和 Vector 各自的 Init/Process
class MyMixOp {
public:
    // AIC 侧初始化（Cube 计算）
    __aicore__ inline void InitAIC(GM_ADDR a, GM_ADDR b, GM_ADDR c,
                                   const MyTilingData& tiling)
    {
        aGm.SetGlobalBuffer((__gm__ half*)a);
        bGm.SetGlobalBuffer((__gm__ half*)b);
        cGm.SetGlobalBuffer((__gm__ float*)c);
        // Cube tiling 初始化
    }
    // AIV 侧初始化（Vector 计算）
    __aicore__ inline void InitAIV(GM_ADDR a, GM_ADDR b, GM_ADDR c,
                                   GM_ADDR workspace, const MyTilingData& tiling)
    {
        // Vector 侧 UB 队列初始化
        pipe.InitBuffer(inQueueA, 2, tileSize * sizeof(half));
        pipe.InitBuffer(inQueueB, 2, tileSize * sizeof(half));
        pipe.InitBuffer(outQueueC, 2, tileSize * sizeof(float));
        // 如需 workspace 暂存：workspaceGm.SetGlobalBuffer((__gm__ uint8_t*)workspace);
    }
    // AIC 侧处理：纯 Cube 逻辑
    __aicore__ inline void ProcessAIC()
    {
        // Mmad / Matmul 等 Cube 计算
        // 结果可通过 Fixpipe 输出到 GM 或 UB 供 AIV 读取
    }
    // AIV 侧处理：Vector 逻辑
    __aicore__ inline void ProcessAIV()
    {
        // 从 GM/UB 读取 AIC 输出，做 Add/Softmax/激活等 Vector 计算
        // 结果写回 GM
    }
private:
    GlobalTensor<half> aGm, bGm;
    GlobalTensor<float> cGm;
    TPipe pipe;
    TQue<TPosition::VECIN, 2> inQueueA, inQueueB;
    TQue<TPosition::VECOUT, 2> outQueueC;
};
// 核函数：MIX_AIC_1_2 版本
extern "C" __global__ __aicore__ void my_mix_op_custom(
    GM_ADDR a, GM_ADDR b, GM_ADDR c,
    GM_ADDR workspace, GM_ADDR tiling)
{
    GET_TILING_DATA(tilingData, tiling);
    // MIX 类型必须设置 workspace（不能为 nullptr / 0）
    if (workspace == nullptr) {
        return;
    }
    SetSysWorkspace(workspace);
    MyMixOp op;
    // 关键：声明为 MIX AIC:AIV = 1:2
    KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2);
    // 关键：按核身份分支隔离
    if ASCEND_IS_AIC {
        // Cube 核执行：矩阵乘
        op.InitAIC(a, b, c, tilingData);
        op.ProcessAIC();
    } else {
        // Vector 核执行：向量计算（注意 MIX_AIC_1_2 下每个 Cube 配 2 个 Vector）
        op.InitAIV(a, b, c, workspace, tilingData);
        op.ProcessAIV();
    }
    // 如有核间数据依赖，在此处做同步
    // AscendC::SyncAll();  // 硬同步（需配套 KERNEL_TYPE_MIX_*_1_0 才带硬件同步）
}

💡 关于 ASCEND_IS_AIC / ASCEND_IS_AIV 宏：这是官方提供的条件编译宏，用于分离模式下 AIV/AIC 代码隔离。注意：当使用高阶 API Matmul 时，其内部已通过 REGIST_MATMUL_OBJ 宏方式实现了 AIV 与 AIC 核代码的隔离，用户无需再使用该宏进行处理。
💡 另一种写法（__mix__ 修饰符 + Async）：CANN 还提供了更简洁的写法，直接在核函数上标注 mix 比例，并用 Async 模板分派：

__global__ __mix__(1,2) void mmad_custom(GM_ADDR a, GM_ADDR b, GM_ADDR c)
{
    AscendC::InitSocState();
    MyMixOp op;
    // Async 自动分派到对应核，无需手写 if ASCEND_IS_AIC
    AscendC::Async<AscendC::EngineType::AIC, cubeProcess>(op, a, b, c);
    AscendC::Async<AscendC::EngineType::AIV, vectorProcess>(op, a, b, c);
}

这种写法避免了手写硬件条件分支。两种方式选其一即可。

5.3 Host 侧 Tiling 完整模板

static ge::graphStatus TilingFunc(gert::TilingContext* context)
{
    // 1. shape 推导
    // ...
    auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
    // 2. 多核切分：MIX_AIC_1_2 的 blockDim = 组合数
    uint32_t aicNum = ascendcPlatform.GetCoreNumAic();   // Cube 核数
    uint32_t aivNum = ascendcPlatform.GetCoreNumAiv();   // Vector 核数
    // 950PR: aicNum=32, aivNum=64，物理配比 1:2
    // 组合数 = min(aicNum, aivNum / 2) = 32
    uint32_t blockDim = std::min(aicNum, aivNum / 2);
    context->SetBlockDim(blockDim);  // 启动 blockDim 组（每组 1 Cube + 2 Vector）
    // 3. workspace 设置：MIX 类型强制要求非零
    size_t usrSize = 256;  // 用户 workspace，必须 > 0（官方建议 16/32 等小值即可）
    uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
    size_t *currentWorkspace = context->GetWorkspaceSizes(1);
    currentWorkspace[0] = usrSize + sysWorkspaceSize;
    // 4. tiling 数据填充
    // ...
    return ge::GRAPH_SUCCESS;
}

5.4 950PR 上的 blockDim 计算实例

950PR 物理核：32 Cube + 64 Vector（配比恰好 1:2）。

Kernel 类型	blockDim 建议值	实际启动的核
`AIC_ONLY`	32（= GetCoreNumAic）	32 Cube
`AIV_ONLY`	64（= GetCoreNumAiv）	64 Vector
`MIX_AIC_1_2`	32（= min(32, 64/2)）	32 Cube + 64 Vector

六、两种策略对比：保持 AIC_ONLY vs 改造成 MIX

对比项	保持 AIC_ONLY	改造成 MIX_AIC_1_2
改动量	零（已就绪）	大（kernel 分支 + workspace + blockDim）
性能	Cube 核满载，无空转核	取决于 Vector 逻辑占比：Vector 重则显著加速，纯 Cube 则无收益甚至倒退
精度风险	无	核间同步/数据传递可能引入新问题
适用场景	纯矩阵乘算子	矩阵乘 + 激活/量化/softmax 融合算子
调试复杂度	低	高（需同时调试两类核的流水）
推荐度	⭐⭐⭐⭐⭐（你当前情况）	⭐⭐（仅当确有 Vector 并行需求时）
给你的明确建议：

当前保持 KERNEL_TYPE_AIC_ONLY 不变——这是正确且性能无损的选择，不要为了“用上 MIX”而强行改造。
性能不会因此变慢：纯 Cube 算子用 AIC_ONLY，所有 Cube 核满负荷工作，没有空转的 Vector 核消耗调度资源。
未来需要切 MIX 的时机：当 kernel 里真正加入 Vector 并行逻辑（比如融合 softmax、激活函数、量化前后处理）时，再按方案 B 的三件套（分支 + workspace + blockDim）一起改。
改 MIX 前先做性能预估：用 profiler 采集当前 AIC_ONLY 版本的流水占比，如果 Vector 类计算占比超过 20-30%，MIX 才有明确的收益空间；否则纯 Cube 路径已经打满，MIX 只会增加复杂度。

七、如需进一步定位 error 95

如果你想在当前基础上精确定位 error 95 的触发环节，请提供以下信息：

完整报错日志：开启 ASCEND_GLOBAL_LOG_LEVEL=0 + ASCEND_SLOG_PRINT_TO_STDOUT=1 后，error 95 前后各 10 行的输出
报错发生的阶段：是在 msopgen / 编译阶段，还是 kernel 运行阶段
当时的完整 kernel 代码（含 KERNEL_TASK_TYPE_DEFAULT 那一行前后的上下文）
Host 侧 tiling 函数中 workspace 和 blockDim 的设置代码
有了这些信息，可以逐行对照官方约束精确定位是哪一条校验失败，而不是靠现象反推。