3.4 共享存储器优化
共享存储器比较靠近SM的,因此其速度相当的快,一般而言可以在一到二个时钟内读写,因此使用共享存储器取代全局存储器会极大的节约带宽,但是共享存储器 的使用是有要求,它要求数据有局部性(一个block内共享)重用。当然有时也可用共享存储器取代寄存器,这会在后面说到。
一般的使用共享存储器的方式是:s_x[tid] = x[id];其中tid是块内线程的块内索引,也就是类似threadIdx,id是线程的网格内索引。然后可以访问s_x取代访问x。这样一般可以提升 程序的效率。
要注意的是,使用共享存储器往往伴随着存储器的一致性问题,此时可求助于__syncthreads()和memory fence大神。
3.4.1 存储体冲突
类似于合并访问,共享存储器一次也能满足 半束线程的要求,条件是这半束线程访问的数据在不同的存储体中。CUDA将共享存储器组织成16列,每一列称为 一个存储体。如果有两个或以上的线程访问的数据在同一个存储体中,此时不能在一次存储体访问中,满足半束线程的要求,这称为存储体冲突。一般而言,存储体 冲突对性能的影响还不是太大,当然具体的影响要看半束内最多有多少线程落入同一存储体内。要提到的编程时要注意的是对于一个SM来说,共享存储器的数量是有限的,如果超出其使用限制,其结果未知,我的经验告诉我,有时没有问题,有时会有大问题。另外, 使用共享存储器有时可能会影响程序的可扩展性,这个在编写库代码的时候要特别注意。
共享存储器比较靠近SM的,因此其速度相当的快,一般而言可以在一到二个时钟内读写,因此使用共享存储器取代全局存储器会极大的节约带宽,但是共享存储器 的使用是有要求,它要求数据有局部性(一个block内共享)重用。当然有时也可用共享存储器取代寄存器,这会在后面说到。
一般的使用共享存储器的方式是:s_x[tid] = x[id];其中tid是块内线程的块内索引,也就是类似threadIdx,id是线程的网格内索引。然后可以访问s_x取代访问x。这样一般可以提升 程序的效率。
要注意的是,使用共享存储器往往伴随着存储器的一致性问题,此时可求助于__syncthreads()和memory fence大神。
3.4.1 存储体冲突
类似于合并访问,共享存储器一次也能满足 半束线程的要求,条件是这半束线程访问的数据在不同的存储体中。CUDA将共享存储器组织成16列,每一列称为 一个存储体。如果有两个或以上的线程访问的数据在同一个存储体中,此时不能在一次存储体访问中,满足半束线程的要求,这称为存储体冲突。一般而言,存储体 冲突对性能的影响还不是太大,当然具体的影响要看半束内最多有多少线程落入同一存储体内。要提到的编程时要注意的是对于一个SM来说,共享存储器的数量是有限的,如果超出其使用限制,其结果未知,我的经验告诉我,有时没有问题,有时会有大问题。另外, 使用共享存储器有时可能会影响程序的可扩展性,这个在编写库代码的时候要特别注意。
来自 “ ITPUB博客 ” ,链接:http://blog.itpub.net/23057064/viewspace-666187/,如需转载,请注明出处,否则将追究法律责任。
转载于:http://blog.itpub.net/23057064/viewspace-666187/