
一、CUDA中图的应用图在算法应用中非常多其实并行编程也可以看作是一种算法处理的方式。在前面学习了CUDA中的图的相关概念后就该进一步了解并掌握CUDA中图如何构建。只有把图的构建的过程和原理弄清楚才能为后面的图的实例化生成和执行做好铺垫。这样才能够真正的把图与并行编程融合在一起。二、CUDA中图的创建在CUDA中图的创建指在定义或创建阶段程序会创建图中操作的描述以及它们之间的依赖关系。说的直白一些其实就是根据预先提供的节点说明和操作的描述对图的节点创建和它们之间的操作依赖生成具体的图。说得再简单一点就是如何在纸上把图画出来而且画出他们之间的关系同时还需要把这种画出的图转到计算机中映射出来。在CUDA中流的创建方式主要有两种形式直接使用API创建在CUDA中提供了如cudaGraphCreate()和cudaGraphAddNode()等接口用来创建图并添加节点。更细节的大家可以查看相关的API文档来学习CUDA如何控制和管理图。// Create the graph - it starts out emptycudaGraphCreate(graph,0);// Create the nodes and their dependenciescudaGraphNode_tnodes[4];cudaGraphNodeParams kParams{cudaGraphNodeTypeKernel};kParams.kernel.func(void*)kernelName;kParams.kernel.gridDim.xkParams.kernel.gridDim.ykParams.kernel.gridDim.z1;kParams.kernel.blockDim.xkParams.kernel.blockDim.ykParams.kernel.blockDim.z1;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应用的尝试都是无效的。向传统流中加入工作将创建对被捕获流的依赖而查询或同步它将等同于查询或同步那些被捕获的流。在这种情况下同步API也是无效的。注意通常情况下当某个依赖关系将已捕获的内容与未捕获且已排队执行的内容相连接时CUDA一般返回错误而不是忽略该依赖。不过有一种特殊情况即将流置入或退出捕获模式。其会移除在模式转换前后立即添加到流中的操作项之间的依赖关系。通过等待来自一个正在被捕获、且与事件关联不同捕获图的流的捕获事件来合并两个独立的捕获图是无效的。在没有指定cudaEventWaitExternal标志的情况下从一个正在被捕获的流中等待非捕获事件是无效的。目前有少数向流中排队异步操作的API不支持在图中使用如果在流被捕获期间调用它们如cudaStreamAttachMemAsync函数时将返回错误。失效在流捕获时尝试无效操作时相关的捕获图都将无效。即其正在使用的流和图关联的事件也都是无效的它会一直返回错误。直到cudaStreamEndCapture函数结束退出捕获模式返回错误值和NULL图。内省捕获利用cudaStreamGetCaptureInfo函数可以检查活动流捕获操作。它允许获得捕获的状态、捕获的唯一每个进程ID、底层图对象以及流中要捕获的下一个节点的依赖关系/边缘数据。此依赖关系信息可用于获得流中最后捕获的节点的句柄。四、应用下面给出一个图的简单例程API创建voidcudaGraphsManual(float*inputVec_h,float*inputVec_d,double*outputVec_d,double*result_d,size_tinputSize,size_tnumOfBlocks){cudaStream_tstreamForGraph;cudaGraph_tgraph;std::vectorcudaGraphNode_tnodeDependencies;cudaGraphNode_tmemcpyNode,kernelNode,memsetNode;doubleresult_h0.0;cudaStreamCreate(streamForGraph));cudaKernelNodeParams kernelNodeParams{0};cudaMemcpy3DParms memcpyParams{0};cudaMemsetParams memsetParams{0};memcpyParams.srcArrayNULL;memcpyParams.srcPosmake_cudaPos(0,0,0);memcpyParams.srcPtrmake_cudaPitchedPtr(inputVec_h,sizeof(float)*inputSize,inputSize,1);memcpyParams.dstArrayNULL;memcpyParams.dstPosmake_cudaPos(0,0,0);memcpyParams.dstPtrmake_cudaPitchedPtr(inputVec_d,sizeof(float)*inputSize,inputSize,1);memcpyParams.extentmake_cudaExtent(sizeof(float)*inputSize,1,1);memcpyParams.kindcudaMemcpyHostToDevice;memsetParams.dst(void*)outputVec_d;memsetParams.value0;memsetParams.pitch0;memsetParams.elementSizesizeof(float);// elementSize can be max 4 bytesmemsetParams.widthnumOfBlocks*2;memsetParams.height1;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.gridDimdim3(numOfBlocks,1,1);kernelNodeParams.blockDimdim3(THREADS_PER_BLOCK,1,1);kernelNodeParams.sharedMemBytes0;kernelNodeParams.kernelParams(void**)kernelArgs;kernelNodeParams.extraNULL;cudaGraphAddKernelNode(kernelNode,graph,nodeDependencies.data(),nodeDependencies.size(),kernelNodeParams);nodeDependencies.clear();nodeDependencies.push_back(kernelNode);memset(memsetParams,0,sizeof(memsetParams));memsetParams.dstresult_d;memsetParams.value0;memsetParams.elementSizesizeof(float);memsetParams.width2;memsetParams.height1;cudaGraphAddMemsetNode(memsetNode,graph,NULL,0,memsetParams);nodeDependencies.push_back(memsetNode);memset(kernelNodeParams,0,sizeof(kernelNodeParams));kernelNodeParams.func(void*)reduceFinal;kernelNodeParams.gridDimdim3(1,1,1);kernelNodeParams.blockDimdim3(THREADS_PER_BLOCK,1,1);kernelNodeParams.sharedMemBytes0;void*kernelArgs2[3]{(void*)outputVec_d,(void*)result_d,numOfBlocks};kernelNodeParams.kernelParamskernelArgs2;kernelNodeParams.extraNULL;cudaGraphAddKernelNode(kernelNode,graph,nodeDependencies.data(),nodeDependencies.size(),kernelNodeParams);nodeDependencies.clear();nodeDependencies.push_back(kernelNode);memset(memcpyParams,0,sizeof(memcpyParams));memcpyParams.srcArrayNULL;memcpyParams.srcPosmake_cudaPos(0,0,0);memcpyParams.srcPtrmake_cudaPitchedPtr(result_d,sizeof(double),1,1);memcpyParams.dstArrayNULL;memcpyParams.dstPosmake_cudaPos(0,0,0);memcpyParams.dstPtrmake_cudaPitchedPtr(result_h,sizeof(double),1,1);memcpyParams.extentmake_cudaExtent(sizeof(double),1,1);memcpyParams.kindcudaMemcpyDeviceToHost;cudaGraphAddMemcpyNode(memcpyNode,graph,nodeDependencies.data(),nodeDependencies.size(),memcpyParams);nodeDependencies.clear();nodeDependencies.push_back(memcpyNode);cudaGraphNode_thostNode;cudaHostNodeParams hostParams{0};hostParams.fnmyHostNodeCallback;callBackData_thostFnData;hostFnData.dataresult_h;hostFnData.fn_namecudaGraphsManual;hostParams.userDatahostFnData;cudaGraphAddHostNode(hostNode,graph,nodeDependencies.data(),nodeDependencies.size(),hostParams);}捕获流创建voidcudaGraphsUsingStreamCapture(float*inputVec_h,float*inputVec_d,double*outputVec_d,double*result_d,size_tinputSize,size_tnumOfBlocks){cudaStream_tstream1,stream2,stream3,streamForGraph;cudaEvent_tforkStreamEvent,memsetEvent1,memsetEvent2;cudaGraph_tgraph;doubleresult_h0.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);reducenumOfBlocks,THREADS_PER_BLOCK,0,stream1(inputVec_d,outputVec_d,inputSize,numOfBlocks);cudaStreamWaitEvent(stream1,memsetEvent2,0);reduceFinal1,THREADS_PER_BLOCK,0,stream1(outputVec_d,result_d,numOfBlocks);cudaMemcpyAsync(result_h,result_d,sizeof(double),cudaMemcpyDefault,stream1);callBackData_thostFnData{0};hostFnData.dataresult_h;hostFnData.fn_namecudaGraphsUsingStreamCapture;cudaHostFn_tfnmyHostNodeCallback;cudaLaunchHostFunc(stream1,fn,hostFnData);cudaStreamEndCapture(stream1,graph);}代码来自官方文档大家可以在上面代码的基础上进行完善并运行。由于图的支持需要在CUDA11以上才可以所以如果想运行上面的代码需要更新至相关的CUDA版本。特别是图的更新迭代一直在进行在编译出现问题时看是否出现了不匹配的API接口。五、总结图的应用创建是基础。只有把握了如何正确合理的创建图才能够对图以后的各种应用有着全面的把控。正所谓万丈高楼平地起掌握图的创建就是在给高楼打地基。