Skip to content

Add CPU support for JAX FFI (GH-1661)#1662

Draft
shi-eric wants to merge 1 commit into
NVIDIA:mainfrom
shi-eric:ershi/jax-cpu-ffi
Draft

Add CPU support for JAX FFI (GH-1661)#1662
shi-eric wants to merge 1 commit into
NVIDIA:mainfrom
shi-eric:ershi/jax-cpu-ffi

Conversation

@shi-eric

@shi-eric shi-eric commented Jul 17, 2026

Copy link
Copy Markdown
Contributor

Description

Add CPU execution support for wp.jax_kernel() and wp.jax_callable()
through JAX FFI. The same wrapper is registered for CPU and CUDA, allowing
JAX to select the implementation from the device used during lowering.

Closes #1661.

Related work: #1446 covers the same feature area. This implementation was
developed independently. Its CPU shaped-tile correction follows the
block_dim=1 issue reported while reviewing that PR:
#1446 (comment)

Changes

  • Register FFI targets for CPU and CUDA without initializing either backend.
  • Execute Warp kernels and callables directly against zero-copy JAX CPU
    buffers, while preserving existing CUDA behavior.
  • Make module preloading device-aware and compile CPU FFI kernels with a
    single-lane block dimension so shaped tile operations cover their full
    extent.
  • Validate CPU graph-mode behavior and report unsupported CUDA execution
    clearly when Warp lacks CUDA support.
  • Add JAX to the relevant CI test environments and cover CPU, mixed-device,
    preloading-disabled, and shaped-tile user workflows.
  • Document automatic device selection and CPU-specific behavior, and add an
    Unreleased changelog entry.

Checklist

  • I am familiar with the Contributing Guidelines.
  • New or existing tests cover these changes.
  • The documentation is up to date with these changes.
  • CHANGELOG.md is updated for any user-facing changes under the Unreleased section.

Validation summary

  • Ran the complete JAX interop test module across CPU and two CUDA devices:
    146 tests passed, with 12 expected skips.
  • Re-ran the CPU subprocess path, including public execution with module
    preloading disabled, plus the CPU shaped-tile, CPU callable graph-mode, and
    mixed CPU/CUDA regression tests. All four passed.
  • Ran pre-commit on every changed source and documentation file; all hooks
    passed.
  • Built the HTML documentation successfully with no Sphinx warnings.

Bug fix

Without the CPU-specific block dimension, shaped tile stores only populate the
first element:

import jax
import numpy as np
import warp as wp

TILE_SIZE = 64

@wp.kernel
def fill(output: wp.array(dtype=float)):
    tile = wp.tile_ones(dtype=float, shape=TILE_SIZE)
    wp.tile_store(output, tile)

run = jax.jit(wp.jax_kernel(fill, launch_dims=1, output_dims=TILE_SIZE))
with jax.default_device(jax.devices("cpu")[0]):
    result = run()[0]

np.testing.assert_array_equal(result, np.ones(TILE_SIZE, dtype=np.float32))

New feature / enhancement

The same JAX FFI wrapper can now execute on CPU or CUDA according to the input
device:

import jax
import numpy as np
import warp as wp

@wp.kernel
def triple(x: wp.array(dtype=float), y: wp.array(dtype=float)):
    i = wp.tid()
    y[i] = 3.0 * x[i]

run = jax.jit(wp.jax_kernel(triple))
x = np.arange(64, dtype=np.float32)
cpu_result = run(jax.device_put(x, jax.devices("cpu")[0]))[0]

Summary by CodeRabbit

  • New Features
    • Enabled CPU execution for wp.jax_kernel() and wp.jax_callable(), with automatic dispatch based on the JAX device.
    • Improved module preloading to support CPU + CUDA (best-effort device mapping).
  • Documentation
    • Expanded the JAX interoperability guide with device-selection behavior and clearer CPU vs CUDA graph-mode/capture rules.
  • Tests
    • Added CPU interop tests, mixed CPU/CUDA validation, sharded CPU execution, and module preload/module-load failure coverage.
  • Chores
    • Updated CI to install the correct JAX extras per platform/GPU job.
  • Changelog
    • Recorded the new CPU support for wp.jax_kernel() and wp.jax_callable().

@coderabbitai

coderabbitai Bot commented Jul 17, 2026

Copy link
Copy Markdown

Review Change Stack

Note

Reviews paused

It looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the reviews.auto_review.auto_pause_after_reviewed_commits setting.

Use the following commands to manage reviews:

  • @coderabbitai resume to resume automatic reviews.
  • @coderabbitai review to trigger a single review.

Use the checkboxes below for quick actions:

  • ▶️ Resume reviews
  • 🔍 Trigger review
📝 Walkthrough

Walkthrough

JAX FFI wrappers now support CPU and CUDA execution selected by JAX, with platform-specific registration, module preloading, graph-mode validation, expanded interoperability tests, updated documentation, and CI JAX requirements.

Changes

JAX CPU FFI support

Layer / File(s) Summary
CI JAX platform requirements
.github/workflows/ci.yml, .gitlab-ci.yml
CI jobs now install platform-specific JAX extras requiring JAX 0.5 or newer.
Platform registration and module preload
warp/_src/jax/ffi.py
FFI targets are registered for CPU and CUDA, with preload handling for current and supported devices.
CPU and CUDA callback execution
warp/_src/jax/ffi.py, CHANGELOG.md, docs/user_guide/interoperability_jax.rst
Kernel and callable callbacks dispatch by platform, execute CPU paths through invoke, enforce CPU graph-mode restrictions, and document the updated behavior.
Interop validation
warp/tests/interop/*
Tests cover CPU execution, preload modes, graph modes, mixed devices, CUDA availability, subprocess execution, and JAX test registration.

Estimated code review effort: 4 (Complex) | ~60 minutes

Sequence Diagram(s)

sequenceDiagram
  participant JAX
  participant Warp FFI
  participant Warp Runtime
  participant CPU or CUDA Device
  JAX->>Warp FFI: Lower and invoke jax_kernel or jax_callable
  Warp FFI->>Warp Runtime: Select platform and load module
  Warp Runtime->>CPU or CUDA Device: Execute kernel or callable
  CPU or CUDA Device-->>JAX: Return output buffers
Loading

Suggested reviewers: nvlukasz, christophercrouzet

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name Status Explanation Resolution
Docstring Coverage ⚠️ Warning Docstring coverage is 0.00% which is insufficient. The required threshold is 80.00%. Write docstrings for the functions missing them to satisfy the coverage threshold.
✅ Passed checks (4 passed)
Check name Status Explanation
Description Check ✅ Passed Check skipped - CodeRabbit’s high-level summary is enabled.
Title check ✅ Passed The title clearly states the main change: adding CPU support for JAX FFI.
Linked Issues check ✅ Passed The code, docs, CI, and tests implement CPU JAX FFI support, preserve CUDA behavior, and cover the required CPU graph-mode and preload cases.
Out of Scope Changes check ✅ Passed The changes stay focused on CPU JAX FFI support and related CI, docs, and tests without introducing unrelated scope.
✨ Finishing Touches
🧪 Generate unit tests (beta)
  • Create PR with unit tests

Comment @coderabbitai help to get the list of available commands.

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 1

🧹 Nitpick comments (1)
warp/tests/interop/test_jax.py (1)

131-131: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Use regular unittest methods for tests without harness-selected devices.

The preload and CPU-subprocess tests ignore the injected device; make them TestJax methods rather than registering them with add_function_test().

As per coding guidelines, use standard unittest.TestCase methods for fixed-device tests and reserve add_function_test() for tests running across multiple devices via get_test_devices().

Also applies to: 184-184, 1290-1292, 2462-2475

🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@warp/tests/interop/test_jax.py` at line 131, Convert the preload and
CPU-subprocess tests, including test_ffi_module_preload_modes and the
additionally referenced tests, into methods on TestJax instead of registering
them through add_function_test(). Remove their harness registration and unused
injected device parameters, while preserving standard unittest behavior for
these fixed-device tests.

Source: Coding guidelines

🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@warp/tests/interop/test_jax.py`:
- Around line 1279-1285: Update the auxiliary subprocess invocation in the test
around subprocess.run to include an appropriate timeout value, ensuring JAX
initialization or compilation hangs terminate before the CI job timeout while
preserving the existing output capture and error-handling behavior.

---

Nitpick comments:
In `@warp/tests/interop/test_jax.py`:
- Line 131: Convert the preload and CPU-subprocess tests, including
test_ffi_module_preload_modes and the additionally referenced tests, into
methods on TestJax instead of registering them through add_function_test().
Remove their harness registration and unused injected device parameters, while
preserving standard unittest behavior for these fixed-device tests.
🪄 Autofix (Beta)

Fix all unresolved CodeRabbit comments on this PR:

  • Push a commit to this branch (recommended)
  • Create a new PR with the fixes

ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yml

Review profile: CHILL

Plan: Enterprise

Run ID: 1c570803-f266-453c-809f-2812af04ea44

📥 Commits

Reviewing files that changed from the base of the PR and between d8ed53d and b480f81.

📒 Files selected for processing (7)
  • .github/workflows/ci.yml
  • .gitlab-ci.yml
  • CHANGELOG.md
  • docs/user_guide/interoperability_jax.rst
  • warp/_src/jax/ffi.py
  • warp/tests/interop/aux_test_jax_cpu_ffi.py
  • warp/tests/interop/test_jax.py

Comment thread warp/tests/interop/test_jax.py
@greptile-apps

greptile-apps Bot commented Jul 17, 2026

Copy link
Copy Markdown

Greptile Summary

This PR adds CPU execution support for wp.jax_kernel() and wp.jax_callable() through JAX FFI by registering a unified callback for both the \"cpu\" and \"CUDA\" platforms, then dispatching inside the callback based on which platform XLA selects at lowering time.

  • Registers per-platform ctypes closures in _register_ffi_targets, compiles CPU kernels with block_dim=1 (fixing shaped-tile store correctness), and adds _load_ffi_module / _preload_ffi_module helpers shared by both FfiKernel and FfiCallable.
  • Refactors FfiCallable to extract _build_arg_list, adds a CPU-specific early-return path calling invoke, and guards WARP/WARP_STAGED/WARP_STAGED_EX graph modes as CUDA-only.
  • Adds comprehensive tests covering CPU-only subprocess execution, shaped-tile correctness, mixed CPU/CUDA dispatch, graph-mode validation, module-load failure surfacing, and preload-mode device selection.

Confidence Score: 5/5

Safe to merge — changes are well-scoped, thoroughly tested, and preserve all existing CUDA behavior.

Platform dispatch is clean: ctypes closures capture platform via default-argument (no late-binding bug), _load_ffi_module correctly derives block_dim from device type, and CPU early-return paths in FfiKernel and FfiCallable are symmetric with CUDA paths. Refactored _build_arg_list matches original inline code exactly. Module deduplication uses id(device) which is stable since Warp devices are singletons. Test coverage spans unit mocks, in-process CPU/CUDA dispatch, subprocess CPU isolation, and graph-mode boundary validation.

No files require special attention.

Important Files Changed

Filename Overview
warp/_src/jax/ffi.py Core CPU FFI implementation — platform dispatch, module loading helpers, CPU kernel invocation, and callable CPU path all look correct and consistent with the existing CUDA pattern.
warp/tests/interop/test_jax.py Test refactoring extends coverage to CPU and mixed-device scenarios; new TestJax tests cover the added surface well.
warp/tests/interop/aux_test_jax_cpu_ffi.py New auxiliary subprocess script correctly sets JAX_PLATFORMS and XLA_FLAGS before JAX initializes.
.github/workflows/ci.yml Adds per-matrix jax-extra values with correct CPU-only vs CUDA extras per platform.
.gitlab-ci.yml Pins jax>=0.5 consistently across all Linux test jobs.
docs/user_guide/interoperability/jax.rst Documentation accurately describes device-selection behavior and CPU-specific restrictions.

Reviews (9): Last reviewed commit: "Add CPU support for JAX FFI (GH-1661)" | Re-trigger Greptile

Comment thread warp/_src/jax/ffi.py
Comment thread warp/_src/jax/ffi.py

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🧹 Nitpick comments (1)
warp/tests/interop/test_jax.py (1)

2365-2382: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Assert that the unavailable CUDA backend is actually exercised.

The expected loaded-device list remains [warp_cpu] even if a refactor stops probing CUDA, so this does not verify the intended CUDA-initialization fallback. Record backend requests and assert that "cuda" was requested.

Proposed test strengthening
+        requested_backends = []
+
         class FakeJax:
             `@staticmethod`
             def local_devices(process_index=None, backend=None, host_id=None):
+                requested_backends.append(backend)
                 if backend == "cpu":
                     return [cpu_0, cpu_1]
                 if backend == "cuda":
                     raise RuntimeError("CUDA backend is unavailable")
@@
                 ffi_module._preload_ffi_module(module, wp.JaxModulePreloadMode.ALL_DEVICES)
             self.assertEqual(module.loaded_devices, [warp_cpu])
+            self.assertIn("cuda", requested_backends)
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@warp/tests/interop/test_jax.py` around lines 2365 - 2382, Strengthen the
all-devices test around FakeJax.local_devices by recording each requested
backend, then assert that "cuda" was requested in addition to preserving the
existing loaded_devices assertion. Keep the CUDA RuntimeError path exercised so
the test verifies fallback behavior rather than only the final CPU device list.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Nitpick comments:
In `@warp/tests/interop/test_jax.py`:
- Around line 2365-2382: Strengthen the all-devices test around
FakeJax.local_devices by recording each requested backend, then assert that
"cuda" was requested in addition to preserving the existing loaded_devices
assertion. Keep the CUDA RuntimeError path exercised so the test verifies
fallback behavior rather than only the final CPU device list.

ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yml

Review profile: CHILL

Plan: Enterprise

Run ID: 7e103a80-5c3e-457d-ab99-2e9ec0ef7577

📥 Commits

Reviewing files that changed from the base of the PR and between b480f81 and 335a120.

📒 Files selected for processing (2)
  • warp/_src/jax/ffi.py
  • warp/tests/interop/test_jax.py
🚧 Files skipped from review as they are similar to previous changes (1)
  • warp/_src/jax/ffi.py

@shi-eric
shi-eric force-pushed the ershi/jax-cpu-ffi branch from 335a120 to cead8ea Compare July 17, 2026 19:16

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 2

🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@warp/_src/jax/ffi.py`:
- Line 166: Update the preloading calls around module.load in the relevant FFI
initialization path to use _load_ffi_module() instead, including the additional
calls at the referenced locations. Preserve the existing device and
block-dimension arguments so cached module build failures propagate rather than
being silently ignored.

In `@warp/tests/interop/test_jax.py`:
- Line 1108: Update every skip message associated with the visible
_jax_version() decorators to capitalize the proper name, replacing “Jax version
too old” with “JAX version too old” consistently across all occurrences.
🪄 Autofix (Beta)

Fix all unresolved CodeRabbit comments on this PR:

  • Push a commit to this branch (recommended)
  • Create a new PR with the fixes

ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yml

Review profile: CHILL

Plan: Enterprise

Run ID: 70802b3e-4b65-4b92-a19a-c6105fb24285

📥 Commits

Reviewing files that changed from the base of the PR and between 335a120 and cead8ea.

📒 Files selected for processing (7)
  • .github/workflows/ci.yml
  • .gitlab-ci.yml
  • CHANGELOG.md
  • docs/user_guide/interoperability_jax.rst
  • warp/_src/jax/ffi.py
  • warp/tests/interop/aux_test_jax_cpu_ffi.py
  • warp/tests/interop/test_jax.py
🚧 Files skipped from review as they are similar to previous changes (5)
  • warp/tests/interop/aux_test_jax_cpu_ffi.py
  • CHANGELOG.md
  • .github/workflows/ci.yml
  • .gitlab-ci.yml
  • docs/user_guide/interoperability_jax.rst

Comment thread warp/_src/jax/ffi.py Outdated
Comment thread warp/tests/interop/test_jax.py Outdated
@shi-eric
shi-eric force-pushed the ershi/jax-cpu-ffi branch 6 times, most recently from a4c54a4 to 726ba2f Compare July 20, 2026 05:24
JAX FFI wrappers previously assumed CUDA execution, blocking CPU-only
applications and forcing mixed-device programs onto backend-specific
paths.

Dispatch callbacks using the device selected by JAX, compile CPU kernels
with a single-lane block dimension, and preserve CUDA graph caching and
module preloading behavior. Report module and backend failures through
clear FFI errors.

Cover CPU, mixed-device, graph replay, module loading, and CI paths, and
document the supported behavior.

Signed-off-by: Eric Shi <ershi@nvidia.com>
@shi-eric
shi-eric force-pushed the ershi/jax-cpu-ffi branch from 726ba2f to cdd26d6 Compare July 20, 2026 05:40
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

[REQ] Support JAX FFI calls on CPU

1 participant