Personal notes I took when learning to use NVIDIA's Warp.
| Warp | JAX analogue |
|---|---|
wp.array |
a GPU array, but mutable. |
arr.requires_grad = True |
gives the array a .grad array of the same shape, where gradients accumulate |
@wp.kernel + wp.launch(k, dim=(B,n), inputs=[...], outputs=[...]) |
a hand-written elementwise op; wp.tid() returns the thread index (like vmap indices) |
wp.Tape() |
a stateful jax.vjp. record with with tape:, write the output cotangent into out.grad, call tape.backward(), then read in.grad. tape.zero() clears every .grad |
mjw.Model / mjw.Data |
structs of wp.array's. Data fields have a leading world dimension: d.pos (nworld, nq), d.qvel, d.ctrl, d.qacc_warmstart |
mjw.step(m, d) |
in place: overwrites d |
with wp.ScopedCapture() as cap: |
tracing under jax.jit: the launches inside are recorded into cap.graph instead of just running |
wp.capture_launch(cap.graph) |
calling the compiled function |
wp.atomic_add(a, i, v) / wp.atomic_max(a, i, v) |
a.at[i].add(v) / a.at[i].max(v), but safe when many threads write the same slot |
wp.jax_callable(fn, num_outputs=k, graph_mode=...) turns a Python function into a JAX op:
fn's arguments need Warp type annotations. Inputs come first and outputs last.output_dims={'name': shape}graph_mode=wp.JaxCallableGraphMode.NONE, the Python body runs on every call, so host-side logic like .numpy() and if works. With the default JAX mode it runs only once, during CUDA graph capture.They control how much of the code actually runs as Python on each call. The closest JAX comparison is eager op-by-op executions versus the JIT.
Taking MJWarp into the context, a step in it is hundreds of small kernels. Launching each one from Python costs CPU time, often more than the kernel itself takes on the GPU. A CUDA graph records a sequence of kernel launches once and then replays the whole sequence with one call. It's roughly the speedup JIT gets over op-by-op dispatch, except nothing is recompiled or fused. The same kernels just run without per-launch Python overhead.
| Mode | What happens on each call | JAX analogy |
|---|---|---|
| NONE | Python function runs every time and launches each kernel on the stream XLA provides. | Eager mode: jax.disable_jit() or pure callback |
| JAX | Python runs once, while Warp records the launches into a CUDA graph. XLA replays that graph on later calls. | jax.jit |
| WARP | Also captures a graph, but Warp manages and replays it instead of XLA. | jax.jit (but different owner) |
In theory JAX or WARP mode is all you need. In practice, with a function that also records a wp.Tape, both failed for me: JAX mode with CUDA_ERROR_NOT_SUPPORTED, WARP mode with an illegal memory access on the first call and allocations during capture.
What worked was staying in NONE mode and capturing the graph by hand:
_GRAPHS = {}
def run(key, body):
if key not in _GRAPHS:
body() # warm up: compile kernels, allocate everything
with wp.ScopedCapture() as cap:
body() # record the launches
_GRAPHS[key] = cap.graph
wp.capture_launch(_GRAPHS[key]) # replay the launches with one callbody() reads and writes a set of staging buffers that are allocated once and never move. Each call copies the JAX inputs into them, replays the graph, and copies the outputs back. key is anything that changes the graph, like the batch size. This is the same reason jit retraces on a new shape.
The speedup was an order of magnitude over launching every kernel from Python.
Outside of jax_callable, wp.from_jax(x) and wp.to_jax(a) convert between the two. Neither copies: both sides end up looking at the same GPU memory.
Inside a jax_callable, you don't even need those. The arguments already arrive as wp.array's that point into JAX's buffers, and the outputs are buffers JAX allocated for you to write into. So the function body is just Warp code.
The catch is that JAX gives you a different buffer on every call. That's fine for eager launches, but a captured graph has its addresses baked in. So the body has three steps:
def fwd(qpos: wp.array2d(dtype=float), ..., qpos_out: wp.array2d(dtype=float)):
wp.copy(stage.qpos, qpos) # JAX buffer => fixed staging buffer
run('fwd', lambda: body(stage)) # replay the graph on the staging buffers
wp.copy(qpos_out, stage.qpos_out) # fixed buffer => JAX's output bufferThese copies are tiny next to a physics step, but proceed with caution and take them in consideration.
Two smaller things:
float in a Warp annotation means float32, which matches JAX's default. With jax_enable_x64 on, the dtypes stop matching.vmap over the Python function. If the rest of your code works per-env with (nq,) arrays, add the batch axis yourself (qpos[None]) before the call.A captured graph is a fixed list of kernel launches with the memory addresses baked in. It's very close to the rules of tracing:
.numpy() or printing a value forces a sync and doesn't get recorded. Same as not being able to print a traced value.if runs once, at capture time, and the branch it took is what gets replayed. Same as if x > 0 on a tracer.wp.atomic_max into a small "peak" array, and you read it once at the end. Don't read the counter after every call.wp.Tape doesn't trace anything. It only keeps a Python list of which kernels were launched, and tape.backward() launches their adjoint kernels in reverse order. Because the list is plain Python and the adjoint launches are ordinary launches, the whole forward + backward can be captured into one graph.
.grad arrays accumulate: each backward adds into them, the same way cotangents sum when a value is used twice. If you forget tape.zero() between calls, you get silently wrong gradients. Nothing errors out.
To differentiate through a loop of in-place steps without keeping a tape over the whole loop, you can write the reverse pass by hand:
jax.checkpoint..gradtape.backward().grad, which becomes the incoming gradient for the step beforeThat's a reverse scan written as a Python loop. Wrap the forward and backward in a jax.custom_vjp and the JAX code calling it never knows Warp is underneath.
XLA_PYTHON_CLIENT_PREALLOCATE=false before importing JAX.jit-ed function, except Warp caches the compiled kernels on disk, so the next run of the script starts fast. With hundreds of kernels in a MuJoCo Warp step, the first run can take minutes. It's also why the warm-up body() comes before the capture: compiling inside a capture would break it. (JAX can also save cache to disk but needs to be activated through a flag).Fixed sizes, but not dense work. Like MJX, MJWarp allocates everything up front: naconmax contacts and njmax constraint rows per batch. The difference is that the kernels read the real counts on the device and skip the empty slots. MJX has to compute over the full fixed shape. This is also why MJWarp is graph-friendly: the number of launches never depends on how many contacts there are.