GPU编程

发表于

学习GPU编程，希望能够用到开发的深度学习框架中

官方教程

CUDA示例

GPU和CPU的区别

GPU基于高吞吐设计的：多ALU，小cashe，多Thread并行，控制单元能够将多个访问和并未较少的访问。

CPU基于低延时设计的：强大ALU，复杂逻辑控制单元，大cache，少register

两种调用GPU方式

  • 驱动API：直接调用底层驱动
  • 运行API：调用封装好的API

通讯模式

  • 映射(MAP):输入输出一对一，例如求数组元素平方。
  • 聚合(GATHER) ：多对一，每相邻三个元素平方。
  • 分散(scatter) ：一对多
  • 模板(stencil) ：以固定的模式读取相邻的内存数值，serveral-to-one
  • 转置(transpose) ：一对一，位置模式的转换
  • 压缩(reduce) ：多对一(all-to-one)
  • 重排(scan/sort) ：多对多

CUDA线程通信

线程通信在CUDA中有三种实现方式：

  1. 共享存储器；
  2. 线程 同步；
  3. 原子操作；

硬件模式

  • SM(流处理器)：包含很多简单处理器和memory，每个simple process就是来处理线程的。每个GPU至少一个SM，每个SM并行而独立运行。
  • Kernel(核)：可以理解为C/C++的一个function。function由内存块来运行
  • thread blocks(线程块)：一些线程block组成线程块，一个线程块包含多个线程。

编程模型

  • 最大特点：对线程块将在何处，合适运行不作保证。
    • 优点：充分利用硬件，灵活；不需要线程互相等待；可扩展性强
    • 缺点：对于线程块在哪个流处理器运行无法进行任何假设；无法获取块之间的明确通信（并行死锁，线程退出）。

CUDA编程模型的原则

  1. 所有在同一个线程块上的线程必然会在同一时间运行在同一个SM上
  2. 同一个内核的所有线程块必然会全部完成了后，才会运行下一个内核

内存模型

  • 每个线程有自己的local memory,局部变量
  • 每个线程块有shared memory，线程块之间共享变量
  • thread blocks有global memory 同步主机内存 host memory
  • 速度local > shared >> global > host
  • 将数据切割的越小越好

线程束

GPU采用单指令多线程(SIMT)架构来管理和执行线程，每32个线程为一组，称为线程束。

线程束分化

同一个线程块必须按照相同的指令执行，当一个线程束中出现分支时，会执行符合条件的线程，禁用不符合条件的线程，相当于分支两侧顺序执行了。

对于短指令，cuda会将断定变量设为1或者0，用断定指令替代分支指令。

同步性synchronisation和屏障barrier

  • 不同的线程在共享和全局内存中读写数据需要有先后的控制，所以引入了同步性的概念。
  • 屏障的作用：用来控制多个线程的停止与等待，当所有线程都到达了屏障点，程序才继续进行。

CUDA编程模CUDA

程序中CPU是主导地位，负责完成以下的事情：

  1. 从CPU同步数据到GPU
  2. 从GPU同步数据到CPU
    （1、2使用cudaMemcpy）
  3. 给GPU分配内存（cudaMalloc）
  4. 加载Kernel到GPU上，launch kernel on GPU
    ​

CUDA代码

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
#include <stdio.h>
//kernel 表示要在GPU上运行
__global__ void square(float* d_out,float* d_in){
  int idx = threadIdx.x;
  float f = d_in[idx];
  d_out[idx] = f * f;
}
int main(int argc,char** argv){
  const int ARRAY_SIZE = 8;
  const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float);
  // generate the input array on the host
  // h_* represents variable in CPU
  float h_in[ARRAY_SIZE];
  for(int i=0;i<ARRAY_SIZE;i++){
	h_in[i] = float(i);
  }
  float h_out[ARRAY_SIZE];
  // declare GPU memory pointers
  float* d_in;
  float* d_out;
  // allocate GPU memory
  cudaMalloc((void**) &d_in,ARRAY_BYTES);
  cudaMalloc((void**) &d_out,ARRAY_BYTES);
  // transfer the array to GPU
  cudaMemcpy(d_in,h_in,ARRAY_BYTES,cudaMemcpyHostToDevice);
  // launch the kernel
  //1 表示加载个1 个线程块上面
  //第二个参数表示每个线程块里面有多少个线程
  square<<<1,ARRAY_SIZE>>>(d_out,d_in);
  // copy back the result array to the GPU
  cudaMemcpy(h_out,d_out,ARRAY_BYTES,cudaMemcpyDeviceToHost);
  // print out the resulting array
  for(int i=0;i<ARRAY_SIZE;i++){
	printf("%f",h_out[i]);
	printf(((i%4) != 3) ? "\t" : "\n");
  }
  // free GPU memory allocation
  cudaFree(d_in);
  cudaFree(d_out);
  return 0;
}
编译命令： nvcc -o square square.cu

CUDA代码的高效策略

高效公式
最大化计算强度 = Math（数学计算量）/Memory(每个线程的内存)

  1. 最大化每个线程的计算量
  2. 最小化每个线程的内存读取速度
    • 每个线程读取的数据量少
    • 每个线程读取的速度快（尽量存在本地内存 > 共享内存 or 合并全局内存）
  3. 避免线程发散
    线程发散：同一个线程块中的线程执行不同内容的代码
    例子： kernel中if语句 或者 长度不一致的循环语句

Kernel加载方式

注意：Kernel的加载中，自定义的线程数，线程块的数量等都不要超过系统本身的设定，否则，会影响机器的效率。

CUDA中的各种内存的代码使用

  • 局部变量：kernel里面定义的变量
  • 全局变量：与GPU进行交互时需要拷入，拷出
  • 共享变量：需要进行barrier的操作。加载kernel时第三个参数可以表述共享变量大小。

CUDA同步操作

原子操作

原子操作：对于有很多线程需要同时读取或写入相同的内存时，保证同一时间只有一个线程能进行操作。

  1. 只支持某些运算(加、减、最小值、异或运算等，不支持求余和求幂等)和数据类型（整型）
  2. 运行顺序不定
  3. 安排不当，会使速度很慢（因为内部是个串行的运行）

同步函数

_syncthreads（）
​ 线程块内线程同步
​ 保证线程块内所有线程都执行到统一位置

_threadfence()
​ 一个线程调用__threadfence后，该线程在该语句前对全局存储器或共享存储器的访问已经全部完成，执行结果对grid中的所有线程可见。

_threadfence_block()
​ 一个线程调用__threadfence_block后，该线程在该语句前对全局存储器或者共享存储器的访问已经全部完成，执行结果对block中的所有线程可见。

以上两个函数的重要作用是，及时通知其他线程，全局内存或者共享内存内的结果已经读入或写入完成了。

CPU/GPU同步

  • cudaStreamSynchronize()/cudaEventSynchronize()
  • 主机端代码中使用cudaThreadSynchronize():实现CPU和GPU线程同步
  • kernel启动后控制权将异步返回，利用该函数可以确定所有设备端线程均已运行结束；

并行化高效策略

  • 规约(Reduce)
  • 扫描(Scan)
  • 压缩（Compact）和 分配（Allocate）
  • 分段扫描：例子，稀疏矩阵与向量相乘的并行化方式
  • 排序：奇偶排序，归并排序，排序网

图像处理的例子

  • RGB转灰度图
  • 图像模糊

CUDA流Streams

  • 流：一系列将在GPU上按顺序执行的操作
    • ​ 定义流：cudaStream_t s1;
      ​ 创建流：cudaStreamCreate(&s1);
      ​ 销毁流：cudaStreamDestory(s1);
  • 多GPU编程
    • ​ 两个问题：确定此时使用的是哪一个GPU，多个device如何与host同行

多GPU编程

  • 两个方面

    Peer-to-peer(P2P) memcopies

  • 使用另一个GPU的地址

    cudaDeviceEnablePeerAccess( peer_device, 0 )
    允许current GPU访问peer_device GPU
cudaDeviceCanAccessPeer( &accessible, dev_X, dev_Y)
    检查是否dev_X可以访问dev_Y的内存返回0/1(第一个参数)

  • Peer-to-peer memcopy

    cudaMemcpyPeerAsync(void* dst_addr, intdst_dev, void* src_addr, intsrc_dev, size_tnum_bytes, cudaStream_tstream)
两个设备之间拷贝字节
1）如果peer-access允许字节在最短的PCIe路径上传输
2）如果peer-access不允许CUDA驱动通过CPU memory传输

纹理内存

纹理内存的优势：
​ 1.它们是被缓存的,如果它们在texture fetch 中将提供更高的带宽
​ 2.它们不会像全局或常驻内存读取时受内存访问模式的约束
​ 3.寻址计算时的延迟更低,从而提高随机访问数据时的性能

​ 在一个操作中,包装的数据可以通过广播到不同的变量中5.8-bit和16-bit的整型输入数据可以被转换成在范围[0.0,1.0]或[-1.0,1.0]的浮点数

CPU与GPU协同

CPU先将数据复制到CPU内的一个区域（staging）,然后再复制到GPU（通过PCLe）
用于优化
​ cudaHostMalloc()
​ cudaHostRegister()

MPI

MPI安装
MPI教程

​
​

-------------本文结束 感谢您的阅读-------------
作者：GonewithGt
有问题请 留言 或者私信我的 微博。
满分是10分的话，这篇文章你给几分