CUDA C++ Programming Guide——编程接口 CUDA Runtime Asynchronous Concurrent Execution

淩亂°似流年 2022-12-17 12:59 214阅读 0赞

### Asynchronous Concurrent Execution ###

CUDA将以下操作公开为可以彼此并发运行的独立任务：

*  主机上的计算；
 *  设备上的计算；
 *  内存从主机传输到设备；
 *  内存从设备传输到主机；
 *  内存在给定设备的内存中传输；
 *  设备之间的内存传输。  
    这些操作之间实现的并发级别将取决于设备的功能集和计算能力，如下所述。

#### Concurrent Execution between Host and Device ####

异步库函数可促进并发主机执行(Concurrent host execution is facilitated through asynchronous library functions)，这些函数在设备完成请求的任务之前将控制权返回给主机线程。 使用异步调用，可以在适当的设备资源可用时将许多设备操作排在一起，由CUDA驱动程序执行(many device operations can be queued up together to be executed by the CUDA driver when appropriate device resources are available)。 这减轻了主机线程管理设备的大部分责任，使它可以自由执行其他任务。 以下设备操作相对于主机是异步的：

*  内核启动；
 *  在单个设备的内存中进行内存复制；
 *  从主机到设备的内存复制，大小为64 KB或更少；
 *  由带有Async后缀的函数执行的内存复制(Memory copies performed by functions that are suffixed with Async)；
 *  内存设置函数调用。  
    程序员可以通过将CUDA\_LAUNCH\_BLOCKING环境变量设置为1来全局禁用系统上运行的所有CUDA应用程序的内核启动异步性。此功能仅用于调试目的，不应用作使生产软件可靠运行。如果通过分析器（Nsight，Visual Profiler）收集硬件计数器hardware counters，则内核启动是同步的，除非启用了并发内核分析。 如果异步内存复制涉及未页面锁定的主机内存not page-locked，则它们也将是同步的。

#### Concurrent Kernel Execution ####

一些具有2.x和更高版本的计算能力的设备可以同时执行多个内核。 应用程序可以通过检查并发内核设备属性（请参阅[设备枚举][Link 1]）来查询此功能，对于支持该功能的设备，该属性等于1。设备可以同时执行的最大内核启动次数取决于其计算能力(The maximum number of kernel launches that a device can execute concurrently)，并在表[15][]中列出。**一个CUDA上下文中的内核无法与另一CUDA上下文中的内核同时执行**。使用许多纹理或大量本地内存的内核不太可能与其他内核同时执行。

#### Overlap of Data Transfer and Kernel Execution ####

一些设备可以在内核执行的同时向GPU或从GPU执行异步内存复制。 应用程序可以通过检查asyncEngineCount设备属性（请参阅[设备枚举][Link 1]）来查询此功能，对于支持该属性的设备，该属性大于零。 如果复制操作中包含主机内存(host memory is involved in the copy)，则必须对其进行页面锁定。

设备内拷贝(intra-device copy)可以与内核执行同时执行（在支持并发内核设备属性的设备上）和/或与设备之间复制同时执行（对于支持asyncEngineCount属性的设备）。 使用标准内存复制功能（目标地址和源地址位于同一设备上）启动设备内复制。(Intra-device copies are initiated using the standard memory copy functions with destination and source addresses residing on the same device.)

#### Concurrent Data Transfers ####

某些具有2.x和更高版本的计算能力的设备可以使设备之间的复制重叠(overlap copies to and from the device)。 应用程序可以通过检查asyncEngineCount设备属性（请参阅[设备枚举][Link 1]）来查询此功能，对于支持该属性的设备，该属性等于2。 为了重叠，传输中涉及的所有主机内存都必须页锁定。

#### Streams ####

应用程序通过流管理上述并行操作。 流是按顺序执行的一系列命令 A stream is a sequence of commands（可能由不同的主机线程发出）。 另一方面，不同的流可能会相对于彼此无序或同时执行。 这种行为无法得到保证，因此不应依赖于它的正确性（例如，内核间通信未定义）。 当满足命令的所有依赖性时，可以执行在流上发出的命令。 依赖关系可以是先前在同一流上启动的命令，也可以是来自其他流的依赖关系。 同步命令调用成功保证了所有启动的命令均已完成(The successful completion of synchronize call guarantees that all the commands launched are completed)。

##### Creation and Destruction #####

通过创建流对象并将其指定为一系列内核启动和主机和设备之间的内存复制的流参数来定义流。 下面的代码示例创建两个流，并在页面锁定的内存中分配一个float数组hostPtr。

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

以下代码示例将这些流中的每一个定义为从主机到设备的一个内存复制，内核启动以及从设备到主机的一个内存复制的序列(Each of these streams is defined by the following code sample as a sequence of one memory copy from host to device, one kernel launch, and one memory copy from device to host)：

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

每个流将其输入数组hostPtr的一部分复制到设备内存中的数组inputDevPtr，通过调用MyKernel处理设备上的inputDevPtr，并将结果outputDevPtr复制回到hostPtr的相同部分。 [重叠][Link 2]overlap行为描述了此示例中的流如何根据设备的功能重叠。 请注意，hostPtr必须指向页面锁定的主机内存，以便发生任何重叠。通过调用cudaStreamDestroy()释放流。

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

如果在调用cudaStreamDestroy()时设备仍在流中工作，则该函数将立即返回，并且一旦设备完成了流中的所有工作，与流相关的资源就会自动释放。

##### Default Stream #####

未指定任何流参数或等效地将流参数设置为零的启动内核和主机设备间内存复制将生成默认流。因此，它们按顺序执行。对于使用–default-stream per-thread编译标志（或在包含CUDA标头（cuda.h和cuda\_runtime.h）之前定义CUDA\_API\_PER\_THREAD\_DEFAULT\_STREAM宏）编译的代码，默认流是常规流，每个主机线程有自己的默认流。  
注意：当代码由nvcc编译时，`＃define CUDA_API_PER_THREAD_DEFAULT_STREAM 1`不能用于启用此行为，因为nvcc隐含在翻译单元顶部的cuda\_runtime.h。在这种情况下，需要使用–default-stream per-thread编译标志，或者需要使用-DCUDA\_API\_PER\_THREAD\_DEFAULT\_STREAM = 1编译器标志定义CUDA\_API\_PER\_THREAD\_DEFAULT\_STREAM宏。  
对于使用–default-stream旧式编译标志编译的代码，默认流被称为NULL流的特殊流，并且每个设备都有一个用于所有主机线程的NULL流。 NULL流是特殊的，因为它会导致隐式同步，如“[隐式同步][Link 3]”中所述。对于未指定–default-stream编译标志进行编译的代码，默认使用–default-stream旧版。

##### Explicit Synchronization #####

有多种方法可以使流彼此显式同步。cudaDeviceSynchronize()等待所有主机线程所有流中的所有先前命令完成。cudaStreamSynchronize()将流作为参数，并等待直到给定流中的所有先前命令都完成为止。 它可用于将主机与特定流同步，从而允许其他流继续在设备上执行。  
cudaStreamWaitEvent()将流和事件作为参数（有关事件的描述，请参见[事件][Link 4]），并使对cudaStreamWaitEvent()的调用之后添加到给定流中的所有命令延迟执行，直到给定事件完成为止。  
cudaStreamQuery()为应用程序提供了一种方法，该方法可以知道流中的所有先前命令是否已完成。

##### Implicit Synchronization #####

如果主机线程在它们之间发出以下任何操作，则来自不同流的两个命令不能同时运行(Two commands from different streams cannot run concurrently if any one of the following operations is issued in-between them by the host thread)：

*  页面锁定的主机内存分配，
 *  设备内存分配，
 *  设备存储集a device memory set，
 *  两个地址之间到同一设备内存的内存复制，
 *  到NULL流的任何CUDA命令，
 *  在“计算能力3.x”和“计算能力7.x”中描述的L1 /共享内存配置之间进行切换。

对于支持并发内核执行且计算能力为3.0或更低版本的设备，需要进行依赖性检查以查看流式内核启动是否完成的任何操作：  
仅当从CUDA上下文中的任何流启动的所有先前内核启动的所有线程块都已开始执行时才可以开始执行；  
从CUDA上下文中的任何流中阻止所有以后的内核启动，直到检查完内核启动为止。

需要依赖检查的操作包括与正在检查的启动相同的流中的任何其他命令，以及对该流上对cudaStreamQuery()的任何调用。因此，应用程序应遵循以下准则，以提高其并发执行内核的潜力：所有独立操作应在相关操作之前发布，任何形式的同步都应尽可能延迟。

##### Overlapping Behavior #####

两个流之间的执行重叠量取决于向每个流发出命令的顺序，以及设备是否支持数据传输和内核执行的重叠（请参见[数据传输和内核执行的重叠][Link 5]），并发内核执行（ 请参阅[并发内核执行][Link 6]）和/或并发数据传输（请参见[并发数据传输][Link 7]）。

例如，在不支持并发数据传输的设备上，[创建和销毁代码示例][Link 8]的两个流完全不重叠，因为从主机到设备的内存复制被发布到到流\[1\]在从设备进行内存复制被发布给stream \[0\]之后，因此只有在从设备到发给stream \[0\]的主机的内存复制完成后，它才能启动。 如果按照以下方式重写代码（并假设设备支持数据传输和内核执行的重叠）

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\]的设备的内存副本与发布到stream \[0\]的内核启动重叠。

在支持并发数据传输的设备上，[创建和销毁代码示例][Link 8]的两个流确实重叠：从主机到设备的内存复制（发布到stream \[1\]）与从设备到主机的内存复制发布到stream \[0\] \]重叠，甚至将内核启动发布到stream \[0\]重叠（假设设备支持数据传输和内核执行的重叠）。但是，对于计算能力为3.0或更低的设备，内核执行可能不会重叠，因为在从设备到主机的内存副本发布到stream \[0\]之后，第二次内核启动发布到stream \[1\]，因此它被阻塞直到根据隐式同步，发布到stream \[0\]的第一个内核启动已完成。如果代码按上述方式重写，则内核执行会重叠（假设设备支持并发内核执行），因为在从设备到主机的内存副本发布到stream \[0\]之前，第二次内核启动已发布到stream \[1\]。但是，在这种情况下，根据隐式同步，从设备到主机的发布到stream \[0\]的内存复制仅与发布到stream \[1\]的内核启动的最后线程块重叠，这仅占总数的一小部分内核的执行时间。

##### Host Functions (Callbacks) #####

运行时提供了一种通过cudaLaunchHostFunc()将CPU函数调用随时插入流中的方法。(The runtime provides a way to insert a CPU function call at any point into a stream via cudaLaunchHostFunc()) 一旦在回调之前向流发出的所有命令都已完成，就会在主机上执行提供的功能。

以下代码示例在向每个流发布主机到设备的内存复制，内核启动和设备到主机的内存复制之后，将主机函数MyCallback添加到两个流中的每个流中。 *每个设备到主机的内存复制完成后，该功能将在主机上开始执行*。

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

主机功能添加之后在流中发出的命令在功能完成之前不会开始执行。排队到流中的主机函数不得（直接或间接）进行CUDA API调用，因为如果进行此类调用导致死锁，它可能最终会等待自身。

##### Stream Priorities #####

可以在创建时使用cudaStreamCreateWithPriority()指定流的相对优先级。 可以使用cudaDeviceGetStreamPriorityRange()函数获得以\[最高优先级，最低优先级\]排序的允许优先级范围。 在运行时，优先级较高的流中的挂起工作优先于优先级较低的流中的挂起工作（ pending work in higher-priority streams takes preference over pending work in low-priority streams）。

下面的代码示例获取当前设备允许的优先级范围，并创建具有最高和最低可用优先级的流。

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

#### CUDA Graphs ####

CUDA图形(CUDA Graphs)为CUDA中的工作提交提供了一种新模型。图是由依赖关系连接的一系列操作（例如内核启动），它们是与执行分开定义的。这样就可以一次定义一个图，然后重复启动。将图形的定义与图形的执行分开可以进行许多优化：首先，与流相比，CPU的启动成本降低了，因为许多设置是预先完成的；其次，将整个工作流程呈现给CUDA可以实现优化，而这些优化可能是流的分段工作提交机制无法实现的。

要查看使用图形进行的优化，请先考虑流中发生的情况：将内核放入流中时，主机驱动程序将执行一系列操作，以准备在GPU上执行内核。这些操作是设置和启动内核所必需的，是必须为发出的每个内核支付的间接费用。对于执行时间短的GPU内核，此开销成本可能占整个端到端执行时间的很大一部分。

使用图提交工作分为三个不同的阶段：定义，实例化和执行。

*  在定义阶段，程序会在图中创建操作的描述以及它们之间的依赖关系。
 *  实例化对图形模板进行快照(Instantiation takes a snapshot of the graph template)，对其进行验证，并执行大部分工作的设置和初始化，以最大程度地减少启动时需要完成的工作。 生成的实例称为可执行图。
 *  可执行图可以启动到流中，类似于任何其他CUDA工作。 它可以启动多次，而无需重复实例化。

##### Graph Structure #####

一个操作在图中形成一个节点。 操作之间的依赖关系是边缘。 这些依赖性限制了操作的执行顺序。一旦操作所依赖的节点完成，就可以在任何时间调度该操作。 调度由CUDA系统负责。  
An operation forms a node in a graph. The dependencies between the operations are the edges. These dependencies constrain the execution sequence of the operations. An operation may be scheduled at any time once the nodes on which it depends are complete. Scheduling is left up to the CUDA system.

###### Node Types ######

图节点可以是以下之一：

*  kernel
 *  CPU function call
 *  memory copy
 *  memset
 *  empty node
 *  child graph: To execute a separate nested graph. See Figure 9.  
    ![Figure 9. Child Graph Example][]

##### Creating a Graph Using Graph APIs #####

可以通过两种机制创建图形：显式API和流捕获。 以下是创建和执行以下图形的示例。  
![Figure 10. Creating a Graph Using Graph APIs Example][]

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

##### Creating a Graph Using Stream Capture #####

流捕获提供了一种从现有的基于流的API创建图形的机制。 可将一部分工作启动到流中的代码（包括现有代码）用cudaStreamBeginCapture()和cudaStreamEndCapture()的调用括起来。 见下文。

cudaGraph_t graph;
    
    cudaStreamBeginCapture(stream);
    
    kernel_A<<< ..., stream >>>(...);
    kernel_B<<< ..., stream >>>(...);
    libraryCall(stream);
    kernel_C<<< ..., stream >>>(...);
    
    cudaStreamEndCapture(stream, &graph);

调用cudaStreamBeginCapture()会将流置于捕获模式。 捕获流时，启动该流的工作不会排队等待执行。 而是将其附加到逐渐建立的内部图上。 然后，通过调用cudaStreamEndCapture()返回此图，这也将结束流的捕获模式。 通过流捕获主动构造的图称为捕获图。

除cudaStreamLegacy（“ NULL流”）外，流捕获可用于任何CUDA流。 请注意，它可以在cudaStreamPerThread上使用。 如果程序正在使用旧式流legacy stream，则可以在不进行功能更改的情况下将流0重新定义为per-thread stream。 请参阅[默认流][Link 9]。可以使用cudaStreamIsCapturing()查询是否正在捕获流。

###### Cross-stream Dependencies and Events ######

流捕获可以处理用cudaEventRecord()和cudaStreamWaitEvent()表示的跨流依赖性，前提是要将等待的事件记录到同一捕获图中。当事件记录在处于捕获模式的流中时，它将导致捕获事件。捕获的事件表示捕获图中的一组节点。When an event is recorded in a stream that is in capture mode, it results in a captured event. A captured event represents a set of nodes in a capture graph.

当流等待捕获的事件时，如果事件还未发生则将流置于捕获模式，则流中的下一项将对捕获的事件中的节点具有其他依赖性。然后将两个流捕获到同一捕获图。When a captured event is waited on by a stream, it places the stream in capture mode if it is not already, and the next item in the stream will have additional dependencies on the nodes in the captured event. The two streams are then being captured to the same capture graph.

当流捕获中存在跨流依赖项时，仍必须在调用cudaStreamBeginCapture()的同一流中调用cudaStreamEndCapture()；这是原始流。由于基于事件的依赖性，正在捕获到同一捕获图的任何其他流也必须重新加入原始流。如下所示。在cudaStreamEndCapture()上，所有捕获到同一捕获图的流都退出捕获模式。无法重新加入原始流将导致整个捕获操作失败。  
When cross-stream dependencies are present in stream capture, cudaStreamEndCapture() must still be called in the same stream where cudaStreamBeginCapture() was called; this is the origin stream. Any other streams which are being captured to the same capture graph, due to event-based dependencies, must also be joined back to the origin stream. This is illustrated below. All streams being captured to the same capture graph are taken out of capture mode upon cudaStreamEndCapture(). Failure to rejoin to the origin stream will result in failure of the overall capture operation.

// 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所示。注意：当流从捕获模式中退出时，尽管中间项已被删除，流中的下一个未捕获项（如果有）仍将依赖于最新的先前未捕获项。

###### Prohibited and Unhandled Operations ######

同步或查询正在捕获的流或捕获的事件的执行状态是无效的，因为它们不代表计划执行的项目。查询或同步包含活动流捕获的更广泛的句柄（例如，当任何关联的流处于捕获模式时的设备或上下文句柄）的执行状态或同步也是无效的。 It is also invalid to query the execution status of or synchronize a broader handle which encompasses an active stream capture, such as a device or context handle when any associated stream is in capture mode.

当捕获相同上下文中的任何流且未使用cudaStreamNonBlocking创建该流时，任何尝试使用旧式流都是无效的。这是因为旧版流在任何时候都包含这些其他流。排队进入遗留流将对捕获的流产生依赖性，对其进行查询或同步将查询或同步捕获的流。enqueueing to the legacy stream would create a dependency on the streams being captured, and querying it or synchronizing it would query or synchronize the streams being captured.因此，在这种情况下，调用同步API也是无效的。同步API（例如cudaMemcpy()）将工作加入旧版流，并在返回之前对其进行同步。

注意：作为一般规则，当依赖性关系将捕获的内容与未捕获的内容连接起来，而不是排队等待执行时，CUDA宁愿返回错误而不是忽略依赖关系。将流放入或退出捕获模式是一个例外An exception is made for placing a stream into or out of capture mode。这在模式转换之前和之后立即切断了添加到流中的项目之间的依赖关系。通过等待来自正在捕获的流中的捕获事件并与该事件不同的捕获图相关联来合并两个单独的捕获图是无效的。等待正在捕获的流中的未捕获事件是无效的。

当前，图形中目前不支持少数将异步操作加入流中的API，如果使用正在捕获的流进行调用，则将返回错误，例如cudaStreamAttachMemAsync()。

###### Invalidation ######

在流捕获期间尝试执行无效操作时，所有关联的捕获图都将无效。 当捕获图无效时，将无法继续使用任何正在捕获的流或与该图关联的捕获事件，并且将返回错误，直到使用cudaStreamEndCapture()结束流捕获为止。 此调用将使关联的流退出捕获模式，但还将返回错误值和NULL图。

##### Updating Instantiated Graphs #####

使用图提交工作分为三个不同的阶段：定义，实例化和执行。在工作流不变的情况下，可以在许多执行中分摊定义和实例化的开销，并且图形相对于流提供了明显的优势。图形是工作流程的快照，包括内核，参数和相关性，以便尽可能快速，高效地重播它。在工作流程更改的情况下，该图将过时，必须进行修改。图结构的重大更改（例如拓扑或节点类型）将需要重新实例化源图，因为必须重新应用各种与拓扑相关的优化技术。  
重复实例化的成本可能降低图形执行带来的总体性能优势，但是通常只有节点参数（例如内核参数和cudaMemcpy地址）发生变化，而图形拓扑保持不变，这是很常见的。对于这种情况，**CUDA提供了一种称为“图形更新”的轻量级机制，该机制允许就地修改某些节点参数，而不必重建整个图形**。这比重新实例化要有效得多。更新将在下一次启动图形时生效，因此即使更新时正在运行，更新也不会影响先前的图形启动。可以重复更新和重新启动图，因此可以在流上将多个更新/发布排队。

CUDA提供了两种更新实例化图的机制：*整个图更新*和*单个节点更新*。整个图更新允许用户提供拓扑相同的cudaGraph\_t对象，该对象的节点包含更新的参数。单个节点更新允许用户显式更新单个节点的参数。当要更新大量节点时，或者调用者不知道图拓扑时（即图是由库调用的流捕获产生的），使用更新的cudaGraph\_t更为方便。当更改数量少并且用户拥有需要更新的节点的句柄时，首选使用单个节点更新。单个节点更新会跳过对未更改节点的拓扑检查和比较，因此在许多情况下可以提高效率。以下各节将详细介绍每种方法。

###### Graph Update Limitations ######

Kernel nodes:该函数不得更改，但是可以更新内核参数和启动参数。  
cudaMemset and cudaMemcpy nodes:分配/映射操作数的CUDA设备无法更改、必须从与原始源/目标存储器相同的上下文中分配源/目标存储器、只能更改一维cudaMemset / cudaMemcpy节点。  
Additional memcpy node restrictions:不支持更改源或目标内存类型（即cudaPitchedPtr，cudaArray\_t等）或更改传输类型（即cudaMemcpyKind）。  
主机节点更新没有任何限制。

###### Whole Graph Update ######

cudaGraphExecUpdate()允许使用拓扑相同的图形（“更新”图形）中的参数更新实例化的图形（“原始图形”）。 更新图的拓扑必须与用于实例化cudaGraphExec\_t的原始图相同。 此外，将节点添加到原始图中或从原始图中删除的顺序必须与将节点添加到更新图中（或从更新图中删除）的顺序相匹配。 因此，在使用流捕获时，必须以相同顺序捕获节点，而在使用显式图节点创建API时，必须以相同顺序添加和/或删除所有节点。  
以下示例显示了如何使用API更新实例化图：

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创建初始cudaGraph\_t。然后将cudaGraph\_t实例化并正常启动。初始启动后，使用与初始图形相同的方法创建新的cudaGraph\_t，并调用cudaGraphExecUpdate（）。如果图形更新成功（如上例中的updateResult参数所示），则将启动更新的cudaGraphExec\_t。如果更新由于任何原因失败，则调用cudaGraphExecDestroy（）和cudaGraphInstantiate（）销毁原始的cudaGraphExec\_t并实例化一个新的cudaGraphExec\_t。  
也可以直接更新cudaGraph\_t节点（即，使用cudaGraphKernelNodeSetParams（）），然后再更新cudaGraphExec\_t，但是使用下一节中介绍的显式节点更新API效率更高。

注意：当前无法使用cudaGraphExecUpdate（）更新包含子图节点或空节点的图。有关使用和当前限制的更多信息，请参见Graph API。

###### Individual node update ######

实例化的图节点参数可以直接更新。 这消除了实例化的开销以及创建新的cudaGraph\_t的开销。 如果需要更新的节点数相对于图中的节点总数少，则最好单独更新节点。 以下方法可用于更新cudaGraphExec\_t节点：  
cudaGraphExecKernelNodeSetParams（）  
cudaGraphExecMemcpyNodeSetParams（）  
cudaGraphExecMemsetNodeSetParams（）  
cudaGraphExecHostNodeSetParams（）  
有关使用和当前限制的更多信息，请参见Graph API。

##### Using Graph APIs #####

cudaGraph\_t对象不是线程安全的。 用户有责任确保多个线程不会同时访问同一cudaGraph\_t。cudaGraphExec\_t不能与其自身并行运行。 先前启动同一可执行图后，将命令启动cudaGraphExec\_t。图形执行是在流中完成的，以便与其他异步工作进行排序。 但是，该流仅用于订购； 它不限制图的内部并行性，也不影响图节点的执行位置。  
cudaGraph\_t objects are not thread-safe. It is the responsibility of the user to ensure that multiple threads do not concurrently access the same cudaGraph\_t. A cudaGraphExec\_t cannot run concurrently with itself. A launch of a cudaGraphExec\_t will be ordered after previous launches of the same executable graph.Graph execution is done in streams for ordering with other asynchronous work. However, the stream is for ordering only; it does not constrain the internal parallelism of the graph, nor does it affect where graph nodes execute.

#### Events ####

运行时还通过允许应用程序在程序中的任何点异步记录事件并查询何时完成这些事件，来提供一种密切监视设备进度以及执行准确定时的方法。 当事件之前的所有任务-或可选地，给定流中的所有命令-已完成时，事件即告完成。 在所有流中的所有先前任务和命令完成之后，将完成流零中的事件。

##### Creation and Destruction #####

The following code sample creates two events:

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

They are destroyed this way:

cudaEventDestroy(start);
    cudaEventDestroy(stop);

##### Elapsed Time #####

在创建和销毁中创建的事件可以通过以下方式用于计时创建和销毁的代码示例：

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

#### Synchronous Calls ####

调用同步函数时，在设备完成请求的任务之前，控制权不会返回给主机线程。 在主机线程执行任何其他CUDA调用之前，可以通过调用带有某些特定标志的cudaSetDeviceFlags()（请参见参考手册以了解详细信息）来指定主机线程是屈服，阻塞还是旋转(yield, block, or spin)。

[Link 1]: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#device-enumeration
[15]: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#features-and-technical-specifications__technical-specifications-per-compute-capability
[Link 2]: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#overlapping-behavior
[Link 3]: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#implicit-synchronization
[Link 4]: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#events
[Link 5]: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#overlap-of-data-transfer-and-kernel-execution
[Link 6]: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#concurrent-kernel-execution
[Link 7]: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#concurrent-data-transfers
[Link 8]: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#creation-and-destruction-streams
[Figure 9. Child Graph Example]: /images/20221122/16a41d49892c46b3b48ed78a60187840.png
[Figure 10. Creating a Graph Using Graph APIs Example]: /images/20221122/7398f321607a40b79058468399e56588.png
[Link 9]: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#default-stream