c++ - 如何在 Vulkan 中跨多个计算队列执行并行计算着色器?

标签 c++ gpu gpgpu vulkan compute-shader

更新:此问题已解决,您可以在此处找到更多详细信息:https://stackoverflow.com/a/64405505/1889253
similar question was asked previously ,但这个问题最初集中在使用多个命令缓冲区,并触发跨不同线程的提交以实现着色器的并行执行。大多数答案表明解决方案是使用多个队列。使用多个队列似乎也是各种博客文章和 Khronos 论坛答案的共识。我已经尝试了跨多个队列运行着色器执行的这些建议,但无法看到并行执行,所以我想问一下我可能做错了什么。正如所建议的,这个问题包括提交到多个队列的多个计算着色器的可运行代码,希望这对希望做同样事情的其他人有用(一旦解决)。
当前的实现是in this pull request / branch ,但是我将介绍 Vulkan 的主要特定点,以确保仅需要 Vulkan 知识来回答这个问题。还值得一提的是,当前的用例专门用于计算队列和计算着色器,而不是图形或传输队列(尽管在这些队列之间实现并行性的见解/经验仍然非常有用,并且很可能也会导致答案)。
更具体地说,我有以下几点:

  • Multiple queues first are "fetched" - 我的设备是 NVIDIA 1650,支持队列族索引 0 中的 16 个图形+计算队列,队列族索引 2 中的 8 个计算队列
  • evalAsync performs the submission (which contains recorded shader commands) - 您应该注意到创建了一个我们可以使用的围栏。此外,提交没有任何 waitStageMasks (PipelineStageFlags)。
  • evalAwait allows us to wait for the fence - 调用evalAwait时,我们可以通过创建的围栏等待提交完成

  • 在上面的示例中看不到但很重要的几点:
  • 所有 evalAsync 都在同一个应用程序、实例和设备上运行
  • 每个 evalAsync 都使用自己单独的 commandBuffer 和缓冲区执行,并在单独的队列中执行
  • 如果你想知道内存屏障是否有什么关系,我们已经尝试通过完全删除所有内存屏障(在着色器执行之前运行的 this on for example),但这对性能没有任何影响

  • 基准测试中使用的测试 can be found here ,但是唯一需要理解的关键是:
  • This is the shader我们用于测试的,如您所见,我们只是添加了一堆 atomicAdd 步骤来增加处理时间
  • 目前测试有small buffer尺寸和 high number of shader loop iterations ,但我们还使用大缓冲区大小(即 100,000 而不是 10)和较小的迭代(1,000 代替 100,000,000)进行了测试。

  • 运行测试时,我们首先在同一个队列上运行一组“同步”着色器执行(数量是可变的,但我们已经用 6-16 进行了测试,后者是队列的最大数量)。然后我们以异步方式运行它们,在那里我们运行所有它们和 evalAwait 直到它们完成。当比较两种方法的结果时间时,即使它们运行在不同的计算队列中,它们也花费相同的时间。
    我的问题是:
  • 获取队列时我目前是否遗漏了什么?
  • vulkan 设置中是否还有其他参数需要配置以确保异步执行?
  • 关于潜在的操作系统进程只能以同步方式向 GPU 提交 GPU 工作负载,是否有任何我可能不知道的限制?
  • 在处理多个队列提交时,是否需要多线程才能使并行执行正常工作?

  • 此外,我在各种 reddit 帖子和 Khronos Group 论坛上在线找到了一些有用的资源,这些资源提供了有关该主题的非常深入的概念和理论概述,但我还没有遇到显示并行执行着色器的端到端代码示例。如果您可以分享任何实际示例,这些示例具有着色器的功能并行执行,那将非常有帮助。
    如果有更多细节或问题可以帮助提供更多背景信息,请告诉我,很乐意回答它们和/或提供更多细节。
    为了完整起见,我的测试使用了:
  • Vulkan SDK 1.2
  • Windows 10
  • 英伟达 1650

  • 在类似帖子中分享的其他相关链接:
  • Similar discussion with suggested link to example but which seems to have disappeared...
  • Post on Leveraging asynchronous queues for concurrent execution (遗憾的是没有示例代码)
  • (相对年龄 - 5 岁)Post that suggests nvidia cards can't do parallel execution of shaders ,但似乎没有一个决定性的答案
  • 英伟达演示 on Vulkan Multithreading with multiple queue execution (因此我上面关于线程的问题)
  • 最佳答案

    您正在获得“异步执行”。你只是不期望它的行为方式。
    在 CPU 上,如果您有一个事件线程,那么您使用的是一个 CPU 内核(或超线程)。该核心的所有执行和计算能力都单独提供给您的线程(忽略抢占)。但同时,如果有其他内核,您的一个线程将无法使用这些内核的任何计算资源。除非你创建另一个线程。
    GPU 不是这样工作的。队列不像 CPU 线程。它并不具体涉及特定数量的计算资源。队列只是执行命令的接口(interface);底层硬件决定如何将命令分给 GPU 整体提供的各种计算资源。
    当您执行命令时,通常会发生硬件尝试使用您的命令使可用着色器执行单元完全饱和的情况。如果可用的着色器单元多于您的操作所需的调用次数,则某些资源可立即用于下一个命令。但如果不是,那么整个 GPU 的计算资源将专用于执行第一个操作;第二个必须等待资源可用才能启动。
    您将工作插入多少个计算队列并不重要;他们都将尝试使用尽可能多的计算资源。因此,它们将主要按某些特定顺序执行。
    队列优先级系统存在,但这些系统主要帮助确定命令的执行顺序。也就是说,如果高优先级队列有一些命令需要执行,那么它们将在下次计算资源可用于新命令时优先。
    因此,在 3 个单独的队列上提交 3 个调度批次不会比在包含 3 个调度操作的一个队列上提交 1 个批次完成得更快。
    存在多个队列(同一家族)的主要原因是能够从多个线程提交工作,而无需让它们进行线程间同步(并提供一些可能的提交优先级)。

    关于c++ - 如何在 Vulkan 中跨多个计算队列执行并行计算着色器?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/64384786/

    相关文章:

    c++ - Chromium Embedded Framework 绑定(bind)按键

    c++ - 将 boost::log 用于具有额外 'channel' 和 'id' 属性的多线程应用程序的最佳方法是什么

    caching - GPU 中的 L1 缓存

    cuda - GPU 的每个多处理器有多少个 'CUDA cores'?

    ios - 关于 Metal 中线程组内存的问题

    c++ - 在 ARM 架构上构建时,使用 boost 的动态库具有 undefined reference

    c++ - 是否可以仅使用PDB调试而没有源代码的DLL?

    c++ - 使用 C++ 在 ffmpeg 中解码为特定像素格式

    c - 如何修改此组合算法以在启用 cuda 的 gpu 上并行运行?

    c - __forceinline__ 对 CUDA C __device__ 函数的影响