一. 圖
圖為CUDA中的工作送出提供了一種新的模型。圖是一系列操作，如核心啟動，由依賴項連接配接，依賴項與執行分開定義。這允許定義一次圖形，然後重複啟動。将圖的定義與其執行分離可以實作許多優化：第一，與流相比，CPU啟動成本降低，因為大部分設定是提前完成的；第二，将整個工作流呈現給CUDA可以實作優化，而流的分段工作送出機制可能無法實作優化。
要檢視圖形可能的優化，請考慮流中發生的情況：将核心放入流中時，主機驅動程式執行一系列操作，以準備在GPU上執行核心。這些操作是設定和啟動核心所必需的，它們是一種開銷，必須為釋出的每個核心支付。對于執行時間較短的GPU核心，這種開銷可能是整個端到端執行時間的一個重要部分。
使用圖的工作送出分為三個不同的階段：定義、執行個體化和執行。
在定義階段，程式将建立對圖形中的操作及其依賴關系的描述。
執行個體化擷取圖形模闆的快照，對其進行驗證，并執行大部分設定和初始化工作，以最小化啟動時需要執行的操作。結果執行個體稱為可執行圖。
一個可執行圖可以被發送到一個流中，類似于任何其他CUDA工作。它可以在不重複執行個體化的情況下啟動任意次數。
二. 圖形結構
操作在圖中形成一個節點。操作之間的依賴關系是邊。這些依賴關系限制操作的執行順序。 一旦操作所依賴的節點完成，就可以随時排程該操作。日程安排由CUDA系統決定。 三. 節點類型
圖形節點可以是：
核心
CPU函數調用
記憶體複制
清零
空節點
子圖：執行單獨的嵌套圖。見圖11。
圖11. 子圖示例
四. 使用圖形api建立圖形
圖形可以通過兩種機制建立：顯式API和流捕獲。下面是建立和執行下圖的示例。
圖12. 用圖形api建立圖形示例
// 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
五. 使用流捕獲建立圖
流捕獲提供了一種從現有的基于流的api建立圖的機制。将工作啟動到流（包括現有代碼）中的一段代碼可以用對cudaStreamBeginCapture（）和cudastreamndcapture（）的調用括起來。見下文。
cudaGraph_t graph; cudaStreamBeginCapture(stream);
kernel_A<<< ..., stream >>>(...);
kernel_B<<< ..., stream >>>(...);
libraryCall(stream); kernel_C<<< ..., stream >>>(...);
cudaStreamEndCapture(stream, &graph);
調用cudaStreamBeginCapture（）會将流置于捕獲模式。捕獲流時，啟動到流中的工作不會排隊執行。它被附加到一個正在逐漸建立的内部圖中。然後通過調用cudastreamndcapture（）傳回此圖，該函數也結束流的捕獲模式。由流捕獲主動構造的圖稱為捕獲圖。 流捕獲可用于除cudaStreamLegacy以外的任何CUDA流（“空流”）。注意，它可以用于cudaStreamPerThread。如果程式正在使用遺留流，則可以将流0重新定義為每個線程的流，而無需更改函數。請參見預設流。
可以使用cudaStreamIsCapturing（）查詢是否正在捕獲流。
六. 跨流依賴項和事件
流捕獲可以處理用cudaEventRecord（）和cudaStreamWaitEvent（）表示的跨流依賴關系，前提是等待的事件被記錄到同一個捕獲圖中。
當事件記錄在處于捕獲模式的流中時，它将導緻捕獲的事件。捕獲的事件表示捕獲圖中的一組節點。
當捕獲的事件被流等待時，如果流尚未處于捕獲模式，則它會将該流置于捕獲模式，流中的下一項将對捕獲的事件中的節點具有額外的依賴關系。然後将這兩個流捕獲到同一個捕獲圖。
當流捕獲中存在跨流依賴項時，仍必須在調用cudaStreamBeginCapture（）的同一流中調用cudastreamndcapture（）；這是源流。由于基于事件的依賴關系，被捕獲到同一捕獲圖的任何其他流也必須連接配接回原始流。如下所示。在cudaStreamEndCapture（）上，所有捕獲到同一捕獲圖的流都将退出捕獲模式。未能重新加入原始流将導緻整個捕獲操作失敗。
// 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
上述代碼傳回的圖如圖12所示。
注意：當流退出捕獲模式時，流中的下一個未捕獲項（如果有）仍将依賴于最新的先前未捕獲項，盡管中間項已被移除。
七. 禁止和未處理的操作
同步或查詢正在捕獲的流或捕獲的事件的執行狀态是無效的，因為它們不表示計劃執行的項。當任何關聯的流處于捕獲模式時，查詢或同步包含活動流捕獲（例如裝置或上下文句柄）的更寬句柄的執行狀态也是無效的。
當捕獲同一上下文中的任何流時，并且該流不是使用cudaStreamNonBlocking建立的，則嘗試使用遺留流的任何操作都是無效的。這是因為遺留流句柄始終包含這些其他流；加入遺留流隊列将建立對正在捕獲的流的依賴關系，查詢或同步它将查詢或同步正在捕獲的流。 是以，在這種情況下調用同步api也是無效的。同步api，例如cudammcpy（），在傳回之前将隊列工作到遺留流并同步它。
注意：一般情況下，當依賴關系将被捕獲的内容與未被捕獲的内容連接配接起來并排隊等待執行時，CUDA甯願傳回錯誤，而不是忽略依賴關系。将流置于捕獲模式或置于捕獲模式之外時會發生異常；這會切斷在模式轉換之前和之後添加到流中的項之間的依賴關系。 通過等待從正在捕獲的流中捕獲的事件來合并兩個單獨的捕獲圖是無效的，該流與事件之外的另一個捕獲圖相關聯。等待正在捕獲的流中的未捕獲事件是無效的。
圖中目前不支援将異步操作排隊到流中的少數API，如果使用正在捕獲的流（如cudastreamattachemasync（））調用這些API，則會傳回錯誤。
八. 無效
在流捕獲期間嘗試無效操作時，任何關聯的捕獲圖都将無效。當捕獲圖失效時，進一步使用正在捕獲的任何流或與該圖相關聯的已捕獲事件是無效的，并且将傳回錯誤，直到流捕獲以cudastreamndcapture（）結束。此調用将使關聯的流退出捕獲模式，但也将傳回一個錯誤值和一個空圖。
九. 使用圖形API
CudaGraph_t對象不是線程安全的。使用者有責任確定多個線程不會同時通路同一個cudaGraph。
cudaGraphExec不能與自身同時運作。cudaGraphExec_t的啟動将在以前啟動同一個可執行圖形之後進行。
圖的執行是在流中完成的，以便與其他異步工作一起排序。但是，流僅用于排序；它不限制圖的内部并行性，也不影響圖節點的執行位置。
請參見圖形API。
十. 事件
運作時還提供了一種方法，通過讓應用程式在程式中的任意點異步記錄事件并查詢這些事件何時完成，可以密切監視裝置的進度，并執行準确的計時。當事件之前的所有任務（或者可選地，給定流中的所有指令）都已完成時，事件即已完成。流0中的事件在所有流中的所有先前任務和指令完成後完成。
十一. 創造與銷毀
下面的代碼示例建立兩個事件：
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
它們是這樣被銷毀的：
cudaEventDestroy(start);
cudaEventDestroy(stop);
十二. 經過的時間
在建立和銷毀中建立的事件可用于按以下方式計時建立和銷毀的代碼示例：
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);
十三. 同步調用
調用同步函數時，在裝置完成請求的任務之前，不會将控件傳回到主機線程。在主機線程執行任何其他CUDA調用之前，可以通過使用某些特定标志（有關詳細資訊，請參閱參考手冊）調用cudaSetDeviceFlags（）來指定主機線程是否會産生、阻塞或旋轉。
十四. 多裝置系統
十六. 裝置辨別
主機系統可以有多個裝置。下面的代碼示例示範如何枚舉這些裝置、查詢它們的屬性以及确定啟用CUDA的裝置的數量。
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);
}
十七 同步調用
十八. 裝置選擇
主機線程可以通過調用cudaSetDevice（）随時設定其操作的裝置。在目前設定的裝置上進行裝置記憶體配置設定和核心啟動；流和事件與目前設定的裝置關聯建立。如果未調用cudastedevice（），則目前裝置為裝置0。
下面的代碼示例示範了設定目前裝置如何影響記憶體配置設定和核心執行。
十九. 多裝置系統
二十. 裝置辨別
主機系統可以有多個裝置。下面的代碼示例示範如何枚舉這些裝置、查詢它們的屬性以及确定啟用CUDA的裝置的數量。
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
二十一. 流和事件行為
如果将核心發送到與目前裝置無關的流，則核心啟動将失敗，如下面的代碼示例所示。
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
即使将記憶體副本發送到與目前裝置無關的流，它也會成功。
如果輸入事件和輸入流與不同的裝置關聯，則cudaEventRecord（）将失敗。
如果兩個輸入事件關聯到不同的裝置，則cudaEventLapsedTime（）将失敗。
即使輸入事件與不同于目前裝置的裝置關聯，cudaEventSynchronize（）和cudaEventQuery（）也将成功。
即使輸入流和輸入事件關聯到不同的裝置，cudaStreamWaitEvent（）也将成功。是以，可以使用cudaStreamWaitEvent（）來同步多個裝置。
每個裝置都有自己的預設流（請參閱預設流），是以，向裝置的預設流發出的指令可能會無序執行，或者與向任何其他裝置的預設流發出的指令同時執行。
二十二. 對等記憶體通路
根據系統屬性，特别是PCIe和/或NVLINK拓撲，裝置能夠尋址彼此的存儲器（即，在一個裝置上執行的核心可以解除對另一個裝置存儲器的指針的引用）。如果這兩個裝置的cudaDeviceCanAccessPeer（）傳回true，則在兩個裝置之間支援此對等記憶體通路功能。 對等記憶體通路僅在64位應用程式中受支援，必須通過調用cudaDeviceEnablePeerAccess（）在兩個裝置之間啟用，如下面的代碼示例所示。在非NVSwitch啟用的系統上，每個裝置最多可支援8個系統範圍的對等連接配接。
兩個裝置都使用統一的位址空間（請參閱統一虛拟位址空間），是以可以使用同一個指針對兩個裝置的記憶體進行尋址，如下面的代碼示例所示。
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);
二十三. Linux上的IOMMU
僅在Linux上，CUDA和顯示驅動程式不支援啟用IOMMU的裸機PCIe對等記憶體複制。但是，CUDA和顯示驅動程式确實通過虛拟機傳遞支援IOMMU。是以，Linux上的使用者在本機裸機系統上運作時，應該禁用IOMMU。應啟用IOMMU，并将VFIO驅動程式用作虛拟機的PCIe直通。
在Windows上不存在上述限制。
另請參閱在64位平台上配置設定DMA緩沖區。