CUDA进阶第三篇:CUDA计时方式
2016-02-04 00:02
543 查看
写在前面
内容分为两部分,第一部分为翻译《Professional CUDA C Programming》 Section 2. CUDA Programming Model中的TIMING YOUR KERNEL;第二部分为自己的经验。经验不足,欢迎各位大大补充。写CUDA,追求的就是加速比,想要得到准确的时间,计时函数就是必不可少
计时通常分为两种情况,(1)直接得到接口函数的时间,一般用于得到加速比;(2)获得接口函数内核函数、内存拷贝函数等所耗时间,一般用于优化代码时。
情况(1)方法有两种,CPU计时函数和GPU计时函数。
情况(2)有三种工具nsight,nvvp,nvprof
本博客会详细介绍情况(1)的两种方法;情况(2),nsight不会用,简单介绍一下nvvp和nvprof的用法。
CPU计时函数
在利用CPU计时函数时,要考虑的一个问题是:核函数的执行是异步执行的,所以必须加上核函数同步函数,才能得到准确的时间。示例代码如下:
double cpuSecond() { struct timeval tp; gettimeofday(&tp,NULL); return ((double)tp.tv_sec + (double)tp.tv_usec*1.e-6); } double iStart = cpuSecond(); function(argument list); cudaDeviceSynchronize(); // 同步函数 double iElaps = cpuSecond() - iStart;
GPU计时函数
GPU计时函数就不需要考虑同步问题,直接用计时事件函数就可以了,示例代码如下:cudaEvent_t start, stop; float elapsedTime = 0.0; cudaEventCreate(&start); cudaEventCreate(&stop); cudaEventRecord(start, 0); function(argument list);; cudaEventRecord(stop, 0); cudaEventSynchronize(stop); cudaEventElapsedTime(&elapsedTime, start, stop); cout << elapsedTime << endl; cudaEventDestroy(start); cudaEventDestroy(stop);
如何获得精确的计时
正常情况下,第一次执行核函数的时间会比第二次慢一些。这是因为GPU在第一次计算时需要warmup。所以想要第一次核函数的执行时间是不精确的。获得精确的计时我总结为三种,如下循环执行一百次所需要计时的部分,求平均值,将第一次的误差缩小100倍。这种方法的优点是简单粗暴。但缺点也很明显:(1)程序的执行时间大大增长,特别是比较大的程序(2)要考虑内存溢出问题,C++的内存需要程序猿自己手动管理。写出执行一次不出内存溢出问题的程序很容易,但是写出循环执行一百次而不出内存溢出问题的代码就有一定难度了(3)计时不是特别准确,虽然误差已经被缩小了100倍。
在计时之前先执行一个warmup函数,warmup函数随便写,比如我从cuda sample里的vectoradd作为warmup函数。这种方法的优点是程序执行时间缩短;缺点是需要在程序中添加一个函数,而且因为GPU乱序并行的执行方式,核函数的两次执行时间并不能完全保持一样。所以推荐使用方法3.
先执行warmup函数,在循环10遍计时部分。
nvvp和nvprof的用法
nvprof是自cuda5.0开始存在的一个命令行Profiler,你可以只用nvprof来你代码的一些执行细节。简单用法如下:$ nvprof ./sumArraysOnGPU-timer
你就可以得到如下:
./sumArraysOnGPU-timer Starting... Using Device 0: Tesla M2070 ==17770== NVPROF is profiling process 17770, command: ./sumArraysOnGPU-timer Vector size 16777216 sumArraysOnGPU <<<16384, 1024>>> Time elapsed 0.003266 sec Arrays match. ......
关于nvprof的更多参数信息,可以使用帮助命令:
$ nvprof --help
The NVIDIA Visual Profiler(nvvp)是一款图形化界面的Profiler,也是我一直在用的Profiler。
简单图文教程见链接 ,我随后会补充一些,先睡了。
写在后面
OpenCUDA:CUDA图像算法开源项目,算法内都有详细的注释,大家一起学习。私人接各种CUDA相关外包(调试、优化、开发图像算法等),有意向请联系,加好友时请注明。
相关文章推荐
- Spring Security安全框架入门篇
- Symfony2在Nginx下的配置方法图文教程
- Symfony2安装第三方Bundles实例详解
- 伸展树的原理及实现源代码(有图文详解和C++实现代码)
- 详解Python编程中基本的数学计算使用
- Python的math模块中的常用数学函数整理
- Python中的字符串类型基本知识学习教程
- Python中常用操作字符串的函数与方法总结
- Node.js重新刷新session过期时间的方法
- Nodejs中session的简单使用及通过session实现身份验证的方法
- 关于JavaScript作用域你想知道的一切
- 分享我对JS插件开发的一些感想和心得
- 你一定会收藏的Nodejs代码片段
- JS获取时间的相关函数及时间戳与时间日期之间的转换
- AngularJs中route的使用方法和配置
- 你需要知道的12个Git高级命令
- 肯德基怒告十个微信公号获赔60万:拿什么拯救营销号与你的朋友圈
- 马云和他的高管们,在内部年会都聊了什么?
- 应用市场正式走入黄昏
- 20元入群费背后的社群经济究竟是什么?