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

CUDA kernel seemingly ignoring "if" statement

本文关键字：if 语句 内核 CUDA      更新时间：2023-10-16

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

__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哪些评估为真，直接跳入设备代码中的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的真正分支仍未执行。

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

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

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

相关文章：