前言

在 Ascend C 的算子开发实践中，不同的业务场景需要不同的开发流程：有的场景追求 “快速验证”（如算法原型测试），有的场景需要 “生产级稳定性”（如集成到商用模型中）。针对这两种需求，Ascend C 提供了 ** 快速流程（Kernel 直调）和标准流程（自定义算子）** 两种工程化开发方式。

很多开发者在选择流程时容易陷入困惑：什么时候用快速流程？什么时候必须用标准流程？两种流程的核心差异在哪里？本文将从适用场景、实现细节、优缺点对比三个维度，详细解析这两种流程的选择逻辑，帮助大家根据实际需求制定最优的开发策略。

一、算子开发工程的核心目标

不管选择哪种流程，算子开发的核心目标都是一致的：

• 实现算子的计算逻辑（正确完成业务需求）；
• 保证算子的性能（充分利用 AI Core 的硬件资源）；
• 确保算子的易用性（能被上层框架或业务代码调用）。

不同流程的差异，本质是对 “开发效率” 和 “工程化程度” 的权衡：快速流程牺牲部分工程化特性换取开发速度，标准流程则通过完整的工程化设计保证算子的稳定性与兼容性。

二、快速流程：Kernel 直调（适合快速验证）

2.1 什么是 Kernel 直调？

Kernel 直调是一种 “轻量级” 的算子开发方式：跳过 Host 侧的 Op 封装，直接在 Host 侧代码中调用 Device 侧的 Kernel 函数，专注于验证 Kernel 的计算逻辑是否正确。

这种流程的核心是 “极简”—— 不需要编写 Op 的参数校验、Tiling 逻辑，也不需要注册算子，只需实现 Kernel 函数并直接调用。

2.2 适用场景

Kernel 直调主要适用于以下场景：

• Kernel 逻辑的快速验证：如测试一个新的计算逻辑是否能在 AI Core 上正确运行；
• 算法原型的 PoC 验证：在算法探索阶段，快速验证算子的计算效果；
• 小规模、临时性的算子开发：如内部工具中的简单算子，不需要对外提供通用接口。

例如，当我们想测试 “向量加法的并行计算逻辑” 时，用 Kernel 直调可以在 10 分钟内完成代码编写与验证，而无需关注 Op 封装的细节。

2.3 实现步骤

Kernel 直调的实现步骤非常简单，以 Add 算子为例：

步骤 1：编写 Kernel 函数

c++

// Add算子的Kernel函数（运行在AI Core）
__global__ void AddKernel(
    const half* global_x1,
    const half* global_x2,
    half* global_y,
    int64_t size
) {
    int block_id = blockIdx.x;
    int block_size = 256;
    int start = block_id * block_size;
    int end = min(start + block_size, size);

    __local half local_x1[256];
    __local half local_x2[256];
    __local half local_y[256];

    memcpy(local_x1, global_x1 + start, (end - start) * sizeof(half));
    memcpy(local_x2, global_x2 + start, (end - start) * sizeof(half));
    vadd(local_y, local_x1, local_x2, end - start);
    memcpy(global_y + start, local_y, (end - start) * sizeof(half));
}

步骤 2：Host 侧直接调用 Kernel

c++

#include <iostream>
#include <vector>
#include "ascend_c_runtime.h"  // Ascend C运行时头文件

int main() {
    // 1. 初始化Ascend C运行时
    AscendCRuntime rt;
    rt.Init();

    // 2. 准备输入输出数据
    int64_t size = 1024;
    std::vector<half> x1_data(size, 1.0f);
    std::vector<half> x2_data(size, 2.0f);
    std::vector<half> y_data(size, 0.0f);

    // 3. 分配Device侧内存
    half* d_x1 = rt.Malloc<half>(size * sizeof(half));
    half* d_x2 = rt.Malloc<half>(size * sizeof(half));
    half* d_y = rt.Malloc<half>(size * sizeof(half));

    // 4. 将Host侧数据拷贝到Device侧
    rt.Memcpy(d_x1, x1_data.data(), size * sizeof(half), HOST_TO_DEVICE);
    rt.Memcpy(d_x2, x2_data.data(), size * sizeof(half), HOST_TO_DEVICE);

    // 5. 配置Kernel的线程块与线程格
    dim3 grid_dim((size + 256 - 1) / 256, 1, 1);
    dim3 block_dim(1, 1, 1);

    // 6. 直接调用Kernel
    AddKernel<<<grid_dim, block_dim>>>(d_x1, d_x2, d_y, size);

    // 7. 将Device侧结果拷贝回Host侧
    rt.Memcpy(y_data.data(), d_y, size * sizeof(half), DEVICE_TO_HOST);

    // 8. 验证结果
    bool success = true;
    for (int i = 0; i < size; ++i) {
        if (y_data[i] != 3.0f) {
            success = false;
            break;
        }
    }
    std::cout << (success ? "验证成功" : "验证失败") << std::endl;

    // 9. 释放资源
    rt.Free(d_x1);
    rt.Free(d_x2);
    rt.Free(d_y);
    rt.Destroy();

    return 0;
}

步骤 3：编译运行

使用 Ascend C 的编译器编译代码，直接运行可执行文件即可验证结果。

2.4 优缺点分析

优点：

• 开发速度快：无需编写 Op 封装、算子注册等代码；
• 调试成本低：直接在 Host 侧代码中控制 Kernel 的调用流程，便于打印日志、调试问题。

缺点：

• 易用性差：无法被上层框架（如 PyTorch、MindSpore）直接调用；
• 工程化程度低：缺少参数校验、Tiling 优化等逻辑，不适合生产环境；
• 复用性差：每个 Kernel 都需要单独编写调用代码，无法统一管理。

三、标准流程：自定义算子（适合生产环境）

3.1 什么是自定义算子？

自定义算子是 Ascend C 推荐的 “生产级” 开发流程：完整实现 Host 侧的 Op 封装（参数校验、Tiling、Shape 推导）与 Device 侧的 Kernel，并将算子注册到昇腾框架中，使其成为框架的 “原生算子”。

这种流程的核心是 “规范”—— 遵循 Ascend C 的算子开发标准，保证算子的兼容性、稳定性与易用性。

3.2 适用场景

自定义算子主要适用于以下场景：

• 需要被上层框架调用的算子：如集成到 PyTorch 模型中的卷积算子、激活函数算子；
• 生产环境中的核心算子：如商用 AI 系统中的关键计算算子；
• 需要高复用性、高可维护性的算子：如企业内部的算子库开发。

例如，当我们需要将 Add 算子集成到一个基于 MindSpore 的图像分类模型中时，必须使用标准流程开发，才能让模型通过框架的 API 调用该算子。

3.3 实现步骤

标准流程的实现步骤相对完整，以 Add 算子为例：

步骤 1：定义算子原型与 Op 类

c++

// 定义Add算子的Tiling类
class AddTiling : public TilingBase {
public:
    static constexpr int BLOCK_SIZE = 256;
    Status ComputeTiling(const std::vector<TensorPtr>& inputs, const std::vector<TensorPtr>& outputs) override {
        int64_t size = inputs[0]->GetShape()[0];
        int block_num = (size + BLOCK_SIZE - 1) / BLOCK_SIZE;
        grid_dim_.x = block_num;
        block_dim_.x = 1;
        tiling_info_.block_size = BLOCK_SIZE;
        tiling_info_.size = size;
        return Status::SUCCESS;
    }
private:
    AddTilingInfo tiling_info_;
};

// 定义Add算子的Op类
class AddOp : public OpBase {
public:
    AddOp() {
        shape_infer_func_ = std::bind(&AddOp::InferShape, this, std::placeholders::_1, std::placeholders::_2);
    }
    Status InferShape(const std::vector<TensorPtr>& inputs, std::vector<TensorPtr>& outputs) override {
        if (inputs.size() != 2) return Status::FAILED("Need 2 inputs");
        if (inputs[0]->GetShape() != inputs[1]->GetShape()) return Status::FAILED("Shape mismatch");
        outputs.resize(1);
        outputs[0] = std::make_shared<Tensor>(inputs[0]->GetShape(), inputs[0]->GetDtype());
        return Status::SUCCESS;
    }
};

步骤 2：实现 Kernel 函数

c++

__global__ void AddKernel(
    const half* global_x1,
    const half* global_x2,
    half* global_y,
    const AddTilingInfo* tiling_info
) {
    int block_id = blockIdx.x;
    int block_size = tiling_info->block_size;
    int size = tiling_info->size;
    int start = block_id * block_size;
    int end = min(start + block_size, size);

    __local half local_x1[256];
    __local half local_x2[256];
    __local half local_y[256];

    memcpy(local_x1, global_x1 + start, (end - start) * sizeof(half));
    memcpy(local_x2, global_x2 + start, (end - start) * sizeof(half));
    vadd(local_y, local_x1, local_x2, end - start);
    memcpy(global_y + start, local_y, (end - start) * sizeof(half));
}

步骤 3：注册算子

c++

REGISTER_OP(Add)
    .INPUT(x1, TensorType::FLOAT16)
    .INPUT(x2, TensorType::FLOAT16)
    .OUTPUT(y, TensorType::FLOAT16)
    .REQUIRE(x1.shape() == x2.shape(), "Shape mismatch")
    .SET_OP_CONSTRUCT_FUNC(AddOp::Construct)
    .SET_TILING_FUNC(AddTiling::ComputeTiling)
    .SET_KERNEL_FUNC(AddKernel);

步骤 4：编译为算子库

bash

运行

msopgen build --op-name Add --source add_op.cpp --output ./output --target ascend310b

步骤 5：框架调用算子

python

运行

# MindSpore调用Add算子示例
import mindspore as ms
from mindspore import ops

# 加载自定义算子
ops.load_custom_op("./output/add_op.so")

# 调用Add算子
x1 = ms.Tensor([1,2,3], dtype=ms.float16)
x2 = ms.Tensor([4,5,6], dtype=ms.float16)
y = ops.Add()(x1, x2)
print(y)  # 输出：[5,7,9]

3.4 优缺点分析

优点：

• 易用性强：能被上层框架直接调用，符合开发者的使用习惯；
• 工程化程度高：包含参数校验、Tiling 优化等逻辑，稳定性强；
• 复用性好：注册到框架后可被多个业务模块复用，便于维护。

缺点：

• 开发周期长：需要编写 Op 封装、算子注册等代码，步骤较多；
• 调试复杂度高：涉及 Host 侧与 Device 侧的交互，调试需要兼顾多个环节。

四、快速流程 vs 标准流程：对比与选择

4.1 核心差异对比

特性	快速流程（Kernel 直调）	标准流程（自定义算子）
开发速度	快（1 天内可完成）	慢（3-5 天）
工程化程度	低（缺少参数校验、Tiling）	高（完整的工程化设计）
框架兼容性	差（无法被上层框架调用）	好（支持 PyTorch/MindSpore）
复用性	差（需单独编写调用代码）	好（注册后全局可用）
适用场景	原型验证、临时算子	生产环境、框架集成

4.2 选择策略

• 做原型验证 / 算法探索：选快速流程，用最短时间验证 Kernel 逻辑；
• 做生产级算子 / 框架集成：选标准流程，保证算子的稳定性与兼容性；
• 混合使用：先用快速流程验证 Kernel 逻辑，再基于验证后的 Kernel 扩展为标准流程的 Op 封装。

例如，在开发一个新的 Transformer 算子时，可以先用快速流程验证 “多头注意力计算” 的 Kernel 逻辑，确认正确后再添加 Op 封装、算子注册等逻辑，转换为标准流程。

五、从快速流程迁移到标准流程的技巧

很多开发者会先用快速流程验证 Kernel，再迁移到标准流程。以下是迁移的关键技巧：

5.1 复用 Kernel 代码

快速流程中验证过的 Kernel 代码可以直接复用，只需调整参数传递方式（如从直接传参改为通过 TilingInfo 传递）。

5.2 补充 Tiling 逻辑

根据 Kernel 的分块策略，补充 Host 侧的 Tiling 类，实现分块数量计算、线程块配置等逻辑。

5.3 完善参数校验

添加 Shape 推导、Dtype 校验等逻辑，确保算子的鲁棒性。

5.4 注册算子并测试

完成算子注册后，通过框架的 API 调用算子，验证兼容性与正确性。

结语

Ascend C 的快速流程与标准流程，是针对不同开发需求设计的两种工程化路径：快速流程追求 “效率”，标准流程追求 “规范”。理解两种流程的核心差异与适用场景，能帮助我们在实际开发中 “对症下药”—— 既不浪费时间在不必要的工程化环节，也不牺牲生产环境的稳定性。

对于开发者而言，掌握两种流程的开发技巧是提升 Ascend C 算子开发能力的关键：既能用快速流程快速验证创意，又能用标准流程将创意落地为生产级算子。在后续的文章中，我们将深入解析 Ascend C 的 API 体系，帮助大家进一步提升算子开发的效率与质量。

 2025年昇腾CANN训练营第二季，基于CANN开源开放全场景，推出0基础入门系列、码力全开特辑、开发者案例等专题课程，助力不同阶段开发者快速提升算子开发技能。获得Ascend C算子中级认证，即可领取精美证书，完成社区任务更有机会赢取华为手机，平板、开发板等大奖。

报名链接:https://www.hiascend.com/developer/activities/cann20252