cuda简介和编程模型

CUDA（Computer Unified Device Architecture）是一种由NVIDIA推出的通用并行计算架构，在底层API的基础上，封装了一层，并扩展出多语言支持，该架构使GPU能够解决复杂的计算问题。

CUDA C 是标准ANSI C语言的扩展，扩展出一些语法和关键字来编写设备端代码，而且CUDA库本身提供了大量API来操作设备完成计算。

对于API也有两种不同的层次，一种相对交高层，一种相对底层。

• CUDA驱动API
• CUDA运行时API

一个异构环境，通常有多个CPU多个GPU，他们都通过PCIe总线相互通信，也是通过PCIe总线分隔开的。所以我们要区分一下两种设备的内存：

• 主机：CPU及其内存，主机端的是CPU执行的程序部
• 设备：GPU及其内存，设备端的是GPU上执行的程序部

一个核函数对应一个grid，一个grid可以有很多个块，每个块可以有很多的线程，不同块线程不相互影响，每个线程有一个编号，依靠下面两个内置结构体确定线程标号：

• blockIdx：线程块在线程网格内的位置索引

• threadIdx：线程在线程块内的位置索引

  两个内置结构体基于 uint3 定义，包含三个无符号整数的结构，包含三个字段x,y,z

blockDim和gridDim表示blockIdx、threadIdx中三个字段的范围，类型为dim3(基于uint3定义的数据结构)的变量，也包含三个字段x,y,z

dim3是手工定义的，主机端可见。uint3是设备端在执行的时候可见的，不可以在核函数运行时修改，初始化完成后uint3值就不变了

一个Grid可以包含多个Blocks，Blocks的组织方式可以是一维的，二维或者三维的。block包含多个Threads，这些Threads的组织方式也可以是一维，二维或者三维的。

一个Grid可以包含多个Blocks，Blocks的组织方式可以是一维的，二维或者三维的。block包含多个Threads，这些Threads的组织方式也可以是一维，二维或者三维的。

Grid和Block不同划分方式下线程索引ID的计算公式：

• grid划分成1维，block划分为1维

int threadId = blockIdx.x *blockDim.x + threadIdx.x;  

• grid划分成1维，block划分为2维

int threadId = blockIdx.x * blockDim.x * blockDim.y+ threadIdx.y * blockDim.x + threadIdx.x;  

• grid划分成2维，block划分为1维

int blockId = blockIdx.y * gridDim.x + blockIdx.x;  
int threadId = blockId * blockDim.x + threadIdx.x;  

• grid划分成2维，block划分为2维

int blockId = blockIdx.x + blockIdx.y * gridDim.x;  
int threadId = blockId * (blockDim.x * blockDim.y)  
                   + (threadIdx.y * blockDim.x) + threadIdx.x;  

CUDA 每个线程执行的函数。CUDA使用扩展的C语言编写内核函数，关键字为global。内核函数返回值只能是void。

声明核函数

__global__ void kernel_name(argument list);

启动核函数，通过的以下的ANSI C 扩展出的CUDA C指令

kernel_name<<>>(argument list);

在C语言函数前没有的限定符global ,CUDA C中还有一些其他我们在C中没有的限定符，如下：

执行

Kernel核函数编写有以下限制

• 只能访问设备内存
• 必须有void返回类型
• 不支持可变数量的参数
• 不支持静态变量
• 显示异步行为

所有编程都需要对错误进行处理，早起的编码错误，编译器会帮我们搞定，内存错误也能观察出来，但是有些逻辑错误很难发现，甚至到了上线运行时才会被发现，而且有些厉害的bug复现会很难，不总出现，但是很致命，而且CUDA基本都是异步执行的，当错误出现的时候，不一定是哪一条指令触发的，这一点非常头疼；这时候我们就需要对错误进行防御性处理了，例如我们代码库头文件里面的这个宏：

#define CHECK(call)
{
  const cudaError_t error=call;
  if(error!=cudaSuccess)
  {
      printf("ERROR: %s:%d,",__FILE__,__LINE__);
      printf("code:%d,reason:%sn",error,cudaGetErrorString(error));
      exit(1);
  }
}

就是获得每个函数执行后的返回结果，然后对不成功的信息加以处理，CUDA C 的API每个调用都会返回一个错误代码，这个代码我们就可以好好利用了，当然在release版本中可以去除这部分，但是开发的时候一定要有的