深入理解华为 CANN 中的 SPMD 模型与核函数：Ascend 算子并行编程实践（训练营深度实战篇）

在昇腾 AI 处理器的算子开发体系中，Ascend C 构建了一套完整的设备端编程模型，让开发者能够以接近传统 C/C++ 的形式，直接操控多核 AI Core 的执行。无论是高性能数学库、深度学习算子，还是自定义算子，都依托于一个核心思想：SPMD（Single Program, Multiple Data）并行模型。

一键难忘

1035人浏览 · 2025-11-26 11:12:48

一键难忘 · 2025-11-26 11:12:48 发布

深入理解华为 CANN 中的 SPMD 模型与核函数：Ascend 算子并行编程实践

在昇腾 AI 处理器的算子开发体系中，Ascend C 构建了一套完整的设备端编程模型，让开发者能够以接近传统 C/C++ 的形式，直接操控多核 AI Core 的执行。无论是高性能数学库、深度学习算子，还是自定义算子，都依托于一个核心思想：SPMD（Single Program, Multiple Data）并行模型。

本文将以开发者视角，深入拆解 CANN 中 SPMD 的执行机制、核函数的结构与设计规范，并结合示例分析多核计算是如何实现数据划分、同步与调度的。文章的目标不是单纯介绍 API，而是帮助你理解“为什么是这样设计的”。

训练营简介

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

报名链接：https://www.hiascend.com/developer/activities/cann20252#cann-camp-2502-intro
在这里插入图片描述

一、为什么 Ascend C 选择 SPMD 模型？

在这里插入图片描述

在 AI 算子计算中，性能通常由两个关键因素决定：

每个核心的计算能力
如何同时让更多 AI Core 参与计算

SPMD 模型很好地解决了第二点。

SPMD 最早用于高性能并行计算，它的核心思想很简单：

所有处理单元执行同一份代码，但处理不同的数据片段。

这意味着我们只需要编写一份算子实现，当算子被调用时，系统会自动启动多个核心实例，每个核心负责处理一部分数据。

在深度学习算子中，这种数据并行方式非常适用：
矩阵加法、卷积、激活函数、Softmax…
只要数据可以被切分，SPMD 就能让几十甚至上百个核心并行工作。

举个直观的例子：

如果一段输入数据需要经历 T1、T2、T3 三个阶段，那么在单核执行模式下就是一条线性 pipeline；而在 SPMD 下，几十个核心同时处理不同的数据片段，整个 pipeline 被完全并行化，吞吐大幅提升。

在这里插入图片描述

二、Ascend AI Core 中 SPMD 的具体落地方式

在 Ascend AI Core 中，“核心（Core）”就对应文档中提到的 SPMD 模型中的“进程”。区别在于：

每个核心共享同一份指令代码
每个核心通过 block_idx 区分身份
block_idx 决定它应该处理哪段数据

从编程角度看，block_idx 就是“我是谁？”的答案。

Ascend C 提供了 GetBlockIdx() 接口来获取当前核心的逻辑 ID。
当算子启动 N 个核心（例：<<<N, ...>>>）时，系统会创建 N 个实例，每个实例在设备端运行同一份核函数代码，只是 block_idx 不同。

数据切片的本质：起始地址加偏移

在典型的算子（如 Add）中，不同核心通过偏移实现数据分片：

xGm.SetGlobalBuffer((__gm__ half*)x + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);

这行代码意味着：

x 是整个输入数组
BLOCK_LENGTH 是每个核心负责的数据长度
GetBlockIdx() 决定当前核心负责第几个分片

数据切分不依赖框架，而是在算子内部进行，这也是 Ascend C 能非常灵活的原因。

三、核函数：Ascend C 算子的执行入口

在这里插入图片描述

所有 Ascend C 设备端执行代码都以**核函数（Kernel Function）**为入口，类似于 CUDA 的 device kernel。

一个标准核函数需要满足以下条件：

1. 使用函数限定符

extern "C" __global__ __aicore__ void add_custom(...)

含义如下：

__global__: 表示该函数是核函数，可被主机用 <<<>>> 调用
__aicore__: 该函数在设备的 AI Core 上执行
extern "C": 禁止 C++ name mangling，方便符号查找

这三者构成了一个合法核函数的基本框架。

2. 入参与变量限定符

所有指向 Global Memory 的指针需要使用 __gm__ 或 GM_ADDR 修饰。

推荐使用 GM_ADDR：

#define GM_ADDR __gm__ uint8_t*

这样核函数声明更简洁：

extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)

注意，“GM_ADDR 只是入口类型”，你仍然需要将其转换为实际的数据类型（如 half*）。

3. 必须是 void 返回值

核函数不允许返回值，一切输出都通过 Global Memory 传输。

四、核函数内部结构：对象式算子设计

Ascend C 通常推荐使用“算子类”封装计算与数据流逻辑，典型模式如下：

extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
    KernelAdd op;     // 创建算子对象
    op.Init(x, y, z); // 初始化，包括数据切片、queue 分配等
    op.Process();     // 执行关键计算
}

这种写法有几个优势：

实现代码更模块化
可以为每个核心创建独立状态
算子逻辑更容易被工程化集成

例如，Init 中通常包含以下功能：

根据 block_idx 计算数据偏移
为每个核心分配 pipe、queue 缓冲区
初始化 Local Memory / Buffer

而 Process 则负责：

Tile 分块循环
GM → Local Memory 数据搬运
Vector / Cube 计算指令
Local Memory → GM 写回

五、如何在主机端调用核函数？

核函数采用扩展语法调用：

kernel<<<blockDim, l2ctrl, stream>>>(args...)

三个参数分别代表：

1. blockDim：启动多少个核心？

这是最重要的配置。它决定 SPMD 启动多少个实例。

不同处理器架构下规则不同，但经验上：

Vector-only 算子设置为 AIV 核数
Cube-only 算子设置为 AIC 核数
Vector/Cube 混合算子根据“物理组合核数”配置

典型调用示例：

add_custom<<<8, nullptr, stream>>>(x, y, z);

表示：启动 8 个核心执行 add_custom。

2. l2ctrl（保留字段）

目前填 nullptr 即可。

3. stream：执行流

与 CUDA stream 类似，用于保证异步执行顺序。

调用结束后若需要等待，可调用：

aclrtSynchronizeStream(stream);

六、典型算子示例：并行 Add 的核心逻辑

下面是 Add 的简化示例，它展示了典型 SPMD 工作方式：

核函数

extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
    KernelAdd op;
    op.Init(x, y, z);
    op.Process();
}

主机调用

void add_custom_do(uint32_t blockDim, void* l2ctrl, void* stream,
                   uint8_t* x, uint8_t* y, uint8_t* z)
{
    add_custom<<<blockDim, l2ctrl, stream>>>(x, y, z);
}