cuda内存种类

编程层面的内存种类

---

前言

要完整而系统的学习

---

一、全局内存（Global Memory）

• 特性：全局内存是GPU中容量最大但延迟最高的一种内存。它可以被所有线程访问，适合存储大量的数据，但访问速度较慢。

• 使用场景：全局内存通常用于存储需要在多个线程块之间共享的大量数据。

• 编程方法：在CUDA中，通过cudaMalloc()分配全局内存，通过cudaMemcpy()将数据从主机（CPU）传输到设备（GPU）全局内存。全局内存访问需要配合内存对齐和coalescing访问（内存访问合并）策略，以减少访问延迟。

全局内存可以根据不同的作用划分出不同的内存区域，以满足不同的计算图形需求(下列内存都算是全局内存的一部分，除了Pinned Memory，只是根据不同的需求特化为了不同的内存)：

局部内存（Local Memory）：存储线程私有的局部变量，实际上是全局内存的一部分。
常量内存（Constant Memory）：用于存储只读的常量数据，性能经过优化。
纹理内存（Texture Memory）：用于存储纹理数据，支持随机访问和插值操作，有缓存加速。
表面内存（Surface Memory）：类似纹理内存，但允许读写操作。
全局内存（Global Memory）：存储需要跨多个线程块共享的大规模数据，可以读写。
页面锁定内存（Pinned Memory）：加速GPU与主机之间的数据传输，位于主机端。
统一内存（Unified Memory）：用于简化CPU和GPU之间的内存共享和数据传输。
显存帧缓冲区：专门用于存储渲染图像的区域，主要用于图形任务。

二、局部内存（Local Memory）

• 特性：局部内存实际上是全局内存的一部分。它用于存储寄存器不足时溢出的局部变量，或者数组等超出寄存器容量的数据。与全局内存一样，本地内存的访问延迟较高。

• 使用场景：本地内存主要在寄存器不足时被自动使用，通常不直接由程序员管理。

• 编程方法：本地内存由CUDA编译器自动管理，程序员无法直接分配或控制本地内存的使用。
  访问本地内存会导致性能下降，因此编写高效代码时需要尽量减少局部变量和数组的大小。

局部内存位于全局内存中，是全局内存中一小块专门分配出来的区域。

三、纹理内存（Texture Memory）

• 特性：纹理内存是专为图形处理而设计的特殊内存，具有专门的缓存机制，适合随机读取。其主要优势是硬件级的纹理过滤和插值能力。
• 使用场景：纹理内存适合处理2D和3D纹理数据，广泛应用于计算机图形学以及一些需要随机访问的科学计算任务。
• 编程方法：在CUDA中，使用cudaBindTexture()将全局内存绑定到纹理内存。
  纹理内存对随机访问的数据提供优化，能更有效地处理图像和矩阵类操作。

纹理内存位于全局内存中，是全局内存中一小块专门分配出来的区域。

四、常量内存（Constant Memory）

• 特性：常量内存是只读的内存，容量非常有限（通常只有64KB）。它可以被所有线程访问，但它的特点是针对所有线程的统一访问进行了优化。
• 使用场景：常量内存通常用于存储在所有线程中相同的、不经常改变的数据，例如全局常量、配置信息等。
• 编程方法：在CUDA中，使用__constant__关键字声明常量内存。
  虽然常量内存的读取速度比全局内存快，但由于其只读特性和容量限制，适合存储少量全局不变的数据。

常量内存位于全局内存中，是全局内存中一小块专门分配出来的区域。它通常用于存储只读数据，并且有专门的缓存机制来优化读取性能。

五、共享内存（Shared Memory）

• 特性：共享内存是GPU中各个线程块内部的高速内存。它的访问速度比全局内存快很多，但其容量较小（每个线程块大约几十KB）。
• 使用场景：共享内存适合在同一线程块内的线程之间快速共享和交换数据。
• 编程方法：在CUDA中，使用__shared__关键字声明共享内存变量。
  共享内存的访问非常快，但必须在一个线程块内操作，跨线程块的访问需要通过全局内存来进行。
  编写程序时，要仔细管理共享内存以避免“银行冲突”（bank conflict），这是共享内存访问时的性能瓶颈之一。

六、寄存器内存（Registers Memory）

• 特性：寄存器是GPU上最快的内存类型，每个线程都有自己的一组寄存器，存储少量的局部变量。寄存器的访问速度极快，但数量有限。
• 使用场景：寄存器用于存储单个线程执行过程中需要频繁使用的局部变量。
• 编程方法：寄存器在CUDA编程中是隐式管理的，程序员通过声明局部变量让编译器决定分配寄存器。
  当寄存器使用过多时，数据可能会溢出到全局内存或本地内存，导致性能下降。

---

总结