CUDA程序的调试总结【不定时更新】
2013-07-21 15:03
225 查看
1 )CUDA的程序,经常犯,但是很难发现的一个错误就是同步问题。
描述下实例
看看上面的代码,好像没问题。
其实当N < BS的时候上面的代码是没有问题的。但是当N大于BS的时候,每个线程会至少循环两次,这样问题就来了。
假设第一个warp的线程已经执行完了out的赋值,但是第二组warp还在计算那个tp,tp依赖于在shared memory中的数据,如果第一个warp开始执行sda那一句话的话,第二个warp就会得到错误的数据。
虽然你有一个同步了!
解决方法很简单,就是在out输出之后加一个同步操作,当然你加到sda前面也是可以的。
补充一点,这个问题如何发现呢?只要比对下两次执行的结果,看看是否一致,如果结果不一致,那么就很有可能犯了同步的错误。
2)CUDA程序第二经常犯的错误就是线程访问显存越界,或者共享存储器访问越界
如何发现这个问题呢。这种情况下,一般你的kernel不会启动成功。如果不会启动成功,也不一定能就是越界问题,如果你的kernel中使用了过多的共享存储器,也不会启动成功的。
遇到启动不成功的时候,你首先要计算下shared memory是否超出了硬件范围,至于硬件的shared memory有多少,你还需要查一下,我正能说,这个跟GPU的核心有关,你只要根据你的设备计算能力取查找就行了。
如果是因为越界,可以将kernel函数一点点注释起来,查看输出结果。步步蚕食。一定会找到越界的位置。找到后自己解决就行了。
。。。未完待续。。。
描述下实例
for (k = 0; k < N; k+=BS) { sda[tx] = gda[tx+index]; __syncthreads(); for (j = 0; j < BS; j++) { tp += sda[j] } out[index+tx] = tp; }
看看上面的代码,好像没问题。
其实当N < BS的时候上面的代码是没有问题的。但是当N大于BS的时候,每个线程会至少循环两次,这样问题就来了。
假设第一个warp的线程已经执行完了out的赋值,但是第二组warp还在计算那个tp,tp依赖于在shared memory中的数据,如果第一个warp开始执行sda那一句话的话,第二个warp就会得到错误的数据。
虽然你有一个同步了!
解决方法很简单,就是在out输出之后加一个同步操作,当然你加到sda前面也是可以的。
补充一点,这个问题如何发现呢?只要比对下两次执行的结果,看看是否一致,如果结果不一致,那么就很有可能犯了同步的错误。
2)CUDA程序第二经常犯的错误就是线程访问显存越界,或者共享存储器访问越界
如何发现这个问题呢。这种情况下,一般你的kernel不会启动成功。如果不会启动成功,也不一定能就是越界问题,如果你的kernel中使用了过多的共享存储器,也不会启动成功的。
遇到启动不成功的时候,你首先要计算下shared memory是否超出了硬件范围,至于硬件的shared memory有多少,你还需要查一下,我正能说,这个跟GPU的核心有关,你只要根据你的设备计算能力取查找就行了。
如果是因为越界,可以将kernel函数一点点注释起来,查看输出结果。步步蚕食。一定会找到越界的位置。找到后自己解决就行了。
。。。未完待续。。。
相关文章推荐
- CUDA程序的调试总结【不定时更新】
- xcode下真机调试cocos2d-x程序资源不更新的解决办法
- 学习struts2建bbs总结六:hibernate分页查询的问题--分页后程序定时无响应
- 利用Console来调试JS程序、Console用法总结
- 如何写出正确CUDA程序(持续更新中)
- 利用Console来调试JS程序、Console用法总结(1)
- VS2010+VMWare 远程调试exe程序的配置总结
- 有趣linux小程序总结(持续更新)
- 个人平时总结(不定时更新)
- visual profiler 调试cuda并行程序:根据行号定位出错行
- 使用NetBeans + Xdebug调试PHP程序[总结]
- NesC学习经验总结:第三篇 如何调试NesC程序
- 最短路模板总结(不定时更新)
- 微信小程序bug总结, 不定期更新
- linux常用命令总结 ---不定时更新
- 在这里记录自己学习中遇到的零碎的容易忽略的知识点,不定时总结更新
- gdb调试多线程程序总结
- [分享]xcode免证书真机调试iphone(ipad)程序的具体方法(总结贴))
- DebugView调试C#程序 学习总结
- php程序调试方法总结