CODEWITHSUNDEEP
cudacmake
BRabbit27
BRabbit27

Reputation: 6633

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

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

Upvotes: 5

Views: 7021

Answers (2)

user6932536
user6932536

Reputation: 1

I got the same problems as you and I found two ways to solve it.

  1. define your struct in C-type, like this:

    typedef struct {} ClassName;

  2. define both constructor and destructor as __device__ type, like this:

    struct ClassName{
public:
    __device__ ClassName(){...}
    __device__ ~ClassName(){...}
};

Upvotes: 0

Robert Crovella
Robert Crovella

Reputation: 152249

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

Upvotes: 13

Related Questions