GPU 编程入门到精通（三）之 第一个 GPU 程序

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. 编译运行

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

[图片上传失败...(image-cd7585-1515573581453)]

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];
}

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

thread512_plus

修改后程序，比之前的又快了 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

运行后结果如下所示：

bandwidth

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

---

<center style="box-sizing: border-box; color: rgb(69, 69, 69); font-family: "PingFang SC", "Microsoft YaHei", SimHei, Arial, SimSun; font-size: 16px; font-style: normal; font-variant-ligatures: normal; font-variant-caps: normal; font-weight: 400; letter-spacing: normal; orphans: 2; text-indent: 0px; text-transform: none; white-space: normal; widows: 2; word-spacing: 0px; -webkit-text-stroke-width: 0px; background-color: rgb(255, 255, 255); text-decoration-style: initial; text-decoration-color: initial;">
这一篇博文介绍了如何通过利用线程达到程序的并行计算，并且通过优化内存读取方式，实现对程序的优化。通过这个程序，可以学会使用 CUDA 线程的一般流程。下一部分，将进一步分析程序可优化的一些细节。</center>