cuda学习笔记（一）存储

1. 一个GPU上有很多的sm（stream Multiprocessor），每个sm中包括了8个sp（stream Processor）标量流处理器，商业宣传中所说的数百个“核”，大多指的是sp的数量。隶属于同一个sm的sp共用同一套取指与发射单元。CUDA中的kernel是以block为单位执行的，一个block必须在一个sm上执行，一个sp执行一个线程，但是一个sm可以同时存在多个block的上下文。一个sm中活动线程块的数量不超过8个；所有活动线程块中的warp数之和在计算能力1.3设备中，不超过32。

 

2.目前一个kernel中只有一个grid，grid只能设定x、y维，z维默认为1，但是block的三维都可以自己设定。

 

3.一个block中最多有512个thread，这些thread会被分割为线程束warp在sm上进行调度。每个warp包含连续的32个线程，访问存储器时是以半线程束为单位的。

 

4.cuda的源文件必须使用nvcc编译器进行编译。这个编译器在编译的时候，会分离出主机端和设备端代码，主机端调用其他高性能编译器编译，如gcc，设备端代码由nvcc编译成ptx代码或者二进制代码。ptx代码是为动态编译器JIT设计的，这样可以应付不同显卡上的不同的机器语言。

 

5.在计算能力1.3的设备中，每个sm的寄存器文件数量为16384，每个寄存器文件大小为32bit，如果每个线程使用的私有变量太多或者大小不定，可能将这样的变量分配到局部寄存器中（local memory）。例如usigned int mt[3]就是局部寄存器的。 变量被分在哪里可以通过编译输出的ptx汇编代码查看，有.local的则为局部寄存器中的。

 

6.共享存储器：extren __shared__声明的变量都开始于同一个地址，因此变量的布局必须通过显式的定义，如

extern __shared__ char array[];

__device__ void func()

{

 short* array0 = (short*)array;

 float* array1 = (float*)&array0[128];

 int array2 = (int*)&array1[64];

}

 

7. cuda主机端的内存被分为两种：可分页内存（pageable memory）（由malloc和new分配）和页锁定（pinned）内存（一定是存储在物理内存，不会存储在虚拟内存，因此速度较快，可以通过DMA加速与设备通信，使用cudaHostAlloc 和 cudaFreeHost分配和释放

）

cuda2.3版本以上支持portable memory，在cudaHostAlloc（）时加上cudaHostAllocPortable标志，可以使多个CPu共享一块页锁定内存以实现CPU间的通信。默认情况下，pinned内存只能由分配的线程访问。

write-combined Memory，在分配cudaHostAlloc()时设定cudahostAllocWriteCombined的标志，可以免除cpu对内存的监视，减少将内存上的数据缓存到L1、L2 cache中，在主机和显存数据传输时节省了时间，这样的作法比较适合于主机端只写的数据。

mapped memory,主机端指针通过cudaHostAlloc()获得，设备端通过cudaHostGetDevicePointer()获得，这样可以从kernel里直接访问主存，省略了在显存上分配空间，实现了对内存的零拷贝访问，不过使用前要通过cudaGetDeviceProperties()判断设备是否有canMapHostMemory的属性。

 

8.常数存储：显存上的只读地址空间，拥有缓存加速，每个sm拥有8KB的常数存储器缓存。

__constant__ char p_hello[];  可以通过cudaMemcpyToSymbol(p_hello, hello_host, sizeof(char)*11)从内存把数据拷进去。

 

9.纹理存储器：只读存储器，具备一些特殊的功能