博主由于工作当中的需要,开始学习 GPU 上面的编程,主要涉及到的是基于 GPU 的深度学习方面的知识,鉴于之前没有接触过 GPU 编程,因此在这里特地学习一下 GPU 上面的编程。有志同道合的小伙伴,欢迎一起交流和学习,我的邮箱: caijinping220@gmail.com 。使用的是自己的老古董笔记本上面的 Geforce 103m 显卡,虽然显卡相对于现在主流的系列已经非常的弱,但是对于学习来说,还是可以用的。本系列博文也遵从由简单到复杂,记录自己学习的过程。
<a name="t0" style="box-sizing: border-box; background: transparent; color: rgb(79, 161, 219); text-decoration: none; margin: 0px; padding: 0px; font-weight: 400; outline: 0px;"></a>0. 目录
- GPU 编程入门到精通(一)之 CUDA 环境安装
- GPU 编程入门到精通(二)之 运行第一个程序
- GPU 编程入门到精通(三)之 第一个 GPU 程序
- GPU 编程入门到精通(四)之 GPU 程序优化
- GPU 编程入门到精通(五)之 GPU 程序优化进阶
<a name="t1" style="box-sizing: border-box; background: transparent; color: rgb(79, 161, 219); text-decoration: none; margin: 0px; padding: 0px; font-weight: 400; outline: 0px;"></a>1. 数组平方和并行化
GPU 编程入门到精通(三)之 第一个 GPU 程序 中讲到了如何利用 CUDA5.5 在 GPU 中运行一个程序。通过程序的运行,我们看到了 GPU 确实可以作为一个运算器,但是,我们在前面的例子中并没有正真的发挥 GPU 并行处理程序的能力,也就是说之前的例子只利用了 GPU 的一个线程,没有发挥程序的并行性。
先来说说 CUDA5.5 中 GPU 的架构。它是由 grid 组成,每个 grid 又可以由 block 组成,而每个 block 又可以细分为 thread。所以,线程是我们处理的最小的单元了。
接下来的例子通过修改前一个例子,把数组分割成若干个组(每个组由一个线程实现),每个组计算出一个和,然后在 CPU 中将分组的这几个和加在一起,得到最终的结果。这种思想叫做归约 。其实和分治思想差不多,就是先将大规模问题分解为小规模的问题,最后这些小规模问题整合得到最终解。
由于我的 GPU 支持的块内最大的线程数是 512 个,即 cudaGetDeviceProperties
中的 maxThreadsPerBlock
属性。如何获取这个属性,请参看 GPU 编程入门到精通(二)之 运行第一个程序这一章节。 我们使用 512 个线程来实现并行加速。
好了,接下来就是写程序的时候了。
<a name="t2" style="box-sizing: border-box; background: transparent; color: rgb(79, 161, 219); text-decoration: none; margin: 0px; padding: 0px; font-weight: 400; outline: 0px;"></a>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);
<a name="t3" style="box-sizing: border-box; background: transparent; color: rgb(79, 161, 219); text-decoration: none; margin: 0px; padding: 0px; font-weight: 400; outline: 0px;"></a>1.2. 编译运行
编译后,运行结果如下所示:
thread512<a name="t4" style="box-sizing: border-box; background: transparent; color: rgb(79, 161, 219); text-decoration: none; margin: 0px; padding: 0px; font-weight: 400; outline: 0px;"></a>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 线程的一般流程。下一部分,将进一步分析程序可优化的一些细节。
欢迎大家和我一起讨论和学习 GPU 编程。
caijinping220@gmail.com
http://blog.csdn.net/xsc_c</center>
网友评论