Use FlashAttention for `multi_query_kv_attention` by WoosukKwon · Pull Request #4 · vllm-project/vllm

WoosukKwon · 2023-03-02T05:05:56Z

This PR is to use FlashAttention kernels for multi_query_kv_attention, which performs masked attention for the prompt inputs.

Pros

FlashAttention is fast and memory-efficient.
FlashAttention supports 1D inputs and only invokes a single kernel to handle multiple sequences with variable lengths.

Cons

FlashAttention does NOT support FP32.
FlashAttention does not support head_size > 128. (This is fine for all models except GPT-J).
- Ref: Support for 256 head dim Dao-AILab/flash-attention#67
FlashAttention does not support attention bias (GPT-J, BLOOM, LLaMA).

Besides, note that FlashAttention does not support cached KV, which is required for interactive generation.

Tested models:

OPT-125M
OPT-350M
OPT-1.3B
OPT-2.7B
OPT-6.7B
OPT-13B

Tested GPUs:

A100

* Init * refine

Support for optimum-intel models

…o-model-executor Adapt OpenVINO CPU plugin implementation

BA-78760: Jamba * Add support for n concat and splitting * change naming * input_metadata is a dict list now in order to pass "n" * clean up code from unecessary changes and prints * Remove kv cache allocation in case of mamba layer * Add the considerations of mamba layer cache into the num of blocks calculation * Delete mamba cache after profile * Remove prints * Cleaning * - and not _ for requirements Approved-by: Tomer Asida

patching for having type su

…ect#4 magic_wand semi_structured_sparse_tensor_linear branch integrates 2:4 semi-structured sparsity into SparseTensor. This PR adds a new sparsity config for 2:4 sparsity to neuralmagic-vllm, using the SparseTensor 2:4 support. This PR also refactors the sparse linear method into a separate file, vllm/model_executor/layers/sparsity/sparse_w16a16_linear_method.py, which supports all sparsity formats.

- Section 4.3: Update single-request benchmarks with post-fix data (Config 3 TPOT 9.88ms->7.86ms, Config 4 TPOT 8.41ms->7.31ms) - Section 5.2.1: Rewrite from "pending fix" to "fixed and verified" with profile data comparing NCCL ops before/after fix - Section 5.2.2: Update concurrent TPOT worsening percentages with new single-request baselines - Section 5.3.2: Update single-request ITL comparison - Section 7: Add config3_fix profile directory and NCCL comparison table - Section 8: Update conclusion vllm-project#3 and vllm-project#4 with fix results Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>

Three layers of work, all driving toward "load real-world IQ-quantized GGUFs and run them coherently with quantized weights kept on GPU." GGUF quantized-storage default - Default load path now keeps GGUF weights compressed on GPU instead of dequant-to-BF16, saving ~2x weight memory. FERRITE_DEQUANT_AT_LOAD=1 opts back into the legacy dense path for debugging. - Fix the load-time host_buf race: pinned H2D for quantized weights was reusing host_buf for the next tensor before the previous DMA finished, leaving subsequent quantized weights corrupted on GPU. One-line stream_synchronize before the host_buf reuse. - Unpin Q6_K from max_ncols_y=1 (cargo cult, fixed by the sync above). - mul_mat_vec_q template: align rows_per_cuda_block with upstream (ncols_y < 4 → 1, else 2). IQ kernel coverage - ggml IQ1_M (codebook + dequant + MMVQ + dispatch). - Expanded IQ* support: IQ1_S / IQ2_XXS / IQ4_NL / IQ4_XS dequant + MMVQ kernels, sign-bit machinery shared across types (kmask_iq2xs / ksigns_iq2xs / unpack_ksigns / __vcmpne4 / __vsub4 / ggml_cuda_dp4a). - Add IQ2_XS (GgufDType tag 17): GgufDType + GgmlDType variants, the 512-entry iq2xs_grid codebook (uint64), block_iq2_xs struct, dequantize_block_iq2_xs (+ DEQUANTIZE_K macro), vec_dot_iq2_xs_q8_1, mul_mat_vec_iq2_xs_q8_1_cuda1, the three launch wrappers, extern "C" declarations and dispatch arms in ggml.rs. Mistral-7B-Instruct-v0.3 -IQ2_S.gguf (which mixes IQ2_XS for attention with IQ2_S for FFN) now loads and both `[ggml dispatch] first matmul via IQ2_XS` and `IQ3_S` notifiers fire during forward. - Batch IQ prefill: 56 new mul_mat_vec_iq*_q8_1_cuda2 through _cuda8 template instantiations across 8 IQ types, each launcher updated to switch on ncols_y exactly like Q4_K, max_ncols_y cap lifted from is_iq_quant() ? 1 : 8 to a flat 8. Reduces prefill kernel launches ~8x for IQ-family weights. Probe infrastructure (stays in the tree) - crates/ferrite-kernels/src/layers.rs has a ggml_probe module behind two env vars, off by default: - FERRITE_PROBE_MMVQ=1 — runs both the quantized kernel AND a dequant+cuBLAS F32 reference per GgmlLinear::forward, prints per-call max_abs_diff / rel / NaN counts. Used to isolate the host_buf race; now used to validate the cuda2..cuda8 IQ kernels. - FERRITE_USE_REFERENCE=1 — replaces ggml_matmul with the F32 cuBLAS reference at every linear. If inference becomes coherent under this flag, kernels are the bug; if it stays broken, look upstream. End-to-end coverage - Llama-3.2-1B chat tests for Q4_K_M / IQ1_S / IQ1_M / IQ2_XXS / IQ4_NL / IQ4_XS pass with FERRITE_PROBE_MMVQ=1 showing rel <= 2.9% at M=42 prefill (IQ4_NL <= 0.82%, IQ4_XS <= 1.1%, IQ1_M <= 0.81%, IQ1_S <= 2.9%, IQ2_XXS <= 1.3%); for every type the batched cuda2..cuda8 path is at least as accurate as the cuda1 baseline. - assert_coherent_text strengthened. - Mistral-7B IQ2_S / IQ3_S regression-marker tests added under #[ignore]. They go green when the tokenizer-from-GGUF fallback (vllm-serve/src/llm.rs:tokenize_text) is fixed — currently every prompt's ASCII bytes get used as token IDs because bartowski's GGUF-only repo ships no tokenizer.json, which is what the documented "x_max=0 from L0 q_proj" symptom actually was. Docs - HANDOFF_QUANTIZED.md captures the current state: items vllm-project#2 (IQ2_XS) and vllm-project#3 (IQ batched prefill) closed, item vllm-project#1 reclassified as a tokenizer-fallback bug (not a quant-path bug), item vllm-project#4 (IQ decode DMMV) reclassified do-not-port — upstream llama.cpp deleted ggml/src/ggml-cuda/dmmv.cu entirely; ik_llama.cpp only ports DMMV for its custom IQ*_KT types not the standard ones. - vllm-parity/parity.csv: GGUF IQ row updated to enumerate all 8 wired types with per-type rel numbers; new rows for "GGUF dequant kernels (IQ types)", "IQ decode DMMV (wontfix)", and "Tokenizer from GGUF metadata". Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

The previous debug commit (`aabb517a9`) printed bindings for ICB compute commands but skipped the Gemm path (which goes through `resolve_gemm_buffers` rather than `record_compute_dispatch`). This left bug vllm-project#4 — the MPS-side `offset: 0u64` hardcoding fixed in the previous commit — invisible to the diagnostic. Add the matching print for `BucketStep::Gemm`: kernel/dims and `(buf, off)` for each of a / b / c. Drop with the rest of `aabb517a9`'s instrumentation when the metal correctness work is merged-ready. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

…oject#4 still open Update FERRITE_METAL_PROGRESS.md to reflect the third bug fix this session: `MPSMatrix.initWithBuffer:offset:` was hardcoded to 0, silently undoing the per-binding offsets the lowering + worker had threaded all the way to the Gemm encoder. Caught by extending `FERRITE_METAL_BAKE_DEBUG=1` to dump Gemm bindings — the offset discrepancy fell out as soon as we could see the bound `(buf, off)` vs MPS' offset=0 invariant. Post-fix the model output is finally a function of the input: "hi" samples token 29966 (`<`) and continues `<assistant>\n…`; different prompts produce different sequences. But the model still loops on chat-template-like tokens — at least one more bug remains, likely in attention math, RoPE, or kernel arithmetic at runtime conditions. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

Adds a Phase 5.K section documenting this session's startup-time work: 16 commits between 79da7db and f9ae0df, dropping warm init engine from 2.91s to 240ms on Llama-3.2-3B (~20% under vllm-mlx's 300ms baseline). Also closes Phase 5.J's bug vllm-project#4 (bf16 KV slot-0 race) which landed in 3658f2b, and updates the Notes block. Cross-references project_metal_unpacked_mlp_next.md for the structural follow-up (real MetalMulImpl + Instruction::ElementwiseMul) that would eliminate the remaining ~200ms gate_up pack memcpy and drop startup to ~50ms. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>

…3 + walker nuke + bug-class-vllm-project#1 ratchet Establishes the typed `MegaTape<S>` lowering pipeline in a new `ferrite-mega-ir` crate and migrates the megakernel codegen off the walker-based `emit_op` design onto a typed IR that the proc-macro constructs at expansion time. Per-op substrate witnesses move runtime checks (page-slot bounds, lifecycle, mbarrier phase math, scratch overlap, scratch budget, warp roles) toward construction- time invariants — the plan's six bug classes. Sprint 0+1a+1b — substrate vocabulary + crate extraction: - `ferrite-mega-ir` crate carved out of the macro crate so the IR is a real type-level artifact, not a TokenStream stub. - Sealed substrate witnesses: `Page<ID, S, State>` with lifecycle states (`Empty`/`Filled`/`Produced`), `ScratchRegion<OFFSET, BYTES, Scope>`, `MbarrierPhase<P>`, `WarpRoleTag<R>`, `IsValidSlot`, `Disjoint`, `WithinBudget`, `ArriveCountToPhase`. - MegaIR decoupled from `W` (the per-canonical Weights type) so the IR can be constructed without a model in scope. Sprint 2 — gate-up + qkv-rope + downproj migrations: - silu_upgate, gelu_upgate, down_proj_residual, fused_qkv_rope_cache lower through the typed IR. Each variant's `emit_*` function takes a typed node and renders CUDA via `ferrite_mega_ir::emit`, replacing the walker's per-variant bespoke `emit_op` arm. Sprint 3 — fused_add_rms_norm + embed migrations. Walker nuke: - 19 per-variant `emit_X` walker fns deleted; the typed pipeline is now the only path. `EmitCtx` + `WalkerLines` + the `emit_op_from_typed` bridge are gone — `ferrite_mega_ir::emit` is the single render surface. - `BaseStage::new(n, op_page_count, num_pages)` validates page budget at IR construction (bug class vllm-project#1). Tautological in the proc-macro today (base_stage always 0), but the typed gate is load-bearing for Phase 4 cross-op pipelining where non-zero `base_stage` will fail at construction instead of corrupting emitted CUDA at runtime. Squashes 123 in-progress commits (Sprints 0+1+2+3 iters, walker deletion, BaseStage ratchet) into one landing commit. Backup of the original history at `backup/worktree-ff-mega-codegen-pre-squash`. Plan: `MEGA_IR_PLAN.md` — Sprint A done (RmsNorm), Sprint B done (FusedQkvRopeCache), Sprint C in progress (FusedAddRmsNorm + SiluUpgate + GeluUpgate + DownProjResidual landed; remaining Sprint D ops pending). Bug class vllm-project#1 caught by typed gate; bug classes vllm-project#2/vllm-project#3/vllm-project#4 are NIGHTLY TODOs (need per- variant const generics on MegaOp values, blocked on stable-Rust limits noted in `BaseStage` type-level docs); bug class vllm-project#5 (scratch budget overrun summed across the tape) is the next ratchet. Signed-off-by: Nick Mitchell <nickm@us.ibm.com>

Per MEGA_IR_PLAN.md §10: substrate-aware lowering for the second Instruction variant. Bug classes targeted: vllm-project#1 (page bounds), vllm-project#2 (lifecycle: in / qkv-weight / cos-sin / q-out / k-out / v-out across loader → consumer → storer), vllm-project#3 (per-iter mbarrier phase math), vllm-project#4 (scratch overlap inside the per-tok loop), vllm-project#5 (within-budget scratch), vllm-project#6 (warp-role pairing). substrate.rs: - New `RopeScope` (sealed `IsScratchScope` + `IsScratchScopePub`) for the per-token Q/K rotation tiles. Sibling regions in this scope are required to be `disjoint_with`-discharged. - New `IterCount(u32)` newtype with `iters > 0` invariant — used to bump the cumulative arrive count by `iters * arrives_per_iter` (per-iter math is variant-internal: with arrives_per_iter == 1, the kernel's `phase ^= 1` toggle is correct iff the boundary phase matches the cumulative count parity at op start, which is what the lowering validates via MbarrierPhase::assert_matches). nodes.rs: - New `RotaryRef` helper newtype (non-empty path). Lets the typed node distinguish global rotary from per-layer-local rotary (`wm.rotary` vs `wm.rotary_local`) per the typed-fanout walker's injection. - New `MegaNode::FusedQkvRopeCache` variant. Substrate-proof fields: 6 PageIds (lifecycle-witnessed), 2 disjoint ScratchRegion<RopeScope> (Q/K rotation tiles), consumer_phase + storer_phase + iters, four WarpRoleTag<R>. Helper fields alongside: layer / qkv_weight / rotary / biased / interleaved. lower.rs: - `MegaTapeBuilder::push_fused_qkv_rope_cache` walks every page through Empty → Filled → Produced → Empty (typestate-burned), packs Q/K rotation tiles back-to-back in scratch, validates phase parity at the op boundary, releases pages, and bumps the cumulative arrive count by `iters` (per-iter math falls out of `(C+t)&1 == (C&1)^(t&1)`). - `lower()` arm dispatches `Instruction::FusedQkvRopeCache` with a per-tape `SlotAllocator` that picks 5 non-aliasing output / weight page ids modulo `num_pages`. Real pipeline-aware allocation lands in Sprint D. - `OpInput::iters: u32` added (default 1 via `OpInput::new`). weight_paths arity for Rope is required to be 2 (qkv-packed, rotary). New `LowerError::WrongWeightArity` and `LowerError::SubstrateBudgetTooSmall` variants. tape.rs: - `impl Debug for MegaTape` (size + substrate budget summary), so test panics can format `Result<MegaTape, LowerError>`. 9 new tests in `lower::tests` covering: well-formed Rope round-trip; arrive-count phase advance across an op chain (5-iter odd flips parity, 4-iter even keeps it); Q/K scratch overlap rejection; within-op page aliasing rejection; iters=0 rejection; empty-rotary-path rejection; lower() round-trip; weight-paths arity error. Existing RmsNorm tests updated for irrefutable-pattern → `let-else` since `MegaNode` now has 2 variants. 30 tests pass on H100 pod (was 21). `cargo clippy -p ferrite-mega-ir --all-features -- -D warnings` clean. No emit yet — Sprint A — D land the typed lowering for every variant first, syntactic emit after. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> Signed-off-by: Nick Mitchell <nickm@us.ibm.com>

…ilu,Gelu}Mul Per MEGA_IR_PLAN.md §10: substrate-aware lowering for the residual fold (`Instruction::Add` — what the plan calls "DownProjResidual"), the fused add+norm (pre-attention / pre-MLP layer entry), and the gate-up MLP fusion in both Silu and Gelu activation forms. Bug classes vllm-project#1–vllm-project#6 all exercised between the four variants. substrate.rs: - New `MlpScope` (sealed `IsScratchScope` + `IsScratchScopePub`) for the gate-up fusion's per-token-iter gate / up activation tiles. Sibling regions in this scope MUST be `disjoint_with`- discharged at lowering time (Sprint C's bug class vllm-project#4 site). nodes.rs: - New `MegaNode::Add`: 2 PageIds (delta + residual, lifecycle-walked), consumer/storer phase, 4 warp-role tags. No weight, no scratch — element-wise residual fold lives entirely in registers. - New `MegaNode::FusedAddRmsNorm`: 3 PageIds (delta + residual + weight) lifecycle-walked, partial_sums in `RmsNormScope` (sibling regions on a future RmsNorm op `disjoint_with`-discharged), phase parity, 4 role tags, layer + weight helpers. - New `MegaNode::FusedGateUpActivateMul` (single variant covering both `FusedGateUpSiluMul` and `FusedGateUpGeluMul` instructions — same substrate shape, distinguished by `activation: GateUpActivation { Silu | Gelu }`): 3 PageIds (in + packed gate-up weight + out), 2 disjoint `MlpScope` tiles (gate_buf, up_buf), boundary phase parity, IterCount, 4 role tags, layer + weight + activation helpers. lower.rs: - `MegaTapeBuilder::push_add`, `push_fused_add_rms_norm`, `push_fused_gate_up_activate_mul` — substrate-proof discharge per call: page allocation + bounds (vllm-project#1), lifecycle walks Empty → Filled → Produced → Empty (vllm-project#2), scratch (vllm-project#4 disjointness + vllm-project#5 within-budget), boundary phase parity (vllm-project#3 — per-iter math for the gate-up multi-iter shape falls out of `(C+t)&1 == (C&1)^(t&1)`), warp-role pinning (vllm-project#6). - `lower()` arms for `Instruction::Add`, `FusedAddRmsNorm`, `FusedGateUpSiluMul`, `FusedGateUpGeluMul`. - New `pick_distinct_slot` helper: picks the next free page id that doesn't collide with caller-supplied input/output slots. Used by Sprint C arms when the Instruction only supplies in/out and the lowering needs to synthesize a non-aliasing weight slot. 13 new tests in `lower::tests` covering: well-formed Add round-trip; self-aliasing rejection; arrive-count phase advance across two Adds; FusedAddRmsNorm well-formed shape; delta/residual aliasing rejection; gate-up Silu and Gelu lowering; gate/up scratch overlap rejection; Instruction-side round-trip for each variant; multi-op chain threading phase advance through Add → FusedAddRmsNorm → FusedGateUp (verifies the cumulative arrive count carries the correct parity at each boundary, including across the multi-iter gate-up bump). `lower_unmigrated_variant_returns_not_yet_lifted` test repointed at `Gemm` (Sprint D scope) since `Add` now lowers. 43 tests pass on H100 pod (was 30); `cargo clippy -p ferrite-mega-ir --all-features -- -D warnings` clean. No emit yet. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> Signed-off-by: Nick Mitchell <nickm@us.ibm.com>

WoosukKwon added 8 commits March 2, 2023 04:20

Add a FlashAttention test

5685bac

Define MAX_SEQ_LEN

320b20c

Minor

4932c71

Use FlashAttention for multi_query_kv_attention

c302754

Add more head sizes for test

95e5c0f

Add error msgs

1c28f4f

Enhance the server script

eb9e9a0

Add flash-attn to README

72052b7

WoosukKwon merged commit 3e9f991 into main Mar 2, 2023

WoosukKwon deleted the flash-attn branch March 2, 2023 05:13

TheBloke mentioned this pull request Jul 20, 2023

Can't launch OpenAI API server on newly installed vLLM in Docker - fastchat not found #537

Closed

tmm1 mentioned this pull request Aug 3, 2023

Fix the rushed out multi-query kernel #44

Closed

xiangyuT added a commit to xiangyuT/vllm that referenced this pull request Oct 24, 2023

Add BigDL Llama worker for batching on decoding (vllm-project#4)

02b4cac

* Init * refine

shanshanpt mentioned this pull request Nov 17, 2023

Run long conetxt error : CUDA error: an illegal memory access was encountered #1700

Closed

junior-zsy mentioned this pull request Nov 20, 2023

Error with 32k Long Text in chatglm2-6b-32k Model #1725

Closed

hongxiayang referenced this pull request in hongxiayang/vllm Feb 13, 2024

Use FlashAttention for multi_query_kv_attention (#4)

e7c912b

luo-cheng2021 pushed a commit to luo-cheng2021/vllm that referenced this pull request Mar 12, 2024

Merge pull request vllm-project#4 from slyalin/optimum_models

8a9862f

Support for optimum-intel models

luo-cheng2021 pushed a commit to luo-cheng2021/vllm that referenced this pull request Mar 25, 2024

Merge pull request vllm-project#4 from luo-cheng2021/luocheng/openvin…

658407a

…o-model-executor Adapt OpenVINO CPU plugin implementation

dlopes78 mentioned this pull request May 8, 2024

[Bug]: VLLM + tritonserver #4695

Closed

linxihui added a commit to linxihui/vllm that referenced this pull request May 14, 2024

Merge pull request vllm-project#4 from beagleski/bapatra/patching-for-su

7646e00

patching for having type su

Alexei-V-Ivanov-AMD mentioned this pull request May 16, 2024

[Speculative decoding][Re-take] Enable TP>1 speculative decoding #4840

Merged

afeldman-nm mentioned this pull request May 21, 2024

[Core] Cross-attention KV caching and memory-management (towards eventual encoder/decoder model support) #4837

Merged

yuhuixu1993 mentioned this pull request Jun 2, 2024

[Bug]: loading squeezellm model #5190

Closed

afeldman-nm mentioned this pull request Jun 3, 2024

[Bug]: VLLM_ATTENTION_BACKEND set to ROCM_FLASH only in GHA environment, overriding automatic backend selection; this breaks other kernel unit tests. #5208

Closed

oliver-li mentioned this pull request Jul 5, 2024

[Bug]: NCCL hangs and causes timeout #5484

Closed

haichuan1221 mentioned this pull request Jul 5, 2024

Support W4A8 quantization for vllm #5218

Merged

SongXiaoMao mentioned this pull request May 13, 2026

[Bug]: MTP speculative decoding crash with illegal memory access on long sequences (Qwen3.6-27B-FP8, v0.19.1) #40756

Open

1 task

DoradusResearch mentioned this pull request May 13, 2026

cumem allocator: double-free and stale error codes during sleep/wake cycles #36651

Open

1 task

maeehart mentioned this pull request May 18, 2026

[ROCm][DSv4] Functional fixes for DeepSeek V4 on MI300X (gfx942) #42893

Closed

russellb mentioned this pull request May 26, 2026

fix(processor): route MiMo-V2-Omni media fetch through MediaConnector #43117

Merged

JartX mentioned this pull request Jun 3, 2026

[Feature] Triton INT4 per-token-head KV cache quantization #40835

Merged

zt1024 mentioned this pull request Jun 4, 2026

[Feature] TRITON_MLA_SPARSE backend for SM8x/11x/12x DSA Sparse MLA Support #38476

Open

jeffye-dev mentioned this pull request Jun 11, 2026

Handle structured output grammar compilation failures #45012

Open

taneem-ibrahim mentioned this pull request Jun 14, 2026

Added real /v1/embeddings support for messages + chat_template_kw #45173

Merged

waynehacking8 mentioned this pull request Jun 15, 2026

[Bug]: Scheduler deadlocks without KV-cache connector #45414

Open

1 task

supercoolgreatcoder mentioned this pull request Jun 16, 2026

Add Nemotron-Labs-Diffusion (93.3% GSM8K @ 2461 tok/s, 2.9× SOL) supercoolgreatcoder/vllm-fork#1

Draft

9 tasks

natema mentioned this pull request Jun 18, 2026

[Feature]: DeepSeek-V4 Flash sm_80 (A100/A800) support #40851

Open

1 task

eastwood-c mentioned this pull request Jun 30, 2026

[Spec][V2] Support MTP speculative decoding under pipeline parallelism #46994

Open

paulbrav mentioned this pull request Jul 5, 2026

[Bugfix][Spec Decode][V1] Guard stale async spec prev rows #47199

Closed

4 tasks

dajinge mentioned this pull request Jul 9, 2026

[RFC]: Add API to restore free_block_queue allocation order for long-running RLHF / rollout sessions #40244

Open

1 task

biswapanda mentioned this pull request Jul 20, 2026

dynamo-prime-rl-integration biswapanda/vllm#1

Open

marcinkuk mentioned this pull request Jul 20, 2026

[Bug][XPU]: zeMemOpenIpcHandle INVALID_ARGUMENT on 2x Arc B50 (Battlemage) TP=2 #48953

Open

JohnQinAMD mentioned this pull request Jul 29, 2026

perf(rocm): Kimi-K3 B1 route/AttnRes/tail/KDA fusion (vLLM stack) JohnQinAMD/vllm-amd#3

Closed

slokesha mentioned this pull request Aug 9, 2026

[XPU] Enable MiniMax-M3 on non-CUDA backends: portable fallbacks for CUDA-only paths #51578

Open

Juhyun-Kim-Memphis mentioned this pull request Aug 10, 2026

[minimax-m3] MiniMax-M3-W4AFP8 serving patch set for H200 (9 patches, live on ganymede) vessl-ai/vllm#1

Merged

Navigation Menu

Uh oh!

Uh oh!

Use FlashAttention for `multi_query_kv_attention` - #4

Use FlashAttention for `multi_query_kv_attention`#4
WoosukKwon merged 8 commits into
mainfrom
flash-attn

WoosukKwon commented Mar 2, 2023 •

edited

Loading

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Uh oh!

Uh oh!

Conversation

WoosukKwon commented Mar 2, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Pros

Cons

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

WoosukKwon commented Mar 2, 2023 •

edited

Loading