基于cuda的mmp的bug调试
程序初稿出來后，有這樣幾個bug：
1、內(nèi)存訪存超出范圍
2、每次循環(huán)后，用于周轉(zhuǎn)的寄存器tmp沒有重置為零
3、將數(shù)據(jù)拷貝到sharedmemory后的第一次迭代數(shù)據(jù)不對。數(shù)據(jù)結(jié)果每次運行不一樣，但是一個有限集。
針對第一個bug，是指針的使用上混亂了。mmp這個kernel中用到的指針有點多，且相互交織，邏輯稍不清晰就容易出岔子。一般指針的使用就是指針不變后面加上索引，或者指針隨著循環(huán)變化而索引不變，而在cuda里，還有一種就是每個線程指針值可能相同也可能不同。要對這幾種情況認(rèn)真分析，如果對自己的邏輯能力沒有把握，可以只讓索引變，指針不變。而我之所以指針超出范圍，是因為計算和傳輸?shù)捻樞虿灰粯?#xff0c;而我卻用了同一指針，使得我的訪存很容易就出錯。所以我就讓計算和傳輸?shù)闹羔樂侄沃?/strong>，各用各的，邏輯就不易出錯。ps內(nèi)存訪存的問題可以用memorychecker來檢測。
第二個bug其實是比較容易找出來的，在寫程序的時候就會注意到這個問題。
對第三個bug，最初我以為是因為存放結(jié)果的指針沒有初始化，使得后面結(jié)果在迭代時出現(xiàn)問題，所以我就把第一次循環(huán)分離出來賦給結(jié)果指針，之后再循環(huán)迭代（+=）。但是這樣做并沒有解決問題。一方面，我在host代碼中初始化了結(jié)果指針并拷貝到了device端；另一方面，cuda對于沒初始化的指針會自動初始化為全0或1，而不是亂碼。所以如果是因為沒有初始化那么每次結(jié)果應(yīng)該是一樣的。而這種數(shù)據(jù)結(jié)果每次運行不一樣，但是一個有限集的情況，99%是因為沒有同步。programming guide里指出了如果在需要同步的地方?jīng)]有同步，就會導(dǎo)致結(jié)果出現(xiàn)意外，但一定是一個有限集的一個解。詳情參見guide。一般來說，在數(shù)據(jù)寫入后，讀取前，需要一個同步保證寫讀依賴。
另外目前寫的mmp只占了峰值性能的20%，是cublas性能的1/3。主要問題有兩個，1是對矩陣b的訪問地址不連續(xù)，不滿足coalesce；2是對矩陣c（global）的訪存次數(shù)還是有點多，只縮小了8倍。所以目前來說可能要考慮第二個方案就是一次性算出c的一個元素，這樣可以減少global的訪存次數(shù)，解決這個主要矛盾。
以后再更一篇關(guān)于nsight的調(diào)試功能技巧的。
?
------------------------------------割-------------------------------------------------------------
另外想起來之前調(diào)試的時候一個現(xiàn)象。當(dāng)我在把數(shù)據(jù)拷貝到shared memory后少了__syncthreads()時，在debug模式下沒有問題，在release模式下卻會出現(xiàn)上文提到的第三個bug，具體原因我還不明白，有待探究。
轉(zhuǎn)載于:https://www.cnblogs.com/bitghost/p/5754719.html
總結(jié)
以上是生活随笔為你收集整理的基于cuda的mmp的bug调试的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。
- 上一篇： C#语句——循环语句（for循环与for
- 下一篇： list array解析(总算清楚一点了