【并行计算-CUDA开发】warp是调度和执行的基本单位而harf-warp为存储器操作基本单位
warning C4819 ......
解决这个警告的方法是打开出现warning的文件,Ctrl+A全选,然后在文件菜单:file->Advanced save options,在弹出的选项中选择新的编码方式为:UNICODE- codepage 1200 ,点确定后重新编译。
为什么会出现这个警告呢?原因在于NvidIA方面,他们的在编写文件的时候用的字符集不通用。
2.关于warp和half-warp
一个warp包含32个threads。warp是调度和执行的基本单位,half-warp是存储器操作的基本单位,这两个非常重要。
在分支的时候,warp大显身手
有合并访问和bank conflict的时候half-warp当仁不让。
每个bank的带宽为32bit = 4byte= 4 char = 1 int = 1float
只要half-warp中的线程访问的数据在同一个段中,就可以满足合并访问条件。
(cudaMalloc分配的存储器,能够保证其首地址至少会按照256byte进行对齐,因此选择合适的线程块大小,如16的整数倍,能使half-warp的访问请求按照段长对齐。使用__align__(8),__align__(16)限定符来定义结构体,可以使对结构体构成的数组进行访问时能够做到对齐到段)
(在sdk的很多例子中,在定义sharedmemory的时候,都用了宽度17或者threaddim.x+1的行数,来避免bank conflict。如__shared__ float bs[16][17])
3 block和grid的维度设计
为了有效的利用sm中的执行资源,block中的线程数量应该为32的整数倍,最好让线程数量保持在64~256之间。
每个block的blockdim.x应该为16或16的整数倍,以提高global和shared memory的访问效率。
4. volatile关键字
如果共享内存用于线程块内的warp间通信,则在共享内存声明前必须使用volatile关键字,以避免有误读缓存数据引起的错误。