Skip to content

Commit aa28f4c

Browse files
liuyao0322liuyao0322
authored
Merge branch 'main' into perf/glm53-h20-fp8-moe
2 parents 878ff03 + a1541f5 commit aa28f4c

923 files changed

Lines changed: 80249 additions & 10977 deletions

File tree

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

.claude/skills/ci-fails-buildkite/SKILL.md renamed to .agents/skills/ci-fails-buildkite/SKILL.md

File renamed without changes.

‎.agents/skills/kernel-microbenchmark/SKILL.md‎

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -42,6 +42,25 @@ description: Build, debug, and interpret vLLM GPU kernel microbenchmarks for CUD
4242
- Keep metadata setup, plan construction, allocation, random input generation,
4343
and logging outside the timed region unless that overhead is the experiment.
4444

45+
## Sanity-Check Reference Numbers
46+
47+
Use these as rough reference points for large, well-shaped workloads, not as
48+
gold standards, guaranteed peaks, or hard limits. Hardware SKU, clocks, shape,
49+
precision conventions, and the FLOP/byte accounting can move the result. A
50+
large gap is a prompt to investigate, not proof that a kernel is poor.
51+
52+
| Kernel regime | Hardware | Rough reference |
53+
| --- | --- | ---: |
54+
| Memory-bound, large batch | Blackwell | 6 TB/s |
55+
| BF16 GEMM | Blackwell | 2 PFLOP/s |
56+
| BF16 attention | B200 | 1.6 PFLOP/s |
57+
| FP8 GEMM | Blackwell | 4 PFLOP/s |
58+
| FP8 attention | B300 | 2.8 PFLOP/s |
59+
60+
The BF16 attention reference is approximately the 1613 TFLOP/s result reported
61+
by the FlashAttention-4 paper. Compare kernels only with matching workload and
62+
throughput conventions.
63+
4564
## Multi-GPU Benchmarks
4665

4766
- State whether the run is local or multi-node and report the GPU topology,

‎.agents/skills/triton-kernel-writing/SKILL.md‎

Lines changed: 62 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,62 @@
1+
---
2+
name: triton-kernel-writing
3+
description: Write or review Triton kernels for vLLM, with practical guidance for generated-code inspection, launch grids, indexing, specialization, tuning, and representative performance validation.
4+
---
5+
6+
# Triton Kernel Writing
7+
8+
## Implementation
9+
10+
- Follow the official
11+
[Triton semantics](https://triton-lang.org/main/python-api/triton-semantics.html).
12+
Check it when behavior may differ from Python or NumPy, especially type
13+
promotion, integer division and modulo, casts, broadcasting, and variable
14+
scoping.
15+
- Use the Triton kernel generated by `torch.compile` as a possible
16+
implementation to inspect. Print Inductor's generated code with
17+
`TORCH_LOGS="output_code" .venv/bin/python <script>` or enable
18+
`torch._logging.set_logs(output_code=True)` before the compiled function
19+
runs. Treat generated code as a reference, not as proof of correctness or
20+
optimality.
21+
- Find reasonable defaults for compile-time knobs such as `BLOCK_SIZE`, or use
22+
a small, legible heuristic when workloads need different choices. Use
23+
`triton.autotune` only when tuning is critical to performance, such as for a
24+
matrix multiplication. Otherwise prioritize simple code and fast startup.
25+
- Be careful to avoid unintended runtime JIT compilation. For example, put
26+
unimportant runtime integer scalars in `do_not_specialize`, especially those
27+
that may alternate between values such as 0 and 1, which can produce
28+
different specialization keys.
29+
- The Triton compiler does not guarantee safe ordering when a kernel writes to
30+
a pointer and subsequently reads from the same pointer. This pattern must
31+
have a `tl.debug_barrier()` between the write and read. The barrier
32+
synchronizes threads in the block; it does not synchronize separate program
33+
instances.
34+
35+
## Launch and Indexing
36+
37+
- `grid[1]` and `grid[2]` must be at most 65,535. Choose or flatten the grid
38+
order so those dimensions cannot exceed the limit for supported shapes.
39+
For example, `num_tokens` is commonly 8K or 16K, but users may configure 32K
40+
or more. If `num_tokens` is a grid dimension, it is safe to put it in
41+
`grid[0]` (or tile it).
42+
- Use `int64` for offset arithmetic when an index can exceed 32-bit range,
43+
especially for KV-cache addressing. Cast operands before multiplication or
44+
addition so an intermediate does not overflow in 32-bit arithmetic.
45+
- A `[num_tokens, num_heads]` grid can be a good low-latency mapping for decode,
46+
but it can be very slow for prefill. If the kernel serves prefill, consider
47+
tiling tokens or otherwise increasing the work and locality per program.
48+
49+
## Validation
50+
51+
- Check correctness at boundary shapes and at sizes that exercise masks and
52+
large offsets.
53+
- Choose accumulation and intermediate dtypes explicitly. Test numerically
54+
difficult inputs, not only random, well-scaled tensors.
55+
- Use `$kernel-microbenchmark` for benchmark construction, measurement, and
56+
interpretation.
57+
- Benchmark a sweep of `num_tokens` covering decode and representative prefill
58+
workloads. Include relevant head counts and dimensions when they affect the
59+
launch shape, and do not select an implementation or tuning heuristic from a
60+
single setup.
61+
- Include compilation or autotuning overhead when evaluating startup behavior;
62+
report steady-state kernel performance separately.

‎.agents/skills/triton-kernel-writing/agents/openai.yaml‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,4 @@
1+
interface:
2+
display_name: "Triton Kernel Writing"
3+
short_description: "Write robust, practical Triton kernels for vLLM."
4+
default_prompt: "Use $triton-kernel-writing to implement or review this Triton kernel."

‎.buildkite/hardware_tests/ascend_npu.yaml‎

Lines changed: 94 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4,8 +4,101 @@ steps:
44
- label: "Ascend NPU Test"
55
key: ascend-npu-test
66
soft_fail: true
7-
timeout_in_minutes: 20
87
no_plugin: true
98
device: ascend_npu
9+
source_file_dependencies:
10+
- vllm/
11+
- .buildkite/scripts/hardware_ci/run-npu-test.sh
12+
- "!vllm/platforms/cuda.py"
13+
- "!vllm/platforms/xpu.py"
14+
- "!vllm/platforms/tpu.py"
15+
- "!vllm/platforms/cpu.py"
16+
- "!vllm/platforms/zen_cpu.py"
17+
- "!vllm/cute_utils/"
18+
- "!vllm/tilelang_utils/"
19+
- "!vllm/device_allocator/"
20+
- "!vllm/vllm_flash_attn/"
21+
- "!vllm/kernels/"
22+
- "!vllm/model_executor/kernels/"
23+
- "!vllm/engine/"
24+
- "!vllm/entrypoints/"
25+
- "!vllm/ray/"
26+
- "!vllm/v1/attention/backends/flash_attn.py"
27+
- "!vllm/v1/attention/backends/flash_attn_diffkv.py"
28+
- "!vllm/v1/attention/backends/flashinfer.py"
29+
- "!vllm/v1/attention/backends/triton_attn.py"
30+
- "!vllm/v1/attention/backends/triton_attn_diffkv.py"
31+
- "!vllm/v1/attention/backends/rocm_attn.py"
32+
- "!vllm/v1/attention/backends/rocm_aiter_fa.py"
33+
- "!vllm/v1/attention/backends/rocm_aiter_unified_attn.py"
34+
- "!vllm/v1/attention/backends/flex_attention.py"
35+
- "!vllm/v1/attention/backends/hpc_attn.py"
36+
- "!vllm/v1/attention/backends/mamba1_attn.py"
37+
- "!vllm/v1/attention/backends/mamba2_attn.py"
38+
- "!vllm/v1/attention/backends/mamba_attn.py"
39+
- "!vllm/v1/attention/backends/short_conv_attn.py"
40+
- "!vllm/v1/attention/backends/turboquant_attn.py"
41+
- "!vllm/v1/attention/backends/cpu_attn.py"
42+
- "!vllm/v1/attention/backends/fa_utils.py"
43+
- "!vllm/distributed/device_communicators/pynccl.py"
44+
- "!vllm/distributed/device_communicators/pynccl_wrapper.py"
45+
- "!vllm/distributed/device_communicators/pynccl_allocator.py"
46+
- "!vllm/distributed/device_communicators/xpu_communicator.py"
47+
- "!vllm/distributed/device_communicators/cpu_communicator.py"
48+
- "!vllm/distributed/device_communicators/cuda_wrapper.py"
49+
- "!vllm/distributed/device_communicators/flashinfer_all_reduce.py"
50+
- "!vllm/distributed/device_communicators/aiter_custom_all_reduce.py"
51+
- "!vllm/distributed/device_communicators/custom_all_reduce.py"
52+
- "!vllm/distributed/device_communicators/quick_all_reduce.py"
53+
- "!vllm/distributed/device_communicators/all2all.py"
54+
- "!vllm/distributed/device_communicators/all_reduce_utils.py"
55+
- "!vllm/distributed/device_communicators/shm_object_storage.py"
56+
- "!vllm/distributed/device_communicators/symm_mem.py"
57+
- "!vllm/distributed/device_communicators/mnnvl_compat.py"
58+
- "!vllm/distributed/device_communicators/ray_communicator.py"
59+
- "!vllm/third_party/flashmla/"
60+
- "!vllm/third_party/pynvml.py"
61+
- "!vllm/_custom_ops.py"
62+
- "!vllm/_xpu_ops.py"
63+
- "!vllm/_aiter_ops"
64+
- "!vllm/collect_env.py"
65+
- "!vllm/connections.py"
66+
- "!vllm/env_override.py"
67+
- "!vllm/logits_process.py"
68+
- "!vllm/logprobs.py"
69+
- "!vllm/model_inspection.py"
70+
- "!vllm/outputs.py"
71+
- "!vllm/scalar_type.py"
72+
- "!vllm/scripts.py"
73+
- "!vllm/version.py"
74+
- "!vllm/assets/"
75+
- "!vllm/benchmarks/"
76+
- "!vllm/inputs/"
77+
- "!vllm/ir/"
78+
- "!vllm/parser/"
79+
- "!vllm/plugins/"
80+
- "!vllm/reasoning/"
81+
- "!vllm/renderers/"
82+
- "!vllm/tokenizers/"
83+
- "!vllm/tool_parsers/"
84+
- "!vllm/tracing/"
85+
- "!vllm/usage/"
86+
- "!tests/"
87+
- "!benchmarks/"
88+
- "!docs/"
89+
- "!examples/"
90+
- "!scripts/"
91+
- "!vllm/model_executor/layers/quantization/auto_awq.py"
92+
- "!vllm/model_executor/layers/quantization/auto_gptq.py"
93+
- "!vllm/model_executor/layers/quantization/inc/"
94+
- "!vllm/model_executor/layers/quantization/awq_triton.py"
95+
- "!vllm/model_executor/layers/quantization/experts_int8.py"
96+
- "!vllm/model_executor/layers/quantization/fbgemm_fp8.py"
97+
- "!vllm/model_executor/layers/quantization/fp_quant.py"
98+
- "!vllm/model_executor/layers/quantization/humming.py"
99+
- "!vllm/model_executor/layers/quantization/moe_wna16.py"
100+
- "!vllm/model_executor/layers/quantization/qutlass_utils.py"
101+
- "!vllm/model_executor/layers/quantization/quark/"
102+
- "!vllm/model_executor/layers/quantization/turboquant/"
10103
commands:
11104
- bash .buildkite/scripts/hardware_ci/run-npu-test.sh

‎.buildkite/hardware_tests/cpu.yaml‎

Lines changed: 16 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -18,6 +18,8 @@ steps:
1818
- vllm/model_executor/layers/attention/mla_attention.py
1919
- tests/kernels/moe/test_cpu_fused_moe.py
2020
- tests/kernels/moe/test_cpu_quant_fused_moe.py
21+
- tests/kernels/moe/test_zen_cpu_int8_moe.py
22+
- vllm/model_executor/layers/fused_moe/experts/cpu_moe.py
2123
- tests/kernels/test_onednn.py
2224
- tests/kernels/test_awq_int4_to_int8.py
2325
- tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
@@ -36,6 +38,7 @@ steps:
3638
pytest -x -v -s tests/kernels/attention/test_amx_mla.py
3739
pytest -x -v -s tests/kernels/moe/test_cpu_fused_moe.py
3840
pytest -x -v -s tests/kernels/moe/test_cpu_quant_fused_moe.py
41+
pytest -x -v -s tests/kernels/moe/test_zen_cpu_int8_moe.py
3942
pytest -x -v -s tests/kernels/mamba/test_cpu_short_conv.py
4043
pytest -x -v -s tests/kernels/test_onednn.py
4144
pytest -x -v -s tests/kernels/test_awq_int4_to_int8.py
@@ -59,7 +62,7 @@ steps:
5962
# bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 20m "
6063
# bash .buildkite/scripts/hardware_ci/run-cpu-compatibility-test.sh"
6164

62-
- label: CPU-Language Generation and Pooling Model Tests
65+
- label: CPU-Language Generation and Pooling Model Tests Shard %N
6366
key: cpu-language-generation-and-pooling-model-tests
6467
depends_on: []
6568
device: intel_cpu
@@ -72,10 +75,9 @@ steps:
7275
- tests/v1/e2e/test_cpu_linear_attn_chunked_prefix.py
7376
commands:
7477
- |
75-
bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 50m "
76-
pytest -x -v -s tests/models/language/generation -m cpu_model
77-
pytest -x -v -s tests/models/language/pooling -m cpu_model
78-
pytest -x -v -s tests/v1/e2e/test_cpu_linear_attn_chunked_prefix.py"
78+
bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 25m "
79+
pytest -x -v -s tests/models/language/generation tests/models/language/pooling tests/v1/e2e/test_cpu_linear_attn_chunked_prefix.py -m cpu_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB"
80+
parallelism: 3
7981

8082
- label: CPU-Quantization Model Tests
8183
key: cpu-quantization-model-tests
@@ -159,11 +161,16 @@ steps:
159161
bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 40m "
160162
VLLM_CI_ENV=0 pytest -x -v -s tests/models/multimodal/generation/test_qwen2_5_vl.py"
161163
162-
- label: "Arm CPU Test"
164+
- label: "Arm CPU Test Shard %N"
163165
key: arm-cpu-test
164-
depends_on: []
166+
depends_on:
167+
- cpu-arm64-image-build
165168
soft_fail: false
166169
device: arm_cpu
167170
no_plugin: true
168-
commands:
169-
- bash .buildkite/scripts/hardware_ci/run-cpu-test-arm.sh
171+
commands:
172+
- >-
173+
bash .buildkite/scripts/hardware_ci/run-cpu-test-arm.sh
174+
"$$BUILDKITE_PARALLEL_JOB"
175+
"$REGISTRY/$REPO:$BUILDKITE_COMMIT-arm64-cpu"
176+
parallelism: 3

‎.buildkite/image_build/image_build.sh‎

Lines changed: 48 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -160,6 +160,42 @@ print_bake_config() {
160160
(cd "$(dirname "${BAKE_CONFIG_FILE}")" && buildkite-agent artifact upload "$(basename "${BAKE_CONFIG_FILE}")")
161161
}
162162

163+
record_buildkit_trace() (
164+
local metadata_file="$1"
165+
local helper=".buildkite/scripts/ci-otel/ci_otel.py"
166+
local trace_file="${BUILD_TMP_DIR}/buildkit-trace.json"
167+
local build_ref record_id
168+
169+
if [[ "${BUILDKITE:-}" != "true" || ! -f "${helper}" ]] ||
170+
! command -v timeout >/dev/null 2>&1 ||
171+
! command -v python3 >/dev/null 2>&1; then
172+
return 0
173+
fi
174+
build_ref="$(timeout 3s python3 "${helper}" build-ref "${metadata_file}" "${TARGET}" 2>/dev/null)" || {
175+
echo "vLLM CI OTel: BuildKit metadata unavailable; trace skipped" >&2
176+
return 0
177+
}
178+
record_id="${build_ref##*/}"
179+
if [[ -z "${record_id}" ]]; then
180+
return 0
181+
fi
182+
183+
echo "--- :mag: Recording BuildKit trace"
184+
if ! timeout 20s docker buildx history trace "${record_id}" >"${trace_file}"; then
185+
echo "vLLM CI OTel: BuildKit history trace unavailable; upload skipped" >&2
186+
return 0
187+
fi
188+
189+
export CI_INFRA_BUILDX_REF="${build_ref}"
190+
export CI_INFRA_OTEL_SPOOL_DIR="${BUILD_TMP_DIR}/otel-spans"
191+
timeout 5s python3 "${helper}" record-buildkit "${trace_file}" || {
192+
echo "vLLM CI OTel: BuildKit trace conversion skipped" >&2
193+
return 0
194+
}
195+
timeout 5s python3 "${helper}" flush ||
196+
echo "vLLM CI OTel: BuildKit trace upload skipped" >&2
197+
)
198+
163199
#################################
164200
# Main Script #
165201
#################################
@@ -269,7 +305,18 @@ export PARENT_COMMIT
269305
print_bake_config
270306

271307
echo "--- :docker: Building ${TARGET}"
272-
docker --debug buildx bake -f "${VLLM_BAKE_FILE_PATH}" -f "${CI_HCL_PATH}" --progress plain "${TARGET}"
308+
BUILD_TMP_DIR="$(mktemp -d)"
309+
trap 'rm -rf -- "${BUILD_TMP_DIR}"' EXIT
310+
BUILD_METADATA_FILE="${BUILD_TMP_DIR}/build-metadata.json"
311+
BUILD_STATUS=0
312+
docker --debug buildx bake -f "${VLLM_BAKE_FILE_PATH}" -f "${CI_HCL_PATH}" \
313+
--progress plain --metadata-file "${BUILD_METADATA_FILE}" "${TARGET}" || BUILD_STATUS=$?
314+
315+
record_buildkit_trace "${BUILD_METADATA_FILE}" || true
316+
317+
if [[ "${BUILD_STATUS}" -ne 0 ]]; then
318+
exit "${BUILD_STATUS}"
319+
fi
273320

274321
echo "--- :white_check_mark: Build complete"
275322

0 commit comments

Comments
 (0)