kepler架构GPU新特性--HyperQ

  对于Kepler架构GPU的新特性——HyperQ，往上的讨论贴子还是比较少的，官方文档中也只是有一个sample，给出了代码，但对于有些情况下，HyperQ不能成功的原因没有过多的涉及，我们今天就来谈一谈。HyperQ允许多个CPU线程或进程同时向一个GPU发射任务，提高了GPU的使用率。
        我们来看一个例子：

for(int i = 0; i < 3; i++)
{
        A<<<gdim, bdim, smem, streams[i]>>>();
        B<<<gdim, bdim, smem, streams[i]>>>();
        C<<<gdim, bdim, smem, streams[i]>>>();
}

        如果是在Fermi架构上运行上述代码，每一个流的三个Kernel函数会串行执行，它们有各自的任务队列，根据流的特性，只有A1 C0和A2 C1会具有并行性，如下图：

        我们再来看一个开发包自带的例子：

for(int i = 0; i < nstreams; ++i)
{
        kernel_A<<<1, 1, 0, streams[i]>>>(&d_a[2 * i], time_clocks);
        total_chocks += time_clocks;
        kernel_B<<<1, 1, 0, streams[i]>>>(&d_a[2 * i + 1], time_clocks);
        total_clocks += time_clocks;
}

        下图是在Fermi架构的时序图：

        而在具有HyperQ特性的Kepler架构的GPU上，时序图是这样子的：

        可见，HyperQ特性将同时执行所有流的Kernel函数。
        下面我们来讨论两个问题：
         1. HyperQ和流的关系：
        HyperQ和流的区别。我们看到上边的例子都有流的参与，流可以实现“数据的传输和kernel函数同步执行”，即传输数据的同时，执行kernel函数；而HyperQ则是更高大上的实现了kernel同时执行。如果程序的函数执行时间远远大于数据的传输时间，那么HyperQ就减少了很多很多的执行时间。二者实现的功能是不一样的，这对于初次接触HyperQ的小伙伴们是容易迷惑的地方。
         2. 只要是Kepler架构的GPU，按照上述例子那样写代码，就一定能实现kernel函数并行吗？
        不一定。这句话正确的描述应该是：“只要是Kepler架构的GPU就可以支持HyperQ特性，而HyperQ能不能达到使kernel函数并行的目的是不一定的，要看GPU的资源状况”。什么意思？就是说：如果GPU还有资源，包括显存、空闲的SMX等，那么是可以并行的，如果一个kernel已经让GPU满载运作了，还怎么加入第二个kernel，更不用说多个并行了。有兴趣的小伙伴们可以自行将CUDA中HyperQ的例子使用的资源量加大，在观察时序图。因此，任何高大上的新特性，都是有使用前提的，不能一味的盲目使用。
        上述两点就是本人使用HyperQ的心得，如果在资源允许、GPU没有满载的情况下，HyperQ还是一个很好的特性，减少了程序运行时间；但如果kernel函数的任务量很繁重，那么HyperQ就体现不出kernel并行的特点了，实际执行还是串行的。大家也来谈谈你们使用HyperQ时遇到的问题吧~