何时在共享CUDA内存中使用volatile

Question

在什么情况下，您应该在CUDA内核的共享内存中使用volatile关键字？我知道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，还是只会缓存元素？

谢谢！

- Taj Morton

Answer 1

如果您不将共享数组声明为volatile，那么编译器可以通过将它们定位在寄存器中（其作用域仅限于单个线程），自由地优化共享内存中的位置，适用于任何线程。这对于您是否仅从一个线程访问该特定共享元素都是正确的。因此，如果您使用共享内存作为块的各个线程之间的通信工具，则最好将其声明为volatile。然而，这种通信模式通常还需要执行障碍以强制读/写的顺序，因此请继续阅读下面有关障碍的内容。

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

在您的情况下，当每个线程都访问其自己的共享内存元素，并且唯一的线程间访问发生在一个明确定义的位置时，您可以使用内存障碍函数来强制编译器将暂存在寄存器中的任何值驱逐回共享数组。因此，您可能认为__threadfence_block()很有用，但在您的情况下，__syncthreads()已经内置了内存障碍功能。所以您的__syncthreads()调用足以强制线程同步，以及强制任何寄存器缓存值在共享内存中被驱逐回共享内存。

顺便说一句，如果您关心代码末尾的那个规约，可以考虑使用并行规约方法来加速它。

- Robert Crovella

---

很棒的答案，我不知道内存栅栏。谢谢！ - Taj Morton

Answer 2

简单来说，对于其他人来说：

调用__syncthreads()比将共享内存声明为volatile更强。 __syncthreads()使给定工作组中的所有线程在1个公共点停止并同步内存。

另一方面，volatile通过防止编译器进行任何缓存优化来保持给定内存缓冲区在线程之间保持一致（因此可能会带来成本），但每个线程都可以按自己的节奏前进，这使得编译器/硬件可以执行各种调度优化。 （请注意，如果写入由多个处理器指令组成，则volatile不能保证数据完整性）

总之，当您需要的只是线程之间的内存一致性，而不是全部在1个点停止时，volatile通常比__syncthreads()提供更好的性能。但是，具体算法甚至输入数据可能会影响结果，因此如果需要挤出每一个性能位，请测试两种方法。

此外，如果工作组中的活动线程数小于SIMD宽度（warp size），那么可以使用volatile代替__synchthreads()，因为同一warp中的所有线程都同步执行指令。例如，参见last wrap unrolling optimization to parallel reduction algorithm（幻灯片21-23），最初使用__synchthreads()，后来仅依赖于volatile当活动线程数量小于warp大小时。

- morgwai

---

使用动态并行性，这不再是真的了。您需要调用__syncwarp()而不是依赖于隐式的warp同步执行。 - Sebastian

问题也在于，如果没有使用 volatile，编译器是否会在线程0中缓存未初始化的 products 值，并且从任何地方都不会读取它。 - Sebastian

@Sebastian，哪一部分不准确呢？你能提供一些链接吗？这将非常有帮助 :) 谢谢！ - morgwai

“由于同一warp中的所有线程都会同步执行指令，因此存在独立线程调度。” - 抱歉，我的意思是独立线程调度，请参见此处的示例 https://docs.nvidia.com/cuda/volta-tuning-guide/index.html#sm-independent-thread-scheduling 和这里的章节关于warp同步的内容 https://developer.nvidia.com/blog/using-cuda-warp-level-primitives/。 - Sebastian