1. 引言：为什么需要 Ascend C？

随着人工智能模型规模的爆炸式增长，通用 GPU 在能效比和专用场景下的性能逐渐遇到瓶颈。华为推出的 昇腾（Ascend）系列 AI 芯片，凭借其独特的 达芬奇架构（Da Vinci Architecture），在推理和训练场景中展现出卓越的性能与能效优势。

然而，要充分发挥昇腾芯片的潜力，仅依赖高层框架（如 MindSpore、TensorFlow）提供的内置算子是远远不够的。当面对 定制化模型结构、新型神经网络层或极致性能优化需求 时，开发者必须深入到底层，编写 自定义高性能算子。

为此，华为推出了 Ascend C —— 一种专为昇腾 AI 处理器设计的 C++ 扩展编程语言。它允许开发者直接操作芯片的计算单元（Cube Unit）、向量单元（Vector Unit）和片上内存（Unified Buffer, UB），实现对硬件资源的精细控制，从而获得接近理论峰值的计算性能。

本文将带您从零开始，系统学习 Ascend C 的核心概念，并通过一个经典的 GEMM（General Matrix Multiplication） 算子实现，掌握其编程范式。

2. Ascend C 开发环境准备

在动手编码前，请确保已配置好以下环境：

  • 硬件：昇腾 910/310 芯片服务器（或 Atlas 系列加速卡）
  • 软件栈：
    • CANN（Compute Architecture for Neural Networks）5.1 或更高版本
    • Ascend-cann-toolkit
    • MindSpore（可选，用于集成测试）
  • 开发工具：支持 C++17 的编译器（如 g++）

安装完成后，可通过 npu-smi info 命令验证设备状态。

3. Ascend C 核心概念解析

3.1 计算单元：Cube 与 Vector

昇腾芯片的核心计算单元分为两类：

  • Cube Unit：专为 矩阵乘加（MatMul） 设计，支持 INT8/FP16 数据类型，单周期可完成 16×16×16 的矩阵乘累加。
  • Vector Unit：处理 向量化操作，如加法、乘法、激活函数等，支持 FP16/FP32/INT8/INT32 等多种数据类型。

Ascend C 提供了 cce::cube::matmul 和 cce::vector::add 等内建函数，但更常见的是使用 指令级 API 进行精细控制。

3.2 内存层次：Global Memory 与 Unified Buffer (UB)

  • Global Memory (GM)：片外 DDR，容量大但延迟高。
  • Unified Buffer (UB)：片上高速缓存，带宽极高，但容量有限（通常几十 KB 到几百 KB）。

关键原则：所有计算必须在 UB 中进行。因此，Ascend C 编程的核心之一就是 高效的数据搬运（Data Movement），即通过 DMA（Direct Memory Access） 指令将数据从 GM 加载到 UB，并在计算完成后写回 GM。

3.3 Block 与 Thread 组织

Ascend C 使用 Block-Thread 模型：

  • Block：对应芯片上的一个 AI Core，每个 Block 可独立执行一段核函数（Kernel）。
  • Thread：Block 内的并行执行单元，通常以 16 线程为一组（称为一个 Warp） 协同工作。

开发者需通过 blockIdx 和 threadIdx 来划分任务。

4. 实战：用 Ascend C 实现 GEMM 算子

我们将实现 C = A * B + bias，其中 A (M×K), B (K×N), bias (N,)，输出 C (M×N)。假设所有张量均为 FP16 格式。

4.1 工程结构

gemm_ascendc/
├── kernel/
│   └── gemm_kernel.cpp
├── host/
│   └── gemm_host.cpp
└── CMakeLists.txt

4.2 Kernel 侧代码（gemm_kernel.cpp）

#include "acl/acl.h"
#include "ascendc.h"
#include "common.h"

using namespace cce;

// 定义常量
constexpr int32_t BLOCK_SIZE_M = 64;
constexpr int32_t BLOCK_SIZE_N = 64;
constexpr int32_t BLOCK_SIZE_K = 64;
constexpr int32_t TILE_M = 16;
constexpr int32_t TILE_N = 16;
constexpr int32_t TILE_K = 16;

// Kernel 函数
extern "C" __global__ void gemm_kernel(
    half* __restrict__ a_gm,
    half* __restrict__ b_gm,
    half* __restrict__ bias_gm,
    half* __restrict__ c_gm,
    int32_t M, int32_t N, int32_t K)
{
    // 获取当前 Block ID
    int32_t blockIdX = blockIdx.x;
    int32_t blockIdY = blockIdx.y;

    // 计算当前 Block 负责的输出块起始位置
    int32_t startM = blockIdX * BLOCK_SIZE_M;
    int32_t startN = blockIdY * BLOCK_SIZE_N;

    // 分配 UB 内存（静态分配，编译期确定大小）
    __shared__ half a_ub[BLOCK_SIZE_M * TILE_K];     // [64, 16]
    __shared__ half b_ub[TILE_K * BLOCK_SIZE_N];     // [16, 64]
    __shared__ half bias_ub[BLOCK_SIZE_N];           // [64]
    __shared__ float c_ub[BLOCK_SIZE_M * BLOCK_SIZE_N]; // 累加用 FP32

    // 初始化累加器为 0
    for (int32_t i = 0; i < BLOCK_SIZE_M * BLOCK_SIZE_N; ++i) {
        c_ub[i] = 0.0f;
    }

    // 加载 bias 到 UB（仅由 blockIdX == 0 的 Block 加载）
    if (blockIdX == 0) {
        for (int32_t n = 0; n < BLOCK_SIZE_N; ++n) {
            int32_t global_n = startN + n;
            if (global_n < N) {
                bias_ub[n] = bias_gm[global_n];
            } else {
                bias_ub[n] = 0.0_h;
            }
        }
    }
    __sync(); // 同步所有线程

    // 主循环：沿 K 维度分块
    for (int32_t k0 = 0; k0 < K; k0 += TILE_K) {
        // 1. 从 GM 加载 A 的切片到 UB
        for (int32_t m = 0; m < BLOCK_SIZE_M; ++m) {
            int32_t global_m = startM + m;
            for (int32_t k = 0; k < TILE_K; ++k) {
                int32_t global_k = k0 + k;
                if (global_m < M && global_k < K) {
                    a_ub[m * TILE_K + k] = a_gm[global_m * K + global_k];
                } else {
                    a_ub[m * TILE_K + k] = 0.0_h;
                }
            }
        }

        // 2. 从 GM 加载 B 的切片到 UB
        for (int32_t k = 0; k < TILE_K; ++k) {
            int32_t global_k = k0 + k;
            for (int32_t n = 0; n < BLOCK_SIZE_N; ++n) {
                int32_t global_n = startN + n;
                if (global_k < K && global_n < N) {
                    b_ub[k * BLOCK_SIZE_N + n] = b_gm[global_k * N + global_n];
                } else {
                    b_ub[k * BLOCK_SIZE_N + n] = 0.0_h;
                }
            }
        }
        __sync();

        // 3. 执行分块矩阵乘：利用 Cube 指令
        // 将 A_ub 转置为 [TILE_K, BLOCK_SIZE_M] 以便与 B_ub 对齐
        // 实际中应使用 tiling 策略和 double buffer 优化，此处简化
        for (int32_t m = 0; m < BLOCK_SIZE_M; m += TILE_M) {
            for (int32_t n = 0; n < BLOCK_SIZE_N; n += TILE_N) {
                // 调用内建 matmul 指令（伪代码，实际需使用 ascendc 提供的 intrinsic）
                // 此处用循环模拟
                for (int32_t tm = 0; tm < TILE_M; ++tm) {
                    for (int32_t tn = 0; tn < TILE_N; ++tn) {
                        float sum = 0.0f;
                        for (int32_t tk = 0; tk < TILE_K; ++tk) {
                            sum += static_cast<float>(a_ub[(m+tm)*TILE_K + tk]) *
                                   static_cast<float>(b_ub[tk*BLOCK_SIZE_N + (n+tn)]);
                        }
                        c_ub[(m+tm)*BLOCK_SIZE_N + (n+tn)] += sum;
                    }
                }
            }
        }
        __sync();
    }

    // 4. 加上 bias 并写回 GM
    for (int32_t m = 0; m < BLOCK_SIZE_M; ++m) {
        int32_t global_m = startM + m;
        if (global_m >= M) continue;
        for (int32_t n = 0; n < BLOCK_SIZE_N; ++n) {
            int32_t global_n = startN + n;
            if (global_n >= N) continue;

            float result = c_ub[m * BLOCK_SIZE_N + n];
            if (blockIdX == 0) {
                result += static_cast<float>(bias_ub[n]);
            }
            c_gm[global_m * N + global_n] = static_cast<half>(result);
        }
    }
}

注意：上述代码为教学简化版。实际 Ascend C 开发中，应使用 cce::dma_copy、cce::cube::mma_sync 等 intrinsic 函数 直接调用硬件指令，并采用 double buffering 隐藏 DMA 延迟。

4.3 Host 侧代码（gemm_host.cpp）

Host 侧负责内存分配、数据拷贝和 Kernel 启动。

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

int main() {
    // 1. 初始化 ACL
    aclInit(nullptr);
    aclrtSetDevice(0);
    aclrtCreateContext(nullptr, 0);

    // 2. 分配 Host 内存
    int M = 1024, N = 1024, K = 1024;
    size_t sizeA = M * K * sizeof(half);
    size_t sizeB = K * N * sizeof(half);
    size_t sizeBias = N * sizeof(half);
    size_t sizeC = M * N * sizeof(half);

    std::vector<half> h_a(M*K), h_b(K*N), h_bias(N), h_c(M*N);

    // 初始化数据（略）

    // 3. 分配 Device 内存
    half *d_a, *d_b, *d_bias, *d_c;
    aclrtMalloc(&d_a, sizeA, ACL_MEM_MALLOC_NORMAL_ONLY);
    aclrtMalloc(&d_b, sizeB, ACL_MEM_MALLOC_NORMAL_ONLY);
    aclrtMalloc(&d_bias, sizeBias, ACL_MEM_MALLOC_NORMAL_ONLY);
    aclrtMalloc(&d_c, sizeC, ACL_MEM_MALLOC_NORMAL_ONLY);

    // 4. 拷贝数据到 Device
    aclrtMemcpy(d_a, sizeA, h_a.data(), sizeA, ACL_MEMCPY_HOST_TO_DEVICE);
    aclrtMemcpy(d_b, sizeB, h_b.data(), sizeB, ACL_MEMCPY_HOST_TO_DEVICE);
    aclrtMemcpy(d_bias, sizeBias, h_bias.data(), sizeBias, ACL_MEMCPY_HOST_TO_DEVICE);

    // 5. 配置 Kernel 启动参数
    dim3 blockDim(1, 1, 1); // Ascend C 中 blockDim 通常为 1
    dim3 gridDim((M + 63) / 64, (N + 63) / 64, 1); // 每个 Block 处理 64x64 输出

    // 6. 加载并启动 Kernel（需提前编译 .o 文件并通过 aclmdlLoad）
    // 此处省略模型加载步骤，实际需使用 aclnn 或自定义算子注册机制

    // 7. 拷贝结果回 Host
    aclrtMemcpy(h_c.data(), sizeC, d_c, sizeC, ACL_MEMCPY_DEVICE_TO_HOST);

    // 8. 清理资源
    aclrtFree(d_a); aclrtFree(d_b); aclrtFree(d_bias); aclrtFree(d_c);
    aclrtDestroyContext(nullptr);
    aclFinalize();

    return 0;
}

4.4 性能优化要点

  1. Tiling 策略：选择合适的 BLOCK_SIZE 和 TILE_SIZE 以匹配 UB 容量。
  2. Double Buffering：在计算当前 tile 的同时，预取下一个 tile 的数据，隐藏 DMA 延迟。
  3. 数据重排（Layout Transformation）：将输入数据预处理为 Cube 友好的格式（如 FRACTAL_ZZ）。
  4. 使用 FP16 累加：若精度允许，避免 FP32 累加以节省 UB 空间和带宽。

5. 与 MindSpore 集成

Ascend C 算子可通过 自定义算子注册机制 集成到 MindSpore：

from mindspore.ops import Custom

gemm_op = Custom(
    "gemm_kernel.so",
    out_shape=lambda a, b, bias: (a.shape[0], b.shape[1]),
    out_dtype=lambda a, b, bias: a.dtype,
    func_type="aot"  # Ahead-of-Time 编译
)

# 使用
output = gemm_op(a, b, bias)

6. 总结

本文介绍了 Ascend C 的基本架构，并通过 GEMM 算子展示了其编程模型。虽然示例代码做了简化，但已涵盖 内存管理、数据搬运、计算调度 三大核心要素。在下一篇文章中，我们将深入 卷积算子 的实现，并探讨 性能分析工具（如 msprof） 的使用。

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

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

# c语言 # 人工智能 # 矩阵
Logo
人工智能6S服务平台

作为“人工智能6S店”的官方数字引擎，为AI开发者与企业提供一个覆盖软硬件全栈、一站式门户。

更多推荐

cover

深度详解华为昇腾Ascend C：从零构建高性能AI算子（附完整代码与图解

avatar 人工智能6S服务平台

昇腾Ascend C实战进阶：手把手实现高性能Softmax算子（支持动态Shape + 多核并行）

场景原因大模型推理默认Softmax按行处理，但某些Attention变体需按列或块处理FP8/INT4支持官方算子未开放低精度接口融合优化如或融合减少访存特殊维度如对非最后一维做Softmax（PyTorch默认仅支持dim=-1）💡本文目标：实现一个通用Softmax算子，支持任意输入Shape、任意axis，并在昇腾NPU上达到接近理论带宽的性能。本文通过实现一个工业级Softmax算子✅

avatar 人工智能6S服务平台
cover

昇腾Ascend C高阶实战：实现带Mask的Attention Score算子（支持大模型推理场景）

avatar 人工智能6S服务平台