摘要

本文作为系列开篇，将深度剖析 Ascend C 算子的两种核心开发模式：Kernel 算子拆解开发(Kernel Operator Disassembled Development) 和 工程化程序开发(Engineering Program Development)。我们将从昇腾（Ascend）硬件基础架构入手，通过对比两种模式的适用场景、工作流程与优劣，并结合详尽的架构示意图、代码示例与性能分析，帮助开发者根据实际需求选择最佳开发路径。本文将为后续深入探讨 Tiling (分块处理)、流水线 (Pipeline) 等高级主题打下坚实的理论基础。

1. 背景介绍：为什么需要特定的算子开发模式？

随着人工智能模型规模的指数级增长，对底层算力的需求愈发迫切。通用处理器（CPU）在处理大规模并行计算时显得力不从心，因此，专为AI计算设计的AI加速器（如NPU）应运而生。华为的昇腾（Ascend）AI处理器就是其中的杰出代表。然而，强大的硬件需要与之匹配的软件栈才能发挥极致性能。Ascend C 正是在这一背景下诞生的，它是一种基于C/C++的编程语言扩展，旨在让开发者能够直接、高效地操控昇腾AI处理器中的计算核心（AI Core）。

一个高效的AI算子（Operator）不仅是算法逻辑的正确实现，更需要对硬件架构有深刻的理解。直接将为CPU设计的算法移植到NPU上，往往无法获得理想的性能提升，有时甚至性能更差。因此，需要一套与硬件特性紧密耦合的编程模型和开发范式。您素材图片中提到的 “Kernel算子拆解开发”和 “工程化程序开发”，正是Ascend C为不同复杂度、不同性能要求的计算任务所提供的两种标准化开发范式。理解它们的本质差异和适用场景，是进行高性能算子定制和优化的第一步，也是避免陷入“高性能硬件，低效能代码”困境的关键。

（此处插入素材图片1：开发模式概览图）

https://via.placeholder.com/800x400.png?text=Material+Image+1+-+Kernel+vs+Engineering+Development图1：素材中所示的两种Ascend C开发模式概览。它清晰地展示了两种路径的选择，是本文讨论的起点。

2. 昇腾AI Core基础架构浅析

要理解开发模式，必须先了解其运行的硬件环境。昇腾AI Core的内部是一个典型的单指令多数据流(SIMD, Single Instruction Multiple Data) 和超长指令字(VLIW, Very Long Instruction Word) 架构，其核心计算资源包括：

• Cube Unit (立方计算单元)：用于执行高密度的矩阵乘法（MatMul）和卷积（Convolution）运算。

• Vector Unit (向量计算单元)：用于执行向量运算，如加法、乘法、激活函数等。

• Scalar Unit (标量计算单元)：负责执行控制逻辑、地址计算等标量任务。

• Local Memory (本地内存)/ Buffer (缓冲区)：位于AI Core内部的高速缓存，用于暂存待计算的数据和中间结果。与片外Global Memory (全局内存)相比，其访问速度极快但容量有限。

graph TD
    A[Global Memory] -->|DMA搬运| B[Local Memory/Buffer];
    B --> C[Vector Unit];
    B --> D[Cube Unit];
    C --> E[计算结果写回Buffer];
    D --> E;
    E -->|DMA搬运| A;
    F[Scalar Unit] --> G[控制指令流];
    G --> C;
    G --> D;

图2：昇腾AI Core简化计算流程。数据通过DMA从Global Memory搬运至Local Memory，计算单元从Buffer中取数计算，结果再写回。

这个架构带来了一个关键挑战：计算速度远高于数据搬运速度。如果计算单元经常等待数据从Global Memory加载，那么大部分时间将花在等待上，形成“内存墙”，无法发挥计算单元的性能。因此，Ascend C开发模式的核心思想之一，就是通过并行化(Parallelism) 和流水线(Pipelining) 技术，将数据搬运与计算过程重叠起来，尽可能掩盖数据搬运的延迟。

3. 原理详解：两种开发模式的深度对比

3.1 Kernel 算子拆解开发模式：标准化的高效流水线

这种模式是Ascend C推荐的首选范式，尤其适用于计算逻辑规整、数据依赖性不强的算子（如Element-wise操作、简单的Reduce等）。

💡 核心思想：将算子的计算过程抽象并“拆解”为一系列可以并行执行的、标准化的流水任务(Pipeline Tasks)。开发者只需像填空一样，在预设的“模板”中实现特定任务的逻辑。

3.1.1 标准三级流水线模型

一个典型的Kernel算子开发遵循如图3所示的三级流水线，这对应了您素材中提到的“拆解”概念。

graph LR
    subgraph Pipeline
        A[Data Move-In<br>数据搬入] --> B[Compute<br>核心计算];
        B --> C[Data Move-Out<br>数据搬出];
    end

    subgraph Memory
        D[Global Memory<br>输入数据] --> A;
        A --> E[Local Buffer];
        E --> B;
        B --> F[Local Buffer<br>结果数据];
        F --> C;
        C --> G[Global Memory<br>输出数据];
    end

    H[Scalar Unit<br>循环控制 & 地址计算] -.->|控制信号| A;
    H -.->|控制信号| B;
    H -.->|控制信号| C;

图3：标准的三级流水线模型。Data Move-In、Compute、Data Move-Out三个任务并行执行，重叠数据搬运与计算。

1. 数据搬入 (Data Move-In / Copy-In)：

   • 任务：通过DMA（直接内存访问）控制器，将当前需要计算的数据块从片外Global Memory(GM) 异步地搬运到片内Local Memory(LM) 或特定的Buffer中。

   • 关键点：此操作由硬件DMA单元执行，不占用计算核心资源。__aicore__关键字确保该函数在AI Core上执行。

2. 核心计算 (Compute)：

   • 任务：计算单元（Vector或Cube）从本地Buffer中读取数据，执行具体的算术或逻辑运算，并将结果写回另一个本地Buffer。

   • 关键点：计算任务与数据搬运任务是并行的。当计算单元在处理第N个数据块时，DMA可能正在搬运第N+1个数据块，同时另一个DMA可能在写回第N-1个数据块的结果。

3. 数据搬出 (Data Move-Out / Copy-Out)：

   • 任务：将计算完成的结果数据从本地Buffer通过DMA异步地搬运回Global Memory。

   • 关键点：搬出操作也是并行的。

3.1.2 优势与局限

• 优势:

  • 高开发效率：模板化开发，降低入门门槛，代码结构清晰统一。

  • 自动并行优化：框架内置了流水线并行机制，开发者无需显式管理并行任务，即可获得不错的性能。

  • 高可维护性：代码遵循统一规范，易于团队协作和后期维护。

• 局限:

  • 灵活性受限：流水线级数和任务类型相对固定，难以处理复杂的数据依赖或非常规的控制流。

  • 性能天花板：对于某些极端优化场景，固定的流水线可能无法完全榨干硬件性能。

3.2 工程化程序开发模式：极致的自由与灵活性

当算子的计算逻辑无法被简单地套入“搬入-计算-搬出”的三级流水线时，就需要采用工程化程序开发模式。这通常发生在以下情况：

• 复杂的数据依赖：计算过程需要多次中间结果的回写和重读。

• 动态计算路径：计算流程需要根据数据内容进行动态分支选择。

• 需要多核协同：单个算子的计算需要多个AI Core之间进行复杂的通信和同步。

• 对性能有极致要求：需要手动控制每一个时钟周期和内存访问。

🛠️ 核心思想：将AI Core视为一个完整的、可编程的计算机，开发者拥有对计算单元、内存系统、DMA控制器等几乎所有资源的完全控制权。开发者需要自行设计数据流、控制流和任务间的同步机制。

3.2.1 工程化模式的关键技术

1. 精细内存管理：开发者需要手动管理Global Memory、Local Memory以及多级Buffer之间的数据流动，可能涉及复杂的数据切片、重叠（Double Buffering甚至Multi-Buffering）以最大化数据复用率。

2. 显式任务并行与同步：可以启动多个并行任务（如多个DMA搬运任务和多个计算任务），并使用信号量(Semaphore) 或事件(Event) 等机制来精确控制任务间的依赖关系和执行顺序。

3. 复杂控制流：直接使用C/C++的控制流语句（如if-else, switch-case）来实现动态逻辑，这在对稀疏矩阵或条件执行（如While循环）的支持中至关重要。

graph TD
    A[GM: Input A] -->|DMA_Task1| B[LM: Buffer A1];
    A -->|DMA_Task2| C[LM: Buffer A2];
    D[GM: Input B] -->|DMA_Task3| E[LM: Buffer B1];
    D -->|DMA_Task4| F[LM: Buffer B2];

    subgraph “并行计算阶段”
        direction LR
        G[Compute_Task1<br>处理A1&B1] --> H[LM: Result1];
        I[Compute_Task2<br>处理A2&B2] --> J[LM: Result2];
    end

    B --> G;
    E --> G;
    C --> I;
    F --> I;

    K[Sync Barrier] --> G;
    K --> I;

    H -->|DMA_Task5| L[GM: Final Result];
    J -->|DMA_Task6| L;

图4：一个简化的工程化模式示例：双缓冲（Double Buffering）技术。当Compute_Task1在处理Buffer A1和B1时，DMA任务正在并行地为下一次计算填充Buffer A2和B2。计算任务之间可能需要同步点（Sync Barrier）。

3.2.2 优势与挑战

• 优势:

  • 终极灵活性：可以实现任何复杂度的算法。

  • 性能潜力最大：通过精细调控，可以逼近硬件的理论性能峰值。

• 挑战:

  • 开发复杂度极高：相当于在汇编层面进行优化，容易出错。

  • 调试困难：并行和异步带来的不确定性使得问题定位非常棘手。

  • 可移植性差：代码高度依赖特定硬件架构，硬件迭代可能带来大量迁移工作。

4. 代码实战：从Hello World到初步优化

让我们通过一个具体的例子——一维向量加法(Add算子)，来感受两种模式在代码实现上的巨大差异。

4.1 基于Kernel模式的向量加法实现

以下是使用Kernel算子拆解开发模式实现向量加法的详细代码。

// ==== 头文件包含 ====
#include "kernel_operator.h" // 包含Ascend C内核操作符的关键宏和接口定义

// ==== 算子类定义 ====
// 类名通常与算子名对应，并继承自某个基类（此处简化）
class AddCustomKernel {
public:
    // 构造函数，通常为空
    __aicore__ inline AddCustomKernel() {}

    // ==== 初始化函数 (Init) ====
    // 作用：从GM中获取输入输出Tensor的地址和Tiling参数。
    // 参数说明：
    // - xGm: 输入向量x在Global Memory中的起始地址
    // - yGm: 输入向量y在Global Memory中的起始地址
    // - zGm: 输出向量z在Global Memory中的起始地址
    // - tiling: Tiling结构体指针，包含了数据总长度、每个Core的分块大小等信息。
    __aicore__ inline void Init(GM_ADDR xGm, GM_ADDR yGm, GM_ADDR zGm, uint32_t tiling) {
        // 将GM地址赋值给类成员变量，供Process函数使用
        this->xGm = xGm;
        this->yGm = yGm;
        this->zGm = zGm;
        // 将Tiling参数从GM拷贝到Local Memory中，以便快速访问
        this->tiling = (uint32_t)tiling;
    }

    // ==== 核心处理函数 (Process) ====
    // 作用：实现标准的三级流水线逻辑。这是算子的执行入口。
    __aicore__ inline void Process() {
        // 步骤1: 解析Tiling参数，确定当前AI Core需要处理的数据块
        // GET_TILING_DATA_LEN 和 GET_TILING_DATA_OFFSET 是宏，用于从tiling参数中提取信息
        uint32_t totalDataLen = GET_TILING_TOTAL_LENGTH(this->tiling); // 数据总长度
        uint32_t currentCoreDataLen = GET_TILING_DATA_LEN(this->tiling); // 当前Core处理的数据长度
        uint32_t currentCoreDataOffset = GET_TILING_DATA_OFFSET(this->tiling); // 当前Core的数据起始偏移

        // 步骤2: 初始化管道 (Pipe)
        // Pipe是Ascend C中管理流水线任务和数据队列的核心抽象
        Pipe pipe;

        // 步骤3: 定义数据队列 (Queue) 及其缓冲区大小 (Buffer Size)
        // TBuffer是模板类，TPosition::VECIN表示输入向量队列
        // BUFFER_SIZE 是一个预定义的常量，表示每次流水线循环处理的数据量（如256个float）
        constexpr uint32_t bufferSize = BUFFER_SIZE;
        TBuffer<TPosition::VECIN, float> xQueue; // 输入x的数据队列
        TBuffer<TPosition::VECIN, float> yQueue; // 输入y的数据队列
        TBuffer<TPosition::VECOUT, float> zQueue; // 输出z的数据队列

        // 步骤4: 主循环 - 以Buffer为单位，循环处理当前Core负责的数据块
        for (uint32_t i = 0; i < currentCoreDataLen; i += bufferSize) {
            // 计算本次循环实际处理的数据长度（最后一次循环可能小于bufferSize）
            uint32_t currentLoopLen = min(bufferSize, currentCoreDataLen - i);

            // ==== 流水线任务开始 ====
            // 任务1: Data Move-In (数据搬入)
            // 将输入数据x和y从GM异步搬运到LM的队列中
            pipe.In(xQueue, this->xGm + currentCoreDataOffset + i, currentLoopLen);
            pipe.In(yQueue, this->yGm + currentCoreDataOffset + i, currentLoopLen);

            // 任务2: Compute (核心计算)
            // 调用自定义的计算函数，从队列中取数，计算，结果放入输出队列
            this->AddCompute(xQueue, yQueue, zQueue, currentLoopLen);

            // 任务3: Data Move-Out (数据搬出)
            // 将计算结果从LM的队列异步搬运回GM
            pipe.Out(this->zGm + currentCoreDataOffset + i, zQueue, currentLoopLen);
            // ==== 流水线任务结束 ====
        }
    }

private:
    // ==== 计算核函数 (AddCompute) ====
    // 作用：执行具体的向量加法运算。
    // 参数：输入队列x, y，输出队列z，以及计算长度len。
    __aicore__ inline void AddCompute(TBuffer<float>& xQueue,
                                      TBuffer<float>& yQueue,
                                      TBuffer<float>& zQueue,
                                      uint32_t len) {
        // 使用循环处理队列中的每一个元素
        for (uint32_t i = 0; i < len; ++i) {
            // 从输入队列中读取数据，执行加法，结果写入输出队列
            // 这是计算发生最密集的地方
            zQueue[i] = xQueue[i] + yQueue[i];
        }
    }

    // ==== 成员变量 ====
    GM_ADDR xGm; // 输入向量x在GM的地址
    GM_ADDR yGm; // 输入向量y在GM的地址
    GM_ADDR zGm; // 输出向量z在GM的地址
    uint32_t tiling; // Tiling参数
};

// ==== 算子调用入口函数 (Kernel入口) ====
// 这是一个extern "C"函数，是主机侧调用该算子的接口。
// 参数：KernelId（内核ID），用于多核并行；接口参数句柄。
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR tiling) {
    // 实例化算子类
    AddCustomKernel addKernel;
    // 调用初始化函数，设置参数
    addKernel.Init(x, y, z, (uint32_t)tiling);
    // 调用处理函数，开始执行
    addKernel.Process();
}

代码清单1：基于Kernel模式的向量加法完整实现。注释详细解释了每个关键部分的作用。

代码分析：

• 结构清晰：Init-> Process-> AddCompute，分工明确，符合模板要求。

• 流水线并行：Pipe.In, AddCompute, Pipe.Out在循环中依次调用，但借助硬件能力，这三个阶段是并行执行的。

• Tiling策略：代码通过tiling参数实现了数据分块，这是支持多核并行计算的基础，我们将在后续文章中详细讲解。

4.2 工程化模式的概念性伪代码

由于工程化模式代码过于复杂且无固定范式，这里给出一个实现向量加法的概念性伪代码，以展示其思路的不同。

// 伪代码，仅用于示意
extern "C" __global__ __aicore__ void add_engineering(GM_ADDR x, GM_ADDR y, GM_ADDR z, ...) {
    // 1. 手动计算本核的数据偏移和长度（可能更复杂）
    int blockId = get_block_id();
    int blockDim = get_block_dim();
    // ... 复杂的地址计算

    // 2. 手动管理多级Buffer（例如双缓冲）
    __local__ float bufferX1[BUFFER_SIZE], bufferX2[BUFFER_SIZE];
    __local__ float bufferY1[BUFFER_SIZE], bufferY2[BUFFER_SIZE];
    __local__ float bufferZ1[BUFFER_SIZE], bufferZ2[BUFFER_SIZE];

    // 3. 显式启动异步DMA任务进行数据搬运
    int dmaTaskId1 = hc::dma_async(bufferX1, x + offset1, size1);
    int dmaTaskId2 = hc::dma_async(bufferY1, y + offset1, size1);

    // 4. 使用信号量进行同步：等待第一批数据搬运完成
    hc::wait_semaphore(dmaTaskId1);
    hc::wait_semaphore(dmaTaskId2);

    for (int i = 0; i < numBlocks; ++i) {
        // 5. 在计算当前块的同时，异步启动下一块数据的DMA搬运（双缓冲）
        if (i + 1 < numBlocks) {
            hc::dma_async(get_next_bufferX(i), x + offset_next, size_next);
            hc::dma_async(get_next_bufferY(i), y + offset_next, size_next);
        }

        // 6. 执行计算（可能使用内联汇编或 intrinsic 函数以追求极致性能）
        vector_add_intrinsic(get_current_bufferX(i), get_current_bufferY(i), get_current_bufferZ(i), size);

        // 7. 异步启动结果回写DMA
        hc::dma_async(z + offset_out, get_current_bufferZ(i), size);

        // 8. 在循环尾部进行复杂的同步，切换Buffer指针
        hc::barrier(); // 等待所有核到达同步点（如果需要核间同步）
        swap_buffers(); // 交换当前缓冲和下一块缓冲
    }
}

代码清单2：工程化模式的概念性伪代码。展示了手动DMA控制、双缓冲、显式同步等关键技术点。

5. 结果分析：性能与开发效率的权衡

为了更直观地对比两种模式，我们构建一个分析矩阵：

评估维度	Kernel 算子拆解开发	工程化程序开发	分析说明
开发速度	⭐⭐⭐⭐⭐ (极快)	⭐⭐ (极慢)	Kernel模式有固定模板，填代码即可；工程化模式需从头设计。
代码可读性/维护性	⭐⭐⭐⭐⭐ (高)	⭐⭐ (低)	Kernel模式结构统一；工程化模式代码像“面条”，难以维护。
初始性能	⭐⭐⭐⭐ (良好)	⭐⭐⭐ (可能较差)	Kernel模式自动流水线能提供良好基础性能；工程化模式若设计不当，性能可能更差。
性能优化上限	⭐⭐⭐ (中等)	⭐⭐⭐⭐⭐ (极高)	Kernel模式受框架限制；工程化模式理论上可逼近硬件峰值。
适用算子范围	规整的Element-wise, 简单Reduce, MatMul等	任何算子，特别是复杂、不规则、动态的算子	Kernel模式覆盖80%的常见场景；工程化模式解决剩余20%的难题。
对开发者要求	中级（理解Ascend C模板和硬件概念）	专家级（需精通硬件架构、并行编程、性能分析）

结论：

对于一维向量加法这类标准算子，Kernel算子拆解开发模式是毫无争议的最佳选择。它能够在极短的开发时间内，交付具有良好、可预测性能的算子。而工程化模式则像一把“手术刀”，只有在面对极其特殊、对性能有变态要求的场景时，才值得投入巨大的开发和研究成本去使用。

（此处插入自绘对比分析图）

quadrantChart
    title Ascend C开发模式选择象限图
    x-axis “低开发灵活性” --> “高开发灵活性”
    y-axis “低性能潜力” --> “高性能潜力”
    quadrant-1 “挑战区：考虑工程化”
    quadrant-2 “优势区：首选Kernel模式”
    quadrant-3 “简单任务：Kernel模式”
    quadrant-4 “创新区：可能需工程化”
    “向量加法”: [0.2, 0.3]
    “矩阵乘法”: [0.4, 0.7]
    “动态稀疏卷积”: [0.8, 0.9]
    “自定义创新算法”: [0.7, 0.6]

图5：开发模式选择象限图。X轴代表灵活性需求，Y轴代表性能需求。像向量加法这样的简单算子落在第三象限，应首选Kernel模式。

6. 总结与讨论

本文深入探讨了Ascend C算子开发的两种根本路径。Kernel算子拆解开发通过标准化的三级流水线模板，在开发效率和性能之间取得了绝佳的平衡，是大多数场景下的推荐做法。而工程化程序开发则提供了终极的灵活性和性能控制力，是解决复杂计算难题的“终极武器”，但代价是极高的复杂度和成本。

理解这两种模式的本质，是成为一名优秀的Ascend C开发者的第一步。在后续的文章中，我们将聚焦于Kernel模式下的关键技术，例如Tiling策略的详细解析、流水线的深度优化等，因为这些是广大开发者最常接触且最能直接产生价值的知识。

讨论点：

1. 在您遇到的实际业务场景中，除了向量加法，您认为哪些算子天然适合Kernel模式？哪些算子您曾考虑过或被迫使用过工程化模式？背后的原因是什么？

2. 您认为能否在Kernel模式的固定流程基础上，引入一些可配置的选项（例如，允许增加一级流水线或自定义同步点），从而在保持易用性的同时提升灵活性？这样做的潜在挑战是什么？

参考链接

1. 昇腾社区官方文档 - Ascend C 编程指南：最权威的参考资料，包含完整的API说明和编程规范。

2. Ascend C 算子开发样例 - GitHub：华为官方提供的算子示例库，包含Add等算子的完整实现。

3. 昇腾AI处理器架构白皮书：深入了解昇腾硬件架构，是进行深度优化的必备知识。

---