c++ - CUDA 内核似乎忽略了 "if"语句

Question

接下来是我的内核中运行不正常的部分，然后是对我在调试时发现的内容的解释。

__global__ void Mangler(float *matrix, int *map)
{
    __shared__ signed int localMap[N];

    if(0 == threadIdx.x) 
    {
        for(int i=0; i<N; i++) 
            localMap[i] = -1;
    }

    __syncthreads();

    int fn = ...; // a lot of code goes into this number, skipped for clarity
    int rnumber = threadIdx.x;

    int X = atomicCAS(&localMap[fn], -1, rnumber); // Spot of bother 1

    if(X == -1) // Spot of bother 2
    {
        // some code
    }
    else 
    {
        // other code
    }
}

我在文档中发现 atomicCAS(*address, compare, value) 基本上返回(并保存到给定地址)(old == compare ? value : old)，其中old是函数执行前地址处的值。

据此，我相信执行 int X = atomicCAS(&localMap[fn], -1, rnumber); 应该有两种可能的结果(根据 NVidia Cuda C 编程指南):

• 如果 localMap[fn] == -1 那么 X 的值应该是 rnumber 和 localMap[fn] 的值应为 rnumber。 这不会发生。
• 如果 localMap[fn] != -1 那么 X 应该被设置为 localMap[fn] 的值并且所述值应该完好无损。

相反，正如使用 NSight 进行调试所显示的那样，X 被赋值为 -1，而 localMap[fn] 被赋值为 r 编号。我不明白，但正如您在我的代码中看到的那样，我更改了 if 以捕捉这种情况。

这让我想到了麻烦点 2:虽然 NSight 将 X 的值显示为 -1，但是 if {} 被完全跳过(内部没有断点随便打)然后执行直接跳到else。

我的问题:

• 我是否完全误解了 atomicCAS？ 是的，我误解了
• 什么会导致和 if 直接跳转到设备代码中的 else？

我使用的是 NVidia CUDA 5.5、Windows 8 上的 Visual Studio 2012 x64、NVidia Nsight Monitor Visual Studio Edition 3.1。机器的GPU是NVidia GeForce GTX 550 Ti。

我尝试将语法更改为 if(X!=-1)； if 的真正分支仍然没有被执行。

Answer 1

根据文档，atomicCAS 返回旧值，这意味着，在您的列表中，您的两个结果是错误的。您的 X 将始终设置为 localMap[fn] 的旧值，无论它具有哪个值。根据与-1的比较设置的，就是localMap[fn]的新值。当它为-1 时，它被设置为rnumber，否则它保持不变。

因此，您看到的 X、rnumber 和 localMap 值的行为符合预期。

我无法解决你的第二个问题，因为我不使用 NSight，也不知道它是如何工作的——根据你的代码，你的真实分支应该被评估(但要小心:你的错误分支也——因为它是多线程的一些线程可以将条件评估为 true，而一些线程可以将条件评估为 false - 我的猜测/假设是您必须以某种方式告诉您的调试器您想要调试哪个线程/warp/ block 并且您查看了 false)。