数组可导致nvlink警告的原因:无法静态确定条目函数的堆栈大小

Question

我正在处理由数组引起的“入口函数的堆栈大小不能被静态确定”的警告，我需要帮助。

我正在处理代码中的“入口函数的堆栈大小不能被静态确定”警告。CUDA ptxas警告(输入堆栈大小)和https://devtalk.nvidia.com/default/topic/524712/a-meaning-of-nvlink-warning-stack-size-for-entry-function-cannot-be-statically-determined/的警告是由递归引起的。

但是，我没有在代码中找到递归，相反，我发现结构数组也会导致这样的警告。

这个问题可以用一个简单的例子来说明。(编辑:我可以通过使用union来消除这些警告，但我仍然不知道为什么。这些代码位于同一个.cu文件中)

class ClassABC {
public:
    __host__ __device__ ClassABC() { ;  }
    int m_iValue;
};

class ClassDEF {
public:
    __host__ __device__ ClassDEF() { ; }

    //Witout warning
    //union 
    //{
    //    ClassABC m_abc[1];
    //    int m_values[1];
    //};

    //With warning
    ClassABC m_abc[1];
};

__global__ void TestFunc()
{
    ClassDEF def[1];
}

int main()
{
    TestFunc << <1, 1 >> > ();
    return 0;
}

它有一个警告：

CUDALINK : nvlink warning : Stack size for entry function '_Z8TestFuncv' cannot be statically determined (target: sm_(35-75))

所以，我的问题是，为什么数组会引起警告，是因为我做错了什么吗？如果我需要使用数组，我能去掉警告吗？它们有害吗？

我在Windows 10和Visual 2017上使用CUDA 10.0.130。警告从sm_35显示到sm_75。

我需要帮助，谢谢！

Answer 1

在我看来，它就像一个bug (意外的行为)，所以它没有得到回答。我可能错了，但是对于那些也遇到这个问题的人来说，这不是一个完美的工作。

• 为什么数组会引起警告，是因为我做错了什么吗？

我不知道。我希望我做错了什么，但我认为这可能是一个错误的库达10.0.130。

• 如果我需要使用数组，我能去掉警告吗？

使用union，参见下面的示例。

• 它们有害吗？

是的，见下面的例子。

这是一个例子：

class ClassABC
{
public:
    __host__ __device__ ClassABC():m_iValue(0){ ;  }
    __device__ void Add(int v)
    {
        m_iValue += v;
    }
    __device__ void DebugPrint() const
    {
        printf("v=%d;", m_iValue);
    }
    int m_iValue;
};

class ClassDEF
{
public:
    __host__ __device__ ClassDEF() { ; }

    __device__ void Add(int v)
    {
        m_abc[10].Add(v);
        //m_values[10] += v; also work
    }

    __device__ void DebugPrint() const
    {
        m_abc[10].DebugPrint();
    }
    //Witout warning
    union 
    {
        ClassABC m_abc[20];
        int m_values[20];
    };

    //With warning
    //Output:
    //ClassABC m_abc[20];
};

__global__ void TestFunc()
{
    ClassDEF def[100];

    for (int i = 0; i < 100; ++i)
    {
        def[i].Add(i);
        def[i].DebugPrint();
    }
}

int main()
{

    //If use the version with warning, must set stack size, or there will be a stackoverflow.
    //checkCudaErrors(cudaDeviceSetLimit(cudaLimitStackSize, 1 << 16));
    TestFunc << <1, 1 >> > ();
    checkCudaErrors(cudaDeviceSynchronize());
    return 0;
}

首先，它是有害的，如果不手动增加堆栈大小，可能会导致堆栈溢出。工会将解决这个问题。

然而，不是是一个很好的工作围绕：

• 如果使用union，请小心对齐ClassABC。

我希望这个工作能帮助那些也有这个问题的人。我仍然怀疑我做错了什么。如果有人知道我做错了什么，请回答这个问题。非常感谢!

Answer 2

与NVCC 10.1.243 (和缩短的示例程序)-我没有收到警告。你也是别在GodBolt上看到。

所以，可能是10.0版本的问题，或者是特定设置的问题。