Skip to content

[envpool] migrate XLA custom calls to OpenXLA FFI - #343

Merged
Trinkle23897 merged 3 commits into
mainfrom
jiayi/xla-openxla-ffi-source
Mar 23, 2026
Merged

[envpool] migrate XLA custom calls to OpenXLA FFI#343
Trinkle23897 merged 3 commits into
mainfrom
jiayi/xla-openxla-ffi-source

Conversation

@Trinkle23897

Copy link
Copy Markdown
Collaborator

Summary

  • Problem: EnvPool's JAX custom call path still relied on the legacy untyped API on GPU, and the new FFI headers were wired through Python-version-specific jaxlib wheel repos.
  • Scope: Switch EnvPool custom calls to typed OpenXLA FFI and pin the exported FFI headers from the GitHub openxla/xla source snapshot that jax-v0.9.2 already references.
  • Outcome: GPU custom calls no longer go through the legacy API path, and Bazel no longer depends on dynamically extracted pip_requirements_*_jaxlib headers for C++ compilation.

This updates the XLA integration to match current JAX FFI expectations while keeping the header dependency sourced from upstream GitHub code.

Technical Details

  • Approach: Replace the legacy custom call ABI with typed FFI handlers in C++, pass the handle as an FFI attr from Python, and fetch only the xla/ffi/api subtree from the XLA source commit pinned by jax-v0.9.2.
  • Code pointers:
    • envpool/core/xla_template.h: rewrites the custom call bridge to use xla::ffi typed handlers for CPU and GPU.
    • envpool/python/xla_template.py: registers API version 1 FFI targets and threads the custom handle through attrs with input/output aliasing.
    • envpool/workspace0.bzl: adds the pinned openxla_ffi_headers source archive.
    • third_party/openxla_ffi/ffi_api.BUILD: exposes only the exported FFI headers with the expected include prefix.
  • Notes: the external FFI headers come from openxla/xla@187a5eb58277a85847d1516bd1e20b7faf03d5ef, which is the XLA revision referenced by jax-v0.9.2.

Test Plan

Automated

  • brix ssh dev-0 -C -- 'cd /root/code/envpool && USE_BAZEL_VERSION=8.6.0 bazel --output_user_root=/tmp/envpool-openxla-source-subtree-bazel build //envpool/core:py_envpool --@rules_python//python/config_settings:python_version=3.11': passed
  • brix ssh dev-0 -C -- 'cd /root/code/envpool && USE_BAZEL_VERSION=8.6.0 bazel --output_user_root=/tmp/envpool-openxla-source-subtree-bazel build //envpool/core:py_envpool --@rules_python//python/config_settings:python_version=3.12': passed
  • brix ssh dev-0 -C -- 'cd /root/code/envpool && USE_BAZEL_VERSION=8.6.0 bazel --output_user_root=/tmp/envpool-openxla-source-subtree-bazel build //envpool/core:py_envpool --@rules_python//python/config_settings:python_version=3.13': passed
  • brix ssh dev-0 -C -- 'cd /root/code/envpool && BAZELOPT="--output_user_root=/tmp/envpool-openxla-subtree-wheel-bazel2" make bazel-build': passed
  • brix ssh dev-0 -C -- 'cd /root/code/envpool && /tmp/envpool-issue303-fix-venv/bin/python examples/xla_step.py': passed earlier during the typed FFI migration validation on dev-0

Suggested Manual

  • pip install <built wheel> "jax[cuda12]" gym packaging: verify install against the same JAX family used in the repro.
  • python examples/xla_step.py: confirm the GPU path runs without the old API_VERSION_ORIGINAL warning.

@Trinkle23897 Trinkle23897 mentioned this pull request Mar 23, 2026
3 tasks
@Trinkle23897 Trinkle23897 linked an issue Mar 23, 2026 that may be closed by this pull request
3 tasks

@chatgpt-codex-connector chatgpt-codex-connector 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.

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: e6fd862740

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

  • Open a pull request for review
  • Mark a draft as ready
  • Comment "@codex review".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".

Comment on lines +78 to +81
handle_value = int.from_bytes(handle, byteorder=sys.byteorder, signed=False)

def call(*args: Any) -> Any:
return xla_func(*args)
return xla_func(*args, handle=handle_value)

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

P1 Badge Pass the FFI handle attribute as an explicit int64

ffi_call() now threads the env handle through a named attribute, and CustomCall::ResolveHandle() decodes that attribute with attrs.get<std::int64_t>("handle"). Here we pass handle_value as a plain Python int; on the default JAX configuration (jax_enable_x64=False), Python ints are 32-bit, so the typed-FFI attribute is either rejected as the wrong type or truncated before the pointer cast on 64-bit hosts. In that configuration every send/recv/step call will fail unless users globally enable x64. Use an explicit np.int64/np.uint64 scalar for the attribute.

Useful? React with 👍 / 👎.

@Trinkle23897
Trinkle23897 merged commit a8c4620 into main Mar 23, 2026
2 of 4 checks passed
@Trinkle23897
Trinkle23897 deleted the jiayi/xla-openxla-ffi-source branch March 23, 2026 16:51
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.

[BUG] Using XLA crashes out

1 participant