前言

昇腾 CANN 作为华为面向 AI 异构计算的核心架构，算子开发是其生态落地的关键能力之一，但入门阶段常因硬件感知特性、核函数语法等门槛让开发者望而却步。本文基于 CANN 训练营初级课程的知识体系，从底层逻辑到代码实战，拆解算子开发的最小闭环 —— 既能帮新手建立 “概念→工具→实践” 的认知链路，也能为有基础的开发者梳理初级阶段的核心要点。

一、算子开发的底层逻辑：从 AI 计算到算子的分层抽象

AI 模型的执行本质是「计算图→算子→硬件指令」的分层映射：

• 计算图：由 TensorFlow/PyTorch 等框架定义的业务逻辑（如 CNN 的卷积层）；
• 算子：计算图的最小执行单元（如 Conv2d、Erf），是连接框架与硬件的桥梁；
• 硬件指令：算子经 CANN 编译后生成的昇腾 NPU（如 Ascend 910）可执行指令。

算子开发的核心目标是：在昇腾硬件上实现高效、高精度的算子计算逻辑。

二、Ascend C：昇腾算子开发的原生语言

Ascend C 是 CANN 提供的算子开发语言，基于 C/C++ 扩展，专为昇腾 NPU 的达芬奇架构设计，核心特性包括：

1. SIMD/SPMD 并行模型：

   • SIMD（单指令多数据）：单个指令同时操作多个数据（适配达芬奇架构的向量计算单元）；
   • SPMD（单程序多数据）：多个线程执行同一程序，通过线程 ID 区分数据分片。

2. 硬件感知的内存模型：昇腾 NPU 内存分为 Global Memory（全局内存）、Local Memory（局部内存）、Register（寄存器），Ascend C 通过限定符实现内存显式管理。

3. 自动化并行与优化：内置算子自动并行、数据类型自动转换等能力，降低开发复杂度。

三、Device 侧核函数开发：从语法到 Hello World 实战

核函数是运行在昇腾 NPU 的 AICore 上的计算函数，是算子开发的核心载体。

3.1 核函数的基础语法规则

Ascend C 核函数需满足以下约束：

• 函数类型限定符：必须用__global__标识，表明是 Device 侧可执行函数；
• 参数类型：仅支持指针类型（指向 Global Memory 的 Tensor 数据）；
• 线程索引：通过blockIdx（线程块 ID）、threadIdx（线程 ID）区分并行线程。

3.2 实战：Device 侧 Hello World 核函数

以下是完整的「Host 侧调用 + Device 侧核函数」示例（基于 CANN 7.0 版本）：

步骤 1：Host 侧代码（负责环境初始化、核函数调用）

c

运行

#include <iostream>
#include "ascendc/ascendc_runtime.h"  // CANN运行时头文件

// 核函数声明（Host侧需提前声明）
__global__ void HelloWorldKernel(const float* input, float* output, uint32_t size);

int main() {
    // 1. 初始化CANN运行时
    aclInit(nullptr);
    aclrtContext context;
    aclrtCreateContext(&context, 0);  // 0为设备ID
    aclrtSetCurrentContext(context);

    // 2. 分配Host/Device内存
    const uint32_t size = 1;
    float hostInput = 0.0f;
    float hostOutput = 0.0f;

    void* deviceInput;
    void* deviceOutput;
    aclrtMalloc(&deviceInput, size * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc(&deviceOutput, size * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST);

    // 3. Host→Device数据拷贝
    aclrtMemcpy(deviceInput, size * sizeof(float), &hostInput, size * sizeof(float), ACL_MEMCPY_HOST_TO_DEVICE);

    // 4. 配置线程块（gridDim/blockDim）
    dim3 gridDim(1);    // 线程块数量：1
    dim3 blockDim(1);   // 每个线程块的线程数：1

    // 5. 启动核函数
    HelloWorldKernel<<<gridDim, blockDim>>>((const float*)deviceInput, (float*)deviceOutput, size);
    aclrtSynchronizeStream(nullptr);  // 同步Stream，等待核函数执行完成

    // 6. Device→Host数据拷贝
    aclrtMemcpy(&hostOutput, size * sizeof(float), deviceOutput, size * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST);

    // 7. 打印结果
    std::cout << "Device侧输出：" << hostOutput << std::endl;  // 输出：Device侧输出：1.0

    // 8. 资源释放
    aclrtFree(deviceInput);
    aclrtFree(deviceOutput);
    aclrtDestroyContext(context);
    aclFinalize();

    return 0;
}

步骤 2：Device 侧核函数代码

c

运行

#include "ascendc/ascendc_kernel.h"  // Ascend C核函数头文件

__global__ void HelloWorldKernel(const float* input, float* output, uint32_t size) {
    // 获取当前线程ID（本例中仅1个线程）
    uint32_t tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid >= size) {
        return;  // 避免线程越界
    }

    // 核心计算逻辑：将输入+1后写入输出
    output[tid] = input[tid] + 1.0f;

    // Device侧打印（需开启NPU日志）
    printf("Hello Ascend C! Thread ID: %u, Input: %f, Output: %f\n", tid, input[tid], output[tid]);
}

3.3 编译与运行说明

• 编译命令：需使用昇腾 CANN 提供的ascendc-cc编译器：

bash

运行

ascendc-cc -o hello_world hello_world_host.cpp hello_world_kernel.cpp -lascendc_runtime -lascendc_kernel

• 运行前提：需部署在昇腾 NPU 服务器（或华为云昇腾实例），并配置 CANN 环境变量。

四、初级算子开发的常见坑点与避坑指南

1. 线程越界：需在核函数中判断tid < size，避免访问超出 Tensor 范围的内存；
2. 内存拷贝方向错误：Host→Device 用ACL_MEMCPY_HOST_TO_DEVICE，反之用ACL_MEMCPY_DEVICE_TO_HOST；
3. 未同步 Stream：核函数是异步执行的，需调用aclrtSynchronizeStream等待执行完成。

 结语

从概念分层到 Device 侧核函数的完整实战，我们完成了昇腾 CANN 初级算子开发的入门闭环。核函数的语法规则、线程管理、内存操作是后续进阶的基础，而避坑指南里的细节（如线程越界、Stream 同步）则是保障算子稳定运行的关键。后续内容将围绕动态 Shape、精度优化等场景深入，也建议大家结合 CANN 官方文档与实际硬件环境多做验证。

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

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