美文网首页
GPU 编程入门到精通(一)之 CUDA 环境安装

GPU 编程入门到精通(一)之 CUDA 环境安装

作者: 04282aba96e3 | 来源:发表于2018-01-10 09:53 被阅读75次

    博主由于工作当中的需要,开始学习 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. 目录

    <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>

    相关文章

      网友评论

          本文标题:GPU 编程入门到精通(一)之 CUDA 环境安装

          本文链接:https://www.haomeiwen.com/subject/ucwnnxtx.html