【CUDA编程】CUDA内存模型 test ok

小硕算法工程师

于 2025-04-28 20:37:18 发布

阅读量942

点赞数 17

CC 4.0 BY-SA版权

文章标签：算法

本文链接：https://blog.csdn.net/hfut31415926/article/details/147595094

1. 内存结构

在CUDA中可编程内存的类型有：

寄存器(Registers)
本地内存(Local Memory)
共享内存(Shared Memory)
常量内存(Constant Memory)
纹理内存(Texture Memory)
全局内存(Global Memory)

CUDA中的内存模型分为以下几个层次：

thread：每个线程都用自己的registers（寄存器）和local memory（局部内存）

block：每个线程块(block)内都有自己的shared memory（共享内存），所有线程块内的所有线程共享这段内存资源

grid：每个grid都有自己的global memory（全局内存），constant memory（常量内存）和texture memory（纹理内存），不同线程块的线程都可使用。其中常量内存和纹理内存为只读内存空间。

线程访问这几类存储器的速度是：register > shared memory >Constant Memory > Texture Memory > Local Memory and Global Memory。下面这幅图表示这些内存在计算机架构中的所在层次。

在这里插入图片描述

2. GPU device内存

2.1 寄存器(Registers)

在内核函数中声明且没有其他修饰符修饰的变量通常是存放在GPU的寄存器中，比如下面代码中的线程索引变量i。寄存器通常用于存放内核函数中需要频繁访问的线程私有变量，这些变量与内核函数的生命周期相同，内核函数执行完毕后，就不能再对它们进行访问了。

特点：每个线程私有，速度快

__global__ void VectorAddGPU(const float *const a, const float *const b,
                             float *const c, const int n) {
  int i = blockDim.x * blockIdx.x + threadIdx.x;  //变量i 在寄存器中
  if (i < n) {
    c[i] = a[i] + b[i]; 
  }
}

寄存器是GPU中访问速度最快的内存空间，但是一个SM中寄存器的数量比较有限，一旦内核函数使用了超过硬件限制的寄存器数量，则会使用本地内存来代替多占用的寄存器，这种寄存器溢出的情况会带来性能上的不利影响，实际编程过程中我们应该避免这种情况。

使用nvcc的编译选项maxrregcount可以控制内核函数使用的寄存器的最大数量：