CUDA 2.0 Beta and device array

Hello.

I tested my CUDA application on CUDA 1.1 version and it compiled and worked flawlessly. Yesterday I upgraded to CUDA 2.0 Beta and some code ceased to compile.

Please, consider the following code:

__device__ __constant__ const int K = 7;

__device__ __constant__ const size_t dec_bits = 192;

__device__ __constant__ const UINT32 nStates = (1<<(K-1));

__device__ UINT32 decoder_matrix[dec_bits][nStates];

and the compiler output is (line 132 is the last line of the code above):

1>file.cu(132) : error C2057: expected constant expression

1>file.cu(132) : error C2466: cannot allocate an array of constant size 0

1>file.cu(132) : error C2057: expected constant expression

1>file.cu(132) : error C2466: cannot allocate an array of constant size 0

1>file.cu(132) : error C2087: '__shadow_decoder_matrix' : missing subscript

1>file.cu(132) : error C2133: '__shadow_decoder_matrix' : unknown size

It compiled on CUDA 1.1 and does not on CUDA 2.0 Beta. What am I missing?

P.S.: Please note, that decoder_matrix array is written from inside the kernel and could not be declared as constant.

constant tells the compiler where to allocate the variable; it’s not a guarantee it won’t change. Presumably the 1.1 compiler was confused about this, and the 2.0 compiler isn’t.

Try this:

#define K_VALUE 7

#define DEC_BITS 192

#define NSTATES (1<<(K_VALUE-1))

__device__ __constant__ const int K = K_VALUE;

__device__ __constant__ const size_t dec_bits = DEC_BITS;

__device__ __constant__ const UINT32 nStates = NSTATES;

__device__ UINT32 decoder_matrix[DEC_BITS][NSTATES];

Thanks for reply. The code you suggested works. But I still can not understand why I’m not able to use constants (variables declared with ‘const’ specifier) in declaration of fixed size array. Wherever variable declared ‘const’ it shall not change after initialization.