[Quantization] Introduce layer re-quantization API for linear layers (MXFP8 linear showcase)

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

@@ -42,6 +42,25 @@ description: Build, debug, and interpret vLLM GPU kernel microbenchmarks for CUD
 - Keep metadata setup, plan construction, allocation, random input generation,
   and logging outside the timed region unless that overhead is the experiment.
 
+## Sanity-Check Reference Numbers
+
+Use these as rough reference points for large, well-shaped workloads, not as
+gold standards, guaranteed peaks, or hard limits. Hardware SKU, clocks, shape,
+precision conventions, and the FLOP/byte accounting can move the result. A
+large gap is a prompt to investigate, not proof that a kernel is poor.
+
+| Kernel regime | Hardware | Rough reference |
+| --- | --- | ---: |
+| Memory-bound, large batch | Blackwell | 6 TB/s |
+| BF16 GEMM | Blackwell | 2 PFLOP/s |
+| BF16 attention | B200 | 1.6 PFLOP/s |
+| FP8 GEMM | Blackwell | 4 PFLOP/s |
+| FP8 attention | B300 | 2.8 PFLOP/s |
+
+The BF16 attention reference is approximately the 1613 TFLOP/s result reported
+by the FlashAttention-4 paper. Compare kernels only with matching workload and
+throughput conventions.
+
 ## Multi-GPU Benchmarks
 
 - State whether the run is local or multi-node and report the GPU topology,

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

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

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

@@ -0,0 +1,4 @@
+interface:
+  display_name: "Triton Kernel Writing"
+  short_description: "Write robust, practical Triton kernels for vLLM."
+  default_prompt: "Use $triton-kernel-writing to implement or review this Triton kernel."

.buildkite/release-pipeline.yaml

@@ -959,6 +959,144 @@ steps:
           DOCKER_BUILDKIT: "1"
           DOCKERHUB_USERNAME: "vllmbot"
 
+  # =============================================================================
+  # ROCk Release Pipeline (x86_64 only)
+  # =============================================================================
+  #
+  # Same shape as the ROCm pipeline above (Jobs 1 and 6 only), but built from
+  # docker/Dockerfile.rock_base and docker/Dockerfile.rock, which install the
+  # ROCm SDK from TheRock wheels instead of the ROCm apt repos.
+  #
+  # Note: ROCm SDK version is determined by ROCM_SDK_VERSION in
+  # docker/Dockerfile.rock_base
+  #
+  # =============================================================================
+
+  - group: "Build ROCk Image"
+    key: "build-rock-image"
+    depends_on: ~
+    steps:
+      # ROCk Job 1: Build ROCk Base Image (with ECR caching)
+      - label: ":rocm: Build ROCk Base Image"
+        id: build-rock-base-image
+        depends_on: ~
+        agents:
+          queue: cpu_queue_release
+        timeout_in_minutes: 360
+        commands:
+          - |
+            set -euo pipefail
+
+            # The SDK version, torch wheels and component branches all live in
+            # the Dockerfile, so its hash is the cache key.
+            CACHE_KEY=$$(sha256sum docker/Dockerfile.rock_base | cut -c1-16)
+            ECR_CACHE_TAG="public.ecr.aws/q9t5s3a7/vllm-release-repo:$${CACHE_KEY}-rock-base"
+
+            echo "========================================"
+            echo "ROCk Base Build Configuration"
+            echo "========================================"
+            echo "  CACHE_KEY: $${CACHE_KEY}"
+            echo "  ECR_CACHE_TAG: $${ECR_CACHE_TAG}"
+            echo "========================================"
+
+            # Login to ECR
+            aws ecr-public get-login-password --region us-east-1 | \
+              docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
+
+            if docker manifest inspect "$${ECR_CACHE_TAG}" > /dev/null 2>&1; then
+              echo "ECR image cache HIT"
+            else
+              echo "CACHE MISS - Building from scratch..."
+
+              DOCKER_BUILDKIT=1 docker buildx build \
+                --file docker/Dockerfile.rock_base \
+                --tag "$${ECR_CACHE_TAG}" \
+                --build-arg USE_SCCACHE=1 \
+                --build-arg SCCACHE_BUCKET_NAME=vllm-build-sccache \
+                --build-arg SCCACHE_REGION_NAME=us-west-2 \
+                --build-arg SCCACHE_S3_NO_CREDENTIALS=0 \
+                --push \
+                .
+            fi
+
+            # Save ECR tag for downstream jobs
+            buildkite-agent meta-data set "rock-base-image-tag" "$${ECR_CACHE_TAG}"
+        env:
+          DOCKER_BUILDKIT: "1"
+
+      # ROCk Job 2: Build ROCk Release Docker Image
+      - label: ":docker: Build release image - x86_64 - ROCk"
+        id: build-rock-release-image
+        depends_on:
+          - step: block-build-release-images
+            allow_failure: true
+          - step: build-rock-base-image
+            allow_failure: false
+        agents:
+          queue: cpu_queue_release
+        timeout_in_minutes: 180
+        commands:
+          - |
+            set -euo pipefail
+
+            # Login to ECR
+            aws ecr-public get-login-password --region us-east-1 | \
+              docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
+
+            ECR_IMAGE_TAG="$$(buildkite-agent meta-data get rock-base-image-tag 2>/dev/null || echo '')"
+            if [ -z "$${ECR_IMAGE_TAG}" ]; then
+              echo "ERROR: rock-base-image-tag metadata not found"
+              echo "This should have been set by the build-rock-base-image job"
+              exit 1
+            fi
+
+            docker pull "$${ECR_IMAGE_TAG}"
+
+            # Pass the base image ECR tag to downstream steps (nightly publish)
+            buildkite-agent meta-data set "rock-base-ecr-tag" "$${ECR_IMAGE_TAG}"
+
+            echo "========================================"
+            echo "Building vLLM ROCk release image with:"
+            echo "  BASE_IMAGE: $${ECR_IMAGE_TAG}"
+            echo "  BUILDKITE_COMMIT: $${BUILDKITE_COMMIT}"
+            echo "========================================"
+
+            DOCKER_BUILDKIT=1 docker build \
+              --build-arg max_jobs=16 \
+              --build-arg BASE_IMAGE="$${ECR_IMAGE_TAG}" \
+              --build-arg USE_SCCACHE=1 \
+              --build-arg SCCACHE_BUCKET_NAME=vllm-build-sccache \
+              --build-arg SCCACHE_REGION_NAME=us-west-2 \
+              --build-arg SCCACHE_S3_NO_CREDENTIALS=0 \
+              --build-arg INSTALL_LMCACHE=true \
+              --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$${BUILDKITE_COMMIT}-rock \
+              --target vllm-openai \
+              --progress plain \
+              -f docker/Dockerfile.rock .
+
+            docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$${BUILDKITE_COMMIT}-rock
+        env:
+          DOCKER_BUILDKIT: "1"
+
+      - label: "Publish nightly ROCk image to DockerHub"
+        depends_on:
+          - build-rock-release-image
+        if: build.env("NIGHTLY") == "1"
+        agents:
+          queue: small_cpu_queue_release
+        commands:
+          - "bash .buildkite/scripts/push-nightly-builds-rocm.sh rock rocm100"
+          # Clean up old nightly builds (keep only last 14)
+          - "bash .buildkite/scripts/cleanup-nightly-builds.sh nightly-rocm100- vllm/vllm-openai-rocm"
+          - "bash .buildkite/scripts/cleanup-nightly-builds.sh base-nightly-rocm100- vllm/vllm-openai-rocm"
+        plugins:
+          - docker-login#v3.0.0:
+              username: vllmbot
+              password-env: DOCKERHUB_TOKEN
+        env:
+          DOCKER_BUILDKIT: "1"
+          DOCKERHUB_USERNAME: "vllmbot"
+
   # =============================================================================
   # Publish to DockerHub and PyPI (at the end so all builds complete first)
   # =============================================================================

.buildkite/scripts/ci-bake-rocm.sh

@@ -17,7 +17,7 @@ DEFAULT_REPO_SLUG="vllm-project/vllm"
 DEFAULT_CI_HCL_SOURCE="docker/ci-rocm.hcl"
 DEFAULT_CI_BASE_CONTENT_FILES=".dockerignore requirements/common.txt requirements/rocm.txt requirements/test/rocm.txt tools/install_torchcodec_rocm.sh rust-toolchain.toml tests/vllm_test_utils"
 DEFAULT_CI_BASE_DOCKERFILE="docker/Dockerfile.rocm"
-DEFAULT_CI_BASE_DOCKERFILE_STAGES="base rust_toolchain_input_0 rust-toolchain-input rust-toolchain build_nixl lmcache_source build_lmcache build_rocshmem build_deepep mooncake_source build_mooncake package_mooncake export_mooncake mori_base ci_base"
+DEFAULT_CI_BASE_DOCKERFILE_STAGES="base rust_toolchain_input_0 rust-toolchain-input rust-toolchain build_nixl lmcache_source build_lmcache build_rocshmem build_deepep mori_base ci_base"
 DEFAULT_CI_BASE_METADATA_VERSION="3"
 # ROCm CI forces REMOTE_VLLM=0, so content identity covers only the selected
 # local-source stages rather than unreachable remote-fetch alternatives.
@@ -1684,12 +1684,9 @@ ci_base_metadata_pairs() {
     metadata_pair "vllm.rocm.deepep_commit" "${DEEPEP_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_BRANCH")}"
     metadata_pair "vllm.rocm.deepep_nic" "$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_NIC")"
     metadata_pair "vllm.rocm.deepep_rocm_arch" "$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_ROCM_ARCH")"
-    metadata_pair "vllm.rocm.mooncake_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "MOONCAKE_REPO")"
-    metadata_pair "vllm.rocm.mooncake_commit" "${MOONCAKE_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "MOONCAKE_BRANCH")}"
     metadata_pair "vllm.rocm.nixl_cache_key" "${NIXL_CACHE_KEY:-}"
     metadata_pair "vllm.rocm.rocshmem_cache_key" "${ROCSHMEM_CACHE_KEY:-}"
     metadata_pair "vllm.rocm.deepep_cache_key" "${DEEPEP_CACHE_KEY:-}"
-    metadata_pair "vllm.rocm.mooncake_cache_key" "${MOONCAKE_CACHE_KEY:-}"
 }
 
 write_ci_base_metadata_annotations() {
@@ -2271,7 +2268,7 @@ extract_dependency_pins() {
         return 0
     fi
 
-    for var in NIXL_BRANCH UCX_BRANCH ROCSHMEM_BRANCH DEEPEP_BRANCH MOONCAKE_BRANCH; do
+    for var in NIXL_BRANCH UCX_BRANCH ROCSHMEM_BRANCH DEEPEP_BRANCH; do
         if [[ -n "${!var:-}" ]]; then
             echo "Using provided ${var}: ${!var}"
             continue
@@ -2295,19 +2292,16 @@ compute_dependency_cache_keys() {
     local ucx_branch=""
     local rocshmem_branch=""
     local deepep_branch=""
-    local mooncake_branch=""
     local nixl_material=""
     local rocshmem_material=""
     local deepep_material=""
-    local mooncake_material=""
 
     bake_dir=$(dirname "${VLLM_BAKE_FILE}")
     dockerfile_rocm="${bake_dir}/Dockerfile.rocm"
     nixl_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "NIXL_BRANCH")
     ucx_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "UCX_BRANCH")
     rocshmem_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "ROCSHMEM_BRANCH")
     deepep_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "DEEPEP_BRANCH")
-    mooncake_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "MOONCAKE_BRANCH")
 
     if [[ -n "${nixl_branch}" && -n "${ucx_branch}" ]]; then
         nixl_material=$(compose_stage_cache_material "${dockerfile_rocm}" "base build_nixl")
@@ -2341,19 +2335,6 @@ compute_dependency_cache_keys() {
         export DEEPEP_CACHE_KEY
         echo "DeepEP dependency cache key: ${DEEPEP_CACHE_KEY}"
     fi
-
-    if [[ -n "${mooncake_branch}" ]]; then
-        mooncake_material=$(compose_stage_cache_material \
-            "${dockerfile_rocm}" \
-            "base mooncake_source build_mooncake package_mooncake export_mooncake")
-        MOONCAKE_CACHE_KEY=$(
-            compose_dependency_cache_key \
-                "${mooncake_branch}" \
-                "${mooncake_material}"
-        )
-        export MOONCAKE_CACHE_KEY
-        echo "Mooncake dependency cache key: ${MOONCAKE_CACHE_KEY}"
-    fi
 }
 
 compose_stage_cache_material() {
@@ -2402,13 +2383,6 @@ dependency_cache_ref_for_target() {
                 printf '%s\n' "${cache_repo}:deepep-rocm-${DEEPEP_BRANCH}-rocshmem-${ROCSHMEM_BRANCH:-}"
             fi
             ;;
-        mooncake-rocm-ci)
-            if [[ -n "${MOONCAKE_CACHE_KEY:-}" ]]; then
-                printf '%s\n' "${cache_repo}:mooncake-rocm-${MOONCAKE_CACHE_KEY}"
-            elif [[ -n "${MOONCAKE_BRANCH:-}" ]]; then
-                printf '%s\n' "${cache_repo}:mooncake-rocm-${MOONCAKE_BRANCH}"
-            fi
-            ;;
     esac
 }
 
@@ -2426,14 +2400,13 @@ resolve_ci_base_dependency_targets() {
     local nixl_ref=""
     local rocshmem_ref=""
     local deepep_ref=""
-    local mooncake_ref=""
 
     [[ "${TARGET}" == "ci-base-rocm-ci-with-deps" ]] || return 0
 
     case "${mode}" in
         always)
             echo "ROCM_DEP_CACHE_EXPORT_MODE=always; exporting all dependency caches serially"
-            for target in nixl-rocm-ci rocshmem-rocm-ci deepep-rocm-ci mooncake-rocm-ci; do
+            for target in nixl-rocm-ci rocshmem-rocm-ci deepep-rocm-ci; do
                 if [[ -n "$(dependency_cache_ref_for_target "${target}")" ]]; then
                     add_dependency_cache_target "${target}"
                 fi
@@ -2483,16 +2456,6 @@ resolve_ci_base_dependency_targets() {
         fi
     fi
 
-    if [[ "${mode}" != "always" && -n "${MOONCAKE_CACHE_KEY:-}" ]]; then
-        mooncake_ref=$(dependency_cache_ref_for_target "mooncake-rocm-ci")
-        if dependency_cache_ref_exists "${mooncake_ref}"; then
-            echo "Mooncake dependency cache exists: ${mooncake_ref}"
-        else
-            echo "Mooncake dependency cache missing; will seed: ${mooncake_ref}"
-            add_dependency_cache_target "mooncake-rocm-ci"
-        fi
-    fi
-
     # DeepEP inherits from ROCShmem. If ROCShmem is being seeded, seed DeepEP too
     # so the pair stays consistent for future ci_base rebuilds.
     if printf '%s\n' "${DEPENDENCY_CACHE_TARGETS[@]}" | grep -qx "rocshmem-rocm-ci" \

.buildkite/scripts/cleanup-nightly-builds.sh

@@ -4,6 +4,7 @@ set -ex
 
 # Clean up old nightly builds from DockerHub, keeping only the last 14 builds
 # This script uses DockerHub API to list and delete old tags with specified prefix
+# Tags starting with "nightly-dev" are always excluded from deletion
 # Usage: cleanup-nightly-builds.sh [TAG_PREFIX] [REPO]
 # Example: cleanup-nightly-builds.sh "nightly-"
 # Example: cleanup-nightly-builds.sh "cu130-nightly-"
@@ -55,7 +56,8 @@ get_all_tags() {
         set -x
 
         # Get both last_updated timestamp and tag name, separated by |
-        local tags=$(echo "$response" | jq -r --arg prefix "$TAG_PREFIX" '.results[] | select(.name | startswith($prefix)) | "\(.last_updated)|\(.name)"')
+        # Exclude nightly-dev tags from cleanup
+        local tags=$(echo "$response" | jq -r --arg prefix "$TAG_PREFIX" '.results[] | select(.name | startswith($prefix)) | select(.name | startswith("nightly-dev") | not) | "\(.last_updated)|\(.name)"')
 
         if [ -z "$tags" ]; then
             break