Skip to content

Opt-in fused triangle kernels (Pallas) for inference - #4

Merged
DimaMolod merged 8 commits into
mainfrom
exp/af3-fused-triangle-v3.0.4
Oct 7, 2026
Merged

DimaMolod merged 8 commits into
mainfrom
exp/af3-fused-triangle-v3.0.4

Conversation

@DimaMolod

Copy link
Copy Markdown
Contributor

Summary

Opt-in fused triangle kernels for inference. They are Anthropic's FlashPairformer Pallas kernels for TriangleMultiplication and GridSelfAttention, from uplifting-biomolecular-modeling @ f4f62fa6 (Apache-2.0). Off by default; with the defaults the model runs exactly the original code.

  • Vendored unchanged (commit 9aedcb1; trimul_pallas.py, triattn_pallas.py and NOTICE are byte-identical to the kit):
    • in src/alphafold3/jax/fused_triangle/, the place AF3 kept its own Pallas kernels before tokamax;
    • provenance in SOURCE.md.
  • Our code:
    • jax/fused_triangle/dispatch.py: per-device tile tables and size and memory limits. The tile rows are copied from the kit and credited.
    • model/network/fused_triangle.py: Haiku adapters that read the stock parameters, so the weights load unchanged. They use DeepMind's ending-node bias convention, not OpenFold's.
    • Two hooks in modules.py. init always runs the stock module.
  • Config, named like flash_attention_implementation:
    • triangle_multiplication_implementation: 'default' or 'pallas';
    • triangle_attention_implementation: 'default', 'pallas', 'pallas_tokamax_core' or 'auto';
    • plus fused_triangle_compute_capability and fused_triangle_memory_gib, which the caller sets.
    • Any shape, card or size the kernels don't serve runs the stock path. An explicit request that falls back logs a warning.
  • Device policy:
    • Pallas attention on compute capability 8.0, 9.0 and 12.0 (12.0 on the 9.0 tiles);
    • the kit's prologue and epilogue around tokamax's attention core on 8.6 and 8.9, where the kit's generic tiles are slower;
    • triangle multiplication on every card from 8.0;
    • sm_100 is not enabled, because nothing was measured there.
  • test_data/model_config.json gains the new fields.

Validation

This branch carries the same commits as #2, replayed onto the v3.0.4 update (#3); the fused code is byte-identical between the two. The GPU numbers below were measured on the v3.0.3 / jax 0.9.1 stack. The v3.0.4 / jax 0.10.2 stack is validated on CPU here, and gets the full GPU campaign on the rebuilt AlphaPulldown image before any release.

  • fused_triangle_test.py (absltest, added to CI): 90 tests, passing on jax 0.9.1 and on 0.10.2 / tokamax 0.0.12 on CPU in Pallas interpret mode. They cover:

    • parity with the stock layers and an f64 reference;
    • both trimul equations and both attention orientations;
    • padded and asymmetric masks, and C = 128 and C = 64;
    • determinism and finiteness;
    • that the kernels actually run;
    • the fallbacks and their warnings;
    • the OpenFold orientation failing.
  • Layer level on GPU, 6 cards:

    • Accuracy: every fused cell passes the gate against an f64 reference (error at most 1.2× stock's), with leak, determinism and finiteness checks.
    • Speed per layer: trimul 1.24–1.71×; attention 1.11–1.54× on 8.0/9.0/12.0, and 1.03–1.36× with the tokamax core on 8.6/8.9.
  • End to end through AlphaPulldown, DeepMind weights, warm calls, fused vs stock on the same card:

    H100 A100 A40 L40S RTX 3090 RTX Pro 6000
    1.15–1.30× 1.18–1.66× 1.14–2.03× 1.11–1.34× 1.08–1.53× 1.12–1.23×
    • The largest complex each card can fold is unchanged.
    • With the flag off, outputs are bit-identical to the original fork on every card (deterministic XLA).
    • 12 heterodimers, paired seeds: Δranking and ΔDockQ are 0.00 on nearly all of them. Protein–RNA, protein–ligand and template inputs match too.

🤖 Generated with Claude Code

DimaMolod and others added 8 commits October 6, 2026 16:22
Anthropic's FlashPairformer Pallas kernels for TriangleMultiplication and
GridSelfAttention, with their serve layer (tile tables, served-shape
checks), copied verbatim from uplifting-biomolecular-modeling @ f4f62fa6
(Apache-2.0). Nothing imports them yet; later commits adapt and wire
them, so each change shows as a diff against the original.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
AlphaFold 3 keeps its own JAX building blocks under alphafold3/jax/, next
to geometry/. Move the two vendored Pallas kernels, the kit's NOTICE and
SOURCE.md there with git mv; the kernel files stay byte-identical to the
kit at f4f62fa6. The package __init__ re-exports nothing so that the
Triton backend is imported only once a kernel is selected. The Haiku
adapters now import the kernels by absolute path.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Replace fused_triangle_multiplication (bool) and fused_triangle_attention
('off'/'auto'/'pallas'/'tokamax') with fields named and valued like
flash_attention_implementation:

  triangle_multiplication_implementation: 'default' | 'pallas'
  triangle_attention_implementation:
      'default' | 'pallas' | 'pallas_tokamax_core' | 'auto'

Old values map off -> default, True/pallas -> pallas, tokamax ->
pallas_tokamax_core, auto -> auto. Dispatch reports the same names, and
an unknown multiplication implementation now raises like the attention
one. fused_triangle_compute_capability and fused_triangle_memory_gib are
unchanged.

Add the four GlobalConfig fields to the model_config.json golden; the
fields added with the adapters had left run_alphafold_data_test's
test_config failing.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
fpf_pallas_serve.py had been cut from the kit's 820 lines to 131 lines of
this fork's own policy while keeping the kit's file name. Remove it from
the vendored set and put the policy in alphafold3/jax/fused_triangle/
dispatch.py, in AlphaFold 3 style: device_policy, select_implementation
and the two tile lookups. The dispatch takes the requested
implementation and a DevicePolicy instead of reading GlobalConfig, so
alphafold3.jax does not depend on alphafold3.model. Only the bf16 tile
rows are copied from the kit, credited in a comment.

Decisions are unchanged; select_implementation now names the original
body 'default' (was 'stock'), with reason 'not_requested' (was 'off').
Checked against the removed file over 1.8M argument combinations.

SOURCE.md now lists the three files vendored unchanged, the file whose
tile rows were copied, the copyright holder, and that NOTICE's paths
are paths of the kit repository.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Move the Haiku adapters out of the fused_triangle package __init__ into
the module alphafold3/model/network/fused_triangle.py (AlphaFold 3
__init__ files only re-export), and restyle them like modules.py: 2-space
indent, absolute imports, licence header, Google docstrings, type hints,
descriptive names, pyink formatting at 80 columns. The adapters take the
module settings they need (equation, use_glu_kernel, num_head) instead of
the module config; the unused key= parameter and the lazy haiku, jnp, hm
and tokamax imports are gone.

modules.py imports the adapter module at the top and formats the hooks
like the surrounding code; the hooks stay where they were and
initialisation stays on the original body. Only the two kernel modules
are imported lazily, so importing modules.py does not import them; the
lazy-import test now checks exactly that. (jax.experimental.pallas.triton
is already imported by tokamax, through model_config.)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Move tests/test_fused_triangle.py beside the module it tests, as
src/alphafold3/model/network/fused_triangle_test.py (run_alphafold_test.py
sits beside run_alphafold.py the same way), and convert it from pytest to
absltest + parameterized with the AlphaFold 3 licence header, 2-space
indent and 80 columns. tests/ is gone.

The adapter test that made seven kinds of assertion per case is split
into focused tests: parameter compatibility at init, finite output,
nonzero-output parity, determinism, padding invariance, the
pallas_tokamax_core path, and the float32 fallback. A new check that the
fused config actually traces pallas_call guards the GlobalConfig to
dispatch wiring, which parity alone cannot see. Shared reference outputs
are cached, so the 80 tests take about as long as the old 28.

CI already runs CPU absltest files with `uv run python`, with the same
JAX, tokamax and Haiku versions the tests were verified with, so add the
file as a step.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
When triangle_multiplication_implementation or
triangle_attention_implementation asks for a fused kernel but a layer
runs the original module body, log an absl warning at trace time that
names the requested implementation, the activation shape and the reason
(dispatch reasons, plus the adapters' glu_kernel_disabled,
mask_dtype_mismatch and unsupported_equation). Each reason is logged
once per process. Nothing raises; the default implementation, and layers
that do run a kernel (including their initialisation), log nothing.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@DimaMolod
DimaMolod merged commit 7809148 into main Oct 7, 2026
1 check passed
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.

1 participant