除了使用volatile正如其他答案中建议的那样，使用__threadfence还需要适当地获得具有安全内存排序的原子加载。
虽然一些评论说只使用普通读取，因为它不能撕裂，但这与原子加载不同。原子不仅仅是撕裂：
正常读取可能会重用寄存器中已有的先前加载，因此可能不会反映其他 SM 具有所需内存排序的更改。例如，int *flag = ...; while (*flag) { ... }只能阅读flag一次并在循环的每次迭代中重用该值。如果您正在等待另一个线程更改标志的值，您将永远不会观察到更改。这volatile修饰符确保每次访问时实际上都是从内存中读取该值。请参阅关于易失性的 CUDA 文档 https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#volatile-qualifier了解更多信息。
此外，您需要使用内存栅栏来在调用线程中强制执行正确的内存排序。如果没有栅栏，您将获得 C++11 术语中的“宽松”语义，而在使用原子进行通信时这可能是不安全的。
例如，假设您的代码（非原子地）将一些大数据写入内存，然后使用正常写入来设置原子标志以指示数据已被写入。指令可能会被重新排序，硬件缓存行可能不会在设置标志之前被刷新等等。结果是这些操作不能保证以任何顺序执行，并且其他线程可能不会按照您期望的顺序观察这些事件：允许写入标志before受保护的数据被写入。
同时，如果读取线程在有条件加载数据之前也使用正常读取来检查标志，则会在硬件级别出现竞争。无序和/或推测执行可能会在标志读取完成之前加载数据。然后使用推测加载的数据，该数据可能无效，因为它是在读取标志之前加载的。
放置得当的内存栅栏可以通过强制执行指令重新排序来防止此类问题，从而不会影响您所需的内存顺序，并且以前的写入对其他线程可见。__threadfence()和朋友也被覆盖在 CUDA 文档中 https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#memory-fence-functions.
将所有这些放在一起，在 CUDA 中编写您自己的原子加载方法看起来像：
// addr must be aligned properly.
__device__ unsigned int atomicLoad(const unsigned int *addr)
{
const volatile unsigned int *vaddr = addr; // volatile to bypass cache
__threadfence(); // for seq_cst loads. Remove for acquire semantics.
const unsigned int value = *vaddr;
// fence to ensure that dependent reads are correctly ordered
__threadfence();
return value;
}
// addr must be aligned properly.
__device__ void atomicStore(unsigned int *addr, unsigned int value)
{
volatile unsigned int *vaddr = addr; // volatile to bypass cache
// fence to ensure that previous non-atomic stores are visible to other threads
__threadfence();
*vaddr = value;
}
对于其他非撕裂加载/存储大小，这可以类似地编写。
通过与一些从事 CUDA 原子工作的 NVIDIA 开发人员交谈，我们似乎应该开始看到 CUDA 中对原子的更好支持，并且 PTX 已经包含具有获取/释放内存顺序的加载/存储指令 https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#release-acquire-patterns语义——但目前无法在不诉诸内联 PTX 的情况下访问它们。他们希望在今年的某个时候将它们添加进来。一旦这些就位，一个完整的std::atomic实施应该不会落后太远。
