CPU和GPU架构
处理器结构主要要考虑的两个指标：延迟和吞吐量。
- 延迟：一条指令从发布到返回结果所经历的时间。
- 吞吐量：单位时间内处理的指令的条数。
CPU： 延迟到向内核
GPU： 吞吐导向内核
CPUs
内存大
- 多级缓存结构提高访问速度
有复杂的控制单元
- 分支预测机制（判断运算分支）
- 流水线数据前送机制（判断经常用到的数据，并将其前送）
运算单元强大
- 整型浮点型运算速度快
GPUs
缓存小
- 提高内存吞吐
控制简单
- 没有分支预测
- 没有数据转发
精简运算单元
- 多长延时流水线以实现高吞吐量
- 需要大量线程来容忍延迟
GPU&CPU特点
CPUs：连续计算部分，延迟优先。CPU比GPU单条复杂指令延迟快10倍以上。
GPUs：并行计算部分，吞吐优先。GPU比CPU单位时间内执行指令数量10倍以上。
什么样的问题适合GPU
计算密集：数值计算的比例要远大于内存操作，因此内存访问的延时可以被计算掩盖。
数据并行：大任务可以被拆解为执行相同指令的小任务，因此对复杂流程控制的需求较低。
- 内存访问次数少
- 控制简单
- 计算简单
- 并行度高
GPU编程与CUDA
CUDA (Compute Unified Device Architecture)，初衷是为GPU增加一个易用的编程接口，让开发者无需学习复杂的着色语言或者图形处理原语。
OpenCL(Open Computing Languge)是2008年发布的异构平台并行编程的开发标准，也是一个编程框架。OpenCL相比于CUDA，支持的平台更多，除了GPU还支持CPU、DSP、FPGA等设备。
GPU编程并行计算整体流程
void GPUkernel(float* A, float* B, float* C, int n){
1. // Allocate device memory for A, B and C
// copy A and B from host memory (CPU) to device memory (GPU)
2. // Kernel launch code - to have the device
// to perform the actual vector addition
3. // copy C from the device memory to host memory
// free device vectors
}
CUDA编程术语
硬件侧
-
Device=GPU
-
Host=CPU
-
Kernel=GPU上运行的函数
-
CUDA中的内存模型分为以下几个层次：
- 每个线程处理器（SP）都用自己的registers（寄存器）
- 每个SP都有自己的local memory（局部内存）， register和local memory只能被线程自己访问
- 每个多核处理器（SM）内都有自己的shared memeory（共享内存），shared memory可以被线程块内所有线程访问，不过shared memory的大小较小
- 一个GPU的所有SM共有一块global memory（全局内存）， 不用线程块的线程都可使用
软件侧
- CUDA中的内存模型分为以下几个层次：
- 线程处理器（SP）对应线程（thread）
- 多核处理器（SM）对应线程块（thread block）
- 设备端（device）对应线程块组合体（grid）
- 一个kernel其实由一个grid来执行
- 一个kernel一次只能在一个GPU上执行
线程块：可扩展的集合体
线程是内存模型在软件侧最基本的执行单位，线程块就是线程的组合体。
将线程数组分成多个块
- 块内的线程通过共享内存、原子操作和屏障同步进行协作（shared memory， atomic operations and barrier synvhronization）
- 不同块中的线程不能协作
特点：
- 线程块内的所有线程的计算和内存都是互相独立的
- 线程块内的共享内存都可访问
- 可以用公有时钟来同步
网格（grid）：并行线程块组合体
CUDA核函数由线程网格（数组）执行
- 每个线程都有一个索引，用于计算内存地址和做出控制决策
特点：
- 网格内的线程块是互相独立的
- 网格内的全局内存可由网格内的所有线程块访问
- 可用公有时钟同步网格内所有线程块
线程块id&线程id：定位独立线程
也是通过索引来确定线程的地址， 但GPU不止需要线程的索引，还需要线程块的索引（可能有多维）。
dim3 dimGrid(M, N);
dim3 dimBlock(P, Q, S);
threadId.x = blockIdx.x * blockDim.x + threadIdx.x;
threadId.y = blockIdx.y * blockDim.y + threadIdx.y;
线程束（warp）
- SM采用的SMIT（Single-Instruction， Multiple-Thread， 单指令多线程）架构，warp（线约束）是最基本的执行单元，一个warp包含32个并行thread，这些thread以不同数据资源执行相同的指令。warp本质上是线程在GPU上运行的最小单元。
- 当一个kernel被执行时，grid中的线程块被分配到SM上，一个线程块的thread只能在一个SM上调度，SM一般可以调度多个线程块，大量的thread可能被分到不同的SM上。每个threade拥有它自己的程序计数器和状态寄存器，并且用该线程自己的数据执行指令，这就是所谓的Single Instruction Multiple Thread（SIMT）。
- 由于warp的大小为32， 所以block包含的thread的大小一般设置为32的倍数。
并行计算实例：向量相加
CPU
// Compute vector sum C=A+B
void vecAdd(float* A, float* B, float *c, int n){
for (i = 0, i < n, i++)
c[i] = A[i] + B[i];
}
int main(){
//Memory alocation for A_h, B_h and C_h
// I/O to read A_h and B_h, N elements
vecAdd(A_h, B_h, C_h, N);
}
GPU
void vecAdd(float* A, float* B, float* C, int n){
int size = n * sizeof(float);
float* A_d, *B_d, *C_d;
// Transfer A and B to device memory
cudaMalloc((void **) &A_d, size);
CUDAMemcpy(A_d, A, size, cudaMemcpyHostToDevice);
cudaMalloc((void **) &B_d, size);
cudaMemcpy(B_d, B, size, cudaMemcpyHostToDevice);
// Allocate device memory for
cudaMalloc((void **) &C_d, size);
// Kernel invocation code - to be shown later
// Transfer C from device to host
cudaMemcpy(C, C_d, size, cudaMemcpyDeviceToHost);
// Free device memory for A, B, C
cudaFree(A_d);
cudaFree(B_d);
cudaFree(C_d);
}