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

Question

假设我有这样的代码:

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 by Example第11.4节，

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

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

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

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

和

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

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

Cuda by Example中的信息是否已经过时?Talonmies对这个问题的回答如下

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

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

提前感谢您的帮助!

Answer 1

对于在特定设备上使用相同上下文的所有线程来说，页面锁定状态应该是明显的。如果你正在使用运行时API(就像你在这里一样)，那么通常每个设备每个进程只有一个上下文，所以该进程中的所有线程应该在特定设备上共享相同的上下文，并且在该上下文中具有相同的指针视图。

cudaHostAllocPortable标志的一个功能在CUDA文档中有描述:

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

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