一、CUDA简介
CUDA是并行计算的平台和类C编程模型，可以实现并行算法。电脑要配备NVIDIA GPU，就可以在许多设备上运行你的并行程序。
二、CUAD编程
CUDA编程允许程序执行在异构系统上，即CPU和GPU，并由PCL-Express总线区分开。
Host：CPU and itsmemory(host memory)
Device：GPU and its memory(device memory)
在编程时，一般用h_前缀代表host memory，d_代表device memory
kernel是CUDA编程的关键，是在GPU上运行的代码，用标识符 __global__注明。
host可以独立于host进行大部分操作。当一个kernel启动后，控制权会立刻返还给CPU来执行其他额外的任务。所以，CUDA编程是异步的。一个典型的CUD程序包含由并行代码补足的串行代码，串行代码由host执行，并行代码在device中执行，host 端代码是标准C，device是CUDA C代码。我们可以把所有代码放到一个单独的源文件，也可以使用多个文件或库。NVIDIA C编译器(nvcc)可以编译host和device生成可执行程序。
CUDA程序的处理流程：
（1）分配host内存，进行数据初始化
（2）分配device内存，将数据从host上复制到device上
（3）调用CUDA的核函数在device上完成运算
（4）将device上的运算结果复制到host上
（5）释放device和host上分配的内存
三、memory操作
CUDA程序将系统分为host和device，二者有各自的memory。kernel可以操作device memory，为了可以很好的控制device端内存，CUDA提供了几个内存操作函数：
CUDA C的风格跟C很接近：
cudaError_cudaMalloc(void** devPtr，size_t size)
我们主要看看cudaMencpy，其函数原型为：
cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count,cudaMemcpyKind kind )
其中cudaMemcpykijnd的可选类型有：
1. cudaMemcpyHostToHost
2. cudaMemcpyHostToDevice
3.cudaMemcpyDeviceToHost
4.cudaMemcpyDeviceToDevice
对于返回类型cudaError_t，如果正确调用，则返回cudaSuccess,否则返回cudaErrorMemoryAllocation,可以使用char* cudaGetErrorString(cudaError_t error)将其转化为易于理解的格式。
四、组织线程
CUDA线程分为Grid和Block两个层次。
一个单独的kernel启动所有线程组成一个grid, grid中所有线程共享global memory, 一个grid由许多block组成，block由许多线程组成，grid和block都可以使一维二维或三维，上图是一个二维grid和二维block.
CUDA内置变量：
（1）blockIdx: block的索引，blockIdx.x表示block的x坐标。
（2）threadIdx: 线程索引，同理blockIdx.
（3）blockDim: block维度，上图blockDim.x=5.
（4）gridDim: grid维度，同理blockDim.
一般会把grid组织成2D,block为3D。grid 和 block都使用dim3作为声明，例如;
dim3 block(3);
dim3 grid((nElem+block.x-1)/block.x);
需要注意的是：dim3仅为host端可见，其对应的device端类型为uint3.
CUDA的线程模型从小往大总结就是：
（1）Thread : 线程，并行的基本单位。
（2）Thread Block：线程块，互相合作的线程组，线程块有如下几个特点：
允许彼此同步
可以通过共享内存快速交换数据。
以1 维、2维、3维组织
Grid: 一组线程块
以1维、2维或3维组织
共享全局内存
Kernel: 在GPU上执行的核心程序，这个kernel函数是运行在Grid上的。
one kernel<-> one Grid
每一个block和每个thread都有自己的ID，我们通过相应的索引找到相应的线程和线程块。
threadIdx, blockIdx
Block ID:1D or 2D
Thread ID: 1D 2D 3D
GPU上很多并行化的轻量级线程。Kernel在device上执行时实际上是启动很多线程，一个kernel所启动的所有线程称为一个网格（grid）,同一个网格上的线程共享相同的全局内存空间。
grid是线程结构的第一层次，而网格可分为很多线程块(block)
第二层次：一个线程块里面包含很多线程
线程两层组织结构：这是一个grid、block均为2-dim的线程组织。grid和block都是定义为dim3类型的变量，dim3可以看成是包含是三个无符号整数（x,y,z）成员的结构体变量，在定义时，缺省值初始化为1.
grid和block可以灵活地定义为1-dim，2-dim以及3-dim结构，kernel调用时也必须通过执行配置<<<grid,block>>>来指定kernel所使用的网格维度和线程块维度。
CUDA的这种<<<grid,block>>>其实就是一个多级索引的方法，
第一级索引是(grid.xIdx,grid.yIdy), 对应上图例子就是（1,1）
通过它我们能找到了这个线程块的位置
第二级索引是(block.xIdx,block.yIdx,block.zIdx)来定位指定的线程
这就是CUDA的线程组织结构。
每个线程由每个线程处理器（SP）、线程块由多核处理器（SM）执行、一个kernel其实由一个grid来执行，一个kernel一次只能在一个GPU上执行
block是软件概念，一个 block 只会由一个sm调整，程序员在开发时，通过设定block的属性，告诉GPU硬件，我有多少线程，线程怎么组织。
而具体怎么调整由sm的warp scheduler负责，block一旦被分配好SM，该block就会一直驻留该SM中，直到执行结束。一个SM可以拥有多个Blocks,但需要序列执行。
下图显示了GPU内部的硬件架构：
3. CUDA内存模型
CUDA中的内存模型分为以下几个层次：
每个线程都用自己的registers（寄存器）
每个线程都有自己的local memory（局部内存）
每个线程块都有自己的shared memory（共享内存），所有线程块内的所有线程共享这 段内存资源。
每个grid都有自己的global memory（全局内存），不同线程块的线程都可使用
每个grid都有自己的constant memory(常量内存）和 texture memory(纹理内存），不 同线程块的线程都可使用。
线程访问这几类存储器的速度是register > local memory > shared memory >global memory
计算机架构中的所在层次：
4.CUDA编程模型
CUDA是怎么写程序的了：
CUDA术语：
重要：
1. 通过关键字就可以表示某个程序在CPU上跑还是GPU上跑！
我们用__global__定义一个kernel函数，就是CPU上调用，GPU上执行，
注：__global__函数的返回值必须设置为void.
2. CPU和GPU间的数据传输：
GPU内存分配回收内存的函数接口：
cudaMemcpyHostToDevice（CPU到GPU）
cudaMemcpyDeviceToHost（GPU到CPU）
cudaMemcpyDeviceToDevice （GPU到GPU）
3. 用代码表示线程组织模型？
可以用dim3类来表示网格和线程块的组织方式，
网格grid可以 表示为一维和二维格式，
线程块block可以表示为一维、二维、三维的数据格式。
dim3 DimGrid(100,50) //5000个线程块，维度是100*50
dim3 DimBlock(4,8,8); //每个线层内包含256个线程，线程块内的维度4*8*8
4. 计算线程号
（1）使用N个线程块，每一个线程块只有一个线程，即：
dim3 dimGrid (N);
dim3 dimBlock (1);
此时的线程号的计算方式是：
threadId = blockIdx.x;
其中threadId的范围为0到N-1。对于这种情况，我们可以将其看做是一个列向量， 列向量中的每一行对应一个线程块。列向量中每一行只有一个元素，对应一个线 程。
（2）使用M*N个线程块，每个线程块1个线程
由于线程块是2维的，故可以看做是一个M*N的2维矩阵，其线程号有两个维 度, 即：
dim3 dimGrid(M,N);
dim3 dimBlock（1）；
其中
blockIdx.x 取值0到M-1
blockIdx.y取值0到N-1
这种情况一般用于处理2维数据结构，比如2维图像，，每一个像素用一个线程来处理处理，此时需要线程号来映射图像像素的对应位置，如：
pos= blockIdx.y * blockDim.x+blockIdx.x;
（3）使用一个线程块，该线程具有N个线程
dim3 dimGrid(1);
dim3 dimblock(N)
此时线程号的计算方式是：
threadId = threadIdx.x;
其中threadId的范围是0到N-1，对于这种情况，可以看做是一个行向量，行向量中的每一个元素对应着一个线程。
（4）使用M个线程块，每个线程块内含有N个线程，即：
dim3 dimGrid（M）；
dim3 dimBlock(N);
这可以想象成二维矩阵，矩阵的行与线程块相对应，矩阵的列与线程编号对应，那线程号的计算方式是：
threadId = threadIdx.x + blockIdx*blockDim.x;
就是把二维的索引空间转换为一维索引空间的过程。
（5）使用M*N的二维线程块，每一个线程块具有P*Q个线程，即
dim3 dimGrid（M， N）；
dim3 dimBlock(P,Q);
这种适合于处理具有二维数据结构的算法，比如图像处理领域。
索引有两个维度：
threadId.x = blockIdx.x* blockDim.x+ threadIx.x;
threadId.y = blockIdx.y*blockDim.y+threadIdx.y;
上述是把线程和线程块的索引映射为图像像素坐标的计算方法.
