不正确的同步的“while”循环（仅在释放模式存在的）

Question

我有一个内核“while”循环，在以下示例代码，其迭代地更新使用关于邻居 （仅一个邻居的信息的阵列的元件）。当在当前迭代中没有元素被改变时，该循环停止。

不幸的是，在线程的某些情况下，部分外出这个循环的过早（例如，如果他们忽略同步屏障）。 一些输入正确每次处理和其它输入（其中许多是）被不正确地每次 （即不存在随机因素）处理。奇怪的是，这个错误只在发布版本中发生，而Debug版本总是 工作正常。更准确地说，CUDA编译器选项“-G（GPU生成调试信息）”确定 处理是否正确。大小为32x32或更小的数组始终可以正确处理。

下面是一个示例代码：

__global__ void kernel(int *source, int size, unsigned char *result, unsigned char *alpha) 
{ 
    int x = threadIdx.x, y0 = threadIdx.y * 4; 
    int i, y; 
    __shared__ bool alpha_changed; 

    // Zero intermediate array using margins for safe access to neighbors 
    const int stride = MAX_SIZE + 2; 
    for (i = threadIdx.x + threadIdx.y * blockDim.x; i < stride * (stride + 3); i += blockDim.x * blockDim.y) 
    { 
     alpha[i] = 0; 
    } 
    __syncthreads(); 

    for (int bit = MAX_BITS - 1; bit >= 0; bit--) 
    { 
     __syncthreads(); 

     // Fill intermediate array with bit values from input array 
     alpha_changed = true; 
     alpha[(x + 1) + (y0 + 1) * stride] = (source[x + (y0 + 0) * size] & (1 << bit)) != 0; 
     alpha[(x + 1) + (y0 + 2) * stride] = (source[x + (y0 + 1) * size] & (1 << bit)) != 0; 
     alpha[(x + 1) + (y0 + 3) * stride] = (source[x + (y0 + 2) * size] & (1 << bit)) != 0; 
     alpha[(x + 1) + (y0 + 4) * stride] = (source[x + (y0 + 3) * size] & (1 << bit)) != 0; 
     __syncthreads(); 

     // The loop in question 
     while (alpha_changed) 
     { 
      alpha_changed = false; 
      __syncthreads(); 
      if (alpha[(x + 0) + (y0 + 1) * stride] != 0 && alpha[(x + 1) + (y0 + 1) * stride] == 0) 
      { 
       alpha_changed = true; 
       alpha[(x + 1) + (y0 + 1) * stride] = 1; 
      } 
      __syncthreads(); 
      if (alpha[(x + 0) + (y0 + 2) * stride] != 0 && alpha[(x + 1) + (y0 + 2) * stride] == 0) 
      { 
       alpha_changed = true; 
       alpha[(x + 1) + (y0 + 2) * stride] = 1; 
      } 
      __syncthreads(); 
      if (alpha[(x + 0) + (y0 + 3) * stride] != 0 && alpha[(x + 1) + (y0 + 3) * stride] == 0) 
      { 
       alpha_changed = true; 
       alpha[(x + 1) + (y0 + 3) * stride] = 1; 
      } 
      __syncthreads(); 
      if (alpha[(x + 0) + (y0 + 4) * stride] != 0 && alpha[(x + 1) + (y0 + 4) * stride] == 0) 
      { 
       alpha_changed = true; 
       alpha[(x + 1) + (y0 + 4) * stride] = 1; 
      } 
      __syncthreads(); 
     } 
     __syncthreads(); 

     // Save result 
     result[x + (y0 + 0) * size + bit * size * size] = alpha[(x + 1) + (y0 + 1) * stride]; 
     result[x + (y0 + 1) * size + bit * size * size] = alpha[(x + 1) + (y0 + 2) * stride]; 
     result[x + (y0 + 2) * size + bit * size * size] = alpha[(x + 1) + (y0 + 3) * stride]; 
     result[x + (y0 + 3) * size + bit * size * size] = alpha[(x + 1) + (y0 + 4) * stride]; 
     __syncthreads(); 
    } 
} 

// Run only 1 thread block, where size equals 64. 
kernel <<< 1, dim3(size, size/4) >>> (source_gpu, size, result_gpu, alpha_gpu); 

该样品内核预期的结果是阵列，其中每行只能包含连续的时间间隔的“1”值 。但是取而代之的是，我得到了一些行，其中“0”和“1”以某种方式交替出现。

此错误被再现于我的移动GPU的GeForce 740M（开普勒），在Windows 7 X64的SP1，在任CUDA 6.0或6.5， 使用Visual C++ 2012或2013我还可以提供样品的Visual Studio项目与样本输入数组（即不正确处理）。

我已经尝试了syncthreads（），fences和“volatile”限定符的不同配置，但仍然存在此错误 。

任何帮助表示赞赏。

Answer 1

我认为这个问题是alpha_changed您的访问。请记住，这只是块中所有线程的一个值。有一个经重置此变量，而另一个经检查循环条件之间的竞争条件：

// The loop in question 
    while (alpha_changed) 
    { 
     alpha_changed = false; 
     // ... 
     // alpha_changed may be set to true here 
     // ... 

     __syncthreads(); 

     // race condition window here. Another warp may already execute 
     // the alpha_changed = false; line before this warp continues. 
    } 

的关键是共享变量设置为false之前做__syncthreads()。

您可以使用一个局部变量内循环，以找出是否该线程所做的任何更改。这避免了不得不在整个地方使用__syncthreads()。然后做在循环结束时减少：

// The loop in question 
    while (alpha_changed) 
    { 
     bool alpha_changed_here = false; 
     // ... 
     // alpha_changed_here may be set to true here 
     // ... 

     __syncthreads(); 
     alpha_changed = false; 
     __syncthreads(); 
     // I think you can get away with a simple if-statement here 
     // instead of a proper reduction 
     if (alpha_changed_here) alpha_changed = true; 
     __syncthreads(); 
    } 

据我所知，在共享内存只用一个变量的这种方法当前工作。如果您想确定，请使用适当的简化算法。你可以使用__any()来减少一个指令中的32个值。要使用的算法取决于块的大小（我不知道确切的行为是大小不是32的倍数）。