CUDA内存管理：全局内存、共享内存、常量内存、纹理内存

核心原理：CUDA的内存层次结构与带宽真相

CUDA程序性能的瓶颈往往不在计算单元，而在数据搬运。GPU拥有庞大的并行线程，但每个线程访问不同内存空间的速度差异可达两个数量级。理解内存层次，是写出高性能kernel的前提。

从硬件视角看，CUDA内存可划分为片外内存与片上内存两大类。全局内存、常量内存、纹理内存、本地内存均位于显存（DRAM）中，访问延迟高达400~800个时钟周期；而共享内存、寄存器位于芯片内部（SRAM），延迟仅为20~30个周期。下图展示了线程块、线程与各内存空间的访问关系：

全局内存是GPU的“主存”，容量最大（数GB至数十GB），但延迟最高。所有线程（包括不同block中的线程）都能读写全局内存。共享内存则是每个block私有的高速缓存，供块内线程协作。常量内存和纹理内存本质上是全局内存的特殊视图，但硬件为它们提供了独立的缓存与加速路径。

方法流程：四种内存的分配、使用与适用场景

1. 全局内存（Global Memory）

全局内存通过cudaMalloc分配，通过cudaMemcpy在主机与设备间传输数据。访问全局内存时，合并访问（coalesced access）是性能关键——同一warp（32个线程）访问连续地址时，硬件将其合并为少数几次事务；若访问跳跃，则事务数倍增，带宽利用率骤降。

float *d_a;
cudaMalloc(&d_a, N * sizeof(float));
cudaMemcpy(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice);
// kernel中使用 d_a[i]（i为线程全局索引）
cudaFree(d_a);

2. 共享内存（Shared Memory）

共享内存在kernel内部用__shared__修饰符声明，生命周期与block一致。它常用于：数据复用（如矩阵分块）、块内归约、避免重复全局访问。使用共享内存需注意bank conflict——当多个线程同时访问同一bank的不同地址时，访问会被串行化。

__global__ void matMulShared(float *A, float *B, float *C, int n) {
    __shared__ float As[16][16], Bs[16][16];
    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    // 协作加载数据到共享内存
    As[ty][tx] = A[by * 16 + ty][bx * 16 + tx];
    Bs[ty][tx] = B[by * 16 + ty][bx * 16 + tx];
    __syncthreads();
    // 计算...
}

3. 常量内存（Constant Memory）

常量内存容量仅64KB，用__constant__声明，通过cudaMemcpyToSymbol初始化。其最大优势是广播机制：当warp内所有线程访问同一地址时，硬件仅需一次内存事务。适合存储权重、系数、查找表等只读且全体线程访问相同地址的数据。

__constant__ float c_coeff[256];
cudaMemcpyToSymbol(c_coeff, h_coeff, 256 * sizeof(float));

4. 纹理内存（Texture Memory）

纹理内存通过纹理对象（cudaTextureObject_t）绑定全局内存，提供硬件插值、边界处理（clamp/wrap）和二维局部性优化。适合图像处理、空间插值、科学计算中的网格采样。其只读缓存针对空间局部性优化，对二维随机访问友好。

cudaTextureObject_t texObj;
cudaResourceDesc resDesc = {};
resDesc.resType = cudaResourceTypeLinear;
resDesc.res.linear.devPtr = d_data;
// 配置纹理参数（归一化坐标、插值模式等）
cudaCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
// kernel中通过 tex1Dfetch(texObj, idx) 访问

案例/场景：矩阵乘法与图像滤波的对比

以经典的矩阵乘法为例，不使用共享内存时，每个线程需读取A的一行和B的一列，全局内存访问量为2N³次；使用16×16分块后，每个block将数据加载到共享内存，全局访问量降至2N³/16次，性能提升可达10倍以上。这正是共享内存“数据复用”价值的直观体现。

再看图像高斯滤波场景：每个输出像素需要读取周围3×3邻域。若直接使用全局内存，相邻线程的访问区域高度重叠，缓存命中率低。改用纹理内存后，硬件自动利用空间局部性缓存，且边界模式（clamp）免去了手动判断越界的开销，代码更简洁、性能更高。

内存类型	容量	延迟	缓存	典型用途
全局内存	数GB	400-800周期	L2	主数据存储
共享内存	48-228KB/SM	20-30周期	无（SRAM）	块内协作、数据复用
常量内存	64KB	广播时≈1次事务	常量缓存	全体线程同址访问的只读数据
纹理内存	同全局	依赖局部性	纹理缓存	图像、插值、二维空间访问

误区/复盘：常见性能陷阱与优化要点

误区1：共享内存越大越好。共享内存占用过多会降低SM上活跃block数量，导致并行度下降。应通过cudaFuncSetAttribute动态调整，或使用__launch_bounds__限制线程数来平衡。

误区2：忽略bank conflict。例如访问shared[threadIdx.x * 32]时，所有线程访问同一bank，产生32路冲突。可通过填充（padding）或改变数据布局消除。

误区3：常量内存随意用。若warp内线程访问不同地址，常量内存会串行化，性能反而不如全局内存。只有“全体同址”才适合常量内存。

误区4：纹理内存万能。纹理内存的插值功能对整数坐标无效，且创建纹理对象有开销。对于纯线性访问，全局内存合并访问往往更快。

复盘建议：使用nvprof或Nsight Compute查看内存吞吐量、共享内存bank conflict计数、全局内存合并率。若全局内存利用率低于80%，优先检查访问模式是否合并；若共享内存利用率高但性能未提升，检查是否存在bank conflict或同步开销过大。

内存优化是CUDA性能调优的核心环节。掌握四种内存的特性与适用边界，结合具体算法的访问模式选择合适的内存空间，往往能获得数倍乃至数量级的性能提升。实践中应遵循“先合并全局访问，再用共享内存消除冗余，最后考虑特殊内存”的优化路径。