cuda - 来自/到统一内存的 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 体系中。)