c++ - 我是否需要在 CUDA 中跨多个 GPU 镜像输入缓冲区/纹理？

Question

TL;DR:在使用 CUDA 进行多 GPU 编程时，我是否需要跨多个设备镜像只读查找纹理和输入缓冲区(无论是严格要求还是为了获得最佳性能)？

我有一个 GPU 内核，它接受两个用于查找的纹理和两个(较小的)输入数据缓冲区。

我扩展了我的代码以允许多个 GPU(我们的系统最多有 8 个，但为了测试我在一个较小的开发系统上只使用 2 个)。我们的系统使用 NVLINK 并且启用了 UVA。

我的设置涉及使设备 0 成为一种“主”或“根”设备，其中存储最终结果并发生最终串行(串行，因为只能在一个 GPU 上执行)操作。所有设备都设置为允许对 dev 0 进行对等访问。内核在每个设备上以以下形式循环调用多次:

for(unsigned int f = 0; f < maxIterations; f++)
{
    unsigned int devNum = f % maxDevices; //maxIterations >> maxDevices
    cudaSetDevice(devNum);
    cudaDeviceSynchronize(); //Is this really needed?
    executeKernel<<<>>>(workBuffers[devNum], luTex1, luTex2, inputBufferA, inputBufferB);
    cudaMemcpyAsync(&bigGiantBufferOnDev0[f * bufferStride],
                     workBuffers[devNum],
                     sizeof(float) * bufferStride,
                     cudaMemcpyDeviceToDevice);
}

可以看到，每个设备都有自己的“工作缓冲区”用于写出中间结果，然后将这些结果存储到设备 0。

工作(输出)缓冲区的大小比输入缓冲区大几个数量级，我注意到当我犯了一个错误并跨设备访问缓冲区时，性能受到了重大影响(大概是因为内核正在访问另一台设备上的内存)。然而，在修复输出缓冲区问题后，我没有注意到只读输入缓冲区有类似的问题。

这让我想到了我的问题:我是否真的需要跨设备镜像这些输入缓冲区和纹理，或者是否有一种缓存机制可以让这变得不必要？为什么我在跨设备访问工作缓冲区时注意到如此巨大的性能下降，但似乎对输入缓冲区/纹理没有这样的影响？

Answer 1

Texturing ，以及普通的全局数据访问，如果您启用了对等访问，则可以“远程”。由于此类访问是通过 NVLink(或具有对等功能的结构)进行的，因此它通常会更慢。

对于“小型”输入缓冲区，GPU 缓存机制可能会减少或减轻与远程访问相关的惩罚。 GPU 具有特定的只读缓存，旨在帮助处理只读/输入数据，当然纹理机制也有自己的缓存。除非使用实际代码进行实际分析，否则不可能有详细的性能陈述。