对于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时遇到的问题吧~