CUDA运行时 Runtime（四）
CUDA運行時 Runtime（四）
一. 圖
圖為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緩沖區。
總結
以上是生活随笔為你收集整理的CUDA运行时 Runtime（四）的全部內容，希望文章能夠幫你解決所遇到的問題。
- 上一篇： CUDA运行时Runtime（三）
- 下一篇： 光刻机技术解析