constexpr array in device code

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello,

could you please tell me, is there any way to use constexpr arrays in device code ? According to “Cuda C programming guide 7.0” I have no problems with constexpr scalars, but arrays seem do not compile. Below some example:

template<unsigned D, unsigned Q>
class LatticeArrangement 
{
} ;

template<>
class LatticeArrangement<3,19> 
{
	public:
		static constexpr double c[19] = { 0,1,2,3,4,5,6,7,8,9,10,
																			11,12,13,14,15,16,17,18 } ;

		static constexpr double d = 19.0 ;

		__host__ __device__ 
		static constexpr double getC( unsigned index ) 
		{
			// FIXME: error: identifier "LatticeArrangement<(unsigned int)3u, (unsigned int)19u> ::c" is undefined in device code 
			return c[ index ] ; 

			//return d * index ; // OK, this one works
		} ;
} ;

constexpr double LatticeArrangement<3,19>::c[] ;

template< class LatticeArrangement >
class FluidModelIncompressible
{
	public:
		__host__ __device__ 
		static double computeSomething(double v, unsigned idx)
		{
			return LatticeArrangement::getC( idx ) * v ;
		}
} ;

// Does nothing useful, we want only to compile this
template< class FluidModel >
__global__ void
kernel1 ( double * data )
{
	data[ threadIdx.x ] = FluidModel::computeSomething( threadIdx.y, threadIdx.z ) ;
}

int main( int argc, char ** argv )
{
	dim3 numBlocks  ( 2 ) ;
	dim3 numThreads ( 4, 4, 4 ) ;	

	double * vptr = NULL ;
	
	kernel1< FluidModelIncompressible< LatticeArrangement<3,19> > > 
		<<< numBlocks, numThreads >>> 	( vptr ) ;
	
	return 0 ;
}

Maybe there is some other way to avoid duplication of code between host and device ?
I am using
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2015 NVIDIA Corporation
Built on Mon_Feb_16_22:59:02_CST_2015
Cuda compilation tools, release 7.0, V7.0.27

Thanks in advance :)