我有一些内核使用一个公共（共享）参数，但是，每个内核都做独立的事情：它们不重叠。如何让它们重叠？
CUDA 代码如下所示（k0 和 k1 在专用流 str[i] 中启动）：
k0<<<..., str[0]>>>(arg1, arg2, arg3);
k1<<<..., str[1]>>>(arg4, arg2, arg5); // k2 uses also arg2
另外，我特意声明了如下内核：
k0(double const * const arg1, double const * const arg2, double * arg3);
k1(double const * const arg1, double const * const arg2, double * arg3);
这是 arg2 的 double const * const，希望这个提示能帮助 nvcc。
nvvp 显示 k1 在 str[0] 中运行完毕后， k2 在 str[1] 中运行。
内核不重叠：这与公共参数（arg2）有关吗？如果是，在这种情况下如何使内核相互重叠？
我使用开普勒 K20m ：
不幸的是，该算法总是访问数据，因此我希望它具有高度的带宽限制（每个线程通常从多个数组读取大量数据，进行一些加法或乘法，然后写回数组 - 没有太多方法可以避免这种情况，或者甚至以不同的方式编码......）。注意：算法中没有“if”，因此发散执行不是问题。
从 v0 开始，我实现了另一个版本 v1 来“使内核更大”：现在，我有内核启动时间（根据 nvvp 为 50 微秒）<< 内核运行时间（根据 nvvp 为 4 毫秒），这对我来说似乎是一件好事。此外，与 v0 相比，V1 的效率（存储：90％，负载：155％）和占用率（实现：52％，理论：62％）都提高了很多。此时，nvvp 内核性能限制器显示内核是“计算限制的”（“功能单元”：85%，“内存”：5%），这对我来说似乎又是一件好事（即使我很惊讶）正如我预期内核是有带宽限制的）。
与 CPU 相比，v1 的速度仍然下降：根据 nvvp，似乎 50% 的“停顿原因”是“执行依赖性”（饼图 - 内核延迟）。因此，从 v1 开始，我尝试实现 v2，即“v1 分割成几个独立的部分”（希望通过并行独立的事情来增加指令级并行性）：执行依赖性跳转到 70% 并且（独立）内核不重叠，这就是导致我来到 StackOverflow 的根本原因...
根据Tom的回答，我不确定是否知道“如何检查GPU是否已满”。v2 的内核有 30% 的占用率：对我来说，这意味着已经有空间供其他内核使用，不是吗？！我刚刚尝试过较小的块/网格大小，但似乎没有帮助（更糟糕的是：占用率下降到 10%）。
从 nvvp（在应用程序级别）提供的总体提示中，我收到一条消息，例如“低并发内核：并行执行 2 个内核的时间百分比较低”。
结论：我在隧道尽头看不到任何光明......任何好主意将不胜感激！我觉得我错过了一些卡住 GPU 的东西，但是，我不知道这是什么。
我已经尝试使用 CUDA 占用计算器电子表格来调整网格和块的大小。事实证明，对于这个算法，大块和大网格似乎比小块和小网格表现得更好。内核不使用共享内存。我猜想没有足够的寄存器供独立内核同时运行：如果是这样，我所观察到的就是合乎逻辑的。我会尝试在这一点上发挥作用。
根据记录，减少寄存器的使用并没有帮助。
内核不重叠：这与公共参数（arg2）有关吗？
不，这并不是因为它们共享输入参数。
通过使用 CUDA 流，程序员明确表示 k0 和 k1 可以并发执行，因此即使您将 arg5 替换为 arg3，硬件仍然可以并发执行它们。
由于这没有发生：
concurrentKernels设备属性来检查（参见deviceQuery示例）根据您提供的附加信息（请不要发布附加信息作为答案，您应该修改您的问题或评论），我们可以得出结论：(1) 就是这种情况。工作调度程序将在开始下一个内核之前调度一个内核的所有块；由于第一个内核需要几毫秒的时间来执行并且由许多线程块组成，因此第二个内核只会在最后一个块开始完成时才开始执行。
占用率是 SM 管理的 warp 数量的衡量标准，占用率为 30%，这意味着某些其他因素限制了可以调度的 warp 数量（例如寄存器使用、共享内存使用） - 这并不意味着我们可以运行第二个内核。您可以使用 CUDA 占用计算器电子表格进一步调查占用情况。
有关并发内核的更多信息，您还可以查看simpleHyperQ 文档。