发帖
雷达卡
大家好，我是继续在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训练营第二季
京公网安备 11010802022788号
