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缓冲区。
