为什么CUDA C++类的主机/设备成员函数不能重载

Question

我有一个3D向量类，其中的成员函数标记为主机和设备函数。以下是其中一个成员函数的片段：

__host__ __device__
double Vector::GetMagReciprocal()
{
    double result = 1/sqrt(x*x + y*y + z*z);
    return result;
}

我想实现的是分别为主机和设备函数定义不同的函数，这样我就可以在设备上执行时使用CUDA数学内置函数rqsrt，从而获得更好的性能。我会通过重载此成员函数来实现主机和设备的区分：

__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;
}

现在，当我使用nvcc（-x cu标志）编译Vector.cpp文件时，我会收到以下错误信息：

“函数“Vector :: GetMagReciprocal”已经定义”

现在我想知道为什么NVIDIA不支持这种重载方式。

我可以想到其他方式来实现分离，但它们也有自己的问题:

1. 在向量类中创建单独的主机和设备成员函数，比如 GetMagReciprocalHost 和 GetMagReciprocalDevice，并在主机/设备代码中调用适当的函数。 2. 有一个单一的成员函数GetMagReciprocal，但是将标志传递给成员函数以在主机代码和设备代码之间选择。

也许有另一种更简单的实现方法。如果有人有任何建议，那就太好了。

修订：我没有提及使用CUDA ARCH标志进行条件编译以生成单独的主机和设备的可能性。这实际上是我修改成员函数时做的第一件事。但是我想到了一些东西，说这不起作用。也许我对使用这个编译标志的理解是错误的。因此，sgarizvi建议的答案是正确的。

- nurabha

---

CUDA_ARCH确实解决了这个问题。您可以使用__host__和__device__两个修饰符装饰一个函数，然后使用CUDA_ARCH宏进行条件编译，以改变设备代码的代码生成方式。 - Robert Crovella

这是我做的第一件事。但是我觉得它似乎不起作用。我必须测试一下，看看是否如此。 - nurabha

如果您提供一个简短、完整的代码，别人可以复制、粘贴、编译和运行，并查看问题所在，我相信它可以被解释清楚。但是一些小的代码片段连同断言说“这个不起作用”，是无法为您排除故障的。如果这个问题不是因为重复而关闭的话，在我看来，也可能会因为缺乏MCVE而关闭。请确保将来提供一个MCVE。 - Robert Crovella

Answer 1

您可以使用条件编译标志__CUDA_ARCH__在__host__ __device__函数中为主机和设备生成不同的代码。 __CUDA_ARCH__仅在设备代码中定义，因此要为主机和设备创建不同的实现，您可以执行以下操作：

__host__ __device__
double Vector::GetMagReciprocal()
{
    double result;
    #ifdef __CUDA_ARCH__
    result = rsqrt(x*x + y*y + z*z);
    #else
    result = 1/sqrt(x*x + y*y + z*z);
    #endif
    return result;
}

- sgarizvi

---

我总是使用 nvcc 进行编译，以便为主机和设备函数生成相同的代码。但是它并没有起作用。 - nurabha

这绝对是正确答案，任何声称它不可行的说法都是完全错误的。 - talonmies

@nurabha 不，它不会为主机和设备生成相同的代码。 nvcc 不是编译器。它是一个编译器驱动程序，将代码传递到后端的实际主机和设备编译器。代码的两个版本都将被生成。一个由主机编译器生成，另一个由设备编译器生成。 - sgarizvi

@sgarizvi：让我在会话中执行一步调试代码，以便我可以验证。我相信你的答案是正确的。实际上，在此之前，我已经在做你所建议的条件编译，但后来我又看了一遍，似乎它不会起作用。 - nurabha

@sgarizvi：嗯，一个月前我改变了整个向量类，并在所有我想要单独的设备/主机定义的成员函数中特别使用了条件标志。我想我只是有点困惑了。谢谢你的回答。如果没有复习我的CUDA概念，我本来就不应该首先问这个愚蠢的问题。 - nurabha