CUDA编程模型

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，用于获得网格块各个维度的大小。