GK110 的 Hyper-Q 与 GK104/GF104 nvidia 卡的并发内核相比有什么优势？

Question

如果我有多线程应用程序和我自己的线程来控制 CUDA 设备并将内核调度到不同的流，我也可以在 Kepler-2 (GK110) 系列之前的设备上实现非常高的 GPU 使用率，例如 Fermi 和 Kepler-1 (GK104) .

所以我认为没有充分的理由渴望购买更昂贵的卡。

流动我的测试示例并在 Fermi 卡上进行分析：

void ut_concurent_kernels()
{
  int i,j;
  cudaEvent_t kernelEvent;
  cudaStream_t    work_stream[14];


  for (i = 0; i < 14;i++)
  {
    cudaStreamCreate( &work_stream[i]);
  }
  cudaEventCreateWithFlags(&kernelEvent, cudaEventDisableTiming);

  for (j = 0; j < 2;j++)
  {
    for (i = 0; i < 14;i++)
    {
      if (i == 13)
      {
        checkCudaErrors(cudaEventRecord(kernelEvent, work_stream[i]));
      }
      Kernel_Work<<<1,256,0,work_stream[i]>>>(100000);
    }
    checkCudaErrors(cudaStreamWaitEvent(work_stream[i-1], kernelEvent,0));
  }
  cudaDeviceSynchronize();

  for (i = 0; i < 14;i++)
  {
    cudaStreamDestroy(work_stream[i]);
  }
  cudaEventDestroy(kernelEvent);
} 

Answer 1

我相信 Hyper-Q 最能被程序员看到的体现是一个名为 CUDA Multi-Process-Server (MPS) 的 Linux 功能。目前 CUDA MPS（在 Hyper-Q 之上）的主要用例是使 MPI 集群计算作业能够在多个 MPI rank 之间共享 GPU，因为 MPI rank 在同一节点上执行，如 here 所讨论的。在这种情况下，多个 MPI 等级通常在独立的 CPU 进程（相对于线程）之外运行，而 CUDA MPS/Hyper-Q 提供了一种方便的共享机制。

如前所述，Hyper-Q 是一种硬件功能，在 cc3.5 设备中可用，但在 cc 1.x - 3.0 设备中不可用。在单个应用程序进程的上下文中，无论是单线程还是多线程，我知道 CUDA MPS（在 Hyper-Q 之上）提供的任何优势或引人注目的用例都不会超过您自己管理多个并发执行。

关于自己管理多个并发执行，Hyper-Q 将（透明地）为程序员提供好处，正如@Greg Smith 在下面的 cmets 中所描述的那样。在 cc3.5 设备中利用它没有特定要求，从概念上讲，用户将使用相同的整体方法管理并发执行，但可能会在 cc3 中体验到改进的并发性或实现改进的并发性的能力。 5 设备，由于 Hyper-Q 的硬件进步。正如 Greg 建议和 OP 所证明的那样，在 cc3.5 之前的设备中实现并发仍然是可能的，但这样做的机会可能更加有限，或者可能需要的编程技术不太明显。要更详细地了解这些问题，请参阅幻灯片 15-22 here