"while" 循环内的不正确同步（仅在 Release 模式下发生）

Question

我有一个带有 "while" 循环的内核，它使用有关邻居的信息迭代更新数组的元素（下面的示例代码中只有一个邻居）。当当前迭代中没有元素发生变化时，此循环停止。

不幸的是，在某些情况下，部分线程过早地退出了这个循环（比如它们忽略了同步屏障）。有些输入每次都正确处理，而其他输入（其中很多）每次都处理不正确（即没有随机因素）。奇怪的是，这个错误只发生在Release版本，而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 (Kepler)、Windows 7 x64 SP1、CUDA 6.0 或 6.5 上重现，使用 Visual C++ 2012 或 2013。我还可以提供一个示例 Visual Studio 项目和示例输入数组（即处理不正确）。

我已经尝试了 syncthreads()、fences 和 "volatile" 限定符的不同配置，但是这个错误留下来了。

感谢任何帮助。

Answer 1

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

    // 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 个值减去一个 warp。使用的算法取决于块的大小（我不知道确切的行为是大小不是 32 的倍数）。

"while" 循环内的不正确同步（仅在 Release 模式下发生）

Incorrect synchronization inside a "while" loop (occuring only in Release mode)

visual-c++

cuda

thread-synchronization

release-mode