CUDA程序的调试总结【不定时更新】

1 )CUDA的程序,经常犯,但是很难发现的一个错误就是同步问题。

描述下实例

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函数一点点注释起来,查看输出结果。步步蚕食。一定会找到越界的位置。找到后自己解决就行了。

。。。未完待续。。。

上一篇:小测试 php代理,nginx代理,直接访问对比


下一篇:启动两个tomcat服务,以及使用nginx代理实现访问