cuda编程入门

cuda编程：面向大众的高性能计算

- GPU架构特点
我们知道，高性能计算的关键利用多核处理器进行并行计算。
当我们求解一个计算机程序任务时，我们很自然的想法就是将该任务分解成一系列小任务，把这些小任务一一完成。在串行计算时，我们的想法就是让我们的处理器每次处理一个计算任务，处理完一个计算任务后再计算下一个任务，直到所有小任务都完成了，那么这个大的程序任务也就完成了。如下图所示，就是我们怎么用串行编程思想求解问题的步骤。

但是串行计算的缺点非常明显，如果我们拥有多核处理器，我们可以利用多核处理器同时处理多个任务时，而且这些小任务并没有关联关系（不需要相互依赖，比如我的计算任务不需要用到你的计算结果），那我们为什么还要使用串行编程呢？为了进一步加快大任务的计算速度，我们可以把一些独立的模块分配到不同的处理器上进行同时计算（这就是并行），最后再将这些结果进行整合，完成一次任务计算。下图就是将一个大的计算任务分解为小任务，然后将独立的小任务分配到不同处理器进行并行计算，最后再通过串行程序把结果汇总完成这次的总的计算任务。

所以，一个程序可不可以进行并行计算，关键就在于我们要分析出该程序可以拆分出哪几个执行模块，这些执行模块哪些是独立的，哪些又是强依赖强耦合的，独立的模块我们可以试着设计并行计算，充分利用多核处理器的优势进一步加速我们的计算任务，强耦合模块我们就使用串行编程，利用串行+并行的编程思路完成一次高性能计算。
- CPU与GPU的区别
首先CPU是专为顺序串行处理而优化的几个核心组成。而GPU则由数以千计的更小、更高效的核心组成，这些核心专门为同时处理多任务而设计，可高效地处理并行任务。也就是，CPU虽然每个核心自身能力极强，处理任务上非常强悍，无奈他核心少，在并行计算上表现不佳；反观GPU，虽然他的每个核心的计算能力不算强，但他胜在核心非常多，可以同时处理多个计算任务，在并行计算的支持上做得很好。
整理一下两者特点就是：
CPU：擅长流程控制和逻辑处理，不规则数据结构，不可预测存储结构，单线程程序，分支密集型算法
GPU：擅长数据并行计算，规则数据结构，可预测存储模式

- cuda线程模型
下面我们介绍CUDA的线程组织结构。首先我们都知道，线程是程序执行的最基本单元，CUDA的并行计算就是通过成千上万个线程的并行执行来实现的。下面的机构图说明了GPU的不同层次的结构。

CUDA的线程模型从小往大来总结就是：
1、Thread：线程，并行的基本单位
2、Thread Block：线程块，互相合作的线程组，线程块有如下几个特点：
允许彼此同步
可以通过共享内存快速交换数据
以1维、2维或3维组织
3、Grid：一组线程块
以1维、2维组织
共享全局内存
4、Kernel：在GPU上执行的核心程序，这个kernel函数是运行在某个Grid上的。
One kernel <-> One Grid，一个kernel其实由一个grid来执行，一个kernel一次只能在一个GPU上执行
每一个block和每个thread都有自己的ID，我们通过相应的索引找到相应的线程和线程块。
threadIdx，blockIdx
Block ID: 1D or 2D
Thread ID: 1D, 2D or 3D
解释：
1、kernel在device上执行时实际上是启动很多线程，一个kernel所启动的所有线程称为一个网格（grid），同一个网格上的线程共享相同的全局内存空间。
2、 grid是线程结构的第一层次，而网格又可以分为很多线程块（block），一个线程块里面包含很多线程，这是第二个层次。
3、线程两层组织结构如上图所示，这是一个gird和block均为2-dim的线程组织。grid和block都是定义为dim3类型的变量，dim3可以看成是包含三个无符号整数（x，y，z）成员的结构体变量，在定义时，缺省值初始化为1。因此grid和block可以灵活地定义为1-dim，2-dim以及3-dim结构，kernel调用时也必须通过执行配置<<<grid, block>>>来指定kernel所使用的网格维度和线程块维度。
4、举个例子，我们以上图为例，分析怎么通过<<<grid,block>>>>这种标记方式索引到我们想要的那个线程。CUDA的这种<<<grid,block>>>其实就是一个多级索引的方法，第一级索引是(grid.xIdx, grid.yIdy)，对应上图例子就是(1, 1)，通过它我们就能找到了这个线程块的位置，然后我们启动二级索引(block.xIdx, block.yIdx, block.zIdx)来定位到指定的线程。这就是我们CUDA的线程组织结构。

1. cuda内存模型
   CUDA中的内存模型分为以下几个层次：
   每个线程都用自己的registers（寄存器）
   每个线程都有自己的local memory（局部内存）
   每个线程块内都有自己的shared memory（共享内存），所有线程块内的所有线程共享这段内存资源
   每个grid都有自己的global memory（全局内存），不同线程块的线程都可使用
   每个grid都有自己的constant memory（常量内存）和texture memory（纹理内存），不同线程块的线程都可使用
   线程访问这几类存储器的速度是register > local memory >shared memory > global memory
   下面这幅图表示就是这些内存在计算机架构中的所在层次。
   

• cuda编程模型
  常见cuda术语
  
  通过关键字就可以表示某个程序在CPU上跑还是在GPU上跑！如下表所示，比如我们用__global__定义一个kernel函数，就是CPU上调用，GPU上执行，注意__global__函数的返回值必须设置为void。
  CPU和GPU间的数据传输怎么写？
  首先介绍在GPU内存分配回收内存的函数接口：
  cudaMalloc(): 在设备端分配global memory
  cudaFree(): 释放存储空间
  CPU的数据和GPU端数据做数据传输的函数接口是一样的，他们通过传递的函数实参（枚举类型）来表示传输方向：
  cudaMemcpy(void dst, void src, size_t nbytes,
  enum cudaMemcpyKind direction)
  enum cudaMemcpyKind:
  cudaMemcpyHostToDevice（CPU到GPU）
  cudaMemcpyDeviceToHost（GPU到CPU）
  cudaMemcpyDeviceToDevice（GPU到GPU）