2. 【CUDA】内存模型：从Global到Shared Memory

GPU 有 10,000 个计算核心，为什么实际吞吐常常只是理论峰值的 10-30%？

答案不在计算，而在内存。

1. 内存墙 —— GPU 的隐藏瓶颈

1.1 算力过剩 vs 带宽不足

想象一个工厂：1000 个工人在流水线上组装零件，但仓库只有 1 个窄门，每秒只能通过 10 箱原材料到达车间。

结果是什么？工人们大部分时间在等原材料，而不是在工作。

瓶颈不是工人不够快，而是门太窄。

GPU 的情况完全一样。看 RTX 4090 的数据：

单次运算的例子

以最简单的向量加法为例：C[i] = A[i] + B[i]

单次运算的需求：
├─ 读取：A[i] (4字节) + B[i] (4字节) = 8字节
├─ 计算：1次浮点加法
└─ 写入：C[i] (4字节)
──────────────────────────────────────
总计：需要搬运 12字节，完成 1次运算

算一下带宽限制下，每秒能做多少次这样的运算：

带宽限制的算力：
1008 GB/s ÷ 12 字节/次 = 84 G次/秒 = 84 GFLOPS

对比：
├─ 带宽能支撑：84 GFLOPS
└─ 芯片能计算：82.6 TFLOPS = 82,600 GFLOPS
──────────────────────────────────────
利用率：84 / 82,600 = 0.1%

核心矛盾：99.9%的计算单元在等内存数据，只有 0.1%在真正工作，这里 GPU 的计算单元比内存快了 1000 倍。

这就是内存墙：即使有 10,000 个核心，如果数据供应不上，大部分核心也只能闲置。

1.2 GPU 如何应对：分层内存系统

既然搬数据这么慢，GPU 的设计思路是什么？

答案是局部性原理：把常用的数据放在快速但小的存储里，不常用的数据放在慢速但大的存储里。

这和人类处理文件的方式一样：

正在处理的文件 → 摊在桌上（最快）
本周要用的文件 → 放办公室书柜（快）
历史档案 → 存档案室（慢但容量大）

GPU 也设计了多层内存，速度差异巨大：

关键数字：

• Register 比 Global Memory 快600 倍
• 但容量小 10 万倍

如果能把数据从 Global Memory 搬到 Shared Memory 或 Register，就能大幅提升性能。

假设某个数据需要被访问 100 次：
├─ 全部从 Global 读取：600 周期 × 100 = 60,000 周期
└─ 搬到 Shared 后访问：600 周期（搬运）+ 30 周期 × 100（访问）= 3,600 周期
────────────────
加速比：16.7×

2. Global Memory —— 合并访问是生命线

为什么相邻线程读相邻地址，能快 32 倍？

2.1 内存事务的真相

GPU 访问 Global Memory 不是一个字节一个字节读的，而是以128 字节为单位。这被称为一次内存事务（Memory Transaction）。

想象超市结账：

情况 1：非合并访问
32 个顾客排队，每人买 1 件商品，分别结账。收银员要扫描 32 次，依次收款找零。

情况 2：合并访问
这 32 个顾客一起结账，收银员一次扫描所有商品，统一收款找零。

第二种情况快 32 倍，因为减少了找零、刷卡等固定开销。

GPU 的内存访问也是如此:

一个 Warp 有 32 个线程，它们会同时发起内存请求。
GPU 的内存控制器会分析这 32 个请求的地址：

- 如果地址连续：32 个线程访问的是连续的 128 字节（32 个 float） → 合并成 1 次事务
- 如果地址分散：可能需要 32 次独立的内存事务

2.2 合并 vs 非合并访问

假设有一个数组，256 个线程访问它：

__global__ void example(float* array) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;

    // 访问模式1：连续访问
    float val = array[tid];

    // 访问模式2：跨步访问
    float val = array[tid * 32];
}

模式 1：连续访问

前 32 个线程（第一个 Warp）访问：

• Thread 0 → array[0]
• Thread 1 → array[1]
• Thread 2 → array[2]
• ...
• Thread 31 → array[31]

这 32 个地址刚好是连续的 128 字节（32 × 4 字节），可以合并成1 次内存事务。

模式 2：跨步访问

前 32 个线程访问：

• Thread 0 → array[0]
• Thread 1 → array[32]
• Thread 2 → array[64]
• ...
• Thread 31 → array[992]

每两个地址之间相隔 128 字节，这 32 个地址分散在 32 个不同的 cache line 里，需要32 次内存事务。

真实测试结果（RTX 4090，向量加法）：

步长为 32 时，性能只有合并访问的 3%——慢了32.8 倍。

2.3 什么情况会破坏合并？

以下是常见的访问模式及其效率：

关键规律：Warp 内 32 个线程访问的地址越分散，需要的内存事务越多，效率越低。

2.4 访问模式的影响

回到最简单的向量加法：

__global__ void vectorAdd(float* A, float* B, float* C, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    C[i] = A[i] + B[i];
}

这个 kernel 在 RTX 4090 上能达到 920 GB/s（91%峰值带宽），因为：

• Warp 内 32 个线程访问A[0]到A[31]（连续）
• Warp 内 32 个线程访问B[0]到B[31]（连续）
• Warp 内 32 个线程写入C[0]到C[31]（连续）

关键：索引计算i = threadIdx.x + blockIdx.x * blockDim.x保证了 Warp 内线程访问的地址是连续的。

但如果算法本身要求非连续访问（如矩阵转置、图遍历），就无法避免性能损失。

3. L1/L2 Cache —— 隐形的加速器

为什么做了合并访问，有时带宽还是只有理论峰值的 80%？

这时候要了解 GPU 自动缓存系统的原理。

3.1 Cache 是什么

Cache 是 GPU 在 Global Memory 和 SM 之间自动插入的缓存层，程序员不需要显式管理。

想象你在图书馆学习：

需要的书：放在书包里（最快）      ← L1 Cache
常用的书：放在书架上（快）        ← L2 Cache
所有的书：在图书馆仓库里（慢）    ← Global Memory

每次你要读一本书：

1. 先看书包里有没有（L1 Cache）—— 1 秒
2. 没有就去书架找（L2 Cache）—— 10 秒
3. 还没有就去仓库取（Global Memory）—— 60 秒

GPU 的 Cache 也是类似的自动机制：

关键特点：

• 自动管理：程序员不需要（也无法）显式地把数据"放入"Cache
• 透明性：访问 Global Memory 时，GPU 自动检查 Cache
• Cache Line：每次从 Global 加载数据到 Cache，以 128 字节为单位（和内存事务大小一致）

3.2 Cache 如何工作

Cache 的核心概念是 Cache Hit（命中） 和 Cache Miss（未命中）。

Cache Hit：要的数据已经在 Cache 里

线程请求：读取 array[100]
GPU检查：
├─ L1有吗？有！ → 30 周期返回 ✅
└─ 结果：快！

Cache Miss：要的数据不在 Cache 里

线程请求：读取 array[100]
GPU检查：
├─ L1有吗？没有
├─ L2有吗？没有
└─ 去 Global 取 → 600 周期 + 加载到 L1/L2 ❌

性能差异有多大？

一个程序如果 Cache 命中率从 90% 降到 50%，性能可能下降 2-3 倍。

Cache Line 的影响

Cache 不是一个字节一个字节加载的，而是以 128 字节的 Cache Line 为单位。

你访问 array[0]（4字节）：
GPU 实际做的事：
├─ 从 Global 加载 array[0:31]（128字节）到 Cache Line
└─ 返回 array[0] 给你

副作用：array[1] 到 array[31] 也在 Cache 里了

这就是为什么合并访问能提高 Cache 效率。

3.3 合并访问为什么能提高 Cache 命中率

我们回到 2.2 节的例子，从 Cache 的角度重新审视：

场景 1：连续访问

float val = array[tid];  // Thread 0→array[0], Thread 1→array[1], ...

Warp 内 32 个线程访问 array[0:31]：

第一次访问 array[0]：
├─ Cache Miss → 从 Global 加载 128 字节到 L1
└─ array[0:31] 都进入 Cache ✅

后续 31 个线程访问 array[1:31]：
├─ Cache Hit！（数据已经在 L1 里）
└─ 只需 30 周期

结果：
├─ 1 次 Global 访问（600 周期）
├─ 31 次 L1 Hit（30 周期 × 31）
└─ 平均：(600 + 30×31) / 32 ≈ 48 周期/线程

场景 2：跨步访问（步长=32）

float val = array[tid * 32];  // Thread 0→array[0], Thread 1→array[32], ...

Warp 内 32 个线程访问的地址：

Thread 0  → array[0]   → Cache Line 0
Thread 1  → array[32]  → Cache Line 1
Thread 2  → array[64]  → Cache Line 2
...
Thread 31 → array[992] → Cache Line 31

结果：
├─ 32 次 Cache Miss（每个线程触发一次 Global 加载）
├─ 每个 Cache Line 只用了 1 个元素（浪费了其余 31 个）
└─ 平均：600 周期/线程

性能对比（相同的 32 次访问）：

关键洞察：

• 合并访问 = 高 Cache 命中率：一次加载，32 个线程受益
• 非合并访问 = 低 Cache 命中率：32 次加载，每次只用 1 个元素

这解释了为什么 2.2 节的测试中，步长=32 的性能只有连续访问的 3%（不仅事务多，Cache 也完全失效）。

3.4 局部性原理

Cache 能发挥作用，依赖于两种局部性：

空间局部性（Spatial Locality）：访问了 array[i]，很可能马上访问 array[i+1]

好的空间局部性：
for (int i = 0; i < N; i++) {
    sum += array[i];  // 连续访问 ✅
}

坏的空间局部性：
for (int i = 0; i < N; i++) {
    sum += array[random()];  // 随机跳跃 ❌
}

时间局部性（Temporal Locality）：访问了 array[i]，很可能在短时间内再次访问

好的时间局部性：
for (int k = 0; k < 100; k++) {
    sum += array[i] * weight[k];  // array[i] 被重复访问 ✅
}

坏的时间局部性：
for (int i = 0; i < N; i++) {
    result[i] = array[i];  // 每个元素只访问 1 次 ❌
}

GPU 程序的 Cache 效率主要看空间局部性（因为单个线程访问的数据量小，时间局部性不明显）。

3.5 实测：Cache 的影响

简单的向量加法，三种数据访问模式：

// 模式1：顺序访问（好的空间局部性）
__global__ void sequential(float* A, float* B, float* C, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    C[i] = A[i] + B[i];
}

// 模式2：步长访问（中等空间局部性）
__global__ void strided(float* A, float* B, float* C, int N, int stride) {
    int i = (threadIdx.x + blockIdx.x * blockDim.x) * stride;
    C[i] = A[i] + B[i];
}

// 模式3：随机访问（无空间局部性）
__global__ void random_access(float* A, float* B, float* C, int* indices, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    int idx = indices[i];  // 随机索引
    C[i] = A[idx] + B[idx];
}

性能测试（RTX 4090，100M 元素）：

关键发现：

• 顺序访问几乎所有数据都在 L1 命中
• 步长访问部分命中 L1，部分走 L2
• 随机访问 Cache 基本失效，大部分走 Global

3.6 程序员能做什么

虽然 Cache 是自动管理的，但程序员可以通过访问模式影响效率：

不能做的事：

• ❌ 不能显式控制哪些数据进入 Cache
• ❌ 不能调整 Cache 大小（L1 固定 128 KB）
• ❌ 不能预取数据到 Cache（GPU 有自动预取）

实用建议：

1. 索引计算要连续

✅ int i = threadIdx.x + blockIdx.x * blockDim.x;
❌ int i = threadIdx.x * stride + blockIdx.x;

2. 避免间接寻址（除非必需）

✅ float val = data[i];           // 直接访问
❌ float val = data[indices[i]];  // 间接寻址，可能破坏局部性

3. 结构体成员访问

// 如果结构体很大，考虑 AoS → SoA 转换
❌ struct Point { float x, y, z; };  // AoS: x,y,z 交错存储
✅ struct Points { float* x; float* y; float* z; };  // SoA: 连续存储

一句话记住：

保证合并访问 = 同时优化了 内存事务 + Cache 效率

4. Shared Memory —— 片上高速缓存

4.1 它是什么

Shared Memory 是位于 SM 芯片上的高速 SRAM，专门用于 Block 内的线程共享数据，比 Global 快 20 倍。

用办公楼的快递柜类比：

- Global Memory = 市区的物流中心（远但容量大）
- Shared Memory = 办公楼 1 层的快递柜（近但容量小）

如果你今天要收 10 个包裹：
├─ 每次都跑市区取 → 耗时 10 趟
└─ 一次性搬到楼下快递柜 → 之后随时取，耗时 1 趟

Shared Memory 就是这个"快递柜"，让 Block 内的线程可以快速共享数据。

和 Global Memory 的对比：

Shared Memory 快在哪里？

• 物理上在 SM 内部，不需要走 HBM 总线
• 带宽是 Global 的 10 倍
• 延迟是 Global 的 1/20

但它也有限制：

• 容量小（每个 SM 只有 48-128 KB）
• 只在 Block 内可见，无法跨 Block 共享
• Kernel 结束后数据消失

4.2 Shared Memory 的标准使用流程

使用 Shared Memory 的典型模式是三步走：

__shared__ float cache[BLOCK_SIZE];  // 声明Shared Memory

// Step 1: 从Global加载到Shared（合并访问）
int tid = threadIdx.x;
int globalIdx = tid + blockIdx.x * blockDim.x;
cache[tid] = input[globalIdx];
__syncthreads();  // 等待所有线程加载完成

// Step 2: 在Shared上多次读写（快！）
float sum = 0;
for (int i = 0; i < BLOCK_SIZE; i++) {
    sum += cache[i] * weight[i];  // 反复访问cache
}

// Step 3: 写回Global（合并访问）
output[globalIdx] = sum;

关键点：

1. __syncthreads()：确保所有线程都完成加载后，才能开始使用 Shared Memory
2. 多次访问：如果数据只用 1 次，不值得搬到 Shared
3. 合并访问 Global：Step 1 和 Step 3 都要保证合并访问

4.3 Bank Conflict —— Shared Memory 的隐藏陷阱

为什么同样用 Shared Memory，有时性能会突然掉到 1/32？

什么是 Bank

Shared Memory 并不是一整块存储，而是被分成了 32 个独立的 Bank，每个 Bank 可以独立响应请求。

想象一个银行有 32 个柜台：

银行大厅（Shared Memory，128 KB）
├─ 柜台 0（Bank 0）：处理账号末尾为 0 的客户
├─ 柜台 1（Bank 1）：处理账号末尾为 1 的客户
├─ 柜台 2（Bank 2）：处理账号末尾为 2 的客户
...
└─ 柜台 31（Bank 31）：处理账号末尾为 31 的客户

地址到 Bank 的映射规则：

Bank ID = (地址 / 4 字节) % 32

示例：
├─ float arr[0]  → 地址 0   → Bank 0
├─ float arr[1]  → 地址 4   → Bank 1
├─ float arr[2]  → 地址 8   → Bank 2
...
├─ float arr[31] → 地址 124 → Bank 31
└─ float arr[32] → 地址 128 → Bank 0（循环）

关键规律：连续的 32 个 float 分布在 32 个不同的 Bank 上。

什么是 Bank Conflict

当一个 Warp 的 32 个线程同时访问 Shared Memory 时：

情况 1：无冲突（理想）

__shared__ float data[32];
int tid = threadIdx.x;
float val = data[tid];  // Thread 0→Bank 0, Thread 1→Bank 1, ...

32 个线程访问 32 个不同的 Bank，可以并行完成，1 个时钟周期。

32个柜台同时服务32个客户
时间：1 周期

情况 2：有冲突（糟糕）

__shared__ float data[32];
int tid = threadIdx.x;
float val = data[tid * 2];  // Thread 0→Bank 0, Thread 1→Bank 2, ...

等等，Thread 0 和 Thread 16 会怎样？

• Thread 0 访问 data[0] → Bank 0
• Thread 16 访问 data[32] → Bank 0（32 % 32 = 0）

两个线程访问同一个 Bank 的不同地址！

2个客户同时找柜台0，但办理不同业务
柜台0只能依次服务 → 需要2个周期

如果有8个线程冲突 → 需要8个周期

Bank Conflict 的性能影响

让我们看真实代码的对比：

__global__ void test_conflict(float* output) {
    __shared__ float cache[1024];
    int tid = threadIdx.x;  // 256个线程

    // 模式1：无冲突
    float val1 = cache[tid];

    // 模式2：2路冲突
    float val2 = cache[tid * 2];

    // 模式3：8路冲突
    float val3 = cache[tid * 8];

    // 模式4：32路冲突（最糟）
    float val4 = cache[0];  // 所有线程访问同一地址
}

性能测试（RTX 4090，每个访问重复 1000 次）：

关键规律：N 路冲突 = 慢 N 倍。最坏情况下，Shared Memory 的速度会降到 Global Memory 的水平！

常见的 Bank Conflict 场景

如何诊断 Bank Conflict？

使用 Nsight Compute 查看：

ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum ./program

关键指标：
- Shared Memory Bank Conflicts: 冲突总次数
- 如果 >0，说明存在 Bank Conflict
- 数值越大，性能损失越严重

如何避免 Bank Conflict？

技巧 1：调整数据布局（Padding）

矩阵转置的经典问题：

// 有冲突的版本
__shared__ float tile[32][32];
tile[threadIdx.x][threadIdx.y] = input[...];  // 写入无冲突
__syncthreads();
output[...] = tile[threadIdx.y][threadIdx.x];  // 读取有32路冲突！

为什么读取有冲突？

• Warp 内 32 个线程的 threadIdx.x 相同（比如都是 0）
• 它们访问 tile[0][0], tile[1][0], tile[2][0], ..., tile[31][0]
• 列访问导致步长 = 32，全部映射到同一个 Bank

解决方法：Padding（加一列）

__shared__ float tile[32][33];  // 注意：33列！
tile[threadIdx.x][threadIdx.y] = input[...];
__syncthreads();
output[...] = tile[threadIdx.y][threadIdx.x];  // 现在无冲突✅

为什么加一列就解决了？

原来（32列）：
Row 0: Bank 0,  1,  2, ..., 31
Row 1: Bank 0,  1,  2, ..., 31  ← 列0都在Bank 0（冲突）
...

现在（33列）：
Row 0: Bank 0,  1,  2, ..., 31, 0
Row 1: Bank 1,  2,  3, ..., 32, 1  ← 列0分布在不同Bank（无冲突）
Row 2: Bank 2,  3,  4, ..., 0,  2
...

性能对比：

技巧 2：改变访问顺序

如果无法改变数据布局，可以改变访问模式：

// 有冲突：步长=16
for (int i = 0; i < 256; i += 16) {
    sum += cache[tid + i];
}

// 无冲突：重新组织循环
for (int i = 0; i < 16; i++) {
    sum += cache[tid * 16 + i];
}

技巧 3：利用广播机制

所有线程读取同一地址时，GPU 会用广播机制，无 Bank Conflict：

float coeff = cache[0];  // 32个线程都读cache[0]
float result = cache[tid] * coeff;  // 无冲突

但如果是写入同一地址，仍需注意数据竞争（用 atomic）。

4.6 Shared Memory vs Cache

Shared Memory 和 Cache 都是快速存储，但用途不同：

何时依赖 Cache？

// 简单的数据处理，无需线程通信
__global__ void process(float* data, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    float val = data[i];  // ✅ Cache 自动优化
    data[i] = val * 2.0f;
}

何时用 Shared Memory？

不是所有场景都适合 Shared Memory。判断标准：

简单的判断公式：

收益 = (节省的Global访问次数) × (Global延迟)
开销 = (加载到Shared的时间) + (从Shared读取的时间)

如果 收益 > 开销，就用Shared Memory

对于访问次数 N：

• N = 1：不用 Shared，直接访问 Global
• N = 2-3：可能持平，视情况而定
• N ≥ 5：应该用 Shared，收益明显

// 需要 Block 内通信或数据被多次访问
__global__ void reduce(float* data, float* result, int N) {
    __shared__ float cache[256];  // ✅ 显式管理
    cache[threadIdx.x] = data[...];
    __syncthreads();
    // Block 内归约...
}

简单判断：

• 数据只访问 1 次 + 顺序访问 → 依赖 Cache
• 数据被多次访问 或 需要线程通信 → 用 Shared Memory

让我们看两个简单的实战案例。

5. 实战案例

实际代码中如何应用这些原则？

5.1 案例 1：向量点积（Block 内归约）

问题：计算两个向量的点积 A·B = Σ(A[i] × B[i])

普通版本：每个线程直接累加

__global__ void dot_naive(float* A, float* B, float* result, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) {
        atomicAdd(result, A[i] * B[i]);  // 所有线程竞争同一个地址
    }
}

问题：上万个线程同时对 result 做原子加法，造成严重的序列化。

优化思路：用 Shared Memory 做两级归约

1. 每个 Block 内部先求和（用 Shared Memory）
2. 每个 Block 只写一次 Global（用 atomic）

__global__ void dot_shared(float* A, float* B, float* result, int N) {
    __shared__ float cache[256];  // 假设Block大小=256

    int tid = threadIdx.x;
    int i = tid + blockIdx.x * blockDim.x;

    // Step 1: 每个线程计算一个乘积，存入Shared
    cache[tid] = (i < N) ? A[i] * B[i] : 0;
    __syncthreads();

    // Step 2: Block内归约（树形求和）
    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) {
            cache[tid] += cache[tid + stride];
        }
        __syncthreads();
    }

    // Step 3: Block 0 代表写回Global
    if (tid == 0) {
        atomicAdd(result, cache[0]);
    }
}

工作原理（以 8 个线程为例）：

初始状态：cache = [1, 2, 3, 4, 5, 6, 7, 8]

第1轮（stride=4）：
  Thread 0: cache[0] += cache[4]  → cache[0] = 1+5 = 6
  Thread 1: cache[1] += cache[5]  → cache[1] = 2+6 = 8
  Thread 2: cache[2] += cache[6]  → cache[2] = 3+7 = 10
  Thread 3: cache[3] += cache[7]  → cache[3] = 4+8 = 12
  结果：cache = [6, 8, 10, 12, 5, 6, 7, 8]

第2轮（stride=2）：
  Thread 0: cache[0] += cache[2]  → cache[0] = 6+10 = 16
  Thread 1: cache[1] += cache[3]  → cache[1] = 8+12 = 20
  结果：cache = [16, 20, 10, 12, 5, 6, 7, 8]

第3轮（stride=1）：
  Thread 0: cache[0] += cache[1]  → cache[0] = 16+20 = 36
  结果：cache = [36, 20, 10, 12, 5, 6, 7, 8]

最终cache[0] = 36 = 1+2+3+4+5+6+7+8 ✅

性能对比（1M 个元素）：

为什么快这么多？

• 普通版本：100 万次 atomic 操作，高度串行化
• Shared 版本：只有约 4000 次 atomic 操作（每个 Block 1 次）

5.2 案例 2：1D 卷积（数据复用）

问题：对数组做简单的 3 点平均

output[i] = (input[i-1] + input[i] + input[i+1]) / 3

直接实现：

__global__ void smooth_naive(float* input, float* output, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x + 1;
    if (i < N - 1) {
        output[i] = (input[i-1] + input[i] + input[i+1]) / 3.0f;
    }
}

看起来很简单，但有个隐藏问题：数据重复读取。

考虑一个 Block 内 256 个线程：

• Thread 0 读取 input[0], input[1], input[2]
• Thread 1 读取 input[1], input[2], input[3]
• Thread 2 读取 input[2], input[3], input[4]
• ...

每个 input[i]被相邻的 3 个线程读取，总共从 Global 读取了 256×3 = 768 次。

Shared Memory 优化：把数据加载到 Shared(用__shared__声明)，每个元素只从 Global 读 1 次

#define BLOCK_SIZE 256

__global__ void smooth_shared(float* input, float* output, int N) {
    __shared__ float tile[BLOCK_SIZE + 2];  // 多加载2个边界元素

    int tid = threadIdx.x;
    int i = tid + blockIdx.x * BLOCK_SIZE + 1;

    // Step 1: 加载到Shared（包括halo区域）
    tile[tid + 1] = input[i];

    // 边界线程负责加载额外的halo元素
    if (tid == 0) {
        tile[0] = input[i - 1];  // 左边界
    }
    if (tid == BLOCK_SIZE - 1) {
        tile[BLOCK_SIZE + 1] = input[i + 1];  // 右边界
    }
    __syncthreads();

    // Step 2: 从Shared计算
    if (i < N - 1) {
        output[i] = (tile[tid] + tile[tid + 1] + tile[tid + 2]) / 3.0f;
    }
}

Halo 区域示意图（8 个线程为例）：

Block处理input[8:16]，但需要访问input[7:17]

Global Memory:
  ... [7] [8] [9] [10] [11] [12] [13] [14] [15] [16] ...
       ↑   ↑                                   ↑    ↑
       |   |        主体数据                   |    |
      halo                                    halo

Shared Memory (tile):
  [0] [1] [2]  [3]  [4]  [5]  [6]  [7]  [8] [9]
   ↑                                          ↑
  左halo                                    右halo

性能对比（1M 个元素）：

为什么快？

• 普通版本：每个元素从 Global 读 3 次
• Shared 版本：每个元素从 Global 读 1 次，从 Shared 读 3 次

Shared 的读取延迟只有 Global 的 1/20，即使多读了几次 Shared，总体仍然更快。

5.3 两个案例的共同点

这两个优化都遵循相同的模式：

1. 识别数据复用：

   • 案例 1：Block 内所有线程都需要访问 cache 数组
   • 案例 2：相邻线程读取重叠的数据

2. 用 Shared 做中转：

   • 从 Global 合并加载到 Shared（1 次）
   • 从 Shared 多次读取（快）

3. 收益计算：

   • 案例 1：减少 100 万次 atomic → 4000 次 atomic
   • 案例 2：减少 3M 次 Global 读取 → 1M 次 Global 读取

核心原则：当数据被多个线程访问 ≥2 次时，考虑 Shared Memory。

6. 实用建议

实际开发中如何选择内存策略？

6.1 决策流程图

开始写Kernel
    ↓
需要线程间通信/同步？
    ├─ 是 → 必须用Shared Memory
    └─ 否 → 继续判断
            ↓
        数据被访问≥3次？
            ├─ 是 → 考虑Shared Memory
            └─ 否 → 继续判断
                    ↓
                能否保证合并访问？
                    ├─ 是 → 直接用Global Memory
                    └─ 否 → 用Shared Memory做中转

6.2 常见错误及修复

6.3 性能分析工具

编译时查看资源使用：

nvcc -Xptxas -v kernel.cu

输出示例：
ptxas info: Used 32 registers, 4096 bytes smem, 0 bytes lmem
                    ↑ 寄存器    ↑ Shared Memory

运行时性能分析：

nsys profile --stats=true ./program

关键指标：
- Memory Throughput: 目标 >80% 峰值
- Achieved Occupancy: 目标 >50%
- Uncoalesced Global Accesses: 目标 <5%

Nsight Compute 查看详细信息：

ncu --set full ./program

重点关注：
- SOL Memory: 内存瓶颈百分比（目标<70%）
- Global Load/Store Throughput: 合并访问效率
- Shared Memory Conflicts: Bank冲突率（目标0%）

6.4 优化的优先级

不要一开始就优化所有细节，按这个顺序逐步优化：

1. 首先：确保合并访问（影响最大，10-32×）
2. 其次：使用 Shared Memory（有数据复用时，2-5×）
3. 然后：调整 Block 大小（影响 Occupancy，1.2-2×）
4. 最后：消除 Bank Conflict（影响较小，1.1-1.5×）

每一步都要用 profiler 验证，不要盲目优化。

7. 总结

三张核心表格

表格 1：内存层次对比

表格 2：优化效果对比

表格 3：何时使用 Shared Memory

五条核心要点

1. 内存带宽是 GPU 性能的真正瓶颈

2. 合并访问是 Global Memory 的生命线
   非合并访问会导致 32 倍性能损失。Warp 内线程应访问连续地址。

3. Shared Memory 是数据复用的关键工具
   当数据被多次访问时，用 Shared Memory 做中转可以加速 2-5 倍。

4. 不是所有场景都需要 Shared Memory
   单次访问的数据直接用 Global 更简单。盲目使用 Shared 可能降低 Occupancy。

5. 优化必须用 profiler 验证
   理论上的优化不一定实际有效。每次改动都要实测性能，避免过早优化。