CANN训练营学习（day2）Ascend C编程语言解析：从入门到高性能算子开发

摘要：AscendC是CANN专为昇腾AI处理器设计的算子开发语言，兼容C/C++标准，提供多层接口抽象和自动并行计算。其核心是核函数编程模型，采用"搬入-计算-搬出"三级流水线，支持多核并行。AscendC提供基础API、高阶API和硬件加速API，并基于Global/Local Tensor实现高效内存管理。开发流程包括环境准备、核函数实现和多核优化，配备CPU孪生调试工具

weixin_43260261

1332人浏览 · 2025-11-05 14:49:31

weixin_43260261 · 2025-11-05 14:49:31 发布

训练营简介

报名链接https://www.hiascend.com/developer/activities/cann20252#cann-camp-2502-intro

1 Ascend C概述与设计理念

Ascend C是CANN（Compute Architecture for Neural Networks）针对算子开发场景推出的专用编程语言，它原生支持C和C++标准规范，通过多层接口抽象、自动并行计算、孪生调试等关键技术，极大提高了算子开发效率，助力AI开发者在昇腾AI处理器上高效实现自定义的创新算法。作为昇腾AI生态系统的重要组成部分，Ascend C充当了连接上层AI框架与底层硬件算力的关键桥梁，使开发者能够充分利用昇腾AI处理器的强大计算能力。

Ascend C的设计理念围绕开发效率与运行性能的双重目标展开。一方面，它最大化匹配用户的现有开发习惯，遵循C/C++编程规范，降低了学习成本；另一方面，它通过结构化编程范式和自动并行调度，确保算子程序能够获得接近硬件极限的执行性能。这种平衡设计使得无论是算法研究人员还是底层系统优化工程师，都能基于Ascend C高效实现其创新想法。

从硬件架构视角看，Ascend C专门针对昇腾AI处理器的达芬奇架构进行了深度优化。昇腾AI处理器包含多个AI Core，每个Core都具备强大的标量、向量和矩阵计算能力。Ascend C通过编程模型抽象了这些硬件的复杂性，使开发者无需关注底层硬件的具体细节，却能通过编译器优化和运行时调度，充分发挥硬件性能。基于Ascend C编写的算子程序，通过编译器编译和运行时调度，最终运行在昇腾AI处理器上，完成各种AI计算任务。

与通用C++相比，Ascend C具有几个显著特点：它提供了专门的内存管理机制，区分全局内存和局部内存，适应AI处理器的存储层次结构；引入了特殊的函数限定符，明确指定函数在主机侧还是设备侧执行；提供了丰富的数据搬运和计算API，简化了并行编程的复杂性；采用结构化的编程范式，将算子实现分解为更小、更易管理的部分。这些特性使Ascend C特别适合开发高性能的AI算子。

2 核函数：Ascend C编程的核心

2.1 核函数的概念与编程模型

在Ascend C中，核函数（Kernel Function）是算子设备侧实现的入口，也是性能优化的核心载体。与传统的C++函数调用时仅执行一次不同，当核函数被调用时，多个核都执行相同的核函数代码，具有相同的函数入参，并行执行。这种单程序多数据（SPMD）的并行模式是Ascend C并行计算的基础，允许开发者编写单一的代码逻辑，却在多个AI Core上并行处理不同数据片段，极大提高了计算效率。

核函数的编程模型采用了独特的矢量编程范式，将算子执行拆解为"搬入、计算、搬出"三级流水任务。这种结构化编程方式帮助开发者搭建清晰的编程框架，使开发者可以聚焦算子的实现逻辑，而不必过度关注底层硬件细节。配合Queue（任务队列）、TPosition（逻辑位置抽象）、LocalTensor/GlobalTensor（内存管理）等组件，实现硬件资源的高效调度。

2.2 核函数的定义与调用规则

定义核函数时需要遵循一系列特定规则。首先，必须使用函数类型限定符__global__和__aicore__。__global__表示这是一个核函数，可以被<<<>>>调用；__aicore__表示该核函数在设备端AI Core上执行。其次，指针入参变量需要增加变量类型限定符__gm__，表明该指针变量指向Global Memory上某处内存地址。为了统一表达，建议使用GM_ADDR宏来修饰入参。

以下是一个完整的Add算子核函数示例：

cpp

// 实现核函数
extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* 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);
}

核函数使用特殊的<<<...>>>内核调用符语法，这种语法形式规定了核函数的执行配置。内核调用符包含3个关键参数：blockDim（调用的核数）、l2ctrl（用于L2缓存控制的参数）和stream（流水句柄）。需要注意的是，核函数的调用是异步的，调用后控制权立刻返回给主机端，不会阻塞主机端程序的继续执行。这种机制使得主机端可以在调用核函数后继续执行其他计算任务，从而实现主机与设备的并行工作。

2.3 三级流水线编程范式

Ascend C引入了结构化的编程范式，将算子核函数的实现程序分解为三个明确的阶段：CopyIn（搬入）、Compute（计算） 和CopyOut（搬出）。这种"搬入、计算、搬出"的三级流水线是Ascend C矢量编程范式的核心，它帮助开发者搭建清晰的编程框架，使开发者可以聚焦算子的实现逻辑，极大提高编程效率。

在CopyIn阶段，负责从Global Memory搬运数据至Local Memory。这一阶段使用数据搬运API将需要处理的数据从外部存储加载到内部存储，为计算阶段做准备。通过合理的数据分块和流水线技术，可以隐藏数据搬运的延迟，提高整体计算效率。

在Compute阶段，在Local Memory数据上执行各种计算操作，如矢量运算、矩阵乘法等。这一阶段充分利用AI Core的强大计算能力，对已加载到局部内存的数据进行加工处理，生成中间结果或最终结果。

在CopyOut阶段，负责将最终计算结果从Local Memory搬运到Global Memory上。计算完成后，需要将结果写回到外部存储中，以便主机端访问或其他算子使用。

这种三级流水线不仅使代码结构清晰，而且通过流水并行技术，可以显著提高计算资源的利用率。当一部分数据正在计算时，下一部分数据可以同时进行搬入操作，上一部分计算结果可以同时进行搬出操作，从而实现数据搬运和计算的并行执行。

3 Ascend C的多层级API设计

3.1 基础API分类与功能

Ascend C提供了一套丰富的基础API，实现对硬件能力的抽象，开放芯片的基础能力，保证完备性和兼容性。基础API是构建算子的基础模块，提供了对硬件资源的直接控制。基础API主要包括以下几类：

计算类API：包括标量计算API、向量计算API、矩阵计算API，分别实现调用标量计算单元、向量计算单元、矩阵计算单元执行计算的功能。根据对数据操作方法的不同，计算API又分为整个Tensor参与计算、Tensor前n个数据计算和Tensor高维切分计算三种方式。
数据搬运API：由于计算API基于本地内存（Local Memory）数据进行计算，所以数据需要先从全局内存（Global Memory）搬运至本地内存，再使用计算接口完成计算，最后从本地内存搬出至全局内存。执行搬运过程的接口称之为数据搬运接口，比如DataCopy接口。高效的数据搬运是提升整体计算性能的关键，特别是在处理大规模数据时。
内存管理API：用于分配板上管理内存，比如AllocTensor、FreeTensor接口。由于板上内存较小，通常无法存储完整数据，因此采用动态内存的方式进行内存管理，实现板上内存的复用。这些API帮助开发者有效管理有限的存储资源，确保内存使用的效率和安全。
任务同步API：完成任务间的通信和同步，比如EnQue、DeQue接口。不同的API指令间有可能存在依赖关系，而不同的指令异步并行执行，为了保证不同指令队列间的指令按照正确的逻辑关系执行，需要向不同的组件发送同步指令。

3.2 高阶API与硬件加速API

高阶API封装常用算法逻辑，通常会调用多种基础API实现常用的计算算法，用于提高开发效率。使用高阶API可以快速的实现相对复杂的算法逻辑，高阶API是对于某种特定算法的表达。以矩阵乘法（Matmul）为例，使用高阶API完成Matmul算子时，需要创建一个矩阵乘法类进行运算，其中入参包含两个相乘的矩阵信息、输出结果矩阵信息、矩阵乘偏置信息。

高阶API特别适用于实现常见的深度学习算子，如卷积、池化、归一化等。对于这些标准操作，使用高阶API可以显著减少代码量，提高代码可读性和可维护性。然而，对于研究性的新算法或特殊计算模式，如果高阶API无法直接满足需求，开发者仍可能需要使用基础API进行自定义实现。

硬件加速API直接对接AI Core硬件特性，如__aicore__（核函数硬件标识）、Pipe（任务管道调度）等。这些API允许开发者更直接地控制硬件行为，实现极致的性能优化，但通常需要开发者对硬件架构有更深入的了解。

3.3 数据存储与Tensor对象

根据Ascend C对于AI Core的硬件抽象设计，AI Core内部的存储统一用Local Memory来表示，AI Core外部的存储统一用Global Memory来表示。Ascend C使用GlobalTensor作为Global Memory的数据基本操作单元，与之对应的，用LocalTensor作为Local Memory的数据基本操作单元。数据的基本操作单元（Tensor，张量）是各种指令API直接处理的对象，也是数据的载体。

GlobalTensor用来存放AI Core外部存储（Global Memory）的全局数据，开发者通过GlobalTensor对象管理在全局内存中的数据，包括设置缓冲区地址和大小等操作。LocalTensor则代表在AI Core内部存储中的数据，用于各种计算操作。由于Local Memory容量有限，通常需要将Global Memory中的数据分块加载到Local Memory中进行处理，然后再将结果写回Global Memory。这种数据流动模式是Ascend C编程的基本范式。

4 Ascend C算子开发工作流

4.1 环境准备与工程创建

开发Ascend C算子的第一步是完成环境准备工作，包括安装CANN软件包、配置开发环境变量以及准备编译和调试工具。环境准备完成后，开发者可以开始创建自定义算子工程。以实现MulAdd算子（功能：z = x * alpha + y，alpha为float类型属性，数据类型为half）为例，完整演示工程创建与配置步骤。

首先需要创建算子原型定义文件（JSON文件），明确输入、输出、属性及硬件实现方式：

json

{
    "op": "MulAdd",
    "input": [
        {"name": "x", "dtype": ["half"], "format": ["ND"]},
        {"name": "y", "dtype": ["half"], "format": ["ND"]}
    ],
    "output": [
        {"name": "z", "dtype": ["half"], "format": ["ND"]}
    ],
    "attr": [
        {"name": "alpha", "dtype": "float", "default_value": 0.5}
    ],
    "op_impl": {
        "ai_core": {
            "kernel": "mul_add_custom",
            "enable_tiling": true
        }
    }
}

通过msOpGen工具生成工程，需先加载Ascend C环境变量，然后执行工程生成命令：

bash

# 加载Ascend C环境变量
source ddk/tools/tools_ascendc/set_ascendc_env.sh
# 生成算子工程
msOpGen -i /path/to/mul_add_custom.json -c ai_core-kirin9020 -out /path/to/MulAdd_Project

生成的工程包含两个核心文件：op_kernel/mul_add_custom.cpp（实现核函数的CopyIn/Compute/CopyOut三级流水逻辑）和op_host/mul_add_custom_tiling.h（定义Tiling数据结构与分片参数计算函数，用于优化数据分片粒度）。

4.2 核函数实现与多核并行

在生成的工程文件中，开发者需要实现核函数的具体逻辑。以下以MulAdd算子的Compute函数为例：

cpp

__aicore__ inline void KernelMulAdd::Compute(int32_t progress) {
    // 从输入队列获取LocalTensor（片上内存数据）
    AscendC::LocalTensor<half> xLocal = inQueueX.DeQue<half>();
    AscendC::LocalTensor<half> yLocal = inQueueY.DeQue<half>();
    // 分配临时与输出LocalTensor
    AscendC::LocalTensor<half> tempLocal = outQueueZ.AllocTensor<half>();
    AscendC::LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>();
    // 执行x*alpha + y的计算（调用Ascend C高阶API）
    half alphaHalf = static_cast<half>(this->alpha);
    AscendC::Mul(tempLocal, xLocal, alphaHalf, this->tileLength);
    AscendC::Add(zLocal, tempLocal, yLocal, this->tileLength);
    // 数据入队与内存释放
    outQueueZ.EnQue<half>(zLocal);
    outQueueZ.FreeTensor(tempLocal);
    inQueueX.FreeTensor(xLocal);
    inQueueY.FreeTensor(yLocal);
}

Ascend C支持多核并行计算，即把数据进行分片，分配到多个核上进行处理。这种并行计算模式能够充分利用昇腾AI处理器的多个计算核心，显著提高计算效率。假设共启用8个核，数据整体长度TOTAL_LENGTH为8 * 2048个元素，平均分配到8个核上运行，每个核上处理的数据大小BLOCK_LENGTH为2048个元素。每个核上处理的数据地址需要在起始地址上增加GetBlockIdx() * BLOCK_LENGTH（每个block处理的数据长度）的偏移来获取。这样也就实现了多核并行计算的数据切分。

在单核内部，数据可以进一步进行切块处理（Tiling）。例如，将单核上的数据（2048个元素）切分成8块（并不意味着8块就是性能最优）。切分后的每个数据块再次切分成2块，即可开启double buffer，实现流水线之间的并行。这样单核上的数据被切分成16块，每块TILE_LENGTH（128）个数据。Pipe为队列分配了两块大小为TILE_LENGTH * sizeof(half)个字节的内存块，每个内存块能容纳TILE_LENGTH（128）个half类型数据。

4.3 调试与调优

Ascend C提供了强大的调试功能，其中最突出的是CPU孪生调试技术，它允许开发者在CPU侧模拟NPU侧的行为，可优先在CPU侧调试。使用ascendebug工具进行CPU孪生调试，命令如下：

bash

ascendebug kernel --backend cpu --chip-version kirin9020 --json-file /path/to/mul_add_custom.json --work-dir /path/to/debug_dir

精度比对配置：在调试配置文件中指定golden_data_path，指向预处理的标杆数据文件（bin格式），工具会自动比对输出结果与标杆数据的误差。若结果存在偏差，可通过以下手段排查：

打印调试：通过printf输出标量参数（如alpha值），验证参数传递逻辑；
张量Dump：调用DumpTensor(zLocal)保存输出张量数据，与标杆数据逐元素比对；
断言检查：用assert(xLocal.GetLength() == this->tileLength)验证数据分片逻辑的正确性。

除了功能正确性调试，性能调优也是Ascend C算子开发的关键环节。通过Profiling工具，开发者可以收集算子的性能数据，如内存带宽利用率（aic_mte2_ratio）、执行时间等指标，从而识别性能瓶颈并实施针对性的优化措施。

5 Ascend C性能优化方法与工具

5.1 性能分析工具与优化流程

Ascend C提供了丰富的性能分析工具，帮助开发者识别性能瓶颈并实施有效的优化措施。性能优化是一个持续迭代的流程，通常包括四个步骤：性能测试、性能分析、性能优化和性能回测。通过这个循环过程，开发者可以逐步提升算子性能，直至达成性能目标。

在进行性能分析时，开发者需要关注几个关键指标：内存带宽利用率（如aic_mte2_ratio）、计算单元利用率、执行时间等。这些指标可以帮助开发者判断算子是受内存带宽限制还是受计算能力限制，从而采取不同的优化策略。例如，如果内存带宽利用率过高，说明内存访问已成为性能瓶颈，需要优化数据访问模式或减少不必要的数据搬运。

5.2 性能优化关键技术

5.2.1 Tiling优化技术

Tiling是一种数据分片优化技术，将大张量切分为匹配AI Core计算能力的小分片。合理的Tiling策略可以显著提高数据局部性，减少内存访问延迟。在Matmul算子优化案例中，通过优化基本块参数（从baseM=64/baseN=64优化到baseM=128/baseN=256），执行时间从2350us降低到810us，性能提升约2.9倍。

Tiling优化的关键是找到适合硬件特性和算法特征的分块大小，这需要综合考虑AI Core的存储层次结构、计算单元数量以及数据访问模式。通常情况下，较大的分块可以提高计算密度，但可能增加内存访问压力；较小的分块则有利于数据复用，但可能无法充分利用并行计算资源。