CUDA 统一内存编程 - 我可以将指针传递给不知道 CUDA 的模块吗？

Question

我会用cudaMallocManaged 分配一些内存。我可以安全地将这个指针传递给另一个不知道 CUDA 的程序模块（它被编译成另一个 .o 文件），并且只会使用普通的 memcpy 或其他东西来操作指针后面的数据吗？

类似的东西

// compiled into A.o
class A{
  void* getMem(int size){
    void* ptr;
    cudaMallocManaged(*ptr, size);
    return ptr;
  }

  // some kernels here

}


// compiled into B.o
class B{
  void manipulateMem(void* ptr, void* source, int size){
    memcpy(ptr, source, size);
  }

}

然后是这样的代码，可能会编译到 main.o：

A a;
B b;
void* mem = a.getMem(100);
b.manipulateMem(mem, source, 100);

我没有发现自动复制/同步在这种情况下不起作用的任何通知。

Answer 1

我可以安全地将这个指针传递给另一个不知道 CUDA 的程序模块（它被编译成另一个 .o 文件），并且只会使用普通的 memcpy 或其他东西来操作指针后面的数据吗？

是的，您可以，但仍必须遵守Unified Memory access rules（当前是：

1. 最初，在分配时（使用托管分配器，例如 cudaMallocManaged），在调用任何内核之前，可以从主机 CPU（运行的代码）访问托管指针。
2. 一旦调用了任何内核，直到调用了后续的cudaDeviceSynchronize()，主机代码中的数据无法访问，任何尝试在主机代码中使用它都会导致 UB，这可能包括段错误。
3. 在同步事件（例如cudaDeviceSynchronize()）之后，主机对托管指针引用的数据的访问被恢复。 （从技术上讲，在当前的实现中，此时主机访问数据通常会导致页面错误，这是一个操作系统可见的事件。这些页面错误基本上会调用 CUDA 运行时，然后在hood，将数据返回给主机，服务页面错误。但括号中的这些 cmets 对于理解一般行为规则不是必需的。）

因此，CUDA 运行时基本上具有明确的标记（内核启动 -> 同步），这些标记明确地向它指明如何管理数据（何时迁移以及迁移的方向）。因此，即使您在某些“模块......不知道 CUDA”中运行代码，如果遵守上述规则，该模块中的代码也将有权访问数据，因为 CUDA 运行时有足够的、显式的标识它如何管理数据的标记。

对于计算能力 6.0 或更高的 GPU，上述规则基本上不适用。对于这些 GPU，主机和设备的并发访问 是可能的，但是对公共资源的多次访问仍然存在竞争条件的可能性，就像在任何多处理器/多线程环境中一样。 目前，CUDA 运行时不会在主机和设备访问同一页内存之间强制执行任何特定的访问顺序规则。

Answer 2

是的，CUDA 统一内存模型明确允许这种访问：GPU 不仅可以直接访问统一内存指针。同样，在主机 CPU 上运行的部分程序将具有统一的访问权限，包括 GPU 内存。但请注意总线带宽瓶颈。