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

// parameters for a basic allocation
cudaMemAllocNodeParams params = {};
params.poolProps.allocType = cudaMemAllocationTypePinned;
params.poolProps.location.type = cudaMemLocationTypeDevice;
// specify device 0 as the resident device
params.poolProps.location.id = 0;
params.bytesize = size;

cudaGraphAddMemAllocNode(&allocNode, graph, NULL, 0, ¶ms);
nodeParams->kernelParams[0] = params.dptr;
cudaGraphAddKernelNode(&a, graph, &allocNode, 1, &nodeParams);
cudaGraphAddKernelNode(&b, graph, &a, 1, &nodeParams);
cudaGraphAddKernelNode(&c, graph, &a, 1, &nodeParams);
cudaGraphNode_t dependencies[2];
// kernel nodes b and c are using the graph allocation, so the freeing node must depend on them.  Since the dependency of node b on node a establishes an indirect dependency, the free node does not need to explicitly depend on node a.
dependencies[0] = b;
dependencies[1] = c;
cudaGraphAddMemFreeNode(&freeNode, graph, dependencies, 2, params.dptr);
// free node does not depend on kernel node d, so it must not access the freed graph allocation.
cudaGraphAddKernelNode(&d, graph, &c, 1, &nodeParams);

// node e does not depend on the allocation node, so it must not access the allocation.  This would be true even if the freeNode depended on kernel node e.
cudaGraphAddKernelNode(&e, graph, NULL, 0, &nodeParams);

G.3.2. Stream Capture

可以通過捕獲相應(yīng)的流序分配和免費調(diào)用 cudaMallocAsync 和 cudaFreeAsync 來創(chuàng)建圖形內(nèi)存節(jié)點。 在這種情況下，捕獲的分配 API 返回的虛擬地址可以被圖中的其他操作使用。 由于流序的依賴關(guān)系將被捕獲到圖中，流序分配 API 的排序要求保證了圖內(nèi)存節(jié)點將根據(jù)捕獲的流操作正確排序（對于正確編寫的流代碼）。

忽略內(nèi)核節(jié)點 d 和 e，為清楚起見，以下代碼片段顯示了如何使用流捕獲來創(chuàng)建上圖中的圖形：

cudaMallocAsync(&dptr, size, stream1);
kernel_A<<< ..., stream1 >>>(dptr, ...);

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

kernel_B<<< ..., stream1 >>>(dptr, ...);
// event dependencies translated into graph dependencies, so the kernel node created by the capture of kernel C will depend on the allocation node created by capturing the cudaMallocAsync call. 
kernel_C<<< ..., stream2 >>>(dptr, ...);

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

// Free depends on all work accessing the memory.
cudaFreeAsync(dptr, stream1);

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

G.3.3. Accessing and Freeing Graph Memory Outside of the Allocating Graph

圖分配不必由分配圖釋放。當(dāng)圖不釋放分配時，該分配會在圖執(zhí)行之后持續(xù)存在，并且可以通過后續(xù) CUDA 操作訪問。這些分配可以在另一個圖中訪問或直接通過流操作訪問，只要訪問操作在分配之后通過 CUDA 事件和其他流排序機(jī)制進(jìn)行排序。隨后可以通過定期調(diào)用 cudaFree、cudaFreeAsync 或通過啟動具有相應(yīng)空閑節(jié)點的另一個圖，或隨后啟動分配圖（如果它是使用 cudaGraphInstantiateFlagAutoFreeOnLaunch 標(biāo)志實例化）來釋放分配。在內(nèi)存被釋放后訪問內(nèi)存是非法的 – 必須在所有使用圖依賴、CUDA 事件和其他流排序機(jī)制訪問內(nèi)存的操作之后對釋放操作進(jìn)行排序。

注意：因為圖分配可能彼此共享底層物理內(nèi)存，所以必須考慮與一致性和一致性相關(guān)的虛擬混疊支持規(guī)則。簡單地說，空閑操作必須在完整的設(shè)備操作（例如，計算內(nèi)核/ memcpy）完成后排序。具體來說，帶外同步——例如，作為訪問圖形內(nèi)存的計算內(nèi)核的一部分，通過內(nèi)存進(jìn)行信號交換——不足以提供對圖形內(nèi)存的寫操作和該圖形內(nèi)存的自由操作之間的排序保證。

以下代碼片段演示了在分配圖之外訪問圖分配，并通過以下方式正確建立順序：使用單個流，使用流之間的事件，以及使用嵌入到分配和釋放圖中的事件。

使用單個流建立的排序：

void *dptr;
cudaGraphAddMemAllocNode(&allocNode, allocGraph, NULL, 0, ¶ms);
dptr = params.dptr;

cudaGraphInstantiate(&allocGraphExec, allocGraph, NULL, NULL, 0);

cudaGraphLaunch(allocGraphExec, stream);
kernel<<< …, stream >>>(dptr, …);
cudaFreeAsync(dptr, stream);

通過記錄和等待 CUDA 事件建立的排序：

void *dptr;

// Contents of allocating graph
cudaGraphAddMemAllocNode(&allocNode, allocGraph, NULL, 0, ¶ms);
dptr = params.dptr;

// contents of consuming/freeing graph
nodeParams->kernelParams[0] = params.dptr;
cudaGraphAddKernelNode(&a, graph, NULL, 0, &nodeParams);
cudaGraphAddMemFreeNode(&freeNode, freeGraph, &a, 1, dptr);

cudaGraphInstantiate(&allocGraphExec, allocGraph, NULL, NULL, 0);
cudaGraphInstantiate(&freeGraphExec, freeGraph, NULL, NULL, 0);

cudaGraphLaunch(allocGraphExec, allocStream);

// establish the dependency of stream2 on the allocation node
// note: the dependency could also have been established with a stream synchronize operation
cudaEventRecord(allocEvent, allocStream)
cudaStreamWaitEvent(stream2, allocEvent);

kernel<<< …, stream2 >>> (dptr, …);

// establish the dependency between the stream 3 and the allocation use
cudaStreamRecordEvent(streamUseDoneEvent, stream2);
cudaStreamWaitEvent(stream3, streamUseDoneEvent);

// it is now safe to launch the freeing graph, which may also access the memory
cudaGraphLaunch(freeGraphExec, stream3);

使用圖外部事件節(jié)點建立的排序：

void *dptr;
cudaEvent_t allocEvent; // event indicating when the allocation will be ready for use.
cudaEvent_t streamUseDoneEvent; // event indicating when the stream operations are done with the allocation.

// Contents of allocating graph with event record node
cudaGraphAddMemAllocNode(&allocNode, allocGraph, NULL, 0, ¶ms);
dptr = params.dptr;
// note: this event record node depends on the alloc node
cudaGraphAddEventRecordNode(&recordNode, allocGraph, &allocNode, 1, allocEvent);
cudaGraphInstantiate(&allocGraphExec, allocGraph, NULL, NULL, 0);

// contents of consuming/freeing graph with event wait nodes
cudaGraphAddEventWaitNode(&streamUseDoneEventNode, waitAndFreeGraph, NULL, 0, streamUseDoneEvent);
cudaGraphAddEventWaitNode(&allocReadyEventNode, waitAndFreeGraph, NULL, 0, allocEvent);
nodeParams->kernelParams[0] = params.dptr;

// The allocReadyEventNode provides ordering with the alloc node for use in a consuming graph.
cudaGraphAddKernelNode(&kernelNode, waitAndFreeGraph, &allocReadyEventNode, 1, &nodeParams);

// The free node has to be ordered after both external and internal users.
// Thus the node must depend on both the kernelNode and the 
// streamUseDoneEventNode.
dependencies[0] = kernelNode;
dependencies[1] = streamUseDoneEventNode;
cudaGraphAddMemFreeNode(&freeNode, waitAndFreeGraph, &dependencies, 2, dptr);
cudaGraphInstantiate(&waitAndFreeGraphExec, waitAndFreeGraph, NULL, NULL, 0);

cudaGraphLaunch(allocGraphExec, allocStream);

// establish the dependency of stream2 on the event node satisfies the ordering requirement
cudaStreamWaitEvent(stream2, allocEvent);
kernel<<< …, stream2 >>> (dptr, …);
cudaStreamRecordEvent(streamUseDoneEvent, stream2);

// the event wait node in the waitAndFreeGraphExec establishes the dependency on the “readyForFreeEvent” that is needed to prevent the kernel running in stream two from accessing the allocation after the free node in execution order.
cudaGraphLaunch(waitAndFreeGraphExec, stream3);

G.3.4. cudaGraphInstantiateFlagAutoFreeOnLaunch

在正常情況下，如果圖有未釋放的內(nèi)存分配，CUDA 將阻止重新啟動圖，因為同一地址的多個分配會泄漏內(nèi)存。使用 cudaGraphInstantiateFlagAutoFreeOnLaunch 標(biāo)志實例化圖允許圖在其仍有未釋放的分配時重新啟動。在這種情況下，啟動會自動插入一個異步釋放的未釋放分配。

啟動時自動對于單生產(chǎn)者多消費者算法很有用。在每次迭代中，生產(chǎn)者圖創(chuàng)建多個分配，并且根據(jù)運(yùn)行時條件，一組不同的消費者訪問這些分配。這種類型的變量執(zhí)行序列意味著消費者無法釋放分配，因為后續(xù)消費者可能需要訪問。啟動時自動釋放意味著啟動循環(huán)不需要跟蹤生產(chǎn)者的分配 – 相反，該信息與生產(chǎn)者的創(chuàng)建和銷毀邏輯保持隔離。通常，啟動時自動釋放簡化了算法，否則該算法需要在每次重新啟動之前釋放圖所擁有的所有分配。

注意： cudaGraphInstantiateFlagAutoFreeOnLaunch 標(biāo)志不會改變圖銷毀的行為。應(yīng)用程序必須顯式釋放未釋放的內(nèi)存以避免內(nèi)存泄漏，即使對于使用標(biāo)志實例化的圖也是如此。

以下代碼展示了使用 cudaGraphInstantiateFlagAutoFreeOnLaunch 來簡化單生產(chǎn)者/多消費者算法：

// Create producer graph which allocates memory and populates it with data
cudaStreamBeginCapture(cudaStreamPerThread, cudaStreamCaptureModeGlobal);
cudaMallocAsync(&data1, blocks * threads, cudaStreamPerThread);
cudaMallocAsync(&data2, blocks * threads, cudaStreamPerThread);
produce<<>>(data1, data2);
...
cudaStreamEndCapture(cudaStreamPerThread, &graph);
cudaGraphInstantiateWithFlags(&producer,
                              graph,
                              cudaGraphInstantiateFlagAutoFreeOnLaunch);
cudaGraphDestroy(graph);

// Create first consumer graph by capturing an asynchronous library call
cudaStreamBeginCapture(cudaStreamPerThread, cudaStreamCaptureModeGlobal);
consumerFromLibrary(data1, cudaStreamPerThread);
cudaStreamEndCapture(cudaStreamPerThread, &graph);
cudaGraphInstantiateWithFlags(&consumer1, graph, 0); //regular instantiation
cudaGraphDestroy(graph);

// Create second consumer graph
cudaStreamBeginCapture(cudaStreamPerThread, cudaStreamCaptureModeGlobal);
consume2<<>>(data2);
...
cudaStreamEndCapture(cudaStreamPerThread, &graph);
cudaGraphInstantiateWithFlags(&consumer2, graph, 0);
cudaGraphDestroy(graph);

// Launch in a loop
bool launchConsumer2 = false;
do {
    cudaGraphLaunch(producer, myStream);
    cudaGraphLaunch(consumer1, myStream);
    if (launchConsumer2) {
        cudaGraphLaunch(consumer2, myStream);
    }
} while (determineAction(&launchConsumer2));

cudaFreeAsync(data1, myStream);
cudaFreeAsync(data2, myStream);

cudaGraphExecDestroy(producer);
cudaGraphExecDestroy(consumer1);
cudaGraphExecDestroy(consumer2);

G.4. Optimized Memory Reuse

CUDA 以兩種方式重用內(nèi)存：

圖中的虛擬和物理內(nèi)存重用基于虛擬地址分配，就像在流序分配器中一樣。

圖之間的物理內(nèi)存重用是通過虛擬別名完成的：不同的圖可以將相同的物理內(nèi)存映射到它們唯一的虛擬地址。

G.4.1. Address Reuse within a Graph

CUDA 可以通過將相同的虛擬地址范圍分配給生命周期不重疊的不同分配來重用圖中的內(nèi)存。 由于可以重用虛擬地址，因此不能保證指向具有不相交生命周期的不同分配的指針是唯一的。

下圖顯示了添加一個新的分配節(jié)點 （2），它可以重用依賴節(jié)點 （1） 釋放的地址。

下圖顯示了添加新的 alloc 節(jié)點（3）。 新的分配節(jié)點不依賴于空閑節(jié)點 （2），因此不能重用來自關(guān)聯(lián)分配節(jié)點 （2） 的地址。 如果分配節(jié)點 （2） 使用由空閑節(jié)點 （1） 釋放的地址，則新分配節(jié)點 3 將需要一個新地址。

G.4.2. Physical Memory Management and Sharing

CUDA 負(fù)責(zé)在按 GPU 順序到達(dá)分配節(jié)點之前將物理內(nèi)存映射到虛擬地址。作為內(nèi)存占用和映射開銷的優(yōu)化，如果多個圖不會同時運(yùn)行，它們可能會使用相同的物理內(nèi)存進(jìn)行不同的分配，但是如果它們同時綁定到多個執(zhí)行圖，則物理頁面不能被重用，或未釋放的圖形分配。

CUDA 可以在圖形實例化、啟動或執(zhí)行期間隨時更新物理內(nèi)存映射。 CUDA 還可以在未來的圖啟動之間引入同步，以防止實時圖分配引用相同的物理內(nèi)存。對于任何 allocate-free-allocate 模式，如果程序在分配的生命周期之外訪問指針，錯誤的訪問可能會默默地讀取或?qū)懭肓硪粋€分配擁有的實時數(shù)據(jù)（即使分配的虛擬地址是唯一的）。使用計算清理工具可以捕獲此錯誤。

下圖顯示了在同一流中按順序啟動的圖形。在此示例中，每個圖都會釋放它分配的所有內(nèi)存。由于同一流中的圖永遠(yuǎn)不會同時運(yùn)行，CUDA 可以而且應(yīng)該使用相同的物理內(nèi)存來滿足所有分配。

G.5. Performance Considerations

當(dāng)多個圖啟動到同一個流中時，CUDA 會嘗試為它們分配相同的物理內(nèi)存，因為這些圖的執(zhí)行不能重疊。 在啟動之間保留圖形的物理映射作為優(yōu)化以避免重新映射的成本。 如果稍后啟動其中一個圖，使其執(zhí)行可能與其他圖重疊（例如，如果它啟動到不同的流中），則 CUDA 必須執(zhí)行一些重新映射，因為并發(fā)圖需要不同的內(nèi)存以避免數(shù)據(jù)損壞 。

一般來說，CUDA中圖內(nèi)存的重新映射很可能是由這些操作引起的

更改啟動圖形的流

圖內(nèi)存池上的修剪操作，顯式釋放未使用的內(nèi)存（在物理內(nèi)存占用中討論）

當(dāng)另一個圖的未釋放分配映射到同一內(nèi)存時重新啟動一個圖將導(dǎo)致在重新啟動之前重新映射內(nèi)存

重新映射必須按執(zhí)行順序發(fā)生，但在該圖的任何先前執(zhí)行完成之后（否則可能會取消映射仍在使用的內(nèi)存）。 由于這種排序依賴性，以及映射操作是操作系統(tǒng)調(diào)用，映射操作可能相對昂貴。 應(yīng)用程序可以通過將包含分配內(nèi)存節(jié)點的圖一致地啟動到同一流中來避免這種成本。

G.5.1. First Launch / cudaGraphUpload

在圖實例化期間無法分配或映射物理內(nèi)存，因為圖將在其中執(zhí)行的流是未知的。 映射是在圖形啟動期間完成的。 調(diào)用 cudaGraphUpload 可以通過立即執(zhí)行該圖的所有映射并將該圖與上傳流相關(guān)聯(lián)，將分配成本與啟動分開。 如果圖隨后啟動到同一流中，它將啟動而無需任何額外的重新映射。

使用不同的流進(jìn)行圖上傳和圖啟動的行為類似于切換流，可能會導(dǎo)致重新映射操作。 此外，允許無關(guān)的內(nèi)存池管理從空閑流中提取內(nèi)存，這可能會抵消上傳的影響。

G.6. Physical Memory Footprint

異步分配的池管理行為意味著銷毀包含內(nèi)存節(jié)點的圖（即使它們的分配是空閑的）不會立即將物理內(nèi)存返回給操作系統(tǒng)以供其他進(jìn)程使用。要顯式將內(nèi)存釋放回操作系統(tǒng)，應(yīng)用程序應(yīng)使用 cudaDeviceGraphMemTrim API。

cudaDeviceGraphMemTrim 將取消映射并釋放由圖形內(nèi)存節(jié)點保留的未主動使用的任何物理內(nèi)存。尚未釋放的分配和計劃或運(yùn)行的圖被認(rèn)為正在積極使用物理內(nèi)存，不會受到影響。使用修剪 API 將使物理內(nèi)存可用于其他分配 API 和其他應(yīng)用程序或進(jìn)程，但會導(dǎo)致 CUDA 在下次啟動修剪圖時重新分配和重新映射內(nèi)存。請注意，cudaDeviceGraphMemTrim 在與 cudaMemPoolTrimTo（） 不同的池上運(yùn)行。圖形內(nèi)存池不會暴露給流序內(nèi)存分配器。 CUDA 允許應(yīng)用程序通過 cudaDeviceGetGraphMemAttribute API 查詢其圖形內(nèi)存占用量。查詢屬性 cudaGraphMemAttrReservedMemCurrent 返回驅(qū)動程序為當(dāng)前進(jìn)程中的圖形分配保留的物理內(nèi)存量。查詢 cudaGraphMemAttrUsedMemCurrent 返回至少一個圖當(dāng)前映射的物理內(nèi)存量。這些屬性中的任何一個都可用于跟蹤 CUDA 何時為分配圖而獲取新的物理內(nèi)存。這兩個屬性對于檢查共享機(jī)制節(jié)省了多少內(nèi)存都很有用。

G.7. Peer Access

圖分配可以配置為從多個 GPU 訪問，在這種情況下，CUDA 將根據(jù)需要將分配映射到對等 GPU。 CUDA 允許需要不同映射的圖分配重用相同的虛擬地址。 發(fā)生這種情況時，地址范圍將映射到不同分配所需的所有 GPU。 這意味著分配有時可能允許比其創(chuàng)建期間請求的更多對等訪問； 然而，依賴這些額外的映射仍然是一個錯誤。

G.7.1. Peer Access with Graph Node APIs

cudaGraphAddMemAllocNode API 接受節(jié)點參數(shù)結(jié)構(gòu)的 accessDescs 數(shù)組字段中的映射請求。 poolProps.location 嵌入式結(jié)構(gòu)指定分配的常駐設(shè)備。 假設(shè)需要來自分配 GPU 的訪問，因此應(yīng)用程序不需要在 accessDescs 數(shù)組中為常駐設(shè)備指定條目。

cudaMemAllocNodeParams params = {};
params.poolProps.allocType = cudaMemAllocationTypePinned;
params.poolProps.location.type = cudaMemLocationTypeDevice;
// specify device 1 as the resident device
params.poolProps.location.id = 1;
params.bytesize = size;

// allocate an allocation resident on device 1 accessible from device 1
cudaGraphAddMemAllocNode(&allocNode, graph, NULL, 0, ¶ms);

accessDescs[2];
// boilerplate for the access descs (only ReadWrite and Device access supported by the add node api)
accessDescs[0].flags = cudaMemAccessFlagsProtReadWrite;
accessDescs[0].location.type = cudaMemLocationTypeDevice;
accessDescs[1].flags = cudaMemAccessFlagsProtReadWrite;
accessDescs[1].location.type = cudaMemLocationTypeDevice;

// access being requested for device 0 & 2.  Device 1 access requirement left implicit.
accessDescs[0].location.id = 0;
accessDescs[1].location.id = 2;

// access request array has 2 entries.
params.accessDescCount = 2;
params.accessDescs = accessDescs;

// allocate an allocation resident on device 1 accessible from devices 0, 1 and 2. (0 & 2 from the descriptors, 1 from it being the resident device).
cudaGraphAddMemAllocNode(&allocNode, graph, NULL, 0, ¶ms);

G.7.2. Peer Access with Stream Capture

對于流捕獲，分配節(jié)點在捕獲時記錄分配池的對等可訪問性。 在捕獲 cudaMallocFromPoolAsync 調(diào)用后更改分配池的對等可訪問性不會影響圖將為分配進(jìn)行的映射。

// boilerplate for the access descs (only ReadWrite and Device access supported by the add node api)
accessDesc.flags = cudaMemAccessFlagsProtReadWrite;
accessDesc.location.type = cudaMemLocationTypeDevice;
accessDesc.location.id = 1;

// let memPool be resident and accessible on device 0

cudaStreamBeginCapture(stream);
cudaMallocAsync(&dptr1, size, memPool, stream);
cudaStreamEndCapture(stream, &graph1);

cudaMemPoolSetAccess(memPool, &accessDesc, 1);

cudaStreamBeginCapture(stream);
cudaMallocAsync(&dptr2, size, memPool, stream);
cudaStreamEndCapture(stream, &graph2);

//The graph node allocating dptr1 would only have the device 0 accessibility even though memPool now has device 1 accessibility.
//The graph node allocating dptr2 will have device 0 and device 1 accessibility, since that was the pool accessibility at the time of the cudaMallocAsync call.