c++ - Cuda:将设备常量声明为模板

Question

我正在开发一个使用模板的 CUDA 程序。内核将使用数据类型 cuComplex 或 cuDoubleComplex 进行实例化。根据实例化内核的数据类型，我需要声明一个驻留在 CUDA 设备的常量内存空间中的常量。为了实现这一点，我做了:

// declared this globally
template <typename T> 
__device__ __constant__ T some_constant;
// set the constant through a function call in main().

// kernel templated
template <typename T>    
__global__ void kernel()
{
    T b_k = cuGet<T>(0.0, 0.0);
    T x_k_1 = cuGet<T>(2.0, 2.0); 
    // cuGet returns a complex no. of type 
    // cuComplex or cuDoubleComplex depending on T.

    b_k = cuAdd(b_k, cuMul(some_constant, x_k_1));
    // cuAdd, cuMul, cuGet are all overloaded functions.
    // They can take cuComplex, or cuDoubleComplex params.

    // Here, some_constant has to cuComplex or cuDoubleComplex, depending 
    // on the datatype of the other arg x_k_1 to cuMul.
    // Therefore, I went about implementing a templated constant.
}

编译时，会出现错误:“some_constant”不是函数或静态数据成员。

解决这个问题的一个选择是定义从 cuDoubleComplex 到 cuComplex 的类型转换，并将常量声明为 cuDoubleComplex而不是将其用作模板，而是在内核中使用常量的任何地方对其进行类型转换。

除了这个还有别的方法吗？

提前致谢。

Answer 1

如果主要要求是避免类型转换，您可以使用具有内联设备函数的模板类来解决它(其余的灵感来自@RobertCrovella 在评论中提出的建议)。最后，一个宏将进行调用(这确实不是一个非常干净的设计，但保持相同的语法)。这是它如何工作的示例:

template <typename T> 
struct holder
{
    static __device__ __inline__ T value () ;
    static __device__ __inline__ void init (T val) ;
} ;

__constant__ char data [16] ;

__device__ __inline__ int holder<int>::value () { return *((int*)data); }
__device__ __inline__ long holder<long>::value () { return *((long*)data); }

#define some_constant holder<T>::value()

template <typename T>
__global__ void kernel(T* res)
{
    *res = some_constant ;
}

int main ()
{
    int *dres ;
    cudaMalloc <> (&dres, sizeof(int)) ;
    int val = 42 ;
    cudaMemcpyToSymbol (data, &val, sizeof(int)) ;

    kernel<int><<<1,1>>>(dres) ;
    int hres ;
    cudaMemcpy (&hres, dres, sizeof(int), cudaMemcpyDeviceToHost) ;
    printf ("RES = %d\n", hres) ;
}

holder<T>::value()调用将被内联，类型转换被优化器删除，并从常量内存中返回适当的类型而不进行转换(这里是生成的 ptx):

// .globl   _Z6kernelIiEvPT_
.const .align 4 .b8 data[16];

.visible .entry _Z6kernelIiEvPT_(
    .param .u32 _Z6kernelIiEvPT__param_0
)
{
    .reg .b32   %r<4>;


    ld.param.u32    %r1, [_Z6kernelIiEvPT__param_0];
    cvta.to.global.u32  %r2, %r1;
    ld.const.u32    %r3, [data];
    st.global.u32   [%r2], %r3;
    ret;
}

主要缺点是期望类型为 T 的宏.