[feat] pack the GenRec training forward through FlashAttention-2 - #663
Conversation
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>
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>
|
Review summary — overall this is a well-built change. I verified the new label-window math against the old left-padded construction: the Findings (details inline):
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 The wheel constraints (cp311-only, torch-ABI-locked, no AUTO/fallback) are already honestly documented in the PR body — nothing to add there. |
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>
| SDPA = 0; | ||
| FLASH_ATTENTION_2 = 1; | ||
| } | ||
| optional AttnKernel attn_kernel = 6 [default = SDPA]; |
There was a problem hiding this comment.
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>
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_lensboundaries, per-rowposition_ids, and response-only logits gathered straight out of the packed layout.Measured
On 1× H20,
GenRecCausalLMModelover Qwen2.5-0.5B: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_kernelpicks the backend, carrying the transformers name straight through toattn_implementation:sdpais the default because it is the backend whose dependency always ships — building a GenRec model then needs neither a GPU nor theflash_attnwheel, which keeps the CPU lane and export working. Training runs opt intoflash_attention_2explicitly. There is noAUTOand 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_forwardbranches — packed forflash_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_attnis declared inrequirements/cu126|cu129|cu130.txtbeside the other self-hosted wheels — three lines each, one per supported Python (3.10/3.11/3.12), matching the coveragedynamicemb,faissandfbgemm_gpu_hstualready have.Test plan
_RowIsolationCaseasserts 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.5.96e-08with identical labels.Known gaps
c10::/at::symbols against unversioned sonames, so a torch bump without a rebuild fails at import, not at install.Requires-Dist: torchcarries no bound, so pip cannot catch the mismatch.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.