Triton monolithic head-loss kernel (#507)#558
Merged
Conversation
jlamypoirier
added a commit
that referenced
this pull request
Jul 10, 2026
Bring main (via the updated compiled base) into the triton PR. The metrics rework and triton work merge cleanly; fix the leftover references to the renamed metrics enum (`GRPOMetricsLevel` -> `PolicyMetricsLevel`) in the triton child classification (`config.py`, `monolithic.py`). GSPO metrics are already rejected for any monolithic composite, so the triton path only rejects GRPO metrics (which it doesn't compute). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
jlamypoirier
force-pushed
the
jlp_monolithic_triton
branch
from
July 10, 2026 20:03
73683f7 to
284fe2f
Compare
jlamypoirier
changed the base branch from
jlp_monolithic_head_loss
to
jlp_monolithic_metrics
July 10, 2026 20:04
jlamypoirier
force-pushed
the
jlp_monolithic_metrics
branch
from
July 11, 2026 01:12
32c3943 to
2368650
Compare
jlamypoirier
force-pushed
the
jlp_monolithic_triton
branch
from
July 11, 2026 01:12
284fe2f to
8c9f052
Compare
jlamypoirier
force-pushed
the
jlp_monolithic_metrics
branch
from
July 13, 2026 15:53
2368650 to
208d509
Compare
jlamypoirier
force-pushed
the
jlp_monolithic_triton
branch
from
July 13, 2026 15:59
8c9f052 to
f3771ce
Compare
Base automatically changed from
jlp_monolithic_metrics
to
jlp_monolithic_head_loss
July 13, 2026 16:30
Add the triton implementation of the monolithic shared-softmax loss and its policy-metric support, on top of the compiled path: - The triton kernel fuses the label-based objective set (CE, z-loss, GRPO, GSPO) over one softmax, superposing their gradients. - Policy metrics: the kernel accumulates the entropy's `Σ exp·logits_norm` for free in the backward re-stream; the layer reduces the family eagerly. Metrics force the reduced forward pass so the softmax is available in every case. - `MonolithicLossConfig.use_triton` selects the path (default: use if available), with per-kind fusion limits validated at config time. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
jlamypoirier
force-pushed
the
jlp_monolithic_triton
branch
from
July 13, 2026 16:38
f3771ce to
2e9e206
Compare
…riton path (#507) - `compute_gspo_metrics`: fuse softmax + predicted-logits + entropy into one compiled pass (new `_gspo_metrics_forward`), leaving only the segment `index_add_` eager. It was fully eager, computing entropy as a separate reduction over a vocab-wide temporary — mirrors the compiled `compute_grpo_metrics`. - GSPO `fused_core`: reduce `entropy_per_token` inside the compiled boundary (only when metrics are requested) and hand it to `finish`, rather than holding `logits_norm` (a vocab-wide tensor) across the eager seam. Fixes the monolithic GSPO single-loss peak memory and the fused-metrics speed. - Monolithic triton kernel: skip the softmax and the backward `exp` for a masked row when only a single label-based loss is active (nothing else needs the softmax on a label-less row). Folded into the existing backward loop rather than an early `return`, which triton compiles into heavy register spills in this large kernel. Brings single-loss cross-entropy and GRPO to parity with the dedicated kernels. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
jlamypoirier
force-pushed
the
jlp_monolithic_triton
branch
from
July 13, 2026 23:32
c66f242 to
b64da2b
Compare
Route the triton monolithic path through per-loss methods instead of a central kind classification, mirroring the compiled `_monolithic_core` loop. Each combinable loss owns `triton_add_inputs` / `triton_seam` / `triton_finish` and a `triton_kind` on its config (the single source of truth consulted by both config validation and the runtime); the driver iterates its children filling a shared `_TritonContext`, with no per-kind branching or duplicated classification. The `@triton.jit` kernel keeps its fixed per-slot signature — the one boundary where the loss kinds are still named. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The ce+z monolithic triton tests never passed `grpo=` / `gspo_coeff=`, so those kernel slots were covered only by the full head tests. Add `_test_monolithic_grpo_loss_triton` and `_test_monolithic_gspo_loss_triton`, each checking a single policy loss through the monolithic kernel against the standalone python reference (GSPO's per-token coefficient produced by the eager segment seam, as the driver runs it), plus matching tensor-parallel subtests in the distributed harness. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The metrics' `Σ exp·logits_norm` is accumulated in the backward re-stream, which previously ran only when a gradient was requested — so a fused policy loss with metrics but no gradient (e.g. a zero-weight child logging diagnostics) crashed the triton path while the compiled path handled it. Run the re-stream when either a gradient or the metrics are requested, and gate only the gradient store on `grad_logits_ptr`, so the two backends agree. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
predicted_logit is already initialized to 0.0, so the else-branch reassigned the same value. No behavioral change. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Claude Opus 4.8 (1M context) note: authored this PR.
Follow-up to #549, stacked on
jlp_monolithic_head_loss. Gives themonolithicloss a single Triton kernel for its shared-softmax path, alongside the existingtorch.compilepath.What it does
The
monolithicloss can now run its shared softmax on one@triton_jitkernel viause_triton, covering the label-based objective set: cross-entropy (label),z_loss, the GRPO objective, and the GSPO objective. Any subset sharing one softmax works (CE+z, GRPO+z, GSPO+z, …), at most one of each kind.Distillation has no Triton fused kernel yet — it is rejected at config time when
use_triton=True, and falls back to the compiled path whenuse_tritonis left at itsNone(auto) default. DPO stays separate as before.Design
Every enabled loss stores its own forward scalar, but the gradients superpose into two per-row coefficients:
so the gradient-write loop is loss-agnostic and each loss's math only contributes to the coefficients (z-loss adds to
prob_coeffonly; CE/GRPO/GSPO share the same label column). The kernel is gated purely by the presence of each loss's input/output pointers — no separate enable flags.Three dispatch cases mirror the existing per-loss Triton kernels:
GSPO's seam-argument assembly is factored out of
finishso the compiled and Triton paths share it.Config
use_triton: bool | NoneonMonolithicLossConfig:True— rejects unsupported children (distillation, duplicate kinds) at validate time.None(default) — Triton on GPU only if every child is supported, else the compiled monolithic.False— always compiled.Validation
Validated on an H100 GPU (
FAST_LLM_SKIP_TRITON_AUTOTUNE=TRUE):test_lm_lossesfunctional — the monolithic kernel against the PyTorch and standalone-Triton references (rms-relative ~1e-7): CE / z / CE+z, plus the GRPO and GSPO objectives through their kernel slots, across masking, logit/grad scaling, fp16, gradient accumulation, and the looped multi-block path.test_lm_headfused_tritonconfigs — single-pass and the GSPO seam, across masking / splits / bf16 / logit scaling / softcap.world_size=2monolithic subtests for the tensor-parallel path (CE+z, GRPO, GSPO).The Triton interpreter (
TRITON_INTERPRET=1) can't run this kernel — the grid trips anInterpreterErroron aconstexpr-shaped dimension (reproduces onmain, so it's pre-existing) — so validation of the monolithic Triton path is GPU-only.Follow-ups: distillation in the Triton kernel, and the GRPO/GSPO metrics rework.
🤖 Generated with Claude Code