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


0. 目录

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 GPUcudaMalloc((void**) &sum, sizeof(int) * THREAD_NUM);// calculate the squares's sumsquaresSum<<<1, THREAD_NUM, 0>>>(gpuData, sum, time);
    
  • 最后,在 CPU 内增加部分和求和的代码

      // print resultint 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

运行后结果如下所示:

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

GPU 编程入门到精通(三)之 第一个 GPU 程序相关推荐

  1. GPU 编程入门到精通(五)之 GPU 程序优化进阶

    版权声明:本文为博主原创文章,未经博主允许不得转载. 目录(?)[+] 博主由于工作当中的需要,开始学习 GPU 上面的编程,主要涉及到的是基于 GPU 的深度学习方面的知识,鉴于之前没有接触过 GP ...

  2. GPU 编程入门到精通(四)之 GPU 程序优化

    版权声明:本文为博主原创文章,未经博主允许不得转载. 目录(?)[+] 博主由于工作当中的需要,开始学习 GPU 上面的编程,主要涉及到的是基于 GPU 的深度学习方面的知识,鉴于之前没有接触过 GP ...

  3. GPU 编程入门到精通(二)之 运行第一个程序

    博主由于工作当中的需要,开始学习 GPU 上面的编程,主要涉及到的是基于 GPU 的深度学习方面的知识,鉴于之前没有接触过 GPU 编程,因此在这里特地学习一下 GPU 上面的编程.有志同道合的小伙伴 ...

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

    GPU 编程入门到精通(一)之 CUDA 环境安装 标签: cudagpunvidia GPU 编程入门到精通(一)之 CUDA 环境安装 标签: cudagpunvidia 2014-04-11 2 ...

  5. socket 编程入门教程(三)TCP原理:5、TCP的三次握手(three-way handshake)

    socket 编程入门教程(三)TCP原理:5.TCP的三次握手(three-way handshake) 前面3个小节介绍了socket机制对TCP协议三次握手的实现,需要强调的是,与协议独立于实现 ...

  6. Kali Linux 从入门到精通(三)-入侵系统定制

    Kali Linux 从入门到精通(三)-入侵系统定制 定制 网络配置 临时IP地址 dhclient eth0 ifconfig eth0 192.168.11/24 route add defau ...

  7. C4D致富经典入门到精通(三)

    C4D样条曲线创建 C4D基础界面的介绍与常用快捷键:  C4D致富经典入门到精通(一) C4D父子关系的理解与创建参数几何体与可编辑对象: C4D致富经典入门到精通(二) C4D样条曲线创建 :  ...

  8. webGL入门(1)创建第一个webGL程序

    webGL入门(1)创建第一个webGL程序 createScense.html文件 <!DOCTYPE html> <html><head><meta ch ...

  9. java编程入门到精通看什么书,详细说明

    Java面试高频题:Spring Boot+Sentinel+Nacos高并发已撸完Java成长笔记1. Java基础复盘2. Web编程初探3. SSM从入门到精通4. SpringBoot快速上手 ...

最新文章

  1. VC删除IE缓存、COOKIE及记录
  2. VMware中怎样克隆虚拟机
  3. 设计模式之_Strategy_03
  4. java双缓存机制_详解JVM类加载机制及类缓存问题的处理方法
  5. BATJ一线大公司需要什么样的前端
  6. 06.Qt菜单栏工具栏学习(一)
  7. 【入门】Spring-Boot项目配置Mysql数据库
  8. 大数据之-Hadoop3.x_MapReduce_Combiner案例---大数据之hadoop3.x工作笔记0119
  9. mysql 使用jdbctemplate_SpringBoot学习笔记-使用jdbcTemplate访问mysql
  10. LVS + Nginx +Tomcat 高可用架构
  11. [math][mathematica] archlinux 下 mathematica 的安装 (科学计算软件 mathematica/matlab/sagemath)...
  12. Day001 20210206
  13. oracle shutdown 很慢,oracle shutdown immediate等待时间很长之思考
  14. 摄像机标定:像素焦距与毫米焦距转换
  15. python123程序设计题说句心里话a_C程序设计基础(2019年春)-中国大学mooc-试题题目及答案...
  16. java编程之数字魔方(N阶数字魔方和数字旋转魔方)
  17. 关于一个非常nice的原生UI框架——layerUI
  18. 共享服务器模式(shared server)和专用服务器模式(dedicated server)
  19. 如何利用XMLSpy实现从多个XML实例生成架构
  20. 数电实验4:彩灯控制器设计

热门文章

  1. 哥大计算机工程申请入口和登陆入口
  2. 怎么在坐标上面找outliner的异常检测方法,非监督学习
  3. 熊掌号指数2.0常见问题汇总
  4. WPF无边框拖动、全屏、缩放
  5. CF A. DZY Loves Hash
  6. CentOS关机大法之shutdown应用实例
  7. Java面向对象抽象类案例分析
  8. 阿里云总裁胡晓明:保护客户数据隐私是阿里云第一原则
  9. JS获取浏览器高度 并赋值给类
  10. HDU 5531 Rebuild