cuda share memory

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