在NPU硬件中，AI Core是AI处理器的核心计算载体，是支撑各类AI算子高效运算的基础单元。主流AI处理器均集成多个AI Core，依托成熟的并行计算架构实现海量数据的并行处理。基于昇腾架构的Ascend C编程体系，对AI Core硬件架构进行了高效抽象，屏蔽了不同硬件型号的底层差异，大幅简化了开发流程、降低了AI算子的开发门槛。开发者无需深度钻研硬件底层实现细节，即可完成高性能算子开发，若需极致的硬件适配优化，可进一步参考官方硬件实现手册。

一、AI Core核心组件架构

AI Core作为独立的计算核心，集成了计算单元、存储单元、数据搬运单元三大核心模块，各模块分工明确、协同工作，共同完成AI运算的数据处理、计算、存储流转全流程，是实现异步并行计算的核心基础。

1.1 计算单元：三类算力资源协同运算

计算单元是AI Core的算力核心，包含Scalar、Vector、Cube三种基础计算资源，分别适配不同类型的计算场景，实现标量、向量、矩阵运算的全覆盖：

• Scalar标量计算单元：作为AI Core的控制中枢，主要负责地址计算、循环控制等基础标量计算工作。同时承担指令调度核心职责，可向向量计算、矩阵计算、数据搬运、同步控制等单元发射对应指令，统筹整体运算流程。
• Vector向量计算单元：专注于各类向量运算场景，适配AI模型中大量的一维数据运算需求，是常规特征计算、数据预处理运算的核心算力支撑。
• Cube矩阵计算单元：专门负责矩阵运算，针对深度学习中的卷积、矩阵乘法等核心高维运算做了硬件级优化，是AI模型训练和推理的核心算力单元。

1.2 存储单元：内外双层存储体系

AI Core采用分层存储架构，分为内部存储和外部存储，两类存储对应专属数据类型，适配不同的数据读写场景，兼顾读写速度和存储容量：

• 内部存储Local Memory：是AI Core的高速本地存储，对应数据类型为LocalTensor，读写延迟极低，主要用于存储当前核计算过程中的临时数据、中间结果，保障计算单元高效读取数据。
• 外部存储Global Memory：是AI Core可访问的全局存储，对应数据类型为GlobalTensor，存储容量大，主要用于存储模型输入输出数据、全局参数等大批量、非实时计算数据。

1.3 搬运单元：DMA数据传输核心

DMA（直接内存访问）搬运单元是数据流转的核心载体，无需占用计算单元资源，独立完成数据搬运工作。核心功能涵盖两大场景：一是实现Global Memory与Local Memory之间的数据双向搬运，完成计算数据的加载和结果回写；二是支持不同层级Local Memory之间的数据流转，保障核内数据高效调度。

二、AI Core异步并行计算核心流程

AI Core的核心优势在于异步并行计算能力，整体运行依托异步指令流、同步信号流、计算数据流三大核心过程协同驱动，实现多单元并行高效运算，三个流程相互独立又相互约束，保障运算的高效性与准确性。

2.1 异步指令流：多单元并行执行

整个指令执行流程由Scalar单元主导，Scalar读取完整指令序列后，将向量计算、矩阵计算、数据搬运等不同类型指令，分别下发至对应单元的指令队列。Vector、Cube、DMA单元接收指令后，无需相互等待，异步并行执行各自任务，最大化挖掘硬件并行算力，该流程对应硬件架构图中蓝色箭头指令流。

2.2 同步信号流：保障指令执行逻辑

实际运算中，不同指令队列的指令存在数据依赖、逻辑先后关系，无序执行会导致运算出错。为此，Scalar单元会同步下发同步指令，通过同步信号约束各单元的执行时序，确保存在依赖关系的指令按照正确逻辑顺序执行，规避运算异常，该同步过程对应绿色箭头同步信号流。

2.3 计算数据流：数据处理全链路

AI Core数据处理遵循“加载-计算-回写”的标准链路：首先由DMA单元将批量数据从大容量的Global Memory搬运至高速Local Memory，供计算单元快速读取；随后Vector或Cube单元根据运算需求完成数据计算，并将临时结果写入Local Memory；最后由DMA单元将最终计算结果从Local Memory搬运回Global Memory，完成一次完整的数据处理流程，该数据流对应红色箭头数据流转链路。

三、Ascend C核函数开发规范与实践

核函数（Kernel Function）是Ascend C算子设备侧实现的核心入口，是连接开发者代码与AI Core硬件执行的关键。Ascend C基于C/C++语法扩展，支持开发者编写设备端运行代码，通过核函数实现算子的数据访问、计算、流转全功能，是昇腾算子开发的核心基础。

与普通C++函数单次调用单次执行不同，核函数具备多核并行执行特性：调用核函数时，多个AI核心会加载同一份核函数代码，使用相同的函数入参，同步并行完成运算，充分发挥AI处理器的多核算力优势。

3.1 核函数基础定义规则

核函数的定义需遵循专属语法规范，包含函数限定符、变量限定符、返回值、入参类型等强制规则，具体要求如下：

3.1.1 函数类型限定符

核函数必须同时添加 __global__ 和 __aicore__ 两个限定符：__global__ 用于标识该函数为核函数，支持通过专属内核调用符调用；__aicore__ 用于标识该函数运行在设备端AI Core上。标准声明格式如下：

\_\_global\_\_ \_\_aicore\_\_ void kernel\_name\(argument list\);

3.1.2 变量类型限定符

核函数的指针入参必须添加 **__gm__**限定符，用于标识指针指向Global Memory的内存地址。为统一代码规范、简化开发，官方推荐使用GM_ADDR宏替代原生限定符，宏定义及使用方式如下：

\#define GM\_ADDR \_\_gm\_\_ uint8\_t\*

3.1.3 强制语法规则与开发建议

核心强制规则：核函数必须为void无返回值类型；入参仅支持指针类型或C/C++内置基础数据类型。开发规范建议：统一使用uint8_t类型指针作为入参，后续根据实际运算需求强制转换为对应数据类型，保证代码统一性。

3.2 三类函数调用关系

在Ascend C编程体系中，函数分为核函数、host侧执行函数、device侧普通执行函数三类，三者调用关系严格受限：

• host侧函数：可调用同类host函数，也可通过内核调用符调用核函数；
• device侧普通函数：仅可调用同类device侧函数；
• 核函数：仅可调用device侧普通函数，不可嵌套调用其他核函数。

3.3 核函数调用方式与执行配置

核函数不支持普通函数调用方式，需通过专属内核调用符 <<<...>>> 配置执行参数并调用，该语法仅支持NPU侧编译，CPU侧无法识别。标准调用格式为：

kernel\_name\<\<\<numBlocks, l2ctrl, stream\>\>\>\(argument list\);

三个核心执行参数详解：

• numBlocks：逻辑核数量，指定核函数并行执行的核心数，取值范围1-65535，直接决定并行算力规模。该参数需适配硬件运行模式：耦合模式下统一设置AI Core核实例数量；分离模式下，向量算子对应Vector核数、矩阵算子对应Cube核数、融合算子对应核组合数量，且不可超过硬件物理核上限，同时需适配设备资源限制配置。开发者可通过GetCoreNumAiv、GetCoreNumAic等接口获取硬件核数，合理配置参数。每个执行核可通过GetBlockIdx获取唯一逻辑ID。
• l2ctrl：保留参数，开发中固定设置为nullptr，无需开发者自定义配置；
• stream：异步流，类型为aclrtStream，用于管控异步任务执行顺序，保障代码逻辑与设备执行顺序一致。

核函数调用为异步执行，主机端下发调用指令后会立即收回控制权，若需同步等待执行完成，可调用aclrtSynchronizeStream函数阻塞等待，确保任务执行完毕后再执行后续代码。

3.4 基础核函数代码示例（Add算子）

以矢量相加算子为例，完整展示核函数的定义与调用逻辑，包含算子初始化、核心计算、内核调用全流程：

// 实现核函数
__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 numBlocks, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y, uint8_t* z)
{
    // 指定8个逻辑核并行执行，完成矢量相加运算
    add_custom<<<8, nullptr, stream>>>(x, y, z);
}

3.5 模板核函数开发

Ascend C支持模板核函数开发，可通过模板参数适配不同数据类型、不同运算维度，提升代码复用性。模板参数支持非类型参数和类型参数两类，定义与调用示例如下：

// 模板核函数定义：a为非类型模板参数，T为类型模板参数
template<int a, typename T>
__global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
    AscendC::printf("Print Template a: %d\n", a);
    // 根据核逻辑ID分配对应数据内存
    xGm.SetGlobalBuffer((__gm__T*)x + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
    yGm.SetGlobalBuffer((__gm__T*)y + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
    zGm.SetGlobalBuffer((__gm__T*)z + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
}

// 模板核函数调用：传入模板参数20、float类型，指定并行核数执行
add_custom<20, float><<<numBlocks, nullptr, stream>>>(x, y, z);

四、总结

AI Core依托“计算+存储+搬运”的模块化架构，结合异步并行、同步约束、数据流转的三维运行机制，构建了高效稳定的AI算力底座。而Ascend C编程体系通过硬件架构抽象和标准化的核函数开发规范，既屏蔽了硬件底层差异，保留了多核并行的高性能优势，又大幅降低了算子开发门槛。开发者严格遵循核函数的语法定义、参数配置、调用规则，即可快速开发出适配昇腾硬件的高性能自定义算子，适配各类AI模型的推理与训练场景。