關於CUDA的bank conflict - 顯卡

Sandy avatar
By Sandy
at 2009-06-18T17:39

Table of Contents

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

需要時間消化一下

--
Tags: 顯卡

All Comments

Odelette avatar
By Odelette
at 2009-06-20T10:28
num1 和num2是在register, 不會conflict
Dorothy avatar
By Dorothy
at 2009-06-25T01:35
你一定在num1, num2前加了__shared__了 哪掉就OK
John avatar
By John
at 2009-06-29T14:43
但是num1 num2必須配置於shared memory中
Ula avatar
By Ula
at 2009-07-03T16:37
@@ vector sum怎麼可能把register用完
Necoo avatar
By Necoo
at 2009-07-04T11:27
你可以把之前用過的register重複利用

3850 開1920*1080 使用hdmi>dvi線

Damian avatar
By Damian
at 2009-06-18T17:13
小弟螢幕是 ASUS VH226H 當我把解析度開到1920*1080後 整個windows無法變成全螢幕 邊邊都是黑黑的 使用dvi線就沒問題 可適用hdmiandgt;dvi線就會出現這樣的情況 有沒有辦法可以調整呢 只有把解析度調低才不會這樣 可是圖示大小又怪怪的~and#34;~ 請大家幫小弟解答一下 ...

老電腦想升級顯示卡

Jack avatar
By Jack
at 2009-06-18T16:07
CPU:AMD Sempron(tm) Processor 2500+ 1.40G 主機板:MSI K8MM-V 記憶體:512*2 電源供應器:300W 大約是4年前組的電腦,使用的是內建的顯卡 平常會看看電影,而我比較沒有遊戲的需求, 最常玩的也只有 EA 的 MLB 或是 NBA 除了有時候 ...

關於CUDA的bank conflict

Susan avatar
By Susan
at 2009-06-18T16:03
近來閱讀了版上a大關於cuda的文章 因此想把自己的程式改寫成可利用GPU執行 但bank conflict卻始終困擾著我 舉例來說 int k=threadIdx.x num1=data1[k] num2=data2[k] sum[k]=num1+num2; 若把num1 ...

9600GT VS 3850

Yuri avatar
By Yuri
at 2009-06-18T15:42
最近朋友要升級顯卡 但是預算不高 本來是想推薦他買1990元版本的3850 不過他似乎偏好nvidia的顯卡 查了一下後發現效能比3850好一些的9600GT有2490的 牌子是旌宇 朋友的預算極點也差不多是2500這邊 不知道這張卡的風評好嗎 因為我看其他牌子的9600GT都要接近3000的 ...

該用哪張顯卡比較好呢?

Megan avatar
By Megan
at 2009-06-18T14:47
最近終於要把六年老電腦淘汰啦, 但是菜單開了以後突然對顯卡有點猶豫, 以下是我的配備: CPU:AMD AM3 PhenomII X2-550 MB:技嘉 MA78GM-MD2H RAM:創見 2G DDRII-800 *2 VGA:HIS ICQ4 HD4850 512M??? HDD:WD SATAII ...