memcpy from/to 统一内存是否表现出同步行为？

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 from/to 托管内存呢？在我的测试中，同步似乎是隐式发生的，但我在文档中找不到。

Answer 1

是的，您需要同步。

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

如果你后续的拷贝操作是希望取到内核修改过的数据，需要强制内核先完成。

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

（在 pre-pascal UM 机制中需要同步。您没有遇到段错误这一事实向我表明您处于demand-paged UM 制度。）

memcpy from/to 统一内存是否表现出同步行为？

Does memcpy from/to unified memory exhibits synchronous behavior?

cuda