欢迎关注我的公众号 [极智视界],获取我的更多经验分享

大家好,我是极智视界,本文分享一下 CUDA Memory内存模型。

邀您加入我的知识星球「极智视界」,星球内有超多好玩的项目实战源码和资源下载,链接:t.zsxq.com/0aiNxERDq

熟悉和了解 CUDA Memory 内存模型对于开发高效率的 CUDA 程序是非常重要的,这里就来介绍一下这方面的东西。GPU 具有多种不同功能的内存,主要包括:

  • Register => 寄存器内存;
  • Shared Memory => 共享内存;
  • Global Memory => 全局内存;
  • Thread Private Local Memory => 线程私有内存;
  • Constant Memory => 常量内存;
  • Texture Memory => 纹理内存;

下面这张图展示了 Nvidia GPU 内存架构,涵盖了上述不同块的 GPU 内存,如下,

极智开发 | CUDA Memory内存模型

来看 GPU 上内存的分布,这个架构是基于 Thread、Block 和 Grid 线程层次结构的,Thread、Block 和 Grid 组成的 Memory Hiererchy 架构如下,大致的关系是 Grid 由 Block 组成、Block 由 Thread 组成,

极智开发 | CUDA Memory内存模型

然后回到第一张图,在 Block 里面有 Shared Memory,这个 Shared Memory 是可以被 Block 中所有的线程块所共享的。这个意思是比如这里的 Thread(0, 0) 是可以去访问 Shared Memory 的,而 Thread(1, 0)、Thread(2, 0) 也是可以去访问同一个 Shared Memory 的,这意味着同一个 BLock 中的不同 Thread 是可以通过 Shared Memory 来交换数据的。在 Block 中还有 Registers,这和 Shared Memory 不太一样,是每一个线程独有自己专有的 Register。也就是说,Thread(0, 0) 可以访问自己的 Register,但是它没有办法访问 Thread(1, 0) 的 Register。接着还有 Local Memory,Local Memory 也是一种线程所独有的一种内存。然后在 Block 之外、Grid 之中,还有 Global Memory (也就是平时咱们说的 GPU 显存),另外还有 Constant Memory、Texture Memory。这三块内存有一些共同点,比如可以被所有线程所共享、都可以与主机内存进行数据交换等。

具体来看,Register 寄存器内存是 GPU 中访问速度最快的内存,它可以提高程序的执行效率。如果变量在内核中声明且没有使用其他修饰符,那么这个变量通常会被存储在寄存器中,寄存器通常是用来存放频繁访问的变量。寄存器内存是每个线程私有的,不同线程无法通过寄存器内存来进行数据交换。

接着是 Local Memory,内核函数中无法存放到寄存器中的变量就会被存放到本地内存中,比如比较大的数据结构、不满足寄存器内存限制的变量、动态数组等。Local Memory 相对于 Register 来说,访问速度会慢一些。

共享内存 Shared Memory,是在内核函数中由 __shared__ 修饰的变量,也就是用 __shared__ 修饰的变量都会存放在共享内存中。Shared Memory 是 on-chip 的,具有高带宽、低延迟的特点。共享内存是在一个 Block 中的所有线程都能够访问的内存。相比 Register,Shared Memory 的访问速度会稍微慢一些,但 Shared Memory 的容量通常会比 Register 大得多,可以存储更多的数据,因此在某些情况下 Shared Memory 更适合用于数据的存储和共享。

接着是常量内存 Constant Memory,是在内核函数中由 __constant__ 修饰的变量,常量内存是专门用于存储常量数据的内存区域,可以被所有内核代码访问到。它具有高速的访问速度,通常被用于存储内核中需要频繁访问的常量数据。另外一块是纹理内存 Texture Memory,这跟常量内存比较类似,是专门用于存储纹理数据的内存区域。这里可能会有同学会问:神马是纹理数据?纹理数据可以是颜色、亮度、alpha值等,通常用于纹理映射和图像处理等任务,每个数据元素称为一个纹理,通过基本格式和数据类型确定。纹理数据可以被映射到任何几何形状的表面,使得图像能够以更加逼真的方式呈现在屏幕上。

最后是全局内存 Global Memory,也就是 “最亲民的”、咱们平时喊得最顺口的 GPU 显存。对于静态分配的全局内存,用 __device__ 关键字进行标识,而动态分配的全局内存是 Host 端采用显存管理函数进行分配的。全局内存是 GPU 中的主要内存区域,用于存储程序中的变量、数组和其他数据结构。全局内存的访问速度相对较慢,但容量较大,可以容纳大量的数据。全局内存是所有线程共享的,因此多个线程可以同时访问全局内存中的数据。

好了,以上分享了 CUDA Memory内存模型,希望我的分享能对你的学习有一点帮助。



 【公众号传送】

《极智开发 | CUDA Memory内存模型》


畅享人工智能的科技魅力,让好玩的AI项目不难玩。邀请您加入我的知识星球, 星球内我精心整备了大量好玩的AI项目,皆以工程源码形式开放使用,涵盖人脸、检测、分割、多模态、AIGC、自动驾驶、工业等。不敢说会对你学习有所帮助,但一定非常好玩,并持续更新更加有趣的项目。 t.zsxq.com/0aiNxERDq

极智开发 | CUDA Memory内存模型