c++ - 主机循环的 CUDA bool 变量

Question

我的代码的一个非常简化的版本如下所示:

do {
    //reset loop variable b to 0/false
    b = 0;

    // execute kernel
    kernel<<<...>>>(b);

   // use the value of b for while condition
} while(b);

bool 变量 b可以设置为 true通过 kernel 中的任何线程它会告诉我们是否继续运行循环。

使用 cudaMalloc , cudaMemset , 和 cudaMemcpy我们可以创建/设置/复制device内存来实现这一点。但是我刚刚发现固定内存的存在。使用 cudaMalloHost分配b并调用 cudaDeviceSynchronize在内核在一个简单的测试程序中提供相当大的加速(~50%)之后。

固定内存是这个 bool 变量的最佳选择吗 b还是有更好的选择？

Answer 1

您没有显示您的初始代码和修改代码，因此没有人知道您在帖子中陈述的改进细节。

你的问题的答案因人而异

1. b 可读写或仅写入 GPU 内核。如果在缓存中找不到 b 导致延迟，读取可能需要直接从主机端获取实际值。另一方面，如果有进一步的操作可以使线程保持忙碌，则可以弥补写入的延迟。
2. 您修改值的频率。如果您在程序中频繁访问该值，GPU 可能会将变量保留在 L2 中，避免主机端访问。
3. b 访问之间的内存操作频率。如果对b的访问之间有很多内存事务，则缓存中的b更有可能被其他内容替换。结果，再次访问时，在缓存中找不到b，需要进行一次耗时的主机访问。

在主机端有 b 的情况下会导致许多主机内存事务，将其保存在 GPU 全局内存中并在每次循环迭代结束时将其传回是合乎逻辑的。您可以在与内核相同的流中使用异步拷贝来相当快地完成此操作，然后立即与主机同步。

以上所有项目均适用于启用缓存的设备。如果您的设备是 pr-Fermi (CC<2.0)，情况就不同了。