1. CPU与GPU
- CPU–延迟导向设计
内存大:多级缓存结构提高访存速度
控制复杂:分之预测机制和流水线数据传送
运算单元强大:整型浮点型复杂运算速度快 - GPU–吞吐导向设计
缓存小:提高内存吞吐
控制简单:没有分支预测和数据前传
精简运算单元:多 长延时流水线以实现高吞吐量 和 小大量线程来容忍延迟 - 各自特点
CPU延迟优先,单条复杂指令比GPU快10倍以上
GPU吞入优先,单位时间内执行指令数量10倍以上
显存和内存的联系:都是用来暂存资料的存储空间,只是显存帮GPU,内存帮CPU
2. CUDA和OpenCL
CUDA(Compute Unified Device Architecture)为GPU增加一个易用的编程接口;OpenCL与CUDA作用一致,只是它支持的平台更多,除了GPU外还支持CPU、DSP、FPGA等
a. CUDA编程并行计算整体流程
- Device:GPU
- Host:CPU
- Kernel:GPU上运行的函数
void GPUkernel(float* A, float* B, float* C, int n)
{
1. // Allocate device memory for A, B, and C
// copy A and B to device memory
2. // Kernel launch code – to have the device
// to perform the actual vector addition
3. // copy C from the device memory
// Free device vectors
}
b. CUDA中的内存模型
硬件上(从内到外):
- 每个线程处理器(SP)都有自己的寄存器和局部内存,并且只能自己访问。
- 每个多核处理器(SM)内都有自己的共享内存,可以被线程块内所有线程访问。
- 一个GPU的所有SM共有一块全局内存,不用线程块的线程都可使用。
上面的术语对应到软件上为:
- SP–线程,
- SM–线程块,
- 设备端(device)–线程块组合体(grid)
并且一个kernel函数只能由一个grid来执行,且一次只能在一个GPU上执行。
块内的线程通过共享内存、原子操作和屏障同步进行协作,并且不同线程块的线程不能协作;
c.线程块id&线程id
一个Kernel函数运行在一个grid上,当我们编程的时候,要指明哪个线程执行哪些,并且在运行的时候也要去调用固定线程的寄存器和内存,于是就需要地址来对每个线程和线程块定唯一。
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;
d.线程束(warp)
- SM采用的SIMT(Single-Instruction, Multiple-Thread,单指令多线程)架构,warp(线程束)是最基本的执行单元,一个warp包含32个并行thread,这些thread以不同数据资源执行相同的指令。warp本质上是线程在GPU上运行的最小单元。
- 当一个Kernel被执行时,grid中的线程块被分配到SM上,一个线程块的thread只能在一个SM上调度,SM一般可以调度多个线程块,大量的thread的可能被分到不同的SM上。每个thread拥有它自己的程序计数器和状态寄存器,并且用该线程自己的数据执行指令,这就是所谓的SIMT。
- 由于warp的大小为32,所以block所含的thread大小一般都设置为32的倍数。
3. 并行计算实例:向量相加
主要要做的事:
设备端代码(计算作用):
- 读写线程的寄存器
- 读写Grid中全局内存
- 读写block中共享内存
主机端代码:
- 申请内存和显存的空间,进行数据的转移和拷贝,最后得到结果到主机端后,进行内存的释放。
a. 使用到的基本函数介绍
#include “cuda.h”
- cudaMalloc():cudaError_t cudaMalloc (void **devPtr, size_t size),申请显存函数,其中参数一是申请地址,参数二是申请内存的大小,常用sizeof(float)×N
- cudaFree():cudaError_t cudaFree ( void* devPtr ),释放显存,参数是指向对象的指针。
- cudaMemcpy():cudaError_t cudaMemcpy (void *dst, const void *src, size_t count, cudaMemcpyKind kind)
- 内存数据复制传递,参数一是目的地,参数二是出发地,参数三是大小(同样使用sizeof()×形式)
- 参数三是拷贝方向:目前支持四种,cudaMemcpyHostToDevice,cudaMemcpyDeviceToHost,cudaMemcpyDeviceToDevice,cudaMemcpyDefault
b.核函数简单介绍及调用
- 在GPU上执行的函数。
- 一般通过标识符__global__修饰。
- 调用通过<<<参数1,参数2>>>,用于说明内核函数中的线程数量,以及线程是如何组织的。
- 以网格(Grid)的形式组织,每个线程格由若干个线程块(block)组成,而每个线程块又由若干个线程(thread)组成。
- 调用时必须声明内核函数的执行参数。
- 在编程时,必须先为kernel函数中用到的数组或变量分配好足够的空间,再调用kernel函数,否则在GPU计算时会发生错误(segmentation fault)。
__ global __ 标志核函数
• 核函数必须返回 void
__ device __ & __ host __ 可以一起用
c. CUDA编译流程
有三种