I have a 3d vector class with member functions marked as host and device functions. Below is snippet of one of the member function:
__host__ __device__
double Vector::GetMagReciprocal()
{
double result = 1/sqrt(x*x + y*y + z*z);
return result;
}
What I want to achieve is to have separate definition for host and device function so that I can get better performance by using CUDA math intrinsic function rqsrt when executing on device. The way I would do it is to overload this member function for host and device:
__host__
double Vector::GetMagReciprocal()
{
double result = 1/sqrt(x*x + y*y + z*z);
return result;
}
__device__
double Vector::GetMagReciprocal()
{
double result = rsqrt(x*x + y*y + z*z);
return result;
}
Now when I compile the Vector.cpp file using nvcc(-x cu flag), I get following error
function "Vector::GetMagReciprocal" has already been defined
Now I wonder why NVIDIA doesn't support this sort of overloading.
I can think of alternate ways of achieving the separation, but they have their own issues:
- create separate member functions for host and device in vector class say GetMagReciprocalHost and GetMagReciprocalDevice and call the appropriate function in host/device code
- Have a single member function GetMagReciprocal but pass a flag to the member function to choose between host code and device code
Maybe there is another easier way to achieve this. If someone has any suggestions, it will be nice.
REEDITED: I had not mentioned about possibility of conditional compilation using CUDA ARCH flag to generate separate host and device. This was actually the first thing I had done when modifying the member function. But something came to my mind which said this won't work. Perhaps I was wrong about my understanding of usage of this compilation flag. So the answer suugested by sgarizvi is the right answer
CUDA_ARCH
does solve the problem. You decorate a single function with both__host__
and__device__
, and then use conditional compilation with theCUDA_ARCH
macro to change code generation for device code.