CUDA入门

CUDA是对C语言的扩展，语法有：

1、五个内建变量：

a. gridDim是一个包含三个元素x,y,z的结构体，分别表示网格在x,y,z三个方向上的尺寸，虽然其有三维，但是目前只能使用二维

b. blockDim也是一个包含三个元素x,y,z的结构体，分别表示块在x,y,z三个方向上的尺寸，对应于执行配置中的第一个参数，对应于执行配置的第二个参数； 

c. blockIdx也是一个包含三个元素x,y,z的结构体，分别表示当前线程所在块在网格中x,y,z三个方向上的索引

d. threadIdx也是一个包含三个元素x,y,z的结构体，分别表示当前线程在其所在块中x,y,z三个方向上的索引； 

e. warpSize表明warp的尺寸，在计算能力为1.0的设备中，这个值是24，在1.0以上的设备中，这个值是32。 

2、三种函数限定符：

a.__device__，表示从GPU上调用，在GPU上执行；也就是说其可以被__global__ 或者__device__修饰的函数调用。此限定符修饰的函数使用有限制，比如在G80/GT200架构上不能使用递归，不能使用函数指针等。 

b. __global__，表示在CPU上调用，在GPU上执行，也就是所谓的内核(kernel)函数；内核只能够被主机调用，内核并不是一个完整的程序，它只是一个数据并行步骤，其指令流由多个线程执行

c. __host__，表示在CPU上调用，在CPU上执行，这是默认时的情况，也就是传统的C函数。CUDA支持__host__和__device__的联用，表示同时为主机和设备编译。 

3、<<<>>>

<<< >>>为执行配置运算符，用来传递内核函数的执行参数。执行配置有四个参数，第一个参数声明网格的大小，第二个参数声明块的大小，第三个参数声明动态分配的共享存储器大小，默认为0，最后一个参数声明执行的流，默认为0.

4、变量类型限定符：

__device__，表示从GPU上调用，在GPU上执行；也就是说其可以被__global__ 或者__device__修饰的函数调用。此限定符修饰的函数使用有限制，比如在G80/GT200架构上不能使用递归，不能使用函数指针等

__shared__，__shared__表示数据存放在共享存储器在，只有在所在的块内的线程可以访问，其它块内的线程不能访问；

__constant__，__constant__表明数据存放在常量存储器中，可以被所有的线程访问，也可以被主机通过运行时库访问；
Texture，texture表明其绑定的数据可以被纹理缓存加速存取，其实数据本身的存放位置并没有改变，纹理是来源于图形学的一介概念，CUDA使用它的原因一部分在于支持图形处理，另一方面也可以利用它的一些特殊功能。
如果在GPU上执行的函数内部的变量没有限定符，那表示它存放在寄存器或者本地存储器中，在寄存器中的数据只归线程所有，其它线程不可见。
如果SM的寄存器用完，那么编译器就会将本应放到寄存器中的变量放到本地存储器中。

CUDA中实现并行计算的是通过引入kernel函数，它可以创建多个线程并行计算，kernel函数的限定符是_global_

一个简单的例子：

__global__ void add(int *a, int *b, int *c){   int index = blockIdx.x * blockDim.x + threadIdx.x;   c[index] = a[index] + b[index];}#include "stdio.h"#define N 64int main(){  int ha[N], hb[N], hc[N];    for(int i=0; i< N; i++){  ha[i] = i;  hb[i] = i;  hc[i] = 0;}  int *da,*db,*dc;    cudaMalloc((void **)&da, 4*N);   cudaMalloc((void **)&db, 4*N);  cudaMalloc((void **)&dc, 4*N);  cudaMemcpy(da, &ha, 4*N, cudaMemcpyHostToDevice);  cudaMemcpy(db, &hb, 4*N, cudaMemcpyHostToDevice);  cudaMemcpy(dc, &hc, 4*N, cudaMemcpyHostToDevice);  add<<<8,8>>>(da,db,dc);  cudaMemcpy(&hc, dc,  4*N, cudaMemcpyDeviceToHost);   cudaFree(da); cudaFree(db); cudaFree(dc);  for (int i =0; i< 64; i++)    printf("hc[%d] is %d\n",i,hc[i]);  printf("hello world\n");  return 0;}

它是实现加法功能，cumaMalloc和cumaFree分配和释放显存空间，cumaMemcpy实现内存(Host)和显存(Device)数据之间的传递。

  add<<<8,8>>>(da,db,dc)即是对kernel函数的调用

Linux下NVIDIA GPU Computing SDK的使用：

找到相应目录：“~\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.2\C”，make以后可以得到很多可执行文件，直接运行之即可，

比如./QueryDevice可以查看GPU的配置情况。