CUDA編程模型的整體性能優(yōu)化策略

5.1 整體性能優(yōu)化策略

性能優(yōu)化圍繞四個基本策略：

最大化并行執(zhí)行以實現(xiàn)最大利用率；

優(yōu)化內(nèi)存使用，實現(xiàn)最大內(nèi)存吞吐量；

優(yōu)化指令使用，實現(xiàn)最大指令吞吐量；

盡量減少內(nèi)存抖動。

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

5.2 最大化利用率

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

5.2.1 應(yīng)用程序?qū)哟?/p>

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

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

5.2.2 設(shè)備層次

在較低級別，應(yīng)用程序應(yīng)該最大化設(shè)備多處理器之間的并行執(zhí)行。

多個內(nèi)核可以在一個設(shè)備上并發(fā)執(zhí)行，因此也可以通過使用流來啟用足夠多的內(nèi)核來實現(xiàn)最大利用率，如異步并發(fā)執(zhí)行中所述。

5.2.3 多處理器層次

在更低的層次上，應(yīng)用程序應(yīng)該最大化多處理器內(nèi)不同功能單元之間的并行執(zhí)行。

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

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

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

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

warp 未準(zhǔn)備好執(zhí)行其下一條指令的最常見原因是該指令的輸入操作數(shù)尚不可用。

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

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

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

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

一個塊所需的共享內(nèi)存總量等于靜態(tài)分配的共享內(nèi)存量和動態(tài)分配的共享內(nèi)存量之和。

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

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

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

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

5.2.3.1 占用率計算

存在幾個 API 函數(shù)來幫助程序員根據(jù)寄存器和共享內(nèi)存要求選擇線程塊大小。

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

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

基于占用率的啟動配置器 API，cudaOccupancyMaxPotentialBlockSize 和 cudaOccupancyMaxPotentialBlockSizeVariableSMem，啟發(fā)式地計算實現(xiàn)最大多處理器級占用率的執(zhí)行配置。

以下代碼示例計算 MyKernel 的占用率。然后，它使用并發(fā)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;
}

下面的代碼示例根據(jù)用戶輸入配置了一個基于占用率的內(nèi)核啟動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 軟件堆棧的用例提供了一個自記錄的獨立占用計算器和啟動配置器實現(xiàn)。 還提供了占用計算器的電子表格版本。 電子表格版本作為一種學(xué)習(xí)工具特別有用，它可以可視化更改影響占用率的參數(shù)（塊大小、每個線程的寄存器和每個線程的共享內(nèi)存）的影響。

5.3 最大化存儲吞吐量

最大化應(yīng)用程序的整體內(nèi)存吞吐量的第一步是最小化低帶寬的數(shù)據(jù)傳輸。

這意味著最大限度地減少主機和設(shè)備之間的數(shù)據(jù)傳輸，如主機和設(shè)備之間的數(shù)據(jù)傳輸中所述，因為它們的帶寬比全局內(nèi)存和設(shè)備之間的數(shù)據(jù)傳輸?shù)偷枚唷?/p>

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

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

將數(shù)據(jù)從設(shè)備內(nèi)存加載到共享內(nèi)存，

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

如有必要，再次同步以確保共享內(nèi)存已使用結(jié)果更新，

將結(jié)果寫回設(shè)備內(nèi)存。

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

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

5.3.1 設(shè)備與主機之間的數(shù)據(jù)傳輸

應(yīng)用程序應(yīng)盡量減少主機和設(shè)備之間的數(shù)據(jù)傳輸。 實現(xiàn)這一點的一種方法是將更多代碼從主機移動到設(shè)備，即使這意味著運行的內(nèi)核沒有提供足夠的并行性以在設(shè)備上全效率地執(zhí)行。 中間數(shù)據(jù)結(jié)構(gòu)可以在設(shè)備內(nèi)存中創(chuàng)建，由設(shè)備操作，并在沒有被主機映射或復(fù)制到主機內(nèi)存的情況下銷毀。

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

在具有前端總線的系統(tǒng)上，主機和設(shè)備之間的數(shù)據(jù)傳輸?shù)母咝阅苁峭ㄟ^使用頁鎖定主機內(nèi)存來實現(xiàn)的，如頁鎖定主機內(nèi)存中所述。

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

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

5.3.2 設(shè)備內(nèi)存訪問

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

全局內(nèi)存

全局內(nèi)存駐留在設(shè)備內(nèi)存中，設(shè)備內(nèi)存通過 32、64 或 128 字節(jié)內(nèi)存事務(wù)訪問。這些內(nèi)存事務(wù)必須自然對齊：只有32字節(jié)、64字節(jié)或128字節(jié)的設(shè)備內(nèi)存段按其大小對齊（即，其第一個地址是其大小的倍數(shù)）才能被內(nèi)存事務(wù)讀取或?qū)懭搿?/p>

當(dāng)一個 warp 執(zhí)行一條訪問全局內(nèi)存的指令時，它會將 warp 內(nèi)的線程的內(nèi)存訪問合并為一個或多個內(nèi)存事務(wù)，具體取決于每個線程訪問的大小以及內(nèi)存地址在整個線程中的分布。線程。一般來說，需要的事務(wù)越多，除了線程訪問的字之外，傳輸?shù)奈词褂米忠苍蕉?，相?yīng)地降低了指令吞吐量。例如，如果為每個線程的 4 字節(jié)訪問生成一個 32 字節(jié)的內(nèi)存事務(wù)，則吞吐量除以 8。

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

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

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

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

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

尺寸和對齊要求

全局內(nèi)存指令支持讀取或?qū)懭氪笮〉扔?1、2、4、8 或 16 字節(jié)的字。 當(dāng)且僅當(dāng)數(shù)據(jù)類型的大小為 1、2、4、8 或 16 字節(jié)并且數(shù)據(jù)為 對齊（即，它的地址是該大小的倍數(shù)）。

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

內(nèi)置矢量類型自動滿足對齊要求。

對于結(jié)構(gòu)，大小和對齊要求可以由編譯器使用對齊說明符 __align__（8） 或 __align__（16） 強制執(zhí)行，例如：

struct __align__（8） {

float x;

float y;

};

struct __align__（16） {

float x;

float y;

float z;

};

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

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

二維數(shù)組

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

BaseAddress + width * ty + tx

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

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

本地內(nèi)存

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

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

會占用過多寄存器空間的大型結(jié)構(gòu)或數(shù)組，

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

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

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

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

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

共享內(nèi)存

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

為了實現(xiàn)高帶寬，共享內(nèi)存被分成大小相等的內(nèi)存模塊，稱為banks，可以同時訪問。因此，可以同時處理由落在 n 個不同存儲器組中的 n 個地址構(gòu)成的任何存儲器讀取或?qū)懭胝埱螅瑥亩a(chǎn)生的總帶寬是單個模塊帶寬的 n 倍。

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

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

常量內(nèi)存

常量內(nèi)存空間駐留在設(shè)備內(nèi)存中，并緩存在常量緩存中。

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

然后在緩存命中的情況下以常量緩存的吞吐量為結(jié)果請求提供服務(wù)，否則以設(shè)備內(nèi)存的吞吐量提供服務(wù)。

紋理和表面記憶

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

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

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

尋址計算由專用單元在內(nèi)核外部執(zhí)行；

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

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

5.4最大化指令吞吐量

為了最大化指令吞吐量，應(yīng)用程序應(yīng)該：

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

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

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

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

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

5.4.1 算數(shù)指令

如下圖所示

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

一些函數(shù)的實現(xiàn)在 CUDA 頭文件（math_functions.h、device_functions.h、…）上很容易獲得。

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

Single-Precision Floating-Point Division

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

Single-Precision Floating-Point Reciprocal Square Root

為了保留 IEEE-754 語義，編譯器可以將 1.0/sqrtf（） 優(yōu)化為 rsqrtf（），僅當(dāng)?shù)箶?shù)和平方根都是近似值時（即 -prec-div=false 和 -prec-sqrt=false）。 因此，建議在需要時直接調(diào)用 rsqrtf（）。

Single-Precision Floating-Point Square Root

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

Sine and Cosine

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

更準(zhǔn)確地說，參數(shù)縮減代碼（參見實現(xiàn)的數(shù)學(xué)函數(shù)）包括兩個代碼路徑，分別稱為快速路徑和慢速路徑。

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

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

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

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

Integer Arithmetic

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

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

__［u］mul24 是不再有任何理由使用的遺留內(nèi)部函數(shù)。

Half Precision Arithmetic

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

提供了內(nèi)在的 __halves2half2 以將兩個半精度值轉(zhuǎn)換為 half2 數(shù)據(jù)類型。

提供了內(nèi)在的 __halves2bfloat162 以將兩個 __nv_bfloat 精度值轉(zhuǎn)換為 __nv_bfloat162 數(shù)據(jù)類型。

Type Conversion

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

對 char 或 short 類型的變量進行操作的函數(shù)，其操作數(shù)通常需要轉(zhuǎn)換為 int，

雙精度浮點常量（即那些沒有任何類型后綴定義的常量）用作單精度浮點計算的輸入（由 C/C++ 標(biāo)準(zhǔn)規(guī)定）。

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

5.4.2 控制流指令

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

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

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

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

5.4.3 同步指令

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

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

5.5最小化內(nèi)存抖動

經(jīng)常不斷地分配和釋放內(nèi)存的應(yīng)用程序可能會發(fā)現(xiàn)分配調(diào)用往往會隨著時間的推移而變慢，直至達(dá)到極限。這通常是由于將內(nèi)存釋放回操作系統(tǒng)供其自己使用的性質(zhì)而預(yù)期的。為了在這方面獲得最佳性能，我們建議如下：

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

嘗試在應(yīng)用程序的早期以適當(dāng)大小分配內(nèi)存，并且僅在應(yīng)用程序沒有任何用途時分配內(nèi)存。減少應(yīng)用程序中的 cudaMalloc+cudaFree 調(diào)用次數(shù)，尤其是在性能關(guān)鍵區(qū)域。

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

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

關(guān)于作者

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

審核編輯：郭婷