CUDA：Compute Unified Device Architecture，是由NVIDIA所推出的一种集成技术，允许使用标准C来进行GPU代码编程，最终转为PTX汇编代码。
CPU与GPU
GPU可以看作是CPU的协助处理器，使用GPU实际指的是基于CPU+GPU的异构计算架构。通过PCle总线连接，CPU端成为Host端，GPU端称为Device端。
基于CPU+GPU的异构计算架构
GPU适合数据并行的计算密集型任务，如大型矩阵运算，而CPU的运算核心较少，但是其可以实现复杂的逻辑运算，因此其适合控制密集型任务。此外，CPU上的线程是重量级的，上下文切换开销大，但是GPU由于存在很多核心，其线程是轻量级的。因此，基于CPU+GPU的异构计算平台可以优势互补，CPU负责处理逻辑复杂的串行程序，而GPU重点处理数据密集型的并行计算程序。
CUDA程序执行流程
1、分配host内存，并进行数据初始化；
2、分配device内存，并从host将数据拷贝到device上；
3、调用CUDA kernel在device上完成指定的运算；（kernel是在device上线程中并行执行的函数，核函数用__global__符号声明，在调用时需要用<<<grid, block>>>来指定kernel要执行的线程数量，在CUDA中，每一个线程都要执行核函数，并且每个线程会分配一个唯一的线程号thread ID，这个ID值可以通过核函数的内置变量threadIdx来获得。）
4、将device上的运算结果拷贝到host上；
5、释放device和host上分配的内存。
CUDA函数类型限定词
1）__global__：在device上执行，从host中调用（一些特定的GPU也可以从device上调用），返回类型必须是void，不支持可变参数参数，不能成为类成员函数。注意用__global__定义的kernel是异步的，这意味着host不会等待kernel执行完就执行下一步。
2）__device__：在device上执行，单仅可以从device中调用，不可以和__global__同时用。
3）__host__：在host上执行，仅可以从host上调用，一般省略不写，不可以和__global__同时用，但可和__device__，此时函数会在device和host都编译。
Kernel线程结构
kernel在device上执行时实际上是启动很多线程，一个kernel所启动的所有线程称为一个网格（grid），同一个网格上的线程共享相同的全局内存空间，grid是线程结构的第一层次，而网格又可以分为很多线程块（block），一个线程块里面包含很多线程，这是第二个层次。grid和block都是定义为dim3类型的变量，dim3可以看成是包含三个无符号整数（x，y，z）成员的结构体变量，在定义时，缺省值初始化为1。因此grid和block可以灵活地定义为1-dim，2-dim以及3-dim结构，对于图中结构（主要水平方向为x轴），定义的grid和block如下所示，kernel在调用时也必须通过执行配置<<<grid, block>>>来指定kernel所使用的线程数及结构。
Kernel线程结构
一个线程需要两个内置的坐标变量（blockIdx，threadIdx）来唯一标识，它们都是dim3类型变量，其中blockIdx指明线程所在grid中的位置，而threaIdx指明线程所在block中的位置。以Thread(1, 1)为例：
threadIdx.x = 1 threadIdx.y = 1
blockIdx.x = 1 blockIdx.y =1
测试GPU硬件配置
查看GPU配置
CUDA内存模型
每个线程有自己的私有本地内存（Local Memory），而每个线程块有包含共享内存（Shared Memory）,可以被线程块中所有线程共享，其生命周期与线程块一致。此外，所有的线程都可以访问全局内存（Global Memory）。还可以访问一些只读内存块：常量内存（Constant Memory）和纹理内存（Texture Memory）。
分配托管内存
分配托管内存
使用统一内存更简洁了，由于托管内存自动进行数据传输，这里要用cudaDeviceSynchronize()函数保证device和host同步，这样后面才可以正确访问kernel计算的结果。
实例：矩阵加法
利用上图2-dim结构实现两个矩阵的加法，每个线程负责处理每个位置的两个元素相加，代码如下所示。线程块大小为(16, 16)，然后将N*N大小的矩阵均分为不同的线程块来执行加法运算。
矩阵加法
实例：矩阵乘法
设输入矩阵为
和
，要得到
。实现思路是每个线程计算
的一个元素值
，对于矩阵运算，应该选用grid和block为2-D的。首先定义矩阵的结构体：
nvprof工具
通过命令nvprof cuda9.exe可以得到kernel的运行情况，
网友评论