模板__host__ __device__调用主机定义的函数

Question

在实现CUDA代码期间,我经常需要一些实用程序函数,这些函数将从设备和主机代码中调用.所以我将这些函数声明为__host__ __device__.这是可以的,#ifdef CUDA_ARCH可以处理可能的设备/主机不兼容性.

当效用函数被模板化时出现问题,即.通过一些仿函数类型.如果模板实例调用__host__函数,我会收到此警告:

calling a __host__ function from a __host__ __device__ function is not allowed
      detected during instantiation of "int foo(const T &) [with T=HostObject]" 

我知道的唯一解决方案是定义函数两次 - 一次用于设备,一次用于具有不同名称的主机代码(我不能重载__host__ __device__).但这意味着存在代码重复和所有其他__host__ __device__将调用它的函数,也必须定义两次(甚至更多的代码重复).

简化示例:

#include <cuda.h>
#include <iostream>

struct HostObject {
    __host__ 
    int value() const { return 42; }
};

struct DeviceObject {
    __device__ 
    int value() const { return 3; }
};

template <typename T> 
__host__ __device__ 
int foo(const T &obj) {
    return obj.value();
}

/*
template <typename T> 
__host__ 
int foo_host(const T &obj) {
    return obj.value();
}

template <typename T> 
__device__ 
int foo_device(const T &obj) {
    return obj.value();
}
*/

__global__ void kernel(int *data) {
    data[threadIdx.x] = foo(DeviceObject());
}

int main() {
    foo(HostObject());

    int *data;
    cudaMalloc((void**)&data, sizeof(int) * 64);
    kernel<<<1, 64>>>(data);
    cudaThreadSynchronize();
    cudaFree(data);
}

警告是由函数foo(HostObject());内部的调用引起的main().

foo_host<>并且foo_device<>可能是有问题的替代品foo<>.

有更好的解决方案吗？我可以阻止foo()设备端的异常吗？

Answer 1

您无法阻止实例化__host__ __device__函数模板实例化的任何一半.如果通过在主机(设备)上调用它来实例化该函数,则编译器还将实例化设备(主机)的一半.

从CUDA 7.0开始,您可以为您的用例做的最好的事情是使用#pragma hd_warning_disable以下示例中的方法来抑制警告,并确保未正确调用该函数.

#include <iostream>
#include <cstdio>

#pragma hd_warning_disable
template<class Function>
__host__ __device__
void invoke(Function f)
{
  f();
}

struct host_only
{
  __host__
  void operator()()
  {
    std::cout << "host_only()" << std::endl;
  }
};

struct device_only
{
  __device__
  void operator()()
  {
    printf("device_only(): thread %d\n", threadIdx.x);
  }
};

__global__
void kernel()
{
  // use from device with device functor
  invoke(device_only());

  // XXX error
  // invoke(host_only());
}

int main()
{
  // use from host with host functor
  invoke(host_only());

  kernel<<<1,1>>>();
  cudaDeviceSynchronize();

  // XXX error
  // invoke(device_only());

  return 0;
}