Geek Answers Handbook

Cuda: declaring a device constant as a template

I am working on a CUDA program that uses templates. Kernels will be created with the data type cuComplexor cuDoubleComplex. Depending on the type of data with which the kernel instance is created, I need to declare a constant that will be in the constant memory space of the CUDA device. To implement this, I did:

// 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.
}

When compiling, this gives an error: "some_constant" is not a function or element of static data.

One solution to this problem is to define a type conversion from cuDoubleComplexto cuComplexand declare a constant in cuDoubleComplexinstead of using it as a template, and resorting to creating a constant wherever it is used in the kernel.

Is there any other way from this?

.

+4
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.

+4

More articles:


All Articles