cuda - 模板 __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;
}