CUDA设备代码中的constexpr数组

Question

请问，有没有办法在设备代码中使用“constexpr”数组？根据“Cuda C programming guide 7.0”，我对“constexpr”标量没有问题，但似乎无法编译数组。以下是一个例子：

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

我希望能够在主机和设备上使用相同的代码，并同时从constexpr表达式的编译器优化中受益。也许还有其他方法可以避免在主机和设备之间重复编写代码？目前，在设备代码中有一个大的case。

我正在使用nvcc：NVIDIA（R）Cuda编译器驱动程序，版权所有（c）2005-2015 NVIDIA Corporation，构建于Mon_Feb_16_22:59:02_CST_2015，CUDA编译工具，版本7.0，V7.0.27。

提前感谢您的帮助 :)

- tadtm

Answer 1

请问，有没有办法在设备代码中使用constexpr数组？

CUDA 7.0不支持此功能。

如果您希望支持此功能，请建议在NVIDIA CUDA开发者门户网站上提交RFE（错误报告）。

- Robert Crovella