训练营简介

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

目录

1 Ascend C概述：专为AI计算设计的编程语言

2 Ascend C语法扩展详解

2.1 函数执行空间限定符

2.2 地址空间限定符

2.3 核函数调用机制

3 Ascend C API体系结构

3.1 基础API分类与功能

3.2 高阶API的优势与应用场景

3.3 数据存储与Tensor对象

4 核函数开发详解

4.1 核函数定义与调用规则

4.2 编程范式与三级流水线

4.3 多核并行与数据切分

5 Ascend C算子开发实战：Add算子实现

5.1 算子分析与设计规格

5.2 核函数与算子类实现

5.3 编译运行与验证

6 总结

1 Ascend C概述：专为AI计算设计的编程语言

Ascend C是华为针对昇腾AI处理器专门设计的一款编程语言，它基于标准C++语法规范，并扩展了一组专用的编程类库API，使开发者能够高效利用昇腾AI处理器的强大计算能力。作为CANN（Compute Architecture for Neural Networks）异构计算架构的关键组成部分，Ascend C允许开发者编写在AI Core上执行的算子核函数，与在GPU上使用CUDA开发有相似之处。其设计目标是在保持C++编程熟悉度的同时，提供对硬件底层架构的高度抽象和直接控制，从而实现高性能AI计算的开发需求。

Ascend C的设计理念围绕几个核心原则：一是原生支持C/C++编程规范，降低开发者的学习成本；二是通过多层接口抽象，平衡开发效率与硬件控制能力；三是引入并行编程范式，充分发挥昇腾AI处理器的并行计算能力；四是提供CPU/NPU孪生调试技术，大幅提升算子的调试效率。这些设计原则使得Ascend C成为一款既强大又易用的AI处理器编程工具。

从软硬件基础角度来看，Ascend C构建在昇腾AI处理器和CANN软件栈之上。昇腾AI处理器采用达芬奇架构，包含多个AI Core，每个Core具备强大的标量、向量和矩阵计算能力。CANN作为软件栈，位于底层硬件与深度学习框架之间，起到承上启下的作用，通过软硬件协同优化，充分发挥昇腾AI处理器的算力潜力。Ascend C开发的算子通过编译器编译和运行时调度，最终运行在昇腾AI处理器上，完成各种AI计算任务。

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

2 Ascend C语法扩展详解

Ascend C的语法体系在标准C++基础上进行了有针对性的扩展，主要引入了函数执行空间限定符、地址空间限定符和内核调用符等新语法元素。这些扩展使开发者能够精确控制代码执行位置和数据存储位置，从而高效利用昇腾AI处理器的计算资源。

2.1 函数执行空间限定符

函数执行空间限定符是Ascend C中最基本的语法扩展之一，用于指示函数是在主机（Host）上执行还是在设备（Device）上执行，以及它可以被主机函数还是设备函数调用。Ascend C主要提供了三种函数执行空间限定符：

• __global__限定符用于声明核函数（Kernel Function），即Ascend C算子的设备侧入口函数。使用__global__修饰的函数具有以下特性：在设备上执行；只能被主机侧函数调用；必须返回void类型，且不能是类的成员函数；主机侧调用时必须使用<<<>>>异构调用语法；调用是异步的，意味着主机侧函数返回并不表示核函数在设备侧已经执行完成。如果需要同步，必须使用运行时提供的同步接口（如aclrtSynchronizeStream）进行显式同步。

• __aicore__限定符用于声明在AI Core设备上执行的函数。这类函数只能被__global__函数或其他AI Core函数调用。在核函数中，除了主处理逻辑外，开发者可以将一些通用功能封装成__aicore__函数，以提高代码复用性和可读性。

• __host__限定符用于声明在主机侧执行的函数（通常不显式声明）。这类函数只能在主机侧执行，也只能被主机侧函数调用。需要注意的是，__global__和__host__不能一起使用。

典型使用函数执行空间限定符的示例代码如下：

cpp

// 定义aicore函数
__aicore__ void bar() {}

// 定义核函数
__global__ __aicore__ void foo() { bar();}

// 定义Host函数
int main() {}

这种函数执行空间的明确划分使得Ascend C能够有效管理主机与设备之间的协作，充分利用异构计算环境的优势。

2.2 地址空间限定符

地址空间限定符在变量声明中使用，用于指定对象分配的区域。AI Core具备多级独立片上存储，各个地址空间独立编址，具备各自的访存指令。如果对象的类型被地址空间名称限定，那么该对象将被分配在指定的地址空间中。

Ascend C中最重要的地址空间限定符是__gm__，它用来表示分配于设备侧全局内存（Global Memory）的对象。全局内存是AI Core外部的存储空间，容量较大但访问速度相对较慢，通常用于存储大量的输入数据和输出结果。在核函数中，指向全局内存的指针参数需要使用__gm__限定符修饰，为了统一表达，Ascend C提供了GM_ADDR宏来简化这一过程，定义如下：#define GM_ADDR __gm__ uint8_t*。

除了__gm__外，Ascend C还有private地址空间，这是大多数变量的默认地址空间，特别是局部变量。代码中不显示标识所在局部地址空间类型。private地址空间对应于AI Core的局部内存（Local Memory），访问速度快但容量较小，适合存储中间计算结果和频繁访问的数据。

2.3 核函数调用机制

核函数调用是Ascend C编程中主机侧与设备侧交互的关键环节。与常规C/C++函数调用不同，核函数使用特殊的<<<...>>>内核调用符语法，这种语法形式规定了核函数的执行配置。

内核调用符包含3个关键参数：

• blockDim：表示调用的核数，即并行执行核函数的AI Core数量。

• l2ctrl：用于L2缓存控制的参数，通常可设置为nullptr。

• stream：流水句柄，用于管理异步操作和同步。

以下是一个名为add_custom的核函数调用示例：

cpp

// blockDim设置为8表示在8个核上调用了add_custom核函数，每个核都会独立且并行地执行该核函数
add_custom<<<8, nullptr, stream>>>(x, y, z);

需要注意的是，核函数的调用是异步的，调用后控制权立刻返回给主机端，不会阻塞主机端程序的继续执行。这种机制使得主机端可以在调用核函数后继续执行其他计算任务，从而实现主机与设备的并行工作。如果主机端需要等待核函数执行完毕，可以调用aclrtSynchronizeStream函数来强制主机端程序等待所有核函数执行完毕。

3 Ascend C API体系结构

Ascend C提供了一套丰富的编程类库API，使开发者能够高效利用昇腾AI处理器的计算能力。这些API根据抽象层次和功能特点，可以分为基础API和高阶API两大类。开发者可以根据自己的需求选择合适的API，在保证功能完备性和兼容性的同时，提高开发效率。

3.1 基础API分类与功能

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

• 计算类API：包括标量计算API、向量计算API、矩阵计算API，分别实现调用标量计算单元、向量计算单元、矩阵计算单元执行计算的功能。根据对数据操作方法的不同，计算API又分为整个Tensor参与计算、Tensor前n个数据计算和Tensor高维切分计算三种方式。整个Tensor参与计算通过运算符重载的方式实现，支持+, -, *, /, |, &, <, >, <=, >=, ==, !=等操作，实现计算的简化表达，例如：dst=src1+src2；Tensor前n个数据计算针对源操作数的连续n个数据进行计算并连续写入目的操作数，解决一维Tensor的连续计算问题，例如：Add(dst, src1, src2, n)；Tensor高维切分计算则是功能最灵活的计算API，充分发挥硬件优势，支持对每个操作数的Repeat times（迭代的次数）、Block stride（单次迭代内不同block间地址步长）、Repeat stride（相邻迭代间相同block的地址步长）、Mask（用于控制参与运算的计算单元）等参数的操作。

• 数据搬运API：由于计算API基于本地内存（Local Memory）数据进行计算，所以数据需要先从全局内存（Global Memory）搬运至本地内存，再使用计算接口完成计算，最后从本地内存搬出至全局内存。执行搬运过程的接口称之为数据搬运接口，比如DataCopy接口。高效的数据搬运是提升整体计算性能的关键，特别是在处理大规模数据时。

• 内存管理API：用于分配板上管理内存，比如AllocTensor、FreeTensor接口。由于板上内存较小，通常无法存储完整数据，因此采用动态内存的方式进行内存管理，实现板上内存的复用。这些API帮助开发者有效管理有限的存储资源，确保内存使用的效率和安全。

• 任务同步API：完成任务间的通信和同步，比如EnQue、DeQue接口。不同的API指令间有可能存在依赖关系，而不同的指令异步并行执行，为了保证不同指令队列间的指令按照正确的逻辑关系执行，需要向不同的组件发送同步指令。任务同步类API内部即完成这个发送同步指令的过程，开发者无需关注内部实现逻辑，使用简单的API接口即可完成。

3.2 高阶API的优势与应用场景

高阶API封装常用算法逻辑，通常会调用多种基础API实现常用的计算算法，用于提高开发效率。使用高阶API可以快速的实现相对复杂的算法逻辑，高阶API是对于某种特定算法的表达。

以矩阵乘法（Matmul）为例，使用高阶API完成Matmul算子时，需要创建一个矩阵乘法类进行运算，其中入参包含两个相乘的矩阵（一般称为A矩阵与B矩阵）信息、输出结果矩阵（一般称为C矩阵）信息、矩阵乘偏置（一般称为Bias）信息。这些信息中包括了对应矩阵数据的内存逻辑位置、数据存储格式、数据类型、转置使能等参数。

创建完这样的一个矩阵乘法类后，使用Ascend C高阶API可以直接完成对左右矩阵A、B和Bias的设置和矩阵乘法操作以及结果的输出，开发者不用再自主实现复杂的数据通路和运算操作。这大大简化了复杂算子的实现难度，提高了开发效率，同时由于高阶API经过深度优化，通常也能提供较好的性能。

高阶API特别适用于实现常见的深度学习算子，如卷积、池化、归一化等。对于这些标准操作，使用高阶API可以显著减少代码量，提高代码可读性和可维护性。然而，对于研究性的新算法或特殊计算模式，如果高阶API无法直接满足需求，开发者仍可能需要使用基础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）的全局数据，其原型定义如下：

cpp

template <typename T> class GlobalTensor {
    void SetGlobalBuffer(__gm__ T* buffer, uint32_t bufferSize);
    // ... 其他成员函数
};

开发者通过GlobalTensor对象管理在全局内存中的数据，包括设置缓冲区地址和大小等操作。

LocalTensor则代表在AI Core内部存储中的数据，用于各种计算操作。由于Local Memory容量有限，通常需要将Global Memory中的数据分块加载到Local Memory中进行处理，然后再将结果写回Global Memory。这种数据流动模式是Ascend C编程的基本范式。

Ascend C的API操作数都是Tensor类型：GlobalTensor（外部数据存储空间）和LocalTensor（核上内存空间）。这种统一的数据抽象简化了API设计，使开发者能够以一致的方式处理不同存储位置的数据，同时也有利于代码的清晰性和可维护性。

4 核函数开发详解

核函数（Kernel Function）是Ascend C算子设备侧实现的入口，它是在AI Core上执行的代码单元。与普通的C++函数调用时仅执行一次不同，当核函数被调用时，多个核都执行相同的核函数代码，具有相同的函数入参，并行执行。这种单程序多数据（SPMD）的并行模式是Ascend C并行计算的基础。

4.1 核函数定义与调用规则

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

其他重要规则包括：核函数必须具有void返回类型；仅支持入参为指针或C/C++内置数据类型（Primitive data types），如：half* s0、float* s1、int32_t c。这些规则确保了核函数能够正确地在AI Core上执行并与主机端进行交互。

以下是一个完整的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);
}

在这个示例中，核函数add_custom调用算子类KernelAdd的Init和Process方法，分别完成初始化处理和核心逻辑执行。主机侧函数add_custom_do使用<<<>>>调用符异步调用核函数。

4.2 编程范式与三级流水线

Ascend C引入了结构化的编程范式，将算子核函数的实现程序分解为更小的、易于理解和管理的部分。以矢量编程范式为例，开发者基于"搬入、计算、搬出"三段式结构进行编程。这种编程方式帮助开发者搭建清晰的编程框架，使开发者可以聚焦算子的实现逻辑，极大提高编程效率。

在三级流水线编程范式中，每个阶段都有明确的功能分工：

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

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

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

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

4.3 多核并行与数据切分

Ascend C支持多核并行计算，即把数据进行分片，分配到多个核上进行处理。这种并行计算模式能够充分利用昇腾AI处理器的多个计算核心，显著提高计算效率。

多核并行处理的关键在于数据切分。假设共启用8个核，数据整体长度TOTAL_LENGTH为8 * 2048个元素，平均分配到8个核上运行，每个核上处理的数据大小BLOCK_LENGTH为2048个元素。每个核上处理的数据地址需要在起始地址上增加GetBlockIdx() * BLOCK_LENGTH（每个block处理的数据长度）的偏移来获取。这样也就实现了多核并行计算的数据切分。

以输入x为例，x + BLOCK_LENGTH * GetBlockIdx()即为单核处理程序中x在Global Memory上的内存偏移地址，获取偏移地址后，使用GlobalTensor类的SetGlobalBuffer接口设定该核上Global Memory的起始地址以及长度。具体实现代码如下：

cpp

xGm.SetGlobalBuffer((__gm__ half*)x + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);
yGm.SetGlobalBuffer((__gm__ half*)y + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);
zGm.SetGlobalBuffer((__gm__ half*)z + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);

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

这种多层次的数据切分策略（多核间切分和单核内切分）使得Ascend C能够高效利用并行计算资源，充分发挥硬件性能，处理大规模数据计算任务。

5 Ascend C算子开发实战：Add算子实现

为了全面理解Ascend C算子开发流程，我们将通过一个具体的Add算子实例，详细讲解从算子分析到核函数实现的完整过程。Add算子作为最简单的矢量运算之一，适合用于展示Ascend C开发的基本原理和流程。

5.1 算子分析与设计规格

在开始编码前，首先需要进行详细的算子分析，明确算子的数学表达式、输入输出的数量、Shape范围以及计算逻辑的实现，确定需要调用的Ascend C接口。对于Add算子，分析过程如下：

首先，明确算子的数学表达式及计算逻辑。Add算子的数学表达式为：$z = x + y$。计算逻辑是：从外部存储Global Memory搬运数据至内部存储Local Memory，然后使用Ascend C计算接口完成两个输入参数相加，得到最终结果，再搬运到Global Memory上。

其次，明确输入和输出。Add算子有两个输入：x与y，输出为z。在本样例中，算子输入支持的数据类型为half（float16），算子输出的数据类型与输入数据类型相同；算子输入支持的shape为（8，2048），输出shape与输入shape相同；算子输入支持的format为：ND。

然后，确定核函数名称和参数。本样例中核函数命名为add_custom；根据对算子输入输出的分析，确定核函数有3个参数x，y，z；x，y为输入参数，z为输出参数。

最后，确定算子实现所需接口。实现涉及外部存储和内部存储间的数据搬运，需要使用DataCopy来实现数据搬移；本样例只涉及矢量计算的加法操作，可使用双目指令Add接口实现x+y；计算中使用到的Tensor数据结构，使用AllocTensor、FreeTensor进行申请和释放；并行流水任务之间使用Queue队列完成同步，会使用到EnQue、DeQue等接口。

通过以上分析，得到Ascend C Add算子的设计规格如下：

算子类型（OpType）	AddCustom
算子输入
name	x
shape	(8, 2048)
data type	half
format	ND
name	y
shape	(8, 2048)
data type	half
format	ND
算子输出
name	z
shape	(8, 2048)
data type	half
format	ND
核函数名称	add_custom
使用的主要接口	DataCopy、Add、AllocTensor、FreeTensor、EnQue、DeQue等
算子实现文件名称	add_custom.cpp

5.2 核函数与算子类实现

基于上述分析，我们可以开始实现Add算子的核函数和算子类。首先实现核函数，它在设备侧作为算子的入口点：

cpp

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();
}

这段代码遵循了核函数的基本规则：使用__global__和__aicore__限定符，参数使用GM_ADDR宏修饰，返回类型为void。在核函数内部，创建了KernelAdd算子类实例，并依次调用其Init和Process方法。

接下来，我们需要实现KernelAdd算子类，这是算子的核心逻辑所在：

cpp

class KernelAdd {
public:
    __aicore__ inline KernelAdd(){}
    // 初始化函数，完成内存初始化相关操作
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z)
    {
        // 设置输入输出Global Tensor的Global Memory内存地址
        xGm.SetGlobalBuffer((__gm__ half*)x + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);
        yGm.SetGlobalBuffer((__gm__ half*)y + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);
        zGm.SetGlobalBuffer((__gm__ half*)z + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);
        
        // 通过Pipe内存管理对象为输入输出Queue分配内存
        pipe.InitBuffer(inQueueX, BUFFER_NUM, TILE_LENGTH * sizeof(half));
        pipe.InitBuffer(inQueueY, BUFFER_NUM, TILE_LENGTH * sizeof(half));
        pipe.InitBuffer(outQueueZ, BUFFER_NUM, TILE_LENGTH * sizeof(half));
    }
    
    // 核心处理函数，实现算子逻辑
    __aicore__ inline void Process()
    {
        // 定义三轮循环，分别处理搬入、计算、搬出三个流水阶段
        for (int32_t i = 0; i < TILE_NUM * BUFFER_NUM; i++) {
            CopyIn(i);
            Compute(i);
            CopyOut(i);
        }
    }

private:
    // 搬入函数
    __aicore__ inline void CopyIn(int32_t progress)
    {
        // 从Global Memory搬运数据到Local Memory
        DataCopy(inQueueX.AllocTensor(), xGm[progress * TILE_LENGTH], TILE_LENGTH);
        DataCopy(inQueueY.AllocTensor(), yGm[progress * TILE_LENGTH], TILE_LENGTH);
        
        // 将输入Tensor入队
        inQueueX.EnQue(inQueueX.AllocTensor());
        inQueueY.EnQue(inQueueY.AllocTensor());
    }
    
    // 计算函数
    __aicore__ inline void Compute(int32_t progress)
    {
        // 从输入Queue中取出Tensor
        LocalTensor<half> xLocal = inQueueX.DeQue<half>();
        LocalTensor<half> yLocal = inQueueY.DeQue<half>();
        LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>();
        
        // 执行矢量加法计算
        Add(zLocal, xLocal, yLocal, TILE_LENGTH);
        
        // 将输入Tensor回收，输出Tensor入队
        inQueueX.FreeTensor(xLocal);
        inQueueY.FreeTensor(yLocal);
        outQueueZ.EnQue<half>(zLocal);
    }
    
    // 搬出函数
    __aicore__ inline void CopyOut(int32_t progress)
    {
        // 从输出Queue中取出Tensor
        LocalTensor<half> zLocal = outQueueZ.DeQue<half>();
        
        // 将结果从Local Memory搬运到Global Memory
        DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH);
        
        // 将输出Tensor回收
        outQueueZ.FreeTensor(zLocal);
    }

private:
    TPipe pipe;
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueZ;
    GlobalTensor<half> xGm, yGm, zGm;
};

上述代码完整实现了KernelAdd类，包括Init初始化方法和CopyIn、Compute、CopyOut三个核心私有方法。Init方法中，通过GetBlockIdx()函数获取当前核的索引，从而计算数据偏移，实现多核间的数据切分。Process方法中，通过循环依次调用CopyIn、Compute和CopyOut，实现三级流水处理。

5.3 编译运行与验证

完成核函数开发后，需要进行编译和验证，确保算子的正确性和性能。Ascend C支持CPU模式运行验证和NPU模式运行验证两种方式。

对于CPU模式运行验证，可以使用ICPU_RUN_KF CPU调测宏完成算子核函数CPU侧运行验证。这种模式下，算子可以在CPU上运行，便于调试和验证逻辑正确性，但无法完全模拟NPU的硬件特性。

对于NPU模式运行验证，使用<<<>>>内核调用符完成算子核函数NPU侧运行验证。这种模式下，算子实际在NPU上运行，能够真实反映算子在硬件上的性能和行为。

开发Ascend C算子的基本流程包括：环境准备、算子分析、核函数开发、CPU模式运行验证和NPU模式运行验证。通过这个完整的流程，开发者可以系统地完成Ascend C算子的开发与部署。

6 总结

Ascend C作为昇腾AI处理器的专用编程语言，通过扩展标准C++语法和提供多层次API，有效平衡了开发效率与硬件控制能力。其核心特性包括结构化编程范式、多层次API体系、CPU/NPU孪生调试等，使开发者能够高效开发高性能AI算子。

核函数开发是Ascend C编程的关键，需要遵循特定的函数限定符和调用规则。通过"搬入-计算-搬出"的三级流水线和多核并行数据切分策略，可以充分发挥昇腾AI处理器的并行计算能力。随着AI技术的不断发展，Ascend C将在更广泛的AI应用场景中发挥重要作用，为AI计算提供强大的底层支持。

对于有志于深入掌握Ascend C的开发者，建议从官方文档和样例代码入手，理解编程模型和API用法，然后通过实际算子开发项目积累经验。同时，关注性能调优技巧和最佳实践，不断提升算子效率和质量，为构建高性能AI应用奠定坚实基础。