探索现代职场中有效沟通的重要性与实践策略
1157
2022-05-29
引言
cuda 将 shared memory 按照 4 字节或 8 字节(默认 4 字节,可以设置为 8 字节)被划分到 32 个 bank (楼)中,不同 bank 之间的内存能同时读写,但是同一个 bank 的不同地址(同一栋楼的不同层)的数据则只能串行读写(如果是同一个 bank 的同一个地址则可以 broadcast,不会出现 bank conflict),因此当同一个 warp 的线程去访问 shared memory 数据时,如果有两个以上线程访问了同一个 bank 的不同地址的数据,就会产生多余的内存事务(transaction)请求(后面有具体实例图示)影响程序的性能.
假设在共享内存上申请了 1024 个 float 数据 —— __shared__ float data[32][32],由于每个 float 正好是 4 字节,且 data 按行存储,则 data[0][0] 就位于第 0 个bank,data[0][1] 位于第 1 个 bank,以此类推 data[row][col] 就被划分在了第 col 个 bank 中,即 col 相同的数据划分至了同一个 bank 的不同地址上。如果一个 warp 的线程按 col 处理 data 那么就会造成 bank conflict.
任务调度
版权声明:本文内容由网络用户投稿,版权归原作者所有,本站不拥有其著作权,亦不承担相应法律责任。如果您发现本站中有涉嫌抄袭或描述失实的内容,请联系我们jiasou666@gmail.com 处理,核实后本网站将在24小时内删除侵权内容。