学习目标学完本节你将能够区分 GPU 上各种内存空间寄存器、局部内存、共享内存、全局内存、常量内存、纹理内存了解每种内存的物理位置、访问速度、作用域和生命周期掌握在 Kernel 中声明和使用这些内存的基本方法理解寄存器溢出到局部内存的现象及其影响初步了解共享内存的 Bank Conflict 概念下一节详细展开1. GPU 存储层次总览GPU 的存储层次从最快到最慢延迟从低到高大致如下内存类型物理位置访问延迟作用域生命周期声明方式寄存器SM 内部最快单个线程线程执行期间自动变量共享内存SM 内部很快单个 BlockBlock 执行期间__shared__常量内存设备内存缓存较快全局只读程序执行期间__constant__纹理内存设备内存缓存较快全局只读程序执行期间纹理对象 / 纹理引用局部内存设备内存L1/L2 缓存较慢单个线程线程执行期间自动变量溢出时全局内存设备内存L2 缓存最慢全局程序执行期间cudaMalloc分配的指针缓存说明全局、局部、常量、纹理内存的物理存储都在设备内存显存上但访问时都会经过片上缓存L1/L2。缓存命中率直接决定了实际访问延迟这也是为什么合并访问和空间局部性如此重要。2. 寄存器Registers特点每个线程私有的最快存储位于 SM 的寄存器文件中。分配Kernel 中的自动变量如int a、float b默认优先分配在寄存器中。生命周期与线程相同线程结束时释放。容量限制每个线程可用的寄存器数量有限通常 255 个取决于架构。如果 Kernel 使用了过多局部变量会导致寄存器溢出部分变量被放到局部内存设备内存性能下降。不可显式控制编译器自动决定哪些变量放在寄存器程序员无法直接指定某个变量必须在寄存器中但可以通过限制变量数量和减少依赖来帮助编译器。编译控制参数可以使用-maxrregcount N编译参数限制单线程最大寄存器用量N 为整数人为制造寄存器溢出方便进行性能对照实验。示例__global__ void regKernel(float *out) { float a 1.0f; // 可能放在寄存器 float b 2.0f; // 可能放在寄存器 out[threadIdx.x] a b; }寄存器溢出示例变量过多__global__ void spillKernel(float *out) { float a0 0, a1 1, a2 2, ...; // 大量变量 // 编译器可能将部分变量放入局部内存 }3. 局部内存Local Memory本质局部内存并不是一种独立的物理存储而是设备内存的一部分用于存储无法放入寄存器的自动变量。触发条件寄存器溢出变量太多编译器决定将其放到局部内存。数组下标是动态的如int arr[N]N 在运行时确定编译器无法确定数组元素能否放在寄存器。较大的结构体或数组。性能局部内存的访问延迟接近全局内存因此性能较差。应尽量避免寄存器溢出。作用域单个线程私有但物理上存储在显存中通过 L1/L2 缓存访问。如何判断编译时加上--ptxas-options-v可以看到寄存器和局部内存使用量。示例__global__ void localMemKernel(int *out, int n) { int arr[64]; // 如果编译器认为可以放在寄存器则不会溢出若太大则部分溢出 for (int i 0; i n; i) { arr[i] i; } out[threadIdx.x] arr[0]; }4. 共享内存Shared Memory特点位于 SM 内部速度接近寄存器比全局内存快得多。作用域同一 Block 内的所有线程共享用于线程间数据交换和协作。声明方式使用__shared__关键字。静态共享内存__shared__ float tile[256];动态共享内存在 Kernel 启动时通过第三个参数指定大小Kernel 内用extern __shared__声明。生命周期随 Block 开始而分配Block 结束而释放。容量限制每个 SM 的共享内存容量有限通常 48KB~164KB 可配置每个 Block 可使用的共享内存大小也有限制。同步要求共享内存的写入和读取之间通常需要__syncthreads()来保证数据一致性。Bank Conflict 前置提示共享内存虽然很快但如果多个线程同时访问属于同一个 Bank 的不同地址会发生 Bank Conflict导致访问串行化性能显著下降。下一节会专门讲解 Bank Conflict 的规则和优化方法。示例向量内积的分块求和简化__global__ void sharedMemKernel(float *in, float *out, int N) { __shared__ float sdata[256]; int tid threadIdx.x; int globalIdx blockIdx.x * blockDim.x tid; sdata[tid] (globalIdx N) ? in[globalIdx] : 0.0f; __syncthreads(); // 归约逻辑后续章节详细展开 out[blockIdx.x] sdata[0]; }5. 全局内存Global Memory特点GPU 上容量最大的内存所有线程都能访问但访问延迟最高。分配在 Host 端通过cudaMalloc分配得到设备指针传入 Kernel 使用。访问优化全局内存的访问效率取决于是否满足合并访问Coalesced Access即同一个 Warp 内的线程访问连续的地址。非合并访问会导致多次内存事务性能大幅下降下一节会详细展开。生命周期从cudaMalloc到cudaFree跨所有 Kernel。使用场景存放大规模输入输出数据。示例__global__ void globalMemKernel(float *in, float *out, int N) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx N) { out[idx] in[idx] * 2.0f; // 访问全局内存 } }6. 常量内存Constant Memory特点只读全局作用域带有专用缓存constant cache当同一 Warp 内所有线程读取相同的地址时速度接近寄存器。声明方式在文件作用域使用__constant__关键字然后在 Host 端通过cudaMemcpyToSymbol初始化。容量限制通常为 64KB。适合存放固定参数数组如果数据超过 64KB不能使用__constant__应改用全局内存。适用场景所有线程都读取相同的数据如算法参数、滤波器系数等。注意如果 Warp 内线程读取常量内存的不同地址性能会急剧下降串行化。示例__constant__ float scaleFactor; __global__ void constantKernel(float *in, float *out, int N) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx N) { out[idx] in[idx] * scaleFactor; // 所有线程读相同的 scaleFactor } } // Host 端 float scale 2.5f; cudaMemcpyToSymbol(scaleFactor, scale, sizeof(float));7. 纹理内存Texture Memory特点只读通过纹理缓存访问针对二维空间局部性做了优化支持硬件插值和边界处理。用途图像处理、采样、查找表等。使用方式CUDA 提供了纹理对象Texture Object和纹理引用Texture Reference两种方式现代 CUDA 推荐使用纹理对象。优势对于二维数据的邻近访问如data[x][y]纹理缓存能提高命中率并且自动处理边界如 clamp、wrap。容量限制纹理缓存的大小有限但底层数据存储在全局内存中纹理只是提供了一种优化的访问方式。本节定位纹理属于拓展知识点图像处理场景专用在通用矩阵 GEMM 算子中几乎不用初学阶段了解即可不必深挖。示例简化仅展示概念// 使用纹理对象需要 cudaTextureObject_t __global__ void textureKernel(cudaTextureObject_t tex, float *out, int width, int height) { int x blockIdx.x * blockDim.x threadIdx.x; int y blockIdx.y * blockDim.y threadIdx.y; if (x width y height) { out[y * width x] tex2Dfloat(tex, x, y); // 纹理采样 } }纹理对象的创建和绑定较复杂本节仅了解其存在详细在后续算子开发或存储优化章节展开。8. 代码演示共享内存与全局内存速度对比下面通过一个简单的归约操作对比使用全局内存和共享内存的性能差异概念演示。#include cstdio #include chrono // 使用全局内存进行块内归约简单但低效 __global__ void reduceGlobal(float *in, float *out, int N) { int tid threadIdx.x; int globalIdx blockIdx.x * blockDim.x tid; float sum 0.0f; if (globalIdx N) sum in[globalIdx]; __syncthreads(); // 简化只让线程0写结果 if (tid 0) out[blockIdx.x] sum; } // 使用共享内存进行块内归约高效 __global__ void reduceShared(float *in, float *out, int N) { __shared__ float sdata[256]; int tid threadIdx.x; int globalIdx blockIdx.x * blockDim.x tid; sdata[tid] (globalIdx N) ? in[globalIdx] : 0.0f; __syncthreads(); // 简单归约示意 for (int stride blockDim.x / 2; stride 0; stride 1) { if (tid stride) { sdata[tid] sdata[tid stride]; } __syncthreads(); } if (tid 0) out[blockIdx.x] sdata[0]; } int main() { const int N 1 20; const int blockSize 256; const int gridSize (N blockSize - 1) / blockSize; float *d_in, *d_out; cudaMalloc(d_in, N * sizeof(float)); cudaMalloc(d_out, gridSize * sizeof(float)); // 初始化省略 // 计时全局内存版本 auto start std::chrono::high_resolution_clock::now(); reduceGlobalgridSize, blockSize(d_in, d_out, N); cudaDeviceSynchronize(); auto end std::chrono::high_resolution_clock::now(); printf(Global memory reduce: %f ms\n, std::chrono::durationfloat, std::milli(end - start).count()); // 计时共享内存版本 start std::chrono::high_resolution_clock::now(); reduceSharedgridSize, blockSize(d_in, d_out, N); cudaDeviceSynchronize(); end std::chrono::high_resolution_clock::now(); printf(Shared memory reduce: %f ms\n, std::chrono::durationfloat, std::milli(end - start).count()); cudaFree(d_in); cudaFree(d_out); return 0; }注意本示例是简化示意版本真实工程中的多级归约会做分层优化这里仅用于对比全局内存和共享内存的访存速度差异不必纠结归约算法的完整性。9. 课后练习练习1区分各类内存写出下列变量分别可能存放在哪种内存中int a 5;Kernel 内局部变量__shared__ float tile[128];__constant__ float coeff;float *ptr指向cudaMalloc分配的内存Kernel 内局部数组int arr[100]如果寄存器不足练习2分析寄存器溢出编写一个 Kernel使用 100 个 float 局部变量并简单累加用nvcc --ptxas-options-v编译查看寄存器使用数量和局部内存使用量spill。尝试减少变量数量观察寄存器使用变化。也可以使用-maxrregcount 32限制寄存器数量观察溢出加剧的现象。练习3共享内存同步在共享内存示例中去掉__syncthreads()会怎样修改代码并运行观察结果是否正确理解为什么需要同步。练习4常量内存验证编写一个 Kernel在常量内存中存储一个系数所有线程读取该系数乘以输入。将系数改为不同地址读取例如根据线程索引读取常量数组的不同元素对比性能变化可使用 cudaEvent 计时验证同一 Warp 内读取不同常量地址时性能下降。练习5合并访问初步思考在全局内存访问中如果同一个 Warp 的线程访问连续的地址如in[threadIdx.x]与访问跨步地址如in[threadIdx.x * 2]哪个更快为什么可以编写两个 Kernel 分别实现并计时验证。详细分析在下一节10. 下一步下一节将深入全局内存的合并访问规则、共享内存的 Bank Conflict 问题以及如何通过合理的内存访问模式提升 Kernel 性能。