Re: [請益] 關於CUDA的bank conflict

看板VideoCard (顯卡板)作者rick209時間16年前 (2009/06/18 17:39)推噓2(2推 0噓 3→)

留言5則, 2人參與討論串2/2 (看更多)

※ 引述《BeelZeBub (天使獵人)》之銘言： : ※ 引述《rick209 ()》之銘言： : : 近來閱讀了版上a大關於cuda的文章 : : 因此想把自己的程式改寫成可利用GPU執行 : : 但bank conflict卻始終困擾著我 : : 舉例來說 : : for(int k=0; k<num; k++){ : : num1=data1[k] : : num2=data2[k] : : sum[k]=num1+num2; : : } : : 若把num1 num2的記憶體配置在shared memory時 : : 會因為不同執行緒存取到同一塊記憶體產生bank conflict的問題 : : 但因為計算複雜所需記憶體大的關係也無法配置到暫存器上 : : 想請教cuda有類似openMp中 for private()的指令嗎 : : 還是就只能完全利用陣列運算如把num1改變成陣列num1[k]等 : 來討論一下 : CUDA的threading的確會有memory conflict的問提存在 : 假設今天開四個threading : [th0 th0 th0 th1 th1 th1 th2 th2 th2 th3 th3 th3] : 根據memory allocation : 每個thread的第一個job(1st th0, 1st th1.....)都會同時卡在第一塊block裡面 : 這樣會造成bottleneck 而使的performance上不來 : 以上如果我沒理解錯應該就是你的問題吧 : 但是今天不要像上面規劃的那樣規劃成 : [th0 th1 th2 th3 th0 th1 th2 th3 th0 th1 th2 th3] : 如此一來每個thread的第一個job就可以錯開同一block的存取 : 概念是這樣如果要寫成程式的話 : 就要改指向頭的header (講index比較好) : 詳細怎樣作要想一想不過概念應該是這樣沒錯 : 歡迎版上大大討論... @@! 抱歉寫得有點錯誤自己在補充一下 int k=threadIdx.x num1=data1[k] num2=data2[k] sum[k]=num1+num2; 因為我的程式中需要大量的迴圈因此最簡單的平行方式就是每個thread負責一個運算 k=1 thread 1負責 k=2 thread 2負責依此類推但是同一個block中共用共享記憶體因此有可能 thread1在存取num1時 thread2也剛好存取num1 如此會造成最終運算結果錯誤由於變數眾多因此需將變數配置於shared memory中感謝B大跟D大的回文需要時間消化一下 -- ※ 發信站: 批踢踢實業坊(ptt.cc) ◆ From: 140.113.152.194 ※ 編輯: rick209 來自: 140.113.152.194 (06/18 17:40)