Skip to content

Basic attention kernel that supports cached KV + (multi-)prompts - #24

Merged
suquark merged 25 commits into
mainfrom
mixed_attn_kernel
Apr 5, 2023
Merged

suquark merged 25 commits into
mainfrom
mixed_attn_kernel

Conversation

@suquark

@suquark suquark commented Apr 3, 2023 •

Copy link
Copy Markdown
Contributor

This PR implements a basic and not highly optimized kernel to support cached KV with multiple import prompts.

@suquark
suquark requested a review from WoosukKwon April 3, 2023 12:28
Comment thread csrc/attention_kernels.cu Resolved

@WoosukKwon WoosukKwon left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks a lot for your effort. Left minor comments.

Comment thread csrc/attention_kernels.cu Resolved
Comment thread csrc/attention_kernels.cu Resolved
Comment thread csrc/attention_kernels.cu Resolved
Comment thread csrc/attention_kernels.cu Resolved
Comment thread csrc/attention_kernels.cu Resolved
@suquark
suquark requested a review from WoosukKwon April 3, 2023 22:37
@suquark

suquark commented Apr 4, 2023

Copy link
Copy Markdown
Contributor Author

@WoosukKwon any comments?

slyalin pushed a commit to slyalin/vllm that referenced this pull request Apr 4, 2024
z103cb referenced this pull request in z103cb/opendatahub_vllm May 9, 2024
`format.sh` now has mypy checks after pulling in upstream changes. This
PR makes the mypy suggested modifications to our code.

---------

Signed-off-by: Travis Johnson <tsjohnso@us.ibm.com>
z103cb referenced this pull request in z103cb/opendatahub_vllm May 9, 2024
`format.sh` now has mypy checks after pulling in upstream changes. This
PR makes the mypy suggested modifications to our code.

---------

Signed-off-by: Travis Johnson <tsjohnso@us.ibm.com>
z103cb referenced this pull request in opendatahub-io/vllm May 9, 2024
`format.sh` now has mypy checks after pulling in upstream changes. This
PR makes the mypy suggested modifications to our code.

---------

Signed-off-by: Travis Johnson <tsjohnso@us.ibm.com>
dtrifiro pushed a commit to dtrifiro/vllm that referenced this pull request May 21, 2024
@alixiaodi alixiaodi mentioned this pull request Aug 2, 2024
njhill pushed a commit to njhill/vllm that referenced this pull request Nov 6, 2024
heheda12345 added a commit to heheda12345/vllm that referenced this pull request Sep 29, 2025
Signed-off-by: Chen Zhang <zhangch99@outlook.com>
yma11 added a commit to yma11/vllm that referenced this pull request Nov 14, 2025
* fix vision attn for whisper/llama4

Signed-off-by: Yan Ma <yan.ma@intel.com>

* address comments

Signed-off-by: Yan Ma <yan.ma@intel.com>

---------

Signed-off-by: Yan Ma <yan.ma@intel.com>
dik654 pushed a commit to dik654/vllm-for-study that referenced this pull request Nov 18, 2025
New Industry Use Cases (vllm-project#21-30):
- vllm-project#21 Game Development: AI game testing + balance tuning
- vllm-project#22 Construction: Vision AI safety inspection
- vllm-project#23 Agriculture/Smart Farm: Crop monitoring + pest detection
- vllm-project#24 Government/Public: Document automation + citizen services
- vllm-project#25 Energy/Utilities: Grid monitoring + anomaly detection
- vllm-project#26 Environment/Sustainability: Carbon tracking + ESG reporting
- vllm-project#27 Fashion/Apparel: Trend analysis + inventory optimization
- vllm-project#28 Sports/Fitness: Performance analytics + tactical analysis
- vllm-project#29 Automotive/Mobility: Autonomous driving simulation
- vllm-project#30 Space/Aerospace: Satellite image analysis

Advanced Architecture Patterns:
1. Event-Driven Pattern: Webhook → Event Bus → Agent triggers
2. Streaming Pattern: Large dataset processing with chunking
3. Batch Processing Pattern: Celery-based parallel processing
4. Circuit Breaker Pattern: Fault tolerance + auto recovery
5. CQRS + Event Sourcing: Command/Query separation
6. Saga Pattern: Distributed transaction management

Guide now covers:
- 30+ industry-specific MCP implementations
- 6 production-ready architecture patterns
- Real-world scalability solutions
- Enterprise integration strategies
- Total: 8,672 lines (from 7,249)
amd-xavierwang pushed a commit to amd-xavierwang/vllm that referenced this pull request Apr 8, 2026
Asymmetric AWQ models (zero_points=True) previously fell back to the
slower ConchLinearKernel. This adds zero-point support to the skinny
GEMM int4 kernels so asymmetric models route through the fast path.

A `bool HAS_ZERO_POINTS` template parameter (default false) is added to
the existing kernel templates. When false, the compiler eliminates all
zero-point code, producing identical codegen to the current path. When
true, per-group zero-point subtraction is inserted after dequantization
using the algebraic identity:
  (nibble - zp) = (nibble - 8) - (zp - 8)

The existing dequant (which always produces nibble-8) is unchanged;
pre-adjusted (zp-8) values are subtracted afterward.
jxmorris12 added a commit to jxmorris12/vllm that referenced this pull request Apr 21, 2026
starpit added a commit to starpit/vllm that referenced this pull request May 26, 2026
prefill + skip-logging eprintlns

- codegen.rs: wire I::RopeAppend in dispatch_instruction_to_render so
  prefill canonicals containing fused_qkv_rope_cache reach
  render_fused_qkv_rope_cache and emit role bodies. Without this arm
  the dispatch fell through to `_ => Ok(None)` which silently dropped
  emit_for_canonical / inventory::submit, while build_mega_tape kept
  emitting — explaining the "build_mega_tape says emitted but cu-emit
  inventory empty" mismatch.
- cuda_emit/render.rs: gate_up_activate_mul prefill rewritten to the
  WGMMA_N + N_MICRO microtile loop pattern matching the gemm and
  gemm_add prefill renders, so the inner mma uses cols ∈ [16, 256]
  step 16 (TK 2.0 wgmma::base specialization range from
  third_party/thunderkittens/include/ops/group/mma/base/base.cuh:28-43).
- ferrite-forward-macro/src/codegen.rs: emit_slice_render_with_loop_expansion
  now eprintln!s the canonical / instruction kind on each render-skip
  return so silent drops are visible in cu-emit logs.

Pod two-pass build green at 4m 57s under the 5-min cicc cap; serve +
curl E2E on Llama-3.2-1B-Instruct returns coherent generation
(\" Paris. The Eiffel Tower is located in Paris...\") for the 5
canonicals currently emitted (m_1_sk_128/192, m_8_sk_128,
m_64_sk_128/512). M=512 / M=4096 prefill canonicals are tracked
separately under tasks vllm-project#23/vllm-project#24.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
heliubj18 added a commit to heliubj18/vllm that referenced this pull request Aug 8, 2026
401 of the step's 402 cases need no device, but
test_scheduler.py::test_async_scheduling_pp_allows_rescheduling_with_output_placeholders
builds a ParallelConfig with pipeline_parallel_size=2. It never touches a
second card - the step allocates 0.0 GiB across its whole run - but on a
CUDA image ParallelConfig validates world size against the visible GPU
count and rejects 2. Upstream never sees this: the step depends on
image-build-cpu, where that check is inert. So the marking upstream gives
a step describes the image it runs on, not what its cases ask for, and
_gpus_for() had no way to say otherwise.

Adds docker.devices_by_step, a per-key override of the device count that
indexes into the existing gpus_by_devices map. Only v1-others-cpu uses
it; the generated pipeline differs from before by exactly one line, and
the other CPU-marked steps still get "device=0".

Two cards costs nothing when the step allocates no VRAM, and leaving it
at one meant a guaranteed failure every run - which is expensive now that
fail_fast is back on. In vllm-project#24 it cancelled Basic Models Test and Torch
Stable ABI Audit, both of which pass (3m19s in vllm-project#17). I had written this
case off as "not worth handling" without accounting for that.

Co-authored-by: Claude Opus 5
wenwzhan pushed a commit to wenwzhan/vllm-kimi-fz that referenced this pull request Aug 12, 2026
…cheduling

[Kimi-K3][ROCm] Do not run DSpark under both data parallelism and async scheduling
Chloe-orz pushed a commit to Chloe-orz/vllm-maintomain that referenced this pull request Nov 5, 2026
Merge vllm-data-comm/vllm main at
928b990 while retaining the local gate
fix b40d5cf as the first parent.
Preserve all existing commits, original SHAs and ancestry without rebasing,
reordering or squashing either history.

Bring in PR vllm-project#54: remove the long-request exclusive prefill batch rule,
its model-specific threshold argument and the associated upstream test
and documentation updates. Keep the incoming resume-first, FCFS picker
within its token budget and native request-slot limit.

The scheduler auto-merges without textual conflicts. Retain the local
explicit num_running_reqs argument to is_gated and compute free_seq from
len(self.running). Waiting requests continue to affect routing scores but
cannot reserve native slots before admission. Preserve shared edge slots
across DP routes, the in-flight resume exemption, per-DP pending caps and
the current Host Staging token budget. Relative to the incoming main,
only the existing four-file gate-fix patch remains.

Pair with vllm-ascend old-main-1026 fast-forwarded to
14602594f485da18702e57d8c10dade8c700b734. That upstream history includes
our graph-placeholder storage fix through PR vllm-project#24, the cloud DSV4 Hash
routing repair in PR vllm-project#26 and the dead FlashComm v1 sharding removal in
PR vllm-project#27. Keep the runner-owned placeholder and both slice-based uses.

Validation: source/diff review only. No tests, lint, pre-commit, CI, syntax
checks, startup or benchmarks were run, as requested. NPU validation is
pending. Backup branches were created before either update.

Signed-off-by: zzyccccccccc <zzyccccccccc@noreply.gitcode.com>
cxywhite added a commit to cxywhite/vllm that referenced this pull request Nov 8, 2026
…ucible config

- vllm submodule -> 6e1011be4 (tau scalar EMA sum unit + TTFT-margin force-admit
  + WT_SIMPLIFY_PREDICT predicted_remaining).
- test_evaluation_scheme4_rr_compare.sh: pass --activation-predict to decode so
  scheme4 (greedy) tracks predictive pressure too (admission-only A/B, pitfall vllm-project#24).
- experiment_log.md: o2000/o1600/o1200/o800 sweep, TTFT-margin finding (o2000
  margin=250 +24%, o1600 margin=100 +4.3%), 复现约定 + 已证伪路径 + result index.
- winning_config.md: reproducible config snapshot (model/PD/gpu_util/dataset/theta/margin).
- admission_ab_base: qwen_o800/o1200/o1600/o2000 (output-mean scaled), burstgpt_real (real trace); traces sanitized sk- -> sk_ for secret scanning.
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