5.1 整體性能優化策略

性能優化圍繞四個基本策略：

最大化并行執行以實現最大利用率；

優化內存使用，實現最大內存吞吐量；

優化指令使用，實現最大指令吞吐量；

盡量減少內存抖動。

哪些策略將為應用程序的特定部分產生最佳性能增益取決于該部分的性能限值； 例如，優化主要受內存訪問限制的內核的指令使用不會產生任何顯著的性能提升。 因此，應該通過測量和監控性能限制來不斷地指導優化工作，例如使用 CUDA 分析器。 此外，將特定內核的浮點運算吞吐量或內存吞吐量（以更有意義的為準）與設備的相應峰值理論吞吐量進行比較表明內核還有多少改進空間。

5.2 最大化利用率

為了最大限度地提高利用率，應用程序的結構應該盡可能多地暴露并行性，并有效地將這種并行性映射到系統的各個組件，以使它們大部分時間都處于忙碌狀態。

5.2.1 應用程序層次

在高層次上，應用程序應該通過使用異步函數調用和異步并發執行中描述的流來最大化主機、設備和將主機連接到設備的總線之間的并行執行。它應該為每個處理器分配它最擅長的工作類型：主機的串行工作負載；設備的并行工作負載。

對于并行工作負載，在算法中由于某些線程需要同步以相互共享數據而破壞并行性的點，有兩種情況： 這些線程屬于同一個塊，在這種情況下，它們應該使用 __syncthreads （） 并在同一個內核調用中通過共享內存共享數據，或者它們屬于不同的塊，在這種情況下，它們必須使用兩個單獨的內核調用通過全局內存共享數據，一個用于寫入，一個用于從全局內存中讀取。第二種情況不太理想，因為它增加了額外內核調用和全局內存流量的開銷。因此，應該通過將算法映射到 CUDA 編程模型以使需要線程間通信的計算盡可能在單個線程塊內執行，從而最大限度地減少它的發生。

5.2.2 設備層次

在較低級別，應用程序應該最大化設備多處理器之間的并行執行。

多個內核可以在一個設備上并發執行，因此也可以通過使用流來啟用足夠多的內核來實現最大利用率，如異步并發執行中所述。

5.2.3 多處理器層次

在更低的層次上，應用程序應該最大化多處理器內不同功能單元之間的并行執行。

如硬件多線程中所述，GPU 多處理器主要依靠線程級并行性來最大限度地利用其功能單元。因此，利用率與常駐warp的數量直接相關。在每個指令發出時，warp 調度程序都會選擇一條準備好執行的指令。該指令可以是同一warp的另一條獨立指令，利用指令級并行性，或者更常見的是另一個warp的指令，利用線程級并行性。如果選擇了準備執行指令，則將其發布到 warp 的活動線程。一個warp準備好執行其下一條指令所需的時鐘周期數稱為延遲，并且當所有warp調度程序在該延遲期間的每個時鐘周期總是有一些指令要為某個warp發出一些指令時，就可以實現充分利用，或者換句話說，當延遲完全“隱藏”時。隱藏 L 個時鐘周期延遲所??需的指令數量取決于這些指令各自的吞吐量（有關各種算術指令的吞吐量，請參見算術指令）。如果我們假設指令具有最大吞吐量，它等于：

4L 用于計算能力 5.x、6.1、6.2、7.x 和 8.x 的設備，因為對于這些設備，多處理器在一個時鐘周期內為每個 warp 發出一條指令，一次四個 warp，如計算能力中所述。

2L 用于計算能力 6.0 的設備，因為對于這些設備，每個周期發出的兩條指令是兩條不同warp的一條指令。

8L 用于計算能力 3.x 的設備，因為對于這些設備，每個周期發出的八條指令是四對，用于四個不同的warp，每對都用于相同的warp。

warp 未準備好執行其下一條指令的最常見原因是該指令的輸入操作數尚不可用。

如果所有輸入操作數都是寄存器，則延遲是由寄存器依賴性引起的，即，一些輸入操作數是由一些尚未完成的先前指令寫入的。在這種情況下，延遲等于前一條指令的執行時間，warp 調度程序必須在此期間調度其他 warp 的指令。執行時間因指令而異。在計算能力 7.x 的設備上，對于大多數算術指令，它通常是 4 個時鐘周期。這意味著每個多處理器需要 16 個活動 warp（4 個周期，4 個 warp 調度程序）來隱藏算術指令延遲（假設 warp 以最大吞吐量執行指令，否則需要更少的 warp）。如果各個warp表現出指令級并行性，即在它們的指令流中有多個獨立指令，則需要更少的warp，因為來自單個warp的多個獨立指令可以背靠背發出。

如果某些輸入操作數駐留在片外存儲器中，則延遲要高得多：通常為數百個時鐘周期。在如此高的延遲期間保持 warp 調度程序繁忙所需的 warp 數量取決于內核代碼及其指令級并行度。一般來說，如果沒有片外存儲器操作數的指令（即大部分時間是算術指令）與具有片外存儲器操作數的指令數量之比較低（這個比例通常是稱為程序的算術強度）。

warp 未準備好執行其下一條指令的另一個原因是它正在某個內存柵欄（內存柵欄函數）或同步點（同步函數）處等待。隨著越來越多的warp等待同一塊中的其他warp在同步點之前完成指令的執行，同步點可以強制多處理器空閑。在這種情況下，每個多處理器擁有多個常駐塊有助于減少空閑，因為來自不同塊的warp不需要在同步點相互等待。

對于給定的內核調用，駐留在每個多處理器上的塊和warp的數量取決于調用的執行配置（執行配置）、多處理器的內存資源以及內核的資源需求，如硬件多線程中所述。使用 --ptxas-options=-v 選項編譯時，編譯器會報告寄存器和共享內存的使用情況。

一個塊所需的共享內存總量等于靜態分配的共享內存量和動態分配的共享內存量之和。

內核使用的寄存器數量會對駐留warp的數量產生重大影響。例如，對于計算能力為 6.x 的設備，如果內核使用 64 個寄存器并且每個塊有 512 個線程并且需要很少的共享內存，那么兩個塊（即 32 個 warp）可以駐留在多處理器上，因為它們需要 2x512x64 個寄存器，它與多處理器上可用的寄存器數量完全匹配。但是一旦內核多使用一個寄存器，就只能駐留一個塊（即 16 個 warp），因為兩個塊需要 2x512x65 個寄存器，這比多處理器上可用的寄存器多。因此，編譯器會盡量減少寄存器的使用，同時保持寄存器溢出（請參閱設備內存訪問）和最少的指令數量?？梢允褂?maxrregcount 編譯器選項或啟動邊界來控制寄存器的使用，如啟動邊界中所述。

寄存器文件組織為 32 位寄存器。因此，存儲在寄存器中的每個變量都需要至少一個 32 位寄存器，例如雙精度變量使用兩個 32 位寄存器。

對于給定的內核調用，執行配置對性能的影響通常取決于內核代碼。因此建議進行實驗。應用程序還可以根據寄存器文件大小和共享內存大小參數化執行配置，這取決于設備的計算能力，以及設備的多處理器數量和內存帶寬，所有這些都可以使用運行時查詢（參見參考手冊）。

每個塊的線程數應選擇為 warp 大小的倍數，以避免盡可能多地在填充不足的 warp 上浪費計算資源。

5.2.3.1 占用率計算

存在幾個 API 函數來幫助程序員根據寄存器和共享內存要求選擇線程塊大小。

占用計算器 API，cudaOccupancyMaxActiveBlocksPerMultiprocessor，可以根據內核的塊大小和共享內存使用情況提供占用預測。此函數根據每個多處理器的并發線程塊數報告占用情況。

請注意，此值可以轉換為其他指標。乘以每個塊的warp數得出每個多處理器的并發warp數；進一步將并發warp除以每個多處理器的最大warp得到占用率作為百分比。

基于占用率的啟動配置器 API，cudaOccupancyMaxPotentialBlockSize 和 cudaOccupancyMaxPotentialBlockSizeVariableSMem，啟發式地計算實現最大多處理器級占用率的執行配置。

以下代碼示例計算 MyKernel 的占用率。然后，它使用并發warp與每個多處理器的最大warp之間的比率報告占用率。

/ Device code
__global__ void MyKernel(int *d, int *a, int *b)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    d[idx] = a[idx] * b[idx];
}

// Host code
int main()
{
    int numBlocks;        // Occupancy in terms of active blocks
    int blockSize = 32;

    // These variables are used to convert occupancy to warps
    int device;
    cudaDeviceProp prop;
    int activeWarps;
    int maxWarps;

    cudaGetDevice(&device);
    cudaGetDeviceProperties(&prop, device);
    
    cudaOccupancyMaxActiveBlocksPerMultiprocessor(
        &numBlocks,
        MyKernel,
        blockSize,
        0);

    activeWarps = numBlocks * blockSize / prop.warpSize;
    maxWarps = prop.maxThreadsPerMultiProcessor / prop.warpSize;

    std::cout << "Occupancy: " << (double)activeWarps / maxWarps * 100 << "%" << std::endl;
    
    return 0;
}

下面的代碼示例根據用戶輸入配置了一個基于占用率的內核啟動MyKernel。

// Device code
__global__ void MyKernel(int *array, int arrayCount)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < arrayCount) {
        array[idx] *= array[idx];
    }
}

// Host code
int launchMyKernel(int *array, int arrayCount)
{
    int blockSize;      // The launch configurator returned block size
    int minGridSize;    // The minimum grid size needed to achieve the
                        // maximum occupancy for a full device
                        // launch
    int gridSize;       // The actual grid size needed, based on input
                        // size

    cudaOccupancyMaxPotentialBlockSize(
        &minGridSize,
        &blockSize,
        (void*)MyKernel,
        0,
        arrayCount);

    // Round up according to array size
    gridSize = (arrayCount + blockSize - 1) / blockSize;

    MyKernel<<>>(array, arrayCount);
    cudaDeviceSynchronize();

    // If interested, the occupancy can be calculated with
    // cudaOccupancyMaxActiveBlocksPerMultiprocessor

    return 0;
}

CUDA 工具包還在 《CUDA_Toolkit_Path》/include/cuda_occupancy.h 中為任何不能依賴 CUDA 軟件堆棧的用例提供了一個自記錄的獨立占用計算器和啟動配置器實現。 還提供了占用計算器的電子表格版本。 電子表格版本作為一種學習工具特別有用，它可以可視化更改影響占用率的參數（塊大小、每個線程的寄存器和每個線程的共享內存）的影響。

5.3 最大化存儲吞吐量

最大化應用程序的整體內存吞吐量的第一步是最小化低帶寬的數據傳輸。

這意味著最大限度地減少主機和設備之間的數據傳輸，如主機和設備之間的數據傳輸中所述，因為它們的帶寬比全局內存和設備之間的數據傳輸低得多。

這也意味著通過最大化片上內存的使用來最小化全局內存和設備之間的數據傳輸：共享內存和緩存（即計算能力 2.x 及更高版本的設備上可用的 L1 緩存和 L2 緩存、紋理緩存和常量緩存 適用于所有設備）。

共享內存相當于用戶管理的緩存：應用程序顯式分配和訪問它。 如 CUDA Runtime 所示，典型的編程模式是將來自設備內存的數據暫存到共享內存中； 換句話說，擁有一個塊的每個線程：

將數據從設備內存加載到共享內存，

與塊的所有其他線程同步，以便每個線程可以安全地讀取由不同線程填充的共享內存位置， 處理共享內存中的數據，

如有必要，再次同步以確保共享內存已使用結果更新，

將結果寫回設備內存。

對于某些應用程序（例如，全局內存訪問模式依賴于數據），傳統的硬件管理緩存更適合利用數據局部性。如 Compute Capability 3.x、Compute Capability 7.x 和 Compute Capability 8.x 中所述，對于計算能力 3.x、7.x 和 8.x 的設備，相同的片上存儲器用于 L1 和共享內存，以及有多少專用于 L1 與共享內存，可針對每個內核調用進行配置。

內核訪問內存的吞吐量可能會根據每種內存類型的訪問模式而變化一個數量級。因此，最大化內存吞吐量的下一步是根據設備內存訪問中描述的最佳內存訪問模式盡可能優化地組織內存訪問。這種優化對于全局內存訪問尤為重要，因為與可用的片上帶寬和算術指令吞吐量相比，全局內存帶寬較低，因此非最佳全局內存訪問通常會對性能產生很大影響。

5.3.1 設備與主機之間的數據傳輸

應用程序應盡量減少主機和設備之間的數據傳輸。 實現這一點的一種方法是將更多代碼從主機移動到設備，即使這意味著運行的內核沒有提供足夠的并行性以在設備上全效率地執行。 中間數據結構可以在設備內存中創建，由設備操作，并在沒有被主機映射或復制到主機內存的情況下銷毀。

此外，由于與每次傳輸相關的開銷，將許多小傳輸批處理為單個大傳輸總是比單獨進行每個傳輸執行得更好。

在具有前端總線的系統上，主機和設備之間的數據傳輸的更高性能是通過使用頁鎖定主機內存來實現的，如頁鎖定主機內存中所述。

此外，在使用映射頁鎖定內存（Mapped Memory）時，無需分配任何設備內存，也無需在設備和主機內存之間顯式復制數據。 每次內核訪問映射內存時都會隱式執行數據傳輸。 為了獲得最佳性能，這些內存訪問必須與對全局內存的訪問合并（請參閱設備內存訪問）。 假設它們映射的內存只被讀取或寫入一次，使用映射的頁面鎖定內存而不是設備和主機內存之間的顯式副本可以提高性能。

在設備內存和主機內存在物理上相同的集成系統上，主機和設備內存之間的任何拷貝都是多余的，應該使用映射的頁面鎖定內存。 應用程序可以通過檢查集成設備屬性（請參閱設備枚舉）是否等于 1 來查詢設備是否集成。

5.3.2 設備內存訪問

訪問可尋址內存（即全局、本地、共享、常量或紋理內存）的指令可能需要多次重新發出，具體取決于內存地址在 warp 內線程中的分布。 分布如何以這種方式影響指令吞吐量特定于每種類型的內存，在以下部分中進行描述。 例如，對于全局內存，一般來說，地址越分散，吞吐量就越低。

全局內存

全局內存駐留在設備內存中，設備內存通過 32、64 或 128 字節內存事務訪問。這些內存事務必須自然對齊：只有32字節、64字節或128字節的設備內存段按其大小對齊（即，其第一個地址是其大小的倍數）才能被內存事務讀取或寫入。

當一個 warp 執行一條訪問全局內存的指令時，它會將 warp 內的線程的內存訪問合并為一個或多個內存事務，具體取決于每個線程訪問的大小以及內存地址在整個線程中的分布。線程。一般來說，需要的事務越多，除了線程訪問的字之外，傳輸的未使用字也越多，相應地降低了指令吞吐量。例如，如果為每個線程的 4 字節訪問生成一個 32 字節的內存事務，則吞吐量除以 8。

需要多少事務以及最終影響多少吞吐量取決于設備的計算能力。 Compute Capability 3.x、Compute Capability 5.x、Compute Capability 6.x、Compute Capability 7.x 和 Compute Capability 8.x 提供了有關如何為各種計算能力處理全局內存訪問的更多詳細信息。

為了最大化全局內存吞吐量，因此通過以下方式最大化合并非常重要：

遵循基于 Compute Capability 3.x、Compute Capability 5.x、Compute Capability 6.x、Compute Capability 7.x 和 Compute Capability 8.x 的最佳訪問模式

使用滿足以下“尺寸和對齊要求”部分中詳述的大小和對齊要求的數據類型，

在某些情況下填充數據，例如，在訪問二維數組時，如下面的二維數組部分所述。

尺寸和對齊要求

全局內存指令支持讀取或寫入大小等于 1、2、4、8 或 16 字節的字。 當且僅當數據類型的大小為 1、2、4、8 或 16 字節并且數據為 對齊（即，它的地址是該大小的倍數）。

如果未滿足此大小和對齊要求，則訪問將編譯為具有交錯訪問模式的多個指令，從而阻止這些指令完全合并。 因此，對于駐留在全局內存中的數據，建議使用滿足此要求的類型。

內置矢量類型自動滿足對齊要求。

對于結構，大小和對齊要求可以由編譯器使用對齊說明符 __align__（8） 或 __align__（16） 強制執行，例如：

struct __align__（8） {

float x;

float y;

};

struct __align__（16） {

float x;

float y;

float z;

};

駐留在全局內存中， 或由驅動程序， 或運行時 API 的內存分配例程之一返回的變量的任何地址始終與至少 256 字節對齊。

讀取非自然對齊的 8 字節或 16 字節字會產生不正確的結果（相差幾個字），因此必須特別注意保持這些類型的任何值或數組值的起始地址對齊。 一個可能容易被忽視的典型情況是使用一些自定義全局內存分配方案時，其中多個數組的分配（多次調用 cudaMalloc（） 或 cuMemAlloc（））被單個大塊內存的分配所取代分區為多個數組，在這種情況下，每個數組的起始地址都與塊的起始地址有偏移。

二維數組

一個常見的全局內存訪問模式是當索引 （tx，ty） 的每個線程使用以下地址訪問一個寬度為 width 的二維數組的一個元素時，位于 type* 類型的地址 BaseAddress （其中 type 滿足最大化中描述的使用要求 ）：

BaseAddress + width * ty + tx

為了使這些訪問完全合并，線程塊的寬度和數組的寬度都必須是 warp 大小的倍數。

特別是，這意味著如果一個數組的寬度不是這個大小的倍數，如果它實際上分配了一個寬度向上舍入到這個大小的最接近的倍數并相應地填充它的行，那么訪問它的效率會更高。 參考手冊中描述的 cudaMallocPitch（） 和 cuMemAllocPitch（） 函數以及相關的內存復制函數使程序員能夠編寫不依賴于硬件的代碼來分配符合這些約束的數組。

本地內存

本地內存訪問僅發生在可變內存空間說明符中提到的某些自動變量上。 編譯器可能放置在本地內存中的變量是：

無法確定它們是否以常數索引的數組，

會占用過多寄存器空間的大型結構或數組，

如果內核使用的寄存器多于可用寄存器（這也稱為寄存器溢出），則為任何變量。

檢查 PTX 匯編代碼（通過使用 -ptx 或 -keep 選項進行編譯）將判斷在第一個編譯階段是否已將變量放置在本地內存中，因為它將使用 .local 助記符聲明并使用 ld 訪問.local 和 st.local 助記符。即使沒有，后續編譯階段可能仍會做出其他決定，但如果他們發現它為目標體系結構消耗了過多的寄存器空間：使用 cuobjdump 檢查 cubin 對象將判斷是否是這種情況。此外，當使用 --ptxas-options=-v 選項編譯時，編譯器會報告每個內核 （lmem） 的總本地內存使用量。請注意，某些數學函數具有可能訪問本地內存的實現路徑。

本地內存空間駐留在設備內存中，因此本地內存訪問與全局內存訪問具有相同的高延遲和低帶寬，并且與設備內存訪問中所述的內存合并要求相同。然而，本地存儲器的組織方式是通過連續的線程 ID 訪問連續的 32 位字。因此，只要一個 warp 中的所有線程訪問相同的相對地址（例如，數組變量中的相同索引，結構變量中的相同成員），訪問就會完全合并。

在某些計算能力 3.x 的設備上，本地內存訪問始終緩存在 L1 和 L2 中，其方式與全局內存訪問相同（請參閱計算能力 3.x）。

在計算能力 5.x 和 6.x 的設備上，本地內存訪問始終以與全局內存訪問相同的方式緩存在 L2 中（請參閱計算能力 5.x 和計算能力 6.x）。

共享內存

因為它是片上的，所以共享內存比本地或全局內存具有更高的帶寬和更低的延遲。

為了實現高帶寬，共享內存被分成大小相等的內存模塊，稱為banks，可以同時訪問。因此，可以同時處理由落在 n 個不同存儲器組中的 n 個地址構成的任何存儲器讀取或寫入請求，從而產生的總帶寬是單個模塊帶寬的 n 倍。

但是，如果一個內存請求的兩個地址落在同一個內存 bank 中，就會發生 bank 沖突，訪問必須串行化。硬件根據需要將具有bank沖突的內存請求拆分為多個單獨的無沖突請求，從而將吞吐量降低等于單獨內存請求數量的總數。如果單獨的內存請求的數量為 n，則稱初始內存請求會導致 n-way bank 沖突。

因此，為了獲得最佳性能，重要的是要了解內存地址如何映射到內存組，以便調度內存請求，從而最大限度地減少內存組沖突。這在計算能力 3.x、計算能力 5.x、計算能力 6.x、計算能力 7.x 和計算能力 8.x 中針對計算能力 3.x、5.x、6.x 7.x 和 8.x 的設備分別進行了描述。

常量內存

常量內存空間駐留在設備內存中，并緩存在常量緩存中。

然后，一個請求被拆分為與初始請求中不同的內存地址一樣多的單獨請求，從而將吞吐量降低等于單獨請求數量的總數。

然后在緩存命中的情況下以常量緩存的吞吐量為結果請求提供服務，否則以設備內存的吞吐量提供服務。

紋理和表面記憶

紋理和表面內存空間駐留在設備內存中并緩存在紋理緩存中，因此紋理提取或表面讀取僅在緩存未命中時從設備內存讀取一次內存，否則只需從紋理緩存讀取一次。 紋理緩存針對 2D 空間局部性進行了優化，因此讀取 2D 中地址靠近在一起的紋理或表面的同一 warp 的線程將獲得最佳性能。 此外，它專為具有恒定延遲的流式提取而設計； 緩存命中會降低 DRAM 帶寬需求，但不會降低獲取延遲。

通過紋理或表面獲取讀取設備內存具有一些優勢，可以使其成為從全局或常量內存讀取設備內存的有利替代方案：

如果內存讀取不遵循全局或常量內存讀取必須遵循以獲得良好性能的訪問模式，則可以實現更高的帶寬，前提是紋理提取或表面讀取中存在局部性；

尋址計算由專用單元在內核外部執行；

打包的數據可以在單個操作中廣播到單獨的變量；

8 位和 16 位整數輸入數據可以選擇轉換為 ［0.0， 1.0］ 或 ［-1.0， 1.0］ 范圍內的 32 位浮點值（請參閱紋理內存）。

5.4最大化指令吞吐量

為了最大化指令吞吐量，應用程序應該：

盡量減少使用低吞吐量的算術指令； 這包括在不影響最終結果的情況下用精度換取速度，例如使用內部函數而不是常規函數（內部函數在內部函數中列出），單精度而不是雙精度，或者將非規范化數字刷新為零；

最大限度地減少由控制流指令引起的發散warp，如控制流指令中所述

減少指令的數量，例如，盡可能優化同步點（如同步指令中所述）或使用受限指針（如 restrict 中所述）。

在本節中，吞吐量以每個多處理器每個時鐘周期的操作數給出。 對于 32 的 warp 大小，一條指令對應于 32 次操作，因此如果 N 是每個時鐘周期的操作數，則指令吞吐量為每個時鐘周期的 N/32 條指令。

所有吞吐量都是針對一個多處理器的。 它們必須乘以設備中的多處理器數量才能獲得整個設備的吞吐量。

5.4.1 算數指令

如下圖所示

其他指令和功能是在本機指令之上實現的。不同計算能力的設備實現可能不同，編譯后的native指令的數量可能會隨著編譯器版本的不同而波動。對于復雜的函數，可以有多個代碼路徑，具體取決于輸入。 cuobjdump 可用于檢查 cubin 對象中的特定實現。

一些函數的實現在 CUDA 頭文件（math_functions.h、device_functions.h、…）上很容易獲得。

通常，使用 -ftz=true 編譯的代碼（非規范化數字刷新為零）往往比使用 -ftz=false 編譯的代碼具有更高的性能。類似地，使用 -prec-div=false（不太精確的除法）編譯的代碼往往比使用 -prec-div=true 編譯的代碼具有更高的性能，使用 -prec-sqrt=false（不太精確的平方根）編譯的代碼往往比使用 -prec-sqrt=true 編譯的代碼具有更高的性能。 nvcc 用戶手冊更詳細地描述了這些編譯標志。

Single-Precision Floating-Point Division

__fdividef（x， y）（參見內部函數）提供比除法運算符更快的單精度浮點除法。

Single-Precision Floating-Point Reciprocal Square Root

為了保留 IEEE-754 語義，編譯器可以將 1.0/sqrtf（） 優化為 rsqrtf（），僅當倒數和平方根都是近似值時（即 -prec-div=false 和 -prec-sqrt=false）。 因此，建議在需要時直接調用 rsqrtf（）。

Single-Precision Floating-Point Square Root

單精度浮點平方根被實現為倒數平方根后跟倒數，而不是倒數平方根后跟乘法，因此它可以為 0 和無窮大提供正確的結果。

Sine and Cosine

sinf（x）、cosf（x）、tanf（x）、sincosf（x） 和相應的雙精度指令更昂貴，如果參數 x 的量級很大，則更是如此。

更準確地說，參數縮減代碼（參見實現的數學函數）包括兩個代碼路徑，分別稱為快速路徑和慢速路徑。

快速路徑用于大小足夠小的參數，并且基本上由幾個乘加運算組成。 慢速路徑用于量級較大的參數，并且包含在整個參數范圍內獲得正確結果所需的冗長計算。

目前，三角函數的參數縮減代碼為單精度函數選擇幅度小于105615.0f，雙精度函數小于2147483648.0的參數選擇快速路徑。

由于慢速路徑比快速路徑需要更多的寄存器，因此嘗試通過在本地內存中存儲一些中間變量來降低慢速路徑中的寄存器壓力，這可能會因為本地內存的高延遲和帶寬而影響性能（請參閱設備內存訪問）。 目前單精度函數使用28字節的本地內存，雙精度函數使用44字節。 但是，確切的數量可能會發生變化。

由于在慢路徑中需要進行冗長的計算和使用本地內存，當需要進行慢路徑縮減時，與快速路徑縮減相比，這些三角函數的吞吐量要低一個數量級。

Integer Arithmetic

整數除法和模運算的成本很高，因為它們最多可編譯為 20 條指令。 在某些情況下，它們可以用按位運算代替：如果 n 是 2 的冪，則 （i/n） 等價于 （i》》log2（n）） 并且 （i%n） 等價于 （i&（n- 1））; 如果 n 是字母，編譯器將執行這些轉換。

__brev 和 __popc 映射到一條指令，而 __brevll 和 __popcll 映射到幾條指令。

__［u］mul24 是不再有任何理由使用的遺留內部函數。

Half Precision Arithmetic

為了實現 16 位精度浮點加法、乘法或乘法加法的良好性能，建議將 half2 數據類型用于半精度，將 __nv_bfloat162 用于 __nv_bfloat16 精度。 然后可以使用向量內在函數（例如 __hadd2、__hsub2、__hmul2、__hfma2）在一條指令中執行兩個操作。 使用 half2 或 __nv_bfloat162 代替使用 half 或 __nv_bfloat16 的兩個調用也可能有助于其他內在函數的性能，例如warp shuffles。

提供了內在的 __halves2half2 以將兩個半精度值轉換為 half2 數據類型。

提供了內在的 __halves2bfloat162 以將兩個 __nv_bfloat 精度值轉換為 __nv_bfloat162 數據類型。

Type Conversion

有時，編譯器必須插入轉換指令，從而引入額外的執行周期。 情況如下：

對 char 或 short 類型的變量進行操作的函數，其操作數通常需要轉換為 int，

雙精度浮點常量（即那些沒有任何類型后綴定義的常量）用作單精度浮點計算的輸入（由 C/C++ 標準規定）。

最后一種情況可以通過使用單精度浮點常量來避免，這些常量使用 f 后綴定義，例如 3.141592653589793f、1.0f、0.5f。

5.4.2 控制流指令

任何流控制指令（if、switch、do、for、while）都可以通過導致相同 warp 的線程發散（即遵循不同的執行路徑）來顯著影響有效指令吞吐量。如果發生這種情況，則必須對不同的執行路徑進行序列化，從而增加為此 warp 執行的指令總數。

為了在控制流取決于線程 ID 的情況下獲得最佳性能，應編寫控制條件以最小化發散warp的數量。這是可能的，因為正如 SIMT 架構中提到的那樣，整個塊的warp分布是確定性的。一個簡單的例子是當控制條件僅取決于 （threadIdx / warpSize） 時，warpSize 是warp大小。在這種情況下，由于控制條件與warp完全對齊，因此沒有warp發散。

有時，編譯器可能會展開循環，或者它可能會通過使用分支預測來優化短 if 或 switch 塊，如下所述。在這些情況下，任何warp都不會發散。程序員還可以使用#pragma unroll 指令控制循環展開（參見#pragma unroll）。

當使用分支預測時，其執行取決于控制條件的任何指令都不會被跳過。相反，它們中的每一個都與基于控制條件設置為真或假的每線程條件代碼或預測相關聯，盡管這些指令中的每一個都被安排執行，但實際上只有具有真預測的指令被執行。帶有錯誤預測的指令不寫入結果，也不評估地址或讀取操作數。

5.4.3 同步指令

對于計算能力為 3.x 的設備，__syncthreads（） 的吞吐量為每個時鐘周期 128 次操作，對于計算能力為 6.0 的設備，每個時鐘周期為 32 次操作，對于計算能力為 7.x 和 8.x 的設備，每個時鐘周期為 16 次操作。 對于計算能力為 5.x、6.1 和 6.2 的設備，每個時鐘周期 64 次操作。

請注意，__syncthreads（） 可以通過強制多處理器空閑來影響性能，如設備內存訪問中所述。

5.5最小化內存抖動

經常不斷地分配和釋放內存的應用程序可能會發現分配調用往往會隨著時間的推移而變慢，直至達到極限。這通常是由于將內存釋放回操作系統供其自己使用的性質而預期的。為了在這方面獲得最佳性能，我們建議如下：

嘗試根據手頭的問題調整分配大小。不要嘗試使用 cudaMalloc / cudaMallocHost / cuMemCreate 分配所有可用內存，因為這會強制內存立即駐留并阻止其他應用程序能夠使用該內存。這會給操作系統調度程序帶來更大的壓力，或者只是阻止使用相同 GPU 的其他應用程序完全運行。

嘗試在應用程序的早期以適當大小分配內存，并且僅在應用程序沒有任何用途時分配內存。減少應用程序中的 cudaMalloc+cudaFree 調用次數，尤其是在性能關鍵區域。

如果應用程序無法分配足夠的設備內存，請考慮使用其他內存類型，例如 cudaMallocHost 或 cudaMallocManaged，它們的性能可能不高，但可以使應用程序取得進展。

對于支持該功能的平臺，cudaMallocManaged 允許超額訂閱，并且啟用正確的 cudaMemAdvise 策略，將允許應用程序保留 cudaMalloc 的大部分（如果不是全部）性能。 cudaMallocManaged 也不會強制分配在

關于作者

Ken He 是 NVIDIA 企業級開發者社區經理 & 高級講師，擁有多年的 GPU 和人工智能開發經驗。自 2017 年加入 NVIDIA 開發者社區以來，完成過上百場培訓，幫助上萬個開發者了解人工智能和 GPU 編程開發。在計算機視覺，高性能計算領域完成過多個獨立項目。并且，在機器人和無人機領域，有過豐富的研發經驗。對于圖像識別，目標的檢測與跟蹤完成過多種解決方案。曾經參與 GPU 版氣象模式GRAPES，是其主要研發者。

審核編輯：郭婷