關於CUDA的bank conflict - 顯卡

※ 引述《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大的回文

需要時間消化一下

--