单个线程内CUDA共享内存读写顺序与写线程自身同步性疑问

线程0在#1处完全可以保证i=10且j=10，不需要使用内存栅栏，也没必要引入局部变量。

原因如下：

• 单线程内的指令执行遵循顺序一致性，CUDA编译器和硬件不会对存在数据依赖的指令做重排序。这里j = i的操作依赖于i = 10的写入结果，因此硬件和编译器会确保i的赋值完成后才执行j的赋值。
• 共享内存的读写在单线程内天然同步，线程自身对共享内存的写入操作，后续自己的读取操作必然能看到最新值。跨线程的共享内存需要栅栏是因为不同线程的执行时序不确定，但单线程内不存在这个问题。

示例代码：

__global__ void kernel()
{
    __shared__ int i, j;

    if(threadIdx.x == 0)
    {
        i = 10;
        j = i;
    }

   // #1
}

内容的提问来源于stack exchange，提问作者rafoo