GPU计算基础知识
- CUDA编程模型是一个异构模型,需要CPU和GPU协同工作
- 在CUDA中,host和device是两个重要的概念,我们用host指代CPU及其内存,而用device指代GPU及其内存。
- CUDA程序中既包含host程序,又包含device程序,它们分别在CPU和GPU上运行。
- host与device直接可以进行通信,这样它们之间可以进行数据拷贝。
CUDA程序执行流程
1.分配host内存,并进行数据初始化;
2.分配device内存,并从host将数据拷贝到device上;
3.调用CUDA的核函数在device上完成指定的运算;
4.将device上的运算结果绪贝到host上(性能)
5.释放device和host上分配的内存。
CUDA程序
- 上面流程中最重要的一个过程是调用CUDA的核橱数来执行并行计算.
- kernel是CUDA中一个重要的概念,kernel是在device上线程中并行执行的函数
- 核函数用_global_符号声明,在调用时需要用<<<grid, block>>>采指定kernel要执行的线程数量
- 在CUDA中,每一个线程都要执行核函数,并且每个线程会分配二个唯一的线程号thread ID,这个ID值可以通过核函数的内置变量threadldx来获得。
GPU代码片段
// Kernel定义
global_void vec_add(doublex,double*y, double*z, int n){
int i= get_tid();// user-defined macro/function
if (i<n)z0]=x[门]+yl;
}
int main(){
int N= 100oo00;// 1M
int bs= 256;
int gs =(N +bs- 1)/bs;
/kernel,call GPU
vec_add<<<gs, bs>>>(x,y,z,N);
}
CUDA程序层次结构
-
GPU上很多并行化的轻量级线程。
-
kernel在device上执行时实际上是后动很多线程,一个kernel所启动的所有线程称为一个网格(grid)
-
回一个网格上的线程共享相同的全局内存空间,grid是线程结构的第一层次
- 网格又可以分为很多线程块(block),一个线程块里面包含很多线程,这是第二个层次。
- warp:32个线程一组,这是第三个层次
- grid和 block都是定义为dim3类型的变量
- dim3可以看成是包含三个无符号整数(x,, y,z)成员的结构体变量,在定义时,缺省值初始化为1。
- grid和block可以灵活地定义为1-dim,2-dim以及3-dim结构
- 定义的grid和block如下所示,kernel在调用时也必须通过执行配置<<<grid,block>>>来指定kernel所使用的线程数及结构。
- 不同GPU 架构, grid和block的维度有限制
CUDA程序调用
dim3 grid(3,2);
dim3 block(5,3);
kernel_fun<< grid, block>>>(prams...);
dim3 grid(128);
dim3 block(256);
kernel_fun<<< grid, block>>>(prams...);
dim3 grid(100,120);
dim3 block(16,16,1);
kernel_fun<<< grid, block >>>(prams...);
CUDA程序层次结构
GPU是异构模型,所以需要区分host和device上的代码,在CUDA中是通过函数类型限定词并区别host和device上的函数,主要的三个函数类型限定词如下:
1)_global_:在device上执行,从host中调用(一些特定的GPU也可以从device上调用),返回类型必须是void,不支持可变参数,不能称为类成员函数。
2)注意用_global_定义的kernel是异步的,这意味着host不会等待kernel执行完就执行下一步。
3)_device_:在device上执行,单仅可以从device中调用,不可以和_global_同时调用。
4)_host_:在host上执行,仅可以从host上调用,一般省略不写,不可以和_global_同时用,但可和_device_,此时函数会在device和host都编译。
CUDA内置变量
-
一个线程需要两个内置的坐标变量(blockldx,threadldx)来唯一标识,它们都是dim3类型变量,其中blockldx指明线程所在grid中的莅置,而threaldx指明线程所在block中的位置:
- threadIdx包含三个值: threadIdx.x,threadldx.y, threadldx.z
- blockldx同样包含三个值:blockldx.x,blockldx.y, blockldx.z
- .一个线程块上的线程是放在同一个流式多处理器(SM)上的
- 单个SM的资源有限,这导致线程块中的线程数是有限制的,现代GPUs的线程块可支持的线程数可达1024个。
- 有时候,我们要知道一个线程在blcok中的全局ID,此时就必须还要知道block的组织结构,这是通过线程的内置变量blockDim来获得。它获取线程各个维度的大小。
- 对于一个2-dim的block,线程的ID值为,如果是3-dim的block,线程的ID值为。另外线程还有内置变量gridDim,用于获得网格块各个维度的大小。