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

Question

如果我的用户有一个兼容 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 回退方式吗？

Answer 1

如果addSingleItem 和addKernel 定义在同一个翻译单元/模块/文件中，则设备到设备的函数调用应该没有成本。编译器会积极地内联该代码，就像您将其编写在单个函数中一样。

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

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

另一种可能的方法是使用编译器宏来协助添加、删除或修改代码，以处理 GPU 情况与非 GPU 情况。这里有无限的可能性，但请参阅here 了解一个简单的想法。例如，您可以在不同的场景中redefine what __host__ __device__ means。我想说这可能只有在您为 GPU 和非 GPU 案例构建单独的二进制文件时才有意义，但您可能会找到一种巧妙的方法来在同一个可执行文件中处理它。

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

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