CUDA 动态并行性：在全局内存中访问子内核结果

Question

我目前正在 CUDA 中尝试我的第一个动态并行代码。这很简单。在父内核中，我正在做这样的事情：

int aPayloads[32];
// Compute aPayloads start values here

int* aGlobalPayloads = nullptr;
cudaMalloc(&aGlobalPayloads, (sizeof(int) *32));
cudaMemcpyAsync(aGlobalPayloads, aPayloads, (sizeof(int)*32), cudaMemcpyDeviceToDevice));

mykernel<<<1, 1>>>(aGlobalPayloads); // Modifies data in aGlobalPayloads
cudaDeviceSynchronize();

// Access results in payload array here

假设到目前为止我做的都是正确的，那么在内核执行后访问 aGlobalPayloads 中的结果的最快方法是什么？（我尝试 cudaMemcpy() 将 aGlobalPayloads 复制回 aPayloads 但设备代码中不允许使用 cudaMemcpy()。

Answer 1

您可以直接从您的父内核代码访问aGlobalPayloads中的数据，无需任何复制：

mykernel<<<1, 1>>>(aGlobalPayloads); // Modifies data in aGlobalPayloads
cudaDeviceSynchronize();
int myval = aGlobalPayloads[0];

我鼓励仔细检查错误（阅读整个接受的答案 here). You do it in device code the same way as in host code. The programming guide 指出：“不得传入本地或共享内存指针”。您对 aPayloads 的使用是本地内存指针。
如果出于某种原因您希望将该数据显式放回本地数组，您可以使用 in-kernel memcpy：
```
memcpy(aPayloads, aGlobalPayloads, sizeof(int)*32);
int myval = aPayloads[0]; // retrieves the same value
```
（这也是我解决第 2 项中提到的问题的方法 - 使用 in-kernel memcpy）

CUDA 动态并行性：在全局内存中访问子内核结果

CUDA dynamic parallelism: Access child kernel results in global memory

memory-management

cuda

dynamic-parallelism