程序初稿出來后，有這樣幾個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，具體原因我還不明白，有待探究。

轉載于:https://www.cnblogs.com/bitghost/p/5754719.html

總結

以上是生活随笔為你收集整理的基于cuda的mmp的bug调试的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。