【问题标题】:How do you keep data in fast GPU memory (l1/shared) across kernel invocations?如何跨内核调用将数据保存在快速 GPU 内存(l1/共享)中?
【发布时间】:2018-10-07 18:22:19
【问题描述】:

如何跨内核调用将数据保存在快速 GPU 内存中?

假设,我需要回答 100 万个查询,每个查询有大约 1.5MB 的数据可以在调用中重复使用,并且每个查询有大约 8KB 的数据。

一种方法是为每个查询启动一个内核,每次将 1.5MB + 8KB 的数据复制到共享内存。但是,我花了很多时间只复制 1.5MB 的数据,这些数据确实可以在查询中持续存在。

另一种方法是“回收”GPU 线程(请参阅https://stackoverflow.com/a/49957384/3738356)。这涉及启动一个内核,该内核会立即将 1.5MB 的数据复制到共享内存。然后内核等待请求进来,等待 8KB 的数据出现,然后再进行每次迭代。看起来 CUDA 不应该以这种方式使用。如果仅使用托管内存和 volatile+单调递增的计数器进行同步,则仍然无法保证计算答案所需的数据在您阅读时会在 GPU 上。您可以使用诸如 -42 之类的虚拟值来为内存中的值播种,这表明该值尚未进入 GPU(通过缓存/托管内存机制),然后忙等待直到这些值变为有效。从理论上讲,这应该有效。但是,我有足够的内存错误,我现在已经放弃了,我一直在追求....

另一种方法仍然使用回收的线程,而是通过cudaMemcpyAsync、cuda 流、cuda 事件和几个 volatile+单调递增的计数器来同步数据。我听说我需要固定每个查询的 8KB 新鲜数据,以便cudaMemcpyAsync 正常工作。但是,异步副本没有被阻止——它的效果是不可观察的。我怀疑只要有足够的勇气,我也可以完成这项工作。

但是,以上所有情况都让我觉得“我做错了”。 您如何在 GPU 缓存中保留高度可重用的数据,以便可以从一个查询访问到下一个查询?

【问题讨论】:

  • 我一直想知道将数据放入纹理内存而不是一般内存是否会有所帮助。纹理内存似乎没有跨内核调用进行 g.c.'d。但是,我认为这样的选择会影响我的 L1 可用性……这在 P100 上并不是很好(在 V100 上情况应该会得到很大改善)。此外,似乎每个 SM 的共享内存比 L1 更适合。
  • 你不能。从技术上讲,所有快速(片上)内存在内核终止时都会失效,主要是出于安全原因。如您所料,您最好的选择是使用循环回收线程,为此您必须手动实现线程分解和线程间块同步。
  • 简而言之,您不能在内核调用之间共享任何内容。因此,您将遇到与 persistent kernel 方法相关的方法(您在上一个关于 CUDA 标签的问题中已经写过。)Here 是另一个示例。 CUDA 9 中合作组的某些方面旨在帮助促进此类事情(尤其是合作启动)。

标签: cuda


【解决方案1】:
  • 首先观察流和异步复制的效果 你肯定需要固定主机内存。然后你可以观察 并发内核调用“几乎”同时发生。 我宁愿使用异步复制,因为它让我感觉可以控制 情况。
  • 其次,您可以只保留全局内存中的数据并加载 在您需要时将其保存在共享内存中。据我所知共享 内存只有内核自己知道并在之后处理 终止。尝试在内核运行时使用异步副本 相应地同步流。不要忘记 __syncthreads() 加载到共享内存后。希望对您有所帮助。

【讨论】:

    猜你喜欢
    • 2018-03-27
    • 1970-01-01
    • 2018-05-31
    • 2017-08-28
    • 2012-06-24
    • 1970-01-01
    • 1970-01-01
    • 2015-10-26
    相关资源
    最近更新 更多