一、CPU和GPU交互
1.各自有自己的物理内存空间，CPU的是内存，GPU的是显存
2.通过PCI-E总线互连(8GB/S~16GB/S)
3.交互开销较大
GPU各存储访存速度：
Register寄存器，最快
Shared Memory，共享存储，很快
Local Memory，本地存储，在显存中，有缓存，相对较慢
Global Memory，全局存储，在显存中，有缓存，相对较慢
Constant Memory，在显存中，多级缓存，1-100时钟周期，比较快
Texture Memory，在显存中，多级缓存，1-100时钟周期，比较快
Instruction Memory，不可见的，在显存中，有缓存
二、GPU线程组织模型
线程组成Block，Block组成Grid。
Warp是几个线程的组合，有一定特殊的规律，用于内部管理。
线程组织架构说明：
1.一个Kernel就是一个要运行的程序，里面有大量的线程。Kernel启动一个Grid，里面有若干个Blocks，由用户设定。Grid可以理解为一个公司。
2.一个Block中包含多个线程，一个Block内部的线程共享Shared Memory，可以同步“_syncthreads()”。Block可以理解为一个部门。
3.线程和线程块具有唯一的标识。
程序对于GPU也有一定的映射关系：
其中，一个线程对应一个CUDA core或ALU，一个Block对应一个SM或SMX，一个Grid对应多个SM，最大为整个设备。
GPU内存和线程的关系：
1.一个线程有自己的存储器，叫做Local Memory，是私有的，只能自己访问。例如私人的办工作，电脑等资源。
2.每个Block，有内部线程可共享的Shared Memory，相当于部门中的打印机等共享资源。
3.每个Grid（Kernal）之间有共享的Global Memory，也就是GPU设备的全局存储。相当于多个公司都可以访问的大楼。
4.主机端的存储器（内存）可以和不同的GPU设备的内存（显存）相互拷贝数据。
如下图所示：
1.线程运算时与寄存器交互最快。
2.线程读取Local Memory时，由于该存储位于外部显存，所以速度相当较慢。
3.一个Block中共享Shared Memory。
4.各个Block中的线程都可以访问Global Memory。
5.Constant和Texture对于线程都是只读的存储。
6.Constant和Texture可以由主机端来读写。
三、CUDA编程模式
CUDA编程语言实际上是扩展的C语言（Extended C）
CUDA提供了许多特定的关键词。例如__device__,__global__,__host__等。
CUDA函数声明：
__device__ float DeviceFunc(); __global__ void KernelFunc(); __host__ float HostFunc();
1.由__device__修饰的函数声明表示该函数的执行位置是在GPU设备上，需要由其他GPU上的函数来调用。
2.由__global__修饰的函数是kernel函数，也是入口函数，在CPU上调用，在GPU上执行，必须返回void。
3.__host__修饰的函数是在主机端调用，也在主机端运行。
4.__device__和__host__可以同时作用于一个函数，说明该函数的操作在CPU和GPU上是一样的。
Kernel：
数据并行处理函数。
通过调用Kernel函数在设备端创建轻量级线程，线程由硬件负责创建并调度。
Kernel函数是在CPU上调用，然后再GPU上执行，是一个入口函数。
// 定义一个Kernel函数用__global__修饰 __global__ void VecAdd(float * A, float *B, float *C) { int i = threadIdx.x; C[i] = A[i] + B[i]; } int main() { //.....需要将A B都拷贝到显存 //.....在显存中分配C的空间 // 使用N个线程来计算 VecAdd<<<1, N>>> (A, B, C); return 0; }
线程层次Thread Hierarchies：
使用一个Block来处理：
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) { // 线程有N*N个，xy代表线程索引 int i = threadIdx.x; int j = threadIdx.y; C[i][j] = A[i][j] + B[i][j]; } int main() { // 使用一个Block int numBlocks = 1; // 每个Block有N*N个线程 dim3 threadPerBlock(N, N); // 这里使用一个Block，每个Block有N*N个线程 MatAdd <<<numBlocks, threadPerBlock>>> (A, B, C); return 0; }
上述代码中，只使用一个Block（一个部门），该Block中有N*N个线程（人员）。这个Block是一个2D的Block。
Block中的线程：
在G80和GT200显卡中，每个Block最多512个线程，而Fermi架构的GPU每个Block可以有1024个线程，可以查阅相关GPU手册。
每个Block相当于一个SM，即核心。所以该Block中的线程都是工作在相同的处理器核心中的。他们共享所在核心的Shared Memory。
使用多个Block处理：
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) { // 遍历每个Block的所有元素，并分别执行加法 int i = blockIdx.x * blockDim.x + threadIdx.x; int j = blockIdx.y * blockDim.y + threadIdx.y; if (i < N && j < N) { C[i][j] = A[i][j] + B[i][j]; } } int main() { // 每个Block有16*16个线程 dim3 threadPerBlock(16, 16); // 使用需要计算矩阵的尺寸来计算需要多少个Block dim3 numBlocks(N / threadPerBlock.x, N / threadPerBlock.y); // 这里使用一个Block，每个Block有N*N个线程 MatAdd <<<numBlocks, threadPerBlock>>> (A, B, C); return 0; }
Block与GPU核心（SM）数量关系：
当GPU只有2个SM（核心）时，程序有需要8个Block，则需要通过2个核心4次运算才能完成。
如果是4个核心，则需要2次运算才能完成。
四、数据传输
使用cudaMalloc在device上申请内存空间：
// 该指针用于存放device上分配空间的首地址 float * Md = 0; // 申请设备内存大小为size int size = 16 * 16 * sizeof(float); // 这里必须传入&Md，即Md指针的地址。 // 因为cudaMalloc会将分配好的设备内存首地址赋值给Md，这个Md只能在Device上使用，不能直接在CPU程序中赋值等 cudaMalloc((void **)&Md, size); // 释放Md指向的设备内存空间 cudaFree(Md);
内存传输：
Host to Host
Host to Device
Device to Host
Device to Device
对应一下四种操作：
// 申请设备内存大小为size int size = 16 * 16 * sizeof(float); // M指向CPU上的空间 float * M = (float *)malloc(size); float * M2 = (float *)malloc(size); // Md指向GPU上的空间 float * Md = 0; float * Md2 = 0; cudaMalloc((void **)&Md, size); cudaMalloc((void **)&Md2, size); // 从主机端内存中拷贝数据到Device的Global Memory中 cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice); // 从设备端拷贝数据到主机端 cudaMemcpy(M, Md, size, cudaMemcpyDeviceToHost); // 从主机端数据拷贝到主机端另一个空间，相当于memcpy cudaMemcpy(M, M2, size, cudaMemcpyHostToHost); // 从设备端拷贝数据到设备端另一个空间 cudaMemcpy(Md2, Md, size, cudaMemcpyDeviceToDevice);
五、矩阵乘法示例
// Md，Nd，Pd都是Width*Width的方阵，使用的Block中线程数也是W*W __global__ void MatMulKernel(float * Md, float * Nd, float * Pd, int Width) { // 横坐标为tx的列索引 int tx = threadIdx.x; // 纵坐标为ty的行索引 int ty = threadIdx.y; float Pvalue = 0; for (int k = 0;k < Width;++k) { // 处于tx的一行 float Mdelement = Md[ty * Width + k]; // 处于ty的一列 float Ndelement = Nd[k * Width + tx]; // Width元素做累加，得到坐标ty,tx的值 Pvalue += Mdelement * Ndelement; } // 将计算得到的ty,tx的值写入相应的位置 Pd[ty * Width + tx] = Pvalue; }
六、GPU上函数需要注意的问题
由于GPU特殊的工作情况和结构，在__Global__和__device__函数中，注意以下几点：
1.尽量少用递归（不鼓励）
2.不要使用静态变量
3.少用malloc（允许但不鼓励，因为并行的使用malloc，空间很快耗光）
4.小心通过指针实现的函数调用（注意指针时CPU端的还是GPU端的）
七、CUDA数据类型
矢量数据类型（同时适用于host和device代码）：
通过函数make_<type name>构造:
int2 i2 = make_int2(1, 2); float4 f4 = make_float4(1.0f, 2.0f, 3.0f, 4.0f); cout << i2.x << i2.y << endl; cout << f4.x << f4.y << f4.z << f4.w << endl;
八、CUDA支持的部分函数
部分函数列表：
面向Device端，更快，精度降低：
九、线程同步
块内线程可以同步（Block内）：
调用__syncthreads 创建一个barrier栅栏
每个线程在调用点等待块内所有线程执行到这个地方，然后再继续执行后续指令
Md[i] = Hd[j]; __syncthreads(); func(Md[i], Md[i + 1]);
如上述代码所示，func中同时需要Md[i]和Md[i+1]，当Md[i]准备好时，Md[i+1]不一定准备好了，所以需要在前面等待Md[i+1]准备好后，再继续执行func函数。
__syncthreads会导致线程的暂停，破坏了线程执行的独立性，并可能由于线程同步的位置不同（条件分支中使用同步）导致同步死锁。所以在使用同步时一定要小心。
十、线程调度
以G80显卡为例：
G80包含以下：
1.有16个核，也就是SM（8个绿色方框为一个SM）
2.每个SM有8个SP，也就是CUDA core或ALU（1个绿色方框）
3.每个SM最多可驻扎768个线程，128 X 6 = 768，每个SM可以保存6个上下文（蓝色部分）
4.总共可以同时驻扎12288个线程
5.但是由于只有128个CUDA core，同时也就只能执行128个线程
对于一个GPU设备来说，最大处理的线程量主要和CUDA core总量以及每个SM的上下文数量有关。但同时执行的线程数只与CUDA core数一致。
Warp：
针对Block中的线程，例如有64个线程（CUDA core），编号是连续的0-63。
假设一个Warp是32个线程组成（Warp的线程数和Block的线程数一般呈倍数关系，warpSize），则该Block中就有2个Warp，都运行在同一个SM上。第一个Warp线程编号为0-31，第二个Warp的线程编号为32-63。
Warp是线程调度的最小单位。
Warp的线程是天生同步的，也就是说他们必须是执行相同的指令流，当遇到分支可能导致执行的程序不同时（例如if else）则会出现串行的可能：
可能出现最差性能，就是1/N的性能。
例子：
1.如果一个SM分配了3个Block，其中每个Block含256个线程，那么总共有24个Warp（每个Warp 32个线程）。
2.GT200的一个SM最多可以驻扎1024个线程，那相当于1024/32=32个Warp。
3.假设每个Warp有32个线程，但每个SM只有8个SPs，如何分配？需要将一个Warp分成4份，然后在一个SM上轮换执行4次。流程如下：
指令已经预备
第一个周期8个线程进入SPs
在第二、三、四周期各进入8个线程
因此，分发一个Warp需要4个周期
4.对于目前的GPU来说，SM中所含的SP数一般都大于Warp含线程数量，所以以上分发流程一般不会再出现。
十一、内存模型
寄存器：
假设每个SM有8K个寄存器，有768个线程。则每个线程可以分到10个寄存器。
当超出限制时，则将因为Block的减少而减少。
例如，当一个线程需要用到11个寄存器，一个Block含256个线程。
本来如果每个线程使用寄存器不超出限制的时候，这个SM可以容纳3个Block（一个Block内的线程只能在同一个SM上执行），也就是刚好768个线程。
但由于寄存器超出限制，这个SM就只能容纳2个Block，即512个线程。所以就造成了资源的浪费。剩下未分配的SP也就只能闲着。
共享存储：
和寄存器类似原理类似。
假设每个SM最多8个Block，一共有16KB共享存储器。如果一个Block需要大于2K的共享存储器，则这个SM就不能容纳8个Block，同样造成资源浪费。
全局存储（显存）：
访存延时（100个周期），访存较慢，片外存储
Host主机可读写
GT200 GPU访存带宽150GB/s，容量4GB，新的显卡的访存带宽已达到300-500GB/s，容量达到8-32GB
位于不同存储的变量定义：
其中register和local存储我们不能操作。
__shared__定义存放在共享存储中的变量，这个变量只能是Block内部线程共享。
使用__device__关键字来定义全局存储（显存）中的变量。
__constant__用来定义常量（例如PI），存放在constant Memory中的。