C++ design for CUDA codes

Question

I have a piece of C++ CUDA code which I have to write declaring the data variable in float. I also have to rewrite the code declaring the data variable in double.

What is a good design to handle a situation like this in CUDA?

I do not want to have two sets of same code because then in the future for any change I will have to have to change two sets of otherwise identical code. I also want to keep the code clean without too many #ifdef to change between float and double within the code.

Can anyone please suggest any good (in terms of maintenance and "easy to read") design?

score 6 · Accepted Answer · answered Jan 20 '13 at 18:28

CUDA supports type templating, and it is without doubt the most efficient way to implement kernel code where you need to handle multiple types in the same code.

As a trivial example, consider a simple BLAS AXPY type kernel:

template<typename Real>
__global__ void axpy(const Real *x, Real *y, const int n, const Real a)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;

    for(; tid<n; tid += stride) {
        Real yval = y[tid];
        yval += a * x[tid];
        y[tid] = yval;
    }
}

This templated kernel can be instantiated for both double and single precision without loss of generality:

template axpy<float>(const float *, float *, const int, const float);
template axpy<double>(const double *, double *, const int, const double);

The thrust template library, which ships with all recent versions of the CUDA toolkit, makes extensive use of this facility for implementing type agnostic algorithms.

wow!! i did not know this existed. thank you very much. is there any tutorial kind of document i can look at. — user1612986, Jan 20 '13 at 18:36
There is a brief discussion of C++ features supported in CUDA in the Appendices of the programming guide. — talonmies, Jan 20 '13 at 18:49
Also see [this question](http://stackoverflow.com/a/6179580/681865) for another very useful template based CUDA design pattern. — talonmies, Jan 20 '13 at 19:05

score 2 · Answer 2 · answered Jan 20 '13 at 18:35

2

In addition to templating, you may be able to achieve what you want with a single typedef:

typedef float mysize;  // or double

Then just use mysize throughout where you would use float or double.

You might be interested in the simpleTemplates sample code, and there are other templatized CUDA examples as well, in addtion to thrust where, as talonmies states, it's used extensively. Thrust provides many other benefits as well to C++ programmers.

answered Jan 20 '13 at 18:35

Robert Crovella

143,785
11
213
257

That can also be a perfectly good option, although it precludes being able to use both versions within the same compilation unit. – talonmies Jan 20 '13 at 18:58

C++ design for CUDA codes

2 Answers2