GPU 并行编程技术介绍 胡光龙 huguanglong@tju

1. GPU并行编程技术介绍 胡光龙 huguanglong@tju.edu.cn

3. GPU简介 GPU，Graphics Processing Unit 的简写，是现代显卡中非常重要的一部分，其地位与CPU在主板上的地位一致，设计的初衷是用于在显示器上渲染计算机图形，加速图形处理速度（现在仍然主要用于这个目的）。 GPU是一个高度并行化的多线程，多核心处理器。

4. GPU简介 GPU极大提升了计算机图形处理的速度，增强了图形的质量，并促进了与计算机图形相关的其他领域的快速发展。 与中央处理器（Central Processing Unit，CPU）的串行设计模式不同，GPU专门为图形处理设计，具有天然的并行特性。

5. GPU简介 GPU计算特点吸引着许多研究人员来探索如何使用GPU在图形处理之外的领域中进行计算。 早期的基于GPU的通用计算非常复杂。标准图形接口（如OpenGL，DirectX），是与GPU交互的唯一方式。要在GPU上执行计算，必然受限于图形API的编程模型。

6. CUDA简介 2006年11月，NVIDIA公布了业界第一个DirectX 10 GPU，即GeForce 8800 GTX。这是第一个基于CUDA架构的GPU。 CUDA，可以理解为NVIDIA公司构建其GPU的模式，极大地方便了我们在GPU中完成通用计算任务。CUDA的出现使用户无需了解图形编程结构也不需要将通用计算问题伪装成图形计算问题

7. CUDA简介 NVIDIA采取了工业标准的C语言，并增加了一部分关键字来支持CUDA架构的特殊功能，这便是CUDA C 语言。它专门用来在GPU上编写通用计算。 CUDA在医学图像处理，流体力学的计算以及环境科学等领域得到广泛的应用，基于GPU的代码甚至比之前的代码在性能上提升了多个数量级

8. CUDA编程模型 CUDA的基本思想是支持大量的线程级并行，并在硬件中动态的调度和执行这些线程。 GPU只有在计算高度数据并行任务时才发挥作用。 异构思想： Host（CPU）+Device（GPU）

9. CUDA编程模型 在 CUDA 的架构下，一个程序分为两个部份：host 端和 device 端。Host 端是指在 CPU 上执行的部份，而 device 端则是在显示芯片上执行的部份。通常 host 端程序会将数据准备好后，复制到显卡的内存中，再由显示芯片执行 device 端程序，完成后再由 host 端程序将结果从显卡的内存中取回。

10. CUDA编程模型 注意： 由于 CPU 存取显卡内存时只能透过 PCI Express 接口，速度较慢（PCI Express x16 的理论带宽是双向各 4GB/s），因此不能太频繁进行这类动作，以免降低效率。

11. CUDA编程模型 CUDA C的关键字： 1 函数类型限定符，用来确定函数是在CPU还是在GPU上执行，以及这个函数是从CPU调用还是从GPU调用。 __global__ （host端调用，device端执行） __device__（device端调用，device端执行）

12. CUDA编程模型 CUDA C的关键字： 2执行配置运算符<<< x , y>>>，用来传递内核函数(__global__修饰的函数)的执行参数。 第一个参数x代表一个线程格（grid）中线程块（block）的个数（一维或者多维，每一维度最大65535个）。 第二个参数y代表每个线程块（block）中线程（thread）的个数（一维或者多维，最多512个）。

13. CUDA编程模型 CUDA C的关键字： 3 几个内建变量： gridDim.x gridDim.y代表线程格中不同维度的线程块（block）的个数。 blockDim.x blockDim.y代表线程块中不同维度的线程（thread）的个数。

14. CUDA编程模型 CUDA C的关键字： 3 几个内置变量： blockIdx.x blockIdx.y表示当前线程块所在线程格中不同方向上的索引。 threadIdx.x threadIdx.y表示当前线程所在线程块中不同方向上的索引。

15. CUDA编程模型 CUDA C的关键字： 4 内存操作函数： cudaMalloc（）: 类似于malloc（），作用是告诉CUDA运行时在设备上分配内存。对应的cudaFree（）来释放内存。 CudaMemcpy（）：类似于memcpy（），在设备和主机之间存取内存。

16. 简单地CUDA程序实例 我们通过一个最简单的矢量求和运算的例子来加深对CUDA并行计算的理解。 1 基于CPU的矢量求和 void add(int *a, int *b, int *c) { for(int tid=0; tid<N; tid++) { C[tid]=a[tid]+b[tid]; } // 要执行N次 }

17. 简单地CUDA程序实例 2 基于GPU的矢量求和 我们可以再GPU中实现相同的加法运算，这需要将add（）编写为一个设备函数。在给出设备代码之前，首先给出main（）。 #define N 256 main（） { int a[N], b[N] ,c[N]; //CPU上的矢量 int * dev_a , *dev_b , *dev_c ; //GPU上的设备指针

18. 简单地CUDA程序实例 //在GPU上分配内存 cudaMalloc((void**)&dev_a, N* sizeof (int)) cudaMalloc((void**)&dev_a, N* sizeof (int)) cudaMalloc((void**)&dev_a, N* sizeof (int)) //为数组‘a’和‘b’赋值 for(int i=0; i<N ;i++) {a[i]= -i; b[i]=i*i; }

19. 简单地CUDA程序实例 //将数组‘a’和‘b’复制到GPU cudaMemcpy ( dev_a, a, N*sizeof(int), cudaMemcpyHostToDevice ); cudaMemcpy ( dev_b, b, N*sizeof(int), cudaMemcpyHostToDevice ); //在GPU中使用设备函数来计算矢量和 //总共启用N个线程块，每个线程块包含一个线程 add<<<N,1>>>(dev_a, dev_b, dev_c) ;

20. 简单地CUDA程序实例 //将结果从GPU复制回CPU cudaMemcpy ( c, dev_c, N*sizeof(int), cudaMemcpyDeviceToHost ); //释放GPU上分配的内存 cudaFree（dev_a）； cudaFree（dev_b）； cudaFree（dev_c）； }

21. 简单地CUDA程序实例 接下来我们来看 add（）函数。 __global__ void add(int *a, int *b, int *c) { int tid = blockIdx.x; c[tid] = a[tid] + b[tid]; } 程序运行时将创建add（）核函数的N个副本。 blockIdx是一个内置变量，变量中包含的值就是当前执行设备代码的线程块的索引。 所以仅需要一步即可完成矢量的相加。

22. Thanks