CUDA編程模型的統(tǒng)一內(nèi)存

N.1. Unified Memory Introduction

統(tǒng)一內(nèi)存是 CUDA 編程模型的一個組件，在 CUDA 6.0 中首次引入，它定義了一個托管內(nèi)存空間，在該空間中所有處理器都可以看到具有公共地址空間的單個連貫內(nèi)存映像。

注意：處理器是指任何具有專用 MMU 的獨立執(zhí)行單元。這包括任何類型和架構(gòu)的 CPU 和 GPU。

底層系統(tǒng)管理 CUDA 程序中的數(shù)據(jù)訪問和位置，無需顯式內(nèi)存復(fù)制調(diào)用。這在兩個主要方面有利于 GPU 編程：

通過統(tǒng)一系統(tǒng)中所有 GPU 和 CPU 的內(nèi)存空間以及為 CUDA 程序員提供更緊密、更直接的語言集成，可以簡化 GPU 編程。

通過透明地將數(shù)據(jù)遷移到使用它的處理器，可以最大限度地提高數(shù)據(jù)訪問速度。

簡單來說，統(tǒng)一內(nèi)存消除了通過 cudaMemcpy*（） 例程進行顯式數(shù)據(jù)移動的需要，而不會因?qū)⑺袛?shù)據(jù)放入零拷貝內(nèi)存而導(dǎo)致性能損失。當(dāng)然，數(shù)據(jù)移動仍然會發(fā)生，因此程序的運行時間通常不會減少；相反，統(tǒng)一內(nèi)存可以編寫更簡單、更易于維護的代碼。

統(tǒng)一內(nèi)存提供了一個“單指針數(shù)據(jù)”模型，在概念上類似于 CUDA 的零拷貝內(nèi)存。兩者之間的一個關(guān)鍵區(qū)別在于，在零拷貝分配中，內(nèi)存的物理位置固定在 CPU 系統(tǒng)內(nèi)存中，因此程序可以快速或慢速地訪問它，具體取決于訪問它的位置。另一方面，統(tǒng)一內(nèi)存將內(nèi)存和執(zhí)行空間解耦，以便所有數(shù)據(jù)訪問都很快。

統(tǒng)一內(nèi)存一詞描述了一個為各種程序提供內(nèi)存管理服務(wù)的系統(tǒng)，從針對運行時 API 的程序到使用虛擬 ISA （PTX） 的程序。該系統(tǒng)的一部分定義了選擇加入統(tǒng)一內(nèi)存服務(wù)的托管內(nèi)存空間。

托管內(nèi)存可與特定于設(shè)備的分配互操作和互換，例如使用 cudaMalloc（） 例程創(chuàng)建的分配。所有在設(shè)備內(nèi)存上有效的 CUDA 操作在托管內(nèi)存上也有效；主要區(qū)別在于程序的主機部分也能夠引用和訪問內(nèi)存。

注意：連接到 Tegra 的離散 GPU 不支持統(tǒng)一內(nèi)存。

N.1.1. System Requirements

統(tǒng)一內(nèi)存有兩個基本要求：

具有 SM 架構(gòu) 3.0 或更高版本（Kepler 類或更高版本）的 GPU

64 位主機應(yīng)用程序和非嵌入式操作系統(tǒng)（Linux 或 Windows） 具有 SM 架構(gòu) 6.x 或更高版本（Pascal 類或更高版本）的 GPU 提供額外的統(tǒng)一內(nèi)存功能，例如本文檔中概述的按需頁面遷移和 GPU 內(nèi)存超額訂閱。 請注意，目前這些功能僅在 Linux 操作系統(tǒng)上受支持。 在 Windows 上運行的應(yīng)用程序（無論是 TCC 還是 WDDM 模式）將使用基本的統(tǒng)一內(nèi)存模型，就像在 6.x 之前的架構(gòu)上一樣，即使它們在具有 6.x 或更高計算能力的硬件上運行也是如此。 有關(guān)詳細信息，請參閱數(shù)據(jù)遷移和一致性。

N.1.2. Simplifying GPU Programming

內(nèi)存空間的統(tǒng)一意味著主機和設(shè)備之間不再需要顯式內(nèi)存?zhèn)鬏?。在托管?nèi)存空間中創(chuàng)建的任何分配都會自動遷移到需要的位置。

程序通過以下兩種方式之一分配托管內(nèi)存： 通過 cudaMallocManaged（） 例程，它在語義上類似于 cudaMalloc（）；或者通過定義一個全局 __managed__ 變量，它在語義上類似于一個 __device__ 變量。在本文檔的后面部分可以找到這些的精確定義。 注意：在具有計算能力 6.x 及更高版本的設(shè)備的支持平臺上，統(tǒng)一內(nèi)存將使應(yīng)用程序能夠使用默認系統(tǒng)分配器分配和共享數(shù)據(jù)。這允許 GPU 在不使用特殊分配器的情況下訪問整個系統(tǒng)虛擬內(nèi)存。有關(guān)更多詳細信息，請參閱系統(tǒng)分配器。 以下代碼示例說明了托管內(nèi)存的使用如何改變主機代碼的編寫方式。首先，一個沒有使用統(tǒng)一內(nèi)存的簡單程序：

__global__ void AplusB(int *ret, int a, int b) {
    ret[threadIdx.x] = a + b + threadIdx.x;
}
int main() {
    int *ret;
    cudaMalloc(&ret, 1000 * sizeof(int));
    AplusB<<< 1, 1000 >>>(ret, 10, 100);
    int *host_ret = (int *)malloc(1000 * sizeof(int));
    cudaMemcpy(host_ret, ret, 1000 * sizeof(int), cudaMemcpyDefault);
    for(int i = 0; i < 1000; i++)
        printf("%d: A+B = %d\n", i, host_ret[i]); 
    free(host_ret);
    cudaFree(ret); 
    return 0;
}

第一個示例在 GPU 上將兩個數(shù)字與每個線程 ID 組合在一起，并以數(shù)組形式返回值。 如果沒有托管內(nèi)存，則返回值的主機端和設(shè)備端存儲都是必需的（示例中為 host_ret 和 ret），使用 cudaMemcpy（） 在兩者之間顯式復(fù)制也是如此。

將此與程序的統(tǒng)一內(nèi)存版本進行比較，后者允許從主機直接訪問 GPU 數(shù)據(jù)。 請注意 cudaMallocManaged（） 例程，它從主機和設(shè)備代碼返回一個有效的指針。 這允許在沒有單獨的 host_ret 副本的情況下使用 ret，大大簡化并減小了程序的大小。

__global__ void AplusB(int *ret, int a, int b) {
    ret[threadIdx.x] = a + b + threadIdx.x;
}
int main() {
    int *ret;
    cudaMallocManaged(&ret, 1000 * sizeof(int));
    AplusB<<< 1, 1000 >>>(ret, 10, 100);
    cudaDeviceSynchronize();
    for(int i = 0; i < 1000; i++)
        printf("%d: A+B = %d\n", i, ret[i]);
    cudaFree(ret); 
    return 0;
}

最后，語言集成允許直接引用 GPU 聲明的__managed__變量，并在使用全局變量時進一步簡化程序。

__device__ __managed__ int ret[1000];
__global__ void AplusB(int a, int b) {
    ret[threadIdx.x] = a + b + threadIdx.x;
}
int main() {
    AplusB<<< 1, 1000 >>>(10, 100);
    cudaDeviceSynchronize();
    for(int i = 0; i < 1000; i++)
        printf("%d: A+B = %d\n", i, ret[i]);
    return 0;
}

請注意沒有明確的 cudaMemcpy（） 命令以及返回數(shù)組 ret 在 CPU 和 GPU 上都可見的事實。

值得一提的是主機和設(shè)備之間的同步。 請注意在非托管示例中，同步 cudaMemcpy（） 例程如何用于同步內(nèi)核（即等待它完成運行）以及將數(shù)據(jù)傳輸?shù)街鳈C。 統(tǒng)一內(nèi)存示例不調(diào)用 cudaMemcpy（），因此需要顯式 cudaDeviceSynchronize（），然后主機程序才能安全地使用 GPU 的輸出。

N.1.3. Data Migration and Coherency

統(tǒng)一內(nèi)存嘗試通過將數(shù)據(jù)遷移到正在訪問它的設(shè)備來優(yōu)化內(nèi)存性能（也就是說，如果 CPU 正在訪問數(shù)據(jù)，則將數(shù)據(jù)移動到主機內(nèi)存，如果 GPU 將訪問它，則將數(shù)據(jù)移動到設(shè)備內(nèi)存）。數(shù)據(jù)遷移是統(tǒng)一內(nèi)存的基礎(chǔ)，但對程序是透明的。系統(tǒng)將嘗試將數(shù)據(jù)放置在可以最有效地訪問而不違反一致性的位置。

數(shù)據(jù)的物理位置對程序是不可見的，并且可以隨時更改，但對數(shù)據(jù)的虛擬地址的訪問將保持有效并且可以從任何處理器保持一致，無論位置如何。請注意，保持一致性是首要要求，高于性能；在主機操作系統(tǒng)的限制下，系統(tǒng)被允許訪問失敗或移動數(shù)據(jù)，以保持處理器之間的全局一致性。

計算能力低于 6.x 的 GPU 架構(gòu)不支持按需將托管數(shù)據(jù)細粒度移動到 GPU。每當(dāng)啟動 GPU 內(nèi)核時，通常必須將所有托管內(nèi)存轉(zhuǎn)移到 GPU 內(nèi)存，以避免內(nèi)存訪問出錯。計算能力 6.x 引入了一種新的 GPU 頁面錯誤機制，可提供更無縫的統(tǒng)一內(nèi)存功能。結(jié)合系統(tǒng)范圍的虛擬地址空間，頁面錯誤提供了幾個好處。首先，頁面錯誤意味著 CUDA 系統(tǒng)軟件不需要在每次內(nèi)核啟動之前將所有托管內(nèi)存分配同步到 GPU。如果在 GPU 上運行的內(nèi)核訪問了一個不在其內(nèi)存中的頁面，它就會出錯，從而允許該頁面按需自動遷移到 GPU 內(nèi)存。或者，可以將頁面映射到 GPU 地址空間，以便通過 PCIe 或 NVLink 互連進行訪問（訪問映射有時可能比遷移更快）。請注意，統(tǒng)一內(nèi)存是系統(tǒng)范圍的：GPU（和 CPU）可以從 CPU 內(nèi)存或系統(tǒng)中其他 GPU 的內(nèi)存中發(fā)生故障并遷移內(nèi)存頁面。

N.1.4. GPU Memory Oversubscription

計算能力低于 6.x 的設(shè)備分配的托管內(nèi)存不能超過 GPU 內(nèi)存的物理大小。

計算能力 6.x 的設(shè)備擴展了尋址模式以支持 49 位虛擬尋址。 這足以覆蓋現(xiàn)代 CPU 的 48 位虛擬地址空間，以及 GPU 自己的內(nèi)存。 大的虛擬地址空間和頁面錯誤能力使應(yīng)用程序可以訪問整個系統(tǒng)的虛擬內(nèi)存，而不受任何一個處理器的物理內(nèi)存大小的限制。 這意味著應(yīng)用程序可以超額訂閱內(nèi)存系統(tǒng)：換句話說，它們可以分配、訪問和共享大于系統(tǒng)總物理容量的數(shù)組，從而實現(xiàn)超大數(shù)據(jù)集的核外處理。 只要有足夠的系統(tǒng)內(nèi)存可用于分配，cudaMallocManaged 就不會耗盡內(nèi)存。

N.1.5. Multi-GPU

對于計算能力低于 6.x 的設(shè)備，托管內(nèi)存分配的行為與使用 cudaMalloc（） 分配的非托管內(nèi)存相同：當(dāng)前活動設(shè)備是物理分配的主站，所有其他 GPU 接收到內(nèi)存的對等映射。這意味著系統(tǒng)中的其他 GPU 將以較低的帶寬通過 PCIe 總線訪問內(nèi)存。請注意，如果系統(tǒng)中的 GPU 之間不支持對等映射，則托管內(nèi)存頁面將放置在 CPU 系統(tǒng)內(nèi)存（“零拷貝”內(nèi)存）中，并且所有 GPU 都會遇到 PCIe 帶寬限制。有關(guān)詳細信息，請參閱 6.x 之前架構(gòu)上的多 GPU 程序的托管內(nèi)存。

具有計算能力 6.x 設(shè)備的系統(tǒng)上的托管分配對所有 GPU 都是可見的，并且可以按需遷移到任何處理器。統(tǒng)一內(nèi)存性能提示（請參閱性能調(diào)優(yōu)）允許開發(fā)人員探索自定義使用模式，例如跨 GPU 讀取重復(fù)數(shù)據(jù)和直接訪問對等 GPU 內(nèi)存而無需遷移。

N.1.6. System Allocator

計算能力 7.0 的設(shè)備支持 NVLink 上的地址轉(zhuǎn)換服務(wù) （ATS）。 如果主機 CPU 和操作系統(tǒng)支持，ATS 允許 GPU 直接訪問 CPU 的頁表。 GPU MMU 中的未命中將導(dǎo)致向 CPU 發(fā)送地址轉(zhuǎn)換請求 （ATR）。 CPU 在其頁表中查找該地址的虛擬到物理映射并將轉(zhuǎn)換提供回 GPU。 ATS 提供 GPU 對系統(tǒng)內(nèi)存的完全訪問權(quán)限，例如使用 malloc 分配的內(nèi)存、在堆棧上分配的內(nèi)存、全局變量和文件支持的內(nèi)存。 應(yīng)用程序可以通過檢查新的 pageableMemoryAccessUsesHostPageTables 屬性來查詢設(shè)備是否支持通過 ATS 一致地訪問可分頁內(nèi)存。

這是一個適用于任何滿足統(tǒng)一內(nèi)存基本要求的系統(tǒng)的示例代碼（請參閱系統(tǒng)要求）：

int *data;
cudaMallocManaged(&data, sizeof(int) * n);
kernel<<>>(data);

具有 pageableMemoryAccess 屬性的系統(tǒng)支持這些新的訪問模式：

int *data = (int*)malloc(sizeof(int) * n);
kernel<<>>(data);

int data[1024];
kernel<<>>(data);

extern int *data;
kernel<<>>(data);

在上面的示例中，數(shù)據(jù)可以由第三方 CPU 庫初始化，然后由 GPU 內(nèi)核直接訪問。 在具有 pageableMemoryAccess 的系統(tǒng)上，用戶還可以使用 cudaMemPrefetchAsync 將可分頁內(nèi)存預(yù)取到 GPU。 這可以通過優(yōu)化數(shù)據(jù)局部性產(chǎn)生性能優(yōu)勢。

注意：目前僅 IBM Power9 系統(tǒng)支持基于 NVLink 的 ATS。

N.1.7. Hardware Coherency

第二代 NVLink 允許從 CPU 直接加載/存儲/原子訪問每個 GPU 的內(nèi)存。結(jié)合新的 CPU 主控功能，NVLink 支持一致性操作，允許從 GPU 內(nèi)存讀取的數(shù)據(jù)存儲在 CPU 的緩存層次結(jié)構(gòu)中。從 CPU 緩存訪問的較低延遲是 CPU 性能的關(guān)鍵。計算能力 6.x 的設(shè)備僅支持對等 GPU 原子。計算能力 7.x 的設(shè)備可以通過 NVLink 發(fā)送 GPU 原子并在目標(biāo) CPU 上完成它們，因此第二代 NVLink 增加了對由 GPU 或 CPU 發(fā)起的原子的支持。

請注意，CPU 無法訪問 cudaMalloc 分配。因此，要利用硬件一致性，用戶必須使用統(tǒng)一內(nèi)存分配器，例如 cudaMallocManaged 或支持 ATS 的系統(tǒng)分配器（請參閱系統(tǒng)分配器）。新屬性 directManagedMemAccessFromHost 指示主機是否可以直接訪問設(shè)備上的托管內(nèi)存而無需遷移。默認情況下，駐留在 GPU 內(nèi)存中的 cudaMallocManaged 分配的任何 CPU 訪問都會觸發(fā)頁面錯誤和數(shù)據(jù)遷移。應(yīng)用程序可以使用帶有 cudaCpuDeviceId 的 cudaMemAdviseSetAccessedBy 性能提示來啟用對受支持系統(tǒng)上 GPU 內(nèi)存的直接訪問。

考慮下面的示例代碼：

__global__ void write(int *ret, int a, int b) {
    ret[threadIdx.x] = a + b + threadIdx.x;
}
__global__ void append(int *ret, int a, int b) {
    ret[threadIdx.x] += a + b + threadIdx.x;
}
int main() {
    int *ret;
    cudaMallocManaged(&ret, 1000 * sizeof(int));
    cudaMemAdvise(ret, 1000 * sizeof(int), cudaMemAdviseSetAccessedBy, cudaCpuDeviceId);  // set direct access hint

    write<<< 1, 1000 >>>(ret, 10, 100);            // pages populated in GPU memory
    cudaDeviceSynchronize();
    for(int i = 0; i < 1000; i++)
        printf("%d: A+B = %d\n", i, ret[i]);        // directManagedMemAccessFromHost=1: CPU accesses GPU memory directly without migrations
                                                    // directManagedMemAccessFromHost=0: CPU faults and triggers device-to-host migrations
    append<<< 1, 1000 >>>(ret, 10, 100);            // directManagedMemAccessFromHost=1: GPU accesses GPU memory without migrations
    cudaDeviceSynchronize();                        // directManagedMemAccessFromHost=0: GPU faults and triggers host-to-device migrations
    cudaFree(ret); 
    return 0;
}

寫內(nèi)核完成后，會在GPU內(nèi)存中創(chuàng)建并初始化ret。 接下來，CPU 將訪問 ret，然后再次使用相同的 ret 內(nèi)存追加內(nèi)核。 此代碼將根據(jù)系統(tǒng)架構(gòu)和硬件一致性支持顯示不同的行為：

在 directManagedMemAccessFromHost=1 的系統(tǒng)上：CPU 訪問托管緩沖區(qū)不會觸發(fā)任何遷移； 數(shù)據(jù)將保留在 GPU 內(nèi)存中，任何后續(xù)的 GPU 內(nèi)核都可以繼續(xù)直接訪問它，而不會造成故障或遷移。

在 directManagedMemAccessFromHost=0 的系統(tǒng)上：CPU 訪問托管緩沖區(qū)將出現(xiàn)頁面錯誤并啟動數(shù)據(jù)遷移； 任何第一次嘗試訪問相同數(shù)據(jù)的 GPU 內(nèi)核都會出現(xiàn)頁面錯誤并將頁面遷移回 GPU 內(nèi)存。

N.1.8. Access Counters

計算能力 7.0 的設(shè)備引入了一個新的訪問計數(shù)器功能，該功能可以跟蹤 GPU 對位于其他處理器上的內(nèi)存進行的訪問頻率。 訪問計數(shù)器有助于確保將內(nèi)存頁面移動到最頻繁訪問頁面的處理器的物理內(nèi)存中。 訪問計數(shù)器功能可以指導(dǎo) CPU 和 GPU 之間以及對等 GPU 之間的遷移。

對于 cudaMallocManaged，訪問計數(shù)器遷移可以通過使用帶有相應(yīng)設(shè)備 ID 的 cudaMemAdviseSetAccessedBy 提示來選擇加入。 驅(qū)動程序還可以使用訪問計數(shù)器來實現(xiàn)更有效的抖動緩解或內(nèi)存超額訂閱方案。

注意：訪問計數(shù)器當(dāng)前僅在 IBM Power9 系統(tǒng)上啟用，并且僅用于 cudaMallocManaged 分配器。

N.2. Programming Model

N.2.1. Managed Memory Opt In

大多數(shù)平臺要求程序通過使用 __managed__ 關(guān)鍵字注釋 __device__ 變量（請參閱語言集成部分）或使用新的 cudaMallocManaged（） 調(diào)用來分配數(shù)據(jù)來選擇自動數(shù)據(jù)管理。

計算能力低于 6.x 的設(shè)備必須始終在堆上分配托管內(nèi)存，無論是使用分配器還是通過聲明全局存儲。 無法將先前分配的內(nèi)存與統(tǒng)一內(nèi)存相關(guān)聯(lián)，也無法讓統(tǒng)一內(nèi)存系統(tǒng)管理 CPU 或 GPU 堆棧指針。

從 CUDA 8.0 和具有計算能力 6.x 設(shè)備的支持系統(tǒng)開始，可以使用相同的指針從 GPU 代碼和 CPU 代碼訪問使用默認 OS 分配器（例如 malloc 或 new）分配的內(nèi)存。 在這些系統(tǒng)上，統(tǒng)一內(nèi)存是默認設(shè)置：無需使用特殊分配器或創(chuàng)建專門管理的內(nèi)存池。

N.2.1.1. Explicit Allocation Using cudaMallocManaged（）

統(tǒng)一內(nèi)存最常使用在語義和語法上類似于標(biāo)準(zhǔn) CUDA 分配器 cudaMalloc（） 的分配函數(shù)創(chuàng)建。 功能說明如下：

    cudaError_t cudaMallocManaged(void **devPtr,
                                  size_t size,
                                  unsigned int flags=0);

cudaMallocManaged（） 函數(shù)保留托管內(nèi)存的 size 字節(jié)，并在 devPtr 中返回一個指針。 請注意各種 GPU 架構(gòu)之間 cudaMallocManaged（） 行為的差異。 默認情況下，計算能力低于 6.x 的設(shè)備直接在 GPU 上分配托管內(nèi)存。 但是，計算能力 6.x 及更高版本的設(shè)備在調(diào)用 cudaMallocManaged（） 時不會分配物理內(nèi)存：在這種情況下，物理內(nèi)存會在第一次觸摸時填充，并且可能駐留在 CPU 或 GPU 上。 托管指針在系統(tǒng)中的所有 GPU 和 CPU 上都有效，盡管程序訪問此指針必須遵守統(tǒng)一內(nèi)存編程模型的并發(fā)規(guī)則（請參閱一致性和并發(fā)性）。 下面是一個簡單的例子，展示了 cudaMallocManaged（） 的使用：

__global__ void printme(char *str) {
    printf(str);
}
int main() {
    // Allocate 100 bytes of memory, accessible to both Host and Device code
    char *s;
    cudaMallocManaged(&s, 100);
    // Note direct Host-code use of "s"
    strncpy(s, "Hello Unified Memory\n", 99);
    // Here we pass "s" to a kernel without explicitly copying
    printme<<< 1, 1 >>>(s);
    cudaDeviceSynchronize();
    // Free as for normal CUDA allocations
    cudaFree(s); 
    return  0;
}

當(dāng) cudaMalloc（） 被 cudaMallocManaged（） 替換時，程序的行為在功能上沒有改變； 但是，該程序應(yīng)該繼續(xù)消除顯式內(nèi)存拷貝并利用自動遷移。 此外，可以消除雙指針（一個指向主機，一個指向設(shè)備存儲器）。

設(shè)備代碼無法調(diào)用 cudaMallocManaged（）。 所有托管內(nèi)存必須從主機或全局范圍內(nèi)分配（請參閱下一節(jié)）。 在內(nèi)核中使用 malloc（） 在設(shè)備堆上的分配不會在托管內(nèi)存空間中創(chuàng)建，因此 CPU 代碼將無法訪問。

N.2.1.2. Global-Scope Managed Variables Using managed

文件范圍和全局范圍的 CUDA __device__ 變量也可以通過在聲明中添加新的 __managed__ 注釋來選擇加入統(tǒng)一內(nèi)存管理。 然后可以直接從主機或設(shè)備代碼中引用它們，如下所示：

釋來選擇加入統(tǒng)一內(nèi)存管理。 然后可以直接從主機或設(shè)備代碼中引用它們，如下所示：

__device__ __managed__ int x[2];
__device__ __managed__ int y;
__global__ void kernel() {
    x[1] = x[0] + y;
}
int main() {
    x[0] = 3;
    y = 5;
    kernel<<< 1, 1 >>>();
    cudaDeviceSynchronize();
    printf("result = %d\n", x[1]); 
    return  0;
}

原始 __device__ 內(nèi)存空間的所有語義，以及一些額外的統(tǒng)一內(nèi)存特定約束，都由托管變量繼承（請參閱使用 NVCC 編譯）。

請注意，標(biāo)記為 __constant__ 的變量可能不會也標(biāo)記為 __managed__； 此注釋僅用于 __device__ 變量。 常量內(nèi)存必須在編譯時靜態(tài)設(shè)置，或者在 CUDA 中像往常一樣使用 cudaMemcpyToSymbol（） 設(shè)置。

N.2.2. Coherency and Concurrency

在計算能力低于 6.x 的設(shè)備上同時訪問托管內(nèi)存是不可能的，因為如果 CPU 在 GPU 內(nèi)核處于活動狀態(tài)時訪問統(tǒng)一內(nèi)存分配，則無法保證一致性。 但是，支持操作系統(tǒng)的計算能力 6.x 的設(shè)備允許 CPU 和 GPU 通過新的頁面錯誤機制同時訪問統(tǒng)一內(nèi)存分配。 程序可以通過檢查新的 concurrentManagedAccess 屬性來查詢設(shè)備是否支持對托管內(nèi)存的并發(fā)訪問。 請注意，與任何并行應(yīng)用程序一樣，開發(fā)人員需要確保正確同步以避免處理器之間的數(shù)據(jù)危險。

N.2.2.1. GPU Exclusive Access To Managed Memory

為了確保 6.x 之前的 GPU 架構(gòu)的一致性，統(tǒng)一內(nèi)存編程模型在 CPU 和 GPU 同時執(zhí)行時對數(shù)據(jù)訪問施加了限制。實際上，GPU 在執(zhí)行任何內(nèi)核操作時對所有托管數(shù)據(jù)具有獨占訪問權(quán)，無論特定內(nèi)核是否正在積極使用數(shù)據(jù)。當(dāng)托管數(shù)據(jù)與 cudaMemcpy*（） 或 cudaMemset*（） 一起使用時，系統(tǒng)可能會選擇從主機或設(shè)備訪問源或目標(biāo)，這將限制并發(fā) CPU 訪問該數(shù)據(jù)，而 cudaMemcpy*（）或 cudaMemset*（） 正在執(zhí)行。有關(guān)更多詳細信息，請參閱使用托管內(nèi)存的 Memcpy（）/Memset（） 行為。

不允許 CPU 訪問任何托管分配或變量，而 GPU 對 concurrentManagedAccess 屬性設(shè)置為 0 的設(shè)備處于活動狀態(tài)。在這些系統(tǒng)上，并發(fā) CPU/GPU 訪問，即使是不同的托管內(nèi)存分配，也會導(dǎo)致分段錯誤，因為該頁面被認為是 CPU 無法訪問的。

__device__ __managed__ int x, y=2;
__global__  void  kernel() {
    x = 10;
}
int main() {
    kernel<<< 1, 1 >>>();
    y = 20;            // Error on GPUs not supporting concurrent access
                       
    cudaDeviceSynchronize();
    return  0;
}

在上面的示例中，當(dāng) CPU 接觸(這里原文中用的是touch這個詞) y 時，GPU 程序內(nèi)核仍然處于活動狀態(tài)。 （注意它是如何在cudaDeviceSynchronize()之前發(fā)生的。）由于 GPU 頁面錯誤功能解除了對同時訪問的所有限制，因此代碼在計算能力 6.x 的設(shè)備上成功運行。 但是，即使 CPU 訪問的數(shù)據(jù)與 GPU 不同，這種內(nèi)存訪問在 6.x 之前的架構(gòu)上也是無效的。 程序必須在訪問 y 之前顯式地與 GPU 同步：

__device__ __managed__ int x, y=2;
__global__  void  kernel() {
    x = 10;
}
int main() {
    kernel<<< 1, 1 >>>();
    cudaDeviceSynchronize();
    y = 20;            //  Success on GPUs not supporing concurrent access
    return  0;
}

如本例所示，在具有 6.x 之前的 GPU 架構(gòu)的系統(tǒng)上，CPU 線程可能不會在執(zhí)行內(nèi)核啟動和后續(xù)同步調(diào)用之間訪問任何托管數(shù)據(jù)，無論 GPU 內(nèi)核是否實際接觸相同的數(shù)據(jù)（或 任何托管數(shù)據(jù)）。 并發(fā) CPU 和 GPU 訪問的潛力足以引發(fā)進程級異常。

請注意，如果在 GPU 處于活動狀態(tài)時使用 cudaMallocManaged（） 或 cuMemAllocManaged（） 動態(tài)分配內(nèi)存，則在啟動其他工作或同步 GPU 之前，內(nèi)存的行為是未指定的。 在此期間嘗試訪問 CPU 上的內(nèi)存可能會也可能不會導(dǎo)致分段錯誤。 這不適用于使用標(biāo)志 cudaMemAttachHost 或 CU_MEM_ATTACH_HOST 分配的內(nèi)存。

N.2.2.2. Explicit Synchronization and Logical GPU Activity

請注意，即使內(nèi)核快速運行并在上例中的 CPU 接觸 y 之前完成，也需要顯式同步。統(tǒng)一內(nèi)存使用邏輯活動來確定 GPU 是否空閑。這與 CUDA 編程模型一致，該模型指定內(nèi)核可以在啟動后的任何時間運行，并且不保證在主機發(fā)出同步調(diào)用之前完成。

任何在邏輯上保證 GPU 完成其工作的函數(shù)調(diào)用都是有效的。這包括 cudaDeviceSynchronize（）; cudaStreamSynchronize（） 和 cudaStreamQuery（）（如果它返回 cudaSuccess 而不是 cudaErrorNotReady），其中指定的流是唯一仍在 GPU 上執(zhí)行的流； cudaEventSynchronize（） 和 cudaEventQuery（） 在指定事件之后沒有任何設(shè)備工作的情況下；以及記錄為與主機完全同步的 cudaMemcpy（） 和 cudaMemset（） 的使用。

將遵循流之間創(chuàng)建的依賴關(guān)系，通過在流或事件上同步來推斷其他流的完成。依賴關(guān)系可以通過 cudaStreamWaitEvent（） 或在使用默認 （NULL） 流時隱式創(chuàng)建。

CPU 從流回調(diào)中訪問托管數(shù)據(jù)是合法的，前提是 GPU 上沒有其他可能訪問托管數(shù)據(jù)的流處于活動狀態(tài)。此外，沒有任何設(shè)備工作的回調(diào)可用于同步：例如，通過從回調(diào)內(nèi)部發(fā)出條件變量的信號；否則，CPU 訪問僅在回調(diào)期間有效。

有幾個重要的注意點：

在 GPU 處于活動狀態(tài)時，始終允許 CPU 訪問非托管零拷貝數(shù)據(jù)。

GPU 在運行任何內(nèi)核時都被認為是活動的，即使該內(nèi)核不使用托管數(shù)據(jù)。如果內(nèi)核可能使用數(shù)據(jù)，則禁止訪問，除非設(shè)備屬性 concurrentManagedAccess 為 1。

除了適用于非托管內(nèi)存的多 GPU 訪問之外，托管內(nèi)存的并發(fā) GPU 間訪問沒有任何限制。

并發(fā) GPU 內(nèi)核訪問托管數(shù)據(jù)沒有任何限制。

請注意最后一點如何允許 GPU 內(nèi)核之間的競爭，就像當(dāng)前非托管 GPU 內(nèi)存的情況一樣。如前所述，從 GPU 的角度來看，托管內(nèi)存的功能與非托管內(nèi)存相同。以下代碼示例說明了這些要點：

int main() {
    cudaStream_t stream1, stream2;
    cudaStreamCreate(&stream1);
    cudaStreamCreate(&stream2);
    int *non_managed, *managed, *also_managed;
    cudaMallocHost(&non_managed, 4);    // Non-managed, CPU-accessible memory
    cudaMallocManaged(&managed, 4);
    cudaMallocManaged(&also_managed, 4);
    // Point 1: CPU can access non-managed data.
    kernel<<< 1, 1, 0, stream1 >>>(managed);
    *non_managed = 1;
    // Point 2: CPU cannot access any managed data while GPU is busy,
    //          unless concurrentManagedAccess = 1
    // Note we have not yet synchronized, so "kernel" is still active.
    *also_managed = 2;      // Will issue segmentation fault
    // Point 3: Concurrent GPU kernels can access the same data.
    kernel<<< 1, 1, 0, stream2 >>>(managed);
    // Point 4: Multi-GPU concurrent access is also permitted.
    cudaSetDevice(1);
    kernel<<< 1, 1 >>>(managed);
    return  0;
}

N.2.2.3. Managing Data Visibility and Concurrent CPU + GPU Access with Streams

到目前為止，假設(shè)對于 6.x 之前的 SM 架構(gòu)：1） 任何活動內(nèi)核都可以使用任何托管內(nèi)存，以??及 2） 在內(nèi)核處于活動狀態(tài)時使用來自 CPU 的托管內(nèi)存是無效的。在這里，我們提出了一個用于對托管內(nèi)存進行更細粒度控制的系統(tǒng)，該系統(tǒng)旨在在所有支持托管內(nèi)存的設(shè)備上工作，包括 concurrentManagedAccess 等于 0 的舊架構(gòu)。

CUDA 編程模型提供流作為程序指示內(nèi)核啟動之間的依賴性和獨立性的機制。啟動到同一流中的內(nèi)核保證連續(xù)執(zhí)行，而啟動到不同流中的內(nèi)核允許并發(fā)執(zhí)行。流描述了工作項之間的獨立性，因此可以通過并發(fā)實現(xiàn)更高的效率。

統(tǒng)一內(nèi)存建立在流獨立模型之上，允許 CUDA 程序顯式地將托管分配與 CUDA 流相關(guān)聯(lián)。通過這種方式，程序員根據(jù)內(nèi)核是否將數(shù)據(jù)啟動到指定的流中來指示內(nèi)核對數(shù)據(jù)的使用。這為基于程序特定數(shù)據(jù)訪問模式的并發(fā)提供了機會?？刂七@種行為的函數(shù)是：

    cudaError_t cudaStreamAttachMemAsync(cudaStream_t stream,
                                         void *ptr,
                                         size_t length=0,
                                         unsigned int flags=0);

前，length 必須始終為 0 以指示應(yīng)該附加整個區(qū)域。）由于這種關(guān)聯(lián)，只要流中的所有操作都已完成，統(tǒng)一內(nèi)存系統(tǒng)就允許 CPU 訪問該內(nèi)存區(qū)域，而不管其他流是否是活躍的。實際上，這將活動 GPU 對托管內(nèi)存區(qū)域的獨占所有權(quán)限制為每個流活動而不是整個 GPU 活動。

最重要的是，如果分配與特定流無關(guān)，則所有正在運行的內(nèi)核都可以看到它，而不管它們的流如何。這是 cudaMallocManaged（） 分配或 __managed__ 變量的默認可見性；因此，在任何內(nèi)核運行時 CPU 不得接觸數(shù)據(jù)的簡單案例規(guī)則。

通過將分配與特定流相關(guān)聯(lián)，程序保證只有啟動到該流中的內(nèi)核才會接觸該數(shù)據(jù)。統(tǒng)一內(nèi)存系統(tǒng)不執(zhí)行錯誤檢查：程序員有責(zé)任確保兌現(xiàn)保證。

除了允許更大的并發(fā)性之外，使用 cudaStreamAttachMemAsync（） 可以（并且通常會）啟用統(tǒng)一內(nèi)存系統(tǒng)內(nèi)的數(shù)據(jù)傳輸優(yōu)化，這可能會影響延遲和其他開銷。

N.2.2.4. Stream Association Examples

將數(shù)據(jù)與流相關(guān)聯(lián)允許對 CPU + GPU 并發(fā)進行細粒度控制，但在使用計算能力低于 6.x 的設(shè)備時，必須牢記哪些數(shù)據(jù)對哪些流可見。 查看前面的同步示例：

__device__ __managed__ int x, y=2;
__global__  void  kernel() {
    x = 10;
}
int main() {
    cudaStream_t stream1;
    cudaStreamCreate(&stream1);
    cudaStreamAttachMemAsync(stream1, &y, 0, cudaMemAttachHost);
    cudaDeviceSynchronize();          // Wait for Host attachment to occur.
    kernel<<< 1, 1, 0, stream1 >>>(); // Note: Launches into stream1.
    y = 20;                           // Success – a kernel is running but “y” 
                                      // has been associated with no stream.
    return  0;
}

在這里，我們明確地將 y 與主機可訪問性相關(guān)聯(lián)，從而始終可以從 CPU 進行訪問。 （和以前一樣，請注意在訪問之前沒有cudaDeviceSynchronize()。）GPU 運行內(nèi)核對 y 的訪問現(xiàn)在將產(chǎn)生未定義的結(jié)果。

請注意，將變量與流關(guān)聯(lián)不會更改任何其他變量的關(guān)聯(lián)。 例如。 將 x 與 stream1 關(guān)聯(lián)并不能確保在 stream1 中啟動的內(nèi)核只能訪問 x，因此此代碼會導(dǎo)致錯誤：

__device__ __managed__ int x, y=2;
__global__  void  kernel() {
    x = 10;
}
int main() {
    cudaStream_t stream1;
    cudaStreamCreate(&stream1);
    cudaStreamAttachMemAsync(stream1, &x);// Associate “x” with stream1.
    cudaDeviceSynchronize();              // Wait for “x” attachment to occur.
    kernel<<< 1, 1, 0, stream1 >>>();     // Note: Launches into stream1.
    y = 20;                               // ERROR: “y” is still associated globally 
                                          // with all streams by default
    return  0;
}

請注意訪問 y 將如何導(dǎo)致錯誤，因為即使 x 已與流相關(guān)聯(lián)，我們也沒有告訴系統(tǒng)誰可以看到 y。 因此，系統(tǒng)保守地假設(shè)內(nèi)核可能會訪問它并阻止 CPU 這樣做。

N.2.2.5. Stream Attach With Multithreaded Host Programs

cudaStreamAttachMemAsync（） 的主要用途是使用 CPU 線程啟用獨立任務(wù)并行性。 通常在這樣的程序中，CPU 線程為它生成的所有工作創(chuàng)建自己的流，因為使用 CUDA 的 NULL 流會導(dǎo)致線程之間的依賴關(guān)系。

托管數(shù)據(jù)對任何 GPU 流的默認全局可見性使得難以避免多線程程序中 CPU 線程之間的交互。 因此，函數(shù) cudaStreamAttachMemAsync（） 用于將線程的托管分配與該線程自己的流相關(guān)聯(lián)，并且該關(guān)聯(lián)通常在線程的生命周期內(nèi)不會更改。

這樣的程序?qū)⒑唵蔚靥砑右粋€對 cudaStreamAttachMemAsync（） 的調(diào)用，以使用統(tǒng)一內(nèi)存進行數(shù)據(jù)訪問：

// This function performs some task, in its own private stream.
void run_task(int *in, int *out, int length) {
    // Create a stream for us to use.
    cudaStream_t stream;
    cudaStreamCreate(&stream);
    // Allocate some managed data and associate with our stream.
    // Note the use of the host-attach flag to cudaMallocManaged();
    // we then associate the allocation with our stream so that
    // our GPU kernel launches can access it.
    int *data;
    cudaMallocManaged((void **)&data, length, cudaMemAttachHost);
    cudaStreamAttachMemAsync(stream, data);
    cudaStreamSynchronize(stream);
    // Iterate on the data in some way, using both Host & Device.
    for(int i=0; i>>(in, data, length);
        cudaStreamSynchronize(stream);
        host_process(data, length);    // CPU uses managed data.
        convert<<< 100, 256, 0, stream >>>(out, data, length);
    }
    cudaStreamSynchronize(stream);
    cudaStreamDestroy(stream);
    cudaFree(data);
}

在這個例子中，分配流關(guān)聯(lián)只建立一次，然后主機和設(shè)備都重復(fù)使用數(shù)據(jù)。 結(jié)果是比在主機和設(shè)備之間顯式復(fù)制數(shù)據(jù)時更簡單的代碼，盡管結(jié)果是相同的。

N.2.2.6. Advanced Topic： Modular Programs and Data Access Constraints

在前面的示例中，cudaMallocManaged（） 指定了 cudaMemAttachHost 標(biāo)志，它創(chuàng)建了一個最初對設(shè)備端執(zhí)行不可見的分配。 （默認分配對所有流上的所有 GPU 內(nèi)核都是可見的。）這可確保在數(shù)據(jù)分配和為特定流獲取數(shù)據(jù)之間的時間間隔內(nèi)，不會與另一個線程的執(zhí)行發(fā)生意外交互。

如果沒有這個標(biāo)志，如果另一個線程啟動的內(nèi)核恰好正在運行，則新分配將被視為在 GPU 上使用。這可能會影響線程在能夠?qū)⑵滹@式附加到私有流之前從 CPU 訪問新分配的數(shù)據(jù)的能力（例如，在基類構(gòu)造函數(shù)中）。因此，為了啟用線程之間的安全獨立性，應(yīng)指定此標(biāo)志進行分配。

注意：另一種方法是在分配附加到流之后在所有線程上放置一個進程范圍的屏障。這將確保所有線程在啟動任何內(nèi)核之前完成其數(shù)據(jù)/流關(guān)聯(lián)，從而避免危險。在銷毀流之前需要第二個屏障，因為流銷毀會導(dǎo)致分配恢復(fù)到其默認可見性。 cudaMemAttachHost 標(biāo)志的存在既是為了簡化此過程，也是因為并非總是可以在需要的地方插入全局屏障。

N.2.2.7. Memcpy（）/Memset（） Behavior With Managed Memory

由于可以從主機或設(shè)備訪問托管內(nèi)存，因此 cudaMemcpy*（） 依賴于使用 cudaMemcpyKind 指定的傳輸類型來確定數(shù)據(jù)應(yīng)該作為主機指針還是設(shè)備指針訪問。

如果指定了 cudaMemcpyHostTo* 并且管理了源數(shù)據(jù)，那么如果在復(fù)制流 （1） 中可以從主機連貫地訪問它，那么它將從主機訪問；否則將從設(shè)備訪問。當(dāng)指定 cudaMemcpy*ToHost 并且目標(biāo)是托管內(nèi)存時，類似的規(guī)則適用于目標(biāo)。

如果指定了 cudaMemcpyDeviceTo* 并管理源數(shù)據(jù)，則將從設(shè)備訪問它。源必須可以從復(fù)制流中的設(shè)備連貫地訪問 （2）；否則，返回錯誤。當(dāng)指定 cudaMemcpy*ToDevice 并且目標(biāo)是托管內(nèi)存時，類似的規(guī)則適用于目標(biāo)。

如果指定了 cudaMemcpyDefault，則如果無法從復(fù)制流中的設(shè)備一致地訪問托管數(shù)據(jù) （2），或者如果數(shù)據(jù)的首選位置是 cudaCpuDeviceId 并且可以從主機一致地訪問，則將從主機訪問托管數(shù)據(jù)在復(fù)制流 （1） 中；否則，它將從設(shè)備訪問。

將 cudaMemset*（） 與托管內(nèi)存一起使用時，始終從設(shè)備訪問數(shù)據(jù)。數(shù)據(jù)必須可以從用于 cudaMemset（） 操作的流中的設(shè)備連貫地訪問 （2）；否則，返回錯誤。

當(dāng)通過 cudaMemcpy* 或 cudaMemset* 從設(shè)備訪問數(shù)據(jù)時，操作流被視為在 GPU 上處于活動狀態(tài)。在此期間，如果 GPU 的設(shè)備屬性 concurrentManagedAccess 為零值，則任何與該流相關(guān)聯(lián)的數(shù)據(jù)或具有全局可見性的數(shù)據(jù)的 CPU 訪問都將導(dǎo)致分段錯誤。在從 CPU 訪問任何相關(guān)數(shù)據(jù)之前，程序必須適當(dāng)同步以確保操作已完成。

（1） 要在給定流中從主機連貫地訪問托管內(nèi)存，必須至少滿足以下條件之一：

給定流與設(shè)備屬性 concurrentManagedAccess 具有非零值的設(shè)備相關(guān)聯(lián)。

內(nèi)存既不具有全局可見性，也不與給定流相關(guān)聯(lián)。

（2） 要在給定流中從設(shè)備連貫地訪問托管內(nèi)存，必須至少滿足以下條件之一：

設(shè)備的設(shè)備屬性 concurrentManagedAccess 具有非零值。

內(nèi)存要么具有全局可見性，要么與給定的流相關(guān)聯(lián)。

###N.2.3. Language Integration

使用 nvcc 編譯主機代碼的 CUDA 運行時 API 用戶可以訪問其他語言集成功能，例如共享符號名稱和通過 《《《。..》》》 運算符啟動內(nèi)聯(lián)內(nèi)核。 統(tǒng)一內(nèi)存為 CUDA 的語言集成添加了一個附加元素：使用 __managed__ 關(guān)鍵字注釋的變量可以直接從主機和設(shè)備代碼中引用。

下面的例子在前面的 Simplifying GPU Programming 中看到，說明了 __managed__ 全局聲明的簡單使用：

// Managed variable declaration is an extra annotation with __device__
__device__ __managed__  int  x;
__global__  void  kernel() {
    // Reference "x" directly - it's a normal variable on the GPU.
    printf( "GPU sees: x = %d\n" , x);
} 
int  main() {
    // Set "x" from Host code. Note it's just a normal variable on the CPU.
    x = 1234;
 
    // Launch a kernel which uses "x" from the GPU.
    kernel<<< 1, 1 >>>(); 
    cudaDeviceSynchronize(); 
    return  0;
}

__managed__ 變量的可用功能是該符號在設(shè)備代碼和主機代碼中都可用，而無需取消引用指針，并且數(shù)據(jù)由所有人共享。這使得在主機和設(shè)備程序之間交換數(shù)據(jù)變得特別容易，而無需顯式分配或復(fù)制。

從語義上講，__managed__ 變量的行為與通過 cudaMallocManaged（） 分配的存儲相同。有關(guān)詳細說明，請參閱使用 cudaMallocManaged（） 進行顯式分配。流可見性默認為 cudaMemAttachGlobal，但可以使用 cudaStreamAttachMemAsync（） 進行限制。

__managed__ 變量的正確操作需要有效的 CUDA 上下文。如果當(dāng)前設(shè)備的上下文尚未創(chuàng)建，則訪問 __managed__ 變量可以觸發(fā) CUDA 上下文創(chuàng)建。在上面的示例中，在內(nèi)核啟動之前訪問 x 會觸發(fā)設(shè)備 0 上的上下文創(chuàng)建。如果沒有該訪問，內(nèi)核啟動將觸發(fā)上下文創(chuàng)建。

聲明為 __managed__ 的 C++ 對象受到某些特定約束，尤其是在涉及靜態(tài)初始化程序的情況下。有關(guān)這些約束的列表，請參閱 CUDA C++ 編程指南中的 C++ 語言支持。

N.2.3.1. Host Program Errors with managed Variables

__managed__ 變量的使用取決于底層統(tǒng)一內(nèi)存系統(tǒng)是否正常運行。 例如，如果 CUDA 安裝失敗或 CUDA 上下文創(chuàng)建不成功，則可能會出現(xiàn)不正確的功能。

當(dāng)特定于 CUDA 的操作失敗時，通常會返回一個錯誤，指出失敗的根源。 使用 __managed__ 變量引入了一種新的故障模式，如果統(tǒng)一內(nèi)存系統(tǒng)運行不正確，非 CUDA 操作（例如，CPU 訪問應(yīng)該是有效的主機內(nèi)存地址）可能會失敗。 這種無效的內(nèi)存訪問不能輕易地歸因于底層的 CUDA 子系統(tǒng)，盡管諸如 cuda-gdb 之類的調(diào)試器會指示托管內(nèi)存地址是故障的根源。

N.2.4. Querying Unified Memory Support

N.2.4.1. Device Properties

統(tǒng)一內(nèi)存僅在具有 3.0 或更高計算能力的設(shè)備上受支持。程序可以通過使用 cudaGetDeviceProperties（） 并檢查新的 managedMemory 屬性來查詢 GPU 設(shè)備是否支持托管內(nèi)存。也可以使用具有屬性 cudaDevAttrManagedMemory 的單個屬性查詢函數(shù) cudaDeviceGetAttribute（） 來確定能力。

如果在 GPU 和當(dāng)前操作系統(tǒng)下允許托管內(nèi)存分配，則任一屬性都將設(shè)置為 1。請注意，32 位應(yīng)用程序不支持統(tǒng)一內(nèi)存（除非在 Android 上），即使 GPU 有足夠的能力。

支持平臺上計算能力 6.x 的設(shè)備無需調(diào)用 cudaHostRegister 即可訪問可分頁內(nèi)存。應(yīng)用程序可以通過檢查新的 pageableMemoryAccess 屬性來查詢設(shè)備是否支持連貫訪問可分頁內(nèi)存。

通過新的缺頁機制，統(tǒng)一內(nèi)存保證了全局?jǐn)?shù)據(jù)的一致性。這意味著 CPU 和 GPU 可以同時訪問統(tǒng)一內(nèi)存分配。這在計算能力低于 6.x 的設(shè)備上是非法的，因為如果 CPU 在 GPU 內(nèi)核處于活動狀態(tài)時訪問統(tǒng)一內(nèi)存分配，則無法保證一致性。程序可以通過檢查 concurrentManagedAccess 屬性來查詢并發(fā)訪問支持。有關(guān)詳細信息，請參閱一致性和并發(fā)性。

N.2.5. Advanced Topics

N.2.5.1. Managed Memory with Multi-GPU Programs on pre-6.x Architectures

在計算能力低于 6.x 的設(shè)備的系統(tǒng)上，托管分配通過 GPU 的對等能力自動對系統(tǒng)中的所有 GPU 可見。

在 Linux 上，只要程序正在使用的所有 GPU 都具有點對點支持，托管內(nèi)存就會在 GPU 內(nèi)存中分配。如果在任何時候應(yīng)用程序開始使用不支持對等支持的 GPU 與任何其他對其進行了托管分配的 GPU，則驅(qū)動程序會將所有托管分配遷移到系統(tǒng)內(nèi)存。

在 Windows 上，如果對等映射不可用（例如，在不同架構(gòu)的 GPU 之間），那么系統(tǒng)將自動回退到使用零拷貝內(nèi)存，無論兩個 GPU 是否都被程序?qū)嶋H使用。如果實際只使用一個 GPU，則需要在啟動程序之前設(shè)置 CUDA_VISIBLE_DEVICES 環(huán)境變量。這限制了哪些 GPU 是可見的，并允許在 GPU 內(nèi)存中分配托管內(nèi)存。

或者，在 Windows 上，用戶還可以將 CUDA_MANAGED_FORCE_DEVICE_ALLOC 設(shè)置為非零值，以強制驅(qū)動程序始終使用設(shè)備內(nèi)存進行物理存儲。當(dāng)此環(huán)境變量設(shè)置為非零值時，該進程中使用的所有支持托管內(nèi)存的設(shè)備必須彼此對等兼容。如果使用支持托管內(nèi)存的設(shè)備并且它與之前在該進程中使用的任何其他托管內(nèi)存支持設(shè)備不兼容，則將返回錯誤 ：：cudaErrorInvalidDevice，即使 ：：cudaDeviceReset 具有在這些設(shè)備上被調(diào)用。這些環(huán)境變量在附錄 CUDA 環(huán)境變量中進行了描述。請注意，從 CUDA 8.0 開始，CUDA_MANAGED_FORCE_DEVICE_ALLOC 對 Linux 操作系統(tǒng)沒有影響。

N.2.5.2. Using fork（） with Managed Memory

統(tǒng)一內(nèi)存系統(tǒng)不允許在進程之間共享托管內(nèi)存指針。 它不會正確管理通過 fork（） 操作復(fù)制的內(nèi)存句柄。 如果子級或父級在 fork（） 之后訪問托管數(shù)據(jù)，則結(jié)果將不確定。

然而，fork（） 一個子進程然后通過 exec（） 調(diào)用立即退出是安全的，因為子進程丟棄了內(nèi)存句柄并且父進程再次成為唯一的所有者。 父母離開并讓孩子接觸句柄是不安全的。

N.3. Performance Tuning

為了使用統(tǒng)一內(nèi)存實現(xiàn)良好的性能，必須滿足以下目標(biāo)：

應(yīng)避免錯誤：雖然可重放錯誤是啟用更簡單的編程模型的基礎(chǔ)，但它們可能嚴(yán)重損害應(yīng)用程序性能。故障處理可能需要幾十微秒，因為它可能涉及 TLB 無效、數(shù)據(jù)遷移和頁表更新。與此同時，應(yīng)用程序某些部分的執(zhí)行將停止，從而可能影響整體性能。

數(shù)據(jù)應(yīng)該位于訪問處理器的本地：如前所述，當(dāng)數(shù)據(jù)位于訪問它的處理器本地時，內(nèi)存訪問延遲和帶寬明顯更好。因此，應(yīng)適當(dāng)遷移數(shù)據(jù)以利用較低的延遲和較高的帶寬。

應(yīng)該防止內(nèi)存抖動：如果數(shù)據(jù)被多個處理器頻繁訪問并且必須不斷遷移以實現(xiàn)數(shù)據(jù)局部性，那么遷移的開銷可能會超過局部性的好處。應(yīng)盡可能防止內(nèi)存抖動。如果無法預(yù)防，則必須進行適當(dāng)?shù)?a target="_blank">檢測和解決。

為了達到與不使用統(tǒng)一內(nèi)存相同的性能水平，應(yīng)用程序必須引導(dǎo)統(tǒng)一內(nèi)存驅(qū)動子系統(tǒng)避免上述陷阱。值得注意的是，統(tǒng)一內(nèi)存驅(qū)動子系統(tǒng)可以檢測常見的數(shù)據(jù)訪問模式并自動實現(xiàn)其中一些目標(biāo)，而無需應(yīng)用程序參與。但是，當(dāng)數(shù)據(jù)訪問模式不明顯時，來自應(yīng)用程序的明確指導(dǎo)至關(guān)重要。 CUDA 8.0 引入了有用的 API，用于為運行時提供內(nèi)存使用提示 （cudaMemAdvise（）） 和顯式預(yù)取 （cudaMemPrefetchAsync（））。這些工具允許與顯式內(nèi)存復(fù)制和固定 API 相同的功能，而不會恢復(fù)到顯式 GPU 內(nèi)存分配的限制。

注意：Tegra 設(shè)備不支持 cudaMemPrefetchAsync（）。

N.3.1. Data Prefetching

數(shù)據(jù)預(yù)取意味著將數(shù)據(jù)遷移到處理器的內(nèi)存中，并在處理器開始訪問該數(shù)據(jù)之前將其映射到該處理器的頁表中。 數(shù)據(jù)預(yù)取的目的是在建立數(shù)據(jù)局部性的同時避免故障。 這對于在任何給定時間主要從單個處理器訪問數(shù)據(jù)的應(yīng)用程序來說是最有價值的。 由于訪問處理器在應(yīng)用程序的生命周期中發(fā)生變化，因此可以相應(yīng)地預(yù)取數(shù)據(jù)以遵循應(yīng)用程序的執(zhí)行流程。 由于工作是在 CUDA 中的流中啟動的，因此預(yù)計數(shù)據(jù)預(yù)取也是一種流操作，如以下 API 所示：

    cudaError_t cudaMemPrefetchAsync(const void *devPtr, 
                                     size_t count, 
                                     int dstDevice, 
                                     cudaStream_t stream);

其中由devPtr指針和count字節(jié)數(shù)指定的內(nèi)存區(qū)域，ptr向下舍入到最近的頁面邊界，count向上舍入到最近的頁面邊界，通過在流中排隊遷移操作遷移到dstDevice。 為dstDevice傳入cudaCpuDeviceId會導(dǎo)致數(shù)據(jù)遷移到 CPU 內(nèi)存。 考慮下面的一個簡單代碼示例：

void foo(cudaStream_t s) {
  char *data;
  cudaMallocManaged(&data, N);
  init_data(data, N);                                   // execute on CPU
  cudaMemPrefetchAsync(data, N, myGpuId, s);            // prefetch to GPU
  mykernel<<<..., s>>>(data, N, 1, compare);            // execute on GPU
  cudaMemPrefetchAsync(data, N, cudaCpuDeviceId, s);    // prefetch to CPU
  cudaStreamSynchronize(s);
  use_data(data, N);
  cudaFree(data);
}

如果沒有性能提示，內(nèi)核 mykernel 將在首次訪問數(shù)據(jù)時出錯，這會產(chǎn)生額外的故障處理開銷，并且通常會減慢應(yīng)用程序的速度。 通過提前預(yù)取數(shù)據(jù)，可以避免頁面錯誤并獲得更好的性能。 此 API 遵循流排序語義，即遷移在流中的所有先前操作完成之前不會開始，并且流中的任何后續(xù)操作在遷移完成之前不會開始。

N.3.2. Data Usage Hints

當(dāng)多個處理器需要同時訪問相同的數(shù)據(jù)時，單獨的數(shù)據(jù)預(yù)取是不夠的。 在這種情況下，應(yīng)用程序提供有關(guān)如何實際使用數(shù)據(jù)的提示很有用。 以下咨詢 API 可用于指定數(shù)據(jù)使用情況：

    cudaError_t cudaMemAdvise(const void *devPtr, 
                              size_t count, 
                              enum cudaMemoryAdvise advice, 
                              int device);

其中，為從 devPtr 地址開始的區(qū)域中包含的數(shù)據(jù)指定的通知和計數(shù)字節(jié)的長度，四舍五入到最近的頁面邊界，可以采用以下值：

cudaMemAdviseSetReadMostly：這意味著數(shù)據(jù)大部分將被讀取并且只是偶爾寫入。 這允許驅(qū)動程序在處理器訪問數(shù)據(jù)時在處理器內(nèi)存中創(chuàng)建數(shù)據(jù)的只讀拷貝。 同樣，如果在此區(qū)域上調(diào)用 cudaMemPrefetchAsync，它將在目標(biāo)處理器上創(chuàng)建數(shù)據(jù)的只讀拷貝。 當(dāng)處理器寫入此數(shù)據(jù)時，相應(yīng)頁面的所有副本都將失效，但發(fā)生寫入的拷貝除外。 此建議忽略設(shè)備參數(shù)。 該建議允許多個處理器以最大帶寬同時訪問相同的數(shù)據(jù)，如以下代碼片段所示：

char *dataPtr;
size_t dataSize = 4096;
// Allocate memory using malloc or cudaMallocManaged
dataPtr = (char *)malloc(dataSize);
// Set the advice on the memory region
cudaMemAdvise(dataPtr, dataSize, cudaMemAdviseSetReadMostly, 0);
int outerLoopIter = 0;
while (outerLoopIter < maxOuterLoopIter) {
    // The data is written to in the outer loop on the CPU
    initializeData(dataPtr, dataSize);
    // The data is made available to all GPUs by prefetching.
    // Prefetching here causes read duplication of data instead
    // of data migration
    for (int device = 0; device < maxDevices; device++) {
        cudaMemPrefetchAsync(dataPtr, dataSize, device, stream);
    }
    // The kernel only reads this data in the inner loop
    int innerLoopIter = 0;
    while (innerLoopIter < maxInnerLoopIter) {
        kernel<<<32,32>>>((const char *)dataPtr);
        innerLoopIter++;
    }
    outerLoopIter++;
}

cudaMemAdviseSetPreferredLocation：此建議將數(shù)據(jù)的首選位置設(shè)置為屬于設(shè)備的內(nèi)存。傳入設(shè)備的 cudaCpuDeviceId 值會將首選位置設(shè)置為 CPU 內(nèi)存。設(shè)置首選位置不會導(dǎo)致數(shù)據(jù)立即遷移到該位置。相反，它會在該內(nèi)存區(qū)域發(fā)生故障時指導(dǎo)遷移策略。如果數(shù)據(jù)已經(jīng)在它的首選位置并且故障處理器可以建立映射而不需要遷移數(shù)據(jù)，那么遷移將被避免。另一方面，如果數(shù)據(jù)不在其首選位置，或者無法建立直接映射，那么它將被遷移到訪問它的處理器。請務(wù)必注意，設(shè)置首選位置不會阻止使用 cudaMemPrefetchAsync 完成數(shù)據(jù)預(yù)取。

cudaMemAdviseSetAccessedBy：這個advice意味著數(shù)據(jù)將被設(shè)備訪問。這不會導(dǎo)致數(shù)據(jù)遷移，并且對數(shù)據(jù)本身的位置沒有影響。相反，只要數(shù)據(jù)的位置允許建立映射，它就會使數(shù)據(jù)始終映射到指定處理器的頁表中。如果數(shù)據(jù)因任何原因被遷移，映射會相應(yīng)更新。此advice在數(shù)據(jù)局部性不重要但避免故障很重要的情況下很有用。例如，考慮一個包含多個啟用對等訪問的 GPU 的系統(tǒng)，其中位于一個 GPU 上的數(shù)據(jù)偶爾會被其他 GPU 訪問。在這種情況下，將數(shù)據(jù)遷移到其他 GPU 并不那么重要，因為訪問不頻繁并且遷移的開銷可能太高。但是防止故障仍然有助于提高性能，因此提前設(shè)置映射很有用。請注意，在 CPU 訪問此數(shù)據(jù)時，由于 CPU 無法直接訪問 GPU 內(nèi)存，因此數(shù)據(jù)可能會遷移到 CPU 內(nèi)存。任何為此數(shù)據(jù)設(shè)置了 cudaMemAdviceSetAccessedBy 標(biāo)志的 GPU 現(xiàn)在都將更新其映射以指向 CPU 內(nèi)存中的頁面。

每個advice也可以使用以下值之一取消設(shè)置：cudaMemAdviseUnsetReadMostly、cudaMemAdviseUnsetPreferredLocation 和 cudaMemAdviseUnsetAccessedBy。

N.3.3. Querying Usage Attributes

程序可以使用以下 API 查詢通過 cudaMemAdvise 或 cudaMemPrefetchAsync 分配的內(nèi)存范圍屬性：

    cudaMemRangeGetAttribute(void *data, 
                             size_t dataSize, 
                             enum cudaMemRangeAttribute attribute, 
                             const void *devPtr, 
                             size_t count);

此函數(shù)查詢從 devPtr 開始的內(nèi)存范圍的屬性，大小為 count 字節(jié)。內(nèi)存范圍必須引用通過 cudaMallocManaged 分配或通過 __managed__ 變量聲明的托管內(nèi)存?？梢圆樵円韵聦傩裕?/p>

cudaMemRangeAttributeReadMostly：如果給定內(nèi)存范圍內(nèi)的所有頁面都啟用了重復(fù)讀取，則返回的結(jié)果將為 1，否則返回 0。

cudaMemRangeAttributePreferredLocation：如果內(nèi)存范圍內(nèi)的所有頁面都將相應(yīng)的處理器作為首選位置，則返回結(jié)果將是 GPU 設(shè)備 ID 或 cudaCpuDeviceId，否則將返回 cudaInvalidDeviceId。應(yīng)用程序可以使用此查詢 API 來決定通過 CPU 或 GPU 暫存數(shù)據(jù)，具體取決于托管指針的首選位置屬性。請注意，查詢時內(nèi)存范圍內(nèi)頁面的實際位置可能與首選位置不同。

cudaMemRangeAttributeAccessedBy： 將返回為該內(nèi)存范圍設(shè)置了該建議的設(shè)備列表。

cudaMemRangeAttributeLastPrefetchLocation：將返回使用 cudaMemPrefetchAsync 顯式預(yù)取內(nèi)存范圍內(nèi)所有頁面的最后位置。請注意，這只是返回應(yīng)用程序請求將內(nèi)存范圍預(yù)取到的最后一個位置。它沒有指示對該位置的預(yù)取操作是否已經(jīng)完成或什至開始。

此外，還可以使用對應(yīng)的 cudaMemRangeGetAttributes 函數(shù)查詢多個屬性。

關(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ā)者。

審核編輯：郭婷