c++ - 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 表达式的编译器优化。也许还有其他方法可以避免主机和设备之间的代码重复？目前我有一个很大的设备代码案例。

我正在使用nvcc:NVIDIA (R) Cuda 编译器驱动程序版权所有 (c) 2005-2015 NVIDIA 公司建立于 Mon_Feb_16_22:59:02_CST_2015Cuda编译工具，release 7.0，V7.0.27

提前致谢:)

Answer 1

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

CUDA 7.0 不支持它。

如果您希望得到支持，我建议在 NVIDIA CUDA developer portal 提交 RFE(错误) .