CUDA内存基础：GPU内存层次、全局内存、共享内存、寄存器、常量内存和纹理内存
在深入内存池设计之前，必须先建立对GPU内存体系的清晰认知。CUDA的内存模型并非单一存储池，而是一个由多种物理存储单元构成的层次化结构。不同内存的容量、延迟、带宽和可见性差异极大，直接决定了内核性能的上限。理解这一层次，是后续实现高效内存复用与TensorRT推理优化的前提。
核心原理：GPU内存层次结构
GPU内存层次可以类比为CPU的缓存体系，但差异更为显著。从硬件物理位置看，寄存器与共享内存在芯片内部（SM内），全局内存、常量内存和纹理内存在显存（DRAM）中，但通过不同的缓存路径访问。下图展示了这一层次关系及关键属性：
从上图可以提炼出几个关键结论：寄存器是速度最快但容量最小的存储，仅当前线程可访问；共享内存是块内线程协作的“手动缓存”，延迟远低于全局内存；全局内存是主存储区，容量大但延迟高，所有线程可见；常量内存和纹理内存是经过特殊缓存优化的只读路径，分别适合广播访问和空间局部性访问。
各内存类型详解与适用场景
| 内存类型 | 作用域 | 生命周期 | 典型用途 | 性能特征 |
|---|---|---|---|---|
| 寄存器 | 单个线程 | 线程执行期间 | 局部变量、中间计算 | 延迟~1 cycle，数量有限（每个线程最多255个） |
| 共享内存 | 线程块 | 块执行期间 | 块内数据共享、数据复用、归约 | 延迟~20 cycles，带宽极高，需避免bank conflict |
| 全局内存 | 所有线程+主机 | 程序运行期间 | 主数据存储、跨块通信 | 延迟400~800 cycles，依赖合并访问(Coalescing) |
| 常量内存 | 所有线程 | 程序运行期间 | 只读参数、系数表 | 广播访问快，不同地址访问会串行化 |
| 纹理内存 | 所有线程 | 程序运行期间 | 图像处理、空间局部性数据 | 通过纹理缓存加速，支持硬件插值 |
方法流程：如何选择与使用内存
在实际CUDA编程中，内存选择应遵循以下决策流程：
- 数据是否只读且所有线程访问相同地址？ → 使用
__constant__常量内存，利用广播机制。 - 数据是否具有空间局部性（相邻线程访问相邻地址）？ → 使用纹理内存或全局内存的合并访问模式。
- 数据是否需要在块内线程间共享或复用？ → 显式使用
__shared__共享内存，将全局数据分块加载。 - 数据是否仅线程私有且访问频繁？ → 尽量放入寄存器，避免数组索引导致的局部内存(Local Memory)溢出。
以下代码展示了共享内存与全局内存的典型配合模式——矩阵分块乘法中的tile加载：
__global__ void matMulShared(const float* A, const float* B, float* C, int N) {
__shared__ float As[TILE][TILE];
__shared__ float Bs[TILE][TILE];
int bx = blockIdx.x, by = blockIdx.y;
int tx = threadIdx.x, ty = threadIdx.y;
int row = by * TILE + ty;
int col = bx * TILE + tx;
float sum = 0.0f;
// 分块循环
for (int t = 0; t < N / TILE; ++t) {
// 协作加载A和B的tile到共享内存
As[ty][tx] = A[row * N + t * TILE + tx];
Bs[ty][tx] = B[(t * TILE + ty) * N + col];
__syncthreads(); // 确保tile加载完成
// 计算当前tile的部分和
for (int k = 0; k < TILE; ++k) {
sum += As[ty][k] * Bs[k][tx];
}
__syncthreads(); // 防止覆盖共享内存
}
C[row * N + col] = sum;
}这段代码的核心思想是：将全局内存中的数据分块（tile）加载到共享内存，块内线程反复从共享内存读取数据，从而将全局内存访问次数从O(N³)降低到O(N³/TILE)。__syncthreads()是块内同步屏障，确保所有线程完成加载后再开始计算。
案例/场景：TensorRT中的内存复用
在TensorRT推理引擎中，内存池设计充分利用了上述层次。以TensorRT 8.x为例，其IExecutionContext在执行时，会将权重存储在常量内存或全局内存的只读区域，而激活值（中间特征图）则通过共享内存和寄存器进行临时计算。TensorRT的setTacticSources()接口允许开发者控制是否启用某些内核策略，这直接影响共享内存的使用量。
一个典型的优化案例：在FP16推理时，TensorRT会将卷积核权重转换为FP16并存入常量内存，利用常量缓存的广播特性加速所有线程对同一权重的读取。同时，对于大尺寸特征图，TensorRT会采用流式内存复用——不同层的输出张量在生命周期不重叠时共享同一块全局内存，从而将峰值内存占用降低40%~60%。
__restrict__关键字声明指针，帮助编译器识别无别名访问，从而更好地利用寄存器缓存。同时，使用cudaFuncSetAttribute()动态调整共享内存大小，以适配不同规模的tile。
误区/复盘：常见内存使用错误
- 误区一：过度使用共享内存而不考虑bank conflict。当多个线程同时访问共享内存中同一bank的不同地址时，访问会被串行化。例如，
shared[tid * 32]会导致所有线程访问同一bank。正确做法是使用shared[tid]或添加padding（如shared[TILE][TILE+1]）。 - 误区二：忽略寄存器溢出。当每个线程使用的寄存器超过硬件限制（通常255个），多余的变量会被“溢出”到局部内存（实际在全局内存中），导致性能骤降。可通过
--maxrregcount编译选项或__launch_bounds__限制寄存器使用。 - 误区三：常量内存不同地址访问。常量内存的广播优化仅在所有线程访问同一地址时生效。若线程访问不同索引（如
coef[threadIdx.x]），访问会被串行化为N次，性能反而低于全局内存。 - 误区四：忽视全局内存合并访问。全局内存的带宽利用率取决于访问模式。线程
tid访问array[tid]是合并的（连续），而访问array[tid * 2]则会产生两倍的事务浪费。