我是否需要在 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);
}

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

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

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

Answer 1

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

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

Answer 2

如果你使用 > Pascal 级别的 gpu，它们有统一的内存。您不需要数据迁移。

当在 CPU 或 GPU 上运行的代码访问以这种方式分配的数据（通常称为 CUDA 托管数据）时，CUDA 系统软件和/或硬件负责将内存页面迁移到访问处理器的内存。

https://devblogs.nvidia.com/unified-memory-cuda-beginners/

如果你使用老派的方式来分配缓冲区（cuMalloc），我认为你确实需要镜像数据。