CUDA下的GPU编程--线程和变量

CUDA的线程是多维的，启动一个线程格，线程格可以是多维的，线程格中分为线程块，线程块也可以是多维的，线程块中包含线程，对于GPU来说，同时启动200万个线程是很轻易的事情。

blockIdx :当前执行设备代码的线程块的索引

dim3 grid(DIM,DIM):声明一个二维的线程格，名字为grid;

kernel<<<grid,1>>(dev_t):线程块为一个二维线程块，每个线程块有一个线程

gridDim：保存线程格每一维的线程块的数量

blockDim：保存的是线程块中每一维的线程数量

若线程块和线程格都是一维的：

blockDim.x * gridDim.x表示的是总共线程数量。

 

CUDA定义的函数和变量方式有 __global__ ­ __device__ 和__shared__，

__global__函数直接告诉编译器，这个函数是在GPU上执行的，不是在CPU上执行，可以在主函数中直接调用（主函数中的代码如果没有特别说明都是在CPU上运行的），此时该函数会自动交由GPU执行。__global__只能写在主函数中。

声明格式：__global__ add(int *a){}

调用格式：add<<<N,M>>>(dev_t);（启动线程块为N，每个线程块有M线程的线程集去执行add(dev_t)函数.

 

__device__可以定义GPU全局变量和函数，表示该变量或函数是保存在GPU上的，代码也是在GPU上执行的，与__global__不同的是，__device__只能在__device__函数或__global__中调用，不能直接写在主函数。

 可以直接在主函数中声明一个普通的变量指针，如int *dev_t，然后再由主函数中调用HANDLE_ERROR(cudaMalloc(  (void**) &dev_t,sizeof(int) ) );就可以给dev_t分配一个GPU内存上的地址，也就是说dev_t指向的是GPU的地址，此后不能在CPU执行的代码中调用该指针，因为CPU的内存跟GPU的内存不是共享的，只能在__global__或者是__device__函数中调用该指针。如果要实现CPU和GPU的通信，只能通过拷贝地址函数将GPU上某段内存拷贝到主存上。比如若GPU中计算好了一个值在指针dev_t所指的int区域，要返回给CPU的t变量话，HANDLE_ERROR(cudaMemcpy( &t , dev_t, sizeof(int),cudaMemcpyDeviceToHost ));

 

__shared__声明的变量，这个变量会驻留在共享内存中，对于在GPU上启动的每个线程块，CUDA编译器都会创建该变量的一个副本，线程块中的每个线程都共享这块内存，但是线程却无法看到也不难修改其他线程块的变量副本。使得一个线程块中的多个线程能在计算机上进行通信和协作，访问共享内存的延迟要远低于访问普通缓冲区的延迟。

 

__constant__ 常量内存保存在设备函数执行期间不会发生变化的数据，nvidia硬件提供了64KB的常量内存。常量内存里的变量声明也要加上 __constant__修饰符，如果是数组的话，要在编译时为这个数组提交一个固定的大小，如__constant__ int a[10]。

当从主机内存复制到GPU上的常量内存时，不使用cudaMemcpy()，用cudaMemcpyToSymbol()。

线程束是指一个包含32个线程的集合，多个线程集合被“编织在一起”并以步调一致“的形式执行，在程序中的每一行，线程束中的每个线程都将在不同的数据上执行相同的指令。如果在半线程束中的每个线程都从常量内存的相同地址上读取数据，那么GPU只会产生一次读取请求并在随后将数据广播到每个线程，如果从常量内存中读取大量的数据，那么这种方式产生的内存流量指示全局时的1/16。但是当半线程束中的所有16个线程需要访问常量内存中不同数据，那么这个16次不同的读取请求操作会被串行化，反而影响效率。

 

设备执行的代码中可以使用__syncthreads()进行对线程块中的线程进行同步，可以确保线程块中的每个线程都执行完__syncthreads()前面的语句后，才会执行下一条语句。