一、CUDA中图的应用
图在各类算法中有着广泛的应用，而并行编程本身也可以被视为一种特定的算法处理模式。在了解了CUDA图的基本概念之后，我们接下来需要深入探讨的是如何在CUDA中构建图。掌握图的构建原理与过程，是为后续图的实例化、生成与执行打下坚实基础的关键一步。只有这样，我们才能将图的概念真正融入到并行编程的实践中。
二、CUDA中图的创建
在CUDA中，图的创建指的是在定义或创建阶段，程序需要定义图中各个操作的描述以及它们之间的依赖关系。更直白地说，就是根据预先设定的节点说明和操作描述，来创建具体的图节点并确立它们之间的依赖。这就像是先在纸上画出流程图及其关联，然后再将这幅“图”映射到计算机中。
CUDA提供了两种主要的图创建方式：
-
直接使用API创建
CUDA提供了诸如 cudaGraphCreate() 和 cudaGraphAddNode() 等接口来直接创建图并添加节点。更详细的API使用方法，建议查阅官方文档以深入学习。
// Create the graph - it starts out empty
cudaGraphCreate(&graph, 0);
// Create the nodes and their dependencies
cudaGraphNode_t nodes[4];
cudaGraphNodeParams kParams = { cudaGraphNodeTypeKernel };
kParams.kernel.func = (void *)kernelName;
kParams.kernel.gridDim.x = kParams.kernel.gridDim.y = kParams.kernel.gridDim.z = 1;
kParams.kernel.blockDim.x = kParams.kernel.blockDim.y = kParams.kernel.blockDim.z = 1;
cudaGraphAddNode(&nodes[0], graph, NULL, NULL, 0, &kParams);
cudaGraphAddNode(&nodes[1], graph, &nodes[0], NULL, 1, &kParams);
cudaGraphAddNode(&nodes[2], graph, &nodes[0], NULL, 1, &kParams);
cudaGraphAddNode(&nodes[3], graph, &nodes[1], NULL, 2, &kParams);
-
通过流捕获的方式
CUDA还支持通过捕获现有的、基于流的API调用来创建图。相比于直接创建，这种方式更像是“重用”已有的工作流来生成图，这更贴近现实场景：我们既可以从头绘制一幅新图，也可以在现有草图的基础上整合完善。主要接口包括 cudaStreamBeginCapture() 和 cudaStreamEndCapture()。
cudaGraph_t graph;
cudaStreamBeginCapture(stream);
kernel_A<<< ..., stream >>>(...);
kernel_B<<< ..., stream >>>(...);
libraryCall(stream);
kernel_C<<< ..., stream >>>(...);
cudaStreamEndCapture(stream, &graph);
上述代码示例来源于官方文档。
三、流捕获创建图的说明
使用流捕获来创建图时，有以下几点需要特别注意：
-
跨流的依赖和事件
当使用 cudaEventRecord 和 cudaStreamWaitEvent 在流捕获中处理跨流依赖时，如果等待的事件记录在同一个被捕获的图中，那么这种处理是有效的。类似地，在捕获模式下记录事件，相当于捕获了图中的一个节点集。
当一个流等待一个捕获事件，但自身尚未进入捕获模式时，它会被置为捕获模式，并在其下一个操作项处处理对该捕获事件节点的依赖。此时，这两个流将被纳入同一张图中。
如果流捕获中存在跨流依赖，则必须在调用 cudaStreamBeginCapture 的同一流中调用 cudaStreamEndCapture。对于基于事件的依赖，事件也必须将任何其他流连接回原始流。调用 cudaStreamEndCapture 后，所有被捕获到同一图中的流都将退出捕获模式。如果无法将流重新加入原始流，将导致整个捕获操作失败。
需要注意的是，当一个流退出捕获模式时，该流中下一个未捕获的操作项（如果存在）仍将依赖于最近一个先前的未捕获操作项，尽管中间的操作项已被移除。
-
特殊的操作与限制
在流捕获期间，某些操作是被禁止或不被处理的。例如，同步或查询正处于捕获状态的流或事件的执行状态是无效的；同样，当流处于捕获模式时，查询或同步包含该活动流捕获的更广泛句柄（如设备或上下文）的执行状态也是无效的。
在捕获同一上下文中非 cudaStreamNonBlocking 创建的流时，尝试操作传统流（legacy stream）是无效的。向传统流添加工作会创建对被捕获流的依赖，而查询或同步传统流则等同于查询或同步那些被捕获的流。
通常，当一个依赖关系试图连接已捕获的内容和未捕获且已排队的内容时，CUDA会返回错误。但在流进入或退出捕获模式的转换时刻，会移除转换前后立即添加到流中的操作项之间的依赖关系，这是一种特殊情况。
试图通过等待一个来自不同捕获图的事件，来合并两个独立的捕获图，是无效的。同样，在没有指定 cudaEventWaitExternal 标志的情况下，从一个被捕获的流中等待一个未捕获的事件，也是无效的。
目前，少数向流中排队异步操作的API不支持在图中使用，如果在流被捕获期间调用它们（例如 cudaStreamAttachMemAsync），将会返回错误。
-
操作失效
在流捕获期间尝试任何无效操作，都会使正在进行的捕获图变为无效状态。与之关联的流和事件也将保持无效，并持续返回错误，直到 cudaStreamEndCapture 被调用以退出捕获模式，届时函数将返回错误值和 NULL 图句柄。
-
内省捕获
可以使用 cudaStreamGetCaptureInfo 函数来检查活动的流捕获操作。它允许获取捕获状态、捕获的唯一（每个进程）ID、底层图对象以及流中要捕获的下一个节点的依赖关系（边）数据。这些依赖信息可用于获取流中最后捕获的节点的句柄。
四、应用实例
下面给出两个创建图的简单例程：
-
API直接创建示例
void cudaGraphsManual(float *inputVec_h,
float *inputVec_d,
double *outputVec_d,
double *result_d,
size_t inputSize,
size_t numOfBlocks)
{
cudaStream_t streamForGraph;
cudaGraph_t graph;
std::vector<cudaGraphNode_t> nodeDependencies;
cudaGraphNode_t memcpyNode, kernelNode, memsetNode;
double result_h = 0.0;
cudaStreamCreate(&streamForGraph);
cudaKernelNodeParams kernelNodeParams = {0};
cudaMemcpy3DParms memcpyParams = {0};
cudaMemsetParams memsetParams = {0};
memcpyParams.srcArray = NULL;
memcpyParams.srcPos = make_cudaPos(0, 0, 0);
memcpyParams.srcPtr = make_cudaPitchedPtr(inputVec_h, sizeof(float) * inputSize, inputSize, 1);
memcpyParams.dstArray = NULL;
memcpyParams.dstPos = make_cudaPos(0, 0, 0);
memcpyParams.dstPtr = make_cudaPitchedPtr(inputVec_d, sizeof(float) * inputSize, inputSize, 1);
memcpyParams.extent = make_cudaExtent(sizeof(float) * inputSize, 1, 1);
memcpyParams.kind = cudaMemcpyHostToDevice;
memsetParams.dst = (void *)outputVec_d;
memsetParams.value = 0;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(float); // elementSize can be max 4 bytes
memsetParams.width = numOfBlocks * 2;
memsetParams.height = 1;
cudaGraphCreate(&graph, 0);
cudaGraphAddMemcpyNode(&memcpyNode, graph, NULL, 0, &memcpyParams);
cudaGraphAddMemsetNode(&memsetNode, graph, NULL, 0, &memsetParams);
nodeDependencies.push_back(memsetNode);
nodeDependencies.push_back(memcpyNode);
void *kernelArgs[4] = {(void *)&inputVec_d, (void *)&outputVec_d, &inputSize, &numOfBlocks};
kernelNodeParams.func = (void *)reduce;
kernelNodeParams.gridDim = dim3(numOfBlocks, 1, 1);
kernelNodeParams.blockDim = dim3(THREADS_PER_BLOCK, 1, 1);
kernelNodeParams.sharedMemBytes = 0;
kernelNodeParams.kernelParams = (void **)kernelArgs;
kernelNodeParams.extra = NULL;
cudaGraphAddKernelNode(
&kernelNode, graph, nodeDependencies.data(), nodeDependencies.size(), &kernelNodeParams);
nodeDependencies.clear();
nodeDependencies.push_back(kernelNode);
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = result_d;
memsetParams.value = 0;
memsetParams.elementSize = sizeof(float);
memsetParams.width = 2;
memsetParams.height = 1;
cudaGraphAddMemsetNode(&memsetNode, graph, NULL, 0, &memsetParams);
nodeDependencies.push_back(memsetNode);
memset(&kernelNodeParams, 0, sizeof(kernelNodeParams));
kernelNodeParams.func = (void *)reduceFinal;
kernelNodeParams.gridDim = dim3(1, 1, 1);
kernelNodeParams.blockDim = dim3(THREADS_PER_BLOCK, 1, 1);
kernelNodeParams.sharedMemBytes = 0;
void *kernelArgs2[3] = {(void *)&outputVec_d, (void *)&result_d, &numOfBlocks};
kernelNodeParams.kernelParams = kernelArgs2;
kernelNodeParams.extra = NULL;
cudaGraphAddKernelNode(
&kernelNode, graph, nodeDependencies.data(), nodeDependencies.size(), &kernelNodeParams);
nodeDependencies.clear();
nodeDependencies.push_back(kernelNode);
memset(&memcpyParams, 0, sizeof(memcpyParams));
memcpyParams.srcArray = NULL;
memcpyParams.srcPos = make_cudaPos(0, 0, 0);
memcpyParams.srcPtr = make_cudaPitchedPtr(result_d, sizeof(double), 1, 1);
memcpyParams.dstArray = NULL;
memcpyParams.dstPos = make_cudaPos(0, 0, 0);
memcpyParams.dstPtr = make_cudaPitchedPtr(&result_h, sizeof(double), 1, 1);
memcpyParams.extent = make_cudaExtent(sizeof(double), 1, 1);
memcpyParams.kind = cudaMemcpyDeviceToHost;
cudaGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(), nodeDependencies.size(), &memcpyParams);
nodeDependencies.clear();
nodeDependencies.push_back(memcpyNode);
cudaGraphNode_t hostNode;
cudaHostNodeParams hostParams = {0};
hostParams.fn = myHostNodeCallback;
callBackData_t hostFnData;
hostFnData.data = &result_h;
hostFnData.fn_name = “cudaGraphsManual”;
hostParams.userData = &hostFnData;
cudaGraphAddHostNode(&hostNode, graph, nodeDependencies.data(), nodeDependencies.size(), &hostParams);
}
-
流捕获创建示例
void cudaGraphsUsingStreamCapture(float *inputVec_h,
float *inputVec_d,
double *outputVec_d,
double *result_d,
size_t inputSize,
size_t numOfBlocks)
{
cudaStream_t stream1, stream2, stream3, streamForGraph;
cudaEvent_t forkStreamEvent, memsetEvent1, memsetEvent2;
cudaGraph_t graph;
double result_h = 0.0;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
cudaStreamCreate(&stream3);
cudaStreamCreate(&streamForGraph);
cudaEventCreate(&forkStreamEvent);
cudaEventCreate(&memsetEvent1);
cudaEventCreate(&memsetEvent2);
cudaStreamBeginCapture(stream1, cudaStreamCaptureModeGlobal);
cudaEventRecord(forkStreamEvent, stream1);
cudaStreamWaitEvent(stream2, forkStreamEvent, 0);
cudaStreamWaitEvent(stream3, forkStreamEvent, 0);
cudaMemcpyAsync(inputVec_d, inputVec_h, sizeof(float) * inputSize, cudaMemcpyDefault, stream1);
cudaMemsetAsync(outputVec_d, 0, sizeof(double) * numOfBlocks, stream2);
cudaEventRecord(memsetEvent1, stream2);
cudaMemsetAsync(result_d, 0, sizeof(double), stream3);
cudaEventRecord(memsetEvent2, stream3);
cudaStreamWaitEvent(stream1, memsetEvent1, 0);
reduce<<<numOfBlocks, THREADS_PER_BLOCK, 0, stream1>>>(inputVec_d, outputVec_d, inputSize, numOfBlocks);
cudaStreamWaitEvent(stream1, memsetEvent2, 0);
reduceFinal<<<1, THREADS_PER_BLOCK, 0, stream1>>>(outputVec_d, result_d, numOfBlocks);
cudaMemcpyAsync(&result_h, result_d, sizeof(double), cudaMemcpyDefault, stream1);
callBackData_t hostFnData = {0};
hostFnData.data = &result_h;
hostFnData.fn_name = “cudaGraphsUsingStreamCapture”;
cudaHostFn_t fn = myHostNodeCallback;
cudaLaunchHostFunc(stream1, fn, &hostFnData);
cudaStreamEndCapture(stream1, &graph);
}
以上代码均来自官方文档，读者可以在此基础上进行完善和运行。需要注意的是，图功能需要CUDA 11及以上版本支持。如果你想运行上面的代码，请确保更新到相应的CUDA版本。由于图API本身也在不断迭代更新，如果在编译时遇到问题，请检查是否使用了不匹配的API接口。
五、总结
在CUDA图的应用中，创建是基础。只有正确、合理地掌握了图的构建方法，才能对图后续的各种高级应用有全面的掌控力。正所谓万丈高楼平地起，熟练掌握图的创建，就是在为构建高性能的并行计算应用打下坚实的根基。对于希望深入GPU并行计算领域的 C++ 开发者而言，理解并熟练运用CUDA图是一个重要的技能提升方向。 |
发表于 19 小时前
|
查看: 5|
回复: 0
