原子地在共享内存中编写结构答案

【问题标题】：Writing an struct in shared memory atomically原子地在共享内存中编写结构
【发布时间】：2016-03-17 20:39:29
【问题描述】：

为了访问结构，nvcc 会生成一个代码来逐个字段地读取/写入结构。具有这种结构：

typedef struct cache_s {
    int tag;
    TYPE data;
} cache_t;

生成以下 PTX 代码以将此类变量写入共享内存：

st.shared.f64   [%rd1+8], %fd53;
st.shared.u32   [%rd1], %r33;

这可能会在程序的执行中引发逻辑错误。如果线程块的两个并发线程在同一共享内存地址写回不同的值，则来自不同结构的字段可能会混淆。 CUDA 编程指南指出：

如果warp执行的非原子指令写入相同的多个线程在全局或共享内存中的位置扭曲的，发生的序列化写入的数量位置因设备的计算能力而异（请参阅计算能力 2.x、计算能力 3.x 和计算能力 5.x)，哪个线程执行最后的写入是未定义的。

据此，我预计其中一个线程会写入其完整的结构（整个字段一起），并且我不希望字段的混合（来自不同的写入）形成未定义的值。有没有办法强制 nvcc 生成我期望的代码？

更多信息：

NVCC 版本：7.5

【问题讨论】：

没有。您的期望不切实际/错误。
From this, I expect one of the threads writes its complete structure (whole the fields together)：我真的很想通过阅读您从 CUDA 编程指南中引用的部分来了解您是如何得出这个结论的。
@RogerDahl 我的印象是该位置可能可以扩展到变量。
@talonmies 我认为在共享内存上使用原子锁是一种解决方法，但这是一个糟糕的解决方案。有更好的解决方案吗？
您真正想说的是“我希望我的硬件支持完全任意的内存事务大小”，而这在 GPU 或我熟悉的任何当前或历史平台上都不会发生。

【解决方案1】：

这可能会在程序的执行中引发逻辑错误。如果一个线程块的两个并发线程在同一个共享内存地址写回不同的值，来自不同结构的字段可能会混淆。

如果您需要块中一个线程的完整结果，同时丢弃其他线程的结果，只需让其中一个线程（线程 0 经常用于此）写出其结果并让其余线程跳过写：

__global__ void mykernel(...)
{
    ...
    if (!threadIdx.x) {
        // store the struct
    }
}

有没有办法强制 nvcc 生成我期望的代码？

您希望看到 NVCC 生成一条指令，对任意大小的完整结构进行原子写入。没有这样的指令，所以，不，你不能让NVCC生成代码。

我认为在共享内存上使用原子锁是一种解决方法，但这是一个糟糕的解决方案。有没有更好的解决方案？

我们无法告诉您什么是更好的解决方案，因为您还没有告诉我们您要解决的问题是什么。在 CUDA 中，原子操作通常仅用于在读-修改-写操作期间锁定单个 32 位或 64 位字，因此不适合保护完整结构。

有一些并行操作，有时称为并行原语，例如“reduce”和“scan”，它们允许在没有锁定的情况下解决许多类型的问题。例如，您可能首先启动一个内核，其中每个线程将其结果写入一个单独的位置，然后启动一个新的内核，该内核执行并行归约以选择您需要的结果。

【讨论】：