Skip to content

Implement LLaMA - #9

Merged
zhuohan123 merged 16 commits into
mainfrom
llama
Mar 30, 2023
Merged

zhuohan123 merged 16 commits into
mainfrom
llama

Conversation

@WoosukKwon

@WoosukKwon WoosukKwon commented Mar 26, 2023 •

Copy link
Copy Markdown
Collaborator

TODO:

  • Test against HF implementation
  • Add TP support (@zhuohan123)

@WoosukKwon

Copy link
Copy Markdown
Collaborator Author

@zhuohan123 Please feel free to approve and merge this PR once you think it's ready.

@zhuohan123
zhuohan123 self-requested a review March 29, 2023 06:37
@zhuohan123
zhuohan123 merged commit 80a2f81 into main Mar 30, 2023
@WoosukKwon
WoosukKwon deleted the llama branch April 12, 2023 03:12
v1nc3nt27 pushed a commit to v1nc3nt27/vllm that referenced this pull request Sep 12, 2023
dont error if user doesnt have kernels installed
bigPYJ1151 pushed a commit to bigPYJ1151/vllm that referenced this pull request Dec 29, 2023
Srinivasoo7 pushed a commit to Srinivasoo7/vllm that referenced this pull request Mar 4, 2026
…Manager

- Add store_threshold >= 2 validation in FilterReusedOffloadingManager
  constructor (mirrors the existing max_tracker_size >= 1 guard)
- Fix cpu.py gate from > 1 to >= 2; update comment to clarify that
  values < 2 disable filtering
- Add internal assertions to test_filter_reused_manager to verify
  tracker eviction and count reset (Comments vllm-project#8 and vllm-project#9)
- Remove tests/v1/kv_offload/__init__.py (not needed for pytest discovery)
- Remove accidentally tracked dev-workflow files (.patch, diff*.txt,
  error.txt, log files, mypy/test output files)

Signed-off-by: Srinivasoo7 <158864704+Srinivasoo7@users.noreply.github.com>
danisereb pushed a commit to de-inf/vllm that referenced this pull request Apr 5, 2026
…-nongated-moe

[Bugfix] Fix BF16 trtllm-gen MoE weight corruption for non-gated models
starpit added a commit to starpit/vllm that referenced this pull request Apr 16, 2026
Seven-commit Gemma2 push landed this session. Compiler stays
architecture-agnostic — the acid test passed for generality:
zero "gemma" strings below the parser, the entire architecture
fits in DSL body + JSON configs + library Impls.

New sections:
- "Next session starts here" rewritten around the Gemma2 state:
  coherent output on common prompts, known hard-fail on
  `test_cuda_correctness_gemma2_2b` prompt 7 (a translation
  prompt with apostrophe-quoted text), ruled-out causes listed
  (not CUDA graphs, not final softcap, not ferrite-vs-hand-written
  drift, not the embed-scale/sliding-flip bugs I already fixed).
- "Commit history" grew a fresh block of 8 commits.
- "Landed since the correctness fix" gained a full item vllm-project#7 for
  the Gemma2 port, enumerating each sub-piece and its commit.
- Gap vllm-project#5 went from "sliding-window / Gemma2" to "Gemma3" (the
  logical next architecture using the same DSL constructs).
- Gap vllm-project#9 repurposed from "per-layer golden diff harness (nice
  to have)" to "per-layer tensor-dump debug harness" — the
  concrete tool needed to unblock the prompt-7 failure. Spelled
  out the approach including the forward_backbone quirk (for
  Gemma2 it returns pre-softcap logits, not pre-lm_head
  hidden_states, because the DSL's terminal tile is
  `tanh_softcap`).
- "Already landed" list gained nine new bullets covering the
  DSL constructs, the scalars mechanism, Gemma2 Impls, kernel
  `weight_offset` param, Gemma2 DSL + configs, and
  `Gemma2Ferrite` cuda_worker wiring.
- Gap vllm-project#4 (Granite multipliers) updated: the DSL now expresses
  scalar multiply via the Gemma embed-scale plumbing.

Honest accounting of the known-fail: correlation r=0.11 between
our post-softcap logits and Python vLLM's on prompt 7. Divergence
is upstream of softcap, specific to this prompt's numerics, and
likely lives in a shared CUDA kernel (flash-attn softcap + sliding,
or cuBLAS gemm on Gemma2's unusual head_dim ratios). Both ferrite
and the pre-existing hand-written Gemma2 path exhibit the same
failure — the fix, once found, applies to both.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
starpit added a commit to starpit/vllm that referenced this pull request Apr 16, 2026
Updates "Next session starts here" with the actual root cause and
fix (commit e2669d6). Marks Gemma2 correctness green alongside
SmolLM/Qwen2 — all three through *Ferrite paths. Replaces the
obsolete per-layer-tensor-dump gap (vllm-project#9) with the newly surfaced
chat-completion double-BOS bug at engine.rs:2857.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
starpit added a commit to starpit/vllm that referenced this pull request May 9, 2026
Add fused_gate_up_silu_mul_decode_f16_specialized: M=1-only
specialization that drops simdgroup_matrix entirely. For a single
query row the MMA tile machinery is overhead — 7/8 of every loaded
A fragment is zero-padded waste, and the partial-M scratch path
inserts a threadgroup_barrier into every K iteration.

The decode kernel uses 256 threads/group = 8 simdgroups x 32 lanes;
each simdgroup owns one output and accumulates K/32 partial products
per lane, reducing via simd_sum. No threadgroup-scratch staging, no
per-K-iter barriers, every device read coalesces across the 32
lanes. Bandwidth-bound at TinyLlama K=2048 (~7ms/forward
back-of-envelope, vs the 8x8 single-simdgroup tile that scheduled
704 tiny threadgroups per layer with one barrier per K-iter).

Pipeline picker in interpreter::metal::pipelines::kernel_msl_names
selects the decode symbol when bucket_m == 1, the matrix symbol
otherwise. Function constants live at indices 3/4/5 (vs 0/1/2 for
matrix) so both kernels compile cleanly into the same library.
Lowering picks dispatch shape per bucket_m: (ceil(N/8), 1, 1)
threadgroups x (256, 1, 1) threads for decode vs the existing 8x8
matrix tile for prefill.

ferrite-forward metal lib suite: 54/54 (was 53, +1 for the new
fused_silu_decode_bag indices test). The fused_mlp_matches_cpu_golden
M=1 case now exercises the decode kernel; fused_mlp_tinyllama_shape
(M=1, N=5632, K=2048) also routes through it and passes at 5e-2.

Prefill (M >= 2) still uses the 8x8 single-simdgroup tile —
larger-tile / multi-simdgroup upgrade tracked as task vllm-project#9.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
almogtavor referenced this pull request in almogtavor/vllm Jun 2, 2026
starpit added a commit to starpit/vllm that referenced this pull request Jun 8, 2026
…ck + rt_alias exhaustive coverage

Closes audit 2026-06-08 findings vllm-project#9, vllm-project#12, vllm-project#13, vllm-project#14 (4 of 21).

== Patch 3 (audit finding vllm-project#9, vllm-project#13): validator's "validator-checked" claims

`rt_alias_pass`'s module doc claimed its postcondition (every
RegTileSlot referenced by an Instr exists in tape.reg_tile_arena)
was "validator-checked". The validator never actually inspected it
— only a unit test in rt_alias.rs did. A dangling slot reference
would reach the player, emit CUDA referencing an undeclared
rt_<slot>, and either fail at nvcc compile or silently UB on id
collision.

Wire the check into validate_tk_tape:

  - New TkValidationError::RegTileSlotNotInArena { slot, at } variant.
  - New `check_reg_tile_arena_membership(tape, errors)` that walks
    every Instr exhaustively (rt-bearing + non-rt-bearing variants
    listed explicitly, no `_ => {}` catch-all) and asserts each
    referenced RegTileSlot is a key in tape.reg_tile_arena.
  - Two new tests: one that injects a dangling slot and asserts the
    error fires, one that mints via the typed RegTileId<...> path
    and asserts validation passes.

The validator's variant list MUST stay in lockstep with
`passes::rt_alias::instr_rt_accesses` — adding a new rt-bearing
Instr variant is now a rustc E0004 at three sites (validator + the
two rt_alias matches) instead of a silent skip.

== Patch 3 documentation drift cleanup (audit finding vllm-project#13)

`page_coalesce_pass` and `split_oversized_loads_pass` module docs
also claimed "validator-checked" postconditions that the validator
never inspected (PageId.0 < NUM_PAGES; LoadAsync.tile.byte_size <=
PAGE_SIZE). The validator is stage-blind — pre-coalesce tapes
legitimately hold PageIds beyond NUM_PAGES, pre-split tapes
legitimately hold oversized External LoadAsyncs. Updated both
docs to say "in-pass-checked, NOT validator-checked" with the
reason. The in-pass defensive walks already enforce these
postconditions; the docs now reflect that.

`rt_alias_pass`'s doc updated to point at the new
`check_reg_tile_arena_membership` site.

== Patch 4 (audit finding vllm-project#12, vllm-project#14): rt_alias exhaustive coverage

Two `_ => {}` catch-alls in `passes::rt_alias` —
`instr_rt_accesses` (line 208) and `rewrite_instr` (line 481) —
silently absorbed any future rt-bearing Instr variant added to
tk_tape::Instr. The variant would be excluded from liveness
tracking AND from rewriting, leaving stale slot ids in the
rewritten tape (silent-wrong-numerics on id collision; nvcc
error if the collision was caught).

Sister pass `passes::page_coalesce` was hardened to be exhaustive
after audit S2 (page_coalesce.rs:187-193 enumerates every
non-PageId-bearing variant, no catch-all). `rt_alias` mirrors that
pattern now: both matches list every Instr variant explicitly. New
rt-bearing variants become rustc E0004 at the pass-add site.

Also removed the fail-soft `slot_remap.insert(slot, slot); continue;`
at lines 335-340 of `coalesce`. An arena-missing slot is a real
bug (lowering minted via a path that bypassed the typed RegTileId
constructor, OR a prior pass evicted the entry while leaving the
Instr reference intact). Replace the silent identity-remap with
a panic carrying full context (slot id, first-def index). Patch 3's
new validator check catches the same gap from a different angle —
belt and suspenders.

While there I also added the four PageBarrierWait* variants
(Static{P0,P1}, LoopStart{0,1}) to rt_alias's non-rt-bearing list
— they were being absorbed by the now-deleted `_ => {}` catch-all.
They have no RegTileSlot fields so the behavioral change is
none, but the explicit listing is the point.

== Test results

ferrite-wavefront --lib: 114 -> 116 (+2 new arena-membership tests).
All 8 pre-existing split_oversized_loads tests + all 6 rt_alias
tests + all 11 tk_tape tests pass on the H100 pod (nick).

== What's left from the audit punch list

Patches 1, 2, 5 still pending (page-shape contract sealing,
AttnDecode KvCache layout, shmem-spill pass). Patches 3+4 are
foundational hardening that closes 4 of 21 findings without
touching the substrate or the lowering.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Signed-off-by: Nick Mitchell <nickm@us.ibm.com>
aoshen02 referenced this pull request in aoshen02/vllm Jul 31, 2026
…ownership

[R3] Move sidecar layout validation out of scheduler
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants