c++ - CUDA 非原子写入冲突结果

Question

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

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

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

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

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

多个线程同时向lbRepeat写入true会是什么结果？

我想将其扩展到复制整数值(特别是无符号 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);

}

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

我认为这是 Not Acceptable ，但我的测试似乎表明它可以按预期工作。这是因为未签名的短拷贝在 CUDA 中是原子的吗？

总结:

如果两个线程将相同的值写入一个 bool 内存位置，结果会怎样？

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

Answer 1

What is the result if two threads write the same value into one boolean memory location?

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

Can reading from a unsigned short memory location as another thread is writing a new value to the same location return a value which is neither the previous value or the new value in that memory location?

假设只有这些是正在进行的两个操作(一个写入，一个读取)，不。读取值将是写入开始前的值或写入完成后的值。如果您正在进行多个 写入，那么当然可以查看第一个问题的答案。实际写入的值是未定义的，只是好像其中一个写入成功而其他所有写入都没有。

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