一、引言：为什么需要 Ascend C？

在人工智能大模型时代，算力成为制约模型训练与推理效率的核心瓶颈。传统 GPU 架构虽已成熟，但其通用性设计在特定 AI 负载下存在能效比不足的问题。为此，全球科技巨头纷纷推出专用 AI 加速器（如 Google TPU、NVIDIA NVDLA、华为昇腾 Ascend 系列等）。

华为昇腾（Ascend）系列 AI 处理器正是在此背景下诞生的国产高性能 AI 芯片，广泛应用于 Atlas 系列服务器、边缘设备及云上 AI 推理服务。然而，要充分发挥昇腾芯片的计算潜力，开发者必须深入其硬件架构，编写高度优化的自定义算子（Custom Operator）。

早期昇腾生态主要依赖 TBE（Tensor Boost Engine），基于 Python + DSL 的方式开发算子。虽然上手简单，但在复杂逻辑、细粒度控制和极致性能方面存在局限。

2023 年，华为正式推出 Ascend C —— 一种基于 C++ 的、面向昇腾 AI 处理器的高性能算子开发语言。它直接映射硬件指令，支持流水线调度、双缓冲、向量化等底层优化手段，使开发者能够像编写 CUDA Kernel 一样精细控制昇腾芯片的计算单元（AI Core）和存储层次。

本文将系统性地介绍 Ascend C 的核心概念、开发流程，并通过一个完整的 Vector Add（向量加法） 示例，带领读者从零构建一个可在昇腾 910B 上运行的高性能算子。

---

二、Ascend C 核心特性与优势

1. 贴近硬件的编程模型

Ascend C 直接操作昇腾 AI Core 的计算单元（Cube Unit、Vector Unit）和存储结构（Global Memory、Unified Buffer、Local L1/L0 Buffer），允许开发者精确控制数据搬运与计算流水。

2. C++ 语法，降低学习门槛

相比 TBE 的 DSL，Ascend C 使用标准 C++ 语法（部分受限），熟悉 CUDA 或 OpenCL 的开发者可快速上手。

3. 自动流水线调度

通过 Pipe 机制，Ascend C 编译器可自动将数据搬运（CopyIn/CopyOut）与计算（Compute）重叠，隐藏内存延迟。

4. 内置高性能模板库

提供 DataCopy、Add、Mul、Reduce 等常用算子模板，支持向量化（Vectorization）和分块（Tiling）策略。

5. 端到端工具链支持

集成于 CANN（Compute Architecture for Neural Networks）软件栈，支持编译、仿真、Profiling 与部署。

---

三、开发环境准备

1. 硬件要求

• 昇腾 910/910B AI 处理器（或使用官方提供的 Docker 镜像进行仿真）
• 推荐：Atlas 800/300I 推理服务器

2. 软件依赖

• CANN Toolkit ≥ 7.0.RC1（建议 7.0.RC2 或更高）
• GCC ≥ 7.3
• CMake ≥ 3.14
• Python ≥ 3.7（用于 Host 侧驱动）

注意：Ascend C 目前仅支持在华为官方提供的 CANN 环境中编译运行。

3. 安装 CANN

# 下载 CANN Toolkit（需华为账号）
wget https://ascend.huawei.com/cann-7.0.RC2-linux-x86_64.tar.gz
tar -zxvf cann-7.0.RC2-linux-x86_64.tar.gz
cd cann-7.0.RC2
bash install.sh --install-mode=toolkit --install-path=/usr/local/Ascend

配置环境变量：

export ASCEND_HOME=/usr/local/Ascend/ascend-toolkit/latest
export PATH=$ASCEND_HOME/bin:$PATH
export PYTHONPATH=$ASCEND_HOME/python/site-packages:$PYTHONPATH

---

四、Ascend C 编程模型详解

Ascend C 程序由两部分组成：

1. Kernel 侧（Device Code）

• 使用 Ascend C 编写，运行在 AI Core 上
• 核心类：Kernel, Pipe, GlobalTensor, LocalTensor
• 关键函数：CopyIn, Compute, CopyOut

2. Host 侧（CPU Code）

• 使用 C++ 或 Python 编写，负责内存分配、参数传递、启动 Kernel
• 通过 acl（Ascend Computing Language）API 与设备交互

---

五、实战：使用 Ascend C 实现 Vector Add

我们将实现一个简单的 A + B = C 向量加法算子，输入两个 float32 向量，输出结果向量。

5.1 目录结构

vector_add/
├── kernel/
│   └── vector_add_kernel.cpp
├── host/
│   └── vector_add_host.cpp
├── CMakeLists.txt
└── run.py

---

5.2 Kernel 侧代码（vector_add_kernel.cpp）

// vector_add_kernel.cpp
#include "kernel_operator.h"

using namespace AscendC;

constexpr int32_t BLOCK_SIZE = 256; // 每个 block 处理 256 个元素
constexpr int32_t TILE_NUM = 8;     // 流水线级数

class VectorAdd {
public:
    __aicore__ inline void Init(GM_ADDR inputA, GM_ADDR inputB, GM_ADDR output, uint32_t totalLength) {
        this->inputA = inputA;
        this->inputB = inputB;
        this->output = output;
        this->totalLength = totalLength;

        // 初始化 pipe
        pipe.InitBuffer(inQueueA, TILE_NUM, BLOCK_SIZE * sizeof(float));
        pipe.InitBuffer(inQueueB, TILE_NUM, BLOCK_SIZE * sizeof(float));
        pipe.InitBuffer(outQueue, TILE_NUM, BLOCK_SIZE * sizeof(float));
    }

    __aicore__ inline void Process() {
        int32_t loopCount = (totalLength + BLOCK_SIZE - 1) / BLOCK_SIZE;
        for (int32_t i = 0; i < loopCount; i++) {
            CopyIn(i);
            Compute(i);
            CopyOut(i);
        }
    }

private:
    __aicore__ inline void CopyIn(int32_t loopIndex) {
        LocalTensor<float> localA = inQueueA.AllocTensor<float>();
        LocalTensor<float> localB = inQueueB.AllocTensor<float>();

        DataCopy(localA, inputA[loopIndex * BLOCK_SIZE], BLOCK_SIZE);
        DataCopy(localB, inputB[loopIndex * BLOCK_SIZE], BLOCK_SIZE);

        inQueueA.EnQue(localA);
        inQueueB.EnQue(localB);
    }

    __aicore__ inline void Compute(int32_t loopIndex) {
        LocalTensor<float> localA = inQueueA.DeQue<float>();
        LocalTensor<float> localB = inQueueB.DeQue<float>();
        LocalTensor<float> localC = outQueue.AllocTensor<float>();

        // 向量化加法
        Add(localC, localA, localB, BLOCK_SIZE);

        outQueue.EnQue(localC);
    }

    __aicore__ inline void CopyOut(int32_t loopIndex) {
        LocalTensor<float> localC = outQueue.DeQue<float>();
        DataCopy(output[loopIndex * BLOCK_SIZE], localC, BLOCK_SIZE);
    }

private:
    GlobalTensor<float> inputA;
    GlobalTensor<float> inputB;
    GlobalTensor<float> output;
    uint32_t totalLength;

    TPipe pipe;
    TQue<QuePosition::VECIN, TILE_NUM> inQueueA;
    TQue<QuePosition::VECIN, TILE_NUM> inQueueB;
    TQue<QuePosition::VECOUT, TILE_NUM> outQueue;
};

extern "C" __global__ __aicore__ void vector_add_kernel(
    Gm<float> inputA,
    Gm<float> inputB,
    Gm<float> output,
    uint32_t totalLength) {
    
    VectorAdd op;
    op.Init(inputA, inputB, output, totalLength);
    op.Process();
}

代码解析：

• __aicore__：标记函数在 AI Core 上执行
• GM_ADDR：全局内存地址（Global Memory）
• TPipe 与 TQue：实现三级流水（CopyIn → Compute → CopyOut）
• DataCopy：高效数据搬运接口
• Add：内置向量化加法模板，自动使用 Vector Unit

---

5.3 Host 侧代码（vector_add_host.cpp）

// vector_add_host.cpp
#include <acl/acl.h>
#include <iostream>
#include <vector>

#define CHECK_ACL(call) do { \
    aclError err = call; \
    if (err != ACL_SUCCESS) { \
        std::cerr << "ACL error at " << #call << ": " << err << std::endl; \
        exit(1); \
    } \
} while(0)

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

    // 2. 分配 Host 内存
    const int N = 1024;
    std::vector<float> h_a(N, 1.0f);
    std::vector<float> h_b(N, 2.0f);
    std::vector<float> h_c(N, 0.0f);

    // 3. 分配 Device 内存
    void *d_a, *d_b, *d_c;
    size_t size = N * sizeof(float);
    CHECK_ACL(aclrtMalloc(&d_a, size, ACL_MEM_MALLOC_HUGE_FIRST));
    CHECK_ACL(aclrtMalloc(&d_b, size, ACL_MEM_MALLOC_HUGE_FIRST));
    CHECK_ACL(aclrtMalloc(&d_c, size, ACL_MEM_MALLOC_HUGE_FIRST));

    // 4. 数据拷贝到 Device
    CHECK_ACL(aclrtMemcpy(d_a, size, h_a.data(), size, ACL_MEMCPY_HOST_TO_DEVICE));
    CHECK_ACL(aclrtMemcpy(d_b, size, h_b.data(), size, ACL_MEMCPY_HOST_TO_DEVICE));

    // 5. 加载自定义算子
    aclopRegisterOpKernel("VectorAdd", "./vector_add_kernel.o");
    
    // 6. 构建 Op Desc
    aclTensorDesc *desc_a = aclCreateTensorDesc(ACL_FLOAT, 1, &N, ACL_FORMAT_ND);
    aclTensorDesc *desc_b = aclCreateTensorDesc(ACL_FLOAT, 1, &N, ACL_FORMAT_ND);
    aclTensorDesc *desc_c = aclCreateTensorDesc(ACL_FLOAT, 1, &N, ACL_FORMAT_ND);

    aclDataBuffer *buf_a = aclCreateDataBuffer(d_a, size);
    aclDataBuffer *buf_b = aclCreateDataBuffer(d_b, size);
    aclDataBuffer *buf_c = aclCreateDataBuffer(d_c, size);

    // 7. 执行算子
    aclrtStream stream;
    CHECK_ACL(aclrtCreateStream(&stream));
    CHECK_ACL(aclopCompileAndExecute(
        "VectorAdd",
        2, 
        desc_a, buf_a,
        desc_b, buf_b,
        desc_c, buf_c,
        ACL_ENGINE_SYS, ACL_COMPILE_SYS, stream
    ));
    CHECK_ACL(aclrtSynchronizeStream(stream));

    // 8. 拷回结果
    CHECK_ACL(aclrtMemcpy(h_c.data(), size, d_c, size, ACL_MEMCPY_DEVICE_TO_HOST));

    // 9. 验证结果
    bool success = true;
    for (int i = 0; i < N; i++) {
        if (std::abs(h_c[i] - 3.0f) > 1e-5) {
            success = false;
            break;
        }
    }
    std::cout << "Vector Add Result: " << (success ? "PASS" : "FAIL") << std::endl;

    // 10. 释放资源
    aclDestroyDataBuffer(buf_a); aclDestroyDataBuffer(buf_b); aclDestroyDataBuffer(buf_c);
    aclDestroyTensorDesc(desc_a); aclDestroyTensorDesc(desc_b); aclDestroyTensorDesc(desc_c);
    aclrtFree(d_a); aclrtFree(d_b); aclrtFree(d_c);
    aclrtDestroyStream(stream);
    aclrtDestroyContext(context);
    aclFinalize();

    return 0;
}

注意：实际部署中，通常使用 Python Host（通过 acl Python API）更便捷。此处用 C++ 展示底层流程。

---

5.4 CMakeLists.txt

cmake_minimum_required(VERSION 3.14)
project(vector_add)

set(CMAKE_CXX_STANDARD 14)

# Ascend C 编译器
set(ASCEND_C_COMPILER ascend-c)
set(KERNEL_FILE ${CMAKE_CURRENT_SOURCE_DIR}/kernel/vector_add_kernel.cpp)
set(KERNEL_OBJ ${CMAKE_CURRENT_BINARY_DIR}/vector_add_kernel.o)

# 编译 Kernel
add_custom_command(
    OUTPUT ${KERNEL_OBJ}
    COMMAND ${ASCEND_C_COMPILER} --host-os=linux --host-arch=x86_64 --device-arch=ascend910b -c ${KERNEL_FILE} -o ${KERNEL_OBJ}
    DEPENDS ${KERNEL_FILE}
)

# 编译 Host
find_package(PkgConfig REQUIRED)
pkg_check_modules(ACL REQUIRED libascendcl)

add_executable(vector_add_host host/vector_add_host.cpp)
target_include_directories(vector_add_host PRIVATE ${ACL_INCLUDE_DIRS})
target_link_libraries(vector_add_host ${ACL_LIBRARIES})

# 依赖 Kernel
add_custom_target(kernel ALL DEPENDS ${KERNEL_OBJ})

---

5.5 编译与运行

mkdir build && cd build
cmake ..
make

# 运行（需在昇腾设备或仿真环境）
./vector_add_host

预期输出：

Vector Add Result: PASS

---

六、性能优化技巧

1. 合理设置 Block Size

• 昇腾 910B 的 Vector Unit 支持 256-bit 向量（float32 下为 8 元素）
• 建议 Block Size 为 8 的倍数，如 256、512

2. 启用双缓冲（Double Buffering）

通过增加 TILE_NUM（如 2 或 4），使数据搬运与计算完全重叠。

3. 使用 Unified Buffer 减少 Global Memory 访问

对于中间结果，尽量复用 Local L1 Buffer，避免写回 Global Memory。

4. 避免分支发散

AI Core 为 SIMD 架构，条件分支会导致性能下降。尽量使用掩码（Mask）操作。

5. 对齐内存访问

确保 Global Memory 地址按 32 字节对齐，提升 DMA 效率。

---

七、调试与 Profiling

1. 使用 msadvisor 分析性能瓶颈

msadvisor -d ./vector_add_host -o report

2. 日志调试

在 Kernel 中使用 printf（仅仿真模式支持）：

printf("Debug: loop=%d, value=%f\n", i, localA[0]);

3. 精度验证

使用 Compare 工具对比 CPU 与 NPU 结果：

from compare import cosine_similarity
assert cosine_similarity(cpu_out, npu_out) > 0.999

---

八、常见问题（FAQ）

Q1：Ascend C 支持哪些数据类型？

• 支持：float16, float32, int8, int16, int32, uint8 等
• 不支持：double, complex 类型

Q2：能否调用第三方库（如 Eigen）？

• 不能。Kernel 侧只能使用 Ascend C 内置 API，禁止 STL、动态内存分配。

Q3：如何处理不规则形状（如 padding）？

• 使用 Mask 操作，例如：

  Add(localC, localA, localB, BLOCK_SIZE, mask);

Q4：与 CUDA 相比，学习曲线如何？

• 更陡峭。需理解昇腾特有的存储层次（L0A/L0B/L1/UB）和计算单元分工。

---

九、未来展望

随着 CANN 8.0 的演进，Ascend C 将支持：

• 自动并行化：跨 AI Core 的分布式算子
• 混合精度模板：FP16 + FP32 累加
• 图算融合：Kernel 内嵌控制流（if/for）
• JIT 编译：运行时生成优化 Kernel

此外，华为正推动 Ascend C 与主流框架（MindSpore、PyTorch via TorchNPU）深度集成，未来开发者可通过 @custom_op 装饰器直接注册 Ascend C 算子。

---

十、结语

Ascend C 是打通昇腾芯片性能“最后一公里”的关键工具。它虽有一定学习成本，但换来的是对硬件资源的极致掌控。在国产 AI 芯片崛起的今天，掌握 Ascend C 不仅是技术能力的体现，更是参与中国 AI 基础设施建设的重要一步。

希望本文能为你打开 Ascend C 的大门。代码已开源，欢迎在评论区交流实践心得！

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

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