Documentation on cudaGraphXXX() Graph Management functions?

Accelerated Computing CUDA CUDA Programming and Performance
1

After a quick skim of the Runtime API’s Graph Management functions, I’m not seeing any information on the lifetimes of cudaGraph_t and cudaGraphExec_t objects.

I’m going to assume that the API supports destroying cudaGraph_t and cudaGraphExec_t objects as soon as they’re no longer needed.

Specifically:

cudaGraphCreate(&cg)          // create the graph
// add nodes...               // build the graph 
cudaGraphInstantiate(&cge,cg) // build the executable
cudaGraphDestroy(cg)          // we're done with the graph so destroy it
cudaGraphLaunch(cge,stream)   // launch the executable
cudaGraphExecDestroy(cge)     // we're done with the executable so destroy it

Please confirm!

2

Yes, you can destroy graphs that are no longer needed.

A fairly complete CUDA graphs example is available in the CUDA sample codes. “simpleCudaGraphs”

3

Hello! I reviewed again the latest version of simpleCudaGraphs.cu example. The relevant portion of the example that I have question about:

    cudaGraphExec_t graphExec;
    checkCudaErrors(cudaGraphInstantiate(&graphExec, graph, NULL, NULL, 0));

    cudaGraph_t     clonedGraph;
    cudaGraphExec_t clonedGraphExec;
    checkCudaErrors(cudaGraphClone(&clonedGraph, graph));
    checkCudaErrors(cudaGraphInstantiate(&clonedGraphExec, clonedGraph, NULL, NULL, 0));

    for (int i = 0; i < GRAPH_LAUNCH_ITERATIONS; i++) {
        checkCudaErrors(cudaGraphLaunch(graphExec, streamForGraph));
    }

    checkCudaErrors(cudaStreamSynchronize(streamForGraph));

    printf("Cloned Graph Output.. \n");
    for (int i = 0; i < GRAPH_LAUNCH_ITERATIONS; i++) {
        checkCudaErrors(cudaGraphLaunch(clonedGraphExec, streamForGraph));
    }

    checkCudaErrors(cudaStreamSynchronize(streamForGraph));

    checkCudaErrors(cudaGraphExecDestroy(graphExec));
    checkCudaErrors(cudaGraphExecDestroy(clonedGraphExec));
    checkCudaErrors(cudaGraphDestroy(graph));
    checkCudaErrors(cudaGraphDestroy(clonedGraph));

In this example we first wait on completion of commands (including launched CUDA graphs) issued on the stream by using

checkCudaErrors(cudaStreamSynchronize(streamForGraph));

My question is, in some way, a repetition of @allanmac . I want to be sure that I can issue

cudaGraphLaunch(clonedGraphExec, streamForGraph);

and then immediately

cudaGraphExecDestroy(graphExec);

even though the actual execution on the streamForGraph could be delayed. In my actual code I use driver API and I wait on some value (acts like a semaphore) in order to release the next batch of CUDA commands on some stream. It is important to be sure that in this way, the resources related to graphExec handle will be properly deallocated/released yet, at the time the GPU get to the actual execution of the graph, the handle is still valid, despite cudaGraphExecDestroy(graphExec); has already been called at that stage.

Please confirm again!

4

It is possible to have a sequence such as you are showing. However the details (surrounding code, exact structure/contents of the graph) matter. I suggest reading this section of the programming guide.