cudaGraphInstantiate() hangs with CUDA device graphs

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello experts

For several months we’ve been experiencing an issue when using CUDA device graphs.
In order for everything to work correctly, we’re forced to create and instantiate all graphs before launching them.

This limitation is quite inconvenient — could you tell us if there’s a specific reason for it?
Could it be related to the use of while() loops that make kernels “persistent” on the GPU?

We’ve created a small project that reproduces the problem; I’ve attached the ZIP file.
The core of the issue is as follows:

graphs[0]->createGraph();
graphs[0]->launchGraph(true);

for (size_t i = 1; i < graphs.size(); ++i)
{
    graphs[i]->createGraph();
}

deviceGraphs_issue.zip (4.4 KB)

The createGraph() call hangs on the second iteration, specifically at cudaGraphInstantiate(), which never returns (confirmed using GDB).

If we reduce the number of kernels per graph (from 20 we had in this example), cudaGraphInstantiate() hangs on the fourth iteration instead of the second.

The issue occurs on multiple GPUs, but definitely also on the NVIDIA RTX 6000 Ada.
If it were a GPU max scheduling issue, I would expect cudaGraphInstantiate() to return an error instead of hanging.

Thanks!

2

I can observe the issue. I discussed this with a few colleagues, and was advised that this generally falls under a documentation item here that says:

Any CUDA API call may block or synchronize for various reasons such as contention for or unavailability of internal resources. Such behavior is subject to change and undocumented behavior should not be relied upon.

You have a kernel (from a graph) that is running and for whatever reason (undocumented) this is interfering with the forward progress of the cudaGraphInstantiate call in this case. When the kernel (or graph) completes, the API call should be able to complete and return.

The condition is not an actual hang - when the kernel or graph completes, eventually whatever internal conditions are holding up the call should be cleared. Furthermore, there is no error here. This is expected behavior.

You can demonstrate this yourself by having the kernel delay for say, a few seconds, rather than waiting for a host interlock. Such a modification to your code results in no hang, and successful application completion.

Of course, its always necessary to create and instantiate a graph before you can launch that graph.

Looking at things from a bulk perspective, as long as your kernels/graphs do eventually complete, it seems the graph creation/instantiation can proceed even as kernels/graphs are getting launched (I’ve already indicated above how you can witness this with a modification to your test case).

Here is an example modification:

# cat t415a.cu
#include <iostream>
#include <memory>
#include <vector>
#include <cstdio>
#include <unistd.h>
#if 0
__global__ void waitKernel(const int *data)
{
    int idx = threadIdx.x + (blockIdx.x * blockDim.x);

    if (idx == 0) printf("GPU - Kernel started, waiting for data...\n");
    int d = __ldcg(data);
    if (idx == 0) printf("GPU - Initial Data[0]: %d\n", d);
    while (d == 0)
    {
        d = __ldcg(data);
        __nanosleep(100);
    }
    __syncthreads();

    if (idx == 0) printf("GPU - Data[0]: %d\n", data[0]);
}
#else
const unsigned long long TDELAY = 0x10000000ULL;
__global__ void waitKernel(const int *data)
{
  unsigned long long start = clock64();
  while (clock64() < (start+TDELAY)) {};
}
#endif


#include <string>

#define GPU_ERR_RET(val) \
    do { \
        cudaError_t err = (val); \
        if (err != cudaSuccess) \
        { \
            std::cerr << "CUDA Error: " << cudaGetErrorString(err) << " at " << __FILE__ << ":" << __LINE__ \
                      << std::endl; \
            return err; \
        } \
    } while (0)

#define GPU_ERR(val) \
    do { \
        cudaError_t err = (val); \
        if (err != cudaSuccess) \
        { \
            std::cerr << "CUDA Error: " << cudaGetErrorString(err) << " at " << __FILE__ << ":" << __LINE__ \
                      << std::endl; \
        } \
    } while (0)

class Graph
{
private:
    cudaGraphExec_t graphExec;
    int *dData;
    int *hData;
    cudaStream_t s1;
    cudaStream_t s2;
    std::string name;

public:
    Graph(const std::string &name);
    virtual ~Graph();

    cudaError_t createGraph();
    cudaError_t launchGraph(const bool loop);
    void stopLoop();
    cudaError_t waitGraph();
};



Graph::Graph(const std::string &name) : name(name) {}

Graph::~Graph()
{
    if (dData)
    {
        std::cout << name << " - Cleaning up resources..." << std::endl;
        GPU_ERR(cudaFree(dData));
        GPU_ERR(cudaFreeHost(hData));
        GPU_ERR(cudaStreamDestroy(s1));
        GPU_ERR(cudaStreamDestroy(s2));
        GPU_ERR(cudaGraphExecDestroy(graphExec));
        std::cout << name << " - Resources cleaned up." << std::endl;
    }
}

cudaError_t Graph::createGraph()
{
    // Allocate device data
    std::cout << name << " - Allocating device memory..." << std::endl;
    GPU_ERR_RET(cudaMalloc(&dData, sizeof(int)));
    std::cout << name << " - Device memory allocated." << std::endl;

    // Allocate host data
    std::cout << name << " - Allocating host memory..." << std::endl;
    GPU_ERR_RET(cudaMallocHost(&hData, sizeof(int)));
    std::cout << name << " - Host memory allocated." << std::endl;

    // Create stream
    std::cout << name << " - Creating stream..." << std::endl;
    GPU_ERR_RET(cudaStreamCreateWithFlags(&s1, cudaStreamNonBlocking));
    GPU_ERR_RET(cudaStreamCreateWithFlags(&s2, cudaStreamNonBlocking));
    std::cout << name << " - Stream created." << std::endl;

    // Create graph
    cudaGraph_t graph;
    std::cout << name << " - Creating graph..." << std::endl;
    GPU_ERR_RET(cudaGraphCreate(&graph, 0));
    std::cout << name << " - Graph created." << std::endl;

    // Kernel launch
    for (int n = 0; n < 20; n++)
    {
        cudaKernelNodeParams kernelNodeParams;
        void *kernelArgs[] = {&dData};
        kernelNodeParams.func = (void *)waitKernel;
        kernelNodeParams.gridDim = dim3(120);
        kernelNodeParams.blockDim = dim3(1024);
        kernelNodeParams.sharedMemBytes = 0;
        kernelNodeParams.kernelParams = (void **)kernelArgs;
        kernelNodeParams.extra = NULL;
        cudaGraphNode_t kernelNode;
        GPU_ERR_RET(cudaGraphAddKernelNode(&kernelNode, graph, NULL, 0, &kernelNodeParams));

        //cudaLaunchAttributeValue attr;
        //attr.cooperative = 1;
        //GPU_ERR_RET(cudaGraphKernelNodeSetAttribute(kernelNode, cudaLaunchAttributeCooperative, &attr));
    }

    // Instantiate graph
    std::cout << name << " - Instantiating graph..." << std::endl;
    GPU_ERR_RET(cudaGraphInstantiate(&graphExec, graph));
    std::cout << name << " - Graph instantiated." << std::endl;

    // Clean up
    std::cout << name << " - Destroying intermediate graph..." << std::endl;
    GPU_ERR_RET(cudaGraphDestroy(graph));
    std::cout << name << " - Intermediate graph destroyed." << std::endl;

    return cudaSuccess;
}

cudaError_t Graph::launchGraph(const bool loop)
{
    *hData = (loop) ? 0 : 1;
    // Copy data to device
    std::cout << "Copying data to device..." << std::endl;
    GPU_ERR_RET(cudaMemcpyAsync(dData, hData, sizeof(int), cudaMemcpyHostToDevice, s1));
    std::cout << "Launching graph..." << std::endl;

    // Launch graph
    GPU_ERR_RET(cudaGraphLaunch(graphExec, s1));
    std::cout << "Graph launched." << std::endl;
    return cudaSuccess;
}

void Graph::stopLoop()
{
    *hData = 1;
    // Copy data to device
    std::cout << "Stopping loop..." << std::endl;
    GPU_ERR(cudaMemcpyAsync(dData, hData, sizeof(int), cudaMemcpyHostToDevice, s2));
    std::cout << "Loop stop signal sent." << std::endl;
}

cudaError_t Graph::waitGraph()
{
    std::cout << "Waiting for graph to complete..." << std::endl;
    GPU_ERR_RET(cudaStreamSynchronize(s1));
    std::cout << "Graph completed." << std::endl;
    return cudaSuccess;
}


cudaError_t runTest()
{
    constexpr int numGraphs = 16;

    std::vector<std::unique_ptr<Graph>> graphs;

    for (int i = 0; i < numGraphs; ++i)
    {
        graphs.push_back(std::make_unique<Graph>("Graph " + std::to_string(i + 1)));
    }
    usleep(10000);

// This works
#if 0
    {
        for (auto &graph : graphs)
        {
            GPU_ERR_RET(graph->createGraph());
        }
        usleep(10000);

        GPU_ERR_RET(graphs[0]->launchGraph(true));
        usleep(10000);
    }
#endif
// This  works
#if 0
    {
        GPU_ERR_RET(graphs[0]->createGraph());
        GPU_ERR_RET(graphs[0]->launchGraph(true));
        usleep(10000);
        for (size_t i = 1; i < 2; ++i)
        {
            GPU_ERR_RET(graphs[i]->createGraph());
        }
    }
#endif
// This works
#if 0
    {
        GPU_ERR_RET(graphs[0]->createGraph());
        usleep(10000);
        for (size_t i = 1; i < graphs.size(); ++i)
        {
            GPU_ERR_RET(graphs[i]->createGraph());
        }
        GPU_ERR_RET(graphs[0]->launchGraph(true));
    }
#endif
// This does not work
#if 1
    {
        GPU_ERR_RET(graphs[0]->createGraph());
        GPU_ERR_RET(graphs[0]->launchGraph(true));
        usleep(10000);
        for (size_t i = 1; i < graphs.size(); ++i)
        {
            GPU_ERR_RET(graphs[i]->createGraph());
        }
    }
#endif

    std::cout << "Waiting 2 seconds to stop loop..." << std::endl;
    usleep(2000000);
    graphs[0]->stopLoop();

    for (size_t i = 1; i < graphs.size(); ++i)
    {
        GPU_ERR_RET(graphs[i]->launchGraph(false));
    }
    usleep(10000);

    std::cout << "Graphs have completed execution." << std::endl;

    return cudaSuccess;
}

int main(int argc, char **argv)
{
#if 0
    int *qq;
    cudaMalloc(&qq, sizeof(int));
    cudaMemset(qq, 1, sizeof(int));
    waitKernel<<<1,1>>>(qq);
    cudaDeviceSynchronize();
#endif
    const auto err = runTest();
    return (err == cudaSuccess) ? 0 : 1;
}

# nvcc -o t415a t415a.cu -arch=sm_89
# compute-sanitizer ./t415a
========= COMPUTE-SANITIZER
Graph 1 - Allocating device memory...
Graph 1 - Device memory allocated.
Graph 1 - Allocating host memory...
Graph 1 - Host memory allocated.
Graph 1 - Creating stream...
Graph 1 - Stream created.
Graph 1 - Creating graph...
Graph 1 - Graph created.
Graph 1 - Instantiating graph...
Graph 1 - Graph instantiated.
Graph 1 - Destroying intermediate graph...
Graph 1 - Intermediate graph destroyed.
Copying data to device...
Launching graph...
Graph launched.
Graph 2 - Allocating device memory...
Graph 2 - Device memory allocated.
Graph 2 - Allocating host memory...
Graph 2 - Host memory allocated.
Graph 2 - Creating stream...
Graph 2 - Stream created.
Graph 2 - Creating graph...
Graph 2 - Graph created.
Graph 2 - Instantiating graph...
Graph 2 - Graph instantiated.
Graph 2 - Destroying intermediate graph...
Graph 2 - Intermediate graph destroyed.
Graph 3 - Allocating device memory...
Graph 3 - Device memory allocated.
Graph 3 - Allocating host memory...
Graph 3 - Host memory allocated.
Graph 3 - Creating stream...
Graph 3 - Stream created.
Graph 3 - Creating graph...
Graph 3 - Graph created.
Graph 3 - Instantiating graph...
Graph 3 - Graph instantiated.
Graph 3 - Destroying intermediate graph...
Graph 3 - Intermediate graph destroyed.
Graph 4 - Allocating device memory...
Graph 4 - Device memory allocated.
Graph 4 - Allocating host memory...
Graph 4 - Host memory allocated.
Graph 4 - Creating stream...
Graph 4 - Stream created.
Graph 4 - Creating graph...
Graph 4 - Graph created.
Graph 4 - Instantiating graph...
Graph 4 - Graph instantiated.
Graph 4 - Destroying intermediate graph...
Graph 4 - Intermediate graph destroyed.
Graph 5 - Allocating device memory...
Graph 5 - Device memory allocated.
Graph 5 - Allocating host memory...
Graph 5 - Host memory allocated.
Graph 5 - Creating stream...
Graph 5 - Stream created.
Graph 5 - Creating graph...
Graph 5 - Graph created.
Graph 5 - Instantiating graph...
Graph 5 - Graph instantiated.
Graph 5 - Destroying intermediate graph...
Graph 5 - Intermediate graph destroyed.
Graph 6 - Allocating device memory...
Graph 6 - Device memory allocated.
Graph 6 - Allocating host memory...
Graph 6 - Host memory allocated.
Graph 6 - Creating stream...
Graph 6 - Stream created.
Graph 6 - Creating graph...
Graph 6 - Graph created.
Graph 6 - Instantiating graph...
Graph 6 - Graph instantiated.
Graph 6 - Destroying intermediate graph...
Graph 6 - Intermediate graph destroyed.
Graph 7 - Allocating device memory...
Graph 7 - Device memory allocated.
Graph 7 - Allocating host memory...
Graph 7 - Host memory allocated.
Graph 7 - Creating stream...
Graph 7 - Stream created.
Graph 7 - Creating graph...
Graph 7 - Graph created.
Graph 7 - Instantiating graph...
Graph 7 - Graph instantiated.
Graph 7 - Destroying intermediate graph...
Graph 7 - Intermediate graph destroyed.
Graph 8 - Allocating device memory...
Graph 8 - Device memory allocated.
Graph 8 - Allocating host memory...
Graph 8 - Host memory allocated.
Graph 8 - Creating stream...
Graph 8 - Stream created.
Graph 8 - Creating graph...
Graph 8 - Graph created.
Graph 8 - Instantiating graph...
Graph 8 - Graph instantiated.
Graph 8 - Destroying intermediate graph...
Graph 8 - Intermediate graph destroyed.
Graph 9 - Allocating device memory...
Graph 9 - Device memory allocated.
Graph 9 - Allocating host memory...
Graph 9 - Host memory allocated.
Graph 9 - Creating stream...
Graph 9 - Stream created.
Graph 9 - Creating graph...
Graph 9 - Graph created.
Graph 9 - Instantiating graph...
Graph 9 - Graph instantiated.
Graph 9 - Destroying intermediate graph...
Graph 9 - Intermediate graph destroyed.
Graph 10 - Allocating device memory...
Graph 10 - Device memory allocated.
Graph 10 - Allocating host memory...
Graph 10 - Host memory allocated.
Graph 10 - Creating stream...
Graph 10 - Stream created.
Graph 10 - Creating graph...
Graph 10 - Graph created.
Graph 10 - Instantiating graph...
Graph 10 - Graph instantiated.
Graph 10 - Destroying intermediate graph...
Graph 10 - Intermediate graph destroyed.
Graph 11 - Allocating device memory...
Graph 11 - Device memory allocated.
Graph 11 - Allocating host memory...
Graph 11 - Host memory allocated.
Graph 11 - Creating stream...
Graph 11 - Stream created.
Graph 11 - Creating graph...
Graph 11 - Graph created.
Graph 11 - Instantiating graph...
Graph 11 - Graph instantiated.
Graph 11 - Destroying intermediate graph...
Graph 11 - Intermediate graph destroyed.
Graph 12 - Allocating device memory...
Graph 12 - Device memory allocated.
Graph 12 - Allocating host memory...
Graph 12 - Host memory allocated.
Graph 12 - Creating stream...
Graph 12 - Stream created.
Graph 12 - Creating graph...
Graph 12 - Graph created.
Graph 12 - Instantiating graph...
Graph 12 - Graph instantiated.
Graph 12 - Destroying intermediate graph...
Graph 12 - Intermediate graph destroyed.
Graph 13 - Allocating device memory...
Graph 13 - Device memory allocated.
Graph 13 - Allocating host memory...
Graph 13 - Host memory allocated.
Graph 13 - Creating stream...
Graph 13 - Stream created.
Graph 13 - Creating graph...
Graph 13 - Graph created.
Graph 13 - Instantiating graph...
Graph 13 - Graph instantiated.
Graph 13 - Destroying intermediate graph...
Graph 13 - Intermediate graph destroyed.
Graph 14 - Allocating device memory...
Graph 14 - Device memory allocated.
Graph 14 - Allocating host memory...
Graph 14 - Host memory allocated.
Graph 14 - Creating stream...
Graph 14 - Stream created.
Graph 14 - Creating graph...
Graph 14 - Graph created.
Graph 14 - Instantiating graph...
Graph 14 - Graph instantiated.
Graph 14 - Destroying intermediate graph...
Graph 14 - Intermediate graph destroyed.
Graph 15 - Allocating device memory...
Graph 15 - Device memory allocated.
Graph 15 - Allocating host memory...
Graph 15 - Host memory allocated.
Graph 15 - Creating stream...
Graph 15 - Stream created.
Graph 15 - Creating graph...
Graph 15 - Graph created.
Graph 15 - Instantiating graph...
Graph 15 - Graph instantiated.
Graph 15 - Destroying intermediate graph...
Graph 15 - Intermediate graph destroyed.
Graph 16 - Allocating device memory...
Graph 16 - Device memory allocated.
Graph 16 - Allocating host memory...
Graph 16 - Host memory allocated.
Graph 16 - Creating stream...
Graph 16 - Stream created.
Graph 16 - Creating graph...
Graph 16 - Graph created.
Graph 16 - Instantiating graph...
Graph 16 - Graph instantiated.
Graph 16 - Destroying intermediate graph...
Graph 16 - Intermediate graph destroyed.
Waiting 2 seconds to stop loop...
Stopping loop...
Loop stop signal sent.
Copying data to device...
Launching graph...
Graph launched.
Copying data to device...
Launching graph...
Graph launched.
Copying data to device...
Launching graph...
Graph launched.
Copying data to device...
Launching graph...
Graph launched.
Copying data to device...
Launching graph...
Graph launched.
Copying data to device...
Launching graph...
Graph launched.
Copying data to device...
Launching graph...
Graph launched.
Copying data to device...
Launching graph...
Graph launched.
Copying data to device...
Launching graph...
Graph launched.
Copying data to device...
Launching graph...
Graph launched.
Copying data to device...
Launching graph...
Graph launched.
Copying data to device...
Launching graph...
Graph launched.
Copying data to device...
Launching graph...
Graph launched.
Copying data to device...
Launching graph...
Graph launched.
Copying data to device...
Launching graph...
Graph launched.
Graphs have completed execution.
Graph 1 - Cleaning up resources...
Graph 1 - Resources cleaned up.
Graph 2 - Cleaning up resources...
Graph 2 - Resources cleaned up.
Graph 3 - Cleaning up resources...
Graph 3 - Resources cleaned up.
Graph 4 - Cleaning up resources...
Graph 4 - Resources cleaned up.
Graph 5 - Cleaning up resources...
Graph 5 - Resources cleaned up.
Graph 6 - Cleaning up resources...
Graph 6 - Resources cleaned up.
Graph 7 - Cleaning up resources...
Graph 7 - Resources cleaned up.
Graph 8 - Cleaning up resources...
Graph 8 - Resources cleaned up.
Graph 9 - Cleaning up resources...
Graph 9 - Resources cleaned up.
Graph 10 - Cleaning up resources...
Graph 10 - Resources cleaned up.
Graph 11 - Cleaning up resources...
Graph 11 - Resources cleaned up.
Graph 12 - Cleaning up resources...
Graph 12 - Resources cleaned up.
Graph 13 - Cleaning up resources...
Graph 13 - Resources cleaned up.
Graph 14 - Cleaning up resources...
Graph 14 - Resources cleaned up.
Graph 15 - Cleaning up resources...
Graph 15 - Resources cleaned up.
Graph 16 - Cleaning up resources...
Graph 16 - Resources cleaned up.
========= ERROR SUMMARY: 0 errors
#

Notes:

  • CUDA 13.0, L4 GPU
  • Because of the kernel time delay, the above code requires a few seconds before it prints out “Graphs have completed execution.”, and requires ~30 seconds to actually complete/finish. You can add a printout to the kernel code if you wish to see this. Many of the kernels will finish after the “Graphs have completed execution.” printout.
3

Thank you for the fast and detailed answer.

Unfortunately we’re using those “persistent” “never-ending” kernels for our real-time application to reduce to the bare minimum the CPU/GPU interactions (every CUDA driver call is taking too much time and it’s not predictable).

I think that cudaGraphInstantiate hanging should be somehow detected by the driver after a certain time deadline to have a more stable SW but anyway thanks for the clarification.

I think we can try to have persistent kernels that are forced to end once every X hundreds of millisecond to allow other CUDA API call to “unlock” and continue…

In the future I hope there would be some official NVIDIA support of these persistent kernels because they’re extremely useful to have a microsecond reaction to events, something absolutely required for real time applications.

I think you Aerial tool would also love to have this feature :-)

Have a nice day