Why am I getting dynamic initialization not supported for __device__, __constant__, __shared__?

Question

I don't understand why am I getting the error dynamic initialization is not supported for __device__, __constant__, __shared__ variables when compiling my code. My code looks like

wrapper.cu

#include "../Params.hpp"

__constant__ Params cparams;

void wrapperFunction(uint& a)
{
   Params ab;
   a = 20;
}

Params.hpp

#include "Utils.hpp"

typedef struct Params
{
   vectorTypef a;
} Params;

Utils.hpp

#include "Vec2.hpp"

typedef unsigned int uint;
typedef Vec2<float> vectorTypef;

Vec2.hpp

template <typename T>
class Vec2
{
public:
   Vec2(){ x = 0.0; y = 0.0;}
   T x, y;
};

Building with cmake with the command

CUDA_ADD_EXECUTABLE(test main.cpp cudasrc/wrapper.cu

Answer 1

Your Params struct is used in the __constant__ memory definition of cparams.

Your Params struct contains an element a of type vectorTypef which is a typedef for the Vec2 class for float. This class has a default constructor, that is assigning elements ultimately of the Params struct. This method of assigning data to a __constant__ region is not legal either in device code or host code.

In device code it's not legal to modify a __constant__ value at all. In host code (which is what is in view here), __constant__ values should be assigned using the appropriate API, i.e. cudaMemcpyToSymbol. I would recommend that you assign these in your host code explicitly, rather than via a constructor.

So, one possible approach to fix this would be to change your default constructor to an empty one:

public:
   __host__ __device__ Vec2(){ }; // change this line
   T x, y;

(you could also just delete the empty default constructor line)

And, in wrapper.cu (perhaps in wrapperFunction), initialize your Params __constant__ struct:

Params hparams;
hparams.a.x = 0.0;
hparams.a.y = 0.0;
cudaMemcpyToSymbol(cparams, &hparams, sizeof(Params));