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

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

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

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

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

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

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

5.2 最大化利用率

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

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

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

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

5.2.2 設(shè)備層次

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

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

5.2.3 多處理器層次

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

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

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

2L 用于計(jì)算能力 6.0 的設(shè)備，因?yàn)閷?duì)于這些設(shè)備，每個(gè)周期發(fā)出的兩條指令是兩條不同warp的一條指令。

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

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

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

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

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

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

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

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

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

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

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

5.2.3.1 占用率計(jì)算

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

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

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

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

以下代碼示例計(jì)算 MyKernel 的占用率。然后，它使用并發(fā)warp與每個(gè)多處理器的最大warp之間的比率報(bào)告占用率。

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

5.3 最大化存儲(chǔ)吞吐量

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

這意味著最大限度地減少主機(jī)和設(shè)備之間的數(shù)據(jù)傳輸，如主機(jī)和設(shè)備之間的數(shù)據(jù)傳輸中所述，因?yàn)樗鼈兊膸挶热謨?nèi)存和設(shè)備之間的數(shù)據(jù)傳輸?shù)偷枚唷?/p>

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

全局內(nèi)存

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

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

需要多少事務(wù)以及最終影響多少吞吐量取決于設(shè)備的計(jì)算能力。 Compute Capability 3.x、Compute Capability 5.x、Compute Capability 6.x、Compute Capability 7.x 和 Compute Capability 8.x 提供了有關(guān)如何為各種計(jì)算能力處理全局內(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 的最佳訪問模式

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

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

尺寸和對(duì)齊要求

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

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

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

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

struct __align__（8） {

float x;

float y;

};

struct __align__（16） {

float x;

float y;

float z;

};

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

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

二維數(shù)組

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

BaseAddress + width * ty + tx

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

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

本地內(nèi)存

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

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

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

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

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

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

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

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

共享內(nèi)存

因?yàn)樗瞧系模怨蚕韮?nèi)存比本地或全局內(nèi)存具有更高的帶寬和更低的延遲。

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

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

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

常量?jī)?nèi)存

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

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

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

紋理和表面記憶

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

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

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

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

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

8 位和 16 位整數(shù)輸入數(shù)據(jù)可以選擇轉(zhuǎn)換為 ［0.0， 1.0］ 或 ［-1.0， 1.0］ 范圍內(nèi)的 32 位浮點(diǎn)值（請(qǐng)參閱紋理內(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)化同步點(diǎn)（如同步指令中所述）或使用受限指針（如 restrict 中所述）。

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

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

5.4.1 算數(shù)指令

如下圖所示

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

一些函數(shù)的實(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 用戶手冊(cè)更詳細(xì)地描述了這些編譯標(biāo)志。

Single-Precision Floating-Point Division

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

Single-Precision Floating-Point Reciprocal Square Root

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

Single-Precision Floating-Point Square Root

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

Sine and Cosine

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

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

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

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

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

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

Integer Arithmetic

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

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

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

Half Precision Arithmetic

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

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

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

Type Conversion

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

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

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

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

5.4.2 控制流指令

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

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

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

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

5.4.3 同步指令

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

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

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

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

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

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

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

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

關(guān)于作者

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

審核編輯：郭婷