何时使用volatile共享CUDA内存

Question

在什么情况下应该使用关键字volatile与CUDA内核的共享内存？据我所知，volatile告诉编译器从未缓存任何值，但我的问题是关于与共享数组的行为：

__shared__ float products[THREADS_PER_ACTION]; 

// some computation 
products[threadIdx.x] = localSum; 

// wait for everyone to finish their computation 
__syncthreads(); 

// then a (basic, ugly) reduction: 
if (threadIdx.x == 0) { 
    float globalSum = 0.0f; 
    for (i = 0; i < THREADS_PER_ACTION; i++) 
     globalSum += products[i]; 
} 

我需要products挥发在这种情况下？每个数组入口只能被单个线程访问，除了最后一行，所有内容都由线程0读取。是否有可能编译器可以缓存整个数组，因此我需要它是volatile，或者它只会缓存元素？

谢谢！

Answer 1

如果您没有声明共享数组为volatile，那么编译器可以通过将它们定位到寄存器（其范围特定于单个线程）中来优化共享内存中的位置，对于任何线程。无论您是否仅从一个线程访问该特定共享元素，情况都是如此。因此，如果使用共享内存作为块间线程间的通信工具，最好声明它为volatile。

显然，如果每个线程只访问自己的共享内存元素，并且从来没有与另一个线程相关联，那么这并不重要，编译器优化不会破坏任何东西。

在你的情况下，如果你有一段代码让每个线程访问它自己的共享内存元素，并且唯一的线程间访问发生在一个很好理解的位置，你可以使用memory fence function来强制编译器驱逐临时存储在寄存器中的任何值，退出到共享阵列。所以你可能认为__threadfence_block()可能是有用的，但在你的情况下，__syncthreads()already has memory-fencing functionality built in。所以你的__syncthreads()调用足以强制线程同步，并强制共享内存中的任何寄存器缓存值被清除回共享内存。顺便说一下，如果代码末尾的减少与性能有关，那么可以考虑使用并行减少方法来加速它。