CUDA实践指南(二十三)

线程和块启发式:
中等优先级:每块的线程数应该是32个线程的倍数,因为这提供了最优的计算效率并促进了合并。
每个网格的块的尺寸和大小以及每个块的线程的尺寸和大小都是重要的因素。 这些参数的多维方面允许更轻松地将多维问题映射到CUDA,并且不影响性能。 因此,本节讨论尺寸而不是维度。
选择执行配置参数应该一起完成; 然而,有一些启发式方法适用于每个参数。 在选择第一个执行配置参数时(每个网格的块数或网格大小),主要关心的是让整个GPU保持繁忙。 网格中的块数应该大于多处理器的数量,以便所有多处理器都具有
至少要执行一个块。 此外,每个多处理器应该有多个活动块,以便不等待__syncthreads()的块可以保持硬件繁忙。 这项建议取决于资源的可用性; 因此,应该在第二个执行参数 - 每个块的线程数或块大小 - 以及共享内存使用情况下确定。 为了扩展到未来的设备,每个内核启动的块数应该是数千。
选择块大小时,务必记住多个并发块可以驻留在多处理器上,因此占用率不仅仅取决于块大小。 特别是,更大的块大小并不意味着更高的占用率。 例如,在计算能力为1.1或更低的设备上,最大块大小为512线程的内核导致66%的占用率,因为此类设备上每个多处理器的最大线程数为768.因此,只有一个 块可以为每个多处理器激活。 但是,在这种设备上每个块有256个线程的内核可以使三个驻留的活动块占用100%。
如占有率所述,较高的占有率并不总是等同于更好的表现。 例如,将占有率从66%提高到100%通常不会转化为性能类似的提高。 较低的占用内核每个线程会拥有比较高的占用内核更多的寄存器,这可能导致寄存器溢出到本地内存的数量减少。 通常情况下,一旦达到50%的占有率,占有率的增加并不意味着改善的性能。 在某些情况下,可以通过更少的翘曲完全覆盖延迟,特别是通过指令级并行(ILP)。
选择块大小涉及很多这样的因素,并且不可避免地需要一些实验。 但是,应遵循一些经验法则:

  • 每个块的线程数应该是warp大小的倍数,以避免在填充不足的warp上浪费计算并促进合并。
  • 每块应该使用至少64个线程,并且每个多处理器只有多个并发块。
  • 每块128到256个线程是更好的选择,并且对于不同块大小的实验来说是一个很好的初始范围。
  • 如果延迟影响性能,则每个多处理器使用几个(3到4个)较小的线程块,而不是一个较大的线程块。 这对经常调用__syncthreads()的内核特别有利。

请注意,当一个线程块分配的寄存器比多处理器上可用的寄存器多时,内核启动失败,因为当请求太多的共享内存或太多线程时,内核启动失败。

上一篇:CUDA实践指南(二十二)


下一篇:C#互操作性入门系列(二):使用平台调用调用Win32 函数