constexpr array in CUDA device code

Question

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

I would like to use the same code on host and device and at the same time benefit from compiler optimisations of constexpr expressions. Maybe there is some other way to avoid duplication of code between host and device ? Currently I have a large case in device code.

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

Answer 1

could you please tell me, is there any way to use constexpr arrays in device code ?

It is not supported in CUDA 7.0.

If you desire this to be supported, I suggest filing an RFE (bug) at the NVIDIA CUDA developer portal.