N.1. Unified Memory Introduction

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

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

底層系統管理 CUDA 程序中的數據訪問和位置，無需顯式內存復制調用。這在兩個主要方面有利于 GPU 編程：

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

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

簡單來說，統一內存消除了通過 cudaMemcpy*（） 例程進行顯式數據移動的需要，而不會因將所有數據放入零拷貝內存而導致性能損失。當然，數據移動仍然會發生，因此程序的運行時間通常不會減少；相反，統一內存可以編寫更簡單、更易于維護的代碼。

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

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

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

注意：連接到 Tegra 的離散 GPU 不支持統一內存。

N.1.1. System Requirements

統一內存有兩個基本要求：

具有 SM 架構 3.0 或更高版本（Kepler 類或更高版本）的 GPU

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

N.1.2. Simplifying GPU Programming

內存空間的統一意味著主機和設備之間不再需要顯式內存傳輸。在托管內存空間中創建的任何分配都會自動遷移到需要的位置。

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

__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 上將兩個數字與每個線程 ID 組合在一起，并以數組形式返回值。 如果沒有托管內存，則返回值的主機端和設備端存儲都是必需的（示例中為 host_ret 和 ret），使用 cudaMemcpy（） 在兩者之間顯式復制也是如此。

將此與程序的統一內存版本進行比較，后者允許從主機直接訪問 GPU 數據。 請注意 cudaMallocManaged（） 例程，它從主機和設備代碼返回一個有效的指針。 這允許在沒有單獨的 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（） 命令以及返回數組 ret 在 CPU 和 GPU 上都可見的事實。

值得一提的是主機和設備之間的同步。 請注意在非托管示例中，同步 cudaMemcpy（） 例程如何用于同步內核（即等待它完成運行）以及將數據傳輸到主機。 統一內存示例不調用 cudaMemcpy（），因此需要顯式 cudaDeviceSynchronize（），然后主機程序才能安全地使用 GPU 的輸出。

N.1.3. Data Migration and Coherency

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

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

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

N.1.4. GPU Memory Oversubscription

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

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

N.1.5. Multi-GPU

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

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

N.1.6. System Allocator

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

這是一個適用于任何滿足統一內存基本要求的系統的示例代碼（請參閱系統要求）：

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

具有 pageableMemoryAccess 屬性的系統支持這些新的訪問模式：

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

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

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

在上面的示例中，數據可以由第三方 CPU 庫初始化，然后由 GPU 內核直接訪問。 在具有 pageableMemoryAccess 的系統上，用戶還可以使用 cudaMemPrefetchAsync 將可分頁內存預取到 GPU。 這可以通過優化數據局部性產生性能優勢。

注意：目前僅 IBM Power9 系統支持基于 NVLink 的 ATS。

N.1.7. Hardware Coherency

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

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

考慮下面的示例代碼：

__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;
}

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

在 directManagedMemAccessFromHost=1 的系統上：CPU 訪問托管緩沖區不會觸發任何遷移； 數據將保留在 GPU 內存中，任何后續的 GPU 內核都可以繼續直接訪問它，而不會造成故障或遷移。

在 directManagedMemAccessFromHost=0 的系統上：CPU 訪問托管緩沖區將出現頁面錯誤并啟動數據遷移； 任何第一次嘗試訪問相同數據的 GPU 內核都會出現頁面錯誤并將頁面遷移回 GPU 內存。

N.1.8. Access Counters

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

對于 cudaMallocManaged，訪問計數器遷移可以通過使用帶有相應設備 ID 的 cudaMemAdviseSetAccessedBy 提示來選擇加入。 驅動程序還可以使用訪問計數器來實現更有效的抖動緩解或內存超額訂閱方案。

注意：訪問計數器當前僅在 IBM Power9 系統上啟用，并且僅用于 cudaMallocManaged 分配器。

N.2. Programming Model

N.2.1. Managed Memory Opt In

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

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

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

N.2.1.1. Explicit Allocation Using cudaMallocManaged（）

統一內存最常使用在語義和語法上類似于標準 CUDA 分配器 cudaMalloc（） 的分配函數創建。 功能說明如下：

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

cudaMallocManaged（） 函數保留托管內存的 size 字節，并在 devPtr 中返回一個指針。 請注意各種 GPU 架構之間 cudaMallocManaged（） 行為的差異。 默認情況下，計算能力低于 6.x 的設備直接在 GPU 上分配托管內存。 但是，計算能力 6.x 及更高版本的設備在調用 cudaMallocManaged（） 時不會分配物理內存：在這種情況下，物理內存會在第一次觸摸時填充，并且可能駐留在 CPU 或 GPU 上。 托管指針在系統中的所有 GPU 和 CPU 上都有效，盡管程序訪問此指針必須遵守統一內存編程模型的并發規則（請參閱一致性和并發性）。 下面是一個簡單的例子，展示了 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;
}

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

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

N.2.1.2. Global-Scope Managed Variables Using managed

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

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

__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__ 內存空間的所有語義，以及一些額外的統一內存特定約束，都由托管變量繼承（請參閱使用 NVCC 編譯）。

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

N.2.2. Coherency and Concurrency

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

N.2.2.1. GPU Exclusive Access To Managed Memory

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

不允許 CPU 訪問任何托管分配或變量，而 GPU 對 concurrentManagedAccess 屬性設置為 0 的設備處于活動狀態。在這些系統上，并發 CPU/GPU 訪問，即使是不同的托管內存分配，也會導致分段錯誤，因為該頁面被認為是 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;
}

在上面的示例中，當 CPU 接觸(這里原文中用的是touch這個詞) y 時，GPU 程序內核仍然處于活動狀態。 （注意它是如何在cudaDeviceSynchronize()之前發生的。）由于 GPU 頁面錯誤功能解除了對同時訪問的所有限制，因此代碼在計算能力 6.x 的設備上成功運行。 但是，即使 CPU 訪問的數據與 GPU 不同，這種內存訪問在 6.x 之前的架構上也是無效的。 程序必須在訪問 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 架構的系統上，CPU 線程可能不會在執行內核啟動和后續同步調用之間訪問任何托管數據，無論 GPU 內核是否實際接觸相同的數據（或 任何托管數據）。 并發 CPU 和 GPU 訪問的潛力足以引發進程級異常。

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

N.2.2.2. Explicit Synchronization and Logical GPU Activity

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

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

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

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

有幾個重要的注意點：

在 GPU 處于活動狀態時，始終允許 CPU 訪問非托管零拷貝數據。

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

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

并發 GPU 內核訪問托管數據沒有任何限制。

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

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

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

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

統一內存建立在流獨立模型之上，允許 CUDA 程序顯式地將托管分配與 CUDA 流相關聯。通過這種方式，程序員根據內核是否將數據啟動到指定的流中來指示內核對數據的使用。這為基于程序特定數據訪問模式的并發提供了機會。控制這種行為的函數是：

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

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

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

通過將分配與特定流相關聯，程序保證只有啟動到該流中的內核才會接觸該數據。統一內存系統不執行錯誤檢查：程序員有責任確保兌現保證。

除了允許更大的并發性之外，使用 cudaStreamAttachMemAsync（） 可以（并且通常會）啟用統一內存系統內的數據傳輸優化，這可能會影響延遲和其他開銷。

N.2.2.4. Stream Association Examples

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

__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 與主機可訪問性相關聯，從而始終可以從 CPU 進行訪問。 （和以前一樣，請注意在訪問之前沒有cudaDeviceSynchronize()。）GPU 運行內核對 y 的訪問現在將產生未定義的結果。

請注意，將變量與流關聯不會更改任何其他變量的關聯。 例如。 將 x 與 stream1 關聯并不能確保在 stream1 中啟動的內核只能訪問 x，因此此代碼會導致錯誤：

__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 將如何導致錯誤，因為即使 x 已與流相關聯，我們也沒有告訴系統誰可以看到 y。 因此，系統保守地假設內核可能會訪問它并阻止 CPU 這樣做。

N.2.2.5. Stream Attach With Multithreaded Host Programs

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

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

這樣的程序將簡單地添加一個對 cudaStreamAttachMemAsync（） 的調用，以使用統一內存進行數據訪問：

// 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);
}

在這個例子中，分配流關聯只建立一次，然后主機和設備都重復使用數據。 結果是比在主機和設備之間顯式復制數據時更簡單的代碼，盡管結果是相同的。

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

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

如果沒有這個標志，如果另一個線程啟動的內核恰好正在運行，則新分配將被視為在 GPU 上使用。這可能會影響線程在能夠將其顯式附加到私有流之前從 CPU 訪問新分配的數據的能力（例如，在基類構造函數中）。因此，為了啟用線程之間的安全獨立性，應指定此標志進行分配。

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

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

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

如果指定了 cudaMemcpyHostTo* 并且管理了源數據，那么如果在復制流 （1） 中可以從主機連貫地訪問它，那么它將從主機訪問；否則將從設備訪問。當指定 cudaMemcpy*ToHost 并且目標是托管內存時，類似的規則適用于目標。

如果指定了 cudaMemcpyDeviceTo* 并管理源數據，則將從設備訪問它。源必須可以從復制流中的設備連貫地訪問 （2）；否則，返回錯誤。當指定 cudaMemcpy*ToDevice 并且目標是托管內存時，類似的規則適用于目標。

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

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

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

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

給定流與設備屬性 concurrentManagedAccess 具有非零值的設備相關聯。

內存既不具有全局可見性，也不與給定流相關聯。

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

設備的設備屬性 concurrentManagedAccess 具有非零值。

內存要么具有全局可見性，要么與給定的流相關聯。

###N.2.3. Language Integration

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

下面的例子在前面的 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__ 變量的可用功能是該符號在設備代碼和主機代碼中都可用，而無需取消引用指針，并且數據由所有人共享。這使得在主機和設備程序之間交換數據變得特別容易，而無需顯式分配或復制。

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

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

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

N.2.3.1. Host Program Errors with managed Variables

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

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

N.2.4. Querying Unified Memory Support

N.2.4.1. Device Properties

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

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

支持平臺上計算能力 6.x 的設備無需調用 cudaHostRegister 即可訪問可分頁內存。應用程序可以通過檢查新的 pageableMemoryAccess 屬性來查詢設備是否支持連貫訪問可分頁內存。

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

N.2.5. Advanced Topics

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

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

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

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

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

N.2.5.2. Using fork（） with Managed Memory

統一內存系統不允許在進程之間共享托管內存指針。 它不會正確管理通過 fork（） 操作復制的內存句柄。 如果子級或父級在 fork（） 之后訪問托管數據，則結果將不確定。

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

N.3. Performance Tuning

為了使用統一內存實現良好的性能，必須滿足以下目標：

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

數據應該位于訪問處理器的本地：如前所述，當數據位于訪問它的處理器本地時，內存訪問延遲和帶寬明顯更好。因此，應適當遷移數據以利用較低的延遲和較高的帶寬。

應該防止內存抖動：如果數據被多個處理器頻繁訪問并且必須不斷遷移以實現數據局部性，那么遷移的開銷可能會超過局部性的好處。應盡可能防止內存抖動。如果無法預防，則必須進行適當的檢測和解決。

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

注意：Tegra 設備不支持 cudaMemPrefetchAsync（）。

N.3.1. Data Prefetching

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

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

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

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);
}

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

N.3.2. Data Usage Hints

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

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

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

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

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：此建議將數據的首選位置設置為屬于設備的內存。傳入設備的 cudaCpuDeviceId 值會將首選位置設置為 CPU 內存。設置首選位置不會導致數據立即遷移到該位置。相反，它會在該內存區域發生故障時指導遷移策略。如果數據已經在它的首選位置并且故障處理器可以建立映射而不需要遷移數據，那么遷移將被避免。另一方面，如果數據不在其首選位置，或者無法建立直接映射，那么它將被遷移到訪問它的處理器。請務必注意，設置首選位置不會阻止使用 cudaMemPrefetchAsync 完成數據預取。

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

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

N.3.3. Querying Usage Attributes

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

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

此函數查詢從 devPtr 開始的內存范圍的屬性，大小為 count 字節。內存范圍必須引用通過 cudaMallocManaged 分配或通過 __managed__ 變量聲明的托管內存。可以查詢以下屬性：

cudaMemRangeAttributeReadMostly：如果給定內存范圍內的所有頁面都啟用了重復讀取，則返回的結果將為 1，否則返回 0。

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

cudaMemRangeAttributeAccessedBy： 將返回為該內存范圍設置了該建議的設備列表。

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

此外，還可以使用對應的 cudaMemRangeGetAttributes 函數查詢多個屬性。

關于作者

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

審核編輯：郭婷