0

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

__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 Programming Guide）：

  • if localMap[fn] == -1thenX的值应该是rnumber并且localMap[fn]应该是rnumber。这不会发生。
  • 如果localMap[fn] != -1thenX应该设置为的值，localMap[fn]并且所述值应该保持不变。

相反，正如使用 NSight 进行的调试向我展示的那样，X它被分配了 -1，而localMap[fn]被分配了rnumber. 我不明白，但正如您在我的代码中看到的那样，我已经更改了if以捕捉这种情况。

这让我想到了第 2 个麻烦点：虽然 NSight 显示的值为X-1，但if {}它被完全跳过（没有任何断点）并且执行直接跳转到else.

我的问题：

  • 我atomicCAS完全误解了吗？ 是的，我做到了
  • 什么可能导致if直接跳入else设备代码，哪些应该评估为真？

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

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

4

1 回答 1

1

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

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

我无法帮助您解决第二个问题，因为我不使用 NSight，并且不知道它是如何工作的 - 根据您的代码，应该评估您的真实分支（但要小心：您的错误分支也是 - 因为它是多线程的，所以一些线程可以将条件评估为真，有些为假-我的猜测/假设是您必须以某种方式告诉您的调试器您要调试哪个线程/扭曲/块，然后您查看了错误）。

于 2013-09-11T09:41:39.207 回答