Why arrays can cause nvlink warning: Stack size for entry function cannot be statically determined

Question

I'm dealing with the "Stack size for entry function cannot be statically determined" warnings caused by arrays, and I need help.

I'm dealing with the "Stack size for entry function cannot be statically determined" warnings in my code. By CUDA ptxas warnings (Stack size for entry) and https://devtalk.nvidia.com/default/topic/524712/a-meaning-of-nvlink-warning-stack-size-for-entry-function-cannot-be-statically-determined/ The warning is caused by recursion.

However, I failed to find recursion in my code, instead, I find that the structure arrays will also cause such warning.

The problem can be shown with a simple example. (Edit: I'm able to get rid of those warnings by using union, but I still don't know why. Those code are in a same .cu file)

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

It has the warning:

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

So, my question is, why the arrays can cause the warning, is it because I did something wrong? If I need to use arrays, can I get rid of the warning? Are they harmful?

I'm using CUDA 10.0.130 on Windows 10, and Visual Studio 2017. The warning show up from sm_35 to sm_75.

I need help, thank you!

Answer 1

To me, it looks like a bug (unexpected behavior), so it is not answered. I might be wrong, but there is a not perfect work around for those who also meet this problem.

• Why the arrays can cause the warning, is it because I did something wrong?

I don't know. I hope I did something wrong, but I think it is possible a bug of cuda 10.0.130.

• If I need to use arrays, can I get rid of the warning?

Use union, see the example below.

• Are they harmful?

Yes, see the example below.

This is the example:

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

First of all, it is harmful, it possibly cause stack overflow if the stack size is not manually increased. And the union will solve this problem.

However, union is NOT a good work around:

• Be careful to align ClassABC if using union.

I hope this work around can help someone who also have this problem. And I still suspect that I did something wrong. If anyone know what I did wrong, please answer this question. Thanks a lot!

Answer 2

With NVCC 10.1.243 (and the shortened example program) - I don't get that warning. You don't get it on GodBolt, either.

So, maybe it was something with version 10.0 or an issue with your specific setup.