如何在 CUDA 中为 CPU 回退重用代码

How to reuse code for CPU fallback in CUDA

如果我的用户有兼容 CUDA 的 GPU，我有一些想要并行化的计算，否则我想在 CPU 上执行相同的代码。我不想有两个版本的算法代码，一个用于CPU，一个用于GPU维护。我正在考虑以下方法，但想知道额外的间接级别是否会影响性能，或者是否有更好的做法。

对于我的测试，我采用了基本的 CUDA 模板，该模板将两个整数数组的元素相加并将结果存储在第三个数组中。我删除了实际的加法操作并将其放入自己的函数中，同时标有设备和主机指令...

__device__ __host__ void addSingleItem(int* c, const int* a, const int* b)
{
    *c = *a + *b;
}

...然后修改内核以在threadIdx标识的元素上调用上述函数...

__global__ void addKernel(int* c, const int* a, const int* b)
{
    const unsigned i = threadIdx.x;
    addSingleItem(c + i, a + i, b + i);
}

现在我的应用程序可以检查是否存在 CUDA 设备。如果找到我可以使用...

addKernel <<<1, size>>> (dev_c, dev_a, dev_b);

...如果不是，我可以放弃并行化并遍历调用函数主机版本的元素...

int* pA = (int*)a;
int* pB = (int*)b;
int* pC = (int*)c;

for (int i = 0; i < arraySize; i++)
{
    addSingleItem(pC++, pA++, pB++);
}

我的小型测试应用程序似乎一切正常，但我担心涉及的额外调用。设备到设备的函数调用是否会导致任何显着的性能下降？有没有更普遍接受的方法来做我应该采用的 CPU 回退？

如果 addSingleItem 和 addKernel 在同一个翻译 unit/module/file 中定义，进行设备到设备的函数调用应该没有成本。编译器会积极地内联该代码，就像您将它写在一个函数中一样。

如果可以管理，这无疑是最好的方法，原因如上所述。

如果仍然需要一些文件级模块化，可以将代码分解成一个单独的文件，并将该文件包含在内核函数的编译中。从概念上讲，这与已经描述的内容没有什么不同。

另一种可能的方法是使用编译器宏来协助添加、删除或修改代码，以处理 GPU 情况与非 GPU 情况。这里有无限的可能性，但例如在不同的场景中看到here for a simple idea. You can 。我想说这可能只有在为 GPU 和非 GPU 情况构建单独的二进制文件时才有意义，但您可能会找到一种聪明的方法来在同一个可执行文件中处理它。

最后，如果您需要这个但必须将 __device__ 函数放在一个单独的翻译单元中，这仍然是可能的，但由于跨模块的设备到设备函数调用可能会有一些性能损失界限。这里的性能损失量很难一概而论，因为它在很大程度上取决于代码结构，但看到 10% 或 20% 的性能损失并不罕见。在这种情况下，您可能希望调查 link-time-optimizations that became available in CUDA 11.

也可能很有趣，尽管这里只是切线相关。