一、线程块内存的核心类型及特性

首先明确各类内存的基本属性，这是优化的基础：

内存类型	访问范围	速度（典型延迟）	容量限制	关键作用
共享内存	线程块内所有线程	~100ns	每个 SM 48KB~16384KB（可配置）	线程间通信、缓存全局内存数据以减少访问次数
寄存器	单个线程私有	~1ns	每个线程 16~255 个 32 位寄存器（SM 总寄存器共享）	存储高频访问的临时变量、中间结果
本地内存	单个线程私有	~1000ns	无固定限制（物理上是全局内存的一部分）	寄存器溢出时的 “备胎”，仅用于大数组或动态索引
全局内存	所有线程 / 主机	~500ns	GB 级（显卡显存）	线程块间数据传递、长期存储

二、共享内存：线程块的 “高速缓存与通信中枢”

共享内存是线程块内性能优化的 “重中之重”，其核心价值是减少全局内存访问（通过缓存复用）和支持线程间协作。优化需围绕 “避免银行冲突”“最大化数据复用” 展开。

1. 分配策略：静态优先，动态按需使用

静态分配：编译时确定大小，性能最优，适合尺寸固定的场景。
优势：编译器可提前优化银行映射，无运行时分配开销，适合矩阵分块、滑动窗口等固定尺寸场景。

__shared__ float s_cache[256]; // 静态分配256个float（1KB）

动态分配：运行时根据线程块维度动态确定大小，适合尺寸可变的场景（如依赖blockDim）。

extern __shared__ float s_cache[]; // 声明外部共享内存 
// 启动内核时指定大小：kernel<<<grid, block, 256*sizeof(float)>>>(...);

注意：动态分配需在内核启动时显式指定字节数，同一内核中多个动态共享内存变量需通过偏移区分（如s_cache用于数据，s_cache + 256用于索引）。

2. 使用技巧：避免冲突 + 高效复用

技巧 1：通过 Padding 消除银行冲突
共享内存被划分为多个 “银行”（硬件存储单元，如 32 个 / 64 个），同一银行的不同地址被多个线程同时访问时会导致序列化访问（冲突）。
例如，32×32 的共享内存矩阵按列访问时，未填充会导致同一列映射到同一银行：

// 未优化：列访问冲突（32个线程访问同一银行） 
__shared__ float s_mat[32][32]; 
float val = s_mat[threadIdx.y][threadIdx.x]; // 列访问时冲突 

// 优化：填充1列，打破地址映射的周期性 
__shared__ float s_mat[32][33]; // 宽度33而非32 
float val = s_mat[threadIdx.y][threadIdx.x]; // 列访问时分散到不同银行

原理：填充后，地址计算公式从row×32 + col变为row×33 + col，模 32 后结果分散，避免集中到同一银行。

技巧 2：分块加载全局内存，实现数据复用
全局内存访问延迟高，通过共享内存分块缓存数据，减少重复访问。例如向量点积中，线程块先加载子数组到共享内存，再重复使用：

__global__ void dotProduct(float *result, float *a, float *b, int N) {
    __shared__ float s_a[256], s_b[256]; // 共享内存缓存子数组
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 加载全局内存到共享内存（1次全局访问）
    s_a[threadIdx.x] = a[idx];
    s_b[threadIdx.x] = b[idx];
    __syncthreads(); // 确保所有线程加载完成
    
    // 在共享内存中计算（多次复用，无全局访问）
    float sum = 0;
    for (int i = 0; i < blockDim.x; i++) {
        sum += s_a[i] * s_b[i]; 
    }
    // ... 汇总结果
}

效果：将全局内存访问次数从blockDim.x次减少到 1 次，带宽利用率提升 256 倍（假设块大小 256）。

技巧 3：最小化同步开销
__syncthreads()是线程块内同步的核心指令，但会打断流水线，需减少使用频率：
合并连续的同步操作（如将两次同步合并为一次）；
避免在循环内同步（可将循环拆解为 “加载→计算→同步” 三阶段）。

三、寄存器：线程私有的 “零延迟存储”

寄存器是线程最快的内存，但其容量有限（单个线程通常最多 255 个 32 位寄存器）。优化核心是 “避免溢出，高频复用”。

1. 分配原则：优先存储高频访问数据

寄存器由编译器自动分配，但可通过代码结构引导优化：
避免定义大尺寸私有数组（如float temp[1024]会直接导致溢出到本地内存）；
将循环计数器、中间结果等高频访问变量存于寄存器，低频数据放入共享内存。
反例：大寄存器数组导致溢出

__global__ void badKernel() {
    float large_arr[512]; // 单个线程使用512个寄存器，远超硬件上限
    // 编译器将超出部分放入本地内存（全局内存），访问延迟从1ns增至1000ns
}

优化：用共享内存存储大数组，寄存器仅存索引和临时结果

__global__ void goodKernel() {
    __shared__ float s_large[512]; // 共享内存存储大数组
    float temp; // 寄存器仅存临时变量
    // 通过threadIdx访问s_large，避免溢出
}

2. 技巧：利用线程束洗牌指令（Shuffle）替代共享内存通信

现代 GPU 支持线程束内寄存器直接通信（无需共享内存），通过__shfl_sync等指令实现，延迟更低且无银行冲突：

__global__ void shuffleExample() {
    int lane = threadIdx.x % 32; // 线程束内ID（0~31）
    int data = lane * 2; // 寄存器数据
    
    // 线程0将数据广播给线程束内所有线程（无需共享内存）
    int broadcast_data = __shfl_sync(0xFFFFFFFF, data, 0); 
}

优势：比 “写入共享内存 + 同步 + 读取” 的流程快 5~10 倍，适合线程束内小数据通信。

四、本地内存：尽量规避，仅作 “最后选择”

本地内存是线程私有的内存，但物理上属于全局内存（访问延迟与全局内存相当），仅在以下情况被使用：

• 寄存器溢出（线程使用的寄存器超过硬件上限）；
• 编译器无法确定地址的数组（如动态索引的大数组，如arr[threadIdx.x * k]）。

规避技巧：

• 用固定索引访问小数组（如arr[3]而非arr[threadIdx.x]），帮助编译器分配到寄存器；
• 大数组改用共享内存（线程块内共享）或分块处理（每次处理寄存器可容纳的部分）。

五、全局内存：线程块访问的 “外部数据源”，优化合并访问

线程块通过全局内存与外部交互，其性能取决于 “合并访问”—— 线程束的连续地址请求被合并为一次内存事务（128 字节），非连续访问会拆分为多次事务，带宽利用率骤降。

优化技巧：

确保访问地址连续：线程i访问global_arr[blockIdx.x * blockDim.x + threadIdx.x]，避免跨步访问（如global_arr[threadIdx.x * 2]）。

// 未优化：跨步访问，32个线程的请求被拆分为16次事务
float val = global_arr[threadIdx.x * 2];

// 优化：连续访问，合并为1次事务
float val = global_arr[blockIdx.x * blockDim.x + threadIdx.x];

匹配数据布局与访问模式：多维数组按 “行优先” 存储（匹配 GPU 内存的线性布局），避免列优先访问导致的非连续地址。

// 二维数组行优先存储：global_mat[row][col] = global_mat[row * cols + col]
float val = global_mat[threadIdx.y * cols + threadIdx.x]; // 行访问（连续）

利用内存对齐：确保全局内存变量按 128 字节对齐（通过__align__(128)），提升合并效率。

六、核心原则总结

1. 共享内存：优先缓存复用，坚决避免冲突
   • 用 Padding 或调整访问模式消除银行冲突；
   • 分块加载全局内存，减少重复访问；
   • 同步操作按需使用，避免冗余。
2. 寄存器：够用即可，高频优先
   • 避免大数组导致溢出到本地内存；
   • 线程束内通信优先用 Shuffle 指令，替代共享内存。
3. 全局内存：合并访问是生命线
   • 确保线程束访问地址连续，避免跨步；
   • 数据布局匹配访问模式（行优先）。
4. 资源平衡：线程块大小与 SM 资源匹配
   • 线程块大小（如 128~256）需使 SM 能同时驻留多个块（2~4 个），提升调度效率；
   • 共享内存和寄存器使用量需留有余地，避免 SM 资源耗尽导致块驻留数下降。

通过以上策略，可充分发挥线程块内各类内存的优势，将性能瓶颈从 “内存访问” 转移到 “计算能力”，实现 CUDA 程序的高效并行。