第4章：计算核心与存储层次

大家好，我是你们的老朋友。今天我们来聊聊芯片里最核心的两个东西——计算核心和存储层次。说白了，就是SM/CU这些计算单元，以及L1/L2 Cache、Shared Memory、Global Memory这些存储结构。

我个人习惯把芯片想象成一个工厂。计算核心就是工人，存储层次就是仓库和传送带。工人干活快不快，很大程度上取决于原材料能不能及时送到手边。嗯，这个比喻虽然简单，但很贴切。

4.1 SM/CU核心详解

先说说计算核心。NVIDIA叫它SM（Streaming Multiprocessor），AMD叫它CU（Compute Unit）。名字不同，本质差不多——都是一组能并行执行线程的硬件单元。

一个SM里通常包含多个CUDA Core（或者叫Shader Core）。比如V100有64个CUDA Core per SM，A100有64个，H100有128个。这些Core共享指令缓存、调度器、寄存器文件等资源。

我在项目中遇到过一个问题：写了一个kernel，理论计算量很大，但实际跑起来就是慢。后来发现是SM里的线程调度出了问题。每个SM能同时调度的线程数有限，你开的线程太多，反而会互相抢资源。

关键指标：

每个SM的最大线程数（比如1024或2048）
每个SM的最大线程块数（比如32或64）
寄存器文件大小（比如256KB）
共享内存大小（比如48KB或96KB，可配置）

你想想看，这些数字决定了你能在单个SM里塞多少计算任务。塞多了，寄存器溢出，性能反而下降。塞少了，SM吃不饱，利用率低。

4.2 存储层次全景

存储层次是另一个大坑。我刚开始做优化时，总觉得显存（Global Memory）够大就行，后来发现延迟才是要命的。

这张图展示了典型的GPU存储层次结构：

看到这个金字塔结构了吗？越往上，延迟越低，但容量越小。越往下，容量越大，但延迟也越高。优化的本质，就是尽量让数据待在金字塔的上层。

4.3 延迟与带宽的真相

我经常跟团队说一句话：延迟是敌人，带宽是朋友。为什么？

Global Memory的延迟是400-800个时钟周期。这意味着你发出一条访存指令，要等400多个周期才能拿到数据。这段时间里，SM只能干等着，或者切换去执行其他线程。

而带宽呢？HBM2e的带宽可以到900 GB/s以上。如果你能连续访问内存，带宽利用率很高。但如果你随机访问，带宽就浪费了。

避坑指南：

我曾经在一个矩阵乘法项目里，把数据放在Global Memory里直接算。结果发现计算单元大部分时间都在等数据。后来我把输入矩阵的tile搬到了Shared Memory里，性能直接提升了5倍。嗯，这就是存储层次优化的威力。

4.4 Shared Memory的正确用法

Shared Memory是每个SM内部的存储空间，延迟只有20-40个周期。比Global Memory快了一个数量级。但它有两个限制：

容量小：通常只有48KB到128KB
SM内共享：不同SM的Shared Memory是隔离的

怎么用好它？我总结了几条经验：

数据复用：如果一个数据会被多个线程用到，就把它搬到Shared Memory里
合并访问：Shared Memory的bank conflict会影响性能，尽量让同一warp的线程访问不同的bank
同步开销：用__syncthreads()同步线程，但别用太多，否则会拖慢速度

举个例子，矩阵乘法里的tiling技术：

// 每个线程块处理一个tile
__shared__ float As[TILE_SIZE][TILE_SIZE];
__shared__ float Bs[TILE_SIZE][TILE_SIZE];

// 循环加载tile到Shared Memory
for (int tile = 0; tile < numTiles; tile++) {
    // 协作加载
    As[tx][ty] = A[row * TILE_SIZE + tx][tile * TILE_SIZE + ty];
    Bs[tx][ty] = B[tile * TILE_SIZE + tx][col * TILE_SIZE + ty];
    __syncthreads();
    
    // 计算
    for (int k = 0; k < TILE_SIZE; k++) {
        sum += As[tx][k] * Bs[k][ty];
    }
    __syncthreads();
}

这段代码里，每个线程只负责加载一个元素，但所有线程共享整个tile的数据。这样，Global Memory的访问次数从O(N^3)降到了O(N^2)。

4.5 L1和L2 Cache的玄机

L1 Cache和Shared Memory在硬件上是同一块SRAM，可以动态配置。比如有的GPU允许你设置48KB Shared Memory + 16KB L1，或者反过来。

L2 Cache是所有SM共享的，延迟200-400周期。它的作用主要是缓存Global Memory的数据。如果你的数据访问有局部性，L2 Cache能帮你省掉很多Global Memory的访问。

注意：

L2 Cache不是你能直接控制的。它由硬件自动管理。但你可以通过调整数据布局来利用它。比如把频繁访问的数据放在连续的内存地址上，这样L2 Cache的命中率会更高。

我记得有一次做图像处理，把输入图像的像素按行优先存储，结果L2 Cache命中率只有60%。改成按块存储后，命中率提升到了90%以上。性能直接翻倍。

4.6 实战中的权衡

说了这么多，到底怎么选？我给大家一个决策树：

数据只用一次：直接放Global Memory，别折腾
数据被多次复用：搬到Shared Memory
数据有空间局部性：靠L1/L2 Cache自动缓存
数据需要跨线程共享：用Shared Memory + __syncthreads()

说白了，没有银弹。每个场景都要具体分析。我习惯先用NVIDIA的Nsight Compute工具跑一遍，看看L1/L2命中率、Shared Memory使用量、Global Memory带宽利用率这些指标。然后针对瓶颈做优化。

嗯，今天就聊到这里。记住一句话：计算核心再快，也快不过数据搬运的速度。优化存储层次，往往比优化计算本身更有效。

本章核心要点：

SM/CU是并行计算的基本单元，注意线程数和资源限制
存储层次从寄存器到Global Memory，延迟差了几个数量级
Shared Memory是手动控制的缓存，适合数据复用场景
L1/L2 Cache自动管理，但数据布局会影响命中率
延迟是敌人，带宽是朋友——尽量让数据靠近计算核心

想要更多资料，联系可可