如何自学GPU编程？

wuchao

如何自学GPU编程？

lqxxql

本文是读书期间的组会汇报材料整理而成，简述了基于CUDA的GPU编程原理和方法，主要内容包括：

• 1 硬件组成
  
• 2 变量类型
  
•    2.1 核函数运行参数
  
•    2.2 函数类型限定符
  
•    2.3 变量类型限定符
  
• 3 线程结构
  
• 4 内存结构
  
• 6 GPU编程难点
  
• 7 学习材料
  

图形处理器（英语：graphics processing unit，GPU），是一种专门在个人电脑和一些移动设备上做图像和图形相关运算工作的微处理器。



GPU编程特点

• 读取慢计算快
  
• 强并行弱串行
  
• 局限：要求不依赖上一次计算的结果，子集计算结果之间互不依赖
  

1 硬件组成



AMD Socket A 主板



AMD Socket A 主板

主板示意图



主板示意图

显卡结构




NVIDIA GPU 架构发展



NVIDIA GPU 架构发展

GP104 SM


GTX Pascal 架构系列






-arch compiler option specifies the compute capability that is assumed when compiling C to PTX code
-code that specifies the targeted architecture: For example, compiling with-code=sm_35 produces binary code for devices of compute capability 3.5





2 变量类型
主机端准备数据 -> 数据复制到GPU内存中 -> GPU执行核函数 -> 数据由GPU取回到主机



2.1 核函数运行参数
在设备端(Device)声明核函数
__global__ void kernel_name (param list) { }
在主机端(Host)调用时采用如下的形式：
kernel_name <<<Dg, Db, Ns, S>>> (param list);

• Dg：int型或者dim3类型(x,y,z),定义一个grid中的block如何组织,int型表示1维结构
  
• Db：int型或者dim3类型(x,y,z),定义一个block中的thread如何组织,int型表示1维结构
  
• Ns：size_t类型，可缺，默认为0。设置最多能分配共享内存大小,单位为byte,0表示不动态分配
  
• S  ：cudaStream_t类型,可缺省，默认为0,表示该核函数位于哪个流
  

例如：kernel_name<<<4, 8>>>(argumentt list);
表明grid为一维，有4个block，block为一维，每个block有8个线程，故此共有4*8=32个线程





2.2 函数类型限定符
函数类型限定符用来标识函数运行在主机还是设备上，函数由主机还是设备调用。
__global__

• 修饰的函数为 核函数。
  
• 运行在设备上；
  
• 可以由主机调用；
  
• 可以由计算能力大于3.2的设备调用；
  
• 必须有void返回类型；
  
• 调用时必须制定运行参数(<<< >>>)
  
• 该函数的调用时异步的，即可以不必等候该函数全部完成，便可以在CPU上继续工作；
  

__device__

• 运行在设备上；
  
• 只能由设备调用，即只能在其他__device__函数或者__global__函数中调用
  
• 编译器会内联所有认为合适的__device__修饰的函数；
  

__host__

• 运行在主机上；
  
• 只能由主机调用；
  
• 效果等同于函数不加任何限定符；
  
• 不能与__global__共同使用， 但可以和__device__联合使用；
  

kernel的限制：

• 仅能获取device memory
  
• 必须返回void类型
  
• 不支持可变数目参数
  
• 不支持静态变量
  
• 不支持函数指针
  
• 异步
  

2.3 变量类型限定符
变量类型限定符用来标识变量在设备上的内存位置
__device__ (单独使用时)

• 位于 global memory space
  
• 生命周期为整个应用期间(即与application同生死)
  
• 可以被grid内的所有threads读取
  

__constant__

• 可以和 __device__ 联合使用
  
• 位于 constant memory space
  
• 生命周期为整个应用期间
  
• 可以被grid内的所有threads读取
  

__shared__

• 可以和 __device__ 联合使用
  
• 位于一个Block的shared memory space
  
• 生命周期为整个Block
  
• 只能被同一block内的threads读写
  

__managed__

• 可以和 __device__ 联合使用
  
• 可以被主机和设备引用，主机或者设备函数可以获取其地址或者读写其值
  
• 生命周期为整个应用期间
  

__restrict__

• 该关键字用来对指针进行限制性说明，目的是为了减少指针别名带来的问题
  
• C99标准，用以缓解C语言中指针二义性的问题
  

3 线程结构


CUDA的线程结构： Grid, Block, Thread


Thread，GPU工作的最小单位,32个一束执行， 可以是一维，二维或三维;
Block，多个thread组成一个 block,最多包含1024个 ；

• 一个block的所有线程最好应当位于同一个处理器核心上，同时共享同一块内存，快速进行同步而不用担心数据通信壁垒
  
• 可以是一维，二维或三维
  

Grid，执行相同程序的多个 block组成grid；
不同 block 中的 thread 无法存取同一块共享的内存，程序不用担心显示芯片实际上能同时执行的 thread 数目限制



不同算力显卡的grid数量不同

从算力表可以看出，不同算力显卡的主要差别是grid数量，其它指标诸如一个grid中blocks数量、一个block中thread最大维度和数量等都是固定的。

4 内存结构
CUDA的线程逻辑结构 Grid, Block, Thread对应不同层次的硬件物理结构。CUDA的SM执行单元是线程束，任何线程数少于32都意味着没有充分利用硬件。CUDA编程的特殊之处在于，硬件的内存结构直接影响着线程分配策略，进行CUDA编程需要了解硬件结构。

• 每一个 thread 都有自己的一份 register 和 local memory 的空间，32个线程组成线程束
  
• 同一个 block 中的每个 thread 则有共享的一份 share memory;
  
• 所有的 thread 都共享一份 global memory、constant memory和 texture memory;
  
• 不同的 grid 则有各自的 global memory、constant memory和 texture memory
  
• global memory：GPU和CPU都可以进行写操作
  

也就是register 空间和share空间分别为单个thread和block私有，在同一个grid里面的所有thread共享global 空间、constant 空间和 texture 空间，不同grid中的这三种空间则互不通气。


5 线程索引
由一个单独的kernel启动的所有线程组成一个grid，几个CUDA内置变量：

• blockIdx ：block的索引，blockIdx.x表示block的x坐标
  
• threadIdx ：线程索引，同理blockIdx
  
• blockDim ：block维度，上图中blockDim.x=5
  
• gridDim ：grid维度，同理blockDim
  

维度相当于说明长宽有多大，索引相当于相应维度的位置。
线程索引相当于盒子套盒子。线程id具有唯一性，所以必须要知道所在grid和block的编号和总体维度



在kernel里，线程索引唯一，为了确定一个线程的索引，以2D为例：

• 线程和block索引
  
• 矩阵中元素坐标
  

首先可以将thread和block索引映射到矩阵坐标：
ix = threadIdx.x + blockIdx.x * blockDim.x
iy = threadIdx.y + blockIdx.y * blockDim.y
之后可以利用上述变量计算线性地址：
idx = iy * nx + ix





•一个Block中的可开辟的线程数量一般不超过1024
•矢量特别长会出现问题，采用多个线程块(Block)来解决线程不足的问题
• 假如每个线程块包含128个线程，则需要的线程块的数量为 size / 128
• 为了避免不能整除带来的问题，可以稍微多开一点 (size + 127) / 128，但需要增加判断条件来避免越界

案例：





6 GPU编程难点

• 需要对硬件结构由一定的了解
  
• 串行编程到并行编程思维转换。循环方式，异步编程
  
• 恰当的内存分配策略，全局内存，共享内存，静态内存
  
• 调试难，需要特殊的编译器
  

7 学习材料