CUDA存储器详解

       为了更好的在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  如果不初始化常量内存，将默认数组各元素为0
	cudaMemcpyToSymbol(constArray, inArray, size);
	//call kernel fun
	dim3 threadPerBlock(16);
	dim3 blockNum((SIZE + threadPerBlock.x - 1) / threadPerBlock.x);
	//计算输入数据和输出数据均采用常量内存
	addKernel << <blockNum , threadPerBlock >> > ();
	//copy  constant data to host data
	cudaMemcpyFromSymbol(result, cosntArray_Result, size);
	
	//show data
	for (int i = 0; i < SIZE; i++) {
		printf("%5d", result[i]);
	}
    return 0;
}

结果：

6、纹理内存

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

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