Summary
HIP currently lacks hipGraphConditionalHandle and conditional graph nodes (if/else/while subgraphs). CUDA 12.4 introduced cudaGraphConditionalHandle which enables dynamic control flow inside captured GPU graphs. This is a significant capability gap that forces framework-level workarounds in several AMD software stack components.
Impact on AMD software stack
1. ROCm/Warp (GPU physics simulation framework) — PR #15
wp.capture_while() requires conditional graph nodes. On HIP, Warp falls back to a fixed-count loop:
# warp/_src/optim/linear.py (PR #15 — feat/hip-graph-capture)
if not wp.is_conditional_graph_supported():
# HIP/ROCm: no conditional nodes — always runs maxiter iterations
fixed_count_loop() # cannot exit early even when convergedCG/CR/BiCGSTAB/GMRES solvers always run maxiter iterations on AMD, even when they converge in 5% of that count.
2. ROCm/mujoco_warp (MuJoCo GPU physics) — PR #4
The Newton solver convergence check cannot be captured inside a graph on HIP. Current workaround: host-side D2H sync every 3 iterations — adds CPU-GPU synchronization that breaks zero-copy graph replay. Without conditional nodes, we pre-compile 5 separate graphs (1, 4, 10, 40, 100 iterations) and dispatch between them from the CPU, partially negating graph capture benefits.
3. General impact
Any iterative convergence-based algorithm (physics solvers, linear solvers, RL training loops, signal processing) that benefits from GPU graph capture is affected on AMD.
Benchmark context (AMD MI325X, gfx942, ROCm 7.2)
- mujoco_warp Newton solver: must run all 100 iterations per step instead of 1-2 (workaround: ~15µs D2H overhead per step for 3 convergence checks)
- Warp linear solvers: always run
maxiter cycles regardless of convergence
- Net effect: physics simulation step time 3.2× slower than H200 at small batch sizes, primarily due to this bottleneck
Requested implementation
// New HIP APIs (matching CUDA 12.4 semantics):
hipError_t hipGraphConditionalHandleCreate(
hipGraphConditionalHandle* pHandle,
hipGraph_t hGraph,
unsigned int defaultLaunchValue,
unsigned int flags);
// Device-side setter:
__device__ void hipGraphSetConditional(
hipGraphConditionalHandle handle,
unsigned int value);References
cc @rtmadduri
Summary
HIP currently lacks
hipGraphConditionalHandleand conditional graph nodes (if/else/while subgraphs). CUDA 12.4 introducedcudaGraphConditionalHandlewhich enables dynamic control flow inside captured GPU graphs. This is a significant capability gap that forces framework-level workarounds in several AMD software stack components.Impact on AMD software stack
1. ROCm/Warp (GPU physics simulation framework) — PR #15
wp.capture_while()requires conditional graph nodes. On HIP, Warp falls back to a fixed-count loop:CG/CR/BiCGSTAB/GMRES solvers always run
maxiteriterations on AMD, even when they converge in 5% of that count.2. ROCm/mujoco_warp (MuJoCo GPU physics) — PR #4
The Newton solver convergence check cannot be captured inside a graph on HIP. Current workaround: host-side D2H sync every 3 iterations — adds CPU-GPU synchronization that breaks zero-copy graph replay. Without conditional nodes, we pre-compile 5 separate graphs (1, 4, 10, 40, 100 iterations) and dispatch between them from the CPU, partially negating graph capture benefits.
3. General impact
Any iterative convergence-based algorithm (physics solvers, linear solvers, RL training loops, signal processing) that benefits from GPU graph capture is affected on AMD.
Benchmark context (AMD MI325X, gfx942, ROCm 7.2)
maxitercycles regardless of convergenceRequested implementation
References
cc @rtmadduri