来自/到统一内存的 memcpy 是否表现出同步行为？

Question

在下面的代码中：

__managed__ int mData[1024];

void foo(int* dataOut)
{
    some_kernel_that_writes_to_mdata<<<...>>>();
    // cudaDeviceSynchronize() // do I need this synch here?
    memcpy(dataOut, mData, sizeof(int) * 1024);

    ...

    cudaDeviceSynchronize();
}

我需要kernel和之间的同步吗memcpy？

cudaMemcpy文档提到该函数在大多数用例中表现出同步行为。但是来自/到托管内存的“正常”情况又如何呢？ memcpy在我的测试中，同步似乎是隐式发生的，但我在文档中找不到这一点。

Answer 1

是的，您需要同步。

内核启动是异步的。因此，在启动内核后，CPU 线程将继续执行下一行代码，但不保证内核完成。

如果您后续的复制操作期望获取内核修改的数据，则有必要先强制内核完成。

cudaMemcpy是一个特例。它被发送到默认流中。它既具有设备同步特性（在开始复制之前强制完成所有先前向该设备发出的工作），也具有 CPU 线程阻塞特性（它不从库调用返回，即允许 CPU 线程继续，直到复制操作完成。）

（在前 Pascal UM 体系中也需要同步。事实上，您没有遇到段错误，这表明我处于请求分页 UM 体系中。）