前言

开发好的 Ascend C 算子，最终需要被上层业务代码或框架调用 —— 有的场景需要在 C++ 中直接调用，有的场景需要在 Python 中通过 PyTorch 调用，还有的场景需要通过 Ascend CL（计算库）调用。不同的调用方式对应不同的业务需求，掌握这些方式能让你的算子适配更多场景。

本文将详细解析 Ascend C 算子的 4 种核心调用方式：Kernel 直调、Ascend CL 调用、PyTorch 调用、MindSpore 调用，并结合实际案例说明每种方式的实现步骤与适用场景，帮助大家打通算子开发与应用的 “最后一公里”。

一、算子调用的核心逻辑

不管采用哪种调用方式，算子调用的核心逻辑都是一致的：

1. 准备输入数据：将业务数据转换为 Tensor 格式，拷贝到 Device 侧内存；
2. 调用算子接口：通过对应方式的 API 调用算子，触发 Device 侧的计算；
3. 获取输出结果：将 Device 侧的计算结果拷贝回 Host 侧，转换为业务数据格式。

不同调用方式的差异，主要在于 “接口形式” 和 “数据流转方式”—— 例如，Kernel 直调需要手动管理 Device 侧内存，而 PyTorch 调用则由框架自动管理内存。

二、方式 1：Kernel 直调（Host 侧 C++ 直接调用）

Kernel 直调是最底层的调用方式，直接在 Host 侧 C++ 代码中调用 Device 侧的 Kernel 函数，无需经过框架封装。

2.1 适用场景

• Kernel 逻辑的快速验证；
• 小型 C++ 项目中的算子调用；
• 需要精细控制内存与执行流程的场景。

2.2 实现步骤（以 Add 算子为例）

步骤 1：编写 Kernel 函数

c++

__global__ void AddKernel(
    const half* d_x1,
    const half* d_x2,
    half* d_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, d_x1 + start, (end - start) * sizeof(half));
    memcpy(local_x2, d_x2 + start, (end - start) * sizeof(half));
    vadd(local_y, local_x1, local_x2, end - start);
    memcpy(d_y + start, local_y, (end - start) * sizeof(half));
}

步骤 2：Host 侧调用 Kernel

c++

#include <iostream>
#include <vector>
#include "ascend_c_runtime.h"

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

    // 2. 准备Host侧数据
    int64_t size = 1024;
    std::vector<half> h_x1(size, 1.0f);
    std::vector<half> h_x2(size, 2.0f);
    std::vector<half> h_y(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, h_x1.data(), size * sizeof(half), HOST_TO_DEVICE);
    rt.Memcpy(d_x2, h_x2.data(), size * sizeof(half), HOST_TO_DEVICE);

    // 5. 配置线程块与线程格
    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(h_y.data(), d_y, size * sizeof(half), DEVICE_TO_HOST);

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

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

    return 0;
}

2.3 优缺点

优点：

• 控制粒度细：可手动管理内存、线程配置，灵活性高；
• 调试方便：直接在 C++ 代码中打印日志，定位问题快。

缺点：

• 开发成本高：需手动处理内存分配、数据拷贝；
• 复用性差：无法被 Python 代码或上层框架调用。

三、方式 2：Ascend CL 调用（C++ 计算库调用）

Ascend CL（Ascend Computing Library）是昇腾提供的 C 语言计算库，支持调用注册好的 Ascend C 算子，是 C++ 项目中集成算子的推荐方式。

3.1 适用场景

• 中大型 C++ 项目中的算子集成；
• 需要与昇腾其他计算库（如 DVPP、MediaProcess）配合使用的场景；
• 追求性能与易用性平衡的场景。

3.2 实现步骤（以 Add 算子为例）

步骤 1：注册算子并编译为动态库

c++

// 注册Add算子
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);

编译命令：

bash

运行

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

步骤 2：通过 Ascend CL 调用算子

c++

#include <iostream>
#include <vector>
#include "acl/acl.h"

int main() {
    // 1. 初始化Ascend CL
    aclInit(nullptr);
    aclSetDevice(0);

    // 2. 加载算子动态库
    aclopLoadFromFile("./output/add_op.so");

    // 3. 准备Host侧数据
    int64_t size = 1024;
    std::vector<half> h_x1(size, 1.0f);
    std::vector<half> h_x2(size, 2.0f);
    std::vector<half> h_y(size, 0.0f);

    // 4. 创建Device侧Tensor
    aclTensorDesc* desc_x1 = aclCreateTensorDesc(ACL_FLOAT16, 1, &size, ACL_FORMAT_ND);
    aclDataBuffer* buffer_x1 = aclCreateDataBuffer(h_x1.data(), size * sizeof(half));
    aclTensor* tensor_x1 = aclCreateTensor(desc_x1, buffer_x1);

    aclTensorDesc* desc_x2 = aclCreateTensorDesc(ACL_FLOAT16, 1, &size, ACL_FORMAT_ND);
    aclDataBuffer* buffer_x2 = aclCreateDataBuffer(h_x2.data(), size * sizeof(half));
    aclTensor* tensor_x2 = aclCreateTensor(desc_x2, buffer_x2);

    aclTensorDesc* desc_y = aclCreateTensorDesc(ACL_FLOAT16, 1, &size, ACL_FORMAT_ND);
    aclDataBuffer* buffer_y = aclCreateDataBuffer(h_y.data(), size * sizeof(half));
    aclTensor* tensor_y = aclCreateTensor(desc_y, buffer_y);

    // 5. 调用算子
    aclopExecute("Add", 2, &tensor_x1, 1, &tensor_y, nullptr);

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

    // 7. 释放资源
    aclDestroyTensor(tensor_x1);
    aclDestroyTensor(tensor_x2);
    aclDestroyTensor(tensor_y);
    aclUnloadOp("Add");
    aclResetDevice(0);
    aclFinalize();

    return 0;
}

3.3 优缺点

优点：

• 易用性好：无需手动管理 Device 侧内存，由 Ascend CL 自动处理；
• 兼容性强：支持与昇腾其他库配合使用，适合复杂项目；
• 性能稳定：经过昇腾官方优化，调用效率高。

缺点：

• 学习成本：需要熟悉 Ascend CL 的 API 体系；
• 灵活性略低：内存与执行流程由库管理，自定义空间较小。

四、方式 3：PyTorch 调用（Python 框架集成）

Ascend C 算子可以通过 PyTorch 的自定义算子接口集成到 Python 项目中，是 AI 模型开发中最常用的调用方式。

4.1 适用场景

• 基于 PyTorch 的 AI 模型开发；
• Python 业务代码中的算子调用；
• 需要与其他 PyTorch 算子配合使用的场景。

4.2 实现步骤（以 Add 算子为例）

步骤 1：编译算子为 PyTorch 扩展

bash

运行

# 编译为PyTorch扩展
python setup.py build_ext --inplace

setup.py 代码：

python

运行

from setuptools import setup, Extension
from torch.utils.cpp_extension import BuildExtension, AscendExtension

setup(
    name='add_op',
    ext_modules=[
        AscendExtension(
            'add_op',
            sources=['add_op.cpp'],
            extra_compile_args=['-std=c++17']
        )
    ],
    cmdclass={'build_ext': BuildExtension}
)

步骤 2：PyTorch 中调用算子

python

运行

import torch
import add_op  # 导入编译后的扩展

# 1. 准备输入Tensor（NPU设备）
device = torch.device('npu:0')
x1 = torch.tensor([1.0, 2.0, 3.0], dtype=torch.float16, device=device)
x2 = torch.tensor([4.0, 5.0, 6.0], dtype=torch.float16, device=device)

# 2. 调用Ascend C算子
y = add_op.Add(x1, x2)

# 3. 输出结果
print("PyTorch调用结果：", y)  # 输出：tensor([5., 7., 9.], device='npu:0', dtype=torch.float16)

4.3 进阶：通过 torch.ops 调用

对于注册到昇腾框架的算子，还可以通过torch.ops直接调用：

python

运行

import torch

# 加载算子
torch.ops.load_library("./output/add_op.so")

# 调用算子
x1 = torch.tensor([1.0, 2.0], dtype=torch.float16, device='npu')
x2 = torch.tensor([3.0, 4.0], dtype=torch.float16, device='npu')
y = torch.ops.ascend.add(x1, x2)
print(y)  # 输出：tensor([4., 6.], device='npu:0', dtype=torch.float16)

4.4 优缺点

优点：

• 易用性极高：符合 PyTorch 开发者的使用习惯；
• 集成性好：可直接嵌入 PyTorch 模型中，与其他算子无缝配合；
• 生态丰富：支持 PyTorch 的自动微分、分布式训练等功能。

缺点：

• 性能损耗：Python 与 C++ 的交互会带来少量性能开销；
• 调试复杂度：涉及 Python、C++、Device 侧三层，调试需要兼顾多个环节。

五、方式 4：MindSpore 调用（全场景框架集成）

MindSpore 是昇腾原生支持的全场景 AI 框架，Ascend C 算子可以通过 MindSpore 的 Custom 算子接口实现无缝调用。

5.1 适用场景

• 基于 MindSpore 的 AI 模型开发；
• 昇腾全栈解决方案中的算子集成；
• 需要端边云协同的场景。

5.2 实现步骤（以 Add 算子为例）

步骤 1：注册算子并编译

bash

运行

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

步骤 2：MindSpore 中调用算子

python

运行

import mindspore as ms
from mindspore import ops, Tensor

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

# 2. 准备输入Tensor
x1 = Tensor([1.0, 2.0, 3.0], dtype=ms.float16)
x2 = Tensor([4.0, 5.0, 6.0], dtype=ms.float16)

# 3. 调用算子（方式1：直接通过算子名称调用）
add_op = ops.Custom("Add", out_shape=lambda x, y: x.shape, out_dtype=lambda x, y: x.dtype)
y = add_op(x1, x2)

# 4. 输出结果
print("MindSpore调用结果：", y)  # 输出：[5. 7. 9.]

# 进阶：方式2：结合MindSpore模型使用
class AddModel(ms.nn.Cell):
    def __init__(self):
        super(AddModel, self).__init__()
        self.add_op = ops.Custom("Add", out_shape=lambda x, y: x.shape, out_dtype=lambda x, y: x.dtype)
    
    def construct(self, x1, x2):
        return self.add_op(x1, x2)

# 实例化模型并运行
model = AddModel()
y_model = model(x1, x2)
print("MindSpore模型调用结果：", y_model)  # 输出：[5. 7. 9.]

5.3 优缺点

优点：

• 原生适配昇腾：与昇腾硬件、CANN 工具链深度协同，性能最优；
• 全场景支持：支持端、边、云多种部署场景，算子可跨设备复用；
• 动态图 / 静态图兼容：同时支持动态图调试与静态图优化。

缺点：

• 框架依赖性强：仅适用于 MindSpore 生态，跨框架兼容性弱；
• 自定义配置复杂：部分高级功能（如动态 Shape 适配）需要额外配置。

六、4 种调用方式的对比与选择策略

6.1 核心特性对比

调用方式	开发成本	性能表现	框架兼容性	适用场景
Kernel 直调	高	最优	无	原型验证、C++ 小型项目
Ascend CL 调用	中	优秀	C++ 生态	中大型 C++ 项目、多库协同
PyTorch 调用	低	良好	PyTorch 生态	AI 模型开发、Python 业务
MindSpore 调用	低	优秀	MindSpore 生态	昇腾全栈方案、端边云协同

6.2 选择策略

• 若需快速验证 Kernel 逻辑，无需框架集成：选择Kernel 直调；
• 若开发 C++ 项目，需与昇腾其他库配合：选择Ascend CL 调用；
• 若基于 PyTorch 开发 AI 模型：选择PyTorch 调用；
• 若使用昇腾全栈解决方案，需跨设备部署：选择MindSpore 调用。

常见问题

1. PyTorch 调用 Ascend C 算子时，提示 “找不到 NPU 设备”，该如何解决？答：首先确认已安装昇腾 NPU 驱动与 CANN 工具链，且环境变量（如ASCEND_HOME、LD_LIBRARY_PATH）配置正确；其次检查 PyTorch 是否支持昇腾 NPU（需安装适配昇腾的 PyTorch 版本）；最后在代码中显式指定设备（device = torch.device('npu:0')），避免设备识别失败。

2. MindSpore 调用算子时，动态 Shape 场景下计算结果错误，是什么原因？答：可能是算子未适配动态 Shape，或out_shape函数配置错误。解决方案：一是在算子注册时开启动态 Shape 支持（通过.DYNAMIC_SHAPE()声明）；二是确保out_shape函数能正确处理所有可能的输入 Shape（如使用lambda x, y: x.shape而非固定 Shape）；三是在 Model 的construct函数中添加 Shape 校验逻辑。

3. Ascend CL 调用算子时，出现 “算子未找到” 错误，如何排查？答：排查步骤：① 确认算子动态库路径正确（aclopLoadFromFile的路径需准确）；② 检查算子注册名称与调用名称一致（如注册时为 “Add”，调用时不能写 “add”）；③ 验证算子编译成功（查看输出目录是否生成.so文件与op_info.json）；④ 检查 CANN 版本与算子编译时的目标版本匹配。

4. Kernel 直调时，如何处理多 Device 场景（如多卡并行）？答：需在代码中手动管理多 Device 资源：① 通过rt.GetDeviceCount()获取设备数量；② 循环切换设备（rt.SetDevice(device_id)）；③ 为每个设备单独分配内存、拷贝数据、调用 Kernel；④ 最后汇总各设备的计算结果。若需更复杂的并行逻辑，建议使用 Ascend CL 或框架的分布式接口。

5. 4 种调用方式中，如何统一实现算子的性能监控？答：可通过昇腾提供的aclprof工具进行统一性能监控，无需修改算子代码：① 启动aclprof采集性能数据（指定设备 ID、采集时长）；② 生成性能报告，查看算子的执行时间、内存带宽、计算效率等指标；③ 针对性能瓶颈（如内存拷贝耗时过长），优化对应的调用逻辑（如调整分块大小、使用异步拷贝）。

结语

Ascend C 算子的 4 种调用方式，覆盖了从底层验证到上层应用的全场景需求：Kernel 直调保证了极致的灵活性与性能，Ascend CL 调用平衡了 C++ 项目的开发效率与兼容性，PyTorch 调用适配了主流 AI 模型生态，MindSpore 调用则最大化发挥了昇腾硬件的原生能力。

对于开发者而言，无需拘泥于单一调用方式 —— 可以根据项目的不同阶段灵活切换：先用 Kernel 直调验证算子逻辑，再通过 Ascend CL 或框架调用集成到业务中。掌握这些调用方式，能让你开发的 Ascend C 算子真正实现 “一次开发，多场景复用”。

在后续的实践中，建议结合具体业务场景，重点深耕 1-2 种核心调用方式，并关注算子的兼容性与性能优化，让算子在实际应用中发挥最大价值。

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

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

编辑分享

在文章中添加一些常见问题的解答

常见问题的添加位置有哪些？

常见问题的数量有要求吗？