CUDA 非原子写入冲突结果

Question

我正在编写一个需要迭代直到完成的函数。我意识到我可以使用原子操作符，但在这个内核中速度很关键，我怀疑它们可能不是必需的。

我已经包含了一小段伪代码来演示我打算做什么

__global__ void TestKernel()
{
  __shared__ bool lbRepeat[1];
  do
  {
    lbRepeat=false;
    __syncthreads();
    if(Condition == true) lbRepeat=true;
    __syncthreads();
  }
  while(lbRepeat);
}

如果没有线程发现条件为真，lbRepeat 将为假。

如果一个线程发现条件为真，lbRepeat 将为真。

如果多个线程同时写入 lbRepeat 会有什么结果？

我想将此扩展到复制整数值（特别是无符号 16 位）。除了检查条件外，我还想复制一个无符号的 16 位整数。

__global__ void TestKernel()
{
  __shared__ unsigned short liValues[32*8];
  __shared__ bool lbRepeat[1];

  unsigned long tid = threadIdx.x+threadIdx.y*blockDim.x;
  do
  {
    lbRepeat=false;

    __syncthreads();
    if(Condition == true)
    {
       liValue[tid] = liValue[Some_Value_In_Range];
       lbRepeat=true;
    }
    __syncthreads();
  }
  while(lbRepeat);

}

如果另一个线程在读取内存时正在写入内存，这是否会导致既不返回先前值也不返回新值？我不介意是否返回先前的值或新的值（两者都将有效），但每个位的混合会导致问题。

我认为这是不可接受的，但我的测试似乎表明它可以按预期工作。这是因为未签名的短副本在 CUDA 中是原子的吗？

总结：

如果两个线程将相同的值写入一个布尔内存位置，结果是什么？

当另一个线程正在将新值写入同一位置时，从无符号短内存位置读取是否可以返回一个既不是该内存位置先前值也不是新值的值？

Answer 1

如果两个线程将相同的值写入一个布尔内存位置，结果是什么？

最终结果将是写入的值之一最终会出现在该内存位置。哪个值未定义。如果所有写入的值都相同，则可以确定该值最终会出现在该位置。

当另一个线程将新值写入同一位置时，是否可以从无符号短内存位置读取返回的值既不是该内存位置的先前值也不是新值？

假设这是仅有的两个操作（一个写，一个读），不。读取值将是写入开始之前的值或写入完成之后的值。如果你有 multiple 正在写，那么当然可以看到第一个问题的答案。实际写入的值是未定义的，只是其中一个写入成功而其他所有写入均未成功。

我在正确对齐的 8、16 或 32 位数据类型的上下文中做出上述陈述，您的示例就是这样。