CUDA存储器详解

来源：互联网发布：山东网络作家协会编辑：程序博客网时间：2024/06/06 06:34

为了更好的在kernel函数中提高数据访问效率，我们需要理解CUDA各个存储器模型的特点，合理的分配存储空间。

-------------------华丽的分割线-------------------

一、存储器比较

每个线程拥有自己的register and loacal memory;(可读可写)

每个线程块拥有一块shared memory;(可读可写)

所有线程都可以访问global memory;(可读可写)

还有，可以被所有线程访问的只读存储器：constant memory and texture memory （只读）

通过上述图，我们理解为什么shared和register拥有同样的访问速度，shared memory 和 register 位于 GPU 片内。全局存储三个：global memory,constant memory,texture memory.

存储器关系图，如图所示：

二、各类存储器细讲

1、寄存器Register

　　寄存器是GPU上的高速缓存器，其基本单元是寄存器文件，每个寄存器文件大小为32bit. Kernel中的局部(简单类型)变量第一选择是被分配到Register中。

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

2、局部存储器 local memory

　　当register耗尽时，数据将被存储到local memory。如果每个线程中使用了过多的寄存器，或声明了大型结构体或数组，或编译器无法确定数组大小，线程的私有数据就会被分配到local memory中。但是local memory 的数据是被保存在显存中的，速度很慢。

　　特点：每个线程私有；没有缓存，慢。

　　注：在声明局部变量时，尽量使变量可以分配到register。如：

　　unsigned int mt[3]; 改为：　unsigned int mt0, mt1, mt2;

3、共享内存

共享存储器也是GPU片内的高速存储器，访问sharedmemory几乎和访问register一样快，是实现线程间通信的延迟最小的方法。shared memory是一块可以被同一block中的所有thread访问的可读写存储器。每个块分配48KB 。对于GPU上启动的每个线程块，CUDA C编译器都将创建该共享变量的一个副本。线程块中的每个线程都共享这块内存，但线程却无法看到也不能修改其他线程块的变量副本。这样使得一个线程块中的多个线程能够在计算上通信和协作。

形式：关键字__shared__添加到变量声明中。如__shared__ float cache[10]。

4、全局内存 global memory

通俗意义上的设备内存。所有线程都可以访问；没有缓存

显存中的全局存储器也称为线性内存，线性内存通常使用cudaMalloc 分配内存

5、常量内存 constant memory

常量内存用于保存在核函数执行期间不会发生变化的数据。变量的访问限制为只读。NVIDIA硬件提供了64KB的常量内存。不再需要cudaMalloc()或者cudaFree(),而是在编译时，静态地分配空间。为了提升性能。常量内存采取了不同于标准全局内存的处理方式。在某些情况下，用常量内存替换全局内存能有效地减少内存带宽。当我们需要拷贝数据到常量内存中应该使用cudaMemcpyToSymbol()，而cudaMemcpy()会复制到全局内存。

形式：关键字__constant__添加到变量声明中。如__constant__ float s[10].

特点：只读；有缓存；空间小(64KB)
使用常量内存性能提升的原因：
I：对常量内存的单次读操作可以广播到其他的“邻近”线程。这将节约15次读取操作。（为什么是15，因为“邻近”指半个线程束，一个线程束包含32个线程的集合。）
II：常量内存的数据将缓存起来，因此对相同地址的连续读操作将不会产生额外的内存通信量。

　　注：定义常数存储器时，需要将其定义在所有函数之外，作用于整个文件

1 __constant__ int devVar;2 cudaMemcpyToSymbol(devVar, hostVar, sizeof(int), 0, cudaMemcpyHostToDevice)3 cudaMemcpyFromSymbol(hostVar, devVar, sizeof(int), 0, cudaMemcpyDeviceToHost)

例子：

#include "cuda_runtime.h"#include "device_launch_parameters.h"#include <stdio.h>#define SIZE 124//***declare the constant value***//两种方法赋值：1-定义直接初始化  2-定义后调用函数初始化cudaMemcpyToSymbol__constant__ int constArray[SIZE];__constant__ int cosntArray_Result[SIZE];__constant__ int constNumber = 10;//********************************__global__ void addKernel(){int idx = threadIdx.x + blockDim.x * blockIdx.x;cosntArray_Result[idx] = constNumber + constArray[idx];}int main(){int inArray[SIZE] = {0};size_t size = SIZE * sizeof(int);int *result = NULL;result = (int*)malloc(size);//初始化数据for (int i = 0; i < SIZE; i++) {inArray[i] = i+1;}//copy host data to constant memory  如果不初始化常量内存，将默认数组各元素为0cudaMemcpyToSymbol(constArray, inArray, size);//call kernel fundim3 threadPerBlock(16);dim3 blockNum((SIZE + threadPerBlock.x - 1) / threadPerBlock.x);//计算输入数据和输出数据均采用常量内存addKernel << <blockNum , threadPerBlock >> > ();//copy  constant data to host datacudaMemcpyFromSymbol(result, cosntArray_Result, size);//show datafor (int i = 0; i < SIZE; i++) {printf("%5d", result[i]);}    return 0;}

结果：

6、纹理内存

和常量内存一样，纹理内存是另一种类型的只读内存，在特定的访问模式中，纹理内存同样能够提升性能并减少内存流量。。纹理缓存是专门为那些在内存访问模式中存在大量空间局部性(Spatial Locality)的图形应用程序而设计的。意味着一个线程读取的位置可能与邻近线程读取的位置“非常接近”。如下图：

纹理变量（引用）必须声明为文件作用域内的全局变量。
形式：分为一维纹理内存：texture<float> texconst 和二维纹理内存 texture<float> texconst 。

阅读全文

1 0