允許圖創(chuàng)建和擁有內(nèi)存分配功能的圖內(nèi)存節(jié)點(diǎn)

// 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

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

忽略內(nèi)核節(jié)點(diǎn) d 和 e，為清楚起見，以下代碼片段顯示了如何使用流捕獲來(lái)創(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)圖不釋放分配時(shí)，該分配會(huì)在圖執(zhí)行之后持續(xù)存在，并且可以通過(guò)后續(xù) CUDA 操作訪問(wèn)。這些分配可以在另一個(gè)圖中訪問(wèn)或直接通過(guò)流操作訪問(wèn)，只要訪問(wèn)操作在分配之后通過(guò) CUDA 事件和其他流排序機(jī)制進(jìn)行排序。隨后可以通過(guò)定期調(diào)用 cudaFree、cudaFreeAsync 或通過(guò)啟動(dòng)具有相應(yīng)空閑節(jié)點(diǎn)的另一個(gè)圖，或隨后啟動(dòng)分配圖（如果它是使用 cudaGraphInstantiateFlagAutoFreeOnLaunch 標(biāo)志實(shí)例化）來(lái)釋放分配。在內(nèi)存被釋放后訪問(wèn)內(nèi)存是非法的 – 必須在所有使用圖依賴、CUDA 事件和其他流排序機(jī)制訪問(wèn)內(nèi)存的操作之后對(duì)釋放操作進(jìn)行排序。

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

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

使用單個(gè)流建立的排序：

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

通過(guò)記錄和等待 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é)點(diǎn)建立的排序：

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 將阻止重新啟動(dòng)圖，因?yàn)橥坏刂返亩鄠€(gè)分配會(huì)泄漏內(nèi)存。使用 cudaGraphInstantiateFlagAutoFreeOnLaunch 標(biāo)志實(shí)例化圖允許圖在其仍有未釋放的分配時(shí)重新啟動(dòng)。在這種情況下，啟動(dòng)會(huì)自動(dòng)插入一個(gè)異步釋放的未釋放分配。

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

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

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

// 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)存重用是通過(guò)虛擬別名完成的：不同的圖可以將相同的物理內(nèi)存映射到它們唯一的虛擬地址。

G.4.1. Address Reuse within a Graph

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

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

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

G.4.2. Physical Memory Management and Sharing

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

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

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

G.5. Performance Considerations

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

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

更改啟動(dòng)圖形的流

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

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

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

G.5.1. First Launch / cudaGraphUpload

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

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

G.6. Physical Memory Footprint

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

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

G.7. Peer Access

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

G.7.1. Peer Access with Graph Node APIs

cudaGraphAddMemAllocNode API 接受節(jié)點(diǎn)參數(shù)結(jié)構(gòu)的 accessDescs 數(shù)組字段中的映射請(qǐng)求。 poolProps.location 嵌入式結(jié)構(gòu)指定分配的常駐設(shè)備。 假設(shè)需要來(lái)自分配 GPU 的訪問(wèn)，因此應(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

對(duì)于流捕獲，分配節(jié)點(diǎn)在捕獲時(shí)記錄分配池的對(duì)等可訪問(wèn)性。 在捕獲 cudaMallocFromPoolAsync 調(diào)用后更改分配池的對(duì)等可訪問(wèn)性不會(huì)影響圖將為分配進(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.