页面锁定状态是否对使用同一设备的所有主机线程可见？答案

【问题标题】：Is pagelocked status visible to all host threads that use the same device?页面锁定状态是否对使用同一设备的所有主机线程可见？
【发布时间】：2017-01-23 22:03:38
【问题描述】：

假设我的代码如下所示：

cudaHostAlloc( (void**)&pagelocked_ptr, size, cudaHostAllocDefault )
#pragma omp parallel num_threads(num_streams)
{ 
  ...
  cudaMemcpyAsync( pagelocked_ptr + offset_thisthread
                 , src
                 , count
                 , kind
                 , stream_thisthread );
  ...
}

请注意，我明确避免在此处设置标志 cudaHostAllocPortable。每个线程都使用自己的流，并且（我相信）隐式选择默认的 Cuda 设备。

根据 Cuda 示例第 11.4 节，

页面只能显示为固定到单个 CPU 线程。也就是说，如果任何线程将它们分配为固定内存，它们将保持页面锁定，但它们只会对分配它们的线程显示为页面锁定。

他们接着说设置cudaHostAllocPortable 可以解决这个问题并允许所有线程将分配识别为固定缓冲区。因此，除非我指定cudaHostAllocPortable 而不是cudaHostAllocDefault，否则我上面的cudaMemcpyAsync 调用将失败。

Cuda C 指南似乎与此信息相冲突。我的假设是 Cuda 上下文会跟踪主机内存的哪些区域是页面锁定的，并且可以在没有中间暂存副本的情况下传输到设备。根据当前的 Cuda C 指南 3.2.1 和 3.2.4.1

此设备的主要上下文...在应用程序的所有主机线程之间共享。

和

默认情况下，上述使用页面锁定内存的好处仅适用于分配块时的当前设备（并且所有设备共享相同的统一地址空间，如果有的话......）

这些似乎暗示分配的页面锁定性质由跨不同线程的 Cuda 调用知道，因为它们都使用设备 0，并且所有线程中对 cudaMemcpyAsync() 的调用都会成功。换句话说，如果我的解释正确，那么设置 cudaHostAllocPortable 仅在尝试在 Cuda 上下文之间共享页面锁定内存时才是必要的（例如，当使用cudaSetDevice 在 GPU 之间切换并卸载一大块页面时） -锁定分配给每个）。

Cuda by Example 中的信息是否已经过时？ Talonmies 对this question states 的回复

在 CUDA 4 之前，上下文不是线程安全的，需要通过上下文迁移 API 显式迁移。

但我不确定这如何影响页面锁定状态对来自不同线程的 Cuda 调用的可见性。

提前感谢您的帮助！

【问题讨论】：

标签： cuda

【解决方案1】：

对于在特定设备上使用相同上下文的所有线程来说，页面锁定状态应该是显而易见的。如果您使用的是运行时 API（就像您在这里一样），那么每个进程的每个设备通常只有一个上下文，因此该进程中的所有线程应该在特定设备上共享相同的上下文，并且对任何指针具有相同的视图在这种情况下。

cudaHostAllocPortable 标志的功能之一在the CUDA documentation 中有描述：

此调用返回的内存将被所有 CUDA 上下文视为固定内存，而不仅仅是执行分配的那个。

这意味着在多上下文设置或多设备设置中（上下文对于特定设备是唯一的），有必要使用此标志在进程可见的所有上下文中从该指针获取固定行为.

【讨论】：