GPU 编程入门到精通（四）之 GPU 程序优化

博主由于工作当中的需要，开始学习 GPU 上面的编程，主要涉及到的是基于 GPU 的深度学习方面的知识，鉴于之前没有接触过 GPU 编程，因此在这里特地学习一下 GPU 上面的编程。有志同道合的小伙伴，欢迎一起交流和学习，我的邮箱: caijinping220@gmail.com 。使用的是自己的老古董笔记本上面的 Geforce 103m 显卡，虽然显卡相对于现在主流的系列已经非常的弱，但是对于学习来说，还是可以用的。本系列博文也遵从由简单到复杂，记录自己学习的过程。

---

0. 目录

• GPU 编程入门到精通（一）之 CUDA 环境安装
• GPU 编程入门到精通（二）之 运行第一个程序
• GPU 编程入门到精通（三）之 第一个 GPU 程序
• GPU 编程入门到精通（四）之 GPU 程序优化
• GPU 编程入门到精通（五）之 GPU 程序优化进阶

1. 数组平方和并行化

GPU 编程入门到精通（三）之 第一个 GPU 程序 中讲到了如何利用 CUDA5.5 在 GPU 中运行一个程序。通过程序的运行，我们看到了 GPU 确实可以作为一个运算器，但是，我们在前面的例子中并没有正真的发挥 GPU 并行处理程序的能力，也就是说之前的例子只利用了 GPU 的一个线程，没有发挥程序的并行性。

先来说说 CUDA5.5 中 GPU 的架构。它是由 grid 组成，每个 grid 又可以由 block 组成，而每个 block 又可以细分为 thread。所以，线程是我们处理的最小的单元了。

接下来的例子通过修改前一个例子，把数组分割成若干个组（每个组由一个线程实现），每个组计算出一个和，然后在 CPU 中将分组的这几个和加在一起，得到最终的结果。这种思想叫做归约 。其实和分治思想差不多，就是先将大规模问题分解为小规模的问题，最后这些小规模问题整合得到最终解。

由于我的 GPU 支持的块内最大的线程数是 512 个，即 cudaGetDeviceProperties 中的 maxThreadsPerBlock 属性。如何获取这个属性，请参看 GPU 编程入门到精通（二）之 运行第一个程序 这一章节。 我们使用 512 个线程来实现并行加速。

好了，接下来就是写程序的时候了。

1.1. 修改代码

• 首先，在程序头部增加一个关于线程数量的宏定义：

    // ======== define area ========
    #define DATA_SIZE 1048576 // 1M
    #define THREAD_NUM 512 // thread num
  

  其中，DATA_SIZE 表示处理数据个数， THREAD_NUM 表示我们将要使用 512 个线程。

• 其次，修改 GPU 部分的内核函数

    const int size = DATA_SIZE / THREAD_NUM;
    const int tid = threadIdx.x;
    int tmp_sum = 0;
  
    for (int i = tid * size; i < (tid + 1) * size; i++) {
    tmp_sum += data[i] * data[i];
    }
    sum[tid] = tmp_sum;
    }
  

  此内核程序的目的是把输入的数据分摊到 512 个线程上去计算部分和，并且 512 个部分和存放到 sum 数组中，最后在 CPU 中对 512 个部分和求和得到最终结果。

  此处对数据的遍历方式请注意一下，我们是根据顺序给每一个线程的，也就是如下表格所示：

  线程编号	数据下标
  0	0 ~ 2047
  … …	… …
  511	1046528 ~ 1048575

• 然后，修改主函数部分
  主函数部分，只需要把 sum 改成数组就可以，并且设置一下调用 GPU 内核函数的方式。

    // malloc space for datas in GPU
    cudaMalloc((void**) &sum, sizeof(int) * THREAD_NUM);
  
    // calculate the squares's sum
    squaresSum<<<1, THREAD_NUM, 0>>>(gpuData, sum, time);
  

• 最后，在 CPU 内增加部分和求和的代码

    // print result
    int tmp_result = 0;
    for (int i = 0; i < THREAD_NUM; ++i) {
    tmp_result += result[i];
    }
    printf("(GPU) sum:%d time:%ld\n", tmp_result, time_used);
  

1.2. 编译运行

编译后，运行结果如下所示：

2. 性能分析

经过修改以后的程序，比之前的快了将近 36 倍（可以参考博文 GPU 编程入门到精通（三）之 第一个 GPU 程序 进行比较），可见并行化处理还是存在优势的。 不过仔细想一下，我们使用了 512 个线程， 可是性能怎么才提升了 36 倍，不应该是 512 倍吗？？？

这里就涉及到内存的存取模式了，显卡上面的内存是 DRAM，是效率最高的存取方式，它是一种连续的存取方式。 前面我们的程序确实的连续读取的呀，都挨个读取了，怎么还是没有达到预期的效果呢？？？

这里还需要考虑 thread 的执行方式，GPU 编程入门到精通（三）之 第一个 GPU 程序 中说到，当一个 thread 在等待内存数据的时候， GPU 就会切换到下一个 thread。所以，实际执行的顺序类似于 thread0 —> thread1 —> … … —> thread511。

这就导致了同一个 thread 在读取内存是连续的， 但是对于整体而言，执行的过程中读取就不是连续的了（这里自己仔细想想，就明白了）。所以，正确的做法如下表格所示：

线程编号	数据下标
0	0 ~ 512
… …	… …
511	511 ~ 1023

根据这个原理，修改内核函数如下：

for (int i = tid; i < DATA_SIZE; i += THREAD_NUM) {
tmp_sum += data[i] * data[i];
}

编译运行后结果如下所示：

修改后程序，比之前的又快了 13 倍左右，可见，对内存的读取方式对于性能的影响很大。
至此，并行化后的程序较未并行化之前的程序，速度上快了 493 倍左右，可见，基本上发挥了 512 个线程的优势。

让我们再来分析一下性能：

此 GPU 消耗的时钟周期： 1595788 cycles
GeForce G 103M 的 clockRate： 1.6 GHz
所以可以计算出 GPU 上运行时间是： 时钟周期 / clockRate = 997.3675 us
1 M 个 int 型数据有 4M Byte 的数据量，实际使用的 GPU 内存带宽是：数据量 / 运行时间 = 4.01 GB/s

再来看看我的 GPU GeForce 103m 的内存带宽：运行 SDK 目录下面 /samples/1_Utilities/bandwidthTest

运行后结果如下所示：

通过与系统参数的对比，可以知道，基本上达到了系统的极限性能。

---

这一篇博文介绍了如何通过利用线程达到程序的并行计算，并且通过优化内存读取方式，实现对程序的优化。通过这个程序，可以学会使用 CUDA 线程的一般流程。下一部分，将进一步分析程序可优化的一些细节。
欢迎大家和我一起讨论和学习 GPU 编程。
caijinping220@gmail.com
http://blog.csdn.net/xsc_c