CUDA实践指南（二十一）

执行配置优化：
良好性能的关键之一是尽可能让设备上的多处理器尽可能地繁忙。 在多处理器中工作平衡不佳的设备将提供不理想的性能。 因此，将应用程序设计为使用线程和块以最大限度地提高硬件利用率并限制妨碍工作免费分配的做法非常重要。 这项工作的一个关键概念是入住率，这在以下几节中进行了解释。
在某些情况下，通过设计应用程序也可以提高硬件利用率，从而可以同时执行多个独立的内核。 同时执行多个内核称为并发内核执行。 并发内核执行如下所述。
另一个重要的概念是分配给特定任务的系统资源的管理。 本章的最后部分将讨论如何管理资源利用率。
占用：
线程指令在CUDA中按顺序执行，因此，当一个warp暂停或停顿时执行其他warp是隐藏延迟和保持硬件繁忙的唯一方法。 因此，在确定硬件保持繁忙的有效程度时，与多处理器上的活动warp数相关的一些度量标准非常重要。 这个度量是占用率。
占用率是每个多处理器的活动弯曲数量与可能的活动弯曲数量的比率。 （要确定后一个数字，请参阅deviceQuery CUDA Sample或参考CUDA C编程指南中的计算功能。）查看占用率的另一种方法是硬件处理正在使用的变形的能力的百分比。
较高的占用率并不总是等同于较高的性能 - 有一点超出额外占用率不会提高性能。 但是，低占用率总是会干扰隐藏内存延迟的能力，从而导致性能下降。
计算占用率：
确定占用率的几个因素之一是注册可用性。 寄存器存储使线程能够将本地变量保持在附近以实现低延迟访问。 但是，这组寄存器（称为寄存器文件）是驻留在多处理器上的所有线程必须共享的有限商品。 寄存器一次分配到整个块。 因此，如果每个线程块使用多个寄存器，则可以驻留在多处理器上的线程块数量减少，从而降低了多处理器的占用率。 可以使用-maxrregcount选项或per-kernel使用__launch_bounds__限定符（请参阅寄存器压力）在每个文件的编译时间手动设置每个线程的最大寄存器数量。
为了计算占用率，每个线程使用的寄存器数量是关键因素之一。例如，具有计算能力1.1的设备每个多处理器具有8,192个32位寄存器，并且最多可以有768个同时驻留的线程（每个warp有24个经线x 32个线程）。这意味着在其中一个设备中，对于多处理器占用率达到100％，每个线程最多可以使用10个寄存器。但是，这种确定寄存器计数如何影响占用的方法并未考虑寄存器分配的粒度。例如，在计算能力1.1的设备上，每个线程使用12个寄存器的具有128个线程块的内核导致占用83％，每个多处理器具有5个活动128线程块，而具有256个线程块的内核每个线程使用相同的12个寄存器会导致占用66％，因为只有两个256线程块可以驻留在多处理器上。此外，在具有计算能力1.1的设备上，寄存器分配四舍五入到每块最近的256个寄存器。
可用寄存器的数量，每个多处理器上同时驻留的最大线程数以及寄存器分配粒度随不同的计算能力而变化。由于寄存器分配的细微差别以及多处理器共享内存也在常驻线程块之间进行分区的事实，寄存器使用情况和占用情况之间的确切关系可能难以确定。 nvcc的--ptxas options = v选项详细说明每个内核每个线程使用的寄存器数量。有关各种计算能力的器件的寄存器分配公式以及CUDA C编程指南的特性和技术规范，请参见CUDA C编程指南的硬件多线程，了解这些器件上可用的寄存器总数。另外，NVIDIA还提供Excel电子表格形式的占用率计算器，使开发人员能够更好地平衡最佳平衡并更轻松地测试不同的可能场景。该电子表格（如图11所示）称为CUDA_Occupancy_Calculator.xls，位于CUDA Toolkit安装的tools子目录中。
除了计算器电子表格之外，可以使用NVIDIA Visual Profiler的“实现占用率”度量标准来确定占位情况。 Visual Profiler也计算占用率，作为应用程序分析的多处理器阶段的一部分。