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关键字，以避免有误读缓存数据引起的错误。