cuda - 如何在 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 回退方式吗？

标签： cuda

Answer 1

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

出于上述原因，如果可以对其进行管理，那无疑是最好的方法。

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

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

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

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