CUDA 设备代码中的 constexpr 数组答案

【问题标题】：constexpr array in CUDA device codeCUDA 设备代码中的 constexpr 数组
【发布时间】：2015-09-09 19:16:34
【问题描述】：

你能告诉我，有没有办法在设备代码中使用 constexpr 数组？根据“Cuda C 编程指南 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 表达式的编译器优化。也许还有其他方法可以避免主机和设备之间的代码重复？目前我在设备代码中有一个大案例。

提前谢谢:)

【问题讨论】：

标签： c++ arrays cuda constexpr

【解决方案1】：

您能告诉我，有什么方法可以在设备代码中使用 constexpr 数组吗？

CUDA 7.0 不支持。

如果您希望得到支持，我建议您在 NVIDIA CUDA developer portal 提交 RFE（错误）。

【讨论】：