CUDA非原子写入冲突结果

Question

我正在编写一个需要迭代直到完成的函数。我意识到我可以使用原子运算符，但是速度在此内核中至关重要，我怀疑可能不需要它们。

我已经包含了一小段伪代码来演示我打算做的事情

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

如果没有线程找到条件为真，则lbRepeat将为false。

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

如果多个线程同时将true写入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

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

最终结果将是其中一个写入值最终将存储在该存储位置中。哪个值未定义。如果所有写入的值都相同，则可以确保该值将在该位置结束。

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

假设只有这两个操作（一次写入和一次读取），否。读取的值将是写入开始之前的值或写入完成之后的值。如果您正在进行多次写入，那么当然可以看到第一个问题的答案。实际的写入值是不确定的，除了它会好像其中一次写入成功而其他所有都不成功一样。

我是在示例正确对齐的8位，16位或32位数据类型的上下文中做出上述声明的。