17

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

谢谢！

4

2 回答 2

23

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

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

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

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

于 2013-03-11T04:20:57.167 回答
-1

简单地说，对于其他会来这里的人：

调用__syncthreads()比将共享内存声明为volatile. __syncthreads()导致来自给定工作组的所有线程一起停止在 1 个公共点并同步内存。

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

总而言之，当您只需要线程之间的内存一致性，而不是在 1 点一起停止时，volatile通常会提供比__syncthreads(). 不过，您的数量可能会有所不同，具体取决于特定的算法甚至输入数据，因此如果您需要压缩最后一点性能，请测试这两种方法。

此外，如果工作组中的活动线程的数量小于 SIMD 宽度（warp 大小），则volatile可以使用，而不是使用__synchthreads()同一 warp 中的所有线程同步执行指令。例如，请参见最后一个包装展开优化到并行缩减算法（幻灯片 21-23），它__synchthreads()首先使用，后来仅在volatile活动线程数小于扭曲大小时使用。

于 2021-10-08T06:18:13.990 回答