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

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

__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]分配了rnumber的值。我不明白，但是正如您在我的代码中所看到的，我已经更改了if以适应这种情况。

这使我想到了麻烦的数字2：尽管NSight将X的值显示为-1，但if {}被完全跳过（命中没有断点），执行直接跳转到else。

我的问题：

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

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

我尝试将语法更改为if(X!=-1)； if的true分支仍未执行。

最佳答案

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

因此，您在X，rnumber和localMap值中看到的行为是预期的。

我不能解决您的第二个问题，因为我不使用NSight，也不知道它是如何工作的-根据您的代码，应该对您的true分支进行评估（但请注意：您的false分支也是如此-因为它是多线程的，所以某些线程可以将条件评估为true，有些评估为false-我的猜测/假设是，您必须以某种方式告诉调试器您想调试哪个线程/线程/块，然后查看false）。

关于c++ - CUDA内核似乎忽略了“if”语句，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/18736863/

localmap

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