Bugfix/139 fa3 schedule full cudagraph prefix caching fa3 fixing by antonpibm · Pull Request #146 · generative-computing/granite-switch

antonpibm · 2026-10-14T19:24:34Z

FA3 FULL-cudagraph correctness (#139) + vLLM 0.26–0.30 support and test-harness fixes

This branch does two things that landed together: it fixes the issue-#139 FA3 correctness bug, and it widens the supported vLLM range to 0.26 through 0.30 in one tree, which in turn flushed out several real src/ bugs and a set of long-standing test-harness defects that only the newer vLLM lines exposed.

A four-leg GPU matrix (one vLLM line each) found unit / composer / hf green everywhere and the vLLM integration suite red on all four legs (7 failures on 0.27.1, 9 on 0.28.0, 16 on 0.29.0, 14 on 0.30.0). Some of those were genuine source bugs live on 0.27+ today; the rest was test scaffolding written against APIs that 0.29 deleted. All of it is addressed here. Where the lines differ, the code gates on the capability (hasattr/getattr) rather than on a version number, so it states the actual requirement and survives the next upstream rename.

Final verification: one vLLM line per pod, rc=0 on all 25 leg/suite combinations (each with a nonzero passed count — checked that way because pytest exits 0 on an all-skipped run).

1. Test fixes — vLLM equivalence timing & integration

Two categories: a real teardown/timing defect in the shared equivalence harness, and the integration-suite breakage from the version widening.

Engine teardown never waited (the "timing" problem)

The equivalence suites start two vLLM engines back to back. The second engine was measuring a card the first still owned: with 0.85 × 79.3 ≈ 67.4 GiB already taken, only ~10 GiB was free, so gpu_memory_utilization resolved to ~0.115 — below the engine's floor — and init failed.

The pre-existing del llm; gc.collect(); torch.cuda.empty_cache() cannot fix this:

The weights and KV cache live in the EngineCore child process, so empty_cache() drains the parent's allocator, which never held them.
Teardown is a weakref.finalize, so del only schedules it — it doesn't actually run.

Fix: a shared shutdown_llm() (now in tests/shared/vllm_gpu_mem.py) calls engine_core.shutdown(), which chains to CoreEngineProcManager.shutdown, terminates and joins the child processes, then polls mem_get_info until the driver has actually reclaimed the memory. The shutdown(timeout=None) signature is byte-identical on 0.26–0.30 and is reached defensively; on timeout it logs and returns rather than raising, leaving gpu_mem_util as the single place that can rule a card too full to use. This is why those ~10 equivalence failures were identical on all four legs — a pre-existing harness defect the newer lines merely exposed, not an API break.

run_vllm_logprobs is the single funnel for both engines in all five runners, so a setdefault there applies symmetrically:

enable_prefix_caching=False — upstream declares IsHybrid and we deliberately do not, so otherwise the two sides resolve different defaults (and it also dodges a 0.28-only "no mamba layers in the model" assertion in upstream's own model).
A free-memory-aware gpu_memory_utilization, measured against total memory the way vLLM does.

`enforce_eager` reverted in the shared runner

An earlier setdefault(enforce_eager=True) — added on the general principle that "an equivalence test must not measure the compiler" — turned TestZeroAdapterNoHiding::test_no_control_tokens red on every line 0.26–0.30. Pinned measurement showed the divergence is real and deterministic, not compiler noise: each engine is bit-exact against a second instance of itself, so there is no noise floor. The divergence appears only when the adapter tier is live (num_adapters=2), where the fused SWITCH kernel plus a zero-valued shrink/expand runs a different op sequence (different float reduction order) than the plain base linear at num_adapters=0. torch.compile normalizes that ordering — which is exactly what an equivalence test wants canonicalized, and is also how vLLM actually serves — so the eager pin was removed from the shared runner. _granite4_fullsize_tests.py keeps enforce_eager=True explicitly (at full size, num_adapters=0, eager is bit-exact and an unpinned autotuner reddened it ~1 run in 7). No tolerance was changed.

Integration-suite API breakage (0.29 deletions)

KV cache shape — AttentionBackend.get_kv_cache_shape is gone at 0.29. New tests/shared/vllm_kv_cache.py prefers compute_layer_kv_cache_shape_bytes (0.29+) and falls back to the old staticmethod, cross-checking the two where both exist (provably equal for dense bf16).
Audio processor test — _apply_uncached grows a second shape for the 0.29 text/MM processing split; the one test asserting on the deleted tri-tuple's flag is skipped there with a reason.
FlashInfer sampler — VLLM_USE_FLASHINFER_SAMPLER=0 is set in tests/conftest.py: on an image shipping CUDA 12.4 the 0.30 engine failed to start on nvcc fatal: Unknown option '--compress-mode=size' (flashinfer JIT-compiles its sampling kernels with an nvcc option added in CUDA 12.8). setdefault, so a CUDA ≥ 12.8 image can opt back in; the real fix is a newer toolkit in the runner image.

2. Audio changes

Audio (ASR cascade) was silently broken on vLLM 0.29+ — not merely untested; no transcript was ever produced — for two independent reasons:

ASR hook deleted at 0.29. 0.29 removed BaseMultiModalProcessor._call_hf_processor (and its only caller _apply_hf_processor_text_mm) and narrowed _apply_hf_processor_main to (mm_items, hf_processor_mm_kwargs) -> BatchFeature. Our _call_hf_processor override is where ASR runs, so on 0.29/0.30 it was simply never called. Fix: factor the ASR work into _audio_features() and add an _apply_hf_processor_main override that delegates to super() while the base class still has the old hook, and otherwise returns the audio tensors alone (on 0.29+ the prompt is vLLM's business — tokenized upstream, then _postprocess_prompt). **kwargs because the two vLLM call sites pass keywords only and no single signature matches both lines.

Prompt-update target type narrowed at 0.29. Audio prompt updates targeted the marker string, but vLLM narrowed UpdateTarget from PromptSeq | PromptIndex (str allowed, 0.26–0.28) to list[int] | PromptIndex (str dropped, 0.29+). A str target on 0.29+ doesn't fail where it's given — token matching finds nothing, drops vLLM into the 0.29-only _apply_prompt_updates_via_text fallback, which hands the raw string to tokenizer.decode(). The error then surfaces far from its cause as a misleading Rust-tokenizer crash (TypeError: Can't extract strtoVec``) that killed EngineCore during init. One bug, two symptoms:
5 `test_audio_processor` failures + all 13 audio integration errors. Fix: target the marker id (`_marker_id()` already existed) — a `list[int]` target is legal on all five lines, so no version branch is need
ed.

The audio runtime deps (av/soundfile/resampy) are also now pulled into the dev* dependency groups so the audio tests can actually run.

3. `_doubled_padding_mask` in the SR decoder

Where: src/granite_switch/vllm/decoder/shadow_residual/decoder.py.

The failure. vLLM ≥ 0.27 began passing an is_padding mask to the topk_softmax/topk_sigmoid MoE router kernels so they can skip padded rows. The router reads it off the forward context and the kernel then requires the lengths to agree. The runner sizes the mask to the padded token count m, but the Shadow-Residual layer routes its whole stacked [2M, H] base+adapter tensor in one expert call — so the kernel saw 2·m = 16384 gating rows against an m = 8192-entry mask and the engine died at init:

RuntimeError: dispatch_topk_launch, topk_softmax_kernels.cu:785,
              is_padding size mismatch, expected: 16384

The fix approach. Double the mask alongside the rows — this is the correct answer, not a workaround: adapter row i is the same token as base row i, so it is padding exactly when the base row is. _doubled_padding_mask(m) is a context manager that concatenates the mask with itself around the router call and restores it afterward.

Capability-gated, inert where there's nothing to match. It reads is_padding via getattr, not attribute access, and is a no-op when:

there is no forward context (unit tests);
the field doesn't exist — ForwardContext gained is_padding at 0.25.1; it's absent at 0.24.0, where attribute access previously raised AttributeError: 'ForwardContext' object has no attribute 'is_padding ' and took EngineCore down. (This was found by a 0.24 floor probe; getattr makes the helper catch up to its own documented "inert where there is nothing to match" contract. Not a regression on any supported l
ine — 0.25.1–0.30 all have the field.)
there's no mask (VLLM_MOE_SKIP_PADDING=0, or 0.25/0.26 where the field exists but no kernel consumes it);
the mask length isn't m — in which case the doubling premise doesn't hold, and failing loudly in the kernel beats silently masking the wrong rows.

4. FA3 correctness on Granite Switch (#139)

Two distinct FA3 faults, one per cudagraph mode, fixed by two composing patches: the FULL-cudagraph divergence (§4, fa3_schedule.py) and the eager AOT-schedule crash (§4a, fa3_aot_schedule.py).

Symptom. On vLLM 0.26 under FlashAttention 3 with cudagraph_mode=FULL, serving a Granite Switch checkpoint produces attention that diverges grossly from the very first prefill position — teacher-forced top-k JSD ≈ 0.78–0.87, tokens outside the adapter's allowed set — versus the eager / FULL_AND_PIECEWISE reference. Turning off FULL (eager), or dropping attention out of the graph (FULL_AND_PIECEWISE / PIECEWISE), matches the reference exactly.

Why FA3 + FULL specifically, and why it was never seen before. FlashAttentionMetadataBuilder.get_cudagraph_support reports AttentionCGSupport.ALWAYS when FA3 is active and UNIFORM_BATCH otherwise. ALWAYS is what lets cudagraph_mode=FULL capture attention for mixed prefill-decode batches. FA3 requires a Hopper GPU (compute 9.x); on any older GPU vLLM uses FA2, which reports UNIFORM_BATCH, so FULL silently downgrades to FULL_AND_PIECEWISE and the bug can't appear. It only surfaced on an FA3-capable (H100) run.

Root cause (and a corrected first guess). An earlier revision of the fix tried to resize FA3's ahead-of-time (AOT) schedule — overwriting num_heads_q for the SR decoder's doubled query heads — on the theory that a mis-sized schedule for SR's doubled-Q geometry was the cause. A GPU reproduction on H100 disproved it: the AOT head counts were already correct in both the broken and the correct paths, and disabling AOT entirely didn't help (attention stayed captured). The schedule shape was a red herring. The actual fault is upstream: _is_uniform_decode mislabels a padded prefill (num_reqs == num_actual_tokens, inconsistent query_start_loc) as a uniform-decode batch, which FULL then captures and FA3 computes wrong (upstream vllm #53051 / #53059, #28207). It is not Granite-Switch-specific in mechanism — a plain Granite base through the same build diverges identically — but Granite Switch is where it bites in production, since FULL was chosen for the latency of tiny prefix-cached decode steps.

The fix. Wrap get_cudagraph_support so that, for a granite_switch model only, the FlashAttention builder reports UNIFORM_BATCH instead of ALWAYS. vLLM then downgrades cudagraph_mode=FULL → FULL_AND_PIECEWISE on its own: the compute-heavy MLP/projection regions stay graph-captured, attention runs eagerly, and output matches the eager reference token-for-token. Every other model keeps FA3's native ALWAYS support untouched. The patch (src/granite_switch/vllm/fa3_schedule.py, wired into register()) targets vLLM 0.26–0.30 internals and is inert on any build without an FA3 ALWAYS promise (e.g. FA2). It fails loudly if the builder symbol or expected attributes have moved, rather than silently installing a no-op (a silent miss would reintroduce wrong tokens).

Cost. The FULL-vs-FULL_AND_PIECEWISE latency delta on attention launches — the price of correctness until the upstream FA3 full-cudagraph path is fixed on this line.

Tests. A subprocess-based harness (tests/vllm/test_fa3_schedule.py + worker, composing a small SR MultiSwitch checkpoint via GraniteSwitchComposer) asserts FULL + prefix caching matches the (eager, prefix-off) reference, and sweeps {FULL, FULL_AND_PIECEWISE, eager} × {prefix on/off} against the same reference. Assertions are on generated tokens (reference-derived allowed-set) and teacher-forced top-k distributions (JSD + Jaccard) over a shared continuation — never on the eager-only debug attributes, since FULL runs real CUDA graphs. A module-scoped require_fa3 fixture probes FA_VERSION in a subprocess and skips both tests unless FA3 is actually active, so the suite can't report a meaningless green on an FA2 box. Verified: with the enforcement reverted the target test fails exactly as #139 reports; with it in place the test passes under FA_VERSION=3.

4a. FA3 AOT-schedule sizing for the switch's own attention heads (#139, eager path)

The downgrade above acts only under FULL cudagraphs. A second, independent FA3 fault lives on the enforce_eager path, where the downgrade is a no-op. It is a distinct bug with a distinct fix; the two patches touch different seams and compose.

Symptom. Under enforce_eager=True on an FA3 (H100) box, the TP-integration arms TestTPRealAdapters::test_tp_logprobs_agree and TestTPMultiSwitch::test_tp_logprobs_agree_multi crashed at the generate TP=1 step with:

RuntimeError: scheduler_metadata must have shape (metadata_size)

inside the MultiSwitch counting head (vllm/switch/multi.py → FA3 flash_attn_varlen_func). Only these arms hit it: they feed a real granite-4.0-micro checkpoint through the chat/intrinsic path (a single small prefill, exactly when the switch runs its own attention). The granitemoe arms use batched synthetic token-prompts and dodged it.

Root cause. vLLM builds one FlashAttentionMetadataBuilder per KV-cache group and sizes FA3's ahead-of-time (AOT) scheduler metadata from the model config (num_attention_heads, num_kv_heads, head_size). But MultiSwitch adds two attention layers whose shapes differ from the backbone — a counting head (1, 1, 32) and a memory head (1, 1, 64) versus the config's (40, 8, 64) on granite-4.0-micro. FA3 built the schedule for the backbone shape; at forward time its CUDA kernel recomputes metadata_size from the switch head's real q/k/v and rejects the mismatched tensor. (This is why the earlier AOT-resize theory, disproven as the cause of the FULL-cudagraph bug in §4, is nonetheless the correct fix here: the resize fixes the eager schedule-shape crash, not the FULL-capture divergence.)

The fix. src/granite_switch/vllm/fa3_aot_schedule.py (patch_flash_attn_aot_schedule, wired into register() alongside the §4 downgrade) wraps FlashAttentionMetadataBuilder.__init__. For a granite_switch model it inspects the group's actual Attention layers:

Single real shape in the group (the normal case — the counting, memory and backbone heads each land in their own KV-cache group): overwrite the builder's num_heads_q/num_heads_kv/headdim with that true shape. AOT stays on (the FULL buffer is sized from batch, not shape), keeping the latency win.
Genuinely mixed shapes in one group (rare): disable AOT consistently (aot_schedule=False, tear down the FULL commitment) so no wrongly-sized schedule is ever built.

Other models are untouched, it fails loudly if the builder symbol/attributes move, and it is inert where FA3 isn't used (non-Hopper → FA2).

Verified on H100 (2026-10-15). An instrumented run confirmed the three groups resize to (1,1,32) / (1,1,64) / (40,8,64). Both previously-crashing arms now pass (2 passed in 8m33s): the TP=1 step returns cleanly and the TP=1-vs-TP=2 logprob distribution-agreement assertion holds — a real correctness result, not just no-crash. (The earlier failing test-logs/after-fix/test-tp.log predates this patch being wired into register(); it was a stale log, not a regression.)

Signed-off-by: itzikvaisman <yizhak.vaisman@ibm.com>

A four-leg GPU matrix (one vLLM line each) found unit/composer/hf green everywhere and the vLLM integration suite red on all four: 7 failures on 0.27.1, 9 on 0.28.0, 16 on 0.29.0, 14 on 0.30.0. Two of those were real bugs in src/, live today on 0.27+ regardless of this effort; the rest was test scaffolding written against APIs that 0.29 deleted. Where the lines differ, gate on the capability (hasattr) rather than on a version number: that states the actual requirement and survives the next rename. src/ fix 1 -- SR MoE router padding mask vLLM 0.27 started passing an is_padding mask to the topk kernels, read off the forward context and sized to the padded token count. The SR layer routes its [2M, H] base+adapter stack in one expert call, so the kernel saw 2M=16384 gating rows against an 8192-entry mask and the engine died at init with "is_padding size mismatch, expected: 16384". Double the mask alongside the rows: adapter row i is the same token as base row i, so it is padding exactly when the base row is. Inert with no forward context, no mask (VLLM_MOE_SKIP_PADDING=0, or 0.26 where no kernel consumes one), or a mask whose length is not m. CI missed this because test_sr_tp_equivalence.py is env-gated on SR_COMPOSED_DIR and never runs -- a follow-up. src/ fix 2 -- audio ASR never ran on 0.29+ 0.29 deleted BaseMultiModalProcessor._call_hf_processor along with _apply_hf_processor_text_mm, its only caller, and narrowed _apply_hf_processor_main to (mm_items, hf_processor_mm_kwargs) -> BatchFeature. Our _call_hf_processor override is where ASR happens, so on 0.29/0.30 it was simply never called and no transcript was ever produced -- silently broken, not merely untested. Factor the ASR work into _audio_features() and add an _apply_hf_processor_main override that delegates to super() while the base class still has the old hook, and otherwise returns the audio tensors alone (on 0.29+ the prompt is vLLM's business: tokenized upstream, then _postprocess_prompt). **kwargs because both vLLM call sites pass keywords only and no single signature matches both lines. _hf_processor_applies_updates stays as-is -- required on 0.26-0.28, inert on 0.29+ where _apply_prompt_updates always runs, which is the False behaviour it was asking for. tests -- KV cache shape AttentionBackend.get_kv_cache_shape is gone at 0.29 (29 defs at 0.28, 0 at 0.29). The three call sites build a real KV cache fed to a real FlashAttention kernel, so the shape has to be genuinely right. New tests/shared/vllm_kv_cache.py prefers compute_layer_kv_cache_shape_bytes (0.29+) and falls back to the old staticmethod, cross-checking the two where both exist. They are provably equal for dense bf16: num_heads == num_kv_heads, one kernel state per token, and content bytes == 2 * head_size * itemsize. tests -- audio processor _apply_uncached grows a second shape for 0.29+, which split text processing away from MM processing; everything after the first call (_get_mm_fields_config -> _get_mm_prompt_updates -> _maybe_apply_prompt_updates) is unchanged on every line. The one test that asserted on the deleted tri-tuple's flag is skipped there with a reason, and the class docstring records that the hook test is vacuous at 0.29+. tests -- equivalence harness run_vllm_logprobs is the single funnel for both engines in all five runners, so setdefault there applies symmetrically: enforce_eager (an equivalence test must not measure the compiler), enable_prefix_caching=False (upstream declares IsHybrid and we deliberately do not, so the two sides otherwise resolve different defaults -- and it dodges a 0.28-only "no mamba layers in the model" assertion in upstream's own model), and a free-memory-aware gpu_memory_utilization. The last reuses the pattern already solved in _multi_adapter_equivalence_worker.py, lifted to tests/shared/vllm_gpu_mem.py: vLLM measures the fraction against TOTAL memory, these suites start engines back to back, and the guard is byte-identical back to 0.26 -- a pre-existing harness defect the newer legs merely exposed. _granite4_mini_tests.py's per-config eager opt-in is now redundant and deleted rather than left to contradict the default. pyproject -- the [vllm] extra claimed 0.30 only while [tool.uv] default-groups said vllm26, which made uv sync unsatisfiable. Widen it to >=0.26,<0.31, prune the duplicate [vllm27] extra (the dependency group covers pinning), and drop the group-vs-extra and extra-vs-extra conflict entries that only existed to describe the contradiction. uv.lock relocked accordingly. Signed-off-by: itzikvaisman <yizhak.vaisman@ibm.com>

Two defects that kept the 0.29/0.30 legs of the support matrix red. Both are latent on the lines already supported, so neither is a cost of widening the range. Audio prompt updates targeted the marker STRING. vLLM narrowed the accepted target type at 0.29: 0.26 - 0.28 UpdateTarget = PromptSeq | PromptIndex # str allowed 0.29 - 0.30 UpdateTarget = list[int] | PromptIndex # str dropped Passing a str on 0.29+ does not fail where it is given. Token matching simply finds nothing, which drops vLLM into _apply_prompt_updates_via_text -- a fallback that only exists from 0.29 -- and that hands the raw string to tokenizer.decode(). The error then surfaces far from its cause and reads as a tokenizer bug: "TypeError: Can't extract `str` to `Vec`" out of the Rust tokenizer, killing EngineCore during init rather than rejecting one request. That accounts for 5 test_audio_processor failures and all 13 audio integration errors -- one bug, two symptoms. A list[int] target is legal on all five lines, so the fix needs no version branch, and _marker_id() already existed. Engine teardown never waited. The suites start two engines back to back and the second was measuring a card the first still owned: 0.85 x 79.3 = 67.4 GiB taken, 10.1 GiB left, gpu_memory_utilization=0.115, below the floor. The existing "del llm; gc.collect(); torch.cuda.empty_cache()" cannot fix that, because the weights and KV cache belong to the EngineCore CHILD process -- empty_cache() drains the parent's allocator, which never held them -- and teardown is a weakref.finalize, so del schedules it rather than performing it. shutdown_llm() calls engine_core.shutdown(), which chains to CoreEngineProcManager.shutdown and terminates AND joins the children, then polls mem_get_info until the driver has reclaimed. The shutdown(timeout=None) signature is byte-identical on 0.26 - 0.30 and is reached defensively. On timeout it prints and returns rather than raising, leaving gpu_mem_util as the single place that rules a card too full to use. This is why those 10 equivalence failures were identical on all four legs: a pre-existing harness defect the newer lines merely exposed, not an API break. Signed-off-by: itzikvaisman <yizhak.vaisman@ibm.com>

_doubled_padding_mask reached ForwardContext.is_padding by attribute access. That field was added at vLLM 0.25.1 and does not exist at 0.24.0, so every SR forward raised AttributeError: 'ForwardContext' object has no attribute 'is_padding' and took EngineCore down during init: RuntimeError: Engine core initialization failed. Measured on itzikv-eval-v24-g5-sr (vllm 0.24.0, torch 2.11.0, transformers 5.17.0 -- the install itself resolved cleanly), 2026-10-04. The docstring already claimed inertness "where there is nothing to match", and a missing field is the clearest such case, so this is the code catching up to its own contract rather than a new behaviour. getattr, not a version compare: the requirement is "this build has the field", which is what the capability check states. Not a regression on any supported line -- 0.25.1 through 0.30 all have the field, and SR scored 0.8960 against its 0.8948 bar on 0.30 with the attribute access in place. It only blocks going BELOW the current floor, which is what the 0.24 probe was submitted to find out. Signed-off-by: itzikvaisman <yizhak.vaisman@ibm.com>

7accc4c added three setdefaults to run_vllm_logprobs. Two fixed observed failures. The third, enforce_eager=True, was added on the general principle that "an equivalence test must not measure the compiler", and it turned TestZeroAdapterNoHiding::test_no_control_tokens red on every vLLM line from 0.26 to 0.30 -- the last failure in the four-leg matrix. Pinned in one pod (vLLM 0.26.0, sha 5df1b8f, same venv, that setdefault the only variable): enforce_eager=True rc=1 3 failed, 12 deselected 4:26 enforce_eager=False rc=0 3 passed, 12 deselected 6:35 and the test was green at 5c0c72d, where the mini suite ran compiled (_UPSTREAM_EAGER_CONFIGS keyed on "4.0-h-350m", which has never been a key of GRANITE4_MINI, so _eager_kwargs_if_needed always returned {} -- it was dead from the initial commit and is not restored here). granite4_equivalence.py is byte-identical over 5c0c72d..5df1b8f, so the tolerance never moved. Measured the floor rather than assuming it, because "the gate is below the eager noise floor" and "there is a real divergence the compiler was hiding" both predict the above and want opposite patches. Eager, on the zero-LoRA pair the test builds: upstream vs upstream switch vs switch switch vs upstream 4.0-1b 0.0000e+00 0.0000e+00 9.9182e-04 4.0-350m 0.0000e+00 0.0000e+00 1.9522e-03 4.0-micro 0.0000e+00 0.0000e+00 5.0116e-04 Each engine is bit-exact against a second instance of itself over all 3840 elements, so there is no noise floor: the divergence is real and deterministic, and the switch/upstream column reproduces the three failures' worst diffs to five digits. It appears only when the adapter tier is live. save_upstream_model builds num_adapters=0, where the projection is the plain base linear and eager is bit-exact; save_zero_adapter_model builds num_adapters=2, where it runs the fused SWITCH kernel plus a zero-valued shrink/expand -- a different op sequence with a different reduction order. get_tolerances' 4.77e-7 basis for the 1e-5 inert gate was measured at num_adapters=0 and compiled, so it never characterized this configuration. torch.compile lowers both sequences to equivalent kernels and normalizes the ordering, which is exactly what an equivalence test wants canonicalized; it preserves semantics, so a genuine kernel bug still trips 1e-5. Compiled is also how vLLM serves. So the choice belongs to each test. _granite4_fullsize_tests.py keeps enforce_eager=True explicitly -- at full size, with num_adapters=0, eager is bit-exact and an unpinned autotuner reddened it ~1 run in 7 -- and setdefault was never what gave it that. enable_prefix_caching=False and gpu_memory_utilization stay: they fix the vLLM 0.28 mamba assertion and the in-pod OOM respectively. No tolerance was changed. The measurements are recorded at all three sites so the next reader does not re-derive them, or repeat the mistake. Signed-off-by: itzikvaisman <yizhak.vaisman@ibm.com>

The two GPU legs now install the two ENDS of the supported range, 0.26 and 0.30, instead of 0.26 and 0.27. WHY THE ENDS. pyproject declares vllm>=0.26,<0.31, and every break this range actually contains sits at a boundary. APIs deleted at 0.29 only fail at the ceiling: AttentionBackend.get_kv_cache_shape (29 defs at 0.26-0.28, zero at 0.29) and BaseMultiModalProcessor._call_hf_processor (the method that ran ASR, so audio was silently producing no transcript). An API that does not exist yet only fails at the floor: ForwardContext.is_padding, added at 0.27, which the SR router slices. A 0.27 leg is bracketed by two tested lines and would have caught none of them. 0.27/0.28/0.29 were verified green out of band by a five-leg matrix across every suite, so this is a per-run cost decision, not a claim that the middle is untested. Measured at 8a4d69d, one vLLM line per pod, VLLM_USE_FLASHINFER_SAMPLER=0: leg vLLM unit composer hf vllm integration v26 0.26.0 1739p 459p 449p 220p/4s 21p/9s v27 0.27.1 1739p 459p 449p 220p/4s 21p/9s v28 0.28.0 1739p 459p 449p 220p/4s 21p/9s v29 0.29.0 1739p 459p 449p 219p/5s 21p/9s v30 0.30.0 1739p 459p 449p 219p/5s 21p/9s rc=0 on all 25 leg/suite combinations, each with a nonzero passed count -- checked that way because pytest exits 0 on an all-skipped run. The 219-vs-220 gap is one deliberate skip: test_uncached_path_reports_updates_not_applied asserts a contract 0.29 removed. The other four skips are identical on every leg and env-gated. VLLM_USE_FLASHINFER_SAMPLER=0 is now set in tests/conftest.py, because raising the ceiling is what exposes the need for it. On an image shipping CUDA 12.4 the 0.30 engine failed to start 171 times on nvcc fatal : Unknown option '--compress-mode=size' flashinfer JIT-compiles its sampling kernels with --compress-mode, an nvcc option added in CUDA 12.8. It is version-dependent: 0.26 is green on that image without the flag and 0.30 is not, so a 0.26-only CI never needed it. The pod template lives in the runner image rather than this repo, so conftest is the only repo-side lever. setdefault, so an image with CUDA >= 12.8 can export VLLM_USE_FLASHINFER_SAMPLER=1 and exercise that path for real; a newer toolkit in the image is the actual fix. Scoped to the sampler and inert for what these suites assert -- the attention backend resolves to FlashAttention independently, and the equivalence tests compare logprobs, produced before any sampling op runs. uv.lock is regenerated: vllm 0.27.1 -> 0.30.0. It resolves transformers 5.17.0 and torch 2.11.0 for the 0.26 leg / 2.13.0 for the 0.30 leg, matching what the matrix pods resolved through pip. Signed-off-by: itzikvaisman <yizhak.vaisman@ibm.com>

The vLLM 0.26-0.30 widening re-locked uv.lock, and the re-lock was done with the local uv (0.7.19, Homebrew). That binary writes `revision = 2` in the lockfile header, so the commit silently DOWNGRADED the field from the `revision = 3` already on the branch. .pre-commit-config.yaml pins uv-pre-commit at rev 0.8.4, whose uv writes `revision = 3`. So CI's `uv-lock` hook re-upgraded the field, reported "pre-commit hook(s) made changes", and exited 1 on a one-line diff: -revision = 2 +revision = 3 Re-locked with `uvx uv@0.8.4 lock` so the committed file matches what the pinned hook produces. The hook is now a no-op. RESOLUTION IS UNCHANGED, which is why this is header-only: the diff against the previous commit is exactly one line. 0.8.4 resolves the same 287 packages (the count CI printed), and the pins this effort cares about all hold -- vllm 0.26.0 for the dev/test/vllm26 groups and 0.30.0 for dev-vllm30/vllm30, transformers 5.17.0, torch 2.11.0 on the 0.26 leg and 2.13.0 on the 0.30 leg. Those are the versions the five-leg GPU matrix actually installed, so the lockfile still describes the configuration that was measured green. `revision` is a lockfile FORMAT marker, not a dependency, so bumping it cannot move a version. The rule it encodes is that the file must be written by a uv at least as new as the hook's pin; writing it with an older one is what regressed it. Verified the rest of the hook list separately, since fail_fast: true means the CI log only proves uv-lock failed: uv-lock is hook 15 of 15, so CI reaching it already implies 1-14 passed. Confirmed directly anyway -- ruff 0.9.0 format (223 files) and check clean, check_headers clean, validate_links clean. Signed-off-by: itzikvaisman <yizhak.vaisman@ibm.com>

Two CI reds on this PR, NEITHER of which was a test failure. CPU TESTS: the suite passed. Step-level conclusions on both 3.11 and 3.12 were `success` for "Run CPU tests" and `failure` for the next step, "Upload coverage". codecov-action@v4 exits nonzero on a token/rate-limit problem of its own BEFORE it consults `fail_ci_if_error: false`, so the flag that already declared "upload problems must not gate CI" could not take effect. Moved to v5 (v4 is deprecated) and added `continue-on-error: true`, which enforces that intent at the workflow level where the action cannot override it. This is not masking a failure: `tests/unit/` is untouched by this branch, and the install step resolved exactly as intended (vllm 0.26.0 / torch 2.11.0 / transformers 5.17.0). The same ci.yaml passed on main on 2026-09-30 and fails now, so the trigger is external drift. GPU TESTS: a chicken-and-egg in the dispatcher. `workflow_dispatch` always executes the workflow file from the DEFAULT BRANCH, so `/gpu-test-multi` ran main's matrix -- legs `vllm26 -> dev` and `vllm27 -> dev-vllm27` -- while checking out this PR's commit (the uploaded artifacts are named for sha 8276623). This branch had deleted `dev-vllm27`, so that leg's `uv sync --group dev-vllm27` had no such group and died before pytest started; both logs are under 1.1 KB, the "suite never started" signature. Restored `vllm27` / `dev-vllm27` so main's dispatched matrix stays resolvable against this tree until the 0.26/0.30 rename lands on main. The group is not a fiction -- the `vllm` extra claims 0.26-0.30, so 0.27 is a supported line that simply is not worth its own per-run leg -- and it is removable once main carries the rename. Conflicts now enumerate every cross-line pair: three lines with a bare and a dev- flavour each give 3 x (2 x 2) = 12. `test` stays absent because it reaches vllm26 through `include-group = "dev"` and uv derives that transitively, which is also why main's two-line version never listed it. DELIBERATELY NOT CHANGING `default-groups = ["vllm26"]`, though `uv sync --group dev-vllm30` reports it as conflicting. That shape is not new: at 6013c7f `default-groups` was `["vllm19"]` against a declared `dev-vllm20`/`vllm19` conflict, and that run's vllm20 leg PASSED -- so the cluster harness already excludes default groups. Changing it here would be a guess against evidence. Verified: `uv lock` is a no-op under the hook's pinned uv 0.8.4 (revision 3); each leg resolves to its own line (dev 0.26.0, dev-vllm27 0.27.1, dev-vllm30 0.30.0); ruff 0.9.0 format and check clean; pyproject and both workflows parse. Signed-off-by: itzikvaisman <yizhak.vaisman@ibm.com>

…hing (#139) On vLLM 0.26, a Shadow-Residual MultiSwitch checkpoint served with cudagraph_mode=FULL and prefix caching ON emits tokens outside the adapter's allowed set and distributions that diverge from the reference, but only on small prefix-cached decode steps. FA3's ahead-of-time schedule is sized from the model config's attention shape, which does not match SR's doubled query heads, so the frozen FULL-cudagraph schedule is wrong for the switch group. Adds a subprocess-based test harness (test + worker, mirroring the generation equivalence convention) that composes a small SR MultiSwitch checkpoint via GraniteSwitchComposer and: - test_fa3_full_cudagraph_prefix_caching_is_correct: asserts FULL + prefix caching matches the (eager, prefix-off) reference. Expected to FAIL until the fa3_schedule fix lands; it is the bug-identifying test. - test_fa3_config_matrix_matches_reference: compares {FULL, FULL_AND_PIECEWISE, eager} x {prefix on/off} against the same reference so a schedule/numerics regression in any setting is caught. Assertions are on generated tokens (allowed-set) and teacher-forced top-k distributions (JSD + Jaccard), never on the eager-only MultiSwitch debug attributes, since FULL runs real CUDA graphs. Full writeup in docs/FA3_SCHEDULE_FULL_CUDAGRAPH_BUG.md. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: yairallouche <yair@il.ibm.com>

…eacher-forcing The first Vela run exposed two flaws that made every config cell fail identically (so the test could not distinguish the bug from noise): - The allowed-set was hardcoded (range(10,40)/range(40,70)); a random-weight base emits arbitrary ids, so every run looked "out of set". The allowed set is now the reference run's own greedy output per adapter — "in-set" means "matches the known-good reference". - Each cell teacher-forced over its OWN continuation, so reference and candidate were scored on different sequences and JSD hit 1.0 even for identical behavior. The reference now generates the continuation once and writes it; every candidate scores distributions over the same [prefix + ctrl + reference-continuation]. Also add an FA_VERSION probe: the bug only exists under FlashAttention 3 + FULL cudagraph, and FA3 requires an SM90 GPU plus a vllm_flash_attn build whose FA3 kernel loads for the installed torch/CUDA. When FA3 is absent vLLM falls back to FA2 and silently downgrades cudagraph_mode=FULL to FULL_AND_PIECEWISE, making the run moot. Printing get_flash_attn_version() makes that visible instead of a misleading pass/fail. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: yairallouche <yair@il.ibm.com>

…heads (#139) vLLM sizes FlashAttention 3's ahead-of-time schedule from the model config's attention shape. Granite Switch's layers differ from it — the Shadow-Residual decoder doubles its query heads (Attention(2*num_heads, ...)) while the config reports the un-doubled count, and the switch's counting/memory heads carry their own head_size. For such a group FA3's schedule is computed for the wrong shape; under cudagraph_mode=FULL the mis-sized schedule is frozen into the captured graph and a small prefix-cached decode produces tokens outside the adapter's allowed set (the token-0 KeyError) and diverged distributions. The patch wraps FlashAttentionMetadataBuilder.__init__ and, for a granite_switch model, resizes the schedule from the group's own layers: - single-shape group: overwrite num_heads_q/num_heads_kv/headdim with the real shape (corrects the SR decoder group; keeps AOT on, so FULL stays fast); - genuinely mixed-shape group: disable AOT consistently (aot_schedule=False, scheduler_metadata=None, max_num_splits=0) so the builder state is self-consistent whatever path build() takes. Version safety: the patch soft-returns only when vLLM is genuinely absent; if vLLM is present but the builder symbol or expected attributes have moved, it logs loudly / raises instead of silently installing a no-op (a silent miss would reintroduce wrong tokens). Wired into register() as its first action. The bug only manifests under FA3, which needs a Hopper GPU (compute 9.x); on older GPUs vLLM uses FA2 and FULL downgrades to FULL_AND_PIECEWISE, so the patch is inert there but still correct to install. See docs/FA3_SCHEDULE_FULL_CUDAGRAPH_BUG.md. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: yairallouche <yair@il.ibm.com>

…notes (#139) Reflects what the Vela diagnostic probe established: - A clear status table and plain-reader summary up top. - A "Reproducing this requires FlashAttention 3" section: FA3 (hence FULL cudagraph) only runs on a Hopper GPU (compute 9.x); on older GPUs vLLM uses FA2 and FULL silently downgrades to FULL_AND_PIECEWISE, so the bug cannot appear. How to tell in one line (FA_VERSION), and why the tests skip when FA3 is absent. - Corrects the earlier torch-version-mismatch guess: vLLM 0.26 pins torch 2.11 (which the pod had); the blocker is the GPU generation. - Verification + caveat updated for the per-vLLM-line meaning of a pass. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: yairallouche <yair@il.ibm.com>

Issue #139 only exists under FlashAttention 3 + cudagraph_mode=FULL. FA3 needs a Hopper GPU (compute capability 9.x); on any other GPU vLLM uses FA2, whose cudagraph support silently downgrades FULL to FULL_AND_PIECEWISE, the setting that is already correct. Every cell then matches the reference and the tests passed without exercising the bug: a green that meant nothing. - Worker: add a `probe` mode that prints FA_VERSION and, when FA3 is absent, vLLM's own fa_version_unsupported_reason(3). It always exits 0, because a missing FA3 is a reason to skip, not a worker failure. - Tests: a module-scoped `require_fa3` fixture runs the probe once in a subprocess and skips both tests unless FA_VERSION == 3, quoting the reason. On an A100 the skip reads "FA3 is only supported on devices with compute capability 9.x". Also move the worker's torch/safetensors imports into the functions that use them. At module level they made every subprocess, including the CPU-only `compare` and `probe` modes, die on a machine without torch. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: yairallouche <yair@il.ibm.com>

Replace the disproven AOT-schedule-resize approach with a get_cudagraph_support downgrade: for a granite_switch model, report AttentionCGSupport.UNIFORM_BATCH (FA2's level) so vLLM downgrades cudagraph_mode=FULL -> FULL_AND_PIECEWISE. Attention then runs eagerly (correct) while the compute-heavy MLP/projection regions stay graph-captured. Why: under FlashAttention 3 + cudagraph_mode=FULL on vLLM 0.26, a granite_switch checkpoint's attention diverges grossly from the eager reference (teacher-forced JSD ~0.78-0.87, out-of-set tokens) from the first prefill position. Root-cause tracing on an H100 this session (see docs/FA3_TESTS.md §4a) showed the fault is NOT a mis-sized AOT schedule for SR's doubled query heads (head counts are identical and correct in the broken and correct paths) but an upstream vLLM batch-classification bug: _is_uniform_decode mislabels a padded prefill (num_reqs == num_actual_tokens, inconsistent query_start_loc) as a uniform-decode batch, which FULL then captures and FA3 computes wrong. The UNIFORM_BATCH downgrade keeps that mislabeled batch out of the FULL graph. Verified with the enforcement reverted the target test fails exactly as #139 reported; with it in place the test passes under FA_VERSION=3. Docs: FA3_TESTS.md rewritten with the confirmed mechanism, the live batch-descriptor dump, and the upstream linkage (vllm #53051/#53059, #28207); FA3_SCHEDULE_FULL_CUDAGRAPH_BUG.md updated. Also includes unrelated in-progress audio/integration test work carried in the working tree. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Anton Puzanov <antonp@il.ibm.com>

Signed-off-by: Anton Puzanov <antonp@il.ibm.com>

Post-rebase cleanup of the #139 branch against the reworked main: - Remove four design/debug docs from version control (FA3_SCHEDULE_FULL_CUDAGRAPH_BUG, FA3_TESTS, FIXING_VLLM, INTEGRATION_FIXES) and the local run_vllm_tests.sh helper. They stay on disk but are no longer tracked. - Scrub the now-dangling doc-path references out of the vLLM backend and the shared/vLLM test helpers, keeping the inline explanations and the issue #139 pointer. - Restore the dev-latest dependency group, dropped during the rebase conflict resolution. It aliases the newest supported line (dev-vllm30) so the ceiling is repointed in one place rather than hardcoded; the GPU workflow's derive_legs.py and CLAUDE.md both ask for it. Add its conflict pairs so uv will not co-resolve two vLLM lines. - Regenerate uv.lock for the resolved pyproject. Signed-off-by: Anton Puzanov <antonp@il.ibm.com>

Signed-off-by: Anton Puzanov <antonp@il.ibm.com>

antonpibm · 2026-10-14T20:34:31Z

/gpu-test-multi

codecov-commenter · 2026-10-14T20:36:43Z

Codecov Report

✅ All modified and coverable lines are covered by tests.

📢 Thoughts on this report? Let us know!

github-actions · 2026-10-15T00:25:02Z

✅ GPU tests passed — `vllm26-multi`

2862 passed, 68 skipped, 20 warnings in 13582.99s (3:46:22)

Commit: bdd9b9411a8a5aff90b74cc73d124b16002b7b77
Full run & artifact log

Last 40 log lines

tests/integration/test_multi_switch_vllm_generate.py::test_different_adapters_produce_different_text SKIPPED

=============================== warnings summary ===============================
<frozen importlib._bootstrap>:488
  <frozen importlib._bootstrap>:488: DeprecationWarning: builtin type SwigPyPacked has no __module__ attribute

<frozen importlib._bootstrap>:488
  <frozen importlib._bootstrap>:488: DeprecationWarning: builtin type SwigPyObject has no __module__ attribute

.venv/lib/python3.12/site-packages/torch/jit/_script.py:365: 14 warnings
  /tmp/granite-switch/.venv/lib/python3.12/site-packages/torch/jit/_script.py:365: DeprecationWarning: `torch.jit.script_method` is deprecated. Please switch to `torch.compile` or `torch.export`.
    warnings.warn(

tests/composer/test_compose_e2e.py:83
  /tmp/granite-switch/tests/composer/test_compose_e2e.py:83: PytestUnknownMarkWarning: Unknown pytest.mark.xdist_group - is this a typo?  You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html
    @pytest.mark.xdist_group("compose_e2e")

tests/composer/test_multi_audio_compose_e2e.py:124
  /tmp/granite-switch/tests/composer/test_multi_audio_compose_e2e.py:124: PytestUnknownMarkWarning: Unknown pytest.mark.xdist_group - is this a typo?  You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html
    @pytest.mark.xdist_group("multi_audio_compose_e2e")

tests/composer/test_multi_audio_compose_e2e.py:130
  /tmp/granite-switch/tests/composer/test_multi_audio_compose_e2e.py:130: PytestUnknownMarkWarning: Unknown pytest.mark.xdist_group - is this a typo?  You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html
    @pytest.mark.xdist_group("multi_audio_compose_e2e")

tests/composer/test_upstream_files.py:218
  /tmp/granite-switch/tests/composer/test_upstream_files.py:218: PytestUnknownMarkWarning: Unknown pytest.mark.xdist_group - is this a typo?  You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html
    @pytest.mark.xdist_group("upstream_build_e2e")

-- Docs: https://docs.pytest.org/en/stable/how-to/capture-warnings.html
========= 2862 passed, 68 skipped, 20 warnings in 13582.99s (3:46:22) ==========
sys:1: DeprecationWarning: builtin type swigvarlink has no __module__ attribute
[rank0]:[W1015 15:07:16.773137390 ProcessGroupNCCL.cpp:1575] Warning: WARNING: destroy_process_group() was not called before program exit, which can leak resources. For more info, please see https://pytorch.org/docs/stable/distributed.html#shutdown (function operator())
===== ALL GPU TESTS PASSED =====
[15:07:28] <job> Succeeded
[15:07:31] verified: found success sentinel in pod log
[15:07:31] cleanup: deleting <job> (exit=0)
<job> "<job>" deleted

Verdict: PASSED

github-actions · 2026-10-15T00:39:30Z

✅ GPU tests passed — `vllm30-multi`

2861 passed, 69 skipped, 23 warnings in 14681.56s (4:04:41)

Commit: bdd9b9411a8a5aff90b74cc73d124b16002b7b77
Full run & artifact log

Last 40 log lines

  /tmp/granite-switch/.venv/lib/python3.12/site-packages/torch/jit/_script.py:365: DeprecationWarning: `torch.jit.script_method` is deprecated. Please switch to `torch.compile` or `torch.export`.
    warnings.warn(

tests/composer/test_compose_e2e.py:83
  /tmp/granite-switch/tests/composer/test_compose_e2e.py:83: PytestUnknownMarkWarning: Unknown pytest.mark.xdist_group - is this a typo?  You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html
    @pytest.mark.xdist_group("compose_e2e")

tests/composer/test_multi_audio_compose_e2e.py:124
  /tmp/granite-switch/tests/composer/test_multi_audio_compose_e2e.py:124: PytestUnknownMarkWarning: Unknown pytest.mark.xdist_group - is this a typo?  You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html
    @pytest.mark.xdist_group("multi_audio_compose_e2e")

tests/composer/test_multi_audio_compose_e2e.py:130
  /tmp/granite-switch/tests/composer/test_multi_audio_compose_e2e.py:130: PytestUnknownMarkWarning: Unknown pytest.mark.xdist_group - is this a typo?  You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html
    @pytest.mark.xdist_group("multi_audio_compose_e2e")

tests/composer/test_upstream_files.py:218
  /tmp/granite-switch/tests/composer/test_upstream_files.py:218: PytestUnknownMarkWarning: Unknown pytest.mark.xdist_group - is this a typo?  You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html
    @pytest.mark.xdist_group("upstream_build_e2e")

tests/unit/test_asr.py::TestMonoAndResample::test_resample_real_vllm_backend[8000-16000]
tests/unit/test_asr.py::TestMonoAndResample::test_resample_real_vllm_backend[44100-16000]
  /tmp/granite-switch/.venv/lib/python3.12/site-packages/torch/backends/mkldnn/__init__.py:96: UserWarning: TF32 acceleration on top of oneDNN is available for Intel GPUs. The current Torch version does not have Intel GPU Support. (Triggered internally at /__w/pytorch/pytorch/aten/src/ATen/Context.cpp:265.)
    torch._C._set_onednn_allow_tf32(_allow_tf32)

tests/integration/test_hf_to_vllm_weights.py::TestMultiSwitchForwardEquivalence::test_forward_logit_equivalence
  /tmp/granite-switch/.venv/lib/python3.12/site-packages/flashinfer/cute_dsl/availability.py:135: DeprecationWarning: Arch from arch.py is deprecated. Use Arch from enums.py instead.
    from cutlass.base_dsl.arch import Arch

-- Docs: https://docs.pytest.org/en/stable/how-to/capture-warnings.html
========= 2861 passed, 69 skipped, 23 warnings in 14681.56s (4:04:41) ==========
sys:1: DeprecationWarning: builtin type swigvarlink has no __module__ attribute
[rank0]:[W1015 15:38:01.336309747 ProcessGroupNCCL.cpp:1624] Warning: WARNING: destroy_process_group() was not called before program exit, which can leak resources. For more info, please see https://pytorch.org/docs/stable/distributed.html#shutdown (function operator())
===== ALL GPU TESTS PASSED =====
===== ALL GPU TESTS PASSED =====
[15:38:28] <job> Succeeded
[15:38:30] verified: found success sentinel in pod log
[15:38:30] cleanup: deleting <job> (exit=0)
<job> "<job>" deleted

Verdict: PASSED

Signed-off-by: Anton Puzanov <antonp@il.ibm.com>

antonpibm · 2026-10-15T07:26:03Z

/gpu-test-multi

Signed-off-by: Anton Puzanov <antonp@il.ibm.com>

antonpibm · 2026-10-15T09:32:08Z

/gpu-test-multi --no-cache

Signed-off-by: Anton Puzanov <antonp@il.ibm.com>

antonpibm · 2026-10-15T10:01:28Z

/gpu-test-multi --no-cache

ItzikVa and others added 13 commits October 14, 2026 19:31

test vLLM 0.30 support

e69ea5b

Signed-off-by: itzikvaisman <yizhak.vaisman@ibm.com>

antonpibm force-pushed the bugfix/139-fa3-schedule-full-cudagraph-prefix-caching-fa3_fixing branch 2 times, most recently from 4d3f1f5 to ed8e3cb Compare October 14, 2026 20:04

antonpibm and others added 4 commits October 14, 2026 20:31

Fixing the timing problems

e947de2

Signed-off-by: Anton Puzanov <antonp@il.ibm.com>

fix tests formatting

457134f

Signed-off-by: Anton Puzanov <antonp@il.ibm.com>

antonpibm force-pushed the bugfix/139-fa3-schedule-full-cudagraph-prefix-caching-fa3_fixing branch from 0f2957e to 457134f Compare October 14, 2026 20:33

fix AOT size problem

38a16e0

Signed-off-by: Anton Puzanov <antonp@il.ibm.com>

antonpibm added 2 commits October 15, 2026 07:33

fix pre-commit

a576285

Signed-off-by: Anton Puzanov <antonp@il.ibm.com>

fix ASR exclusive process issue

a673047

Signed-off-by: Anton Puzanov <antonp@il.ibm.com>

antonpibm marked this pull request as ready for review October 15, 2026 08:53

antonpibm requested a review from aviv1ron1 as a code owner October 15, 2026 08:53

antonpibm requested review from freunda and yairallouche as code owners October 15, 2026 08:53

antonpibm added 2 commits October 15, 2026 09:12

fix compose test

2109b39

Signed-off-by: Anton Puzanov <antonp@il.ibm.com>

match uv formatting

01714d0

Signed-off-by: Anton Puzanov <antonp@il.ibm.com>

fix schedule signature

bdd9b94

Signed-off-by: Anton Puzanov <antonp@il.ibm.com>

freunda approved these changes Oct 15, 2026

View reviewed changes

freunda merged commit 7da02b0 into main Oct 15, 2026
4 checks passed

antonpibm deleted the bugfix/139-fa3-schedule-full-cudagraph-prefix-caching-fa3_fixing branch October 15, 2026 10:04

Navigation Menu

Repository navigation

Bugfix/139 fa3 schedule full cudagraph prefix caching fa3 fixing - #146

Bugfix/139 fa3 schedule full cudagraph prefix caching fa3 fixing#146
freunda merged 23 commits into
mainfrom
bugfix/139-fa3-schedule-full-cudagraph-prefix-caching-fa3_fixing

antonpibm commented Oct 14, 2026 •

edited

Loading

Uh oh!

antonpibm commented Oct 14, 2026

Uh oh!

codecov-commenter commented Oct 14, 2026

Uh oh!

github-actions Bot commented Oct 15, 2026 •

edited

Loading

Uh oh!

github-actions Bot commented Oct 15, 2026 •

edited

Loading

Uh oh!

antonpibm commented Oct 15, 2026

Uh oh!

antonpibm commented Oct 15, 2026

Uh oh!

antonpibm commented Oct 15, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

Conversation

antonpibm commented Oct 14, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

FA3 FULL-cudagraph correctness (#139) + vLLM 0.26–0.30 support and test-harness fixes

1. Test fixes — vLLM equivalence timing & integration

Engine teardown never waited (the "timing" problem)

enforce_eager reverted in the shared runner

Integration-suite API breakage (0.29 deletions)

2. Audio changes

3. _doubled_padding_mask in the SR decoder

4. FA3 correctness on Granite Switch (#139)

4a. FA3 AOT-schedule sizing for the switch's own attention heads (#139, eager path)

Uh oh!

antonpibm commented Oct 14, 2026

Uh oh!

codecov-commenter commented Oct 14, 2026

Codecov Report

Uh oh!

github-actions Bot commented Oct 15, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ GPU tests passed — vllm26-multi

Uh oh!

github-actions Bot commented Oct 15, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ GPU tests passed — vllm30-multi

Uh oh!

antonpibm commented Oct 15, 2026

Uh oh!

antonpibm commented Oct 15, 2026

Uh oh!

antonpibm commented Oct 15, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

antonpibm commented Oct 14, 2026 •

edited

Loading

`enforce_eager` reverted in the shared runner

3. `_doubled_padding_mask` in the SR decoder

github-actions Bot commented Oct 15, 2026 •

edited

Loading

✅ GPU tests passed — `vllm26-multi`

github-actions Bot commented Oct 15, 2026 •

edited

Loading

✅ GPU tests passed — `vllm30-multi`