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

。。。未完待续。。。