performance - 两个连续的内核还是全网格协作组同步?

标签 performance cuda gpu-cooperative-groups

假设我有两个任务要在 GPU 上运行,其中第二个任务基本上依赖于第一个任务的所有工作。传统上,我基本上必须将这些任务编写为两个单独的内核,并安排第二个任务在第一个任务之后的某个时刻运行。但是 - 借助 CUDA 9,我现在可以在整个网格上同步,完成第一个任务 - using the cooperative groups feature ,然后继续让网格执行其第二个任务工作。

我的问题是:

  • 我们能否提供一个经验法则,说明何时编写两个内核的性能更好,以及何时使用全网格同步?
  • 如果是的话,会是什么?
  • 如果不是 - 为什么很难确定哪种情况下哪种更可取?

最佳答案

将此设为 CW 答案,以便其他人可以随意添加他们的意见并进行编辑。

grid-wide sync feature in cooperative groups它要求将线程补充(网格的大小)限制为您所运行的 GPU 的承载能力。这不是主要的性能限制因素,但它要求您编写可以灵活使用不同网格大小的代码,同时仍能实现最大性能。 grid-stride loops是此类编码策略的典型组成部分。

因此,网格范围同步功能通常需要仔细编码和额外的代码开销(例如使用占用 API)才能实现最大性能,特别是与简单或幼稚的内核相比。

为了抵消程序员生产力可能下降的影响,一些可能的好处是:

  1. 在启动开销占整个运行时间很大一部分的情况下,协作式网格范围同步可能会带来显着的好处。除了 2 个独立内核的融合之外,可以在循环中调用内核的算法(例如雅可比迭代/松弛或其他时间步长模拟算法)可能会显着受益,因为启动循环可以有效地“移至 GPU 上”,用单个内核调用替换内核启动循环。

  2. 在存在大量片上“状态”(例如寄存器内容、共享内存内容)的情况下,需要在网格范围同步之前加载,并在网格范围同步之后使用网格范围同步,那么合作组可能是一个重大胜利,节省了网格范围同步之后的内核时间,这些时间将用于重新加载状态。这似乎就是动机here例如(参见第 4.3 节)。我并不是建议他们使用合作团体(他们没有)。我建议他们寻求网格范围的同步,使用当时可用的临时方法,以消除状态重新加载的成本,以及可能的内核启动开销的成本。

关于performance - 两个连续的内核还是全网格协作组同步?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/54083673/

相关文章:

cmake - cooperative_groups::this_grid() 导致任何 CUDA API 调用返回 'unknown error'

MySQL 没有在带有 "IN"的子选择的外部查询上使用键

cuda - Nvidia 设备上的 OpenCL 与 CUDA 性能

cuda - 线程 block 数与计算机上的cuda内核之间的关系(在CUDA C中)

search - ArrayFire帧搜索算法崩溃

cuda - 未解析的 extern 函数 'cudaCGGetIntrinsicHandle' 反击

python:有效地构建和维护一个按降序排列的字符串列表

javascript - 使用事件监听器添加多个元素的最快方法?监听点击窗口不好吗?

performance - 为什么这些定点 cata/ana 态射定义优于递归定义?