大家好，依然是在CANN训练营里持续“升级打怪”的我。实现了第一个向量加法算子后，我一度信心爆棚，觉得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的内存分层，其核心思想是：让数据离计算单元越近越好。这就引出了我们的三位主角：

1. Global Memory：中央大仓库
2. Local Memory：车间的工作台
3. Register：工人手中的工具台

第二章：三足鼎立——三大内存的特性与职责

为了彻底理解，我为他们各自建立了一份“人物档案”。

GM - 全局内存

• 定位：片外DRAM，容量最大（GB级别），但速度最慢，是所有核函数共享的“中央仓库”。
• 访问方式：在核函数中，通过 __gm__ 指针来访问。直接、频繁地在核函数内部访问GM是性能的头号杀手！
• 职责：存放计算的输入和最终输出。Host与Device的数据交换也在这里进行。

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];
    }

    // 4. 【核心计算阶段】在LM上进行所有计算
    //    此时，所有数据都在高速的LM上，计算单元全速运转！
    for (int32_t i = 0; i < currentLength; ++i) {
        // 计算过程中的中间变量，如temp，由编译器自动放入Register
        uint8_t temp = localIn[i] * 2; // temp -> Register
        localOut[i] = temp + 1;        // 读写都在LM上，极快
    }

    // 5. 结果回写：将结果从LM（工作台）批量搬回GM（仓库）
    for (int32_t i = 0; i < currentLength; ++i) {
        globalOut[i] = localOut[i];
    }
}

看，流程非常清晰：GM -> LM -> (计算，配合Register) -> LM -> GM。通过一次批量搬运，代替了之前成千上万次零散的GM访问，性能得到了数量级的提升。

第四章：常见迷思与“翻车”现场

在实践的路上，我也没少踩坑：

• 迷思一：“LM越大越好”：我曾试图把LM数组开得巨大，结果编译失败。老师指出，LM是稀缺资源，一个核用多了，其他核就没得用。需要根据数据块大小精细规划。
• 迷思二：“Register不用管”：虽然编译器自动管理，但如果我们写一个超大的循环体，内部声明大量局部变量，可能会导致寄存器溢出，编译器被迫将一些变量“塞”到速度慢得多的LM甚至GM上，性能急剧下降。
• 翻车现场：忘记LM初始化：有一次，我搬运了数据到LM，但在计算时不小心用了未初始化的LM数组部分，导致结果出现随机值。必须确保LM中用于计算的数据都是已定义的。

结语：从迷宫到通途

曾经，GM、LM、Register对我来说是一个令人困惑的迷宫。但现在，我学会了像一位建筑师一样思考它们的关系。我不再是那个只会拷贝代码的新手，而是开始理解每一个 DataCopy 指令背后的战略意义，开始斟酌 TILE_LENGTH 这个数字背后的权衡。

这套内存体系，是Ascend C能够释放昇腾硬件澎湃算力的基石。理解了它，你才算是真正拿到了打开高性能算子开发大门的钥匙。

在训练营的下一阶段，我们将学习如何让GM到LM的搬运与LM上的计算“同时进行”，这就是更高级的双缓冲流水线技术。我已经迫不及待了。

---

掌握内存模型，是性能优化的第一步。想系统解锁更多Ascend C高阶技能？>> 立即报名2025年CANN训练营第二季