基于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，具体原因我还不明白，有待探究。