为什么在使用同步线程时不需要使用volatile变量?

根据 the programming guide , __syncthreads() 都是执行障碍 和 记忆围栏:

等待线程块中的所有线程都达到此点 以及这些线程在 α同步码() 对块中的所有线程都可见 .

内存隔离功能(即“可见性”)“强制”对共享内存和全局内存的所有更新对其他线程可见。

我想这就是你要问的。我不认为像“使用时不需要使用volatile”这样的笼统说法 α同步码() “这是一个明智的想法。这取决于代码。但在某些情况下,例如 classical parallel reduction ,使用 α同步码() 在块范围缩减的每个步骤中,都意味着用于这种缩减的共享内存不需要标记为 volatile .

自从 α同步码() 是执行屏障还是内存屏障,我们可以做一些声明 α同步码() 不适用于 __threadfence() 独自一人。

假设我有这个代码:

__global__ void k(int *data){
  ...
  *data = 1;
  __syncthreads();
  if (*data == 1){
    ...}
  ...
}

在这种情况下,执行if语句的特定块中的任何线程都可以确保看到 *data 为1。这有两个组成部分:

1. α同步码() 是一个(设备范围)内存围栏。它强制已写入值的任何线程使该值可见。这实际上意味着,由于这是一个设备范围的内存边界,因此写入的值至少已填充了二级缓存(二级缓存是全局内存的设备范围中介器,实际上是全局内存的代理)。

2. α同步码() 是一个(粗制滥造的)执行屏障。它强制所有线程在任何线程继续之前到达屏障。这种执行顺序行为意味着,当任何线程执行上述if语句时,上面第1项中的保证就生效了。

注意这里有一个微妙的区别。其他线程,在其他块中,在代码的其他点上,可能看到或可能看不到由不同块写入的值。

只有当我们将执行同步和内存隔离结合在一起时,我们才能确定由一个线程填充的值对另一个线程是真正可见的。在不使用合作组的情况下,CUDA不提供跨单独块同步执行的机制。

αththffcess() 本身,使价值 最后 可见,但如果不理解写入线程和读取线程之间的相对执行顺序,就不可能仅仅基于代码检查来作出保证。

同样地 不稳定的 保证类似于 αththffcess() (对于书写线),但也有些不同。 αththffcess() 确保写入线程最终将其数据推送到L2(即使其可见)。 不稳定的 执行类似的操作,但也保证读取线程不会读取l1中的“过时副本”,而是在代码中读取该值时转到l2(至少)以获取当前值。

注意,在另一个SM上,设备代码活动不会触发一级缓存数据的“无效”。 不稳定的 有效地保证负载将绕过L1。 不稳定的 同时保证商店将直接进入L2。 αththffcess() 做类似于后者的事情(至少在线程超过 αththffcess() ,但不保证其他SMS中的L1状态,也不保证其他SMS中的线程将如何读取该值。