cuda-gen: Add CUDA Graph capture and replay for composite operators by Nafees01 · Pull Request #1987 · CEED/libCEED

Nafees01 · 2026-07-08T18:28:55Z

Summary

Adds CUDA Graph capture and replay to the 'cuda-gen' backend for composite operators. On the second apply, the kernel sequence is captured into a cudaGraph; all subsequent applies replay the instantiated graph. QFunction contexts are re-synced to device before each replay so time- and load-dependent parameters stay correct.

Sub-operators that cannot be captured, for example, contact operators that allocate memory (cudaMalloc) during apply, which CUDA doesn't allow inside a capture. When this happens, the capture is properly closed to leave the stream in a clean state, any leftover CUDA error is cleared, and that operator falls back to running normally on every call.

Validated against both /gpu/cuda/ref and cuda-gen without graphs across static, quasistatic, contact, multi-material, and dynamic elasticity examples. Strain energy and displacements match to ~1e-13 in every case, with zero graph fallbacks on capturable operators.

Performance-wise, graphs are within ±1% of cuda-gen without graphs across all sizes, so no regression, but no real speedup either. The reason is that cuda-gen already fuses everything into a few large kernels, so there's very little launch overhead left for graphs to remove.

Validated correctness against /gpu/cuda/ref and cuda-gen (CEED_DISABLE_GRAPH=1) on the following Ratel examples:

ex01-static-elasticity-linear-mms
ex02-quasistatic-elasticity-linear-platen (contact)
ex02-quasistatic-elasticity-neo-hookean-current-sinker (multi-material)
ex03-dynamic-elasticity-mooney-rivlin-current
Strain energy matches to ~1e-13 in all cases. Zero graph fallbacks on capturable operators.

zatkins-dev · 2026-07-08T19:30:08Z

Sub-operators that cannot be captured, for example, contact operators that allocate memory (cudaMalloc) during apply

Are there specific operators you were unable to use in Ratel? None of the QFunctions there should alloc during operator application.

zatkins-dev · 2026-07-08T19:31:53Z

There's some style errors here, be sure to make format each commit

zatkins-dev

I'm sure you're still working on this, but here's an initial review to correct some high-level design issues and removed functionality that I think should be fixed.

zatkins-dev · 2026-07-08T19:34:28Z

+    return CEED_ERROR_SUCCESS;
  }
-  if (is_sequential) CeedCallCuda(ceed, cudaStreamDestroy(stream));
-  if (input_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorRestoreArrayRead(input_vec, &input_arr));


This removed the is_sequential checks, which are there to ensure that all suboperators run in a single stream.

zatkins-dev · 2026-07-08T19:37:21Z

-      CeedCallBackend(CeedOperatorApplyAddCore_Cuda_gen(sub_operators[i], stream, input_arr, output_arr, &is_run_good[i], request));
-      if (!is_sequential) CeedCallCuda(ceed, cudaStreamDestroy(stream));
+  // CEED_DISABLE_GRAPH=1: run suboperators directly, no capture/replay (for benchmarking).
+  static bool disable_graph = false;


This should be stored on a per-operator (or at least per-Ceed) basis, not as static variables. It's okay to use the environment variable to set a default value, but there should be a setter (e.g. CeedSetUseCudaGraph(Ceed ceed, bool use_graph))

zatkins-dev · 2026-07-08T19:39:59Z

-    if (!is_run_good[i]) {
-      CeedOperator op_fallback;
+  // No real I/O buffers to capture; just run directly.
+  if (input_vec == CEED_VECTOR_NONE || output_vec == CEED_VECTOR_NONE) {


Should this be?

Suggested change

if (input_vec == CEED_VECTOR_NONE || output_vec == CEED_VECTOR_NONE) {

if (input_vec == CEED_VECTOR_NONE && output_vec == CEED_VECTOR_NONE) {

It's possible to have a passive input and active output

zatkins-dev · 2026-07-08T19:43:51Z

+  }
+  if (disable_graph) {
+    for (CeedInt i = 0; i < num_suboperators; i++) {
+      CeedCallBackend(CeedOperatorApplyAdd(sub_operators[i], input_vec, output_vec, request));


This call pattern prevents each suboperator from running on separate streams. You won't see a perf difference until you have large, multioperator applications, but the perf difference can be significant.

The disabled graph case should restore the previous implementation, with multi-stream operators, sequential checks, etc

Given you do this a bunch of places, I'd recommend a helper function which looks like:

static int CeedOperatorApplyAddComposite_NoGraph_Cuda_gen(CeedOperator op, CeedVector input_vec, CeedVector output_vec, CeedRequest *request) { bool is_run_good[CEED_COMPOSITE_MAX] = {false}, is_sequential; CeedInt num_suboperators; const CeedScalar *input_arr = NULL; CeedScalar *output_arr = NULL; Ceed ceed; CeedOperator *sub_operators; cudaStream_t stream = NULL; CeedCallBackend(CeedOperatorGetCeed(op, &ceed)); CeedCall(CeedOperatorCompositeGetNumSub(op, &num_suboperators)); CeedCall(CeedOperatorCompositeGetSubList(op, &sub_operators)); CeedCall(CeedOperatorCompositeIsSequential(op, &is_sequential)); if (input_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArrayRead(input_vec, CEED_MEM_DEVICE, &input_arr)); if (output_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArray(output_vec, CEED_MEM_DEVICE, &output_arr)); if (is_sequential) CeedCallCuda(ceed, cudaStreamCreate(&stream)); for (CeedInt i = 0; i < num_suboperators; i++) { CeedInt num_elem = 0; CeedCall(CeedOperatorGetNumElements(sub_operators[i], &num_elem)); if (num_elem > 0) { if (!is_sequential) CeedCallCuda(ceed, cudaStreamCreate(&stream)); CeedCallBackend(CeedOperatorApplyAddCore_Cuda_gen(sub_operators[i], stream, input_arr, output_arr, &is_run_good[i], request)); if (!is_sequential) CeedCallCuda(ceed, cudaStreamDestroy(stream)); } } if (is_sequential) CeedCallCuda(ceed, cudaStreamDestroy(stream)); if (input_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorRestoreArrayRead(input_vec, &input_arr)); if (output_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorRestoreArray(output_vec, &output_arr)); CeedCallCuda(ceed, cudaDeviceSynchronize()); // Fallback on unsuccessful run for (CeedInt i = 0; i < num_suboperators; i++) { if (!is_run_good[i]) { CeedOperator op_fallback; CeedDebug(ceed, "\nFalling back to /gpu/cuda/ref CeedOperator for ApplyAdd\n"); CeedCallBackend(CeedOperatorGetFallback(sub_operators[i], &op_fallback)); CeedCallBackend(CeedOperatorApplyAdd(op_fallback, input_vec, output_vec, request)); } } CeedCallBackend(CeedDestroy(&ceed)); return CEED_ERROR_SUCCESS; }

zatkins-dev · 2026-07-08T19:56:53Z

+  // No real I/O buffers to capture; just run directly.
+  if (input_vec == CEED_VECTOR_NONE || output_vec == CEED_VECTOR_NONE) {
+    for (CeedInt i = 0; i < num_suboperators; i++) {
+      CeedCallBackend(CeedOperatorApplyAdd(sub_operators[i], input_vec, output_vec, request));


Need to use the same helper here

zatkins-dev · 2026-07-08T20:05:39Z

+  CeedOperator op_fallback;
+  CeedCallBackend(CeedOperatorGetFallback(op, &op_fallback));
+  CeedCallBackend(CeedOperatorApplyAdd(op_fallback, input_vec, output_vec, request));


Same here -- use the helper, not fall back.

zatkins-dev · 2026-07-08T20:06:37Z

+  // Phase 3: replay.
+  {
+    // graph_instance == NULL means capture failed; direct-apply every call.
+    if (!impl->graph_instance) {


Use helper here

zatkins-dev · 2026-07-08T20:38:51Z

+    {
+      const CeedScalar *in_ptr;
+      bool              ptr_ok;
+
+      CeedCallBackend(CeedVectorGetArrayRead(input_vec, CEED_MEM_DEVICE, &in_ptr));
+      ptr_ok = (in_ptr == impl->captured_input_ptr);
+      CeedCallBackend(CeedVectorRestoreArrayRead(input_vec, &in_ptr));
+      if (!ptr_ok) goto use_fallback;
+    }


This check seems like it should happen before the graph is built, that way you can rebuild the graph if needed instead of waiting until the next apply call

zatkins-dev · 2026-07-08T20:40:22Z

+// Push each suboperator's QFunction context to device. Replay skips the normal
+// apply path, so we do this by hand to keep time/load parameters current.
+static int CeedCompositeRefreshContexts_Cuda_gen(CeedOperator *sub_operators, CeedInt num_suboperators) {
+  for (CeedInt i = 0; i < num_suboperators; i++) {
+    CeedQFunction qf  = NULL;
+    void         *d_c = NULL;
+
+    CeedCallBackend(CeedOperatorGetQFunction(sub_operators[i], &qf));
+    CeedCallBackend(CeedQFunctionGetInnerContextData(qf, CEED_MEM_DEVICE, &d_c));
+    CeedCallBackend(CeedQFunctionRestoreInnerContextData(qf, &d_c));
+    CeedCallBackend(CeedQFunctionDestroy(&qf));
+  }
+  return CEED_ERROR_SUCCESS;
+}


This same issue is going to happen for any passive inputs which have been updated on the CPU -- there should be a more robust way to handle this

zatkins-dev · 2026-07-08T20:42:24Z

+  // CEED_FORCE_BASELINE=1: skip graphs, run via /gpu/cuda/ref
+  static bool force_baseline = false;
+  static bool force_baseline_checked = false;
+  if (!force_baseline_checked) {
+    char *env_val = getenv("CEED_FORCE_BASELINE");
+    force_baseline = (env_val != NULL && strcmp(env_val, "1") == 0);
+    force_baseline_checked = true;
+  }
+  if (force_baseline) {
+    CeedOperator op_fallback;
+    CeedCallBackend(CeedOperatorGetFallback(op, &op_fallback));
+    CeedCallBackend(CeedOperatorApplyAdd(op_fallback, input_vec, output_vec, request));
+    return CEED_ERROR_SUCCESS;
+  }


This isn't the right way to benchmark against ref -- use the /gpu/cuda/ref backend directly.

Suggested change

// CEED_FORCE_BASELINE=1: skip graphs, run via /gpu/cuda/ref

static bool force_baseline = false;

static bool force_baseline_checked = false;

if (!force_baseline_checked) {

char *env_val = getenv("CEED_FORCE_BASELINE");

force_baseline = (env_val != NULL && strcmp(env_val, "1") == 0);

force_baseline_checked = true;

}

if (force_baseline) {

CeedOperator op_fallback;

CeedCallBackend(CeedOperatorGetFallback(op, &op_fallback));

CeedCallBackend(CeedOperatorApplyAdd(op_fallback, input_vec, output_vec, request));

return CEED_ERROR_SUCCESS;

}

zatkins-dev · 2026-07-08T22:58:21Z

I would be interested to see if this speeds up e.g. /gpu/cuda/ref, when recording the graph over the entire CeedOperatorApplyCore_Cuda_Ref function

Nafees01 added 14 commits July 25, 2025 17:17

CUDA graph implemented

6261385

Implement CUDA graphs with fallback system in libCEED

b156a00

Cleaner and Simplified Implementation of CUDA graph

83c1d9f

add vector pointer tracking to detect memory changes in Graph

bbad4e2

Add PETSc vector setup for CUDA Graph compatibility in CUDA-gen backend

285b661

Clean CUDA-graph implementation

10d6dd2

Simple and clean implementation of CUDA graph but numerically incorrect

4abc7ad

per-operator CUDA Graph implementation

1fa80ed

replace cudaMemset with cudaMemsetAsync

5b8e14f

Auto-detect graph capture and use async memset with cudaStreamPerThread

94bdeca

CUDA Graph support for composite operators in cuda-gen backend

a00af53

CUDA Graph is working fine for composite operators

1ca2b5f

cuda-gen: CUDA Graph capture and replay working for composite operators

719027b

Merge remote-tracking branch 'upstream/main' into cuda-graph-dev

40f8c3e

Nafees01 force-pushed the cuda-graph-dev branch from 17c97a3 to 40f8c3e Compare July 8, 2026 18:54

zatkins-dev requested changes Jul 8, 2026

View reviewed changes

style: apply clang-format-22

91c8d6d

	if (input_vec == CEED_VECTOR_NONE \|\| output_vec == CEED_VECTOR_NONE) {
	if (input_vec == CEED_VECTOR_NONE && output_vec == CEED_VECTOR_NONE) {

Uh oh!

Conversation

Nafees01 commented Jul 8, 2026

Summary

Uh oh!

zatkins-dev commented Jul 8, 2026

Uh oh!

zatkins-dev commented Jul 8, 2026

Uh oh!

zatkins-dev left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

zatkins-dev commented Jul 8, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants