快速上手一下cuda.
bg
GPU并不是一个独立运行的计算平台,而需要与CPU协同工作,可以看成是CPU的协处理器,因此当我们在说GPU并行计算时,其实是指的基于CPU+GPU的异构计算架构。
在异构计算架构中,GPU与CPU通过PCIe总线连接在一起来协同工作,CPU所在位置称为为主机端(host),而GPU所在位置称为设备端(device)。
GPU包括更多的运算核心,其特别适合数据并行的计算密集型任务,如大型矩阵运算,而CPU的运算核心较少,但是其可以实现复杂的逻辑运算,因此其适合控制密集型任务。
另外,CPU上的线程是重量级的,上下文切换开销大,但是GPU由于存在很多核心,其线程是轻量级的。因此,基于CPU+GPU的异构计算平台可以优势互补,CPU负责处理逻辑复杂的串行程序,而GPU重点处理数据密集型的并行计算程序,从而发挥最大功效。
支持的编程接口:py,c,c++,fortran,llvm.
编程基础
重要概念:
- host:指代CPU及其内存
- device:指代GPU及其内存
程序中包含host和device程序,分别在CPU和GPU上运行,两者之间可以通信,进行数据拷贝。
执行流程
- 分配host内存,数据初始化
- 分配device内存,从host上拷贝数据到device
- 调用CUDA核函数到device完成指定运算
- device上运算结果拷贝到host
- 释放内存
最重要的过程是3.调用CUDA核函数在device执行并行运算。
- kernel:在device上线程中并行执行的函数。
- 核函数:
__global__
声明,调用时指定<<<grid,block>>>
来制定kernel执行的线程数量。 - cuda中,每个线程都要执行核函数,所以每个线程分配一个唯一线程号threadID,这个值可以用
threadIdx
来获得。
CUDA通过函数类型区分host和device上的函数。
__global__
:device执行,host调用(一些特定GPU可以从device调用),返回类型必须void
,不支持可变参数(必须const),不能成为类成员函数。__global__
定义的kernel是异步的,host不会等待kernel执行完就执行下一步。__device__
:device执行,只能从device调用,不能与__global__
同时用。__host__
:host上执行,只能从host调用,一般省略,不能与__global__
同时用,但能和__device__
同时,此时函数会在device和host都编译。
kernel线程层次结构
所有线程:网格(grid)——>分为线程块(block)——>分为多个线程(thread)
grid和block都是定义为dim3
类型的变量,dim3
可以看成是包含三个无符号整数(x,y,z)成员的结构体变量,在定义时,缺省值初始化为1。也就是说最高3维。
所以grid/block可以灵活的改变为1/2/3维度。
一个线程块上的线程是放在同一个流式多处理器(SM)上的,但是单个SM的资源有限,这导致线程块中的线程数是有限制的,现代GPUs的线程块可支持的线程数可达1024个。有时候,我们要知道一个线程在blcok中的全局ID,此时就必须还要知道block的组织结构,这是通过线程的内置变量blockDim
来获得。它获取线程块各个维度的大小。另外线程还有内置变量gridDim
,用于获得网格块各个维度的大小。
CUDA内存模型
- 每个线程:私有本地内存
- 每个线程块:共享内存:线程块中所有线程共享,生命周期与线程块一致。
- 每个网格:全局内存:所有线程可以访问。
- 此外还有:常量内存(constant)、纹理内存(texture)等。