gpu和主机之间的cuda统一内存答案

【问题标题】：Cuda unified memory between gpu and hostgpu和主机之间的cuda统一内存
【发布时间】：2014-06-18 20:42:21
【问题描述】：

我正在编写一个基于 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 继续推动它并继续。

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

【问题讨论】：

代码中的任何内容都不能保证缓存的一致性。如果没有某种内存栅栏，这种方法就行不通。而是考虑每次都启动一个内核，无论如何与统一内存访问相比，这相当便宜。
您的示例代码不起作用，因为在内核执行期间无法保证跨 PCI-e 总线的内存一致性。这个游戏的基本规则是不要尝试设计任何依赖于 GPU 和主机设备之间显式主机驱动程序级别同步以外的任何东西的执行模型。
您没有使用Unified Memory.。您正在使用零拷贝主机内存。如果您只想查看有效的计数测试，请查看here。除了关于您的方法的所有其他 cmets，今天的统一内存实现并非旨在为主机和当前执行的内核提供对内存区域的同时一致访问。

标签： c++ c cuda

【解决方案1】：

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

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

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

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

【讨论】：

我喜欢您的回答，并尝试在成功的同时自己实施生产者-消费者模型（您的第二段）。我想知道您是否可以举一个例子，或者指出一个例子，说明如何进行您在第 3 段中描述的双缓冲。谢谢。

【解决方案2】：

我会尝试添加

__threadfence_system();

之后

*h_i = 2*i + 1;

详情请参阅here。没有它，修改完全有可能永远留在 GPU 缓存中。但是，您最好听听其他答案：要针对多个线程/块改进它，您必须处理其他“问题”才能使类似的方案可靠地工作。

正如 Tom 建议的 (+1)，最好使用双缓冲。流对这样的方案有很大帮助，您可以找到描述的 here。

【讨论】：