CUDA 图形通过将用户操作定义为任务图（可以在单个操作中启动），显著减少了启动大量用户操作的开销。提前了解工作流使 CUDA 驱动程序能够应用各种优化，这在通过流模型启动时无法执行。
然而，这种性能是以灵活性为代价的：如果事先不知道整个工作流，则 GPU 执行必须中断，才能返回 CPU 做出决定。
CUDA 设备图启动通过基于在运行时确定的数据使任务图能够从正在运行的 GPU 内核高效地启动来解决这个问题。 CUDA 设备图形发射提供了两种不同的发射模式：发射和忘记，以及尾部发射，以实现广泛的应用和使用。
这篇文章演示了如何使用设备图形启动和两种启动模式。它以设备端工作调度程序为例，它解压缩文件以进行数据处理。
设备图初始化
执行任务图涉及以下四个步骤：
- 创建图形
- 将图形实例化为可执行图形
- 将可执行图形的工作描述符上载到 GPU
- 启动可执行图形
通过将启动步骤与其他步骤分离， CUDA 能够优化工作流，并使图形启动尽可能轻。为了方便起见，如果没有明确调用上传步骤， CUDA 也将在第一次启动图形时将上传步骤与启动步骤相结合。
为了从 CUDA 内核启动图形，必须首先在实例化步骤中初始化图形以用于设备启动。此外，在可以从设备启动设备之前，设备图必须已通过手动上传步骤或通过主机启动隐式上传到设备。下面的代码执行主机端步骤以设置设备调度程序示例，显示了两个选项：
// This is the signature of our scheduler kernel
// The internals of this kernel will be outlined later
__global__ void schedulerKernel(
fileData *files,
int numFiles,
int *currentFile,
void **currentFileData,
cudaGraphExec_t zipGraph,
cudaGraphExec_t lzwGraph,
cudaGraphExec_t deflateGraph);
void setupAndLaunchScheduler() {
cudaGraph_t zipGraph, lzwGraph, deflateGraph, schedulerGraph;
cudaGraphExec_t zipExec, lzwExec, deflateExec, schedulerExec;
// Create the source graphs for each possible operation we want to perform
// We pass the currentFileData ptr to this setup, as this ptr is how the scheduler will
// indicate which file to decompress
create_zip_graph(&zipGraph, currentFileData);
create_lzw_graph(&lzwGraph, currentFileData);
create_deflate_graph(&deflateGraph, currentFileData);
// Instantiate the graphs for these operations and explicitly upload
cudaGraphInstantiate(&zipExec, zipGraph, cudaGraphInstantiateFlagDeviceLaunch);
cudaGraphUpload(zipExec, stream);
cudaGraphInstantiate(&lzwExec, lzwGraph, cudaGraphInstantiateFlagDeviceLaunch);
cudaGraphUpload(lzwExec, stream);
cudaGraphInstantiate(&deflateExec, deflateGraph, cudaGraphInstantiateFlagDeviceLaunch);
cudaGraphUpload(deflateExec, stream);
// Create and instantiate the scheduler graph
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
schedulerKernel<<<1, 1, 0, stream>>>(files, numFiles, currentFile, currentFileData, zipExec, lzwExec, deflateExec);
cudaStreamEndCapture(stream, &schedulerGraph);
cudaGraphInstantiate(&schedulerExec, schedulerGraph, cudaGraphInstantiateFlagDeviceLaunch);
// Launch the scheduler graph - this will perform an implicit upload
cudaGraphLaunch(schedulerExec, stream);
}
这里需要注意的是，设备图形可以从主机或设备启动。因此，可以向调度器传递与在主机上启动相同的cudaGraphExec_t句柄，以便在设备上启动。
点火忘记发射
调度器内核根据传入的数据分派工作。对于工作调度，首选的启动方法是“点火后忘记启动”。
当使用 fire and forget launch 启动图形时，将立即发送该图形。它独立于启动图和使用 fire and forget 模式启动的后续图执行。因为工作会立即执行，所以对于调度程序调度的工作，最好是“启动即忘”，因为它会尽快开始运行。 CUDA 引入了一个名为 stream 的新设备端，以执行一个图的点火和忘记启动。请参见下面的简单分派器示例。
enum compressionType {
zip = 1,
lzw = 2,
deflate = 3
};
struct fileData {
compressionType comprType;
void *data;
};
__global__ void schedulerKernel(
fileData *files,
int numFiles
int *currentFile,
void **currentFileData,
cudaGraphExec_t zipGraph,
cudaGraphExec_t lzwGraph,
cudaGraphExec_t deflateGraph)
{
// Set the data ptr to the current file so the dispatched graph
// is operating on the correct file data
*currentFileData = files[currentFile].data;
switch (files[currentFile].comprType) {
case zip:
cudaGraphLaunch(zipGraph, cudaStreamGraphFireAndForget);
break;
case lzw:
cudaGraphLaunch(lzwGraph, cudaStreamGraphFireAndForget);
break;
case deflate:
cudaGraphLaunch(deflateGraph, cudaStreamGraphFireAndForget);
break;
default:
break;
}
}
还需要注意的是，图形启动可以是嵌套的和递归的，因此可以从 fire 和 forget 启动中调度其他设备图形。虽然在本示例中未示出，但是正在解压缩文件数据的图形可以在数据完全解压缩后分派更多的图形来对该数据进行进一步处理（例如，图像处理）。设备图流是分层的，就像图本身一样。
尾部发射
CUDA 工作异步启动到 GPU ，这意味着启动线程必须明确等待工作完成，然后才能使用任何结果或输出。这通常是使用诸如cudaStreamSynchronize或cudaStreamSynchronize的同步操作从 CPU 线程完成的。
GPU 上的启动线程不可能通过cudaDeviceSynchronize等传统方法同步设备图形启动。相反，当需要操作顺序时，应使用尾部发射。
当一个图被提交用于尾部发射时，它不会立即执行，而是在发射图完成后执行。 CUDA 将所有动态生成的工作封装为父图的一部分，因此尾部启动也将在执行之前等待所有生成的 fire 和 forget 工作。
无论尾部发射是在任何发射之前还是之后发出的，这都是正确的。尾部发射本身按照它们排队的顺序执行。一个特殊情况是自动重新启动，其中当前运行的设备图被排队以通过尾部启动重新启动。一次只允许一次待定的自动重新启动。
使用 tail launch ，您可以通过反复重新启动前一个调度器内核，从而有效地在执行流中创建循环，从而将其升级为完整的调度器内核：
__global__ void schedulerKernel(
fileData *files,
int numFiles,
int *currentFile,
void **currentFileData,
cudaGraphExec_t zipGraph,
cudaGraphExec_t lzwGraph,
cudaGraphExec_t deflateGraph)
{
// Set the data ptr to the current file so the dispatched graph
// is operating on the correct file data
*currentFileData = files[currentFile].data;
switch (files[currentFile].comprType) {
case zip:
cudaGraphLaunch(zipGraph, cudaStreamGraphFireAndForget);
break;
case lzw:
cudaGraphLaunch(lzwGraph, cudaStreamGraphFireAndForget);
break;
case deflate:
cudaGraphLaunch(deflateGraph, cudaStreamGraphFireAndForget);
break;
default:
break;
}
// If we have not finished iterating over all the files, relaunch
if (*currentFile < numFiles) {
// Query the current graph handle so we can relaunch it
cudaGraphExec_t currentGraph = cudaGetCurrentGraphExec();
cudaGraphLaunch(currentGraph, cudaStreamGraphTailLaunch);
*currentFile++;
}
}
请注意，重新启动操作如何使用cudaGetCurrentGraphExec检索当前正在执行的图形的句柄。它可以重新启动自己，而不需要自己的可执行图的句柄。
在自动重新启动时使用尾部启动具有额外的效果，即在下一次调度程序内核重新启动开始之前同步（等待）调度的 fire 和 forget 工作。一个设备图一次只能有一次待启动（加上一次自动重新启动）。为了重新启动刚刚发送的图形，您需要确保先前的启动首先完成。执行自我重新启动可以实现这一目标，这样您就可以为下一次迭代调度所需的任何图形。
设备与主机启动性能的比较
此示例与主机启动的图形相比如何？图 1 比较了各种拓扑的启动延迟、尾部启动延迟和主机启动延迟。
该图表显示，不仅设备端启动延迟比主机启动延迟低 2 倍以上，而且还不受图形结构的影响。每个给定拓扑的延迟都是相同的。
如图 2 所示，设备启动也可以更好地扩展到图形的宽度。
与主机启动相比，无论图表中的并行度如何，设备启动延迟几乎保持不变。
结论
CUDA 设备图启动提供了一种在 CUDA 内核内实现动态控制流的高效方式。虽然本文中给出的示例提供了一种开始使用该功能的方法，但它只是该功能使用方式的一个小表示。
有关更多文档，请参阅编程指南的 device graph launch 部分。要尝试设备图形启动， download CUDA Toolkit 12.0 。
