新房装修设计东莞网站建设 乐云seo
CPU 和 GPU 由于结构的不同，具有不同的特点：
- CPU：擅长流程控制和逻辑处理，不规则数据结构，不可预测存储结构，单线程程序，分支密集型算法
- GPU：擅长数据并行计算，规则数据结构，可预测存储模式
在现在的计算机体系架构中，要完成 CUDA 并行计算，单靠 GPU 一人之力是不能完成计算任务的，必须借助 CPU 来协同配合完成一次高性能的并行计算任务。
一般而言，并行部分在 GPU 上运行，串行部分在 CPU 运行，这就是异构计算。
异构计算的意思就是不同体系结构的处理器相互协作完成计算任务。CPU 负责总体的程序流程，GPU 负责具体的计算任务，当 GPU 各个线程完成计算任务后，将 GPU 那边计算得到的结果拷贝到 CPU 端，完成一次计算任务。
Definitions:
- Device --> GPU
- Host --> CPU
- Kernel --> function that runs on the devcie
1 CUDA 线程模型
线程是程序执行的基本单元，CUDA 的并行计算是通过成千上万个线程的并行执行来实现的。
CUDA的线程模型从小往大：
- Thread：线程，并行的基本单位
- Thread Block：线程块，互相合作的线程组，线程块有如下几个特点：
- 允许彼此同步
- 可以通过共享内存快速交换数据
- 以 1 维、2 维或 3 维组织
- Grid：一组线程块
- 以 1 维、2 维组织
- 共享全局内存
- Kernel：在 GPU 上执行的核心程序，这个 kernel 函数是运行在某个 Grid 上的。
每一个 block 和每个 thread 都有自己的 ID，通过相应的索引找到相应的线程和线程块。
- threadIdx，blockIdx
- Block ID: 1D or 2D
- Thread ID: 1D, 2D or 3D
GPU 上很多并行化的轻量级线程。kernel 在 device 上执行时实际上是启动很多线程，一个 kernel 所启动的所有线程称为一个网格（grid），同一个网格上的线程共享相同的全局内存空间。
grid 是线程结构的第一层次，而网格又可以分为很多线程块（block），一个线程块里面包含很多线程，这是第二个层次。
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)，二级索引(block.xIdx, block.yIdx, block.zIdx)，可以用来定位到指定的线程。这就是 CUDA 的线程组织结构。
SP 和 SM 的联系与区别：
- SP：最基本的处理单元，streaming processor，也称为CUDA core。最后具体的指令和任务都是在 SP 上处理。GPU进行并行计算，也就是很多个 SP 同时做处理。
- SM：多个 SP 加上其他的一些资源组成一个 streaming multiprocessor。也叫 GPU 大核，其他资源如：warp scheduler，register，shared memory 等。register 和 shared memory 是 SM 的稀缺资源，CUDA 将这些资源分配给所有驻留在 SM 中的 threads。因此，这些有限的资源就使每个 SM 中 active warps 有非常严格的限制，也就限制了并行能力。
每个 SM 包含的 SP 数量依据 GPU 架构而不同，Fermi架构 GF100 是 32 个，GF10X 是 48 个，Kepler 架构都是 192 个，Maxwell 都是 128 个。
简而言之，SP 是线程执行的硬件单位，SM 中包含多个 SP，一个 GPU 可以有多个 SM（比如16个），最终一个 GPU 可能包含有上千个 SP。软件逻辑上所有 SP 是并行的，但是物理上并不是所有 SP 都能同时执行计算（比如我们只有 8 个 SM 却有 1024 个线程块需要调度处理），因为有些会处于挂起，就绪等其他状态。
从硬件角度和软件角度理解 CUDA 的线程模型：
- 每个线程由每个线程处理器（SP）执行
- 线程块由多核处理器（SM）执行
- 一个 kernel 其实由一个 grid 来执行，一个 kernel 一次只能在一个 GPU 上执行
block 是软件概念，一个 block 只会由一个 sm 调度，程序员在开发时，通过设定 block 的属性，告诉 GPU 硬件，我有多少个线程，线程怎么组织。而具体怎么调度由 sm 的 warps scheduler 负责，block 一旦被分配好 SM，该 block 就会一直驻留在该 SM 中，直到执行结束。一个 SM 可以同时拥有多个 blocks，但需要序列执行。
2 CUDA 内存模型
CUDA 中的内存分为以下几个层次：
- 每个线程都用自己的 registers（寄存器）
- 每个线程都有自己的 local memory（局部内存）
- 每个线程块内都有自己的 shared memory（共享内存），所有线程块内的所有线程共享这段内存资源
- 每个 grid 都有自己的 global memory（全局内存），不同线程块的线程都可使用
- 每个 grid 都有自己的 constant memory（常量内存）和 texture memory（纹理内存），不同线程块的线程都可使用
线程访问这几类存储器的速度: register > local memory >shared memory > global memory
3 CUDA 编程模型
3.1 指定代码在哪里跑
通过关键字可以表示某个程序在 CPU 上跑还是在 GPU 上跑。比如用__global__定义一个 kernel 函数，就是 CPU 上调用，GPU 上执行，注意__global__函数的返回值必须设置为void。
| excuted on | only called from | |
| device float DeviceFunc() | device | device |
| global void KernelFunc() | device | host |
| host HostFunc() | host | host |
3.2 CPU 和 GPU 间的数据传输
在 GPU 内存分配回收内存的函数接口：
cudaMalloc(): 在设备端分配 global memorycudaFree(): 释放存储空间
CPU 的数据和 GPU 端数据做数据传输的函数接口是一样的，他们通过传递的函数实参（枚举类型）来表示传输方向：
cudaMemcpy(void dst, void src, size_t nbytes,
enum cudaMemcpyKind direction)
enum cudaMemcpyKind:
cudaMemcpyHostToDevice（CPU 到 GPU）cudaMemcpyDeviceToHost（GPU 到 CPU）cudaMemcpyDeviceToDevice（GPU 到 GPU）
3.3 用代码表示线程组织模型
可以用dim3类来表示网格和线程块的组织方式，网格 grid 可以表示为一维和二维格式，线程块 block 可以表示为一维、二维和三维的数据格式。( Dim3类型: cuda 的内置类型在定义类型为 dim3 的变量时，未指定的任何组件都将初始化为 1。)
dim3 DimGrid(100, 50); //5000个线程块，维度是100*50
dim3 DimBlock(4, 8, 8); //每个线层块内包含256个线程，线程块内的维度是4*8*8
3.4 计算线程编号
- 使用 N 个线程块，每一个线程块只有一个线程
dim3 dimGrid(N);
dim3 dimBlock(1);
此时计算线程编号：
threadID = blockIdx.x;
其中threadId的取值范围为 0 到 N-1。对于这种情况，可以将其看作是一个列向量，列向量中的每一行对应一个线程块。列向量中每一行只有1个元素，对应一个线程。
- 使用 M×N 个线程块，每个线程块 1 个线程
线程块是2维的，故可以看做是一个M*N的2维矩阵，其线程号有两个维度，即：
dim3 dimGrid(M, N);
dim3 dimBlock(1);
这里，blockIdx.x 取值 0 到 M-1, blcokIdx.y 取值 0 到 N-1。
这种情况一般用于处理 2 维数据结构，比如 2 维图像。每一个像素用一个线程来处理，此时需要线程号来映射图像像素的对应位置，
pos = blockIdx.y * blcokDim.x + blockIdx.x; //其中gridDim.x等于M
- 使用一个线程块，该线程具有 N 个线程
dim3 dimGrid(1);
dim3 dimBlock(N);
此时线程号的计算方式为:
threadID = threadIdx.x;
其中 threadId 的范围是 0 到 N-1，对于这种情况，可以看做是一个行向量，行向量中的每一个元素的每一个元素对应着一个线程。
- 使用 M 个线程块，每个线程块内含有 N 个线程
dim3 dimGrid(M);
dim3 dimBlock(N);
这种情况，可以把它想象成二维矩阵，矩阵的行与线程块对应，矩阵的列与线程编号对应，那线程号的计算方式为:
threadId = threadIdx.x + blcokIdx * blockDim.x;
这里就是把二维的索引空间转换为一维索引空间的过程。
- 使用 M×N 的二维线程块，每一个线程块具有 P×Q 个线程
dim3 dimGrid(M, N);
dim3 dimBlock(P, Q);
其索引有两个维度:
threadId.x = blockIdx.x * blockDim.x + threadIdx.x;
threadId.y = blockIdx.y * blockDim.y + threadIdx.y;