#include <cstdio>
#include <cuda.h>

__host__ __device__
void where() {
#if defined __CUDA_ARCH__
  printf(" on the device");
#else
  printf(" on the host");
#endif
}

struct HostOnlyType {
  __host__ /* __device__ */
  void call()
  {
    printf("__host__ function");
    where();
    printf("\n");
  }
};

struct DeviceOnlyType {
  /* __host__ */ __device__
  void call()
  {
    printf("__device__ function");
    where();
    printf("\n");
  }
};

struct HostDeviceType {
  __host__ __device__
  void call()
  {
    printf("__host__ __device__ function");
    where();
    printf("\n");
  }
};

template <typename T>
struct Dispatcher {
  __host__    // only if T::call() is a __host__ function ?
  __device__  // only if T::call() is a __device__ function ?
  static void dispatch() {
    T t;
    t.call();
  }
};

template <typename T>
__global__
void kernel(void) {
  Dispatcher<T>::dispatch();
}

int main(void) {
  Dispatcher<HostOnlyType>::dispatch();
//Dispatcher<DeviceOnlyType>::dispatch();
  Dispatcher<HostDeviceType>::dispatch();

//kernel<HostOnlyType><<<1,1>>>();
  kernel<DeviceOnlyType><<<1,1>>>();
  kernel<HostDeviceType><<<1,1>>>();

  cudaDeviceSynchronize();
}