bank conflict的疑惑

zyt0303 2009-04-01 04:57:09
如果有以下的声明:
__shared__ int a[12];
__shared__ int b[4];
__shared__ int c[17];
1 那么数据和banks的对应关系是怎样的,是这样的吗:a[0]是bank0, a[1]是bank1, a[2]是bank2, ... a[11]是bank11, b[0]是bank12, b[1]是bank13, b[2]是bank14, b[3]是bank15, c[0]是bank0, c[1]是bank1, ... , c[15]是bank15, c[16]是bank0 ?
2 如果两个线程thread0和thread1同时存取a[0], 则会产生bank conflict; 如果一个线程thread0同时存取a[0]和c[0]或者c[0]和c[16](如果对应同样的bank)(而其他线程都不存取这些地址)会不会产生bank conflict, 会不会产生延迟?
...全文
264 2 打赏 收藏 转发到动态 举报
写回复
用AI写文章
2 条回复
切换为时间正序
请发表友善的回复…
发表回复
  • 打赏
  • 举报
回复
这个问题要看编译器是如何分配smem的了。如果是紧致分配内存,则如LZ的1)中所提的(猜想是如此)。
如有空洞(开始地址对齐什么的),则不一定。
资料太少,不好有肯定的回答(不保证一定如此......)。
另外,要注意对warp中只有1个访问冲突的地方,不形成bank conflict。而是启用广播机制。
在编程指南中有一段SM如何处理warp中smem访问的详细描述,不妨仔细阅读一下:5.1.2.5(特别是图5-8)
同一thread读取数据之间该不存在bank conflict问题,本来就是串行的......
sunpiqi 2009-04-02
  • 打赏
  • 举报
回复
这个问题我也在研究中,期望高手帮忙,顶一下!

580

社区成员

发帖
与我相关
我的任务
社区描述
CUDA™是一种由NVIDIA推出的通用并行计算架构,该架构使GPU能够解决复杂的计算问题。 它包含了CUDA指令集架构(ISA)以及GPU内部的并行计算引擎。
社区管理员
  • CUDA编程社区
加入社区
  • 近7日
  • 近30日
  • 至今
社区公告
暂无公告

试试用AI创作助手写篇文章吧