大家好，我是继续在CANN训练营中不断挑战自我的学员。当我成功实现首个向量加法算子时，自信满满，认为Ascend C也不过如此。然而，当我尝试优化一个较为复杂的算子时，发现性能非常差。这时，训练营的导师一针见血地指出：“你的数据在‘仓库’里停留的时间比在‘车间’里还要长！”

这一句话让我豁然开朗。此前我只是机械地知道需要复制数据，却未能深入理解Ascend C内存体系的设计理念。

在Ascend C的内存体系中，主要有三大内存类型：GM（全局内存）、LM（本地内存）和Register（寄存器）。这三种内存各司其职，又相互依赖。不了解它们之间的关系，就如同在迷宫中徘徊，难以达到高效能的目标。

今天，我将利用在[2025年昇腾CANN训练营第二季]学到的知识，结合个人实践中的经验教训，为大家提供一份走出内存迷宫的指南。

>> 系统化的知识体系是解决问题的关键，欢迎加入训练营共同学习：
点击加入

第一章：因“慢”而起的案例分析——为何要区分不同的内存类型？

我遇到的一个“慢”算子，原始代码如下（简化版）：

__global__ __aicore__ void my_slow_kernel(__gm__ uint8_t* input, __gm__ uint8_t* output) {
    for (int i = 0; i < length; i++) {
        // 直接访问Global Memory进行计算
        output[i] = input[i] * 2 + 1;
    }
}

这段代码逻辑上没有问题，但执行效率极低。导师的评价是：“你让宝贵的AI Core计算单元像搬运工一样，频繁往返于远端的大仓库（GM）取数据，再返回进行处理，时间都浪费在了路上。”

这句话使我茅塞顿开。原来，Ascend C的内存层次设计的核心在于：使数据尽可能接近计算单元。基于此，我们引入了三个关键概念：

• Global Memory：中央大仓库
• Local Memory：车间的工作台
• Register：工人手中的工具台

第二章：三足鼎立——三大内存类型的特征与作用

为了更好地理解这些内存类型，我对它们进行了详细的描述。

GM - 全局内存

• 定位：片外DRAM，具有最大的容量（GB级），但访问速度最慢，是所有核函数共享的“中央仓库”。
• 访问方式：通过核函数中的指针访问。
  

  __gm__

• 职责：存储计算的输入和最终输出，同时也是主机与设备间数据交换的地方。

直接且频繁地在核函数中访问GM会严重拖累性能。

LM - 本地内存

• 定位：位于AI Core上的SRAM，容量较小（MB级），但访问速度极快，是每个AI Core核心独享的“车间工作台”。
• 访问方式：通过声明固定大小的数组来使用，例如：
  

  uint8_t localBuffer[BUFFER_SIZE];

• 职责：作为高速数据缓存区。计算前，将GM中的数据批量转移到LM；计算过程中，在LM上执行；计算完成后，将结果从LM批量写回GM。这是性能优化的主要战场。

Register - 寄存器

• 定位：位于AI Core计算单元旁的极小容量高速存储，具有最快的访问速度，但容量最小（KB级），可以视为每个计算指令直接使用的“工具和材料”。
• 访问方式：由编译器自动管理和分配。通常，我们定义的局部变量（非数组）和循环索引等会存储在寄存器中。
• 职责：存储即将参与计算的标量数据和中间结果。

为了更直观地理解，我制作了一张“访问速度与容量”对比图，放在笔记中随时提醒自己：

[访问速度]： Register >> LM > GM
[存储容量]： GM >> LM > Register

第三章：协同作战——最佳合作模式实践

了解了各自的特性后，正确的“工作流程”应如何操作？训练营的“码力全开特辑”提供了以下黄金法则：批量搬运，本地计算。

让我们重写上述慢速算子，看看标准流程是如何实施的：

__global__ __aicore__ void my_fast_kernel(__gm__ uint8_t* input, __gm__ uint8_t* output, uint32_t totalLength) {
    // 0. 初始化与任务划分 (略)
    // ...

    // 1. 指针绑定：定义指向GM的“提货单”
    __gm__ uint8_t* globalIn = input + currentOffset;
    __gm__ uint8_t* globalOut = output + currentOffset;

    // 2. 申请LM：在“工作台”上开辟两块固定区域
    constexpr int32_t TILE_LENGTH = 256; // 定义每次搬运的数据块大小
    uint8_t localIn[TILE_LENGTH];
    uint8_t localOut[TILE_LENGTH]; // 注意：输出也需要LM缓冲

    // 3. 数据搬运：从GM（仓库）批量搬运数据到LM（工作台）
    for (int32_t i = 0; i < currentLength; ++i) {
        localIn[i] = globalIn[i];
    }

核心计算阶段在LM上执行所有计算任务。此时，所有数据均位于高速的LM上，计算单元处于全速运行状态。

for (int32_t i = 0; i < currentLength; ++i) {
  // 中间变量，例如temp，由编译器自动分配至寄存器
  uint8_t temp = localIn[i] * 2; // temp -> 寄存器
  localOut[i] = temp + 1;        // 读写操作均在LM上进行，速度极快
}
  

随后，计算结果需从LM（工作台）批量传输回GM（仓库）。

for (int32_t i = 0; i < currentLength; ++i) {
  globalOut[i] = localOut[i];
}
  

整个流程简明扼要：从GM到LM，利用计算过程中的寄存器支持，再次返回LM，最后回到GM。通过减少GM的频繁访问，采用一次性批量转移的方式，显著提升了性能效率。

第四章：常见误解与实际案例

在实践中，我也遇到了不少挑战：

误解一：“LM容量越大越好”

我曾尝试扩大LM数组的规模，最终导致编译失败。指导老师解释说，LM是一种有限资源，过度使用会限制其他内核的可用性。因此，需要根据实际数据块大小进行细致规划。

误解二：“无需关注寄存器管理”

尽管编译器能自动管理寄存器，但如果在循环体内声明过多局部变量，可能导致寄存器溢出。这迫使编译器将某些变量转移到较慢的LM或GM上，从而严重影响性能。

实例：忽略LM初始化

有一次，在将数据移至LM后，我在计算过程中不慎使用了未初始化的LM数组部分，导致结果出现随机值。务必保证LM中所有参与计算的数据都已被正确定义。

总结：从复杂到清晰

起初，GM、LM、寄存器对我来说就像一个复杂的迷宫。现在，我学会了像建筑师一样思考这些组件之间的关系。我不再是那个仅复制粘贴代码的新手，而是开始理解每个指令背后的战略价值，开始评估这些数字背后的平衡点。

DataCopy

这一套内存架构是Ascend C发挥昇腾硬件强大计算能力的基础。掌握了这一点，你才算真正获得了开启高性能算子开发之门的钥匙。

TILE_LENGTH

在接下来的训练营阶段，我们将探讨如何实现GM到LM的数据传输与LM上的计算同步进行——即更先进的双缓冲流水线技术。对此，我已经充满期待。

掌握内存模型是性能优化的第一步。想要系统地学习更多Ascend C的高级技巧？>> 立即报名2025年CANN训练营第二季