CUDA實踐指南（二十一）

執行配置優化：

良好性能的關鍵之一是盡可能讓裝置上的多處理器盡可能地繁忙。 在多處理器中工作平衡不佳的裝置将提供不理想的性能。 是以，将應用程式設計為使用線程和塊以最大限度地提高硬體使用率并限制妨礙工作免費配置設定的做法非常重要。 這項工作的一個關鍵概念是入住率，這在以下幾節中進行了解釋。

在某些情況下，通過設計應用程式也可以提高硬體使用率，進而可以同時執行多個獨立的核心。 同時執行多個核心稱為并發核心執行。 并發核心執行如下所述。

另一個重要的概念是配置設定給特定任務的系統資源的管理。 本章的最後部分将讨論如何管理資源使用率。

占用：

線程指令在CUDA中按順序執行，是以，當一個warp暫停或停頓時執行其他warp是隐藏延遲和保持硬體繁忙的唯一方法。 是以，在确定硬體保持繁忙的有效程度時，與多處理器上的活動warp數相關的一些度量标準非常重要。 這個度量是占用率。

占用率是每個多處理器的活動彎曲數量與可能的活動彎曲數量的比率。 （要确定後一個數字，請參閱deviceQuery CUDA Sample或參考CUDA C程式設計指南中的計算功能。）檢視占用率的另一種方法是硬體處理正在使用的變形的能力的百分比。

較高的占用率并不總是等同于較高的性能 - 有一點超出額外占用率不會提高性能。 但是，低占用率總是會幹擾隐藏記憶體延遲的能力，進而導緻性能下降。

計算占用率：

确定占用率的幾個因素之一是注冊可用性。 寄存器存儲使線程能夠将本地變量保持在附近以實作低延遲通路。 但是，這組寄存器（稱為寄存器檔案）是駐留在多處理器上的所有線程必須共享的有限商品。 寄存器一次配置設定到整個塊。 是以，如果每個線程塊使用多個寄存器，則可以駐留在多處理器上的線程塊數量減少，進而降低了多處理器的占用率。 可以使用-maxrregcount選項或per-kernel使用<code>__launch_bounds__</code>限定符（請參閱寄存器壓力）在每個檔案的編譯時間手動設定每個線程的最大寄存器數量。

為了計算占用率，每個線程使用的寄存器數量是關鍵因素之一。例如，具有計算能力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也計算占用率，作為應用程式分析的多處理器階段的一部分。