作者rick209 ()
看板VideoCard
标题Re: [请益] 关於CUDA的bank conflict
时间Thu Jun 18 17:39:34 2009
※ 引述《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)
1F：推 dulcet:num1 和num2是在register, 不会conflict 06/18 17:41
2F：→ dulcet:你一定在num1, num2前加了__shared__了 哪掉就OK 06/18 17:42
※ 编辑: rick209 来自: 140.113.152.194 (06/18 17:44)
3F：→ rick209:但是num1 num2必须配置於shared memory中 06/18 17:44
4F：推 dulcet:@@ vector sum怎麽可能把register用完 06/18 17:46
5F：→ dulcet:你可以把之前用过的register重复利用 06/18 17:47