CUDA graphs#

CUDA graphs (sim only) — stream capture replayed as a recorded op list Capture records the deferred ops a stream would have queued; instantiating snapshots them; launching re-enqueues them onto the target stream. Ordering across captured streams is not modelled (a captured graph replays in the order it was recorded), which is enough for the “record a pipeline once, replay it many times” pattern the samples demonstrate.

Defined in code/simulator.h.

This page documents 12 symbols: 2 enums, 10 functions.

Enums#

Signature

Description

Availability

Location

Example

enum cudaStreamCaptureMode

Capture records the deferred ops a stream would have queued; instantiating snapshots them; launching re-enqueues them onto the target stream.

sim only

simulator.h:3459

enum cudaStreamCaptureStatus

sim only

simulator.h:3463

Functions#

Signature

Description

Availability

Location

Example

inline cudaError_t cudaGraphCreate(cudaGraph_t* g, unsigned )

sim only

simulator.h:3492

View

inline cudaError_t cudaGraphDestroy(cudaGraph_t g)

sim only

simulator.h:3516

View

inline cudaError_t cudaGraphExecDestroy(cudaGraphExec_t e)

sim only

simulator.h:3517

View

inline cudaError_t cudaGraphGetNodes(cudaGraph_t g, cudaGraphNode_t* , size_t* numNodes)

sim only

simulator.h:3518

View

inline cudaError_t cudaGraphInstantiate(cudaGraphExec_t* e, cudaGraph_t g, cudaGraphNode_t* = nullptr, char* = nullptr, size_t = 0)

sim only

simulator.h:3499

View

inline cudaError_t cudaGraphInstantiateWithFlags(cudaGraphExec_t* e, cudaGraph_t g, unsigned long long = 0)

sim only

simulator.h:3507

inline cudaError_t cudaGraphLaunch(cudaGraphExec_t e, cudaStream_t s)

sim only

simulator.h:3510

View

inline cudaError_t cudaStreamBeginCapture(cudaStream_t s, cudaStreamCaptureMode = cudaStreamCaptureModeGlobal)

sim only

simulator.h:3473

View

inline cudaError_t cudaStreamEndCapture(cudaStream_t s, cudaGraph_t* g)

sim only

simulator.h:3481

View

inline cudaError_t cudaStreamIsCapturing(cudaStream_t s, cudaStreamCaptureStatus* st)

sim only

simulator.h:3495

Examples#

Code from the repository’s example corpus (code\*.cu); the highlighted line is the call site. These blocks are what the View links above open.

cudaGraphCreate (test_web_olcf_hw13_graph.cu:100)
 90    cudaMalloc((void**)&d_x, N * sizeof(float));
 91    cudaMalloc((void**)&d_y, N * sizeof(float));
 92    check("cudaMalloc failed");
 93    cudaMemcpy(d_x, h_x.data(), N * sizeof(float), cudaMemcpyHostToDevice);  // [SIM][FIX] byte count
 94    cudaMemcpy(d_y, h_y.data(), N * sizeof(float), cudaMemcpyHostToDevice);  // [SIM][FIX]
 95    check("cudaMemcpy failed");
 96
 97    bool            graphCreated = false;
 98    cudaGraph_t     graph    = nullptr;
 99    cudaGraphExec_t instance = nullptr;   // [SIM] initialised (/sdl warns otherwise)
100    cudaGraphCreate(&graph, 0);
101
102    const int threads = 128;
103    const int blocks  = (N + threads - 1) / threads;   // [SIM][FIX] was N + (threads-1)/threads
104
105    const int nReplays = 3;                            // [SIM] was 100
106    for (int i = 0; i < nReplays; ++i) {
107        if (!graphCreated) {
108            cudaStreamBeginCapture(streams[0], cudaStreamCaptureModeGlobal);
109            check("Stream begin capture failed");
110
111            LAUNCH_STREAM(kernel_a, blocks, threads, streams[0], d_x, d_y);
cudaGraphDestroy (test_web_olcf_hw13_graph.cu:168)
158    }
159    for (int i = 0; i < 8; ++i) std::printf("%2.0f ", h_y[i]);
160    std::printf("\n");
161
162    cudaFree(d_x);
163    cudaFree(d_y);
164    for (int i = 0; i < num_streams; ++i) cudaStreamDestroy(streams[i]);
165    cudaEventDestroy(event1);
166    cudaEventDestroy(event2);
167    cudaGraphExecDestroy(instance);
168    cudaGraphDestroy(graph);
169
170    std::printf("CUDA graph captured over 2 streams, replayed %d times: %s\n",
171                nReplays, g_ok ? "PASSED" : "FAILED");
172    return g_ok ? 0 : 1;
173}
cudaGraphExecDestroy (test_web_olcf_hw13_graph.cu:167)
157        }
158    }
159    for (int i = 0; i < 8; ++i) std::printf("%2.0f ", h_y[i]);
160    std::printf("\n");
161
162    cudaFree(d_x);
163    cudaFree(d_y);
164    for (int i = 0; i < num_streams; ++i) cudaStreamDestroy(streams[i]);
165    cudaEventDestroy(event1);
166    cudaEventDestroy(event2);
167    cudaGraphExecDestroy(instance);
168    cudaGraphDestroy(graph);
169
170    std::printf("CUDA graph captured over 2 streams, replayed %d times: %s\n",
171                nReplays, g_ok ? "PASSED" : "FAILED");
172    return g_ok ? 0 : 1;
173}
cudaGraphGetNodes (test_web_olcf_hw13_graph.cu:136)
126            check("instantiating graph failed");
127            graphCreated = true;
128        }
129        cudaGraphLaunch(instance, streams[0]);
130        check("Launching graph failed");
131        cudaStreamSynchronize(streams[0]);
132    }
133
134    cudaGraphNode_t* nodes    = NULL;
135    size_t           numNodes = 0;
136    cudaGraphGetNodes(graph, nodes, &numNodes);
137    check("Graph get nodes failed");
138    std::printf("Number of nodes in the graph = %zu\n", numNodes);
139    if (numNodes != 4) {                       // kernel_a, _b, _c, _d
140        std::printf("expected 4 captured kernel nodes\n");
141        g_ok = false;
142    }
143
144    cudaDeviceSynchronize();
145    cudaMemcpy(h_y.data(), d_y, N * sizeof(float), cudaMemcpyDeviceToHost);  // [SIM][FIX]
146    cudaDeviceSynchronize();
cudaGraphInstantiate (test_web_olcf_hw13_graph.cu:125)
115            cudaStreamWaitEvent(streams[1], event1);   // streams[1] joins the capture
116            LAUNCH_STREAM(kernel_c, blocks, threads, streams[1], d_x, d_y);
117            cudaEventRecord(event2, streams[1]);
118
119            cudaStreamWaitEvent(streams[0], event2);
120            LAUNCH_STREAM(kernel_d, blocks, threads, streams[0], d_x, d_y);
121
122            cudaStreamEndCapture(streams[0], &graph);
123            check("Stream end capture failed");
124
125            cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);
126            check("instantiating graph failed");
127            graphCreated = true;
128        }
129        cudaGraphLaunch(instance, streams[0]);
130        check("Launching graph failed");
131        cudaStreamSynchronize(streams[0]);
132    }
133
134    cudaGraphNode_t* nodes    = NULL;
135    size_t           numNodes = 0;
136    cudaGraphGetNodes(graph, nodes, &numNodes);
cudaGraphLaunch (test_web_olcf_hw13_graph.cu:129)
119            cudaStreamWaitEvent(streams[0], event2);
120            LAUNCH_STREAM(kernel_d, blocks, threads, streams[0], d_x, d_y);
121
122            cudaStreamEndCapture(streams[0], &graph);
123            check("Stream end capture failed");
124
125            cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);
126            check("instantiating graph failed");
127            graphCreated = true;
128        }
129        cudaGraphLaunch(instance, streams[0]);
130        check("Launching graph failed");
131        cudaStreamSynchronize(streams[0]);
132    }
133
134    cudaGraphNode_t* nodes    = NULL;
135    size_t           numNodes = 0;
136    cudaGraphGetNodes(graph, nodes, &numNodes);
137    check("Graph get nodes failed");
138    std::printf("Number of nodes in the graph = %zu\n", numNodes);
139    if (numNodes != 4) {                       // kernel_a, _b, _c, _d
140        std::printf("expected 4 captured kernel nodes\n");
cudaStreamBeginCapture (test_web_olcf_hw13_graph.cu:108)
 98    cudaGraph_t     graph    = nullptr;
 99    cudaGraphExec_t instance = nullptr;   // [SIM] initialised (/sdl warns otherwise)
100    cudaGraphCreate(&graph, 0);
101
102    const int threads = 128;
103    const int blocks  = (N + threads - 1) / threads;   // [SIM][FIX] was N + (threads-1)/threads
104
105    const int nReplays = 3;                            // [SIM] was 100
106    for (int i = 0; i < nReplays; ++i) {
107        if (!graphCreated) {
108            cudaStreamBeginCapture(streams[0], cudaStreamCaptureModeGlobal);
109            check("Stream begin capture failed");
110
111            LAUNCH_STREAM(kernel_a, blocks, threads, streams[0], d_x, d_y);
112            cudaEventRecord(event1, streams[0]);
113            LAUNCH_STREAM(kernel_b, blocks, threads, streams[0], d_x, d_y);
114
115            cudaStreamWaitEvent(streams[1], event1);   // streams[1] joins the capture
116            LAUNCH_STREAM(kernel_c, blocks, threads, streams[1], d_x, d_y);
117            cudaEventRecord(event2, streams[1]);
118
119            cudaStreamWaitEvent(streams[0], event2);
cudaStreamEndCapture (test_web_olcf_hw13_graph.cu:122)
112            cudaEventRecord(event1, streams[0]);
113            LAUNCH_STREAM(kernel_b, blocks, threads, streams[0], d_x, d_y);
114
115            cudaStreamWaitEvent(streams[1], event1);   // streams[1] joins the capture
116            LAUNCH_STREAM(kernel_c, blocks, threads, streams[1], d_x, d_y);
117            cudaEventRecord(event2, streams[1]);
118
119            cudaStreamWaitEvent(streams[0], event2);
120            LAUNCH_STREAM(kernel_d, blocks, threads, streams[0], d_x, d_y);
121
122            cudaStreamEndCapture(streams[0], &graph);
123            check("Stream end capture failed");
124
125            cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);
126            check("instantiating graph failed");
127            graphCreated = true;
128        }
129        cudaGraphLaunch(instance, streams[0]);
130        check("Launching graph failed");
131        cudaStreamSynchronize(streams[0]);
132    }
On this page