c++ - cuda在gpu和主机之间统一内存

Question

我正在编写一个基于 cuda 的程序，需要定期将一组项目从 GPU 传输到主机内存。为了保持进程异步，我希望使用 cuda 的 UMA 在主机内存中有一个内存缓冲区和标志(这样 GPU 和 CPU 都可以访问它)。 GPU 将确保标志已清除，将其项目添加到缓冲区，然后设置标志。 CPU 等待设置标志，从缓冲区中复制内容，然后清除标志。据我所知，这不会产生任何竞争条件，因为它会强制 GPU 和 CPU 轮流，始终读取和写入彼此相对的标志。

到目前为止，我还没有能够让它工作，因为似乎确实存在某种竞争条件。我想出了一个具有类似问题的更简单的示例:

#include <stdio.h>

__global__
void uva_counting_test(int n, int *h_i);

int main() {
    int *h_i;
    int n;

    cudaMallocHost(&h_i, sizeof(int));

    *h_i = 0;
    n = 2;

    uva_counting_test<<<1, 1>>>(n, h_i);

    //even numbers
    for(int i = 1; i <= n; ++i) {
        //wait for a change to odd from gpu
        while(*h_i == (2*(i - 1)));

        printf("host h_i: %d\n", *h_i);
        *h_i = 2*i;
    }

    return 0;
}

__global__
void uva_counting_test(int n, int *h_i) {
    //odd numbers
    for(int i = 0; i < n; ++i) {
        //wait for a change to even from host
        while(*h_i == (2*(i - 1) + 1));

        *h_i = 2*i + 1;
    }
}

对我来说，这种情况总是在 CPU 发出第一个打印语句后挂起(host h_i: 1)。真正不寻常的事情(这可能是一个线索)是我可以让它在 cuda-gdb 中工作。如果我在 cuda-gdb 中运行它，它会像以前一样挂起。如果我按下 ctrl+C，它将把我带到内核中的 while() 循环行。从那里，令人惊讶的是，我可以告诉它继续并且它会完成。对于 n > 2，它会在每个内核之后再次卡住在内核中的 while() 循环上，但我可以继续使用 ctrl+C 将其向前推进并继续。

如果有更好的方法来完成我正在尝试做的事情，那也会有所帮助。

Answer 1

您描述的是生产者-消费者模型，其中 GPU 正在生成一些数据，而 CPU 会不时地使用这些数据。

实现这一点的最简单方法是让 CPU 成为主控。 CPU 在 GPU 上启动内核，当它准备好使用数据时(即示例中的 while 循环)，它与 GPU 同步，从 GPU 复制数据，启动内核再次生成更多数据，并对复制的数据做任何与它有关的事情。这允许您在 CPU 处理前一批处理时让 GPU 填充固定大小的缓冲区(因为有两个拷贝，一个在 GPU 上，一个在 CPU 上)。

这可以通过对数据进行双缓冲来改进，这意味着您可以通过在将另一个缓冲区复制到 CPU 时在缓冲区之间进行乒乓来保持 GPU 100% 的时间忙于生成数据。这假设回写速度比生产速度快，但如果不是这样，那么您将饱和复制带宽，这也很好。

这些都不是您实际描述的内容。您要求的是让 GPU 掌握数据。我强烈建议您谨慎对待这一点，因为您需要仔细管理缓冲区大小，并且需要仔细考虑时间安排和通信问题。这样做当然是可能的，但在探索该方向之前，您应该阅读有关内存栅栏、原子操作和 volatile 的内容。