CUDA存储类型

CUDA存储器类型：

每个线程拥有自己的register and loacal memory;

每个线程块拥有一块shared memory;

所有线程都可以访问global memory;

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

设备内存要么被分配为线性内存，要么被分配为cuda数组，

还有一种内存是主机锁页内存，锁页有以下特点：

a 在某些设备上，设备存储器和主机锁页存储器之间的数据拷贝和内核函数可以并发执行(只有锁页内存可以实现内存拷贝与核函数的并发执行，一般的主机内存不行)
b 在某些设备上，可以将主机的锁页内存映射到设备地址空间，减少主机和设备之间的数据拷贝，要访问数据的时候不是像上面那那样将数据拷贝过来，而是直接通过主机总线到主机上访问 ，使用cudaHostAlloc分配时传入cudaHostAllocMapped，或者使用cudaHostRegister时传入cudaHostRegisterMapped标签
c 默认情况下，锁页内存是可以缓存的。在使用cudaHostAlloc分配时传入cudaHostAllocWriteCombined标签，将其标定为写结合，这意味着该内存没有一级二级缓存，这样有利用主机写该内存，而如果主机读取的话，速度将会极其慢，所以这种情况下的内存应当只用于那些主机只写的存储器

d 不能分配太多，太多的话会降低系统整体性能
e 锁页内存和显存之间的拷贝速度是6G/s，普通的内存和显存之间的拷贝速度是3G/s（显存之间的拷贝速度是30G/s,CPU之间的速度是10G/s）
f 使用cudaHostAlloc函数分配内存，其内的内容需要从普通内存拷贝到锁页内存中，因此会存在：这种拷贝会带来额外的CPU内存拷贝时间开销，CPU需要把数据从可分页内存拷贝到锁页，但是采用cudaHostRegister把普通内存改为锁页内存，则不会带来额外的cpu内存拷贝时间开销，因为cudaHostAlloc的做法是先分配锁页内存，这时里面是没有数据的，那么需要将一般的内存拷贝过来，而对于cudaHostRegister内存，他是之间就使用malloc分配好的，cudaHostRegister只是设置一些内部标志位以确保其不被换出，相当于只是更改了一些标志位，就不存在前面说的数据拷贝

对锁页内存精简总结：

1. 位置：主机内存。
2. 概念：也称为页锁定内存或者不可分页内存，操作系统将不会对这块内存分页并交换到磁盘上，从而确保了该内存始终驻留在物理内存中。因此操作系统能够安全地使某个应用程序访问该内存的物理地址，因为这块内存将不会破坏或者重新定位。
3. 目的：提高访问速度。由于GPU知道主机内存的物理地址，因此可以通过“直接内存访问DMA（Direct Memory Access)技术来在GPU和主机之间复制数据。由于DMA在执行复制时无需CPU介入。因此DMA复制过程中使用固定内存是非常重要的。
4. 缺点：使用固定内存，将失去虚拟内存的所有功能；系统将更快的耗尽内存。
5. 建议：对cudaMemcpy()函数调用中的源内存或者目标内存，才使用固定内存，并且在不再需要使用它们时立即释放。
6. 形式：通过cudaHostAlloc()函数来分配；通过cudaFreeHost()释放。
7. 只能以异步方式对固定内存进行复制操作。

注意！：对于现在比较新的架构Fermi、Kepler，以及pascal，对local，global存储器存在L1，L2缓存！

当一个核函数使用大量纹理内存或大量局部内存时，其与其他核函数进行并发执行的可能性就会变小

寄存器和局部存储器，对应opencl中的私有内存（private memory）

1、  寄存器Register

　　寄存器是GPU上的高速缓存器，其基本单元是寄存器文件，每个寄存器文件大小为32bit.

　　Kernel中的局部(简单类型)变量第一选择是被分配到Register中。

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

2、  局部存储器 local memory

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

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

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

　　unsigned int mt[3];

　　改为：　unsigned int mt0, mt1, mt2;

3、  共享存储器 shared memory，对应opencl中的局部内存（local memory）

　　可以被同一block中的所有线程读写

　　特点：block中的线程共有；访问共享存储器几乎与register一样快.

[cpp]   view plain  copy

1. //u(i)= u(i)^2 + u(i-1)  
2. //Static  
3. __global__ example(float* u) {  
4.     int i=threadIdx.x;  
5.     __shared__ int tmp[4];  
6.      tmp[i]=u[i];  
7.      u[i]=tmp[i]*tmp[i]+tmp[3-i];  
8. }  
9.   
10. int main() {  
11.     float hostU[4] = {1, 2, 3, 4};  
12.     float* devU;  
13.     size_t size = sizeof(float)*4;  
14.     cudaMalloc(&devU, size);  
15.     cudaMemcpy(devU, hostU, size,  
16.     cudaMemcpyHostToDevice);  
17.     example<<<1,4>>>(devU, devV);  
18.     cudaMemcpy(hostU, devU, size,  
19.     cudaMemcpyDeviceToHost);  
20.     cudaFree(devU);  
21.     return 0;  
22. }  
23.   
24. //Dynamic  
25. extern __shared__ int tmp[];  
26.   
27. __global__ example(float* u) {  
28.     int i=threadIdx.x;  
29.      tmp[i]=u[i];  
30.      u[i]=tmp[i]*tmp[i]+tmp[3-i];  
31. }  
32.   
33. int main() {  
34.     float hostU[4] = {1, 2, 3, 4};  
35.     float* devU;  
36.     size_t size = sizeof(float)*4;  
37.     cudaMalloc(&devU, size);  
38.     cudaMemcpy(devU, hostU, size, cudaMemcpyHostToDevice);  
39.     example<<<1,4,size>>>(devU, devV);  
40.     cudaMemcpy(hostU, devU, size, cudaMemcpyDeviceToHost);  
41.     cudaFree(devU);  
42.     return 0;  
43. }  

 

 4、  全局存储器 global memory，对应opencl中的global memory

　　特点：所有线程都可以访问；没有缓存

[cpp]   view plain  copy

1. //Dynamic  
2. __global__ add4f(float* u, float* v) {  
3. int i=threadIdx.x;  
4.  u[i]+=v[i];  
5. }  
6. int main() {  
7.     float hostU[4] = {1, 2, 3, 4};  
8.     float hostV[4] = {1, 2, 3, 4};  
9.     float* devU, devV;  
10.     size_t size = sizeof(float)*4;  
11.     cudaMalloc(&devU, size);  
12.     cudaMalloc(&devV, size);  
13.     cudaMemcpy(devU, hostU, size,  
14.     cudaMemcpyHostToDevice);  
15.     cudaMemcpy(devV, hostV, size,  
16.     cudaMemcpyHostToDevice);  
17.     add4f<<<1,4>>>(devU, devV);  
18.     cudaMemcpy(hostU, devU, size,  
19.     cudaMemcpyDeviceToHost);  
20.     cudaFree(devV);  
21.     cudaFree(devU);  
22.     return 0;  
23. }  
24.   
25. //static  
26. __device__ float devU[4];  
27. __device__ float devV[4];  
28.   
29. __global__ addUV() {  
30. int i=threadIdx.x;  
31.  devU[i]+=devV[i];  
32. }  
33.   
34. int main() {  
35.     float hostU[4] = {1, 2, 3, 4};  
36.     float hostV[4] = {1, 2, 3, 4};  
37.     size_t size = sizeof(float)*4;  
38.     cudaMemcpyToSymbol(devU, hostU, size, 0, cudaMemcpyHostToDevice);  
39.     cudaMemcpyToSymbol(devV, hostV, size, 0, cudaMemcpyHostToDevice);  
40.      addUV<<<1,4>>>();  
41.     cudaMemcpyFromSymbol(hostU, devU, size, 0, cudaMemcpyDeviceToHost);  
42.     return 0;  
43. }  

   5、  常数存储器constant memory，对应opencl中的constant memory

 　　用于存储访问频繁的只读参数

 　　特点：只读；有缓存；空间小(64KB)

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

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

 6、  纹理存储器 texture memory，与opencl无对应

     是一种只读存储器，其中的数据以一维、二维或者三维数组的形式存储在显存中。在通用计算中，其适合实现图像处理和查找，对大量数据的随机访问和非对齐访问也有良好的加速效果。

     特点：具有纹理缓存，只读。

TNE END

CUDA硬件的并行：对于一个给定的核函数，其能够在多核上驻存的线程块数和线程束数，取决于以下几点：
a 核函数所使用的寄存器数量和共享存储器的大小
b 多核上寄存器和共享存储器的的大小

但是每个多核也存在一个最大的驻存线程块的块数，和最大的驻存的线程束的个数：

1060的块内寄存器大小：64k
每个多核（cuda中的流）的寄存器大小：64k
总的常量存储区：64k
总的内存：6G
每个块内最大的共享存储器大小：48k
每个多核中（个人理解成cuda中的流处理）的最大共享存储器大小：96k
每个块支持的最大线程数：1024
每个多核（cuda中的流）支持的最大线程数：2048
多核的个数（个人认为是支持并行的流处理的个数）：10

CUDA_C_Programming_Guide中3.2.5.2：
一个设备能够并发执行的核函数的最大个数取决于设备的计算能力：在5.3的设备上是16
一个cuda上下文中的核函数不能和另一个cuda上下文的核函数并发执行
当一个核函数使用大量纹理内存或大量局部内存时，其与其他核函数进行并发执行的可能性就会变小