cuda share memory 声明时最好加上 volatile 关键字。
__shared__ volatile T sdata[blockSize];
volatile 表示这个变量会被外部程序读取或者修改。比如变量所在地址是一个端口,外部程序可以读取或者修改端口的数据。为防止编译器优化时使用寄存器缓存这个变量, 我们必须加上volatile关键字。
一个例子是cuda 中的warp reduce 如果不使用volatile关键字,会产生错误结果。
if (threadIdx.x < 32)
{
if (blockSize >= 64) { sdata[threadIdx.x] = mySum = op(mySum, sdata[threadIdx.x + 32]); }
if (blockSize >= 32) { sdata[threadIdx.x] = my