Skip to content

[Feat] Added JAX-Triton bridge for ROCm - #649

Merged
AllenFarcas merged 12 commits into
devfrom
alfarcas/jax-triton-bridge
Aug 21, 2026
Merged

[Feat] Added JAX-Triton bridge for ROCm#649
AllenFarcas merged 12 commits into
devfrom
alfarcas/jax-triton-bridge

Conversation

@AllenFarcas

@AllenFarcas AllenFarcas commented Jun 24, 2026

Copy link
Copy Markdown
Contributor

Description

Extend TE's JAX Triton custom-call bridge to compile and dispatch AMD ROCm (HSACO) and Gluon kernels. This PR enables AMD's layout-explicit Gluon kernels to be called from JAX, mirroring NVIDIA's existing support.

Fixes https://github.com/ROCm/frameworks-internal/issues/16044

Type of change

  • Documentation change (change only to the documentation, either a fix or a new content)
  • Bug fix (non-breaking change which fixes an issue)
  • New feature (non-breaking change which adds functionality)
  • Breaking change (fix or feature that would cause existing functionality to not work as expected)
  • Infra/Build change
  • Code refactoring

Changes

Please list the changes introduced in this PR:

  • Add ROCm/HIP backend path emitting HSACO alongside the existing CUDA/PTX path.
  • Support @gluon.jit kernels via GluonASTSource with a full constexpr-marked signature.
  • Pass HSACO as a temp-file path (nanobind std::string), not raw bytes.
  • Add optional num_warps/num_stages for non-autotuned Gluon layout matching.
  • Add version-guarded Gluon binding test to test_triton_custom_calls.py (requires Triton base > 3.4.0).

Checklist:

  • I have read and followed the contributing guidelines
  • The functionality is complete
  • I have commented my code, particularly in hard-to-understand areas
  • I have made corresponding changes to the documentation
  • My changes generate no new warnings
  • I have added tests that prove my fix is effective or that my feature works
  • New and existing unit tests pass locally with my changes

@AllenFarcas
AllenFarcas marked this pull request as ready for review June 29, 2026 15:57
Comment thread transformer_engine/jax/triton_extensions/utils.py Outdated
Comment thread transformer_engine/jax/triton_extensions/utils.py Outdated
Comment thread transformer_engine/jax/triton_extensions/utils.py Outdated
@AllenFarcas
AllenFarcas requested a review from ipanfilo July 6, 2026 15:41
@AllenFarcas AllenFarcas self-assigned this Jul 6, 2026
@AllenFarcas AllenFarcas added ci-level 3 CI test level 3 and removed ci-level 3 CI test level 3 labels Jul 6, 2026
@AllenFarcas
AllenFarcas requested a review from Micky774 July 8, 2026 16:18
@github-actions

github-actions Bot commented Jul 9, 2026

Copy link
Copy Markdown

Claude Walkthrough

Intent. Extend TE's JAX/Triton custom-call bridge in transformer_engine/jax/triton_extensions/utils.py so it also compiles and launches ROCm kernels (emitting HSACO instead of PTX) and accepts Gluon (@gluon.jit) sources alongside plain @triton.jit. This unlocks AMD's layout-explicit Gluon kernels through the same JAX lowering path NVIDIA already uses.

Key changes.

  • Dual-backend compile in compile_triton at transformer_engine/jax/triton_extensions/utils.py:294 — ROCm branch uses triton.runtime.driver.active.get_current_target() + make_backend(...).parse_options(...) and the CUDA branch keeps the existing CUDAOptions/GPUTarget("cuda", ...) path.
  • Gluon source path: when kernel_fn.is_gluon() is true, swaps tc.ASTSource for triton.experimental.gluon._runtime.GluonASTSource and, because Gluon requires every constexpr in the signature, back-fills any missing constants (utils.py:369).
  • HSACO delivery via temp file: because TritonKernel's binary field is a std::string and nanobind won't coerce raw bytes, HSACO is written to a process-scoped TemporaryDirectory (_hsaco_dir, utils.py:216) and the path is passed instead of the blob (utils.py:410).
  • triton_call_lowering now takes optional num_warps / num_stages (utils.py:453) so callers matching a Gluon BlockedLayout can pin the launch geometry; default num_warps becomes 4 on ROCm (wavefronts are 32 or 64 lanes, so the CUDA-tuned default of 32 would exceed 1024 threads/block).
  • Cache key now folds in is_hip and is_gluon so CUDA/ROCm and Triton/Gluon variants don't collide in _TRITON_KERNEL_CACHE.
  • Adds two Gluon test classes to tests/jax/test_triton_custom_calls.py behind a HAS_GLUON guard keyed on Triton > 3.4.0.

Walkthrough.

transformer_engine/jax/triton_extensions/utils.py — the compile function was previously CUDA-only (cb.CUDAOptionstc.compile(..., target=GPUTarget("cuda", ...))compiled.asm["ptx"]). It's now split by is_hip_extension(). On ROCm the target is discovered from the active driver so warp_size (32 or 64 depending on gfx arch) can be threaded into parse_options; binary_key becomes backend.binary_ext (i.e. "hsaco"). Source construction is orthogonal to backend: Gluon detection (hasattr(kernel_fn, "is_gluon") and kernel_fn.is_gluon()) selects GluonASTSource, which unlike ASTSource requires every constexpr — including those already handled as constants — to appear in the signature, hence the explicit back-fill loop over kernel_fn.arg_names. The GluonASTSource import is wrapped in a try/except ImportError that raises a targeted message pointing users at a Triton upgrade. After tc.compile, the ROCm branch persists HSACO bytes to a temp file (_HSACO_TMPDIR, created lazily and torn down at interpreter exit) and swaps the path in where PTX text used to go, so both TritonKernel(...) signatures downstream (the pre- and post-3.6 shapes) work unchanged. triton_call_lowering grows two optional launch knobs; if unset, num_warps defaults to 4 on ROCm and 32 on CUDA, matching each backend's "warps × warp_size ≤ 1024 threads" ceiling.

tests/jax/test_triton_custom_calls.py — adds a Gluon _double_kernel that doubles its input using an explicit BlockedLayout whose warps_per_cta must match the launch num_warps. Two test classes exercise the bridge: TestGluonBinding for the JIT path (passes num_warps=NUM_WARPS) and TestGluonAutotunedBinding for the autotune path (each triton.Config carries BLOCK_SIZE/NUM_WARPS as constexpr kwargs plus a matching num_warps=, so the layout stays consistent across configs). Both are gated by HAS_GLUON, computed by parsing triton.__version__ with packaging.Version and comparing .release against (3, 4, 0) — using .release rather than raw string comparison so a +rocm.git local version segment can't spuriously satisfy the check. WARP_SIZE is read from the active Triton target, keeping the test agnostic to gfx942 vs. gfx950 vs. NVIDIA.

Testing. New TestGluonBinding (JIT) and TestGluonAutotunedBinding (autotune) classes cover the Gluon end-to-end path; each parametrizes on shapes (1024, 1024) and (1000, 1000) (the latter exercises the trailing masked block) and checks bitwise-close equality against x * 2.0. Both are silently skipped when Gluon is missing or the Triton base version is <= 3.4.0. Existing TestTritonBinding cases are unchanged, so the CUDA/PTX path continues to be exercised.

Notes for reviewers.

  • The HSACO temp file is not deleted eagerly — _HSACO_TMPDIR lives for the process lifetime. Deliberate (the plugin loads the file at launch time, potentially long after compile_triton returns), but long-running processes accumulate one file per unique cache key.
  • Gluon detection uses duck typing (hasattr(kernel_fn, "is_gluon")); autotuned kernels forward is_gluon to the underlying JITFunction, so autotuned Gluon kernels are correctly identified too.
  • The ROCm num_warps default changed from 32 to 4. Existing ROCm callers that relied on the implicit 32 will now launch with fewer warps unless they pass num_warps explicitly — worth flagging in the CHANGELOG if any downstream users depend on the old value.
  • compute_capability is still passed through to TritonKernel(...) on ROCm even though it's meaningless there; the JAX-side plugin ignores it for HIP targets, but future cleanup could carry the ROCm arch string instead.
  • Gluon on ROCm requires an upstream Triton make_ir fix that landed after 3.4.0 — the HAS_GLUON guard encodes that requirement.

Generated by Claude. To request a code review, comment /claude review.

@github-actions

github-actions Bot commented Jul 9, 2026

Copy link
Copy Markdown

Claude review

Reviewed the PR-specific diff (git diff <merge-base>..HEAD, two files):

  • transformer_engine/jax/triton_extensions/utils.py
  • tests/jax/test_triton_custom_calls.py

Verdict: looks clean. The ROCm/HSACO path, Gluon GluonASTSource handling, and the temp-file HSACO plumbing all read correctly, and the PR-level changes address the earlier reviewer feedback in-place (warp-size comment reworded, num_warps=4 default justified in a code comment).

Copyright headers: OK — both modified files carry an AMD 2026 line above the preserved NVIDIA 2022-2026 line, matching the existing repo convention (This file was modified for portability to AMDGPU prefix + AMD copyright + untouched NVIDIA copyright).

No new inline findings.

@wenchenvincent

Copy link
Copy Markdown
Collaborator

@AllenFarcas Please resolve conflict and test against the latest CI image.

Comment thread transformer_engine/jax/triton_extensions/utils.py
Comment thread transformer_engine/jax/triton_extensions/utils.py Outdated
Comment thread tests/jax/test_triton_custom_calls.py
Comment thread tests/jax/test_triton_custom_calls.py Outdated
@github-actions

Copy link
Copy Markdown

Claude review (re-review)

Re-reviewed the full PR diff against the current base (HEAD^1...HEAD^2, two files: transformer_engine/jax/triton_extensions/utils.py, tests/jax/test_triton_custom_calls.py). The PR head is unchanged since the last review, but the base has moved, so this pass re-read the whole diff rather than an incremental slice — it surfaced four findings the earlier pass missed.

Verdict: the ROCm/HSACO and Gluon compile paths are structurally sound and the CUDA path stays guarded and unchanged. Four issues worth addressing before merge, posted inline: one ROCm correctness risk around how the HSACO binary is delivered, one import-time fragility on ROCm-only Triton builds, one Black violation that will fail the lint gate, and one test-visibility gap that lets the new Gluon coverage silently not run.

Copyright headers: OK — both modified files carry an AMD 2026 line above the preserved NVIDIA 2022-2026 line, with the AMDGPU portability marker.

Note: the PR is currently in a conflicting state against dev.

AllenFarcas and others added 3 commits August 12, 2026 15:42
Write HSACO to a persistent, content-addressed path. The blob path is
serialized into the custom call, so a JAX persistent-cache hit in a later
process replayed a deleted temp dir. Naming the file by its own digest also
avoids the kernel cache key, which covers neither kernel source nor Triton
version and would serve a stale binary after an edit.

Import the NVIDIA Triton backend lazily. A Triton built for AMD only ships no
triton/backends/nvidia, so the module-scope import failed and the ROCm path
never ran.

Collect the Gluon tests always and skip them with a reason. Defined inside
"if HAS_GLUON" they were never collected, so the suite passed having run
nothing. Narrow the probe's except for the same reason. Guard the triton
import: require_triton_or_skip_test_file only checks the JAX version.

Run test_triton_custom_calls.py in CI; it was in no leg.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The persistent content-addressed store was written for a failure that does not
occur. Measured on gfx950: the ROCm plugin unlinks the blob once it has loaded
it, and JAX derives its persistent-cache key from the lowered HLO, so lowering
reruns in every process and rewrites the file before launch. A cached
executable therefore never replays a stale path. Persisting the blobs only
leaked the ones no kernel loaded, which the temp dir had been reclaiming.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
@AllenFarcas AllenFarcas added the ci-level 3 CI test level 3 label Aug 12, 2026
AllenFarcas and others added 3 commits August 17, 2026 14:57
The bridge's own tests covered a synthetic kernel; the code that ships through
it, transformer_engine.jax.permutation, ran in no CI leg. Wiring those two
files in exposed a launch failure with multiple local devices:

  INTERNAL: /tmp/te_jax_hsaco_<rand>/<rand>.hsaco; No such file or directory

The plugin's HIP branch reads the blob from its path and unlinks it before
taking the lock guarding its module cache, so two device threads reaching an
unseen kernel together race and the loser opens a file the winner deleted.
Launching each freshly compiled kernel once from a single device inserts the
cache entry before any sharded executable can race for it. Measured on the
distributed test at L1: 3 failures in 3 runs without it, none in 3 runs with
it. NVTE_JAX_TRITON_PREWARM=0 disables it; the whole block goes away once the
plugin does its read and unlink under the writer lock.

The prewarm reuses the lowering closure rather than rebuilding the call, so
the kernel cache key cannot drift, and it zero-fills operands because the
gather kernels index with them.

Adds the L1 parameter sets both files lacked, without which the multi-GPU leg
fails collection.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The prewarm guarded itself with a single _PREWARM_ACTIVE flag and fired only
when the kernel cache had grown during this lowering. Both fail once more than
one thread lowers at a time: the flag is a reentrancy guard rather than mutual
exclusion, so a thread lowering one kernel suppressed the load for another, and
a thread that lost the race to populate the cache saw no growth and skipped the
load for the kernel it was about to launch.

Key the load on a digest of the serialized call proto, record loaded keys in a
process-wide set, and hold a per-key lock across the launch so a second lowering
of the same kernel waits for it rather than racing the plugin. The key is
recorded in a finally: a persistent failure would never succeed on retry and
would repeat its warning once per lowering.

_hsaco_dir and _prewarm_primitive had the same unguarded lazy-init shape; two
threads could build two temp dirs, whichever lost taking the winner's blobs with
it at finalization. Both now initialize under a lock, and the locks are replaced
in a forked child, where one held at fork would never be released. The child
detaches the inherited TemporaryDirectory before dropping it, so the parent's
blobs survive, and clears the loaded set, since the plugin's module cache does
not cross the fork.

Rebuild the FFI lowering inside the rule. ffi_lowering captured an ir.StringAttr
belonging to whichever MLIR context was current when the rule was built, and the
prewarm replays that rule in a context of its own.

Call _require_triton_permutation from the token_dispatch, token_combine and
sort_chunks_by_index forward rules. The guard existed but nothing invoked it, so
a Triton-less install reached the kernel wrappers with the names bound to None.
Widen both import guards to RuntimeError and ValueError, which a mismatched
Triton raises at import, and chain the original cause onto the message.

Import triton inside the ROCm branch of compile_triton, its only use, which also
clears the two pylint messages the module-scope import was adding.

Measured with the prewarm on: the permutation suite passes 24 of 24 at L1 in
each of three runs and 72 of 72 at L2; with it off, all three L1 runs fail. A
four-thread concurrent lowering is clean in five runs. Eight threads on one
kernel produce a single load and no early return; on eight distinct kernels they
do not serialize.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

@wangye805 wangye805 left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

LGTM with minor comments

Comment thread transformer_engine/jax/triton_extensions/utils.py
Comment thread transformer_engine/jax/triton_extensions/utils.py
Comment thread transformer_engine/jax/triton_extensions/utils.py Outdated
Comment thread transformer_engine/jax/triton_extensions/utils.py
Comment thread transformer_engine/jax/triton_extensions/utils.py
Import the NVIDIA Triton backend at module scope again, under
`if not is_hip_extension()`. Moving it into compile_triton kept a Triton built
for AMD only importable, but it also deferred the failure a CUDA build gets from
a broken backend to the first lowering; the guard keeps the import-time error
while leaving the AMD case alone. pylint cannot see that the name is bound
wherever it is read, hence the suppression at the use site.

Register the fork handler only on ROCm, so a CUDA build no longer installs a
callback for state it never creates.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
@AllenFarcas
AllenFarcas merged commit 15c8608 into dev Aug 21, 2026
4 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

ci-level 3 CI test level 3

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants