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