从/到统一内存的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文档提到函数对大多数用例显示同步行为。但是"normal">memcpy从/到托管内存呢?在我的测试中，同步似乎是隐式发生的，但我在文档中找不到。

Answer 1

是的，你需要同步。

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

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

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

(同步也是必需的在pre-pascal UM制度中。在我看来，您没有出现分段错误的事实表明，您处于需求页面的UM机制中。