How to easily switch between single and double precision in CUDA?

Question

When debugging developing and debugging, I would like to run my code with double precision. However, once I know it's working, I'd like the option to run my code using single precision (i.e. floats) only. So I'd like to be able to easily switch between these precisions without a large rewrite of code. I was thinking of some #define flag like

#define PRECISION double
...

thrust::device_vector<PRECISION> myVec;

but the compiler doesn't seem to like that. Any ideas on how to implement this?

I know this question is very similar in that it solves the problem with compiler flags. But I was hoping to be able to directly set a flag directly from within my source code.

Answer 1

You can do it like this:

#ifdef MY_USE_DOUBLE_PRECISION
typedef double Real;
#else
typedef float Real;
#endif

....

thrust::device_vector<Real> myVec;

using MY_USE_DOUBLE_PRECISION to control the definition of the floating point type Real. If you have your own kernels you can also use Real in place of either float or double ie. :

__global__ void kernel (Real *input, Real *output)
{
   ...
}

If you want to have both single and double precision versions of the kernel code compiled and select which one to use outside of the compilation unit where they are defined (in a library, for example), you can template the kernel:

 template<typename T>
__global__ void kernel (T *input, T *output)
{
   ...
}

template __global__ void kernel<float>(float *, float *);
template __global__ void kernel<double>(double *, double *);

and then in another source file

#ifdef MY_USE_DOUBLE_PRECISION
typedef double Real;
#else
typedef float Real;
#endif

....

kernel<Real><<<griddim, blockdim>>>(....);