何时使用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.

显然,如果每个线程只访问自己的共享内存元素,而不是那些与另一个线程相关的元素,那么这无关紧要,编译器优化不会破坏任何东西.

在你的情况下,你有一段代码,其中每个线程访问它自己的共享内存元素,并且唯一的线程间访问发生在一个很好理解的位置,你可以使用内存栅栏函数 强制编译器驱逐临时存储在寄存器中的任何值,退回到共享阵列.所以你可能认为这__threadfence_block()可能有用,但在你的情况下,__syncthreads() 已经内置了内存防护功能.因此,您的__syncthreads()调用足以强制线程同步以及强制共享内存中的任何寄存器缓存值被逐出回共享内存.

顺便说一句,如果代码末尾的减少是性能问题,您可以考虑使用并行缩减方法来加速它.