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

[rick209 ()]

※ 引述《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)
推 dulcet:num1 和num2是在register, 不會conflict
→ dulcet:你一定在num1, num2前加了__shared__了 哪掉就OK
※ 編輯: rick209         來自: 140.113.152.194      (06/18 17:44)
→ rick209:但是num1 num2必須配置於shared memory中
推 dulcet:@@ vector sum怎麼可能把register用完
→ dulcet:你可以把之前用過的register重複利用