Skip to content

[feat] pack the GenRec training forward through FlashAttention-2 - #663

Merged
tiankongdeguiji merged 18 commits into
alibaba:masterfrom
WhiteSwan1:feat/genrec_flash_attention
Sep 22, 2026
Merged

tiankongdeguiji merged 18 commits into
alibaba:masterfrom
WhiteSwan1:feat/genrec_flash_attention

Conversation

@WhiteSwan1

@WhiteSwan1 WhiteSwan1 commented Sep 8, 2026 •

Copy link
Copy Markdown
Collaborator

GenRec padded every variable-length prompt before the teacher-forced forward, so every row paid for the longest row in its batch. This runs that path on packed embeddings through the FlashAttention-2 varlen kernel instead: one concatenated stream, explicit cu_seq_lens boundaries, per-row position_ids, and response-only logits gathered straight out of the packed layout.

Measured

On 1× H20, GenRecCausalLMModel over Qwen2.5-0.5B:

batch 20, real length SDPA (the path this replaces) FLASH_ATTENTION_2
step time 448.8 ms 233.4 ms — 1.92×
peak memory 31,515 MiB 20,147 MiB — −36%

The margin grows with sequence length: 2.55× at 2048 tokens, where SDPA peaks at 76.8 GB on a 96 GB card against FA2's 41.1 GB.

Worth reading the breakdown before optimising further: after packing, attention is only 7.0% of the step at production length, and the linear layers are 60% running at ~90% of H20 BF16 peak. The win here is mostly not paying for padding — 43% of tokens at batch 20 — rather than a faster attention kernel.

Choosing the kernel

GenRecModelConfig.common.attn_kernel picks the backend, carrying the transformers name straight through to attn_implementation:

optional string attn_kernel = 6 [default = "sdpa"];

sdpa is the default because it is the backend whose dependency always ships — building a GenRec model then needs neither a GPU nor the flash_attn wheel, which keeps the CPU lane and export working. Training runs opt into flash_attention_2 explicitly. There is no AUTO and no silent fallback: a misconfiguration surfaces as the library's own error rather than a quietly slower path.

Each kernel gets the layout it can actually serve. Only the varlen kernel reads cu_seq_lens; every other kernel rebuilds the boundaries as a dense (1, 1, T, T) mask, which costs (ΣL)² where left-padding costs Σ(B·L²). So _forward branches — packed for flash_attention_2, left-padded for everything else. The supervised window is the same absolute positions either way, so the labels are built once and only the logits differ. Measured at 8 rows of 128: packed is 1,048,576 score-pairs behind a 1 MiB mask, padded is 131,072 behind no mask at all.

flash_attn is declared in requirements/cu126|cu129|cu130.txt beside the other self-hosted wheels — three lines each, one per supported Python (3.10/3.11/3.12), matching the coverage dynamicemb, faiss and fbgemm_gpu_hstu already have.

Test plan

  • Unit coverage for the packed varlen metadata, the suffix window and its labels, response-only loss and gradients, the short-row guard, and backbone construction under both kernels.
  • _RowIsolationCase asserts the invariant that matters — rewrite row 0 and row 1's logits must stay bit-identical — and runs under both kernels: the sdpa arm on the CPU lane on every PR, the flash arm on the GPU lane.
  • Verified the two branches agree to 5.96e-08 with identical labels.
  • GPU coverage for qwen2/qwen3 packed-vs-solo forwards, projected-slot gradients under BF16/FP16 and under fp32 masters with BF16 autocast, and ragged beam decoding.

Known gaps

  • The wheel is locked to the torch minor version through libtorch's C++ ABI, not just to CUDA: it links 70 c10::/at:: symbols against unversioned sonames, so a torch bump without a rebuild fails at import, not at install. Requires-Dist: torch carries no bound, so pip cannot catch the mismatch.
  • cu126 has no sm100/sm120 build, unlike cu129 and cu130. Blackwell on a CUDA 12.6 install falls back to SDPA.

FA3 was measured too and is not proposed here: +2.5% at production length, no memory saving, and no wheel — it needs a from-source Hopper build against the same torch minor ABI.

WhiteSwan1 and others added 10 commits September 3, 2026 07:54
Master's genrec/prompt rework (c3e9df8 "export the front-end and serving
contract for genrec") renamed the packed-prompt stream keys and re-homed the
prompt tests, so the packed FlashAttention work is carried over rather than
merged textually:

- assembler.py / assembler_test.py take master's rewrite verbatim. The PR's
  empty-prompt-body guard moved to GenRecCausalLMModel._forward, which is
  where the suffix window is indexed; the assembler is scripted for serving
  and cannot raise.
- genrec_causal_lm_model.py keeps the packed varlen forward, renamed onto
  master's stream keys (PROMPT_CU_SEQLENS -> CU_SEQLENS, PROMPT_INPUT_IDS ->
  INPUT_IDS, PROMPT_MAX_SEQLEN -> MAX_SEQLEN, PROMPT_RESPONSE_LENGTHS ->
  RESPONSE_LENGTHS) and on master's 2-tuple _left_pad_packed_inputs.
- tests/prompt_integration_test.py and tests/prompt_test_util.py are deleted
  with master. The one PR-only test that had no home there,
  test_packed_rows_match_solo_runs_and_backpropagate, is ported into
  models/genrec_model_test.py as PackedFlashAttentionTest.
- adds a short-row case for the relocated guard, whose coverage the move
  would otherwise have dropped.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
GenRec builds its backbone with attn_implementation="flash_attention_2",
which needs an NVIDIA GPU and the flash_attn wheel, so the tests that build
a model cannot run on the CPU lane (run.py with no --scope runs everything
not skipped, so a scope marker alone does not exclude them).

- declare flash_attn in requirements/cu126|cu129|cu130.txt, cp311 only,
  beside fbgemm_gpu_hstu
- mark BaseGenRecModelTest, GenRecFrontEndTest,
  test_training_forward_builds_no_cache and test_genrec_train_eval_export
  gpu-scoped and skip them without an NVIDIA GPU
- raise a clear error when the backbone runs in fp32 with no autocast,
  instead of letting the flash kernel fail on the dtype

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
flash_attention_2 has no CPU kernel, so a test that builds the model and
runs a forward has to put both the model and the batch on cuda; five did
not, and failed with "Could not run 'flash_attn::_flash_attn_varlen_forward'
with arguments from the 'CPU' backend" even on a GPU host.

Two of them also ran the backbone in fp32 with no autocast, which the new
dtype check rejects: the loss test now runs under bf16 autocast, matching
the fp32-masters setup lm_parameter_dtype documents, and the no-cache test
narrows its backbone to BF16.

Verified on 2x H20 (torch 2.13.0+cu126, transformers 5.17.0,
flash_attn 2.8.3.post1): 41 tests OK, including PackedFlashAttentionTest,
which asserts atol=0 row isolation and had never executed before.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
nv_gpu_unavailable only reports whether CUDA is present, so a GPU host
running a python the flash_attn wheel is not built for -- it ships cp311
only, while the project supports 3.10/3.11/3.12 -- runs these tests and
errors in init_backbone instead of skipping.

Add flash_attn_unavailable beside the other optional-wheel probes, in the
same find_spec idiom as cutlass_hstu_unavailable and faiss_unavailable, and
guard the four sites that build a flash_attention_2 backbone. Teach the
ci-scope lint the new token so a flash-only skip still has to declare a
gpu lane.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
flash_attention_2 was hard-coded, so building a GenRec model needed an
NVIDIA GPU and the flash_attn wheel -- which ships cp311 only -- even to
construct one for export or on a python the wheel is not built for.

Add common.attn_implementation, SDPA or FLASH_ATTENTION_2, defaulting to
SDPA, in the shape of the existing Kernel enum: the backend whose
dependency is always present is the default, and the optional-wheel one is
explicit opt-in. There is no AUTO and no silent fallback -- configuring
FLASH_ATTENTION_2 without the wheel raises and names the install path,
since a quiet downgrade would hide the regression this path exists to fix.

The packed forward is unchanged: sdpa carries the same row boundaries
through position_ids, and the varlen kwargs pass through it inert. Selecting
SDPA on CUDA also disables the cuDNN sdpa backend, which is eligible for the
packed mask and returns NaN losses on it.

The fp32 check now fires only on the flash path, which is the only one that
rejects fp32, and the packed-vs-solo row-isolation test runs on both
kernels.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
With SDPA the default, building a GenRec model needs neither a GPU nor the
flash_attn wheel, so the tests that only construct and run a model belong
back on the lane that runs every PR. This reverts the cuda moves and the
gpu markers on BaseGenRecModelTest, GenRecFrontEndTest,
test_training_forward_builds_no_cache and test_genrec_train_eval_export.

The packed-vs-solo row-isolation case splits in two: PackedSdpaAttentionTest
runs it in fp32 on CPU, and PackedFlashAttentionTest runs the same body in
bf16 on cuda behind the GPU and wheel probes. The leakage probe now guards
the feature on every PR rather than only where the wheel exists.

That makes the transformers floor load-bearing rather than advisory:
masking_utils.find_packed_sequence_indices, which carries the row
boundaries on every non-flash kernel, lands in 4.56. On 4.51.2 the sdpa
case fails -- rows attend across each other -- so requirements now say
transformers>=4.56.

Verified with no GPU and no wheel: 54 tests, OK, 4 skipped (the two flash
cases and the two flash beam cases).

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

The fp32 check asked self.lm.config._attn_implementation, which is a
private property over _attn_implementation_internal and the only thing
transformers exposes -- on a dependency with no upper bound. A rename there
would make the comparison silently False and the check would stop firing.

init_backbone already knows which kernel it was asked for, and since
FLASH_ATTENTION_2 raises rather than falling back when the wheel is absent,
what was asked for is what was built. Record it and read that instead.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
_forward packed every row into one sequence regardless of the configured
kernel, but only the flash varlen kernel reads cu_seq_lens. On sdpa the
boundaries come back as a dense (1, 1, T, T) mask, so the batch costs
(sum L)^2 where the left-padded path it replaced cost sum(B * L^2) -- B
times more attention work than before this branch, on the default kernel.
Measured at 8 rows of 128: packed 1,048,576 score-pairs and a 1 MiB mask,
padded 131,072 and no mask at all.

So branch: FLASH_ATTENTION_2 keeps the packed forward, every other kernel
gets _left_pad_packed_inputs, which _generate already used. The supervised
window is the same absolute positions either way, so the labels are built
once and only the logits split. Verified the two branches agree to 5.96e-08
with identical labels.

Drop both flash guards. The fp32 check and the wheel probe each replaced an
error transformers already raises -- "FlashAttention only support fp16 and
bf16 data type" and "the package for FlashAttention2 doesn't seem to be
installed" -- so they bought a nicer message and cost a per-step check and
an import probe. lm_parameter_dtype is left alone rather than silently
narrowed to BF16 when flash is configured.

Name the enum after the Kernel it mirrors: AttnImpl was the only *Impl of
the thirteen in tzrec/protos, where eight are bare nouns. AttnKernel, field
attn_kernel, and _attn_kernel now holds that enum instead of an HF string.

Alongside, in the forward: keep the window 2-D rather than flattening and
reshaping back, fold `columns` into `suffix_offsets`, drop a .contiguous()
that cannot copy, and replace a process-wide enable_cudnn_sdp(False) --
which disabled the backend for every other module in the process -- with a
scoped sdpa_kernel exclusion at the call that needs it.

And in the tests: move the row-isolation suite beside the forward it
exercises; delete three assertions that only pinned a mock, the file's own
stub, or a use_cache kwarg asserted next door; give create_tiny_causal_lm a
model_type, split create_genrec_test_prompt out of create_genrec_test_model,
and collapse _packed_model, _compiled_prompt and six copied arrange blocks
onto them.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
@WhiteSwan1 WhiteSwan1 changed the title [perf] run GenRec with packed FlashAttention [perf] pack the GenRec training forward through FlashAttention-2 Sep 18, 2026
@WhiteSwan1
WhiteSwan1 marked this pull request as ready for review September 18, 2026 03:42
@WhiteSwan1 WhiteSwan1 changed the title [perf] pack the GenRec training forward through FlashAttention-2 [feat] pack the GenRec training forward through FlashAttention-2 Sep 18, 2026
@tiankongdeguiji tiankongdeguiji added the claude-review Let Claude Review label Sep 21, 2026
@github-actions github-actions Bot removed the claude-review Let Claude Review label Sep 21, 2026
The sdpa_kernel block excluded CUDNN_ATTENTION around the LM call, to keep
the NaN losses cuDNN returns on a packed mask. It defended a configuration
this branch no longer produces: sdpa takes the left-padded path now, and the
NaN was only ever reported against packed-sdpa. It was also dead on the
flash branch, which never enters sdpa dispatch at all. An environment that
still needs it can set enable_cudnn_sdp(False) itself.

Delete test_loss_is_finite_and_backpropagates_into_the_backbone: its four
assertions are a subset of SdpaRowIsolationTest's, on the same production
path. Checked against 34 production mutations -- 11 kill the loss test and
all 11 kill the sibling too, including two shaped to exploit the only axes
where they differ (a batch-of-one gradient kill and a train-mode zeroing,
both caught by the projected-slot tests). Three mutations kill the sibling
while the loss test passes, so it was the weaker of the two.

Also drop scaffolding nothing exercises: _stub_model's attn_kernel keyword,
which no call site passes, and a cu_dtype override on a test that raises
before the cast it was meant to observe. Two docstrings pointed at code that
has since moved.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Comment thread tzrec/models/genrec_causal_lm_model.py Outdated
Comment thread tzrec/models/genrec_causal_lm_model.py
Comment thread tzrec/models/genrec_causal_lm_model.py Outdated
Comment thread tzrec/models/genrec_causal_lm_model_test.py Outdated
Comment thread tzrec/utils/test_util.py
@github-actions

Copy link
Copy Markdown
Contributor

Review summary — overall this is a well-built change. I verified the new label-window math against the old left-padded construction: the keep positions, the masked_fill condition, and the shift-based loss land on exactly the same (row, response) pairs in both layouts, and the short-row guard is the right call for a packed stream. The bit-exact rewrite-invariance check in _RowIsolationCase is a particularly strong guard, and _check_backbone_interfaces catching unbounded responses up front covers the logits_suffix_len cast.

Findings (details inline):

  1. Per-step host sync — the short-row guard's bool(...any()) blocks on GPU training every step; the collator could validate the same invariant host-side for free.
  2. Cross-kernel equivalence is not pinned by a test — the PR's headline claim (packed FA2 ≡ padded SDPA) is only backed by one-off manual verification; each CI arm self-compares within a single kernel.
  3. Two docstrings contradict the code — _left_pad_packed_inputs is not decode-only, and _padded_logits is not servable by "every kernel" given the CUDNN_ATTENTION exclusion.
  4. Import order — the TYPE_CHECKING block placement in test_util.py likely trips ruff I001.

One behavior change worth spelling out for release notes: the short-row guard now applies to the SDPA path too. The old padded forward trained through rows shorter than logits_suffix_len (with partially ignored labels); now a single such row fails the whole batch with a ValueError. Loud and well-messaged, but existing SDPA users with empty-prompt rows will hit it on upgrade.

The wheel constraints (cp311-only, torch-ABI-locked, no AUTO/fallback) are already honestly documented in the PR body — nothing to add there.

Comment thread tzrec/models/genrec_causal_lm_model.py
WhiteSwan1 and others added 6 commits September 21, 2026 09:26
Five follow-ups from review.

_padded_logits sent no position_ids, so HF assigned arange(max_seqlen) over
the padding too and a length-5 row in a width-12 batch started at position 7.
Both sibling paths spell the positions out -- _packed_logits subtracts the
row starts, dynamic_beam uses the mask cumsum -- and training was the only
one that did not. RoPE is relative so the outputs were identical, but a
backbone with absolute positions would have trained shifted and decoded from
zero. Use the same mask cumsum the decode path uses.

_attn_kernel moves to __init__, beside the config read it comes from.
init_backbone is a documented override point whose purpose is letting a
subclass supply its own self.lm; such a subclass never reached the
assignment, and the forward branches on it.

Add an integration case that runs the pipeline on FLASH_ATTENTION_2. The
kernel now defaults to SDPA, so train_eval -> eval -> export only ever
exercised the padded layout -- nothing drove the packed one through the
train pipeline, FX tracing or export. The config override threads through
_prepare_config, so the sdpa case keeps the CPU lane to itself.

Drop _varlen_batch's cu_dtype, which no caller passes, and hoist the SID
code fixtures into test_util beside the prompt helper that fixes their
offsets -- they were duplicated verbatim in two modules, each with its own
copy of the (4, 4, 4) codebook comment.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
No conflicts. The only file both sides touched is requirements/cu126|cu129|
cu130.txt, where master bumped dynamicemb to 20260920.9643985 and the
flash_attn line this branch adds sits beside it unchanged.

Master's proto changes (loss.proto, model.proto, multi_task_rank.proto) do
not reach GenRec, which scores through the backbone's own loss_function.
model.proto's Kernel enum -- the one AttnKernel mirrors -- is unchanged.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The 20260921 build covers all three supported Pythons on every CUDA minor,
so the pins no longer leave 3.10 and 3.12 without a flash path. Also moves
the variant into the version's local segment the way the sibling wheels
write it: build date, CUDA minor, torch minor, SM list, upstream commit.

cu129 and cu130 gain sm100/sm120; cu126 still builds sm80/sm90 only.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The first response token is predicted by the column before it, so a row whose
prompt assembles to nothing has no predictor for it: the packed arm took that
logit from the row before, the padded arm from a pad position, and neither
label was masked, because masking keys off response_length rather than row
length. The loss stayed finite and plausible, so nothing surfaced it.

The guard covering this compared row length against logits_suffix_len, which is
response_max + 1. That rejected any row below the response segment's maximum --
including a row with a full prompt and an empty answer, which is harmless -- and
it read a device tensor back to the host on every step. PromptAssembler now
zeroes such a row's response_length on the host, which masks its whole window
and leaves the packed stream aligned with the batch's features; a batch where
nothing survives still raises, because the loss returns NaN when it averages
over no tokens. _padded_logits keeps a narrower check of its own, on the padded
width, since logits_to_keep clamps silently and fails later inside the loss.

Renames _packed_logits to _varlen_logits, which only the flash kernel reaches,
and adds the cross-kernel case pinning that the two layouts agree.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
_left_pad_packed_inputs put its loudest noun on the input, so a call site
read `padded = ..._packed_inputs(...)`. The summary line now names the
rectangle it produces and why the padding is on the left, leaving the
packed input to the Args.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Comment thread tzrec/protos/models/genrec_model.proto Outdated
SDPA = 0;
FLASH_ATTENTION_2 = 1;
}
optional AttnKernel attn_kernel = 6 [default = SDPA];

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.

Could we use a string for attn_kernel instead of introducing a new protobuf enum here?

This value is essentially passed through to Hugging Face's attn_implementation ("sdpa", "flash_attention_2", etc.), so using the upstream string directly would avoid adding another enum to the proto namespace and also remove the extra _ATTN_KERNEL mapping.

For example:

optional string attn_kernel = 6 [default = "sdpa"];

The value is passed through to attn_implementation unchanged, so the enum only
bought a second spelling of names transformers already defines, plus a dict to
translate between them. A string keeps the proto namespace clear and lets any
implementation transformers accepts be configured without another proto change.

The cost is that a misspelling now parses and fails when the backbone is built,
rather than at config load; transformers names the accepted values there.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
@tiankongdeguiji
tiankongdeguiji merged commit 91f7f45 into alibaba:master Sep 22, 2026
7 checks 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.

2 participants