为什么数组会引起nvlink警告：入口函数的堆栈大小不能静态确定答案

【问题标题】：Why arrays can cause nvlink warning: Stack size for entry function cannot be statically determined为什么数组会引起nvlink警告：入口函数的堆栈大小不能静态确定
【发布时间】：2019-02-05 09:29:20
【问题描述】：

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

我正在处理我的代码中的“无法静态确定入口函数的堆栈大小”警告。经过 CUDA ptxas warnings (Stack size for entry) 和 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 Studio 2017 上使用 CUDA 10.0.130。警告显示从 sm_35 到 sm_75。

我需要帮助，谢谢！

【问题讨论】：

所有这些都在同一个文件中吗？因为我无法以这种方式重现问题。
@RobertCrovella 是的，它们在同一个文件中。而且我可以通过使用联合来摆脱这些警告（但我仍然不知道为什么......），我将编辑我的问题以包含整个文件。
我假设您在项目中打开了“生成可重定位设备代码”选项。当您关闭它时，您是否看到相同的行为？
@RobertCrovella 是的，关闭它时，警告变为 CUDACOMPILE ：ptxas 警告：无法静态确定入口函数 '_Z8TestFuncv' 的堆栈大小
我编辑了很多未使用的代码。请仔细检查我没有删除重要内容。

标签： cuda

【解决方案1】：

对我来说，它看起来像一个错误（意外行为），所以没有回答。我可能错了，但对于那些也遇到这个问题的人来说，没有完美的解决方法。

为什么数组会导致警告，是因为我做错了什么吗？

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

如果我需要使用数组，我可以摆脱警告吗？

使用联合，参见下面的示例。

它们有害吗？

是的，请参见下面的示例。

这是一个例子：

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;
}

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

但是，联合不是一个很好的解决方法：

如果使用联合，请注意对齐 ClassABC。

我希望这个解决方法可以帮助有这个问题的人。而且我仍然怀疑我做错了什么。如果有人知道我做错了什么，请回答这个问题。非常感谢！

【讨论】：

【解决方案2】：

使用 NVCC 10.1.243（和缩短的示例程序） - 我没有收到该警告。你don't get it on GodBolt，也可以。

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

【讨论】：