5

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

0

1 Answer 1

4

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.

Sign up to request clarification or add additional context in comments.

Comments

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service and acknowledge you have read our privacy policy.

Start asking to get answers

Find the answer to your question by asking.

Ask question

Explore related questions

See similar questions with these tags.