基于cuda的mmp的bug调试
程序初稿出来后,有这样几个bug:
1、内存访存超出范围
2、每次循环后,用于周转的寄存器tmp没有重置为零
3、将数据拷贝到sharedmemory后的第一次迭代数据不对。数据结果每次运行不一样,但是一个有限集。
针对第一个bug,是指针的使用上混乱了。mmp这个kernel中用到的指针有点多,且相互交织,逻辑稍不清晰就容易出岔子。一般指针的使用就是指针不变后面加上索引,或者指针随着循环变化而索引不变,而在cuda里,还有一种就是每个线程指针值可能相同也可能不同。要对这几种情况认真分析,如果对自己的逻辑能力没有把握,可以只让索引变,指针不变。而我之所以指针超出范围,是因为计算和传输的顺序不一样,而我却用了同一指针,使得我的访存很容易就出错。所以我就让计算和传输的指针分而治之,各用各的,逻辑就不易出错。ps内存访存的问题可以用memorychecker来检测。
第二个bug其实是比较容易找出来的,在写程序的时候就会注意到这个问题。
对第三个bug,最初我以为是因为存放结果的指针没有初始化,使得后面结果在迭代时出现问题,所以我就把第一次循环分离出来赋给结果指针,之后再循环迭代(+=)。但是这样做并没有解决问题。一方面,我在host代码中初始化了结果指针并拷贝到了device端;另一方面,cuda对于没初始化的指针会自动初始化为全0或1,而不是乱码。所以如果是因为没有初始化那么每次结果应该是一样的。而这种数据结果每次运行不一样,但是一个有限集的情况,99%是因为没有同步。programming guide里指出了如果在需要同步的地方没有同步,就会导致结果出现意外,但一定是一个有限集的一个解。详情参见guide。一般来说,在数据写入后,读取前,需要一个同步保证写读依赖。
另外目前写的mmp只占了峰值性能的20%,是cublas性能的1/3。主要问题有两个,1是对矩阵b的访问地址不连续,不满足coalesce;2是对矩阵c(global)的访存次数还是有点多,只缩小了8倍。所以目前来说可能要考虑第二个方案就是一次性算出c的一个元素,这样可以减少global的访存次数,解决这个主要矛盾。
以后再更一篇关于nsight的调试功能技巧的。
------------------------------------割-------------------------------------------------------------
另外想起来之前调试的时候一个现象。当我在把数据拷贝到shared memory后少了__syncthreads()时,在debug模式下没有问题,在release模式下却会出现上文提到的第三个bug,具体原因我还不明白,有待探究。