Compiling __host__ __device__ functions

Accelerated Computing CUDA CUDA Programming and Performance
1

I’m trying to understand how host device functions get compiled. My understanding is that two separate versions of the function will be compiled, one as host and one as device. So, I expect the following bit of code to produce two compiler warnings, but compile into two functions, the host version having an additional printf statement:

extern "C"

__host__ __device__ void func()

{

  ...

  #ifndef __CUDA_ARCH__

        #warning __CUDA_ARCH__ Undefined

	printf("Compiler error\n");

  #else

	#warning __CUDA_ARCH__ defined

  #endif

  ...

}

But, there is a compiler error complaining about using printf in device code.

warning: #warning __CUDA_ARCH__ defined

warning: #warning __CUDA_ARCH__ Undefined

error: calling a host function from a __device__/__global__ function is only allowed in device emulation mode

1 error detected in the compilation of "/tmp/tmpxft_000008f9_00000000-10_intEV.cpp4.ii".

To me, it looks like the device version compiles properly without error, then the host version of the function produces an error that should only occur in device or global functions. Shouldn’t printf be fine in the host version of the function?

The CUDA programming guide says these function qualifiers can be used together, and the function will be separately compiled for both the host and the device. Am I thinking about host device functions incorrectly?

2

I have a similar question as well. I am trying to create code that has exception handling on the host side at least and does something else on the device side, but the code doesn’t compile even while implimenting CUDA_ARCH

__host__ __device__ int test(int a, int b)

{

	#ifndef  __CUDA_ARCH__

		if (a > b) throw std::exception("a cannot be less than b");

	#else

		if (a > b) return -1;

	#endif

	return b - a;

}

which generates the following error message:

“error: device code does not support exception handling”