CCUDA編程模型和接口介紹

本項(xiàng)目為CUDA官方手冊的中文翻譯版，有個(gè)人翻譯并添加自己的理解。主要介紹CUDA編程模型和接口。

第三章編程接口

CUDA C++ 為熟悉 C++ 編程語言的用戶提供了一種簡單的途徑，可以輕松編寫由設(shè)備執(zhí)行的程序。

它由c++語言的最小擴(kuò)展集和運(yùn)行時(shí)庫組成。

編程模型中引入了核心語言擴(kuò)展。它們允許程序員將內(nèi)核定義為 C++ 函數(shù)，并在每次調(diào)用函數(shù)時(shí)使用一些新語法來指定網(wǎng)格和塊的維度。所有擴(kuò)展的完整描述可以在 C++ 語言擴(kuò)展中找到。任何包含這些擴(kuò)展名的源文件都必須使用 nvcc 進(jìn)行編譯，如使用NVCC編譯中所述。

運(yùn)行時(shí)在 CUDA Runtime 中引入。它提供了在主機(jī)上執(zhí)行的 C 和 C++ 函數(shù)，用于分配和釋放設(shè)備內(nèi)存、在主機(jī)內(nèi)存和設(shè)備內(nèi)存之間傳輸數(shù)據(jù)、管理具有多個(gè)設(shè)備的系統(tǒng)等。運(yùn)行時(shí)的完整描述可以在 CUDA 參考手冊中找到。

運(yùn)行時(shí)構(gòu)建在較低級(jí)別的 C API（即 CUDA 驅(qū)動(dòng)程序 API）之上，應(yīng)用程序也可以訪問該 API。驅(qū)動(dòng)程序 API 通過公開諸如 CUDA 上下文（類似于設(shè)備的主機(jī)進(jìn)程）和 CUDA 模塊（類似于設(shè)備的動(dòng)態(tài)加載庫）等較低級(jí)別的概念來提供額外的控制級(jí)別。大多數(shù)應(yīng)用程序不使用驅(qū)動(dòng)程序 API，因?yàn)樗鼈儾恍枰@種額外的控制級(jí)別，并且在使用運(yùn)行時(shí)時(shí)，上下文和模塊管理是隱式的，從而產(chǎn)生更簡潔的代碼。由于運(yùn)行時(shí)可與驅(qū)動(dòng)程序 API 互操作，因此大多數(shù)需要驅(qū)動(dòng)程序 API 功能的應(yīng)用程序可以默認(rèn)使用運(yùn)行時(shí) API，并且僅在需要時(shí)使用驅(qū)動(dòng)程序 API。 Driver API 中介紹了驅(qū)動(dòng)API并在參考手冊中進(jìn)行了全面描述。

3.1利用NVCC編譯

內(nèi)核可以使用稱為 PTX 的 CUDA 指令集架構(gòu)來編寫，PTX 參考手冊中對此進(jìn)行了描述。 然而，使用高級(jí)編程語言（如 C++）通常更有效。 在這兩種情況下，內(nèi)核都必須通過 nvcc 編譯成二進(jìn)制代碼才能在設(shè)備上執(zhí)行。

nvcc 是一種編譯器驅(qū)動(dòng)程序，可簡化編譯 C++ 或 PTX 代碼：它提供簡單且熟悉的命令行選項(xiàng)，并通過調(diào)用實(shí)現(xiàn)不同編譯階段的工具集合來執(zhí)行它們。 本節(jié)概述了 nvcc 工作流程和命令選項(xiàng)。 完整的描述可以在 nvcc 用戶手冊中找到。

3.1.1編譯流程

3.1.1.1 離線編譯

使用 nvcc 編譯的源文件可以包含主機(jī)代碼（即在host上執(zhí)行的代碼）和設(shè)備代碼（即在device上執(zhí)行的代碼。 nvcc 的基本工作流程包括將設(shè)備代碼與主機(jī)代碼分離，然后：

將設(shè)備代碼編譯成匯編形式（PTX 代碼）或二進(jìn)制形式（cubin 對象）

并通過CUDA運(yùn)行時(shí)函數(shù)的調(diào)用來替換 《《《…》》》 語法對主機(jī)代碼進(jìn)行修改，以從 PTX 代碼或 cubin 對象加載和啟動(dòng)每個(gè)編譯的內(nèi)核。

修改后的主機(jī)代碼要么作為 C++ 代碼輸出，然后使用另一個(gè)工具編譯，要么直接作為目標(biāo)代碼輸出，方法是讓 nvcc 在最后編譯階段調(diào)用主機(jī)編譯器。

然后應(yīng)用程序可以：

鏈接到已編譯的主機(jī)代碼（這是最常見的情況），

或者忽略修改后的主機(jī)代碼（如果有）并使用 CUDA 驅(qū)動(dòng)程序 API（請參閱驅(qū)動(dòng)程序 API）來加載和執(zhí)行 PTX 代碼或 cubin 對象。

3.1.1.2 即時(shí)編譯

應(yīng)用程序在運(yùn)行時(shí)加載的任何 PTX 代碼都由設(shè)備驅(qū)動(dòng)程序進(jìn)一步編譯為二進(jìn)制代碼。這稱為即時(shí)編譯。即時(shí)編譯增加了應(yīng)用程序加載時(shí)間，但允許應(yīng)用程序受益于每個(gè)新設(shè)備驅(qū)動(dòng)程序帶來的任何新編譯器改進(jìn)。它也是應(yīng)用程序能夠運(yùn)行在編譯時(shí)不存在的設(shè)備上的唯一方式，如應(yīng)用程序兼容性中所述。

當(dāng)設(shè)備驅(qū)動(dòng)程序?yàn)槟承?yīng)用程序?qū)崟r(shí)編譯一些 PTX 代碼時(shí)，它會(huì)自動(dòng)緩存生成二進(jìn)制代碼的副本，以避免在應(yīng)用程序的后續(xù)調(diào)用中重復(fù)編譯。緩存（稱為計(jì)算緩存）在設(shè)備驅(qū)動(dòng)程序升級(jí)時(shí)自動(dòng)失效，因此應(yīng)用程序可以從設(shè)備驅(qū)動(dòng)程序中內(nèi)置的新即時(shí)編譯器的改進(jìn)中受益。

環(huán)境變量可用于控制即時(shí)編譯，如 CUDA 環(huán)境變量中所述

作為使用 nvcc 編譯 CUDA C++ 設(shè)備代碼的替代方法，NVRTC 可用于在運(yùn)行時(shí)將 CUDA C++ 設(shè)備代碼編譯為 PTX。 NVRTC 是 CUDA C++ 的運(yùn)行時(shí)編譯庫；更多信息可以在 NVRTC 用戶指南中找到。

3.1.2 Binary 兼容性

二進(jìn)制代碼是特定于體系結(jié)構(gòu)的。 使用指定目標(biāo)體系結(jié)構(gòu)的編譯器選項(xiàng) -code 生成 cubin 對象：例如，使用 -code=sm_35 編譯會(huì)為計(jì)算能力為 3.5 的設(shè)備生成二進(jìn)制代碼。 從一個(gè)次要修訂版到下一個(gè)修訂版都保證了二進(jìn)制兼容性，但不能保證從一個(gè)次要修訂版到前一個(gè)修訂版或跨主要修訂版。 換句話說，為計(jì)算能力 X.y 生成的 cubin 對象只會(huì)在計(jì)算能力 X.z 且 z≥y 的設(shè)備上執(zhí)行。

注意：僅桌面支持二進(jìn)制兼容性。 Tegra 不支持它。 此外，不支持桌面和 Tegra 之間的二進(jìn)制兼容性。

3.1.3 PTX 兼容性

某些 PTX 指令僅在具有更高計(jì)算能力的設(shè)備上受支持。 例如，Warp Shuffle Functions 僅在計(jì)算能力 3.0 及以上的設(shè)備上支持。 -arch 編譯器選項(xiàng)指定將 C++ 編譯為 PTX 代碼時(shí)假定的計(jì)算能力。 因此，例如，包含 warp shuffle 的代碼必須使用 -arch=compute_30（或更高版本）進(jìn)行編譯。

為某些特定計(jì)算能力生成的 PTX 代碼始終可以編譯為具有更大或相等計(jì)算能力的二進(jìn)制代碼。 請注意，從早期 PTX 版本編譯的二進(jìn)制文件可能無法使用某些硬件功能。 例如，從為計(jì)算能力 6.0 （Pascal） 生成的 PTX 編譯的計(jì)算能力 7.0 （Volta） 的二進(jìn)制目標(biāo)設(shè)備將不會(huì)使用 Tensor Core 指令，因?yàn)檫@些指令在 Pascal 上不可用。 因此，最終二進(jìn)制文件的性能可能會(huì)比使用最新版本的 PTX 生成的二進(jìn)制文件更差。

3.1.4 應(yīng)用程序兼容性

要在具有特定計(jì)算能力的設(shè)備上執(zhí)行代碼，應(yīng)用程序必須加載與此計(jì)算能力兼容的二進(jìn)制或 PTX 代碼，如二進(jìn)制兼容性和 PTX 兼容性中所述。 特別是，為了能夠在具有更高計(jì)算能力的未來架構(gòu)上執(zhí)行代碼（尚無法生成二進(jìn)制代碼），應(yīng)用程序必須加載將為這些設(shè)備實(shí)時(shí)編譯的 PTX 代碼（參見即時(shí)編譯）。

哪些 PTX 和二進(jìn)制代碼嵌入到 CUDA C++ 應(yīng)用程序中由 -arch 和 -code 編譯器選項(xiàng)或 -gencode 編譯器選項(xiàng)控制，詳見 nvcc 用戶手冊。 例如：

nvcc x.cu
        -gencode arch=compute_50,code=sm_50
        -gencode arch=compute_60,code=sm_60
        -gencode arch=compute_70,code="compute_70,sm_70"

嵌入與計(jì)算能力 5.0 和 6.0（第一和第二-gencode選項(xiàng)）兼容的二進(jìn)制代碼以及與計(jì)算能力 7.0（第三-gencode選項(xiàng)）兼容的 PTX 和二進(jìn)制代碼。

生成主機(jī)代碼以在運(yùn)行時(shí)自動(dòng)選擇最合適的代碼來加載和執(zhí)行，在上面的示例中，這些代碼將是：

• 具有計(jì)算能力 5.0 和 5.2 的設(shè)備的 5.0 二進(jìn)制代碼，
• 具有計(jì)算能力 6.0 和 6.1 的設(shè)備的 6.0 二進(jìn)制代碼，
• 具有計(jì)算能力 7.0 和 7.5 的設(shè)備的 7.0 二進(jìn)制代碼，
• PTX 代碼在運(yùn)行時(shí)編譯為具有計(jì)算能力 8.0 和 8.6 的設(shè)備的二進(jìn)制代碼。

例如，x.cu可以有一個(gè)優(yōu)化代碼的方法，使用 warp shuffle 操作，這些操作僅在計(jì)算能力 3.0 及更高版本的設(shè)備中受支持。__CUDA_ARCH__宏可用于根據(jù)計(jì)算能力區(qū)分各種代碼方案。 它僅為設(shè)備代碼定義。 例如，當(dāng)使用-arch=compute_35編譯時(shí)，__CUDA_ARCH__等于 350。

使用驅(qū)動(dòng) API 的應(yīng)用程序必須編譯代碼以分離文件并在運(yùn)行時(shí)顯式加載和執(zhí)行最合適的文件。

Volta 架構(gòu)引入了獨(dú)立線程調(diào)度，它改變了在 GPU 上調(diào)度線程的方式。 對于依賴于以前架構(gòu)中SIMT 調(diào)度的特定行為的代碼，獨(dú)立線程調(diào)度可能會(huì)改變參與線程的集合，從而導(dǎo)致不正確的結(jié)果。 為了在實(shí)現(xiàn)獨(dú)立線程調(diào)度中詳述的糾正措施的同時(shí)幫助遷移，Volta 開發(fā)人員可以使用編譯器選項(xiàng)組合 -arch=compute_60 -code=sm_70 選擇加入 Pascal 的線程調(diào)度。

nvcc 用戶手冊列出了-arch、-code和-gencode編譯器選項(xiàng)的各種簡寫。 例如，-arch=sm_70是-arch=compute_70 -code=compute_70,sm_70的簡寫（與-gencode arch=compute_70,code="compute_70,sm_70"相同）。

3.1.5 C++兼容性

編譯器前端根據(jù) C++ 語法規(guī)則處理 CUDA 源文件。 主機(jī)代碼支持完整的 C++。 但是，設(shè)備代碼僅完全支持 C++ 的一個(gè)子集，如C++ 語言支持中所述。

3.1.6 64位支持

64 位版本的nvcc以 64 位模式編譯設(shè)備代碼（即指針是 64 位的）。 以 64 位模式編譯的設(shè)備代碼僅支持以 64 位模式編譯的主機(jī)代碼。

同樣，32 位版本的nvcc以 32 位模式編譯設(shè)備代碼，而以 32 位模式編譯的設(shè)備代碼僅支持以 32 位模式編譯的主機(jī)代碼。

32 位版本的nvcc也可以使用 -m64 編譯器選項(xiàng)以 64 位模式編譯設(shè)備代碼。

64 位版本的nvcc也可以使用 -m32 編譯器選項(xiàng)以 32 位模式編譯設(shè)備代碼。

3.2 CUDA運(yùn)行時(shí)

運(yùn)行時(shí)在cudart庫中實(shí)現(xiàn)，該庫鏈接到應(yīng)用程序，可以通過cudart.lib或libcudart.a靜態(tài)鏈接，也可以通過cudart.dll或libcudart.so動(dòng)態(tài)鏈接。 需要cudart.dll或cudart.so進(jìn)行動(dòng)態(tài)鏈接的應(yīng)用程序通常將它們作為應(yīng)用程序安裝包的一部分。 只有在鏈接到同一 CUDA 運(yùn)行時(shí)實(shí)例的組件之間傳遞 CUDA 運(yùn)行時(shí)符號(hào)的地址才是安全的。

它的所有入口都以cuda為前綴。

如異構(gòu)編程中所述，CUDA 編程模型假設(shè)系統(tǒng)由主機(jī)和設(shè)備組成，每個(gè)設(shè)備都有自己獨(dú)立的內(nèi)存。設(shè)備內(nèi)存概述了用于管理設(shè)備內(nèi)存的運(yùn)行時(shí)函數(shù)。

共享內(nèi)存說明了使用線程層次結(jié)構(gòu)中引入的共享內(nèi)存來最大化性能。

Page-Locked Host Memory引入了 page-locked 主機(jī)內(nèi)存，它需要將內(nèi)核執(zhí)行與主機(jī)設(shè)備內(nèi)存之間的數(shù)據(jù)傳輸重疊。

異步并發(fā)執(zhí)行描述了用于在系統(tǒng)的各個(gè)級(jí)別啟用異步并發(fā)執(zhí)行的概念和 API。

多設(shè)備系統(tǒng)展示了編程模型如何擴(kuò)展到具有多個(gè)設(shè)備連接到同一主機(jī)的系統(tǒng)。

錯(cuò)誤檢查描述了如何正確檢查運(yùn)行時(shí)生成的錯(cuò)誤。

調(diào)用堆棧提到了用于管理 CUDA C++ 調(diào)用堆棧的運(yùn)行時(shí)函數(shù)。

Texture and Surface Memory呈現(xiàn)了紋理和表面內(nèi)存空間，它們提供了另一種訪問設(shè)備內(nèi)存的方式；它們還公開了 GPU 紋理硬件的一個(gè)子集。

圖形互操作性介紹了運(yùn)行時(shí)提供的各種功能，用于與兩個(gè)主要圖形 API（OpenGL 和 Direct3D）進(jìn)行互操作。

3.2.1 初始化

運(yùn)行時(shí)沒有顯式的初始化函數(shù)；它在第一次調(diào)用運(yùn)行時(shí)函數(shù)時(shí)進(jìn)行初始化（更具體地說，除了參考手冊的錯(cuò)誤處理和版本管理部分中的函數(shù)之外的任何函數(shù)）。在計(jì)時(shí)運(yùn)行時(shí)函數(shù)調(diào)用以及將第一次調(diào)用的錯(cuò)誤代碼解釋到運(yùn)行時(shí)時(shí)，需要牢記這一點(diǎn)。

運(yùn)行時(shí)為系統(tǒng)中的每個(gè)設(shè)備創(chuàng)建一個(gè) CUDA 上下文（有關(guān) CUDA 上下文的更多詳細(xì)信息，請參閱上下文）。此context是此設(shè)備的主要上下文，并在需要此設(shè)備上的活動(dòng)上下文的第一個(gè)運(yùn)行時(shí)函數(shù)中初始化。它在應(yīng)用程序的所有主機(jī)線程之間共享。作為此上下文創(chuàng)建的一部分，設(shè)備代碼會(huì)在必要時(shí)進(jìn)行即時(shí)編譯（請參閱即時(shí)編譯）并加載到設(shè)備內(nèi)存中。這一切都是透明地發(fā)生的。如果需要，例如對于驅(qū)動(dòng)程序 API 互操作性，可以從驅(qū)動(dòng)程序 API 訪問設(shè)備的主要上下文，如運(yùn)行時(shí)和驅(qū)動(dòng)程序 API 之間的互操作性中所述。

當(dāng)主機(jī)線程調(diào)用 cudaDeviceReset() 時(shí)，這會(huì)破壞主機(jī)線程當(dāng)前操作的設(shè)備的主要上下文（即設(shè)備選擇中定義的當(dāng)前設(shè)備）。 任何將此設(shè)備作為當(dāng)前設(shè)備的主機(jī)線程進(jìn)行的下一個(gè)運(yùn)行時(shí)函數(shù)調(diào)用將為該設(shè)備創(chuàng)建一個(gè)新的主上下文。

注意：CUDA接口使用全局狀態(tài)，在主機(jī)程序初始化時(shí)初始化，在主機(jī)程序終止時(shí)銷毀。 CUDA 運(yùn)行時(shí)和驅(qū)動(dòng)程序無法檢測此狀態(tài)是否無效，因此在程序啟動(dòng)或 main 后終止期間使用任何這些接口（隱式或顯式）將導(dǎo)致未定義的行為。

3.2.2 設(shè)備存儲(chǔ)

如異構(gòu)編程中所述，CUDA 編程模型假設(shè)系統(tǒng)由主機(jī)和設(shè)備組成，每個(gè)設(shè)備都有自己獨(dú)立的內(nèi)存。 內(nèi)核在設(shè)備內(nèi)存之外運(yùn)行，因此運(yùn)行時(shí)提供了分配、解除分配和復(fù)制設(shè)備內(nèi)存以及在主機(jī)內(nèi)存和設(shè)備內(nèi)存之間傳輸數(shù)據(jù)的功能。

設(shè)備內(nèi)存可以分配為線性內(nèi)存或CUDA 數(shù)組。

CUDA 數(shù)組是針對紋理獲取優(yōu)化的不透明內(nèi)存布局。 它們在紋理和表面內(nèi)存中有所描述。

線性內(nèi)存分配在一個(gè)統(tǒng)一的地址空間中，這意味著單獨(dú)分配的實(shí)體可以通過指針相互引用，例如在二叉樹或鏈表中。 地址空間的大小取決于主機(jī)系統(tǒng) (CPU) 和所用 GPU 的計(jì)算能力：

Table 1. Linear Memory Address Space

注意：在計(jì)算能力為 5.3 (Maxwell) 及更早版本的設(shè)備上，CUDA 驅(qū)動(dòng)程序會(huì)創(chuàng)建一個(gè)未提交的 40 位虛擬地址預(yù)留，以確保內(nèi)存分配（指針）在支持的范圍內(nèi)。 此預(yù)留顯示為預(yù)留虛擬內(nèi)存，但在程序?qū)嶋H分配內(nèi)存之前不會(huì)占用任何物理內(nèi)存。

線性內(nèi)存通常使用cudaMalloc()分配并使用cudaFree()釋放，主機(jī)內(nèi)存和設(shè)備內(nèi)存之間的數(shù)據(jù)傳輸通常使用cudaMemcpy()完成。 在Kernels的向量加法代碼示例中，需要將向量從主機(jī)內(nèi)存復(fù)制到設(shè)備內(nèi)存：

// Device code
__global__ void VecAdd(float* A, float* B, float* C, int N)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < N)
        C[i] = A[i] + B[i];
}
            
// Host code
int main()
{
    int N = ...;
    size_t size = N * sizeof(float);

    // Allocate input vectors h_A and h_B in host memory
    float* h_A = (float*)malloc(size);
    float* h_B = (float*)malloc(size);
    float* h_C = (float*)malloc(size);

    // Initialize input vectors
    ...

    // Allocate vectors in device memory
    float* d_A;
    cudaMalloc(&d_A, size);
    float* d_B;
    cudaMalloc(&d_B, size);
    float* d_C;
    cudaMalloc(&d_C, size);

    // Copy vectors from host memory to device memory
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    // Invoke kernel
    int threadsPerBlock = 256;
    int blocksPerGrid =
            (N + threadsPerBlock - 1) / threadsPerBlock;
    VecAdd<<>>(d_A, d_B, d_C, N);

    // Copy result from device memory to host memory
    // h_C contains the result in host memory
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    // Free device memory
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
            
    // Free host memory
    ...
}

線性內(nèi)存也可以通過cudaMallocPitch()和cudaMalloc3D()分配。 建議將這些函數(shù)用于 2D 或 3D 數(shù)組的分配，因?yàn)樗_保分配被適當(dāng)?shù)靥畛湟詽M足設(shè)備內(nèi)存訪問中描述的對齊要求，從而確保在訪問行地址或在 2D 數(shù)組和其他區(qū)域設(shè)備內(nèi)存之間執(zhí)行復(fù)制時(shí)獲得最佳性能（使用 cudaMemcpy2D() 和 cudaMemcpy3D() 函數(shù)）。 返回的間距（或步幅）必須用于訪問數(shù)組元素。 以下代碼示例分配一個(gè)width x height的2D浮點(diǎn)數(shù)組，并顯示如何在設(shè)備代碼中循環(huán)遍歷數(shù)組元素：

// Host code
int width = 64, height = 64;
float* devPtr;
size_t pitch;
cudaMallocPitch(&devPtr, &pitch,
                width * sizeof(float), height);
MyKernel<<<100, 512>>>(devPtr, pitch, width, height);

// Device code
__global__ void MyKernel(float* devPtr,
                         size_t pitch, int width, int height)
{
    for (int r = 0; r < height; ++r) {
        float* row = (float*)((char*)devPtr + r * pitch);
        for (int c = 0; c < width; ++c) {
            float element = row[c];
        }
    }
}

以下代碼示例分配了一個(gè)width x height x depth的3D浮點(diǎn)數(shù)組，并展示了如何在設(shè)備代碼中循環(huán)遍歷數(shù)組元素：

// Host code
int width = 64, height = 64, depth = 64;
cudaExtent extent = make_cudaExtent(width * sizeof(float),
                                    height, depth);
cudaPitchedPtr devPitchedPtr;
cudaMalloc3D(&devPitchedPtr, extent);
MyKernel<<<100, 512>>>(devPitchedPtr, width, height, depth);

// Device code
__global__ void MyKernel(cudaPitchedPtr devPitchedPtr,
                         int width, int height, int depth)
{
    char* devPtr = devPitchedPtr.ptr;
    size_t pitch = devPitchedPtr.pitch;
    size_t slicePitch = pitch * height;
    for (int z = 0; z < depth; ++z) {
        char* slice = devPtr + z * slicePitch;
        for (int y = 0; y < height; ++y) {
            float* row = (float*)(slice + y * pitch);
            for (int x = 0; x < width; ++x) {
                float element = row[x];
            }
        }
    }
}

注意：為避免分配過多內(nèi)存從而影響系統(tǒng)范圍的性能，請根據(jù)問題大小向用戶請求分配參數(shù)。 如果分配失敗，您可以回退到其他較慢的內(nèi)存類型（cudaMallocHost()、cudaHostRegister() 等），或者返回一個(gè)錯(cuò)誤，告訴用戶需要多少內(nèi)存被拒絕。 如果您的應(yīng)用程序由于某種原因無法請求分配參數(shù)，我們建議對支持它的平臺(tái)使用 cudaMallocManaged()。

參考手冊列出了用于在使用cudaMalloc()分配的線性內(nèi)存、使用cudaMallocPitch()或cudaMalloc3D()分配的線性內(nèi)存、CUDA 數(shù)組以及為在全局或常量內(nèi)存空間中聲明的變量分配的內(nèi)存之間復(fù)制內(nèi)存的所有各種函數(shù)。

以下代碼示例說明了通過運(yùn)行時(shí) API 訪問全局變量的各種方法：

__constant__ float constData[256];
float data[256];
cudaMemcpyToSymbol(constData, data, sizeof(data));
cudaMemcpyFromSymbol(data, constData, sizeof(data));

__device__ float devData;
float value = 3.14f;
cudaMemcpyToSymbol(devData, &value, sizeof(float));

__device__ float* devPointer;
float* ptr;
cudaMalloc(&ptr, 256 * sizeof(float));
cudaMemcpyToSymbol(devPointer, &ptr, sizeof(ptr));

cudaGetSymbolAddress()用于檢索指向?yàn)槿謨?nèi)存空間中聲明的變量分配的內(nèi)存的地址。 分配內(nèi)存的大小是通過cudaGetSymbolSize()獲得的。

3.2.3 L2級(jí)設(shè)備內(nèi)存管理

當(dāng)一個(gè) CUDA 內(nèi)核重復(fù)訪問全局內(nèi)存中的一個(gè)數(shù)據(jù)區(qū)域時(shí)，這種數(shù)據(jù)訪問可以被認(rèn)為是持久化的。 另一方面，如果數(shù)據(jù)只被訪問一次，那么這種數(shù)據(jù)訪問可以被認(rèn)為是流式的。

從 CUDA 11.0 開始，計(jì)算能力 8.0 及以上的設(shè)備能夠影響 L2 緩存中數(shù)據(jù)的持久性，從而可能提供對全局內(nèi)存的更高帶寬和更低延遲的訪問。

3.2.3.1 為持久訪問預(yù)留L2緩存

可以留出一部分 L2 緩存用于持久化對全局內(nèi)存的數(shù)據(jù)訪問。 持久訪問優(yōu)先使用 L2 緩存的這個(gè)預(yù)留部分，而對全局內(nèi)存的正常訪問或流式訪問只能在持久訪問未使用 L2 的這一部分使用。

可以在以下限制內(nèi)調(diào)整用于持久訪問的 L2 緩存預(yù)留大?。?/p>

cudaGetDeviceProperties(&prop, device_id);                
size_t size = min(int(prop.l2CacheSize * 0.75), prop.persistingL2CacheMaxSize);
cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, size); /* set-aside 3/4 of L2 cache for persisting accesses or the max allowed*/ 

在多實(shí)例 GPU (MIG) 模式下配置 GPU 時(shí)，L2 緩存預(yù)留功能被禁用。

使用多進(jìn)程服務(wù) (MPS) 時(shí)，cudaDeviceSetLimit無法更改 L2 緩存預(yù)留大小。 相反，只能在 MPS 服務(wù)器啟動(dòng)時(shí)通過環(huán)境變量CUDA_DEVICE_DEFAULT_PERSISTING_L2_CACHE_PERCENTAGE_LIMIT指定預(yù)留大小。

3.2.3.2 L2持久化訪問策略

訪問策略窗口指定全局內(nèi)存的連續(xù)區(qū)域和L2緩存中的持久性屬性，用于該區(qū)域內(nèi)的訪問。

下面的代碼示例顯示了如何使用 CUDA 流設(shè)置L2持久訪問窗口。

cudaStreamAttrValue stream_attribute;                                         // Stream level attributes data structure
stream_attribute.accessPolicyWindow.base_ptr  = reinterpret_cast(ptr); // Global Memory data pointer
stream_attribute.accessPolicyWindow.num_bytes = num_bytes;                    // Number of bytes for persistence access.
                                                                              // (Must be less than cudaDeviceProp::accessPolicyMaxWindowSize)
stream_attribute.accessPolicyWindow.hitRatio  = 0.6;                          // Hint for cache hit ratio
stream_attribute.accessPolicyWindow.hitProp   = cudaAccessPropertyPersisting; // Type of access property on cache hit
stream_attribute.accessPolicyWindow.missProp  = cudaAccessPropertyStreaming;  // Type of access property on cache miss.

//Set the attributes to a CUDA stream of type cudaStream_t
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &stream_attribute); 

當(dāng)內(nèi)核隨后在 CUDA 流中執(zhí)行時(shí)，全局內(nèi)存范圍 [ptr..ptr+num_bytes) 內(nèi)的內(nèi)存訪問比對其他全局內(nèi)存位置的訪問更有可能保留在 L2 緩存中。

也可以為 CUDA Graph Kernel Node節(jié)點(diǎn)設(shè)置 L2 持久性，如下例所示：

cudaKernelNodeAttrValue node_attribute;                                     // Kernel level attributes data structure
node_attribute.accessPolicyWindow.base_ptr  = reinterpret_cast(ptr); // Global Memory data pointer
node_attribute.accessPolicyWindow.num_bytes = num_bytes;                    // Number of bytes for persistence access.
                                                                            // (Must be less than cudaDeviceProp::accessPolicyMaxWindowSize)
node_attribute.accessPolicyWindow.hitRatio  = 0.6;                          // Hint for cache hit ratio
node_attribute.accessPolicyWindow.hitProp   = cudaAccessPropertyPersisting; // Type of access property on cache hit
node_attribute.accessPolicyWindow.missProp  = cudaAccessPropertyStreaming;  // Type of access property on cache miss.
                                    
//Set the attributes to a CUDA Graph Kernel node of type cudaGraphNode_t
cudaGraphKernelNodeSetAttribute(node, cudaKernelNodeAttributeAccessPolicyWindow, &node_attribute); 

hitRatio參數(shù)可用于指定接收hitProp屬性的訪問比例。 在上面的兩個(gè)示例中，全局內(nèi)存區(qū)域 [ptr..ptr+num_bytes) 中 60% 的內(nèi)存訪問具有持久屬性，40% 的內(nèi)存訪問具有流屬性。 哪些特定的內(nèi)存訪問被歸類為持久（hitProp）是隨機(jī)的，概率大約為hitRatio； 概率分布取決于硬件架構(gòu)和內(nèi)存范圍。

例如，如果 L2 預(yù)留緩存大小為 16KB，而 accessPolicyWindow 中的 num_bytes 為 32KB：

• hitRatio為 0.5 時(shí)，硬件將隨機(jī)選擇 32KB 窗口中的 16KB 指定為持久化并緩存在預(yù)留的 L2 緩存區(qū)域中。
• hitRatio為 1.0 時(shí)，硬件將嘗試在預(yù)留的 L2 緩存區(qū)域中緩存整個(gè) 32KB 窗口。 由于預(yù)留區(qū)域小于窗口，緩存行將被逐出以將 32KB 數(shù)據(jù)中最近使用的 16KB 保留在 L2 緩存的預(yù)留部分中。

因此，hitRatio可用于避免緩存的破壞，并總體減少移入和移出 L2 高速緩存的數(shù)據(jù)量。

低于 1.0 的hitRatio值可用于手動(dòng)控制來自并發(fā) CUDA 流的不同accessPolicyWindows可以緩存在 L2 中的數(shù)據(jù)量。 例如，讓 L2 預(yù)留緩存大小為 16KB； 兩個(gè)不同 CUDA 流中的兩個(gè)并發(fā)內(nèi)核，每個(gè)都有一個(gè) 16KB 的accessPolicyWindow，并且兩者的hitRatio值都為 1.0，在競爭共享 L2 資源時(shí)，可能會(huì)驅(qū)逐彼此的緩存。 但是，如果兩個(gè)accessPolicyWindows的hitRatio值都為 0.5，則它們將不太可能逐出自己或彼此的持久緩存。

3.2.3.3 L2訪問屬性

為不同的全局內(nèi)存數(shù)據(jù)訪問定義了三種類型的訪問屬性：

1. cudaAccessPropertyStreaming：使用流屬性發(fā)生的內(nèi)存訪問不太可能在 L2 緩存中持續(xù)存在，因?yàn)檫@些訪問優(yōu)先被驅(qū)逐。
2. cudaAccessPropertyPersisting：使用持久屬性發(fā)生的內(nèi)存訪問更有可能保留在 L2 緩存中，因?yàn)檫@些訪問優(yōu)先保留在 L2 緩存的預(yù)留部分中。
3. cudaAccessPropertyNormal：此訪問屬性強(qiáng)制將先前應(yīng)用的持久訪問屬性重置為正常狀態(tài)。來自先前 CUDA 內(nèi)核的具有持久性屬性的內(nèi)存訪問可能會(huì)在其預(yù)期用途之后很長時(shí)間保留在 L2 緩存中。這種使用后的持久性減少了不使用持久性屬性的后續(xù)內(nèi)核可用的 L2 緩存量。使用cudaAccessPropertyNormal屬性重置訪問屬性窗口會(huì)刪除先前訪問的持久（優(yōu)先保留）狀態(tài)，就像先前訪問沒有訪問屬性一樣。

3.2.3.4 L2持久性示例

以下示例顯示如何為持久訪問預(yù)留 L2 緩存，通過 CUDA Stream 在 CUDA 內(nèi)核中使用預(yù)留的 L2 緩存，然后重置 L2 緩存。

cudaStream_t stream;
cudaStreamCreate(&stream);                                                                  // Create CUDA stream

cudaDeviceProp prop;                                                                        // CUDA device properties variable
cudaGetDeviceProperties( &prop, device_id);                                                 // Query GPU properties
size_t size = min( int(prop.l2CacheSize * 0.75) , prop.persistingL2CacheMaxSize );
cudaDeviceSetLimit( cudaLimitPersistingL2CacheSize, size);                                  // set-aside 3/4 of L2 cache for persisting accesses or the max allowed

size_t window_size = min(prop.accessPolicyMaxWindowSize, num_bytes);                        // Select minimum of user defined num_bytes and max window size.

cudaStreamAttrValue stream_attribute;                                                       // Stream level attributes data structure
stream_attribute.accessPolicyWindow.base_ptr  = reinterpret_cast(data1);               // Global Memory data pointer
stream_attribute.accessPolicyWindow.num_bytes = window_size;                                // Number of bytes for persistence access
stream_attribute.accessPolicyWindow.hitRatio  = 0.6;                                        // Hint for cache hit ratio
stream_attribute.accessPolicyWindow.hitProp   = cudaAccessPropertyPersisting;               // Persistence Property
stream_attribute.accessPolicyWindow.missProp  = cudaAccessPropertyStreaming;                // Type of access property on cache miss

cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &stream_attribute);   // Set the attributes to a CUDA Stream

for(int i = 0; i < 10; i++) {
    cuda_kernelA<<>>(data1);                                 // This data1 is used by a kernel multiple times
}                                                                                           // [data1 + num_bytes) benefits from L2 persistence
cuda_kernelB<<>>(data1);                                     // A different kernel in the same stream can also benefit
                                                                                            // from the persistence of data1

stream_attribute.accessPolicyWindow.num_bytes = 0;                                          // Setting the window size to 0 disable it
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &stream_attribute);   // Overwrite the access policy attribute to a CUDA Stream
cudaCtxResetPersistingL2Cache();                                                            // Remove any persistent lines in L2 

cuda_kernelC<<>>(data2);                                     // data2 can now benefit from full L2 in normal mode

3.2.3.5 將L2 Access重置為Normal

來自之前CUDA內(nèi)核的L2緩存在被使用后可能會(huì)長期保存在L2中。因此，L2緩存重設(shè)為正常狀態(tài)對于流或正常內(nèi)存訪問很重要，以便以正常優(yōu)先級(jí)使用L2緩存。有三種方法可以將持久訪問重置為正常狀態(tài)。

1. 使用訪問屬性cudaAccessPropertyNormal重置之前的持久化內(nèi)存區(qū)域。
2. 通過調(diào)用cudaCtxResetPersistingL2Cache()將所有持久L2緩存線重置為正常。
3. 最終，未觸及的空間會(huì)自動(dòng)重置為正常。對自動(dòng)復(fù)位的依賴性很強(qiáng)

3.2.3.6 管理L2預(yù)留緩存的利用率

在不同 CUDA 流中同時(shí)執(zhí)行的多個(gè) CUDA 內(nèi)核可能具有分配給它們的流的不同訪問策略窗口。 但是，L2 預(yù)留緩存部分在所有這些并發(fā) CUDA 內(nèi)核之間共享。 因此，這個(gè)預(yù)留緩存部分的凈利用率是所有并發(fā)內(nèi)核單獨(dú)使用的總和。 將內(nèi)存訪問指定為持久訪問的好處會(huì)隨著持久訪問的數(shù)量超過預(yù)留的 L2 緩存容量而減少。

要管理預(yù)留 L2 緩存部分的利用率，應(yīng)用程序必須考慮以下事項(xiàng)：

• L2 預(yù)留緩存的大小。
• 可以同時(shí)執(zhí)行的 CUDA 內(nèi)核。
• 可以同時(shí)執(zhí)行的所有 CUDA 內(nèi)核的訪問策略窗口。
• 何時(shí)以及如何需要 L2 重置以允許正常或流式訪問以同等優(yōu)先級(jí)利用先前預(yù)留的 L2 緩存。

3.2.3.7 查詢L2緩存屬性

與 L2 緩存相關(guān)的屬性是cudaDeviceProp結(jié)構(gòu)的一部分，可以使用 CUDA 運(yùn)行時(shí) APIcudaGetDeviceProperties進(jìn)行查詢

CUDA 設(shè)備屬性包括：

• l2CacheSize：GPU 上可用的二級(jí)緩存數(shù)量。
• persistingL2CacheMaxSize：可以為持久內(nèi)存訪問留出的 L2 緩存的最大數(shù)量。
• accessPolicyMaxWindowSize：訪問策略窗口的最大尺寸。

3.2.3.8 控制L2緩存預(yù)留大小用于持久內(nèi)存訪問

使用 CUDA 運(yùn)行時(shí) APIcudaDeviceGetLimit查詢用于持久內(nèi)存訪問的 L2 預(yù)留緩存大小，并使用 CUDA 運(yùn)行時(shí) APIcudaDeviceSetLimit作為cudaLimit進(jìn)行設(shè)置。 設(shè)置此限制的最大值是cudaDeviceProp::persistingL2CacheMaxSize。

enum cudaLimit {
    /* other fields not shown */
    cudaLimitPersistingL2CacheSize
}; 

3.2.4共享內(nèi)存

如可變內(nèi)存空間說明中所述，共享內(nèi)存是使用__shared__內(nèi)存空間說明符分配的。

正如線程層次結(jié)構(gòu)中提到的和共享內(nèi)存中詳述的那樣，共享內(nèi)存預(yù)計(jì)比全局內(nèi)存快得多。 它可以用作暫存器內(nèi)存（或軟件管理的緩存），以最大限度地減少來自 CUDA 塊的全局內(nèi)存訪問，如下面的矩陣乘法示例所示。

以下代碼示例是不利用共享內(nèi)存的矩陣乘法的簡單實(shí)現(xiàn)。 每個(gè)線程讀取 A 的一行和 B 的一列，并計(jì)算 C 的相應(yīng)元素，如圖所示。因此，從全局內(nèi)存中讀取 A 為 B.width 次，而 B 為讀取 A.height 次。

// Matrices are stored in row-major order:
// M(row, col) = *(M.elements + row * M.width + col)
typedef struct {
    int width;
    int height;
    float* elements;
} Matrix;

// Thread block size
#define BLOCK_SIZE 16

// Forward declaration of the matrix multiplication kernel
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);

// Matrix multiplication - Host code
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE
void MatMul(const Matrix A, const Matrix B, Matrix C)
{
    // Load A and B to device memory
    Matrix d_A;
    d_A.width = A.width; d_A.height = A.height;
    size_t size = A.width * A.height * sizeof(float);
    cudaMalloc(&d_A.elements, size);
    cudaMemcpy(d_A.elements, A.elements, size,
               cudaMemcpyHostToDevice);
    Matrix d_B;
    d_B.width = B.width; d_B.height = B.height;
    size = B.width * B.height * sizeof(float);
    cudaMalloc(&d_B.elements, size);
    cudaMemcpy(d_B.elements, B.elements, size,
               cudaMemcpyHostToDevice);

    // Allocate C in device memory
    Matrix d_C;
    d_C.width = C.width; d_C.height = C.height;
    size = C.width * C.height * sizeof(float);
    cudaMalloc(&d_C.elements, size);

    // Invoke kernel
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
    dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);
    MatMulKernel<<>>(d_A, d_B, d_C);

    // Read C from device memory
    cudaMemcpy(C.elements, d_C.elements, size,
               cudaMemcpyDeviceToHost);

    // Free device memory
    cudaFree(d_A.elements);
    cudaFree(d_B.elements);
    cudaFree(d_C.elements);
}

// Matrix multiplication kernel called by MatMul()
__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C)
{
    // Each thread computes one element of C
    // by accumulating results into Cvalue
    float Cvalue = 0;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    for (int e = 0; e < A.width; ++e)
        Cvalue += A.elements[row * A.width + e]
                * B.elements[e * B.width + col];
    C.elements[row * C.width + col] = Cvalue;
}

以下代碼示例是利用共享內(nèi)存的矩陣乘法實(shí)現(xiàn)。在這個(gè)實(shí)現(xiàn)中，每個(gè)線程塊負(fù)責(zé)計(jì)算C的一個(gè)方形子矩陣Csub，塊內(nèi)的每個(gè)線程負(fù)責(zé)計(jì)算Csub的一個(gè)元素。如圖所示，Csub 等于兩個(gè)矩形矩陣的乘積：維度 A 的子矩陣 (A.width, block_size) 與 Csub 具有相同的行索引，以及維度 B 的子矩陣(block_size, A.width ) 具有與 Csub 相同的列索引。為了適應(yīng)設(shè)備的資源，這兩個(gè)矩形矩陣根據(jù)需要被分成多個(gè)尺寸為 block_size 的方陣，并且 Csub 被計(jì)算為這些方陣的乘積之和。這些乘積中的每一個(gè)都是通過首先將兩個(gè)對應(yīng)的方陣從全局內(nèi)存加載到共享內(nèi)存中的，一個(gè)線程加載每個(gè)矩陣的一個(gè)元素，然后讓每個(gè)線程計(jì)算乘積的一個(gè)元素。每個(gè)線程將這些乘積中的每一個(gè)的結(jié)果累積到一個(gè)寄存器中，并在完成后將結(jié)果寫入全局內(nèi)存。

通過以這種方式將計(jì)算分塊，我們利用了快速共享內(nèi)存并節(jié)省了大量的全局內(nèi)存帶寬，因?yàn)?A 只從全局內(nèi)存中讀取 (B.width / block_size) 次，而 B 被讀取 (A.height / block_size) 次.

前面代碼示例中的 Matrix 類型增加了一個(gè) stride 字段，因此子矩陣可以用相同的類型有效地表示。__device__函數(shù)用于獲取和設(shè)置元素并從矩陣構(gòu)建任何子矩陣。

// Matrices are stored in row-major order:
// M(row, col) = *(M.elements + row * M.stride + col)
typedef struct {
    int width;
    int height;
    int stride; 
    float* elements;
} Matrix;

// Get a matrix element
__device__ float GetElement(const Matrix A, int row, int col)
{
    return A.elements[row * A.stride + col];
}

// Set a matrix element
__device__ void SetElement(Matrix A, int row, int col,
                           float value)
{
    A.elements[row * A.stride + col] = value;
}

// Get the BLOCK_SIZExBLOCK_SIZE sub-matrix Asub of A that is
// located col sub-matrices to the right and row sub-matrices down
// from the upper-left corner of A
 __device__ Matrix GetSubMatrix(Matrix A, int row, int col) 
{
    Matrix Asub;
    Asub.width    = BLOCK_SIZE;
    Asub.height   = BLOCK_SIZE;
    Asub.stride   = A.stride;
    Asub.elements = &A.elements[A.stride * BLOCK_SIZE * row
                                         + BLOCK_SIZE * col];
    return Asub;
}

// Thread block size
#define BLOCK_SIZE 16

// Forward declaration of the matrix multiplication kernel
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);

// Matrix multiplication - Host code
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE
void MatMul(const Matrix A, const Matrix B, Matrix C)
{
    // Load A and B to device memory
    Matrix d_A;
    d_A.width = d_A.stride = A.width; d_A.height = A.height;
    size_t size = A.width * A.height * sizeof(float);
    cudaMalloc(&d_A.elements, size);
    cudaMemcpy(d_A.elements, A.elements, size,
               cudaMemcpyHostToDevice);
    Matrix d_B;
    d_B.width = d_B.stride = B.width; d_B.height = B.height;
    size = B.width * B.height * sizeof(float);
    cudaMalloc(&d_B.elements, size);
    cudaMemcpy(d_B.elements, B.elements, size,
    cudaMemcpyHostToDevice);

    // Allocate C in device memory
    Matrix d_C;
    d_C.width = d_C.stride = C.width; d_C.height = C.height;
    size = C.width * C.height * sizeof(float);
    cudaMalloc(&d_C.elements, size);

    // Invoke kernel
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
    dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);
    MatMulKernel<<>>(d_A, d_B, d_C);

    // Read C from device memory
    cudaMemcpy(C.elements, d_C.elements, size,
               cudaMemcpyDeviceToHost);

    // Free device memory
    cudaFree(d_A.elements);
    cudaFree(d_B.elements);
    cudaFree(d_C.elements);
}

// Matrix multiplication kernel called by MatMul()
 __global__ void MatMulKernel(Matrix A, Matrix B, Matrix C)
{
    // Block row and column
    int blockRow = blockIdx.y;
    int blockCol = blockIdx.x;

    // Each thread block computes one sub-matrix Csub of C
    Matrix Csub = GetSubMatrix(C, blockRow, blockCol);

    // Each thread computes one element of Csub
    // by accumulating results into Cvalue
    float Cvalue = 0;

    // Thread row and column within Csub
    int row = threadIdx.y;
    int col = threadIdx.x;

    // Loop over all the sub-matrices of A and B that are
    // required to compute Csub
    // Multiply each pair of sub-matrices together
    // and accumulate the results
    for (int m = 0; m < (A.width / BLOCK_SIZE); ++m) {

        // Get sub-matrix Asub of A
        Matrix Asub = GetSubMatrix(A, blockRow, m);

        // Get sub-matrix Bsub of B
        Matrix Bsub = GetSubMatrix(B, m, blockCol);

        // Shared memory used to store Asub and Bsub respectively
        __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
        __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];

        // Load Asub and Bsub from device memory to shared memory
        // Each thread loads one element of each sub-matrix
        As[row][col] = GetElement(Asub, row, col);
        Bs[row][col] = GetElement(Bsub, row, col);

        // Synchronize to make sure the sub-matrices are loaded
        // before starting the computation
        __syncthreads();
        // Multiply Asub and Bsub together
        for (int e = 0; e < BLOCK_SIZE; ++e)
            Cvalue += As[row][e] * Bs[e][col];

        // Synchronize to make sure that the preceding
        // computation is done before loading two new
        // sub-matrices of A and B in the next iteration
        __syncthreads();
    }

    // Write Csub to device memory
    // Each thread writes one element
    SetElement(Csub, row, col, Cvalue);
}

3.2.5 Page-Locked主機(jī)內(nèi)存

運(yùn)行時(shí)提供的函數(shù)允許使用鎖頁（也稱為固定）主機(jī)內(nèi)存（與 malloc() 分配的常規(guī)可分頁主機(jī)內(nèi)存相反）：

• cudaHostAlloc()和cudaFreeHost()分配和釋放鎖頁主機(jī)內(nèi)存；
• cudaHostRegister()將malloc()分配的內(nèi)存范圍變?yōu)殒i頁內(nèi)存（有關(guān)限制，請參閱參考手冊）。

使用頁面鎖定的主機(jī)內(nèi)存有幾個(gè)好處：

• 鎖頁主機(jī)內(nèi)存和設(shè)備內(nèi)存之間的復(fù)制可以與異步并發(fā)執(zhí)行中提到的某些設(shè)備的內(nèi)核執(zhí)行同時(shí)執(zhí)行。
• 在某些設(shè)備上，鎖頁主機(jī)內(nèi)存可以映射到設(shè)備的地址空間，從而無需將其復(fù)制到設(shè)備內(nèi)存或從設(shè)備內(nèi)存復(fù)制，如映射內(nèi)存中所述。
• 在具有前端總線的系統(tǒng)上，如果主機(jī)內(nèi)存被分配為頁鎖定，則主機(jī)內(nèi)存和設(shè)備內(nèi)存之間的帶寬更高，如果另外分配為合并訪存，則它甚至更高，如合并寫入內(nèi)存中所述。

然而，鎖頁主機(jī)內(nèi)存是一種稀缺資源，因此鎖頁內(nèi)存中的分配將在可分頁內(nèi)存中分配之前很久就開始失敗。 此外，通過減少操作系統(tǒng)可用于分頁的物理內(nèi)存量，消耗過多的頁面鎖定內(nèi)存會(huì)降低整體系統(tǒng)性能。

注意：頁面鎖定的主機(jī)內(nèi)存不會(huì)緩存在非 I/O 一致的 Tegra 設(shè)備上。 此外，非 I/O 一致的 Tegra 設(shè)備不支持 cudaHostRegister()。

簡單的零拷貝 CUDA 示例附帶關(guān)于頁面鎖定內(nèi)存 API 的詳細(xì)文檔。

3.2.5.1 Portable Memory

一塊鎖頁內(nèi)存可以與系統(tǒng)中的任何設(shè)備一起使用（有關(guān)多設(shè)備系統(tǒng)的更多詳細(xì)信息，請參閱多設(shè)備系統(tǒng)），但默認(rèn)情況下，使用上述鎖頁內(nèi)存的好處只是與分配塊時(shí)當(dāng)前的設(shè)備一起可用（并且所有設(shè)備共享相同的統(tǒng)一地址空間，如果有，如統(tǒng)一虛擬地址空間中所述）。塊需要通過將標(biāo)志cudaHostAllocPortable傳遞給cudaHostAlloc()來分配，或者通過將標(biāo)志cudaHostRegisterPortable傳遞給cudaHostRegister()來鎖定頁面。

3.2.5.2 寫合并內(nèi)存

默認(rèn)情況下，鎖頁主機(jī)內(nèi)存被分配為可緩存的。它可以選擇分配為寫組合，而不是通過將標(biāo)志cudaHostAllocWriteCombined傳遞給cudaHostAlloc()。 寫入組合內(nèi)存釋放了主機(jī)的 L1 和 L2 緩存資源，為應(yīng)用程序的其余部分提供更多緩存。 此外，在通過 PCI Express 總線的傳輸過程中，寫入組合內(nèi)存不會(huì)被窺探，這可以將傳輸性能提高多達(dá) 40%。

從主機(jī)讀取寫組合內(nèi)存非常慢，因此寫組合內(nèi)存通常應(yīng)用于僅主機(jī)寫入的內(nèi)存。

應(yīng)避免在 WC 內(nèi)存上使用 CPU 原子指令，因?yàn)椴⒎撬?CPU 實(shí)現(xiàn)都能保證該功能。

3.2.5.3 Mapped Memory

通過將標(biāo)志cudaHostAllocMapped傳遞給cudaHostAlloc()或通過將標(biāo)志cudaHostRegisterMapped傳遞給cudaHostRegister()，也可以將鎖頁主機(jī)內(nèi)存塊映射到設(shè)備的地址空間。因此，這樣的塊通常有兩個(gè)地址：一個(gè)在主機(jī)內(nèi)存中，由cudaHostAlloc()或malloc()返回，另一個(gè)在設(shè)備內(nèi)存中，可以使用cudaHostGetDevicePointer()檢索，然后用于從內(nèi)核中訪問該塊。唯一的例外是使用cudaHostAlloc()分配的指針，以及統(tǒng)一虛擬地址空間中提到的主機(jī)和設(shè)備使用統(tǒng)一地址空間。

直接從內(nèi)核中訪問主機(jī)內(nèi)存不會(huì)提供與設(shè)備內(nèi)存相同的帶寬，但確實(shí)有一些優(yōu)勢：

• 無需在設(shè)備內(nèi)存中分配一個(gè)塊，并在該塊和主機(jī)內(nèi)存中的塊之間復(fù)制數(shù)據(jù)；數(shù)據(jù)傳輸是根據(jù)內(nèi)核的需要隱式執(zhí)行的；
• 無需使用流（請參閱并發(fā)數(shù)據(jù)傳輸）將數(shù)據(jù)傳輸與內(nèi)核執(zhí)行重疊；內(nèi)核發(fā)起的數(shù)據(jù)傳輸自動(dòng)與內(nèi)核執(zhí)行重疊。

然而，由于映射的鎖頁內(nèi)存在主機(jī)和設(shè)備之間共享，因此應(yīng)用程序必須使用流或事件同步內(nèi)存訪問（請參閱異步并發(fā)執(zhí)行）以避免任何潛在的 read-after-write、write-after-read 或 write-after-write危險(xiǎn)。

為了能夠檢索到任何映射的鎖頁內(nèi)存的設(shè)備指針，必須在執(zhí)行任何其他 CUDA 調(diào)用之前通過使用cudaDeviceMapHost標(biāo)志調(diào)用cudaSetDeviceFlags()來啟用頁面鎖定內(nèi)存映射。否則，cudaHostGetDevicePointer()將返回錯(cuò)誤。

如果設(shè)備不支持映射的鎖頁主機(jī)內(nèi)存，cudaHostGetDevicePointer()也會(huì)返回錯(cuò)誤。應(yīng)用程序可以通過檢查canMapHostMemory設(shè)備屬性（請參閱[設(shè)備枚舉](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#device-enumeration）來查詢此功能，對于支持映射鎖頁主機(jī)內(nèi)存的設(shè)備，該屬性等于1。

請注意，從主機(jī)或其他設(shè)備的角度來看，在映射的鎖頁內(nèi)存上運(yùn)行的原子函數(shù)（請參閱原子函數(shù)）不是原子的。

另請注意，CUDA 運(yùn)行時(shí)要求從主機(jī)和其他設(shè)備的角度來看，從設(shè)備啟動(dòng)到主機(jī)內(nèi)存的 1 字節(jié)、2 字節(jié)、4 字節(jié)和 8 字節(jié)自然對齊的加載和存儲(chǔ)保留為單一訪問設(shè)備。在某些平臺(tái)上，內(nèi)存的原子操作可能會(huì)被硬件分解為單獨(dú)的加載和存儲(chǔ)操作。這些組件加載和存儲(chǔ)操作對保留自然對齊的訪問具有相同的要求。例如，CUDA 運(yùn)行時(shí)不支持 PCI Express 總線拓?fù)?，其?PCI Express 橋?qū)?8 字節(jié)自然對齊的寫入拆分為設(shè)備和主機(jī)之間的兩個(gè) 4 字節(jié)寫入。

3.2.6 異步并發(fā)執(zhí)行

CUDA 將以下操作公開為可以彼此同時(shí)操作的獨(dú)立任務(wù)：

• 在主機(jī)上計(jì)算；
• 設(shè)備上的計(jì)算；
• 從主機(jī)到設(shè)備的內(nèi)存?zhèn)鬏敚?/li>
• 從設(shè)備到主機(jī)的內(nèi)存?zhèn)鬏敚?/li>
• 在給定設(shè)備的內(nèi)存中進(jìn)行內(nèi)存?zhèn)鬏敚?/li>
• 設(shè)備之間的內(nèi)存?zhèn)鬏敗?/li>

這些操作之間實(shí)現(xiàn)的并發(fā)級(jí)別將取決于設(shè)備的功能和計(jì)算能力，如下所述。

3.2.6.1 主機(jī)和設(shè)備之間的并發(fā)執(zhí)行

在設(shè)備完成請求的任務(wù)之前，異步庫函數(shù)將控制權(quán)返回給宿主線程，從而促進(jìn)了主機(jī)的并發(fā)執(zhí)行。使用異步調(diào)用，許多設(shè)備操作可以在適當(dāng)?shù)脑O(shè)備資源可用時(shí)排隊(duì)，由CUDA驅(qū)動(dòng)程序執(zhí)行。這減輕了主機(jī)線程管理設(shè)備的大部分責(zé)任，讓它自由地執(zhí)行其他任務(wù)。以下設(shè)備操作對主機(jī)是異步的:

• 內(nèi)核啟動(dòng);
• 內(nèi)存復(fù)制在單個(gè)設(shè)備的內(nèi)存中;
• 從主機(jī)到設(shè)備內(nèi)存拷貝的內(nèi)存塊大小不超過64kb的;
• 由帶有Async后綴的函數(shù)執(zhí)行的內(nèi)存拷貝;
• 內(nèi)存設(shè)置函數(shù)調(diào)用。 程序員可以通過將CUDA_LAUNCH_BLOCKING環(huán)境變量設(shè)置為1來全局禁用系統(tǒng)上運(yùn)行的所有CUDA應(yīng)用程序的內(nèi)核啟動(dòng)的異步性。此特性僅用于調(diào)試目的，不應(yīng)用作使生產(chǎn)軟件可靠運(yùn)行的一種方法。

如果通過分析器（Nsight、Visual Profiler）收集硬件計(jì)數(shù)器，則內(nèi)核啟動(dòng)是同步的，除非啟用了并發(fā)內(nèi)核分析。如果異步內(nèi)存復(fù)制涉及非頁面鎖定的主機(jī)內(nèi)存，它們也將是同步的。

3.2.6.2 并行執(zhí)行內(nèi)核

某些計(jì)算能力 2.x 及更高版本的設(shè)備可以同時(shí)執(zhí)行多個(gè)內(nèi)核。 應(yīng)用程序可以通過檢查concurrentKernels設(shè)備屬性（請參閱設(shè)備枚舉）來查詢此功能，對于支持它的設(shè)備，該屬性等于 1。

設(shè)備可以同時(shí)執(zhí)行的內(nèi)核啟動(dòng)的最大數(shù)量取決于其計(jì)算能力，并在表15中列出。

來自一個(gè) CUDA 上下文的內(nèi)核不能與來自另一個(gè) CUDA 上下文的內(nèi)核同時(shí)執(zhí)行。

使用許多紋理或大量本地內(nèi)存的內(nèi)核不太可能與其他內(nèi)核同時(shí)執(zhí)

3.2.6.3 數(shù)據(jù)傳輸和內(nèi)核執(zhí)行的重疊

一些設(shè)備可以在內(nèi)核執(zhí)行的同時(shí)執(zhí)行與 GPU 之間的異步內(nèi)存復(fù)制。 應(yīng)用程序可以通過檢查asyncEngineCount設(shè)備屬性（請參閱設(shè)備枚舉）來查詢此功能，對于支持它的設(shè)備，該屬性大于零。 如果復(fù)制中涉及主機(jī)內(nèi)存，則它必須是頁鎖定的。

還可以與內(nèi)核執(zhí)行（在支持concurrentKernels設(shè)備屬性的設(shè)備上）或與設(shè)備之間的拷貝（對于支持asyncEngineCount屬性的設(shè)備）同時(shí)執(zhí)行設(shè)備內(nèi)復(fù)制。 使用標(biāo)準(zhǔn)內(nèi)存復(fù)制功能啟動(dòng)設(shè)備內(nèi)復(fù)制，目標(biāo)地址和源地址位于同一設(shè)備上。

3.2.6.4 并行數(shù)據(jù)傳輸

某些計(jì)算能力為 2.x 及更高版本的設(shè)備可以重疊設(shè)備之間的數(shù)據(jù)拷貝。 應(yīng)用程序可以通過檢查asyncEngineCount設(shè)備屬性（請參閱設(shè)備枚舉）來查詢此功能，對于支持它的設(shè)備，該屬性等于 2。 為了重疊，傳輸中涉及的任何主機(jī)內(nèi)存都必須是頁面鎖定的。

3.2.6.5 流

應(yīng)用程序通過流管理上述并發(fā)操作。 流是按順序執(zhí)行的命令序列（可能由不同的主機(jī)線程發(fā)出）。 另一方面，不同的流可能會(huì)彼此亂序或同時(shí)執(zhí)行它們的命令； 不能保證此行為，因此不應(yīng)依賴其正確性（例如，內(nèi)核間通信未定義）。 當(dāng)滿足命令的所有依賴項(xiàng)時(shí)，可以執(zhí)行在流上發(fā)出的命令。 依賴關(guān)系可以是先前在同一流上啟動(dòng)的命令或來自其他流的依賴關(guān)系。 同步調(diào)用的成功完成保證了所有啟動(dòng)的命令都完成了。

3.2.6.5.1 創(chuàng)建與銷毀

流是通過創(chuàng)建一個(gè)流對象并將其指定為一系列內(nèi)核啟動(dòng)和主機(jī) <-> 設(shè)備內(nèi)存拷貝的流參數(shù)來定義的。 以下代碼示例創(chuàng)建兩個(gè)流并在鎖頁內(nèi)存中分配一個(gè)浮點(diǎn)數(shù)組hostPtr。

cudaStream_t stream[2];
for (int i = 0; i < 2; ++i)
    cudaStreamCreate(&stream[i]);
float* hostPtr;
cudaMallocHost(&hostPtr, 2 * size);

這些流中的每一個(gè)都由以下代碼示例定義為從主機(jī)到設(shè)備的一次內(nèi)存復(fù)制、一次內(nèi)核啟動(dòng)和從設(shè)備到主機(jī)的一次內(nèi)存復(fù)制的序列：

for (int i = 0; i < 2; ++i) {
    cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
                    size, cudaMemcpyHostToDevice, stream[i]);
    MyKernel <<<100, 512, 0, stream[i]>>>
          (outputDevPtr + i * size, inputDevPtr + i * size, size);
    cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,
                    size, cudaMemcpyDeviceToHost, stream[i]);
}

每個(gè)流將其輸入數(shù)組hostPtr的部分復(fù)制到設(shè)備內(nèi)存中的數(shù)組inputDevPtr，通過調(diào)用MyKernel()處理設(shè)備上的inputDevPtr，并將結(jié)果outputDevPtr復(fù)制回hostPtr的同一部分。 重疊行為描述了此示例中的流如何根據(jù)設(shè)備的功能重疊。 請注意，hostPtr必須指向鎖頁主機(jī)內(nèi)存才能發(fā)生重疊。

通過調(diào)用cudaStreamDestroy()釋放流:

for (int i = 0; i < 2; ++i)
    cudaStreamDestroy(stream[i]);

如果調(diào)用cudaStreamDestroy()時(shí)設(shè)備仍在流中工作，則該函數(shù)將立即返回，并且一旦設(shè)備完成流中的所有工作，與流關(guān)聯(lián)的資源將自動(dòng)釋放。

3.2.6.5.2 默認(rèn)流

未指定任何流參數(shù)或等效地將流參數(shù)設(shè)置為零的內(nèi)核啟動(dòng)和主機(jī) <-> 設(shè)備內(nèi)存拷貝將發(fā)布到默認(rèn)流。因此它們按順序執(zhí)行。

對于使用--default-stream per-thread編譯標(biāo)志編譯的代碼（或在包含 CUDA 頭文件（cuda.h 和 cuda_runtime.h）之前定義CUDA_API_PER_THREAD_DEFAULT_STREAM宏），默認(rèn)流是常規(guī)流，并且每個(gè)主機(jī)線程有自己的默認(rèn)流。

注意：當(dāng)代碼由 nvcc 編譯時(shí)，#define CUDA_API_PER_THREAD_DEFAULT_STREAM 1不能用于啟用此行為，因?yàn)?nvcc 在翻譯單元的頂部隱式包含 cuda_runtime.h。在這種情況下，需要使用--default-stream每個(gè)線程編譯標(biāo)志，或者需要使用-DCUDA_API_PER_THREAD_DEFAULT_STREAM=1編譯器標(biāo)志定義CUDA_API_PER_THREAD_DEFAULT_STREAM宏。

對于使用--default-stream legacy編譯標(biāo)志編譯的代碼，默認(rèn)流是稱為 NULL 流的特殊流，每個(gè)設(shè)備都有一個(gè)用于所有主機(jī)線程的 NULL 流。 NULL 流很特殊，因?yàn)樗鼤?huì)導(dǎo)致隱式同步，如隱式同步中所述。

對于在沒有指定--default-stream編譯標(biāo)志的情況下編譯的代碼，--default-stream legacy被假定為默認(rèn)值。

3.2.6.5.3 顯式同步

有多種方法可以顯式地同步流。

cudaDeviceSynchronize()一直等待，直到所有主機(jī)線程的所有流中的所有先前命令都完成。

cudaStreamSynchronize()將流作為參數(shù)并等待，直到給定流中的所有先前命令都已完成。 它可用于將主機(jī)與特定流同步，允許其他流繼續(xù)在設(shè)備上執(zhí)行。

cudaStreamWaitEvent()將流和事件作為參數(shù)（有關(guān)事件的描述，請參閱事件），并在調(diào)用cudaStreamWaitEvent()后使添加到給定流的所有命令延遲執(zhí)行，直到給定事件完成。

cudaStreamQuery()為應(yīng)用程序提供了一種方法來了解流中所有前面的命令是否已完成。

3.2.6.5.4 隱式同步

如果主機(jī)線程在它們之間發(fā)出以下任一操作，則來自不同流的兩個(gè)命令不能同時(shí)運(yùn)行：

• 頁面鎖定的主機(jī)內(nèi)存分配，
• 設(shè)備內(nèi)存分配，
• 設(shè)備內(nèi)存設(shè)置，
• 兩個(gè)地址之間的內(nèi)存拷貝到同一設(shè)備內(nèi)存，
• 對 NULL 流的任何 CUDA 命令，
• 計(jì)算能力 3.x?和計(jì)算能力 7.x中描述的 L1/共享內(nèi)存配置之間的切換。

對于支持并發(fā)內(nèi)核執(zhí)行且計(jì)算能力為 3.0 或更低的設(shè)備，任何需要依賴項(xiàng)檢查以查看流內(nèi)核啟動(dòng)是否完成的操作：

• 僅當(dāng)從 CUDA 上下文中的任何流啟動(dòng)的所有先前內(nèi)核的所有線程塊都已開始執(zhí)行時(shí)，才能開始執(zhí)行；
• 阻止所有以后從 CUDA 上下文中的任何流啟動(dòng)內(nèi)核，直到檢查內(nèi)核啟動(dòng)完成。

需要依賴檢查的操作包括與正在檢查的啟動(dòng)相同的流中的任何其他命令以及對該流的任何cudaStreamQuery()調(diào)用。 因此，應(yīng)用程序應(yīng)遵循以下準(zhǔn)則來提高并發(fā)內(nèi)核執(zhí)行的潛力：

• 所有獨(dú)立操作都應(yīng)該在依賴操作之前發(fā)出，
• 任何類型的同步都應(yīng)該盡可能地延遲。

3.2.6.5.5 重疊行為

兩個(gè)流之間的執(zhí)行重疊量取決于向每個(gè)流發(fā)出命令的順序以及設(shè)備是否支持?jǐn)?shù)據(jù)傳輸和內(nèi)核執(zhí)行的重疊（請參閱數(shù)據(jù)傳輸和內(nèi)核執(zhí)行的重疊）、并發(fā)內(nèi)核執(zhí)行（ 請參閱并發(fā)內(nèi)核執(zhí)行）和并發(fā)數(shù)據(jù)傳輸（請參閱并發(fā)數(shù)據(jù)傳輸）。

例如,在設(shè)備不支持并行數(shù)據(jù)傳輸,這兩個(gè)流的代碼示例創(chuàng)建和銷毀不重疊,因?yàn)橛蓅tream[1]發(fā)起的內(nèi)存復(fù)制會(huì)在stream[0]發(fā)起的內(nèi)存復(fù)制之后執(zhí)行。如果代碼以以下方式重寫(并且假設(shè)設(shè)備支持?jǐn)?shù)據(jù)傳輸和內(nèi)核執(zhí)行的重疊)

for (int i = 0; i < 2; ++i)
    cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
                    size, cudaMemcpyHostToDevice, stream[i]);
for (int i = 0; i < 2; ++i)
    MyKernel<<<100, 512, 0, stream[i]>>>
          (outputDevPtr + i * size, inputDevPtr + i * size, size);
for (int i = 0; i < 2; ++i)
    cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,
                    size, cudaMemcpyDeviceToHost, stream[i]);

那么在stream[1]上從主機(jī)到設(shè)備的內(nèi)存復(fù)制 與stream[0]上內(nèi)核啟動(dòng)重疊。

在支持并發(fā)數(shù)據(jù)傳輸?shù)脑O(shè)備上，Creation 和 Destruction的代碼示例的兩個(gè)流確實(shí)重疊：在stream[1]上從主機(jī)到設(shè)備的內(nèi)存復(fù)制 與在stream[0]上從設(shè)備到主機(jī)的內(nèi)存復(fù)制甚至在stream[0]上內(nèi)核啟動(dòng)（假設(shè)設(shè)備支持?jǐn)?shù)據(jù)傳輸和內(nèi)核執(zhí)行的重疊）。但是，對于計(jì)算能力為 3.0 或更低的設(shè)備，內(nèi)核執(zhí)行不可能重疊，因?yàn)樵趕tream[0]上從設(shè)備到主機(jī)的內(nèi)存復(fù)制之后，第二次在stream[1]上內(nèi)核啟動(dòng)，因此它被阻塞，直到根據(jù)隱式同步，在stream[0]上第一個(gè)內(nèi)核啟動(dòng)已完成。如果代碼如上重寫，內(nèi)核執(zhí)行重疊（假設(shè)設(shè)備支持并發(fā)內(nèi)核執(zhí)行），因?yàn)樵趕tream[0]上從設(shè)備到主機(jī)的內(nèi)存復(fù)制之前，第二次在stream[1]上內(nèi)核啟動(dòng)被。但是，在這種情況下，根據(jù)隱式同步，在stream[0]上從設(shè)備到主機(jī)的內(nèi)存復(fù)制僅與在stream[1]上內(nèi)核啟動(dòng)的最后一個(gè)線程塊重疊，這只能代表總數(shù)的一小部分內(nèi)核的執(zhí)行時(shí)間。

3.2.6.5.6 Host函數(shù)(回調(diào))

運(yùn)行時(shí)提供了一種通過cudaLaunchHostFunc()在任何點(diǎn)將 CPU 函數(shù)調(diào)用插入到流中的方法。 在回調(diào)之前向流發(fā)出的所有命令都完成后，在主機(jī)上執(zhí)行提供的函數(shù)。

以下代碼示例在向每個(gè)流發(fā)出主機(jī)到設(shè)備內(nèi)存副本、內(nèi)核啟動(dòng)和設(shè)備到主機(jī)內(nèi)存副本后，將主機(jī)函數(shù) MyCallback 添加到兩個(gè)流中的每一個(gè)。 每個(gè)設(shè)備到主機(jī)的內(nèi)存復(fù)制完成后，該函數(shù)將在主機(jī)上開始執(zhí)行。

void CUDART_CB MyCallback(cudaStream_t stream, cudaError_t status, void *data){
    printf("Inside callback %d\n", (size_t)data);
}
...
for (size_t i = 0; i < 2; ++i) {
    cudaMemcpyAsync(devPtrIn[i], hostPtr[i], size, cudaMemcpyHostToDevice, stream[i]);
    MyKernel<<<100, 512, 0, stream[i]>>>(devPtrOut[i], devPtrIn[i], size);
    cudaMemcpyAsync(hostPtr[i], devPtrOut[i], size, cudaMemcpyDeviceToHost, stream[i]);
    cudaLaunchHostFunc(stream[i], MyCallback, (void*)i);
}

在主機(jī)函數(shù)之后在流中發(fā)出的命令不會(huì)在函數(shù)完成之前開始執(zhí)行。

在流中的主機(jī)函數(shù)不得進(jìn)行 CUDA API 調(diào)用（直接或間接），因?yàn)槿绻M(jìn)行這樣的調(diào)用導(dǎo)致死鎖，它可能最終會(huì)等待自身。

3.2.6.5.7 流優(yōu)先級(jí)

可以在創(chuàng)建時(shí)使用cudaStreamCreateWithPriority()指定流的相對優(yōu)先級(jí)。 可以使用cudaDeviceGetStreamPriorityRange()函數(shù)獲得允許的優(yōu)先級(jí)范圍，按 [最高優(yōu)先級(jí)，最低優(yōu)先級(jí)] 排序。 在運(yùn)行時(shí)，高優(yōu)先級(jí)流中的待處理工作優(yōu)先于低優(yōu)先級(jí)流中的待處理工作。

以下代碼示例獲取當(dāng)前設(shè)備允許的優(yōu)先級(jí)范圍，并創(chuàng)建具有最高和最低可用優(yōu)先級(jí)的流。

// get the range of stream priorities for this device
int priority_high, priority_low;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);
// create streams with highest and lowest available priorities
cudaStream_t st_high, st_low;
cudaStreamCreateWithPriority(&st_high, cudaStreamNonBlocking, priority_high);
cudaStreamCreateWithPriority(&st_low, cudaStreamNonBlocking, priority_low);

3.2.6.6 CUDA圖

CUDA Graphs 為 CUDA 中的工作提交提供了一種新模型。圖是一系列操作，例如內(nèi)核啟動(dòng)，由依賴關(guān)系連接，獨(dú)立于其執(zhí)行定義。這允許一個(gè)圖被定義一次，然后重復(fù)啟動(dòng)。將圖的定義與其執(zhí)行分開可以實(shí)現(xiàn)許多優(yōu)化：首先，與流相比，CPU 啟動(dòng)成本降低，因?yàn)榇蟛糠衷O(shè)置都是提前完成的；其次，將整個(gè)工作流程呈現(xiàn)給 CUDA 可以實(shí)現(xiàn)優(yōu)化，這可能無法通過流的分段工作提交機(jī)制實(shí)現(xiàn)。

要查看圖形可能的優(yōu)化，請考慮流中發(fā)生的情況：當(dāng)您將內(nèi)核放入流中時(shí)，主機(jī)驅(qū)動(dòng)程序會(huì)執(zhí)行一系列操作，以準(zhǔn)備在 GPU 上執(zhí)行內(nèi)核。這些設(shè)置和啟動(dòng)內(nèi)核所必需的操作是必須為發(fā)布的每個(gè)內(nèi)核支付的間接成本。對于執(zhí)行時(shí)間較短的 GPU 內(nèi)核，這種開銷成本可能是整個(gè)端到端執(zhí)行時(shí)間的很大一部分。

使用圖的工作提交分為三個(gè)不同的階段：定義、實(shí)例化和執(zhí)行。

• 在定義階段，程序創(chuàng)建圖中操作的描述以及它們之間的依賴關(guān)系。
• 實(shí)例化獲取圖模板的快照，對其進(jìn)行驗(yàn)證，并執(zhí)行大部分工作的設(shè)置和初始化，目的是最大限度地減少啟動(dòng)時(shí)需要完成的工作。 生成的實(shí)例稱為可執(zhí)行圖。
• 可執(zhí)行圖可以啟動(dòng)到流中，類似于任何其他 CUDA 工作。 它可以在不重復(fù)實(shí)例化的情況下啟動(dòng)任意次數(shù)。

3.2.6.6.1圖架構(gòu)

一個(gè)操作在圖中形成一個(gè)節(jié)點(diǎn)。 操作之間的依賴關(guān)系是邊。 這些依賴關(guān)系限制了操作的執(zhí)行順序。

一個(gè)操作可以在它所依賴的節(jié)點(diǎn)完成后隨時(shí)調(diào)度。 調(diào)度由 CUDA 系統(tǒng)決定。

3.2.6.6.1.1 節(jié)點(diǎn)類型

圖節(jié)點(diǎn)可以是以下之一：

• 核函數(shù)
• CPU函數(shù)調(diào)用
• 內(nèi)存拷貝
• 內(nèi)存設(shè)置
• 空節(jié)點(diǎn)
• 等待事件
• 記錄事件
• 發(fā)出外部信號(hào)量的信號(hào)
• 等待外部信號(hào)量
• 子圖：執(zhí)行單獨(dú)的嵌套圖。 請參下圖。

3.2.6.6.2利用API創(chuàng)建圖

可以通過兩種機(jī)制創(chuàng)建圖：顯式 API 和流捕獲。 以下是創(chuàng)建和執(zhí)行下圖的示例。

// Create the graph - it starts out empty
cudaGraphCreate(&graph, 0);

// For the purpose of this example, we'll create
// the nodes separately from the dependencies to
// demonstrate that it can be done in two stages.
// Note that dependencies can also be specified 
// at node creation. 
cudaGraphAddKernelNode(&a, graph, NULL, 0, &nodeParams);
cudaGraphAddKernelNode(&b, graph, NULL, 0, &nodeParams);
cudaGraphAddKernelNode(&c, graph, NULL, 0, &nodeParams);
cudaGraphAddKernelNode(&d, graph, NULL, 0, &nodeParams);

// Now set up dependencies on each node
cudaGraphAddDependencies(graph, &a, &b, 1);     // A->B
cudaGraphAddDependencies(graph, &a, &c, 1);     // A->C
cudaGraphAddDependencies(graph, &b, &d, 1);     // B->D
cudaGraphAddDependencies(graph, &c, &d, 1);     // C->D

3.2.6.6.3 使用流捕獲創(chuàng)建圖

流捕獲提供了一種從現(xiàn)有的基于流的 API 創(chuàng)建圖的機(jī)制。 將工作啟動(dòng)到流中的一段代碼，包括現(xiàn)有代碼，可以等同于用與cudaStreamBeginCapture()和cudaStreamEndCapture()的調(diào)用。

cudaGraph_t graph;

cudaStreamBeginCapture(stream);

kernel_A<<< ..., stream >>>(...);
kernel_B<<< ..., stream >>>(...);
libraryCall(stream);
kernel_C<<< ..., stream >>>(...);

cudaStreamEndCapture(stream, &graph);

對cudaStreamBeginCapture()的調(diào)用將流置于捕獲模式。 捕獲流時(shí)，啟動(dòng)到流中的工作不會(huì)排隊(duì)執(zhí)行。 相反，它被附加到正在逐步構(gòu)建的內(nèi)部圖中。 然后通過調(diào)用cudaStreamEndCapture()返回此圖，這也結(jié)束了流的捕獲模式。 由流捕獲主動(dòng)構(gòu)建的圖稱為捕獲圖(capture graph)。

流捕獲可用于除cudaStreamLegacy（“NULL 流”）之外的任何 CUDA 流。 請注意，它可以在cudaStreamPerThread上使用。 如果程序正在使用legacy stream，則可以將stream 0 重新定義為不更改功能的每線程流。 請參閱默認(rèn)流。

可以使用cudaStreamIsCapturing()查詢是否正在捕獲流。

3.2.6.6.3.1 跨流依賴性和事件

流捕獲可以處理用cudaEventRecord()和cudaStreamWaitEvent()表示的跨流依賴關(guān)系，前提是正在等待的事件被記錄到同一個(gè)捕獲圖中。

當(dāng)事件記錄在處于捕獲模式的流中時(shí)，它會(huì)導(dǎo)致捕獲事件。捕獲的事件表示捕獲圖中的一組節(jié)點(diǎn)。

當(dāng)流等待捕獲的事件時(shí)，如果尚未將流置于捕獲模式，則它會(huì)將流置于捕獲模式，并且流中的下一個(gè)項(xiàng)目將對捕獲事件中的節(jié)點(diǎn)具有額外的依賴關(guān)系。然后將兩個(gè)流捕獲到同一個(gè)捕獲圖。

當(dāng)流捕獲中存在跨流依賴時(shí)，仍然必須在調(diào)用cudaStreamBeginCapture()的同一流中調(diào)用cudaStreamEndCapture()；這是原始流。由于基于事件的依賴關(guān)系，被捕獲到同一捕獲圖的任何其他流也必須連接回原始流。如下所示。在cudaStreamEndCapture()時(shí)，捕獲到同一捕獲圖的所有流都將退出捕獲模式。未能重新加入原始流將導(dǎo)致整個(gè)捕獲操作失敗。

// stream1 is the origin stream
cudaStreamBeginCapture(stream1);

kernel_A<<< ..., stream1 >>>(...);

// Fork into stream2
cudaEventRecord(event1, stream1);
cudaStreamWaitEvent(stream2, event1);

kernel_B<<< ..., stream1 >>>(...);
kernel_C<<< ..., stream2 >>>(...);

// Join stream2 back to origin stream (stream1)
cudaEventRecord(event2, stream2);
cudaStreamWaitEvent(stream1, event2);

kernel_D<<< ..., stream1 >>>(...);

// End capture in the origin stream
cudaStreamEndCapture(stream1, &graph);

// stream1 and stream2 no longer in capture mode  

上述代碼返回的圖如圖 10所示。

注意：當(dāng)流退出捕獲模式時(shí)，流中的下一個(gè)未捕獲項(xiàng)（如果有）仍將依賴于最近的先前未捕獲項(xiàng)，盡管已刪除中間項(xiàng)。

3.2.6.6.3.2 禁止和未處理的操作

同步或查詢正在捕獲的流或捕獲的事件的執(zhí)行狀態(tài)是無效的，因?yàn)樗鼈儾淮碛?jì)劃執(zhí)行的項(xiàng)目。當(dāng)任何關(guān)聯(lián)流處于捕獲模式時(shí)，查詢包含活動(dòng)流捕獲的更廣泛句柄（例如設(shè)備或上下文句柄）的執(zhí)行狀態(tài)或同步也是無效的。

當(dāng)捕獲同一上下文中的任何流時(shí)，并且它不是使用cudaStreamNonBlocking創(chuàng)建的，任何使用舊流的嘗試都是無效的。這是因?yàn)閘egacy stream句柄始終包含這些其他流；legacy stream將創(chuàng)建對正在捕獲的流的依賴，并且查詢它或同步它會(huì)查詢或同步正在捕獲的流。

因此在這種情況下調(diào)用同步 API 也是無效的。同步 API，例如 cudaMemcpy()，將工作legacy stream并在返回之前對其進(jìn)行同步。

注意：作為一般規(guī)則，當(dāng)依賴關(guān)系將捕獲的內(nèi)容與未捕獲的內(nèi)容聯(lián)系起來并排隊(duì)執(zhí)行時(shí)，CUDA 更喜歡返回錯(cuò)誤而不是忽略依賴關(guān)系。將流放入或退出捕獲模式時(shí)會(huì)出現(xiàn)異常；這切斷了在模式轉(zhuǎn)換之前和之后添加到流中的項(xiàng)目之間的依賴關(guān)系。

通過等待來自正在捕獲并且與與事件不同的捕獲圖相關(guān)聯(lián)的流中的捕獲事件來合并兩個(gè)單獨(dú)的捕獲圖是無效的。等待正在捕獲的流中的未捕獲事件是無效的。

圖中當(dāng)前不支持將異步操作排入流的少量 API，如果使用正在捕獲的流調(diào)用，則會(huì)返回錯(cuò)誤，例如cudaStreamAttachMemAsync()。

3.2.6.6.3.3失效

在流捕獲期間嘗試無效操作時(shí)，任何關(guān)聯(lián)的捕獲圖都將失效。 當(dāng)捕獲圖無效時(shí)，進(jìn)一步使用正在捕獲的任何流或與該圖關(guān)聯(lián)的捕獲事件將無效并將返回錯(cuò)誤，直到使用cudaStreamEndCapture()結(jié)束流捕獲。 此調(diào)用將使關(guān)聯(lián)的流脫離捕獲模式，但也會(huì)返回錯(cuò)誤值和 NULL 圖。

3.2.6.6.4 更新實(shí)例化圖

使用圖的工作提交分為三個(gè)不同的階段：定義、實(shí)例化和執(zhí)行。在工作流不改變的情況下，定義和實(shí)例化的開銷可以分?jǐn)偟皆S多執(zhí)行中，并且圖提供了明顯優(yōu)于流的優(yōu)勢。

圖是工作流的快照，包括內(nèi)核、參數(shù)和依賴項(xiàng)，以便盡可能快速有效地重放它。在工作流發(fā)生變化的情況下，圖會(huì)過時(shí)，必須進(jìn)行修改。對圖結(jié)構(gòu)（例如拓?fù)浠蚬?jié)點(diǎn)類型）的重大更改將需要重新實(shí)例化源圖，因?yàn)楸仨氈匦聭?yīng)用各種與拓?fù)湎嚓P(guān)的優(yōu)化技術(shù)。

重復(fù)實(shí)例化的成本會(huì)降低圖執(zhí)行帶來的整體性能優(yōu)勢，但通常只有節(jié)點(diǎn)參數(shù)（例如內(nèi)核參數(shù)和 cudaMemcpy 地址）發(fā)生變化，而圖拓?fù)浔３植蛔?。對于這種情況，CUDA 提供了一種稱為“圖形更新”的輕量級(jí)機(jī)制，它允許就地修改某些節(jié)點(diǎn)參數(shù)，而無需重建整個(gè)圖形。這比重新實(shí)例化要有效得多。

更新將在下次啟動(dòng)圖時(shí)生效，因此它們不會(huì)影響以前的圖啟動(dòng)，即使它們在更新時(shí)正在運(yùn)行。一個(gè)圖可能會(huì)被重復(fù)更新和重新啟動(dòng)，因此多個(gè)更新/啟動(dòng)可以在一個(gè)流上排隊(duì)。

CUDA 提供了兩種更新實(shí)例化圖的機(jī)制，全圖更新和單個(gè)節(jié)點(diǎn)更新。整個(gè)圖更新允許用戶提供一個(gè)拓?fù)湎嗤?code style="font-size:inherit;color:inherit;margin:0px;padding:0px;border:0px;font-style:inherit;font-variant:inherit;font-weight:inherit;line-height:inherit;vertical-align:baseline;background-color:rgb(244,244,244);">cudaGraph_t對象，其節(jié)點(diǎn)包含更新的參數(shù)。單個(gè)節(jié)點(diǎn)更新允許用戶顯式更新單個(gè)節(jié)點(diǎn)的參數(shù)。當(dāng)大量節(jié)點(diǎn)被更新時(shí)，或者當(dāng)調(diào)用者不知道圖拓?fù)鋾r(shí)（即，圖是由庫調(diào)用的流捕獲產(chǎn)生的），使用更新的cudaGraph_t會(huì)更方便。當(dāng)更改的數(shù)量很少并且用戶擁有需要更新的節(jié)點(diǎn)的句柄時(shí)，首選使用單個(gè)節(jié)點(diǎn)更新。單個(gè)節(jié)點(diǎn)更新跳過未更改節(jié)點(diǎn)的拓?fù)錂z查和比較，因此在許多情況下它可以更有效。以下部分更詳細(xì)地解釋了每種方法。

3.2.6.6.4.1 圖更新限制

內(nèi)核節(jié)點(diǎn)：

• 函數(shù)的所屬上下文不能改變。
• 其功能最初未使用 CUDA 動(dòng)態(tài)并行性的節(jié)點(diǎn)無法更新為使用 CUDA 動(dòng)態(tài)并行性的功能。

cudaMemset 和 cudaMemcpy 節(jié)點(diǎn)：

• 操作數(shù)分配/映射到的 CUDA 設(shè)備不能更改。
• 源/目標(biāo)內(nèi)存必須從與原始源/目標(biāo)內(nèi)存相同的上下文中分配。
• 只能更改一維 cudaMemset/cudaMemcpy 節(jié)點(diǎn)。

額外的 memcpy 節(jié)點(diǎn)限制：

• 不支持更改源或目標(biāo)內(nèi)存類型（即 cudaPitchedPtr、cudaArray_t 等）或傳輸類型（即 cudaMemcpyKind）。

外部信號(hào)量等待節(jié)點(diǎn)和記錄節(jié)點(diǎn)：

• 不支持更改信號(hào)量的數(shù)量。
• 對主機(jī)節(jié)點(diǎn)、事件記錄節(jié)點(diǎn)或事件等待節(jié)點(diǎn)的更新沒有限制。

3.2.6.6.4.2全圖更新

cudaGraphExecUpdate()允許使用相同拓?fù)鋱D（“更新”圖）中的參數(shù)更新實(shí)例化圖（“原始圖”）。 更新圖的拓?fù)浔仨毰c用于實(shí)例化cudaGraphExec_t的原始圖相同。 此外，將節(jié)點(diǎn)添加到原始圖或從中刪除的順序必須與將節(jié)點(diǎn)添加到更新圖（或從中刪除）的順序相匹配。 因此，在使用流捕獲時(shí)，必須以相同的順序捕獲節(jié)點(diǎn)，而在使用顯式圖形節(jié)點(diǎn)創(chuàng)建 API 時(shí)，必須以相同的順序添加或刪除所有節(jié)點(diǎn)。

以下示例顯示了如何使用 API 更新實(shí)例化圖：

cudaGraphExec_t graphExec = NULL;

for (int i = 0; i < 10; i++) {
    cudaGraph_t graph;
    cudaGraphExecUpdateResult updateResult;
    cudaGraphNode_t errorNode;

    // In this example we use stream capture to create the graph.
    // You can also use the Graph API to produce a graph.
    cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);

    // Call a user-defined, stream based workload, for example
    do_cuda_work(stream);

    cudaStreamEndCapture(stream, &graph);

    // If we've already instantiated the graph, try to update it directly
    // and avoid the instantiation overhead
    if (graphExec != NULL) {
        // If the graph fails to update, errorNode will be set to the
        // node causing the failure and updateResult will be set to a
        // reason code.
        cudaGraphExecUpdate(graphExec, graph, &errorNode, &updateResult);
    }

    // Instantiate during the first iteration or whenever the update
    // fails for any reason
    if (graphExec == NULL || updateResult != cudaGraphExecUpdateSuccess) {

        // If a previous update failed, destroy the cudaGraphExec_t
        // before re-instantiating it
        if (graphExec != NULL) {
            cudaGraphExecDestroy(graphExec);
        }   
        // Instantiate graphExec from graph. The error node and
        // error message parameters are unused here.
        cudaGraphInstantiate(&graphExec, graph, NULL, NULL, 0);
    }   

    cudaGraphDestroy(graph);
    cudaGraphLaunch(graphExec, stream);
    cudaStreamSynchronize(stream);
}

典型的工作流程是使用流捕獲或圖 API 創(chuàng)建初始cudaGraph_t。 然后cudaGraph_t被實(shí)例化并正常啟動(dòng)。 初始啟動(dòng)后，使用與初始圖相同的方法創(chuàng)建新的cudaGraph_t，并調(diào)用cudaGraphExecUpdate()。 如果圖更新成功，由上面示例中的updateResult參數(shù)指示，則啟動(dòng)更新的cudaGraphExec_t。 如果由于任何原因更新失敗，則調(diào)用cudaGraphExecDestroy()和cudaGraphInstantiate()來銷毀原始的cudaGraphExec_t并實(shí)例化一個(gè)新的。

也可以直接更新cudaGraph_t節(jié)點(diǎn)（即，使用cudaGraphKernelNodeSetParams()）并隨后更新cudaGraphExec_t，但是使用下一節(jié)中介紹的顯式節(jié)點(diǎn)更新 API 會(huì)更有效。

有關(guān)使用情況和當(dāng)前限制的更多信息，請參閱Graph API。

3.2.6.6.4.3 單個(gè)節(jié)點(diǎn)更新

實(shí)例化的圖節(jié)點(diǎn)參數(shù)可以直接更新。 這消除了實(shí)例化的開銷以及創(chuàng)建新cudaGraph_t的開銷。 如果需要更新的節(jié)點(diǎn)數(shù)相對于圖中的總節(jié)點(diǎn)數(shù)較小，則最好單獨(dú)更新節(jié)點(diǎn)。 以下方法可用于更新cudaGraphExec_t節(jié)點(diǎn)：

• cudaGraphExecKernelNodeSetParams()
• cudaGraphExecMemcpyNodeSetParams()
• cudaGraphExecMemsetNodeSetParams()
• cudaGraphExecHostNodeSetParams()
• cudaGraphExecChildGraphNodeSetParams()
• cudaGraphExecEventRecordNodeSetEvent()
• cudaGraphExecEventWaitNodeSetEvent()
• cudaGraphExecExternalSemaphoresSignalNodeSetParams()
• cudaGraphExecExternalSemaphoresWaitNodeSetParams()

有關(guān)使用情況和當(dāng)前限制的更多信息，請參閱Graph API。

3.2.6.6.5 使用圖API

cudaGraph_t對象不是線程安全的。 用戶有責(zé)任確保多個(gè)線程不會(huì)同時(shí)訪問同一個(gè)cudaGraph_t。

cudaGraphExec_t不能與自身同時(shí)運(yùn)行。cudaGraphExec_t的啟動(dòng)將在之前啟動(dòng)相同的可執(zhí)行圖之后進(jìn)行。

圖形執(zhí)行在流中完成，以便與其他異步工作進(jìn)行排序。 但是，流僅用于排序； 它不限制圖的內(nèi)部并行性，也不影響圖節(jié)點(diǎn)的執(zhí)行位置。

請參閱圖API。

3.2.6.7 事件

運(yùn)行時(shí)還提供了一種密切監(jiān)視設(shè)備進(jìn)度以及執(zhí)行準(zhǔn)確計(jì)時(shí)的方法，方法是讓應(yīng)用程序異步記錄程序中任何點(diǎn)的事件，并查詢這些事件何時(shí)完成。 當(dāng)事件之前的所有任務(wù)（或給定流中的所有命令）都已完成時(shí)，事件已完成。 空流中的事件在所有流中的所有先前任務(wù)和命令都完成后完成。

3.2.6.7.1 創(chuàng)建和銷毀

以下代碼示例創(chuàng)建兩個(gè)事件：

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

它們以這種方式被銷毀：

cudaEventDestroy(start);
cudaEventDestroy(stop);

3.2.6.7.2 計(jì)算時(shí)間

可以用以下方式來計(jì)時(shí):

cudaEventRecord(start, 0);
for (int i = 0; i < 2; ++i) {
    cudaMemcpyAsync(inputDev + i * size, inputHost + i * size,
                    size, cudaMemcpyHostToDevice, stream[i]);
    MyKernel<<<100, 512, 0, stream[i]>>>
               (outputDev + i * size, inputDev + i * size, size);
    cudaMemcpyAsync(outputHost + i * size, outputDev + i * size,
                    size, cudaMemcpyDeviceToHost, stream[i]);
}
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);

3.2.6.8同步調(diào)用

調(diào)用同步函數(shù)時(shí)，在設(shè)備完成請求的任務(wù)之前，控制不會(huì)返回給主機(jī)線程。 在主機(jī)線程執(zhí)行任何其他 CUDA 調(diào)用之前，可以通過調(diào)用帶有一些特定標(biāo)志的cudaSetDeviceFlags()來指定主機(jī)線程是否會(huì)產(chǎn)生、阻塞或自旋（有關(guān)詳細(xì)信息，請參閱參考手冊）。

3.2.7 多設(shè)備系統(tǒng)

3.2.7.1設(shè)備枚舉

一個(gè)主機(jī)系統(tǒng)可以有多個(gè)設(shè)備。 以下代碼示例顯示了如何枚舉這些設(shè)備、查詢它們的屬性并確定啟用 CUDA 的設(shè)備的數(shù)量。

int deviceCount;
cudaGetDeviceCount(&deviceCount);
int device;
for (device = 0; device < deviceCount; ++device) {
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, device);
    printf("Device %d has compute capability %d.%d.\n",
           device, deviceProp.major, deviceProp.minor);
}

3.2.7.2 設(shè)備選擇

主機(jī)線程可以通過調(diào)用cudaSetDevice()隨時(shí)設(shè)置它所操作的設(shè)備。 設(shè)備內(nèi)存分配和內(nèi)核啟動(dòng)在當(dāng)前設(shè)置的設(shè)備上進(jìn)行； 流和事件是與當(dāng)前設(shè)置的設(shè)備相關(guān)聯(lián)的。 如果未調(diào)用cudaSetDevice()，則當(dāng)前設(shè)備為設(shè)備0。

以下代碼示例說明了設(shè)置當(dāng)前設(shè)備如何影響內(nèi)存分配和內(nèi)核執(zhí)行。

size_t size = 1024 * sizeof(float);
cudaSetDevice(0);            // Set device 0 as current
float* p0;
cudaMalloc(&p0, size);       // Allocate memory on device 0
MyKernel<<<1000, 128>>>(p0); // Launch kernel on device 0
cudaSetDevice(1);            // Set device 1 as current
float* p1;
cudaMalloc(&p1, size);       // Allocate memory on device 1
MyKernel<<<1000, 128>>>(p1); // Launch kernel on device 1

3.2.7.3 流和事件行為

如果在與當(dāng)前設(shè)備無關(guān)的流上啟動(dòng)內(nèi)核將失敗，如以下代碼示例所示。

cudaSetDevice(0);               // Set device 0 as current
cudaStream_t s0;
cudaStreamCreate(&s0);          // Create stream s0 on device 0
MyKernel<<<100, 64, 0, s0>>>(); // Launch kernel on device 0 in s0
cudaSetDevice(1);               // Set device 1 as current
cudaStream_t s1;
cudaStreamCreate(&s1);          // Create stream s1 on device 1
MyKernel<<<100, 64, 0, s1>>>(); // Launch kernel on device 1 in s1

// This kernel launch will fail:
MyKernel<<<100, 64, 0, s0>>>(); // Launch kernel on device 1 in s0

即使將內(nèi)存復(fù)制運(yùn)行在與當(dāng)前設(shè)備無關(guān)的流，它也會(huì)成功。

如果輸入事件和輸入流關(guān)聯(lián)到不同的設(shè)備，cudaEventRecord()將失敗。

如果兩個(gè)輸入事件關(guān)聯(lián)到不同的設(shè)備，cudaEventElapsedTime()將失敗。

即使輸入事件關(guān)聯(lián)到與當(dāng)前設(shè)備不同的設(shè)備，cudaEventSynchronize()和cudaEventQuery()也會(huì)成功。

即使輸入流和輸入事件關(guān)聯(lián)到不同的設(shè)備，cudaStreamWaitEvent()也會(huì)成功。 因此，cudaStreamWaitEvent()可用于使多個(gè)設(shè)備相互同步。

每個(gè)設(shè)備都有自己的默認(rèn)流（請參閱默認(rèn)流），因此向設(shè)備的默認(rèn)流發(fā)出的命令可能會(huì)亂序執(zhí)行或與向任何其他設(shè)備的默認(rèn)流發(fā)出的命令同時(shí)執(zhí)行。

3.2.7.4 Peer-to-Peer的內(nèi)存訪問

根據(jù)系統(tǒng)屬性，特別是 PCIe 或 NVLINK 拓?fù)浣Y(jié)構(gòu)，設(shè)備能夠相互尋址對方的內(nèi)存（即，在一個(gè)設(shè)備上執(zhí)行的內(nèi)核可以取消引用指向另一設(shè)備內(nèi)存的指針）。 如果cudaDeviceCanAccessPeer()為這兩個(gè)設(shè)備返回 true，則在兩個(gè)設(shè)備之間支持這種對等內(nèi)存訪問功能。

對等內(nèi)存訪問僅在 64 位應(yīng)用程序中受支持，并且必須通過調(diào)用cudaDeviceEnablePeerAccess()在兩個(gè)設(shè)備之間啟用，如以下代碼示例所示。 在未啟用NVSwitch的系統(tǒng)上，每個(gè)設(shè)備最多可支持系統(tǒng)范圍內(nèi)的八個(gè)對等連接。

兩個(gè)設(shè)備使用統(tǒng)一的地址空間（請參閱統(tǒng)一虛擬地址空間），因此可以使用相同的指針來尋址兩個(gè)設(shè)備的內(nèi)存，如下面的代碼示例所示。

cudaSetDevice(0);                   // Set device 0 as current
float* p0;
size_t size = 1024 * sizeof(float);
cudaMalloc(&p0, size);              // Allocate memory on device 0
MyKernel<<<1000, 128>>>(p0);        // Launch kernel on device 0
cudaSetDevice(1);                   // Set device 1 as current
cudaDeviceEnablePeerAccess(0, 0);   // Enable peer-to-peer access
                                    // with device 0

// Launch kernel on device 1
// This kernel launch can access memory on device 0 at address p0
MyKernel<<<1000, 128>>>(p0);

3.2.7.4.1 Linux上的IOMMU

僅在 Linux 上，CUDA 和顯示驅(qū)動(dòng)程序不支持啟用 IOMMU 的裸機(jī) PCIe 對等內(nèi)存復(fù)制。 但是，CUDA 和顯示驅(qū)動(dòng)程序確實(shí)支持通過 VM 傳遞的 IOMMU。 因此，Linux 上的用戶在本機(jī)裸機(jī)系統(tǒng)上運(yùn)行時(shí)，應(yīng)禁用 IOMMU。 如啟用 IOMMU，將 VFIO 驅(qū)動(dòng)程序用作虛擬機(jī)的 PCIe 通道。

在 Windows 上，上述限制不存在。

另請參閱在 64 位平臺(tái)上分配 DMA 緩沖區(qū)。

3.2.7.5 Peer-to-Peer內(nèi)存拷貝

可以在兩個(gè)不同設(shè)備的內(nèi)存之間執(zhí)行內(nèi)存復(fù)制。

當(dāng)兩個(gè)設(shè)備使用統(tǒng)一地址空間時(shí)（請參閱統(tǒng)一虛擬地址空間），這是使用設(shè)備內(nèi)存中提到的常規(guī)內(nèi)存復(fù)制功能完成的。

否則，這將使用cudaMemcpyPeer()、cudaMemcpyPeerAsync()、cudaMemcpy3DPeer() 或cudaMemcpy3DPeerAsync()完成，如以下代碼示例所示。

cudaSetDevice(0);                   // Set device 0 as current
float* p0;
size_t size = 1024 * sizeof(float);
cudaMalloc(&p0, size);              // Allocate memory on device 0
cudaSetDevice(1);                   // Set device 1 as current
float* p1;
cudaMalloc(&p1, size);              // Allocate memory on device 1
cudaSetDevice(0);                   // Set device 0 as current
MyKernel<<<1000, 128>>>(p0);        // Launch kernel on device 0
cudaSetDevice(1);                   // Set device 1 as current
cudaMemcpyPeer(p1, 1, p0, 0, size); // Copy p0 to p1
MyKernel<<<1000, 128>>>(p1);   

兩個(gè)不同設(shè)備的內(nèi)存之間的拷貝（在隱式 NULL 流中）：

• 直到之前向任一設(shè)備發(fā)出的所有命令都完成后才會(huì)啟動(dòng)，并且
• 在復(fù)制到任一設(shè)備之后發(fā)出的任何命令（請參閱異步并發(fā)執(zhí)行）可以開始之前運(yùn)行完成。

與流的正常行為一致，兩個(gè)設(shè)備的內(nèi)存之間的異步拷貝可能與另一個(gè)流中的拷貝或內(nèi)核重疊。

請注意，如果通過cudaDeviceEnablePeerAccess()在兩個(gè)設(shè)備之間啟用Peer-to-Peer訪問，如Peer-to-Peer內(nèi)存訪問中所述，這兩個(gè)設(shè)備之間的Peer-to-Peer內(nèi)存復(fù)制不再需要通過主機(jī), 因此速度更快。

統(tǒng)一虛擬地址空間

當(dāng)應(yīng)用程序作為 64 位進(jìn)程運(yùn)行時(shí)，單個(gè)地址空間用于主機(jī)和計(jì)算能力 2.0 及更高版本的所有設(shè)備。通過 CUDA API 調(diào)用進(jìn)行的所有主機(jī)內(nèi)存分配以及受支持設(shè)備上的所有設(shè)備內(nèi)存分配都在此虛擬地址范圍內(nèi)。作為結(jié)果：

• 通過 CUDA 分配的主機(jī)或使用統(tǒng)一地址空間的任何設(shè)備上的任何內(nèi)存的位置都可以使用cudaPointerGetAttributes()從指針的值中確定。
• 當(dāng)復(fù)制到或從任何使用統(tǒng)一地址空間的設(shè)備的內(nèi)存中復(fù)制時(shí)，可以將cudaMemcpy*()的cudaMemcpyKind參數(shù)設(shè)置為cudaMemcpyDefault以根據(jù)指針確定位置。只要當(dāng)前設(shè)備使用統(tǒng)一尋址，這也適用于未通過 CUDA 分配的主機(jī)指針。
• 通過cudaHostAlloc()進(jìn)行的分配可以在使用統(tǒng)一地址空間的所有設(shè)備之間自動(dòng)移植（請參閱可移植內(nèi)存），并且 cudaHostAlloc() 返回的指針可以直接在這些設(shè)備上運(yùn)行的內(nèi)核中使用（即，沒有需要通過cudaHostGetDevicePointer()獲取設(shè)備指針，如映射內(nèi)存中所述。

應(yīng)用程序可以通過檢查UnifiedAddressing設(shè)備屬性（請參閱設(shè)備枚舉）是否等于 1 來查詢統(tǒng)一地址空間是否用于特定設(shè)備。

3.2.9 進(jìn)程間通信

由主機(jī)線程創(chuàng)建的任何設(shè)備內(nèi)存指針或事件句柄都可以被同一進(jìn)程中的任何其他線程直接引用。然而，它在這個(gè)進(jìn)程之外是無效的，因此不能被屬于不同進(jìn)程的線程直接引用。

要跨進(jìn)程共享設(shè)備內(nèi)存指針和事件，應(yīng)用程序必須使用進(jìn)程間通信 API，參考手冊中有詳細(xì)描述。 IPC API 僅支持 Linux 上的 64 位進(jìn)程以及計(jì)算能力 2.0 及更高版本的設(shè)備。請注意，cudaMallocManaged 分配不支持 IPC API。

使用此 API，應(yīng)用程序可以使用cudaIpcGetMemHandle()獲取給定設(shè)備內(nèi)存指針的 IPC 句柄，使用標(biāo)準(zhǔn) IPC 機(jī)制（例如，進(jìn)程間共享內(nèi)存或文件）將其傳遞給另一個(gè)進(jìn)程，并使用cudaIpcOpenMemHandle()檢索設(shè)備來自 IPC 句柄的指針，該指針是其他進(jìn)程中的有效指針。可以使用類似的入口點(diǎn)共享事件句柄。

請注意，出于性能原因，由cudaMalloc()進(jìn)行的分配可能會(huì)從更大的內(nèi)存塊中進(jìn)行子分配。在這種情況下，CUDA IPC API 將共享整個(gè)底層內(nèi)存塊，這可能導(dǎo)致其他子分配被共享，這可能導(dǎo)致進(jìn)程之間的信息泄露。為了防止這種行為，建議僅共享具有 2MiB 對齊大小的分配。

使用 IPC API 的一個(gè)示例是單個(gè)主進(jìn)程生成一批輸入數(shù)據(jù)，使數(shù)據(jù)可用于多個(gè)輔助進(jìn)程，而無需重新生成或復(fù)制。

使用 CUDA IPC 相互通信的應(yīng)用程序應(yīng)使用相同的 CUDA 驅(qū)動(dòng)程序和運(yùn)行時(shí)進(jìn)行編譯、鏈接和運(yùn)行。

注意：自 CUDA 11.5 起，L4T 和具有計(jì)算能力 7.x 及更高版本的嵌入式 Linux Tegra 設(shè)備僅支持事件共享 IPC API。 Tegra 平臺(tái)仍然不支持內(nèi)存共享 IPC API。

3.2.10 錯(cuò)誤檢查

所有運(yùn)行時(shí)函數(shù)都返回錯(cuò)誤代碼，但對于異步函數(shù)（請參閱異步并發(fā)執(zhí)行），此錯(cuò)誤代碼不可能報(bào)告任何可能發(fā)生在設(shè)備上的異步錯(cuò)誤，因?yàn)楹瘮?shù)在設(shè)備完成任務(wù)之前返回；錯(cuò)誤代碼僅報(bào)告執(zhí)行任務(wù)之前主機(jī)上發(fā)生的錯(cuò)誤，通常與參數(shù)驗(yàn)證有關(guān)；如果發(fā)生異步錯(cuò)誤，會(huì)被后續(xù)一些不相關(guān)的運(yùn)行時(shí)函數(shù)調(diào)用報(bào)告。

因此，在某些異步函數(shù)調(diào)用之后檢查異步錯(cuò)誤的唯一方法是在調(diào)用之后通過調(diào)用cudaDeviceSynchronize()（或使用異步并發(fā)執(zhí)行中描述的任何其他同步機(jī)制）并檢查cudaDeviceSynchronize()。

運(yùn)行時(shí)為每個(gè)初始化為cudaSuccess的主機(jī)線程維護(hù)一個(gè)錯(cuò)誤變量，并在每次發(fā)生錯(cuò)誤時(shí)被錯(cuò)誤代碼覆蓋（無論是參數(shù)驗(yàn)證錯(cuò)誤還是異步錯(cuò)誤）。cudaPeekAtLastError()返回此變量。cudaGetLastError()返回此變量并將其重置為cudaSuccess。

內(nèi)核啟動(dòng)不返回任何錯(cuò)誤代碼，因此必須在內(nèi)核啟動(dòng)后立即調(diào)用cudaPeekAtLastError()或cudaGetLastError()以檢索任何啟動(dòng)前錯(cuò)誤。為了確保cudaPeekAtLastError()或cudaGetLastError()返回的任何錯(cuò)誤不是源自內(nèi)核啟動(dòng)之前的調(diào)用，必須確保在內(nèi)核啟動(dòng)之前將運(yùn)行時(shí)錯(cuò)誤變量設(shè)置為cudaSuccess，例如，通過調(diào)用cudaGetLastError()在內(nèi)核啟動(dòng)之前。內(nèi)核啟動(dòng)是異步的，因此要檢查異步錯(cuò)誤，應(yīng)用程序必須在內(nèi)核啟動(dòng)和調(diào)用cudaPeekAtLastError()或cudaGetLastError()之間進(jìn)行同步。

請注意，cudaStreamQuery()和cudaEventQuery()可能返回的cudaErrorNotReady不被視為錯(cuò)誤，因此cudaPeekAtLastError()或cudaGetLastError()不會(huì)報(bào)告。

3.2.11 調(diào)用棧

在計(jì)算能力 2.x 及更高版本的設(shè)備上，調(diào)用堆棧的大小可以使用cudaDeviceGetLimit()查詢并使用cudaDeviceSetLimit()設(shè)置。

當(dāng)調(diào)用堆棧溢出時(shí)，如果應(yīng)用程序通過 CUDA 調(diào)試器（cuda-gdb、Nsight）運(yùn)行，內(nèi)核調(diào)用將失敗并出現(xiàn)堆棧溢出錯(cuò)誤，否則會(huì)出現(xiàn)未指定的啟動(dòng)錯(cuò)誤。

3.2.12 紋理內(nèi)存和表面內(nèi)存(surface memory)

CUDA 支持 GPU 用于圖形訪問紋理和表面內(nèi)存的紋理硬件子集。 如設(shè)備內(nèi)存訪問中所述，從紋理或表面內(nèi)存而不是全局內(nèi)存讀取數(shù)據(jù)可以帶來多項(xiàng)性能優(yōu)勢。

有兩種不同的 API 可以訪問紋理和表面內(nèi)存：

• 所有設(shè)備都支持的紋理引用 API，
• 僅在計(jì)算能力 3.x 及更高版本的設(shè)備上支持的紋理對象 API。 紋理引用 API 具有紋理對象 API 沒有的限制。 它們在[[DEPRECATED]] 紋理引用 API?中被提及。

3.2.12.1紋理內(nèi)存

使用紋理函數(shù)中描述的設(shè)備函數(shù)從內(nèi)核讀取紋理內(nèi)存。 調(diào)用這些函數(shù)之一讀取紋理的過程稱為紋理提取。 每個(gè)紋理提取指定一個(gè)參數(shù)，稱為紋理對象 API 的紋理對象或紋理引用 API 的紋理引用。

紋理對象或紋理引用指定：

• 紋理，即提取的紋理內(nèi)存。 紋理對象在運(yùn)行時(shí)創(chuàng)建，并在創(chuàng)建紋理對象時(shí)指定紋理，如紋理對象 API?中所述。 紋理引用是在編譯時(shí)創(chuàng)建的，紋理是在運(yùn)行時(shí)通過[[DEPRECATED]] Texture Reference API?中描述的運(yùn)行時(shí)函數(shù)將紋理引用綁定到紋理來指定的； 幾個(gè)不同的紋理引用可能綁定到相同的紋理或內(nèi)存中重疊的紋理。 紋理可以是線性內(nèi)存的任何區(qū)域或 CUDA 數(shù)組（在CUDA 數(shù)組中描述）。
• 它的維數(shù)指定紋理是使用一個(gè)紋理坐標(biāo)的一維數(shù)組、使用兩個(gè)紋理坐標(biāo)的二維數(shù)組還是使用三個(gè)紋理坐標(biāo)的三維數(shù)組。數(shù)組的元素稱為texels，是紋理元素的縮寫。紋理的寬度、高度和深度是指數(shù)組在每個(gè)維度上的大小。表 15列出了取決于設(shè)備計(jì)算能力的最大紋理寬度、高度和深度。
• texels的類型，僅限于基本整數(shù)和單精度浮點(diǎn)類型以及從基本向量類型派生的內(nèi)置向量類型中定義的任何 1、2 和 4 分量向量類型整數(shù)和單精度浮點(diǎn)類型。
• 讀取模式，等同于cudaReadModeNormalizedFloat或cudaReadModeElementType。如果是cudaReadModeNormalizedFloat并且 texel 的類型是 16 位或 8 位整數(shù)類型，則紋理獲取返回的值實(shí)際上是作為浮點(diǎn)類型返回的，并且整數(shù)類型的全范圍映射到 [0.0 , 1.0] 表示無符號(hào)整數(shù)類型，[-1.0, 1.0] 表示有符號(hào)整數(shù)類型；例如，值為 0xff 的無符號(hào) 8 位紋理元素讀取為 1。如果是cudaReadModeElementType，則不執(zhí)行轉(zhuǎn)換。
• 紋理坐標(biāo)是否標(biāo)準(zhǔn)化。默認(rèn)情況下，使用 [0, N-1] 范圍內(nèi)的浮點(diǎn)坐標(biāo)（通過 Texture Functions 的函數(shù)）引用紋理，其中 N 是與坐標(biāo)對應(yīng)的維度中紋理的大小。例如，大小為 64×32 的紋理將分別使用 x 和 y 維度的 [0, 63] 和 [0, 31] 范圍內(nèi)的坐標(biāo)進(jìn)行引用。標(biāo)準(zhǔn)化紋理坐標(biāo)導(dǎo)致坐標(biāo)被指定在[0.0,1.0-1/N]范圍內(nèi)，而不是[0,N-1]，所以相同的64×32紋理將在x和y維度的[0,1 -1/N]范圍內(nèi)被標(biāo)準(zhǔn)化坐標(biāo)定位。如果紋理坐標(biāo)獨(dú)立于紋理大小，則歸一化紋理坐標(biāo)自然適合某些應(yīng)用程序的要求。
• 尋址方式。使用超出范圍的坐標(biāo)調(diào)用 B.8 節(jié)的設(shè)備函數(shù)是有效的。尋址模式定義了在這種情況下會(huì)發(fā)生什么。默認(rèn)尋址模式是將坐標(biāo)限制在有效范圍內(nèi)：[0, N) 用于非歸一化坐標(biāo)，[0.0, 1.0) 用于歸一化坐標(biāo)。如果指定了邊框模式，則紋理坐標(biāo)超出范圍的紋理提取將返回零。對于歸一化坐標(biāo)，還可以使用環(huán)繞模式和鏡像模式。使用環(huán)繞模式時(shí)，每個(gè)坐標(biāo) x 都轉(zhuǎn)換為 frac(x)=x – floor(x)，其中 floor(x) 是不大于 x 的最大整數(shù)。使用鏡像模式時(shí)，如果 floor(x) 為偶數(shù)，則每個(gè)坐標(biāo) x 轉(zhuǎn)換為 frac(x)，如果 floor(x) 為奇數(shù)，則轉(zhuǎn)換為 1-frac(x)。尋址模式被指定為一個(gè)大小為 3 的數(shù)組，其第一個(gè)、第二個(gè)和第三個(gè)元素分別指定第一個(gè)、第二個(gè)和第三個(gè)紋理坐標(biāo)的尋址模式；尋址模式為cudaAddressModeBorder、cudaAddressModeClamp、cudaAddressModeWrap和cudaAddressModeMirror；cudaAddressModeWrap和cudaAddressModeMirror僅支持標(biāo)準(zhǔn)化紋理坐標(biāo)
• 過濾模式指定如何根據(jù)輸入紋理坐標(biāo)計(jì)算獲取紋理時(shí)返回的值。線性紋理過濾只能對配置為返回浮點(diǎn)數(shù)據(jù)的紋理進(jìn)行。它在相鄰紋素之間執(zhí)行低精度插值。啟用后，將讀取紋理提取位置周圍的texels，并根據(jù)紋理坐標(biāo)落在texels之間的位置對紋理提取的返回值進(jìn)行插值。對一維紋理進(jìn)行簡單線性插值，對二維紋理進(jìn)行雙線性插值，對三維紋理進(jìn)行三線性插值。Texture Fetching提供了有關(guān)紋理獲取的更多細(xì)節(jié)。過濾模式等于cudaFilterModePoint或cudaFilterModeLinear。如果是cudaFilterModePoint，則返回值是紋理坐標(biāo)最接近輸入紋理坐標(biāo)的texel。如果是cudaFilterModeLinear，則返回值是紋理坐標(biāo)最接近的兩個(gè)（一維紋理）、四個(gè)（二維紋理）或八個(gè)（三維紋理）texel的線性插值輸入紋理坐標(biāo)。cudaFilterModeLinear僅對浮點(diǎn)類型的返回值有效。

紋理對象 API?。

[[DEPRECATED]] Texture Reference API

16位浮點(diǎn)紋理解釋了如何處理16位浮點(diǎn)紋理。

紋理也可以分層，如分層紋理中所述。

立方體貼圖紋理和立方體貼圖分層紋理描述了一種特殊類型的紋理，立方體貼圖紋理。

Texture Gather?描述了一種特殊的紋理獲取，紋理收集。

3.2.12.1.1 紋理對象API

使用cudaCreateTextureObject()從指定紋理的struct cudaResourceDesc類型的資源描述和定義如下的紋理描述創(chuàng)建紋理對象：

struct cudaTextureDesc
{
    enum cudaTextureAddressMode addressMode[3];
    enum cudaTextureFilterMode  filterMode;
    enum cudaTextureReadMode    readMode;
    int                         sRGB;
    int                         normalizedCoords;
    unsigned int                maxAnisotropy;
    enum cudaTextureFilterMode  mipmapFilterMode;
    float                       mipmapLevelBias;
    float                       minMipmapLevelClamp;
    float                       maxMipmapLevelClamp;
};

• addressMode指定尋址模式；
• filterMode指定過濾模式；
• readMode指定讀取模式；
• normalizedCoords指定紋理坐標(biāo)是否被歸一化；
• sRGB、maxAnisotropy、mipmapFilterMode、mipmapLevelBias、minMipmapLevelClamp和maxMipmapLevelClamp請參閱的參考手冊。

以下代碼示例將一些簡單的轉(zhuǎn)換內(nèi)核應(yīng)用于紋理。

// Simple transformation kernel
__global__ void transformKernel(float* output,
                                cudaTextureObject_t texObj,
                                int width, int height,
                                float theta) 
{
    // Calculate normalized texture coordinates
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    float u = x / (float)width;
    float v = y / (float)height;

    // Transform coordinates
    u -= 0.5f;
    v -= 0.5f;
    float tu = u * cosf(theta) - v * sinf(theta) + 0.5f;
    float tv = v * cosf(theta) + u * sinf(theta) + 0.5f;

    // Read from texture and write to global memory
    output[y * width + x] = tex2D(texObj, tu, tv);
}
// Host code
int main()
{
    const int height = 1024;
    const int width = 1024;
    float angle = 0.5;

    // Allocate and set some host data
    float *h_data = (float *)std::malloc(sizeof(float) * width * height);
    for (int i = 0; i < height * width; ++i)
        h_data[i] = i;

    // Allocate CUDA array in device memory
    cudaChannelFormatDesc channelDesc =
        cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
    cudaArray_t cuArray;
    cudaMallocArray(&cuArray, &channelDesc, width, height);

    // Set pitch of the source (the width in memory in bytes of the 2D array pointed
    // to by src, including padding), we dont have any padding
    const size_t spitch = width * sizeof(float);
    // Copy data located at address h_data in host memory to device memory
    cudaMemcpy2DToArray(cuArray, 0, 0, h_data, spitch, width * sizeof(float),
                        height, cudaMemcpyHostToDevice);

    // Specify texture
    struct cudaResourceDesc resDesc;
    memset(&resDesc, 0, sizeof(resDesc));
    resDesc.resType = cudaResourceTypeArray;
    resDesc.res.array.array = cuArray;

    // Specify texture object parameters
    struct cudaTextureDesc texDesc;
    memset(&texDesc, 0, sizeof(texDesc));
    texDesc.addressMode[0] = cudaAddressModeWrap;
    texDesc.addressMode[1] = cudaAddressModeWrap;
    texDesc.filterMode = cudaFilterModeLinear;
    texDesc.readMode = cudaReadModeElementType;
    texDesc.normalizedCoords = 1;

    // Create texture object
    cudaTextureObject_t texObj = 0;
    cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL);

    // Allocate result of transformation in device memory
    float *output;
    cudaMalloc(&output, width * height * sizeof(float));

    // Invoke kernel
    dim3 threadsperBlock(16, 16);
    dim3 numBlocks((width + threadsperBlock.x - 1) / threadsperBlock.x,
                    (height + threadsperBlock.y - 1) / threadsperBlock.y);
    transformKernel<<>>(output, texObj, width, height,
                                                    angle);
    // Copy data from device back to host
    cudaMemcpy(h_data, output, width * height * sizeof(float),
                cudaMemcpyDeviceToHost);

    // Destroy texture object
    cudaDestroyTextureObject(texObj);

    // Free device memory
    cudaFreeArray(cuArray);
    cudaFree(output);

    // Free host memory
    free(h_data);

    return 0;
}

3.2.12.1.2 [[已棄用]] 紋理引用 API

紋理參考 API 已棄用。

紋理引用的某些屬性是不可變的，必須在編譯時(shí)知道； 它們是在聲明紋理引用時(shí)指定的。 紋理引用在文件范圍內(nèi)聲明為紋理類型的變量：

texture texRef;

• DataType指定紋素的類型；
• Type指定紋理參考的類型，等于cudaTextureType1D、cudaTextureType2D或cudaTextureType3D，分別用于一維、二維或三維紋理，或cudaTextureType1DLayered或cudaTextureType2DLayered用于一維或二維 分別分層紋理；Type是一個(gè)可選參數(shù)，默認(rèn)為cudaTextureType1D；
• ReadMode指定讀取模式； 它是一個(gè)可選參數(shù)，默認(rèn)為cudaReadModeElementType。

紋理引用只能聲明為靜態(tài)全局變量，不能作為參數(shù)傳遞給函數(shù)。

紋理引用的其他屬性是可變的，并且可以在運(yùn)行時(shí)通過主機(jī)運(yùn)行時(shí)進(jìn)行更改。 如參考手冊中所述，運(yùn)行時(shí) API 具有低級(jí) C 樣式接口和高級(jí) C++ 樣式接口。 紋理類型在高級(jí) API 中定義為公開派生自低級(jí) API 中定義的textureReference類型的結(jié)構(gòu)，如下所示：

struct textureReference {
    int                          normalized;
    enum cudaTextureFilterMode   filterMode;
    enum cudaTextureAddressMode  addressMode[3];
    struct cudaChannelFormatDesc channelDesc;
    int                          sRGB;
    unsigned int                 maxAnisotropy;
    enum cudaTextureFilterMode   mipmapFilterMode;
    float                        mipmapLevelBias;
    float                        minMipmapLevelClamp;
    float                        maxMipmapLevelClamp;
}

• normalized指定紋理坐標(biāo)是否被歸一化；
• filterMode指定過濾模式；
• addressMode指定尋址模式；
• channelDesc描述了texel的格式； 它必須匹配紋理引用聲明的DataType參數(shù)；channelDesc屬于以下類型：

struct cudaChannelFormatDesc {
  int x, y, z, w;
  enum cudaChannelFormatKind f;
};
其中 x、y、z 和 w 等于返回值的每個(gè)分量的位數(shù)，f 為：

*cudaChannelFormatKindSigned 如果這些組件是有符號(hào)整數(shù)類型，
*cudaChannelFormatKindUnsigned 如果它們是無符號(hào)整數(shù)類型，
*cudaChannelFormatKindFloat 如果它們是浮點(diǎn)類型。

• sRGB、maxAnisotropy、mipmapFilterMode、mipmapLevelBias、minMipmapLevelClamp 和 maxMipmapLevelClamp請參閱參考手冊

normalized、addressMode和filterMode可以直接在主機(jī)代碼中修改。

在紋理內(nèi)存中讀取之前內(nèi)核可以使用紋理引用，紋理引用必須綁定到紋理，使用cudaBindTexture()或cudaBindTexture2D()用于線性內(nèi)存，或cudaBindTextureToArray()用于 CUDA 數(shù)組。cudaUnbindTexture()用于取消綁定紋理引用。 一旦紋理引用被解除綁定，它可以安全地重新綁定到另一個(gè)數(shù)組，即使使用之前綁定的紋理的內(nèi)核還沒有完成。 建議使用cudaMallocPitch()在線性內(nèi)存中分配二維紋理，并使用cudaMallocPitch()返回的間距作為cudaBindTexture2D()的輸入?yún)?shù)。

以下代碼示例將 2D 紋理引用綁定到 devPtr 指向的線性內(nèi)存：

• 使用低層次API:

texture texRef;
textureReference* texRefPtr;
cudaGetTextureReference(&texRefPtr, &texRef);
cudaChannelFormatDesc channelDesc =
                             cudaCreateChannelDesc();
size_t offset;
cudaBindTexture2D(&offset, texRefPtr, devPtr, &channelDesc,
                  width, height, pitch);

• 使用高層次API:

texture texRef;
cudaChannelFormatDesc channelDesc =
                             cudaCreateChannelDesc();
size_t offset;
cudaBindTexture2D(&offset, texRef, devPtr, channelDesc,
                  width, height, pitch);

以下代碼示例將 2D 紋理引用綁定到 CUDA 數(shù)組cuArray：

• 使用低層次API:

texture texRef;
textureReference* texRefPtr;
cudaGetTextureReference(&texRefPtr, &texRef);
cudaChannelFormatDesc channelDesc;
cudaGetChannelDesc(&channelDesc, cuArray);
cudaBindTextureToArray(texRef, cuArray, &channelDesc);

• 使用高層次API:

texture texRef;
cudaBindTextureToArray(texRef, cuArray);

將紋理綁定到紋理引用時(shí)指定的格式必須與聲明紋理引用時(shí)指定的參數(shù)匹配； 否則，紋理提取的結(jié)果是未定義的。

如表 15中指定的，可以綁定到內(nèi)核的紋理數(shù)量是有限的。

以下代碼示例將一些簡單的轉(zhuǎn)換內(nèi)核應(yīng)用于紋理。

// 2D float texture
texture texRef;

// Simple transformation kernel
__global__ void transformKernel(float* output,
                                int width, int height,
                                float theta) 
{
    // Calculate normalized texture coordinates
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    float u = x / (float)width;
    float v = y / (float)height;

    // Transform coordinates
    u -= 0.5f;
    v -= 0.5f;
    float tu = u * cosf(theta) - v * sinf(theta) + 0.5f;
    float tv = v * cosf(theta) + u * sinf(theta) + 0.5f;


    // Read from texture and write to global memory
    output[y * width + x] = tex2D(texRef, tu, tv);
}

// Host code
int main()
{
    // Allocate CUDA array in device memory
    cudaChannelFormatDesc channelDesc =
               cudaCreateChannelDesc(32, 0, 0, 0,
                                     cudaChannelFormatKindFloat);
    cudaArray* cuArray;
    cudaMallocArray(&cuArray, &channelDesc, width, height);

    // Copy to device memory some data located at address h_data
    // in host memory 
    cudaMemcpyToArray(cuArray, 0, 0, h_data, size,
                      cudaMemcpyHostToDevice);

    // Set texture reference parameters
    texRef.addressMode[0] = cudaAddressModeWrap;
    texRef.addressMode[1] = cudaAddressModeWrap;
    texRef.filterMode     = cudaFilterModeLinear;
    texRef.normalized     = true;

    // Bind the array to the texture reference
    cudaBindTextureToArray(texRef, cuArray, channelDesc);

    // Allocate result of transformation in device memory
    float* output;
    cudaMalloc(&output, width * height * sizeof(float));

    // Invoke kernel
    dim3 dimBlock(16, 16);
    dim3 dimGrid((width  + dimBlock.x - 1) / dimBlock.x,
                 (height + dimBlock.y - 1) / dimBlock.y);
    transformKernel<<>>(output, width, height,
                                           angle);

    // Free device memory
    cudaFreeArray(cuArray);
    cudaFree(output);

    return 0;
}

3.2.12.1.3 16位浮點(diǎn)類型紋理

CUDA 數(shù)組支持的 16 位浮點(diǎn)或half格式與 IEEE 754-2008 binary2 格式相同。

CUDA C++ 不支持匹配的數(shù)據(jù)類型，但提供了通過unsigned short類型與 32 位浮點(diǎn)格式相互轉(zhuǎn)換的內(nèi)在函數(shù)：__float2half_rn(float)和__half2float(unsigned short)。 這些功能僅在設(shè)備代碼中受支持。 例如，主機(jī)代碼的等效函數(shù)可以在 OpenEXR 庫中找到。

在執(zhí)行任何過濾之前，在紋理提取期間，16 位浮點(diǎn)組件被提升為 32 位浮點(diǎn)。

可以通過調(diào)用cudaCreateChannelDescHalf*()函數(shù)來創(chuàng)建 16 位浮點(diǎn)格式的通道描述。

3.2.12.1.4 分層紋理

一維或二維分層紋理（在 Direct3D 中也稱為紋理數(shù)組，在 OpenGL 中也稱為數(shù)組紋理）是由一系列層組成的紋理，這些層都是具有相同維度、大小和數(shù)據(jù)類型的常規(guī)紋理.

使用整數(shù)索引和浮點(diǎn)紋理坐標(biāo)來尋址一維分層紋理；索引表示序列中的層，坐標(biāo)表示該層中的texel。使用整數(shù)索引和兩個(gè)浮點(diǎn)紋理坐標(biāo)來尋址二維分層紋理；索引表示序列中的層，坐標(biāo)表示該層中的texel。

分層紋理只能是一個(gè) CUDA 數(shù)組，方法是使用cudaArrayLayered標(biāo)志調(diào)用的cudaMalloc3DArray()（一維分層紋理的高度為零）。

使用tex1DLayered()、tex1DLayered()、tex2DLayered() 和 tex2DLayered()中描述的設(shè)備函數(shù)獲取分層紋理。紋理過濾（請參閱紋理提?。﹥H在層內(nèi)完成，而不是跨層。

分層紋理僅在計(jì)算能力 2.0 及更高版本的設(shè)備上受支持。

3.2.12.1.5 立方體紋理(Cubemap Textures)

Cubemap Textures是一種特殊類型的二維分層紋理，它有六層代表立方體的面：

• 層的寬度等于它的高度。
• 立方體貼圖使用三個(gè)紋理坐標(biāo) x、y 和 z 進(jìn)行尋址，這些坐標(biāo)被解釋為從立方體中心發(fā)出并指向立方體的一個(gè)面和對應(yīng)于該面的層內(nèi)的texel的方向矢量。 更具體地說，面部是由具有最大量級(jí) m 的坐標(biāo)選擇的，相應(yīng)的層使用坐標(biāo)(s/m+1)/2和(t/m+1)/2來尋址，其中 s 和 t 在表中定義 .

通過使用cudaArrayCubemap標(biāo)志調(diào)用cudaMalloc3DArray()，立方體貼圖紋理只能是 CUDA 數(shù)組。

立方體貼圖紋理是使用texCubemap()和texCubemap()中描述的設(shè)備函數(shù)獲取的。

Cubemap 紋理僅在計(jì)算能力 2.0 及更高版本的設(shè)備上受支持。

3.2.12.1.6 分層的立方體紋理內(nèi)存(Cubemap Layered Textures)

立方體貼圖分層紋理是一種分層紋理，其層是相同維度的立方體貼圖。

使用整數(shù)索引和三個(gè)浮點(diǎn)紋理坐標(biāo)來處理立方體貼圖分層紋理； 索引表示序列中的立方體貼圖，坐標(biāo)表示該立方體貼圖中的紋理元素。

通過使用cudaArrayLayered和cudaArrayCubemap標(biāo)志調(diào)用的cudaMalloc3DArray()，立方體貼圖分層紋理只能是 CUDA 數(shù)組。

立方體貼圖分層紋理是使用texCubemapLayered()和texCubemapLayered()中描述的設(shè)備函數(shù)獲取的。 紋理過濾（請參閱紋理提取）僅在層內(nèi)完成，而不是跨層。

Cubemap 分層紋理僅在計(jì)算能力 2.0 及更高版本的設(shè)備上受支持。

3.2.12.1.7 紋理收集(Texture Gather)

紋理聚集是一種特殊的紋理提取，僅適用于二維紋理。它由tex2Dgather()函數(shù)執(zhí)行，該函數(shù)具有與tex2D()相同的參數(shù)，外加一個(gè)等于 0、1、2 或 3 的附加 comp 參數(shù)（參見tex2Dgather()和tex2Dgather()）。它返回四個(gè) 32 位數(shù)字，對應(yīng)于在常規(guī)紋理提取期間用于雙線性過濾的四個(gè)texel中每一個(gè)的分量 comp 的值。例如，如果這些紋理像素的值是 (253, 20, 31, 255), (250, 25, 29, 254), (249, 16, 37, 253), (251, 22, 30, 250)，并且comp 為 2，tex2Dgather()返回 (31, 29, 37, 30)。

請注意，紋理坐標(biāo)僅使用 8 位小數(shù)精度計(jì)算。因此，對于tex2D()將使用 1.0 作為其權(quán)重之一（α 或 β，請參閱線性過濾）的情況，tex2Dgather()可能會(huì)返回意外結(jié)果。例如，x 紋理坐標(biāo)為 2.49805：xB=x-0.5=1.99805，但是 xB 的小數(shù)部分以 8 位定點(diǎn)格式存儲(chǔ)。由于 0.99805 比 255.f/256.f 更接近 256.f/256.f，因此 xB 的值為 2。因此，在這種情況下，tex2Dgather()將返回 x 中的索引 2 和 3，而不是索引1 和 2。

紋理收集僅支持使用cudaArrayTextureGather標(biāo)志創(chuàng)建的 CUDA 數(shù)組，其寬度和高度小于表 15中為紋理收集指定的最大值，該最大值小于常規(guī)紋理提取。

紋理收集僅在計(jì)算能力 2.0 及更高版本的設(shè)備上受支持。

3.2.12.2 表面內(nèi)存(Surface Memory)

對于計(jì)算能力 2.0 及更高版本的設(shè)備，可以使用Surface Functions中描述的函數(shù)通過表面對象或表面引用來讀取和寫入使用cudaArraySurfaceLoadStore標(biāo)志創(chuàng)建的 CUDA 數(shù)組（在 Cubemap Surfaces 中描述）。

表 15列出了最大表面寬度、高度和深度，具體取決于設(shè)備的計(jì)算能力。

3.2.12.2.1 表面內(nèi)存對象API

使用cudaCreateSurfaceObject()從struct cudaResourceDesc類型的資源描述中創(chuàng)建表面內(nèi)存對象。

以下代碼示例將一些簡單的轉(zhuǎn)換內(nèi)核應(yīng)用于紋理。

// Simple copy kernel
__global__ void copyKernel(cudaSurfaceObject_t inputSurfObj,
                           cudaSurfaceObject_t outputSurfObj,
                           int width, int height) 
{
    // Calculate surface coordinates
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < width && y < height) {
        uchar4 data;
        // Read from input surface
        surf2Dread(&data,  inputSurfObj, x * 4, y);
        // Write to output surface
        surf2Dwrite(data, outputSurfObj, x * 4, y);
    }
}

// Host code
int main()
{
    const int height = 1024;
    const int width = 1024;

    // Allocate and set some host data
    unsigned char *h_data =
        (unsigned char *)std::malloc(sizeof(unsigned char) * width * height * 4);
    for (int i = 0; i < height * width * 4; ++i)
        h_data[i] = i;

    // Allocate CUDA arrays in device memory
    cudaChannelFormatDesc channelDesc =
        cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned);
    cudaArray_t cuInputArray;
    cudaMallocArray(&cuInputArray, &channelDesc, width, height,
                    cudaArraySurfaceLoadStore);
    cudaArray_t cuOutputArray;
    cudaMallocArray(&cuOutputArray, &channelDesc, width, height,
                    cudaArraySurfaceLoadStore);

    // Set pitch of the source (the width in memory in bytes of the 2D array
    // pointed to by src, including padding), we dont have any padding
    const size_t spitch = 4 * width * sizeof(unsigned char);
    // Copy data located at address h_data in host memory to device memory
    cudaMemcpy2DToArray(cuInputArray, 0, 0, h_data, spitch,
                        4 * width * sizeof(unsigned char), height,
                        cudaMemcpyHostToDevice);

    // Specify surface
    struct cudaResourceDesc resDesc;
    memset(&resDesc, 0, sizeof(resDesc));
    resDesc.resType = cudaResourceTypeArray;

    // Create the surface objects
    resDesc.res.array.array = cuInputArray;
    cudaSurfaceObject_t inputSurfObj = 0;
    cudaCreateSurfaceObject(&inputSurfObj, &resDesc);
    resDesc.res.array.array = cuOutputArray;
    cudaSurfaceObject_t outputSurfObj = 0;
    cudaCreateSurfaceObject(&outputSurfObj, &resDesc);

    // Invoke kernel
    dim3 threadsperBlock(16, 16);
    dim3 numBlocks((width + threadsperBlock.x - 1) / threadsperBlock.x,
                    (height + threadsperBlock.y - 1) / threadsperBlock.y);
    copyKernel<<>>(inputSurfObj, outputSurfObj, width,
                                                height);

    // Copy data from device back to host
    cudaMemcpy2DFromArray(h_data, spitch, cuOutputArray, 0, 0,
                            4 * width * sizeof(unsigned char), height,
                            cudaMemcpyDeviceToHost);

    // Destroy surface objects
    cudaDestroySurfaceObject(inputSurfObj);
    cudaDestroySurfaceObject(outputSurfObj);

    // Free device memory
    cudaFreeArray(cuInputArray);
    cudaFreeArray(cuOutputArray);

    // Free host memory
    free(h_data);

  return 0;
}

3.2.12.2.3 立方體表面內(nèi)存

使用surfCubemapread()和surfCubemapwrite()（surfCubemapread和surfCubemapwrite）作為二維分層表面來訪問立方體貼圖表面內(nèi)存，即，使用表示面的整數(shù)索引和尋址對應(yīng)于該面的層內(nèi)的紋素的兩個(gè)浮點(diǎn)紋理坐標(biāo) . 面的順序如表 2所示。

3.2.12.2.4 立方體分層表面內(nèi)存

使用surfCubemapLayeredread()和surfCubemapLayeredwrite()（surfCubemapLayeredread()和surfCubemapLayeredwrite()）作為二維分層表面來訪問立方體貼圖分層表面，即，使用表示立方體貼圖之一的面和兩個(gè)浮點(diǎn)紋理的整數(shù)索引 坐標(biāo)尋址對應(yīng)于該面的層內(nèi)的紋理元素。 面的順序如表 2 所示，因此例如 index ((2 * 6) + 3) 會(huì)訪問第三個(gè)立方體貼圖的第四個(gè)面。

3.2.12.3 CUDA Array

CUDA Array是針對紋理獲取優(yōu)化的不透明內(nèi)存布局。 它們是一維、二維或三維，由元素組成，每個(gè)元素有 1、2 或 4 個(gè)分量，可以是有符號(hào)或無符號(hào) 8 位、16 位或 32 位整數(shù)、16 位浮點(diǎn)數(shù)、 或 32 位浮點(diǎn)數(shù)。 CUDA Array只能由內(nèi)核通過紋理內(nèi)存中描述的紋理獲取或表面內(nèi)存中描述的表面讀取和寫入來訪問。

3.2.12.4 讀寫一致性

紋理和表面內(nèi)存被緩存（請參閱設(shè)備內(nèi)存訪問），并且在同一個(gè)內(nèi)核調(diào)用中，緩存在全局內(nèi)存寫入和表面內(nèi)存寫入方面并不保持一致，因此任何紋理獲取或表面內(nèi)存讀取到一個(gè)地址 ,在同一個(gè)內(nèi)核調(diào)用中通過全局寫入或表面寫入寫入會(huì)返回未定義的數(shù)據(jù)。 換句話說，線程可以安全地讀取某個(gè)紋理或表面內(nèi)存位置，前提是該內(nèi)存位置已被先前的內(nèi)核調(diào)用或內(nèi)存拷貝更新，但如果它先前已由同一個(gè)線程或來自同一線程的另一個(gè)線程更新，則不能內(nèi)核調(diào)用。

3.2.13圖形一致性

來自 OpenGL 和 Direct3D 的一些資源可能會(huì)映射到 CUDA 的地址空間中，以使 CUDA 能夠讀取 OpenGL 或 Direct3D 寫入的數(shù)據(jù)，或者使 CUDA 能夠?qū)懭霐?shù)據(jù)以供 OpenGL 或 Direct3D 使用。

資源必須先注冊到 CUDA，然后才能使用OpenGL 互操作和Direct3D 互操作中提到的函數(shù)進(jìn)行映射。這些函數(shù)返回一個(gè)指向struct cudaGraphicsResource類型的 CUDA 圖形資源的指針。注冊資源可能會(huì)產(chǎn)生高開銷，因此通常每個(gè)資源只調(diào)用一次。使用cudaGraphicsUnregisterResource()取消注冊 CUDA 圖形資源。每個(gè)打算使用該資源的 CUDA 上下文都需要單獨(dú)注冊它。

將資源注冊到 CUDA 后，可以根據(jù)需要使用cudaGraphicsMapResources()和cudaGraphicsUnmapResources()多次映射和取消映射。可以調(diào)用cudaGraphicsResourceSetMapFlags()來指定 CUDA 驅(qū)動(dòng)程序可以用來優(yōu)化資源管理的使用提示（只寫、只讀）。

內(nèi)核可以使用cudaGraphicsResourceGetMappedPointer()返回的設(shè)備內(nèi)存地址來讀取或?qū)懭胗成涞馁Y源，對于緩沖區(qū)，使用cudaGraphicsSubResourceGetMappedArray()的 CUDA 數(shù)組。

在映射時(shí)通過 OpenGL、Direct3D 或其他 CUDA 上下文訪問資源會(huì)產(chǎn)生未定義的結(jié)果。OpenGL 互操作和Direct3D 互操作為每個(gè)圖形 API 和一些代碼示例提供了細(xì)節(jié)。SLI 互操作給出了系統(tǒng)何時(shí)處于 SLI 模式的細(xì)節(jié)。

3.2.13.1. OpenGL 一致性

可以映射到 CUDA 地址空間的 OpenGL 資源是 OpenGL 緩沖區(qū)、紋理和渲染緩沖區(qū)對象。

使用cudaGraphicsGLRegisterBuffer()注冊緩沖區(qū)對象。在 CUDA 中，它顯示為設(shè)備指針，因此可以由內(nèi)核或通過cudaMemcpy()調(diào)用讀取和寫入。

使用cudaGraphicsGLRegisterImage()注冊紋理或渲染緩沖區(qū)對象。在 CUDA 中，它顯示為 CUDA 數(shù)組。內(nèi)核可以通過將數(shù)組綁定到紋理或表面引用來讀取數(shù)組。如果資源已使用cudaGraphicsRegisterFlagsSurfaceLoadStore標(biāo)志注冊，他們還可以通過表面寫入函數(shù)對其進(jìn)行寫入。該數(shù)組也可以通過cudaMemcpy2D()調(diào)用來讀取和寫入。cudaGraphicsGLRegisterImage()支持具有 1、2 或 4 個(gè)分量和內(nèi)部浮點(diǎn)類型（例如，GL_RGBA_FLOAT32）、標(biāo)準(zhǔn)化整數(shù)（例如，GL_RGBA8、GL_INTENSITY16）和非標(biāo)準(zhǔn)化整數(shù)（例如，GL_RGBA8UI）的所有紋理格式（請注意，由于非標(biāo)準(zhǔn)化整數(shù)格式需要 OpenGL 3.0，它們只能由著色器編寫，而不是固定函數(shù)管道）。

正在共享資源的 OpenGL 上下文對于進(jìn)行任何 OpenGL 互操作性 API 調(diào)用的主機(jī)線程來說必須是最新的。

請注意：當(dāng) OpenGL 紋理設(shè)置為無綁定時(shí)（例如，通過使用glGetTextureHandle*/glGetImageHandle*API 請求圖像或紋理句柄），它不能在 CUDA 中注冊。應(yīng)用程序需要在請求圖像或紋理句柄之前注冊紋理以進(jìn)行互操作。

以下代碼示例使用內(nèi)核動(dòng)態(tài)修改存儲(chǔ)在頂點(diǎn)緩沖區(qū)對象中的 2D width x height 網(wǎng)格：

GLuint positionsVBO;
struct cudaGraphicsResource* positionsVBO_CUDA;

int main()
{
    // Initialize OpenGL and GLUT for device 0
    // and make the OpenGL context current
    ...
    glutDisplayFunc(display);

    // Explicitly set device 0
    cudaSetDevice(0);

    // Create buffer object and register it with CUDA
    glGenBuffers(1, &positionsVBO);
    glBindBuffer(GL_ARRAY_BUFFER, positionsVBO);
    unsigned int size = width * height * 4 * sizeof(float);
    glBufferData(GL_ARRAY_BUFFER, size, 0, GL_DYNAMIC_DRAW);
    glBindBuffer(GL_ARRAY_BUFFER, 0);
    cudaGraphicsGLRegisterBuffer(&positionsVBO_CUDA,
                                 positionsVBO,
                                 cudaGraphicsMapFlagsWriteDiscard);

    // Launch rendering loop
    glutMainLoop();

    ...
}

void display()
{
    // Map buffer object for writing from CUDA
    float4* positions;
    cudaGraphicsMapResources(1, &positionsVBO_CUDA, 0);
    size_t num_bytes; 
    cudaGraphicsResourceGetMappedPointer((void**)&positions,
                                         &num_bytes,  
                                         positionsVBO_CUDA));

    // Execute kernel
    dim3 dimBlock(16, 16, 1);
    dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
    createVertices<<>>(positions, time,
                                          width, height);

    // Unmap buffer object
    cudaGraphicsUnmapResources(1, &positionsVBO_CUDA, 0);

    // Render from buffer object
    glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
    glBindBuffer(GL_ARRAY_BUFFER, positionsVBO);
    glVertexPointer(4, GL_FLOAT, 0, 0);
    glEnableClientState(GL_VERTEX_ARRAY);
    glDrawArrays(GL_POINTS, 0, width * height);
    glDisableClientState(GL_VERTEX_ARRAY);

    // Swap buffers
    glutSwapBuffers();
    glutPostRedisplay();
}
void deleteVBO()
{
    cudaGraphicsUnregisterResource(positionsVBO_CUDA);
    glDeleteBuffers(1, &positionsVBO);
}

__global__ void createVertices(float4* positions, float time,
                               unsigned int width, unsigned int height)
{
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    // Calculate uv coordinates
    float u = x / (float)width;
    float v = y / (float)height;
    u = u * 2.0f - 1.0f;
    v = v * 2.0f - 1.0f;

    // calculate simple sine wave pattern
    float freq = 4.0f;
    float w = sinf(u * freq + time)
            * cosf(v * freq + time) * 0.5f;

    // Write positions
    positions[y * width + x] = make_float4(u, w, v, 1.0f);
}

在 Windows 和 Quadro GPU 上，cudaWGLGetDevice()可用于檢索與wglEnumGpusNV()返回的句柄關(guān)聯(lián)的 CUDA 設(shè)備。 Quadro GPU 在多 GPU 配置中提供比 GeForce 和 Tesla GPU 更高性能的 OpenGL 互操作性，其中 OpenGL 渲染在 Quadro GPU 上執(zhí)行，CUDA 計(jì)算在系統(tǒng)中的其他 GPU 上執(zhí)行。

3.2.13.2. Direct3D 一致性

Direct3D 9Ex、Direct3D 10 和 Direct3D 11 支持 Direct3D 互操作性。

CUDA 上下文只能與滿足以下條件的 Direct3D 設(shè)備互操作： Direct3D 9Ex 設(shè)備必須使用設(shè)置為D3DDEVTYPE_HAL的DeviceType和使用D3DCREATE_HARDWARE_VERTEXPROCESSING標(biāo)志的BehaviorFlags創(chuàng)建； Direct3D 10 和 Direct3D 11 設(shè)備必須在DriverType設(shè)置為D3D_DRIVER_TYPE_HARDWARE的情況下創(chuàng)建。

可以映射到 CUDA 地址空間的 Direct3D 資源是 Direct3D 緩沖區(qū)、紋理和表面。 這些資源使用cudaGraphicsD3D9RegisterResource()、cudaGraphicsD3D10RegisterResource()和cudaGraphicsD3D11RegisterResource()注冊。

以下代碼示例使用內(nèi)核動(dòng)態(tài)修改存儲(chǔ)在頂點(diǎn)緩沖區(qū)對象中的 2D width x height網(wǎng)格。

Direct3D 9 Version:

IDirect3D9* D3D;
IDirect3DDevice9* device;
struct CUSTOMVERTEX {
    FLOAT x, y, z;
    DWORD color;
};
IDirect3DVertexBuffer9* positionsVB;
struct cudaGraphicsResource* positionsVB_CUDA;

int main()
{
    int dev;
    // Initialize Direct3D
    D3D = Direct3DCreate9Ex(D3D_SDK_VERSION);

    // Get a CUDA-enabled adapter
    unsigned int adapter = 0;
    for (; adapter < g_pD3D->GetAdapterCount(); adapter++) {
        D3DADAPTER_IDENTIFIER9 adapterId;
        g_pD3D->GetAdapterIdentifier(adapter, 0, &adapterId);
        if (cudaD3D9GetDevice(&dev, adapterId.DeviceName)
            == cudaSuccess)
            break;
    }

     // Create device
    ...
    D3D->CreateDeviceEx(adapter, D3DDEVTYPE_HAL, hWnd,
                        D3DCREATE_HARDWARE_VERTEXPROCESSING,
                        ¶ms, NULL, &device);

    // Use the same device
    cudaSetDevice(dev);

    // Create vertex buffer and register it with CUDA
    unsigned int size = width * height * sizeof(CUSTOMVERTEX);
    device->CreateVertexBuffer(size, 0, D3DFVF_CUSTOMVERTEX,
                               D3DPOOL_DEFAULT, &positionsVB, 0);
    cudaGraphicsD3D9RegisterResource(&positionsVB_CUDA,
                                     positionsVB,
                                     cudaGraphicsRegisterFlagsNone);
    cudaGraphicsResourceSetMapFlags(positionsVB_CUDA,
                                    cudaGraphicsMapFlagsWriteDiscard);

    // Launch rendering loop
    while (...) {
        ...
        Render();
        ...
    }
    ...
}
void Render()
{
    // Map vertex buffer for writing from CUDA
    float4* positions;
    cudaGraphicsMapResources(1, &positionsVB_CUDA, 0);
    size_t num_bytes; 
    cudaGraphicsResourceGetMappedPointer((void**)&positions,
                                         &num_bytes,  
                                         positionsVB_CUDA));

    // Execute kernel
    dim3 dimBlock(16, 16, 1);
    dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
    createVertices<<>>(positions, time,
                                          width, height);

    // Unmap vertex buffer
    cudaGraphicsUnmapResources(1, &positionsVB_CUDA, 0);

    // Draw and present
    ...
}

void releaseVB()
{
    cudaGraphicsUnregisterResource(positionsVB_CUDA);
    positionsVB->Release();
}

__global__ void createVertices(float4* positions, float time,
                               unsigned int width, unsigned int height)
{
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    // Calculate uv coordinates
    float u = x / (float)width;
    float v = y / (float)height;
    u = u * 2.0f - 1.0f;
    v = v * 2.0f - 1.0f;

    // Calculate simple sine wave pattern
    float freq = 4.0f;
    float w = sinf(u * freq + time)
            * cosf(v * freq + time) * 0.5f;

    // Write positions
    positions[y * width + x] =
                make_float4(u, w, v, __int_as_float(0xff00ff00));
}

Direct3D 10 Version

ID3D10Device* device;
struct CUSTOMVERTEX {
    FLOAT x, y, z;
    DWORD color;
};
ID3D10Buffer* positionsVB;
struct cudaGraphicsResource* positionsVB_CUDA;
            
int main()
{
    int dev;
    // Get a CUDA-enabled adapter
    IDXGIFactory* factory;
    CreateDXGIFactory(__uuidof(IDXGIFactory), (void**)&factory);
    IDXGIAdapter* adapter = 0;
    for (unsigned int i = 0; !adapter; ++i) {
        if (FAILED(factory->EnumAdapters(i, &adapter))
            break;
        if (cudaD3D10GetDevice(&dev, adapter) == cudaSuccess)
            break;
        adapter->Release();
    }
    factory->Release();
            
    // Create swap chain and device
    ...
    D3D10CreateDeviceAndSwapChain(adapter, 
                                  D3D10_DRIVER_TYPE_HARDWARE, 0, 
                                  D3D10_CREATE_DEVICE_DEBUG,
                                  D3D10_SDK_VERSION, 
                                  &swapChainDesc, &swapChain,
                                  &device);
    adapter->Release();

    // Use the same device
    cudaSetDevice(dev);

    // Create vertex buffer and register it with CUDA
    unsigned int size = width * height * sizeof(CUSTOMVERTEX);
    D3D10_BUFFER_DESC bufferDesc;
    bufferDesc.Usage          = D3D10_USAGE_DEFAULT;
    bufferDesc.ByteWidth      = size;
    bufferDesc.BindFlags      = D3D10_BIND_VERTEX_BUFFER;
    bufferDesc.CPUAccessFlags = 0;
    bufferDesc.MiscFlags      = 0;
    device->CreateBuffer(&bufferDesc, 0, &positionsVB);
    cudaGraphicsD3D10RegisterResource(&positionsVB_CUDA,
                                      positionsVB,
                                      cudaGraphicsRegisterFlagsNone);
                                      cudaGraphicsResourceSetMapFlags(positionsVB_CUDA,
                                      cudaGraphicsMapFlagsWriteDiscard);

    // Launch rendering loop
    while (...) {
        ...
        Render();
        ...
    }
    ...
}
void Render()
{
    // Map vertex buffer for writing from CUDA
    float4* positions;
    cudaGraphicsMapResources(1, &positionsVB_CUDA, 0);
    size_t num_bytes; 
    cudaGraphicsResourceGetMappedPointer((void**)&positions,
                                         &num_bytes,  
                                         positionsVB_CUDA));

    // Execute kernel
    dim3 dimBlock(16, 16, 1);
    dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
    createVertices<<>>(positions, time,
                                          width, height);

    // Unmap vertex buffer
    cudaGraphicsUnmapResources(1, &positionsVB_CUDA, 0);

    // Draw and present
    ...
}

void releaseVB()
{
    cudaGraphicsUnregisterResource(positionsVB_CUDA);
    positionsVB->Release();
}

__global__ void createVertices(float4* positions, float time,
                               unsigned int width, unsigned int height)
{
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    // Calculate uv coordinates
    float u = x / (float)width;
    float v = y / (float)height;
    u = u * 2.0f - 1.0f;
    v = v * 2.0f - 1.0f;

    // Calculate simple sine wave pattern
    float freq = 4.0f;
    float w = sinf(u * freq + time)
            * cosf(v * freq + time) * 0.5f;
            
    // Write positions
    positions[y * width + x] =
                make_float4(u, w, v, __int_as_float(0xff00ff00));
}

Direct3D 11 Version

ID3D11Device* device;
struct CUSTOMVERTEX {
    FLOAT x, y, z;
    DWORD color;
};
ID3D11Buffer* positionsVB;
struct cudaGraphicsResource* positionsVB_CUDA;

int main()
{
    int dev;
    // Get a CUDA-enabled adapter
    IDXGIFactory* factory;
    CreateDXGIFactory(__uuidof(IDXGIFactory), (void**)&factory);
    IDXGIAdapter* adapter = 0;
    for (unsigned int i = 0; !adapter; ++i) {
        if (FAILED(factory->EnumAdapters(i, &adapter))
            break;
        if (cudaD3D11GetDevice(&dev, adapter) == cudaSuccess)
            break;
        adapter->Release();
    }
    factory->Release();

    // Create swap chain and device
    ...
    sFnPtr_D3D11CreateDeviceAndSwapChain(adapter, 
                                         D3D11_DRIVER_TYPE_HARDWARE,
                                         0, 
                                         D3D11_CREATE_DEVICE_DEBUG,
                                         featureLevels, 3,
                                         D3D11_SDK_VERSION, 
                                         &swapChainDesc, &swapChain,
                                         &device,
                                         &featureLevel,
                                         &deviceContext);
    adapter->Release();

    // Use the same device
    cudaSetDevice(dev);

    // Create vertex buffer and register it with CUDA
    unsigned int size = width * height * sizeof(CUSTOMVERTEX);
    D3D11_BUFFER_DESC bufferDesc;
    bufferDesc.Usage          = D3D11_USAGE_DEFAULT;
    bufferDesc.ByteWidth      = size;
    bufferDesc.BindFlags      = D3D11_BIND_VERTEX_BUFFER;
    bufferDesc.CPUAccessFlags = 0;
    bufferDesc.MiscFlags      = 0;
    device->CreateBuffer(&bufferDesc, 0, &positionsVB);
    cudaGraphicsD3D11RegisterResource(&positionsVB_CUDA,
                                      positionsVB,
                                      cudaGraphicsRegisterFlagsNone);
    cudaGraphicsResourceSetMapFlags(positionsVB_CUDA,
                                    cudaGraphicsMapFlagsWriteDiscard);

    // Launch rendering loop
    while (...) {
        ...
        Render();
        ...
    }
    ...
}
void Render()
{
    // Map vertex buffer for writing from CUDA
    float4* positions;
    cudaGraphicsMapResources(1, &positionsVB_CUDA, 0);
    size_t num_bytes; 
    cudaGraphicsResourceGetMappedPointer((void**)&positions,
                                         &num_bytes,  
                                         positionsVB_CUDA));

    // Execute kernel
    dim3 dimBlock(16, 16, 1);
    dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
    createVertices<<>>(positions, time,
                                          width, height);

    // Unmap vertex buffer
    cudaGraphicsUnmapResources(1, &positionsVB_CUDA, 0);

    // Draw and present
    ...
}

void releaseVB()
{
    cudaGraphicsUnregisterResource(positionsVB_CUDA);
    positionsVB->Release();
}

    __global__ void createVertices(float4* positions, float time,
                          unsigned int width, unsigned int height)
{
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

// Calculate uv coordinates
    float u = x / (float)width;
    float v = y / (float)height;
    u = u * 2.0f - 1.0f;
    v = v * 2.0f - 1.0f;

    // Calculate simple sine wave pattern
    float freq = 4.0f;
    float w = sinf(u * freq + time)
            * cosf(v * freq + time) * 0.5f;

    // Write positions
    positions[y * width + x] =
                make_float4(u, w, v, __int_as_float(0xff00ff00));
}

3.2.13.3 SLI一致性

在具有多個(gè) GPU 的系統(tǒng)中，所有支持 CUDA 的 GPU 都可以通過 CUDA 驅(qū)動(dòng)程序和運(yùn)行時(shí)作為單獨(dú)的設(shè)備進(jìn)行訪問。然而，當(dāng)系統(tǒng)處于 SLI 模式時(shí)，有如下所述的特殊注意事項(xiàng)。

首先，在一個(gè) GPU 上的一個(gè) CUDA 設(shè)備中的分配將消耗其他 GPU 上的內(nèi)存，這些 GPU 是 Direct3D 或 OpenGL 設(shè)備的 SLI 配置的一部分。因此，分配可能會(huì)比預(yù)期的更早失敗。

其次，應(yīng)用程序應(yīng)該創(chuàng)建多個(gè) CUDA 上下文，一個(gè)用于 SLI 配置中的每個(gè) GPU。雖然這不是嚴(yán)格要求，但它避免了設(shè)備之間不必要的數(shù)據(jù)傳輸。應(yīng)用程序可以將cudaD3D[9|10|11]GetDevices()用于 Direct3D 和cudaGLGetDevices()用于 OpenGL 調(diào)用，以識(shí)別當(dāng)前執(zhí)行渲染的設(shè)備的 CUDA 設(shè)備句柄和下一幀。鑒于此信息，應(yīng)用程序通常會(huì)選擇適當(dāng)?shù)脑O(shè)備并將 Direct3D 或 OpenGL 資源映射到由cudaD3D[9|10|11]GetDevices()或當(dāng)deviceList參數(shù)設(shè)置為cudaD3D[9|10 |11]DeviceListCurrentFrame或cudaGLDeviceListCurrentFrame。

請注意，從cudaGraphicsD9D[9|10|11]RegisterResource和cudaGraphicsGLRegister[Buffer|Image]返回的資源只能在發(fā)生注冊的設(shè)備上使用。因此，在 SLI 配置中，當(dāng)在不同的 CUDA 設(shè)備上計(jì)算不同幀的數(shù)據(jù)時(shí)，有必要分別為每個(gè)設(shè)備注冊資源。

有關(guān) CUDA 運(yùn)行時(shí)如何分別與 Direct3D 和 OpenGL 互操作的詳細(xì)信息，請參閱Direct3D 互操作性和OpenGL 互操作性。

3.2.14 擴(kuò)展資源一致性

這里待定(實(shí)際上是作者不熟悉)

3.2.15 CUDA用戶對象

CUDA 用戶對象可用于幫助管理 CUDA 中異步工作所使用的資源的生命周期。 特別是，此功能對于CUDA 圖和流捕獲非常有用。

各種資源管理方案與 CUDA 圖不兼容。 例如，考慮基于事件的池或同步創(chuàng)建、異步銷毀方案。

// Library API with pool allocation
void libraryWork(cudaStream_t stream) {
    auto &resource = pool.claimTemporaryResource();
    resource.waitOnReadyEventInStream(stream);
    launchWork(stream, resource);
    resource.recordReadyEvent(stream);
}
// Library API with asynchronous resource deletion
void libraryWork(cudaStream_t stream) {
    Resource *resource = new Resource(...);
    launchWork(stream, resource);
    cudaStreamAddCallback(
        stream,
        [](cudaStream_t, cudaError_t, void *resource) {
            delete static_cast(resource);
        },
        resource,
        0);
    // Error handling considerations not shown
}

由于需要間接或圖更新的資源的非固定指針或句柄，以及每次提交工作時(shí)需要同步 CPU 代碼，這些方案對于 CUDA 圖來說是困難的。如果這些注意事項(xiàng)對庫的調(diào)用者隱藏，并且由于在捕獲期間使用了不允許的 API，它們也不適用于流捕獲。存在各種解決方案，例如將資源暴露給調(diào)用者。 CUDA 用戶對象提供了另一種方法。

CUDA 用戶對象將用戶指定的析構(gòu)函數(shù)回調(diào)與內(nèi)部引用計(jì)數(shù)相關(guān)聯(lián)，類似于 C++shared_ptr。引用可能歸 CPU 上的用戶代碼和 CUDA 圖所有。請注意，對于用戶擁有的引用，與 C++ 智能指針不同，沒有代表引用的對象；用戶必須手動(dòng)跟蹤用戶擁有的引用。一個(gè)典型的用例是在創(chuàng)建用戶對象后立即將唯一的用戶擁有的引用移動(dòng)到 CUDA 圖。

當(dāng)引用關(guān)聯(lián)到 CUDA 圖時(shí)，CUDA 將自動(dòng)管理圖操作?？寺〉?code style="font-size:inherit;color:inherit;margin:0px;padding:0px;border:0px;font-style:inherit;font-variant:inherit;font-weight:inherit;line-height:inherit;vertical-align:baseline;background-color:rgb(244,244,244);">cudaGraph_t保留源cudaGraph_t擁有的每個(gè)引用的副本，具有相同的多重性。實(shí)例化的cudaGraphExec_t保留源cudaGraph_t中每個(gè)引用的副本。當(dāng)cudaGraphExec_t在未同步的情況下被銷毀時(shí)，引用將保留到執(zhí)行完成。

這是一個(gè)示例用法。

cudaGraph_t graph;  // Preexisting graph

Object *object = new Object;  // C++ object with possibly nontrivial destructor
cudaUserObject_t cuObject;
cudaUserObjectCreate(
    &cuObject,
    object,  // Here we use a CUDA-provided template wrapper for this API,
             // which supplies a callback to delete the C++ object pointer
    1,  // Initial refcount
    cudaUserObjectNoDestructorSync  // Acknowledge that the callback cannot be
                                    // waited on via CUDA
);
cudaGraphRetainUserObject(
    graph,
    cuObject,
    1,  // Number of references
    cudaGraphUserObjectMove  // Transfer a reference owned by the caller (do
                             // not modify the total reference count)
);
// No more references owned by this thread; no need to call release API
cudaGraphExec_t graphExec;
cudaGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0);  // Will retain a
                                                               // new reference
cudaGraphDestroy(graph);  // graphExec still owns a reference
cudaGraphLaunch(graphExec, 0);  // Async launch has access to the user objects
cudaGraphExecDestroy(graphExec);  // Launch is not synchronized; the release
                                  // will be deferred if needed
cudaStreamSynchronize(0);  // After the launch is synchronized, the remaining
                           // reference is released and the destructor will
                           // execute. Note this happens asynchronously.
// If the destructor callback had signaled a synchronization object, it would
// be safe to wait on it at this point.

子圖節(jié)點(diǎn)中的圖所擁有的引用與子圖相關(guān)聯(lián)，而不是與父圖相關(guān)聯(lián)。如果更新或刪除子圖，則引用會(huì)相應(yīng)更改。如果使用cudaGraphExecUpdate或cudaGraphExecChildGraphNodeSetParams更新可執(zhí)行圖或子圖，則會(huì)克隆新源圖中的引用并替換目標(biāo)圖中的引用。在任何一種情況下，如果先前的啟動(dòng)不同步，則將保留任何將被釋放的引用，直到啟動(dòng)完成執(zhí)行。

目前沒有通過 CUDA API 等待用戶對象析構(gòu)函數(shù)的機(jī)制。用戶可以從析構(gòu)代碼中手動(dòng)發(fā)出同步對象的信號(hào)。另外，從析構(gòu)函數(shù)調(diào)用 CUDA API 是不合法的，類似于對cudaLaunchHostFunc的限制。這是為了避免阻塞 CUDA 內(nèi)部共享線程并阻止前進(jìn)。如果依賴是一種方式并且執(zhí)行調(diào)用的線程不能阻止 CUDA 工作的前進(jìn)進(jìn)度，則向另一個(gè)線程發(fā)出執(zhí)行 API 調(diào)用的信號(hào)是合法的。

用戶對象是使用cudaUserObjectCreate創(chuàng)建的，這是瀏覽相關(guān) API 的一個(gè)很好的起點(diǎn)。

3.3 版本和兼容性

開發(fā)人員在開發(fā) CUDA 應(yīng)用程序時(shí)應(yīng)該關(guān)注兩個(gè)版本號(hào)：描述計(jì)算設(shè)備的一般規(guī)范和特性的計(jì)算能力（請參閱計(jì)算能力）和描述受支持的特性的 CUDA 驅(qū)動(dòng)程序 API 的版本。驅(qū)動(dòng)程序 API 和運(yùn)行時(shí)。

驅(qū)動(dòng)程序 API 的版本在驅(qū)動(dòng)程序頭文件中定義為CUDA_VERSION。它允許開發(fā)人員檢查他們的應(yīng)用程序是否需要比當(dāng)前安裝的設(shè)備驅(qū)動(dòng)程序更新的設(shè)備驅(qū)動(dòng)程序。這很重要，因?yàn)轵?qū)動(dòng) API 是向后兼容的，這意味著針對特定版本的驅(qū)動(dòng) API 編譯的應(yīng)用程序、插件和庫（包括 CUDA 運(yùn)行時(shí)）將繼續(xù)在后續(xù)的設(shè)備驅(qū)動(dòng)版本上工作，如下圖所示. 驅(qū)動(dòng) API 不向前兼容，這意味著針對特定版本的驅(qū)動(dòng) API 編譯的應(yīng)用程序、插件和庫（包括 CUDA 運(yùn)行時(shí)）將不適用于以前版本的設(shè)備驅(qū)動(dòng)。

需要注意的是，支持的版本的混合和匹配存在限制：

• 由于系統(tǒng)上一次只能安裝一個(gè)版本的 CUDA 驅(qū)動(dòng)程序，因此安裝的驅(qū)動(dòng)程序必須與必須在已建成的系統(tǒng)其上運(yùn)行的任何應(yīng)用程序、插件或庫所依據(jù)的最大驅(qū)動(dòng)程序 API 版本相同或更高版本 。
• 應(yīng)用程序使用的所有插件和庫必須使用相同版本的 CUDA 運(yùn)行時(shí)，除非它們靜態(tài)鏈接到運(yùn)行時(shí)，在這種情況下，運(yùn)行時(shí)的多個(gè)版本可以共存于同一進(jìn)程空間中。 請注意，如果使用 nvcc 鏈接應(yīng)用程序，則默認(rèn)使用靜態(tài)版本的 CUDA Runtime 庫，并且所有 CUDA Toolkit 庫都針對 CUDA Runtime 靜態(tài)鏈接。
• 應(yīng)用程序使用的所有插件和庫必須使用與使用運(yùn)行時(shí)的任何庫（例如 cuFFT、cuBLAS…）相同的版本，除非靜態(tài)鏈接到這些庫。

對于 Tesla GPU 產(chǎn)品，CUDA 10 為 CUDA 驅(qū)動(dòng)程序的用戶模式組件引入了新的向前兼容升級(jí)路徑。 此功能在CUDA 兼容性中進(jìn)行了描述。 此處描述的對 CUDA 驅(qū)動(dòng)程序版本的要求適用于用戶模式組件的版本。

3.4 Compute Modes

在運(yùn)行 Windows Server 2008 及更高版本或 Linux 的 Tesla 解決方案上，可以使用 NVIDIA 的系統(tǒng)管理接口 (nvidia-smi) 將系統(tǒng)中的任何設(shè)備設(shè)置為以下三種模式之一，這是作為驅(qū)動(dòng)程序一部分分發(fā)的工具：

• 默認(rèn)計(jì)算模式：多個(gè)主機(jī)線程可以同時(shí)使用該設(shè)備（通過在此設(shè)備上調(diào)用cudaSetDevice()，當(dāng)使用運(yùn)行時(shí) API 時(shí)，或通過使 current 成為與設(shè)備關(guān)聯(lián)的上下文，當(dāng)使用驅(qū)動(dòng)程序 API 時(shí)）。
• 獨(dú)占進(jìn)程計(jì)算模式：在設(shè)備上只能在系統(tǒng)中的所有進(jìn)程中創(chuàng)建一個(gè) CUDA 上下文。 在創(chuàng)建該上下文的進(jìn)程中，該上下文可以是當(dāng)前任意數(shù)量的線程。
• 禁止的計(jì)算模式：不能在設(shè)備上創(chuàng)建 CUDA 上下文。

這尤其意味著，如果設(shè)備 0 處于禁止模式或獨(dú)占進(jìn)程模式并被另一個(gè)設(shè)備使用，則使用運(yùn)行時(shí) API 而不顯式調(diào)用cudaSetDevice()的主機(jī)線程可能與設(shè)備 0 以外的設(shè)備相關(guān)聯(lián)過程。cudaSetValidDevices()可用于從設(shè)備的優(yōu)先級(jí)列表中設(shè)置設(shè)備。

另請注意，對于采用 Pascal 架構(gòu)（具有主要修訂號(hào) 6 及更高版本的計(jì)算能力）的設(shè)備，存在對計(jì)算搶占的支持。這允許計(jì)算任務(wù)在指令級(jí)粒度上被搶占，而不是像以前的 Maxwell 和 Kepler GPU 架構(gòu)中那樣以線程塊粒度進(jìn)行搶占，其好處是可以防止具有長時(shí)間運(yùn)行內(nèi)核的應(yīng)用程序壟斷系統(tǒng)或超時(shí)。但是，將存在與計(jì)算搶占相關(guān)的上下文切換開銷，它會(huì)在支持的設(shè)備上自動(dòng)啟用。具有屬性cudaDevAttrComputePreemptionSupported的單個(gè)屬性查詢函數(shù)cudaDeviceGetAttribute()可用于確定正在使用的設(shè)備是否支持計(jì)算搶占。希望避免與不同進(jìn)程相關(guān)的上下文切換開銷的用戶可以通過選擇獨(dú)占進(jìn)程模式來確保在 GPU 上只有一個(gè)進(jìn)程處于活動(dòng)狀態(tài)。

應(yīng)用程序可以通過檢查computeMode設(shè)備屬性來查詢設(shè)備的計(jì)算模式。

3.5 模式切換

具有顯示輸出的 GPU 將一些 DRAM 內(nèi)存專用于所謂的主畫面，用于刷新用戶查看其輸出的顯示設(shè)備。當(dāng)用戶通過更改顯示器的分辨率或位深度（使用 NVIDIA 控制面板或 Windows 上的顯示控制面板）來啟動(dòng)顯示器的模式切換時(shí)，主表面所需的內(nèi)存量會(huì)發(fā)生變化。例如，如果用戶將顯示分辨率從 1280x1024x32 位更改為 1600x1200x32 位，則系統(tǒng)必須將 7.68 MB 專用于主畫面，而不是 5.24 MB。 （在啟用抗鋸齒的情況下運(yùn)行的全屏圖形應(yīng)用程序可能需要更多的主畫面顯示內(nèi)存。）在 Windows 上，可能會(huì)啟動(dòng)顯示模式切換的其他事件包括啟動(dòng)全屏 DirectX 應(yīng)用程序，按 Alt+Tab 來完成任務(wù)從全屏 DirectX 應(yīng)用程序切換，或按 Ctrl+Alt+Del 鎖定計(jì)算機(jī)。

如果模式切換增加了主畫面所需的內(nèi)存量，系統(tǒng)可能不得不蠶食專用于 CUDA 應(yīng)用程序的內(nèi)存分配。因此，模式切換會(huì)導(dǎo)致對 CUDA 運(yùn)行時(shí)的任何調(diào)用失敗并返回?zé)o效的上下文錯(cuò)誤。

3.6 在Windows上的Tesla計(jì)算集群

使用 NVIDIA 的系統(tǒng)管理界面 (nvidia-smi)，可以將 Windows 設(shè)備驅(qū)動(dòng)程序置于 Tesla 和 Quadro 系列設(shè)備的 TCC（Tesla Compute Cluster）模式。

關(guān)于作者

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

審核編輯：郭婷