diff --git a/.github/scripts/build_pr_wheel.sh b/.github/scripts/build_pr_wheel.sh
new file mode 100755
index 0000000000..3fbf123b55
--- /dev/null
+++ b/.github/scripts/build_pr_wheel.sh
@@ -0,0 +1,84 @@
+#!/usr/bin/env bash
+# SPDX-License-Identifier: Apache-2.0
+
+set -euo pipefail
+
+: "${PR_WHEEL_NUMBER:?PR_WHEEL_NUMBER is required}"
+: "${PLATFORM_BACKEND:?PLATFORM_BACKEND is required}"
+: "${PLATFORM_ARCHITECTURE:?PLATFORM_ARCHITECTURE is required}"
+: "${PLATFORM_VERSION_LOCAL:?PLATFORM_VERSION_LOCAL is required}"
+: "${B2_S3_BUCKET:?B2_S3_BUCKET is required}"
+: "${B2_PUBLIC_BASE_URL:?B2_PUBLIC_BASE_URL is required}"
+: "${RUNNER_TEMP:?RUNNER_TEMP is required}"
+
+output_dir="${1:?usage: build_pr_wheel.sh OUTPUT_DIR INDEX_DIR}"
+index_root="${2:?missing index output directory}"
+rclone="${RUNNER_TEMP}/rclone-bin/rclone"
+remote_root="b2:${B2_S3_BUCKET}/pull-requests/${PR_WHEEL_NUMBER}/${PLATFORM_BACKEND}/${PLATFORM_ARCHITECTURE}"
+public_root="${B2_PUBLIC_BASE_URL%/}/pull-requests/${PR_WHEEL_NUMBER}/${PLATFORM_BACKEND}/${PLATFORM_ARCHITECTURE}"
+target_commit="${TARGET_COMMIT:-${GITHUB_SHA:?GITHUB_SHA is required}}"
+
+if [[ ! -x "$rclone" ]]; then
+  echo "::error::rclone is not installed at ${rclone}"
+  exit 1
+fi
+
+tag="$(git describe --tags --abbrev=0 --match 'v[0-9]*')"
+release="${tag#v}"
+IFS=. read -r major minor patch <<<"$release"
+timestamp="$(date -u +%Y%m%d%H%M%S)"
+short_commit="$(git rev-parse --short=9 "$target_commit")"
+version="${major}.${minor}.$((patch + 1)).dev${timestamp}+${PLATFORM_VERSION_LOCAL}.pr${PR_WHEEL_NUMBER}.g${short_commit}"
+
+mkdir -p "$output_dir"
+find "$output_dir" -maxdepth 1 -type f -name '*.whl' -delete
+if [[ "$PLATFORM_BACKEND" == "cuda" ]]; then
+  APHRODITE_VERSION_OVERRIDE="$version" ./docker/export_wheels.sh
+  wheel="$(find wheels/main -maxdepth 1 -type f -name '*.whl' -print -quit)"
+else
+  : "${PLATFORM_BUILD_SCRIPT:?PLATFORM_BUILD_SCRIPT is required for non-CUDA wheels}"
+  APHRODITE_VERSION_OVERRIDE="$version" "$PLATFORM_BUILD_SCRIPT" "$output_dir"
+  wheel="$(find "$output_dir" -maxdepth 1 -type f -name '*.whl' -print -quit)"
+fi
+if [[ -z "$wheel" ]]; then
+  echo "::error::The PR build did not export a ${PLATFORM_BACKEND} wheel"
+  exit 1
+fi
+
+wheel_name="$(basename "$wheel")"
+if command -v sha256sum >/dev/null 2>&1; then
+  digest="$(sha256sum "$wheel" | cut -d ' ' -f 1)"
+else
+  digest="$(shasum -a 256 "$wheel" | cut -d ' ' -f 1)"
+fi
+
+"$rclone" purge "$remote_root" 2>/dev/null || true
+"$rclone" copyto "$wheel" "${remote_root}/wheels/${wheel_name}" \
+  --s3-upload-cutoff 16M \
+  --s3-chunk-size 16M \
+  --s3-upload-concurrency 8 \
+  --retries 5 \
+  --low-level-retries 10 \
+  --retries-sleep 5s
+
+encoded_name="${wheel_name//+/%2B}"
+entries="${RUNNER_TEMP}/sonar-pr-wheel.tsv"
+printf '%s\t%s\t%s\n' \
+  "$wheel_name" "${public_root}/wheels/${encoded_name}" "$digest" >"$entries"
+
+index_dir="${index_root}/whl/pr/${PR_WHEEL_NUMBER}/${PLATFORM_BACKEND}/${PLATFORM_ARCHITECTURE}"
+package_dir="${index_dir}/simple/aphrodite-engine"
+mkdir -p "$package_dir"
+python3 .github/scripts/generate_nightly_index.py \
+  --entry-file "$entries" \
+  --commit "$target_commit" \
+  --title "Sonar PR #${PR_WHEEL_NUMBER} ${PLATFORM_BACKEND} wheel" \
+  --description "Temporary ${PLATFORM_ARCHITECTURE} test wheel for PR #${PR_WHEEL_NUMBER}." \
+  --install-command \
+  "uv pip install aphrodite-engine --extra-index-url ${public_root}/simple --index-strategy first-index" \
+  --output "${index_dir}/index.html"
+cp "${index_dir}/index.html" "${package_dir}/index.html"
+"$rclone" copyto "${index_dir}/index.html" "${remote_root}/index.html"
+"$rclone" copyto "${package_dir}/index.html" "${remote_root}/simple/aphrodite-engine/index.html"
+
+echo "PR wheel index: ${public_root}/simple/aphrodite-engine/"
diff --git a/.github/workflows/platform-wheel.yml b/.github/workflows/platform-wheel.yml
index bf98863d8c..76b7dbc7f0 100644
--- a/.github/workflows/platform-wheel.yml
+++ b/.github/workflows/platform-wheel.yml
@@ -5,20 +5,149 @@ on:
     branches: [main]
   schedule:
     - cron: "30 3 * * *"
+  issue_comment:
+    types: [created]
+  pull_request_target:
+    types: [closed]
   workflow_dispatch:
+    inputs:
+      pr_number:
+        description: "Open PR number to build temporary wheels for; leave empty for the normal nightly reconciliation"
+        required: false
+        type: string
+      platform:
+        description: "Platform to build (manual dispatch only)"
+        required: false
+        default: all
+        type: choice
+        options:
+          - all
+          - cuda
+          - rocm
+          - cpu
+          - metal
 
 permissions:
   contents: read
+  issues: read
+  pull-requests: read
   pages: write
   id-token: write
 
 concurrency:
-  group: platform-wheels-${{ github.ref }}
+  group: platform-wheels-${{ github.event.issue.number || inputs.pr_number || github.ref }}
   cancel-in-progress: false
 
 jobs:
+  authorize:
+    name: Authorize wheel build
+    runs-on: ubuntu-24.04
+    permissions:
+      contents: read
+      issues: read
+      pull-requests: read
+    outputs:
+      build: ${{ steps.authorize.outputs.build }}
+      pr_number: ${{ steps.authorize.outputs.pr_number }}
+      target_sha: ${{ steps.authorize.outputs.target_sha }}
+      temporary: ${{ steps.authorize.outputs.temporary }}
+    steps:
+      - name: Validate trigger and resolve revision
+        id: authorize
+        uses: actions/github-script@v7
+        env:
+          DISPATCH_PR: ${{ inputs.pr_number || '' }}
+        with:
+          script: |
+            const event = context.eventName;
+            core.setOutput("build", "false");
+            core.setOutput("temporary", "false");
+            core.setOutput("pr_number", "");
+            core.setOutput("target_sha", context.sha);
+
+            if (event === "pull_request_target") {
+              return;
+            }
+
+            let prNumber = "";
+            if (event === "issue_comment") {
+              if (!context.payload.issue?.pull_request ||
+                  context.payload.comment?.body?.trim() !== "/build-wheels") {
+                return;
+              }
+              prNumber = String(context.issue.number);
+            } else if (event === "workflow_dispatch") {
+              prNumber = process.env.DISPATCH_PR.trim();
+            }
+
+            if (!prNumber) {
+              core.setOutput("build", "true");
+              return;
+            }
+            if (!/^\d+$/.test(prNumber)) {
+              core.setFailed(`Invalid PR number: ${prNumber}`);
+              return;
+            }
+
+            const actor = context.actor;
+            const {data: permission} = await github.rest.repos.getCollaboratorPermissionLevel({
+              owner: context.repo.owner,
+              repo: context.repo.repo,
+              username: actor,
+            });
+            if (!["admin", "maintain", "write"].includes(permission.permission)) {
+              core.setFailed(`${actor} does not have maintainer write access`);
+              return;
+            }
+
+            const {data: pr} = await github.rest.pulls.get({
+              owner: context.repo.owner,
+              repo: context.repo.repo,
+              pull_number: Number(prNumber),
+            });
+            if (pr.state !== "open") {
+              core.setFailed(`PR #${prNumber} is not open`);
+              return;
+            }
+            if (pr.head.repo.full_name !== `${context.repo.owner}/${context.repo.repo}`) {
+              core.setFailed(
+                "PR wheel builds require a branch in this repository; fork code is not executed on self-hosted runners"
+              );
+              return;
+            }
+
+            core.setOutput("build", "true");
+            core.setOutput("temporary", "true");
+            core.setOutput("pr_number", prNumber);
+            core.setOutput("target_sha", pr.head.sha);
+
+  cleanup-pr-wheels:
+    name: Remove closed PR wheels
+    if: ${{ github.event_name == 'pull_request_target' }}
+    runs-on: ubuntu-24.04
+    env:
+      RCLONE_CONFIG_B2_TYPE: s3
+      RCLONE_CONFIG_B2_PROVIDER: Other
+      RCLONE_CONFIG_B2_ACCESS_KEY_ID: ${{ secrets.B2_S3_KEY_ID }}
+      RCLONE_CONFIG_B2_SECRET_ACCESS_KEY: ${{ secrets.B2_S3_APPLICATION_KEY }}
+      RCLONE_CONFIG_B2_REGION: ${{ vars.B2_S3_REGION }}
+      RCLONE_CONFIG_B2_ENDPOINT: ${{ vars.B2_S3_ENDPOINT }}
+      RCLONE_CONFIG_B2_NO_CHECK_BUCKET: "true"
+      B2_S3_BUCKET: ${{ vars.B2_S3_BUCKET }}
+    steps:
+      - name: Check out trusted cleanup script
+        uses: actions/checkout@v4
+      - name: Install rclone
+        run: ./.github/scripts/install-rclone.sh "$RUNNER_TEMP/rclone-bin"
+      - name: Remove temporary wheel prefix
+        run: |
+          "$RUNNER_TEMP/rclone-bin/rclone" purge \
+            "b2:${B2_S3_BUCKET}/pull-requests/${{ github.event.pull_request.number }}"
+
   cuda-wheel:
     name: CUDA nightly (x86_64)
+    needs: authorize
+    if: ${{ needs.authorize.outputs.build == 'true' && (inputs.platform == '' || inputs.platform == 'all' || inputs.platform == 'cuda') }}
     runs-on: [self-hosted, linux, x64, aphrodite-wheel-builder]
     timeout-minutes: 720
     env:
@@ -32,11 +161,13 @@ jobs:
       RCLONE_CONFIG_B2_NO_CHECK_BUCKET: "true"
       B2_S3_BUCKET: ${{ vars.B2_S3_BUCKET }}
       B2_PUBLIC_BASE_URL: ${{ vars.B2_PUBLIC_BASE_URL }}
+      PR_WHEEL_NUMBER: ${{ needs.authorize.outputs.pr_number }}
+      TARGET_COMMIT: ${{ needs.authorize.outputs.target_sha }}
     steps:
       - name: Check out target revision
         uses: actions/checkout@v4
         with:
-          ref: ${{ github.sha }}
+          ref: ${{ needs.authorize.outputs.target_sha }}
           fetch-depth: 0
 
       - name: Install rclone
@@ -68,6 +199,7 @@ jobs:
             "$public_test_url")" = "nightly storage preflight"
 
       - name: Build and publish missing CUDA wheels
+        if: ${{ needs.authorize.outputs.temporary != 'true' }}
         env:
           PUSH_BEFORE: ${{ github.event.before }}
           MAX_JOBS: "64"
@@ -76,6 +208,18 @@ jobs:
           TORCH_CUDA_ARCH_LIST: "8.0 8.6 8.9 9.0 10.0 12.0+PTX"
         run: ./.github/scripts/build_nightly_commits.sh page-index
 
+      - name: Build and publish PR CUDA wheel
+        if: ${{ needs.authorize.outputs.temporary == 'true' }}
+        env:
+          MAX_JOBS: "64"
+          NVCC_THREADS: "4"
+          CUDA_VERSION: "13.0.2"
+          TORCH_CUDA_ARCH_LIST: "8.0 8.6 8.9 9.0 10.0 12.0+PTX"
+          PLATFORM_BACKEND: cuda
+          PLATFORM_ARCHITECTURE: x86_64
+          PLATFORM_VERSION_LOCAL: cu130
+        run: ./.github/scripts/build_pr_wheel.sh wheels/main page-index
+
       - name: Upload CUDA index
         uses: actions/upload-artifact@v4
         with:
@@ -95,6 +239,8 @@ jobs:
 
   cpu-wheel:
     name: CPU nightly (${{ matrix.architecture }})
+    needs: authorize
+    if: ${{ needs.authorize.outputs.build == 'true' && (inputs.platform == '' || inputs.platform == 'all' || inputs.platform == 'cpu') }}
     runs-on: ${{ matrix.runner }}
     timeout-minutes: 720
     env:
@@ -107,6 +253,8 @@ jobs:
       RCLONE_CONFIG_B2_NO_CHECK_BUCKET: "true"
       B2_S3_BUCKET: ${{ vars.B2_S3_BUCKET }}
       B2_PUBLIC_BASE_URL: ${{ vars.B2_PUBLIC_BASE_URL }}
+      PR_WHEEL_NUMBER: ${{ needs.authorize.outputs.pr_number }}
+      TARGET_COMMIT: ${{ needs.authorize.outputs.target_sha }}
     strategy:
       fail-fast: false
       matrix:
@@ -121,7 +269,7 @@ jobs:
       - name: Check out target revision
         uses: actions/checkout@v4
         with:
-          ref: ${{ github.sha }}
+          ref: ${{ needs.authorize.outputs.target_sha }}
           fetch-depth: 0
 
       - name: Configure Docker Buildx
@@ -149,6 +297,7 @@ jobs:
         run: ./.github/scripts/install-rclone.sh "$RUNNER_TEMP/rclone-bin"
 
       - name: Build and publish missing CPU wheels
+        if: ${{ needs.authorize.outputs.temporary != 'true' }}
         env:
           PUSH_BEFORE: ${{ github.event.before }}
           TARGET_PLATFORM: ${{ matrix.platform }}
@@ -162,6 +311,20 @@ jobs:
           .github/scripts/reconcile_platform_wheels.sh \
             "dist/cpu-${{ matrix.architecture }}" page-index
 
+      - name: Build and publish PR CPU wheel
+        if: ${{ needs.authorize.outputs.temporary == 'true' }}
+        env:
+          TARGET_PLATFORM: ${{ matrix.platform }}
+          PYTHON_VERSION: "3.13"
+          CACHE_SCOPE: cpu-${{ matrix.architecture }}
+          PLATFORM_BACKEND: cpu
+          PLATFORM_ARCHITECTURE: ${{ matrix.architecture }}
+          PLATFORM_VERSION_LOCAL: cpu
+          PLATFORM_BUILD_SCRIPT: .github/scripts/build_cpu_wheel.sh
+        run: |
+          .github/scripts/build_pr_wheel.sh \
+            "dist/cpu-${{ matrix.architecture }}" page-index
+
       - name: Upload CPU index
         uses: actions/upload-artifact@v4
         with:
@@ -172,6 +335,8 @@ jobs:
 
   metal-wheel:
     name: Metal nightly (Apple Silicon)
+    needs: authorize
+    if: ${{ needs.authorize.outputs.build == 'true' && (inputs.platform == '' || inputs.platform == 'all' || inputs.platform == 'metal') }}
     runs-on: macos-15
     timeout-minutes: 720
     env:
@@ -185,11 +350,13 @@ jobs:
       B2_S3_BUCKET: ${{ vars.B2_S3_BUCKET }}
       B2_PUBLIC_BASE_URL: ${{ vars.B2_PUBLIC_BASE_URL }}
       UV_CACHE_DIR: ${{ github.workspace }}/.cache/uv
+      PR_WHEEL_NUMBER: ${{ needs.authorize.outputs.pr_number }}
+      TARGET_COMMIT: ${{ needs.authorize.outputs.target_sha }}
     steps:
       - name: Check out target revision
         uses: actions/checkout@v4
         with:
-          ref: ${{ github.sha }}
+          ref: ${{ needs.authorize.outputs.target_sha }}
           fetch-depth: 0
 
       - name: Install build tools
@@ -235,6 +402,7 @@ jobs:
             "$RUNNER_TEMP/build_metal_wheel.sh"
 
       - name: Build and publish missing Metal wheels
+        if: ${{ needs.authorize.outputs.temporary != 'true' }}
         env:
           PUSH_BEFORE: ${{ github.event.before }}
           MACOSX_DEPLOYMENT_TARGET: "14.0"
@@ -247,6 +415,17 @@ jobs:
           .github/scripts/reconcile_platform_wheels.sh \
             dist/metal-aarch64 page-index
 
+      - name: Build and publish PR Metal wheel
+        if: ${{ needs.authorize.outputs.temporary == 'true' }}
+        env:
+          MACOSX_DEPLOYMENT_TARGET: "14.0"
+          PLATFORM_BACKEND: metal
+          PLATFORM_ARCHITECTURE: aarch64
+          PLATFORM_VERSION_LOCAL: metal
+          PLATFORM_BUILD_SCRIPT: ${{ runner.temp }}/build_metal_wheel.sh
+          METAL_BUILD_PYTHON: ${{ github.workspace }}/.metal-build-venv/bin/python
+        run: .github/scripts/build_pr_wheel.sh dist/metal-aarch64 page-index
+
       - name: Upload Metal index
         uses: actions/upload-artifact@v4
         with:
@@ -257,8 +436,8 @@ jobs:
 
   rocm-wheel:
     name: ROCm nightly (x86_64)
-    needs: cuda-wheel
-    if: ${{ always() }}
+    needs: [authorize, cuda-wheel]
+    if: ${{ always() && needs.authorize.outputs.build == 'true' && (inputs.platform == '' || inputs.platform == 'all' || inputs.platform == 'rocm') && (needs.cuda-wheel.result == 'success' || needs.cuda-wheel.result == 'skipped') }}
     runs-on: [self-hosted, linux, x64, aphrodite-wheel-builder]
     timeout-minutes: 720
     env:
@@ -274,11 +453,13 @@ jobs:
       RCLONE_CONFIG_B2_NO_CHECK_BUCKET: "true"
       B2_S3_BUCKET: ${{ vars.B2_S3_BUCKET }}
       B2_PUBLIC_BASE_URL: ${{ vars.B2_PUBLIC_BASE_URL }}
+      PR_WHEEL_NUMBER: ${{ needs.authorize.outputs.pr_number }}
+      TARGET_COMMIT: ${{ needs.authorize.outputs.target_sha }}
     steps:
       - name: Check out target revision
         uses: actions/checkout@v4
         with:
-          ref: ${{ github.sha }}
+          ref: ${{ needs.authorize.outputs.target_sha }}
           fetch-depth: 0
 
       - name: Start RAM-backed ROCm builder
@@ -297,6 +478,7 @@ jobs:
             "$RUNNER_TEMP/Dockerfile.rocm"
 
       - name: Build and publish missing ROCm wheels
+        if: ${{ needs.authorize.outputs.temporary != 'true' }}
         env:
           PUSH_BEFORE: ${{ github.event.before }}
           PLATFORM_BACKEND: rocm
@@ -308,6 +490,16 @@ jobs:
           .github/scripts/reconcile_platform_wheels.sh \
             dist/rocm-x86_64 page-index
 
+      - name: Build and publish PR ROCm wheel
+        if: ${{ needs.authorize.outputs.temporary == 'true' }}
+        env:
+          PLATFORM_BACKEND: rocm
+          PLATFORM_ARCHITECTURE: x86_64
+          PLATFORM_VERSION_LOCAL: rocm723
+          PLATFORM_BUILD_SCRIPT: ${{ runner.temp }}/build_rocm_wheel.sh
+          ROCM_BUILD_DOCKERFILE: ${{ runner.temp }}/Dockerfile.rocm
+        run: .github/scripts/build_pr_wheel.sh dist/rocm-x86_64 page-index
+
       - name: Upload ROCm index
         uses: actions/upload-artifact@v4
         with:
@@ -325,8 +517,8 @@ jobs:
 
   build-pages:
     name: Build documentation with nightly indexes
-    needs: [cuda-wheel, cpu-wheel, metal-wheel, rocm-wheel]
-    if: ${{ always() && !cancelled() }}
+    needs: [authorize, cuda-wheel, cpu-wheel, metal-wheel, rocm-wheel]
+    if: ${{ always() && !cancelled() && needs.authorize.outputs.temporary != 'true' }}
     runs-on: ubuntu-24.04
     steps:
       - name: Check out target revision
diff --git a/aphrodite/_custom_ops.py b/aphrodite/_custom_ops.py
index 10f2a63eb7..d59355ff7e 100644
--- a/aphrodite/_custom_ops.py
+++ b/aphrodite/_custom_ops.py
@@ -3589,8 +3589,9 @@ def sm89_sparse_mla_fwd(
 ) -> None:
     """Sparse MLA forward on sm89, gathering directly from an fp8_ds_mla pool by slot.
 
-    q is [T, h, 576] bf16; kv_cache_pool [S, 656] uint8 (fp8_ds_mla rows); indices
-    [T, topk] int32 physical slots, -1 padded; out [T, h, 512] bf16; lse [T, h] fp32.
+    q is [T, h, 576] bf16 with a [S, 656] uint8 V3.2 pool, or [T, h, 512]
+    bf16 with a [B, block_size, 584] uint8 V4 pool. Indices are [T, topk]
+    int32 physical slots, -1 padded; out is [T, h, 512] bf16 and lse is [T, h] fp32.
     Rows whose indices are all -1 produce zero output and -inf LSE.
 
     topk_lens is an optional [T] int32 of per-token valid counts. It requires index
diff --git a/aphrodite/model_executor/kernels/mhc/tilelang.py b/aphrodite/model_executor/kernels/mhc/tilelang.py
index ab8798c06c..1006c0a448 100644
--- a/aphrodite/model_executor/kernels/mhc/tilelang.py
+++ b/aphrodite/model_executor/kernels/mhc/tilelang.py
@@ -333,7 +333,13 @@ def mhc_pre_broadcast_tilelang(
     residual_flat = residual
     num_tokens = residual.shape[0]
 
-    n_splits = compute_num_split(64, hidden_size, cdiv(num_tokens, 64))
+    from aphrodite.utils.deep_gemm import is_deep_gemm_supported
+
+    use_deep_gemm = is_deep_gemm_supported()
+    if use_deep_gemm:
+        n_splits = compute_num_split(64, hidden_size, cdiv(num_tokens, 64))
+    else:
+        n_splits = 1
 
     residual_out = torch.empty(
         num_tokens,
@@ -359,15 +365,25 @@ def mhc_pre_broadcast_tilelang(
         device=residual.device,
     )
 
-    from aphrodite.utils.deep_gemm import tf32_hc_prenorm_gemm
+    if use_deep_gemm:
+        from aphrodite.utils.deep_gemm import tf32_hc_prenorm_gemm
 
-    tf32_hc_prenorm_gemm(
-        residual_flat,
-        fn_broadcast,
-        gemm_out_mul,
-        gemm_out_sqrsum,
-        n_splits,
-    )
+        tf32_hc_prenorm_gemm(
+            residual_flat,
+            fn_broadcast,
+            gemm_out_mul,
+            gemm_out_sqrsum,
+            n_splits,
+        )
+    else:
+        _tilelang_hc_prenorm_gemm(
+            residual_flat,
+            fn_broadcast,
+            gemm_out_mul,
+            gemm_out_sqrsum,
+            hidden_size,
+            1,
+        )
     mhc_pre_big_fuse_broadcast_with_norm_tilelang(
         gemm_out_mul,
         gemm_out_sqrsum,
diff --git a/aphrodite/model_executor/layers/quantization/utils/fp8_utils.py b/aphrodite/model_executor/layers/quantization/utils/fp8_utils.py
index 5c307e30bd..c633cf831b 100644
--- a/aphrodite/model_executor/layers/quantization/utils/fp8_utils.py
+++ b/aphrodite/model_executor/layers/quantization/utils/fp8_utils.py
@@ -841,14 +841,13 @@ def w8a8_triton_block_scaled_mm(
     assert len(block_size) == 2
     block_n, block_k = block_size[0], block_size[1]
 
-    # Triton cannot currently bind E8M0 scale tensors directly. On ROCm,
-    # DeepSeek-V4 checkpoints store block scales in exponent-only E8M0 format,
-    # so decode them to fp32 before launching the kernel.
-    if current_platform.is_rocm() or current_platform.is_xpu():
-        if As.dtype == torch.float8_e8m0fnu:
-            As = _upcast_e8m0_to_fp32(As).contiguous()
-        if Bs.dtype == torch.float8_e8m0fnu:
-            Bs = _upcast_e8m0_to_fp32(Bs).contiguous()
+    # Triton cannot currently bind E8M0 scale tensors directly. Checkpoints
+    # may store block scales in exponent-only E8M0 format on any platform, so
+    # decode them to fp32 before launching the kernel.
+    if As.dtype == torch.float8_e8m0fnu:
+        As = _upcast_e8m0_to_fp32(As).contiguous()
+    if Bs.dtype == torch.float8_e8m0fnu:
+        Bs = _upcast_e8m0_to_fp32(Bs).contiguous()
 
     assert A.shape[-1] == B.shape[-1]
     assert A.shape[:-1] == As.shape[:-1] and A.is_contiguous()
diff --git a/aphrodite/models/deepseek_v4/common/ops/cache_utils.py b/aphrodite/models/deepseek_v4/common/ops/cache_utils.py
index 8131519dda..619ec14d26 100644
--- a/aphrodite/models/deepseek_v4/common/ops/cache_utils.py
+++ b/aphrodite/models/deepseek_v4/common/ops/cache_utils.py
@@ -404,7 +404,7 @@ def dequantize_and_gather_k_cache(
     ``current_platform.is_fp8_fnuz()`` for ``swa_k_cache`` (C++ encoder
     writes FNUZ on gfx942 and OCP on gfx950).
     """
-    if has_cutedsl():
+    if has_cutedsl() and current_platform.has_device_capability(90):
         # lazily import, otherwise some tests fail due to CUDA driver init failure.
         from aphrodite.models.deepseek_v4.nvidia.ops.dequant_gather_k_cutedsl import (
             dequantize_and_gather_k_cache_cutedsl,
diff --git a/aphrodite/models/deepseek_v4/nvidia/flashmla.py b/aphrodite/models/deepseek_v4/nvidia/flashmla.py
index b71e85d80e..1662898085 100644
--- a/aphrodite/models/deepseek_v4/nvidia/flashmla.py
+++ b/aphrodite/models/deepseek_v4/nvidia/flashmla.py
@@ -169,6 +169,22 @@ def _forward_decode(
         swa_indices = swa_metadata.decode_swa_indices
         swa_lens = swa_metadata.decode_swa_lens
 
+        from aphrodite.v1.attention.backends.mla.sm89_mla_sparse import (
+            use_sm89_dsa,
+        )
+
+        if use_sm89_dsa():
+            self._forward_decode_sm89(
+                q,
+                kv_cache,
+                swa_indices,
+                swa_lens,
+                topk_indices,
+                topk_lens,
+                output,
+            )
+            return
+
         # We treat queries in the same seq as different queries
         # and later we only attend by generated indices.
         # q arrives pre-padded to self.padded_heads by the outer wrapper.
@@ -220,6 +236,75 @@ def _forward_decode(
             out=output.unsqueeze(1),
         )
 
+    def _forward_decode_sm89(
+        self,
+        q: torch.Tensor,
+        compressed_cache: torch.Tensor | None,
+        swa_indices: torch.Tensor,
+        swa_lens: torch.Tensor,
+        compressed_indices: torch.Tensor | None,
+        compressed_lens: torch.Tensor | None,
+        output: torch.Tensor,
+    ) -> None:
+        """Run V4 sparse decode through the native Ada kernel."""
+        from aphrodite import _custom_ops as ops
+
+        # Decode receives token slices of the runner's padded buffers.  The
+        # native kernel requires dense storage, so materialize only this small
+        # query view.  Do not materialize the KV caches here.
+        q = q.contiguous()
+
+        def run_cache(
+            cache: torch.Tensor,
+            indices: torch.Tensor,
+            lens: torch.Tensor | None,
+        ) -> tuple[torch.Tensor, torch.Tensor]:
+            indices = indices.view(q.shape[0], -1).to(torch.int32).contiguous()
+            if lens is not None:
+                lens = lens.view(-1).to(torch.int32).contiguous()
+            partial = torch.empty(
+                output.shape,
+                dtype=output.dtype,
+                device=output.device,
+            )
+            lse = torch.empty(
+                (q.shape[0], q.shape[1]),
+                dtype=torch.float32,
+                device=q.device,
+            )
+            ops.sm89_sparse_mla_fwd(
+                q,
+                cache,
+                indices,
+                partial,
+                lse,
+                self.scale,
+                lens,
+            )
+            return partial, lse
+
+        swa_output, swa_lse = run_cache(
+            self.swa_cache_layer.kv_cache,
+            swa_indices,
+            swa_lens,
+        )
+        if compressed_cache is None or compressed_indices is None:
+            output.copy_(swa_output)
+            return
+
+        compressed_output, compressed_lse = run_cache(
+            compressed_cache,
+            compressed_indices,
+            compressed_lens,
+        )
+        ops.merge_attn_states(
+            output,
+            swa_output,
+            swa_lse.transpose(0, 1).contiguous(),
+            compressed_output,
+            compressed_lse.transpose(0, 1).contiguous(),
+        )
+
     def _forward_prefill(
         self,
         q: torch.Tensor,
@@ -316,12 +401,36 @@ def _forward_prefill(
                 chunk_M,
                 chunk_N,
             )
-            flash_mla_sparse_fwd(
-                q=q[query_start:query_end],
-                kv=kv.view(-1, 1, q.shape[-1]),
-                indices=combined_indices.unsqueeze(1),
-                sm_scale=self.scale,
-                attn_sink=self.attn_sink,
-                topk_length=combined_lens,
-                out=output[query_start:query_end],
+            from aphrodite.v1.attention.backends.mla.sm89_mla_sparse import (
+                use_sm89_dsa,
             )
+
+            if use_sm89_dsa():
+                # FlashMLA sparse prefill requires SM90a+.  Its Triton
+                # counterpart is device-neutral and supports the same V4
+                # query/KV layout on Ada.
+                from aphrodite.v1.attention.ops.rocm_aiter_mla_sparse import (
+                    _rocm_sparse_attn_prefill_triton,
+                )
+
+                prefill_out = _rocm_sparse_attn_prefill_triton(
+                    q=q[query_start:query_end],
+                    kv=kv.view(-1, q.shape[-1]),
+                    indices=combined_indices,
+                    scale=self.scale,
+                    attn_sink=self.attn_sink,
+                    nope_head_dim=448,
+                    rope_head_dim=64,
+                    topk_length=combined_lens,
+                )
+                output[query_start:query_end].copy_(prefill_out)
+            else:
+                flash_mla_sparse_fwd(
+                    q=q[query_start:query_end],
+                    kv=kv.view(-1, 1, q.shape[-1]),
+                    indices=combined_indices.unsqueeze(1),
+                    sm_scale=self.scale,
+                    attn_sink=self.attn_sink,
+                    topk_length=combined_lens,
+                    out=output[query_start:query_end],
+                )
diff --git a/aphrodite/models/deepseek_v4/nvidia/ops/o_proj.py b/aphrodite/models/deepseek_v4/nvidia/ops/o_proj.py
index 196a767aa7..c89a5a7b58 100644
--- a/aphrodite/models/deepseek_v4/nvidia/ops/o_proj.py
+++ b/aphrodite/models/deepseek_v4/nvidia/ops/o_proj.py
@@ -7,7 +7,56 @@
     fused_inv_rope_fp8_quant,
 )
 from aphrodite.platforms import current_platform
-from aphrodite.utils.deep_gemm import fp8_einsum
+from aphrodite.utils.deep_gemm import fp8_einsum, is_deep_gemm_supported
+
+
+def _expand_block_scales(scale: torch.Tensor, rows: int, cols: int) -> torch.Tensor:
+    if scale.dtype == torch.float8_e8m0fnu:
+        from aphrodite.model_executor.layers.quantization.utils.fp8_utils import (
+            _upcast_e8m0_to_fp32,
+        )
+
+        scale = _upcast_e8m0_to_fp32(scale)
+    else:
+        scale = scale.float()
+    row_blocks, col_blocks = scale.shape[-2:]
+    row_block = (rows + row_blocks - 1) // row_blocks
+    col_block = (cols + col_blocks - 1) // col_blocks
+    scale = torch.repeat_interleave(scale, row_block, dim=-2)[..., :rows, :]
+    return torch.repeat_interleave(scale, col_block, dim=-1)[..., :, :cols]
+
+
+def _bf16_o_proj(
+    o: torch.Tensor,
+    positions: torch.Tensor,
+    cos_sin_cache: torch.Tensor,
+    wo_a: nn.Module,
+    *,
+    n_groups: int,
+    heads_per_group: int,
+    rope_dim: int,
+    o_lora_rank: int,
+) -> torch.Tensor:
+    """Portable inverse-RoPE and dequantized WO_A fallback."""
+    half_rope = rope_dim // 2
+    rope = o[..., -rope_dim:].float()
+    even = rope[..., 0::2]
+    odd = rope[..., 1::2]
+    cos = cos_sin_cache[positions, :half_rope].unsqueeze(1)
+    sin = cos_sin_cache[positions, half_rope:].unsqueeze(1)
+    inv_rope = torch.stack((even * cos + odd * sin, odd * cos - even * sin), dim=-1).flatten(-2)
+    o_ref = torch.cat((o[..., :-rope_dim].float(), inv_rope), dim=-1)
+    o_ref = o_ref.view(o.shape[0], n_groups, -1)
+
+    weight = getattr(wo_a, "_dsv4_wo_a_bf16", None)
+    if weight is None:
+        hidden_dim = heads_per_group * o.shape[-1]
+        weight = wo_a.weight.view(n_groups, o_lora_rank, hidden_dim).float()
+        weight_scale = wo_a.weight_scale if hasattr(wo_a, "weight_scale") else wo_a.weight_scale_inv
+        weight_scale = weight_scale.view(n_groups, -1, weight_scale.shape[-1])
+        weight = (weight * _expand_block_scales(weight_scale, o_lora_rank, hidden_dim)).to(torch.bfloat16)
+        wo_a._dsv4_wo_a_bf16 = weight
+    return torch.einsum("tgd,grd->tgr", o_ref.to(torch.bfloat16), weight)
 
 
 def compute_fp8_einsum_recipe() -> tuple[tuple[int, int, int], bool]:
@@ -45,6 +94,19 @@ def deep_gemm_fp8_o_proj(
     Shared by the FlashMLA and FlashInfer CUDA backends. ``einsum_recipe`` /
     ``tma_aligned_scales`` come from ``compute_fp8_einsum_recipe``.
     """
+    if not is_deep_gemm_supported():
+        z = _bf16_o_proj(
+            o,
+            positions,
+            cos_sin_cache,
+            wo_a,
+            n_groups=n_groups,
+            heads_per_group=heads_per_group,
+            rope_dim=rope_dim,
+            o_lora_rank=o_lora_rank,
+        )
+        return wo_b(z.flatten(1))
+
     o_fp8, o_scale = fused_inv_rope_fp8_quant(
         o,
         positions,
diff --git a/csrc/libtorch_stable/attention/sm89_dsa/sm89_fp8_paged_mqa_logits.cu b/csrc/libtorch_stable/attention/sm89_dsa/sm89_fp8_paged_mqa_logits.cu
index 678eb25294..975c2b3c12 100644
--- a/csrc/libtorch_stable/attention/sm89_dsa/sm89_fp8_paged_mqa_logits.cu
+++ b/csrc/libtorch_stable/attention/sm89_dsa/sm89_fp8_paged_mqa_logits.cu
@@ -72,7 +72,8 @@ paged_mqa_logits_kernel(const uint8_t* __restrict__ q,           // [B, NEXT_N,
                         const int32_t* __restrict__ block_table, // [B, max_pages]
                         const int32_t* __restrict__ sched,       // [(P+1), 2]
                         float* __restrict__ logits,              // [B*NEXT_N, max_model_len]
-                        int max_pages, int64_t max_model_len, int clean_logits) {
+                        int max_pages, int64_t max_model_len, int64_t pool_page_stride,
+                        int clean_logits) {
   static_assert(NUM_HEADS == 32 || NUM_HEADS == 64, "unsupported indexer head count");
   constexpr int HGROUPS = NUM_HEADS / 16;  // head groups of 16 (one m16 mma tile each)
   constexpr int NT = 8 * HGROUPS / 4;      // 8-key n-tiles per warp
@@ -131,7 +132,7 @@ paged_mqa_logits_kernel(const uint8_t* __restrict__ q,           // [B, NEXT_N,
 
     const int32_t* bt = block_table + (int64_t)req * max_pages;
     auto issue_page = [&](int p, int stage) {
-      const uint8_t* src = pool + (int64_t)bt[p] * PAGE_BYTES;
+      const uint8_t* src = pool + (int64_t)bt[p] * pool_page_stride;
       uint8_t* dst = smem + stage * PAGE_BYTES;
       for (int c = tid; c < PAGE_BYTES / 16; c += NEXT_N * 128) {
         int off = (c < PAGE_KEY_BYTES / 16) ? swizzle_chunk(c >> 3, c & 7)
@@ -294,8 +295,11 @@ void sm89_fp8_paged_mqa_logits(const torch::stable::Tensor& q,
                   "q must be [B, next_n, {32|64}, 128]");
   STD_TORCH_CHECK(weights.dim() == 2 && weights.size(1) == q.size(2),
                   "weights must be [B*next_n, num_heads]");
-  STD_TORCH_CHECK(q.is_contiguous() && pool.is_contiguous() && logits.is_contiguous(),
-                  "q, pool, and logits must be contiguous");
+  STD_TORCH_CHECK(q.is_contiguous() && logits.is_contiguous(),
+                  "q and logits must be contiguous");
+  STD_TORCH_CHECK(pool.dim() == 2 && pool.size(1) == sm89_dsa::PAGE_BYTES &&
+                      pool.stride(1) == 1,
+                  "pool must have dense 8448-byte pages");
   STD_TORCH_CHECK(weights.is_contiguous() && seq_lens.is_contiguous() &&
                       block_table.is_contiguous() && sched.is_contiguous(),
                   "weights, seq_lens, block_table, and sched must be contiguous");
@@ -316,7 +320,8 @@ void sm89_fp8_paged_mqa_logits(const torch::stable::Tensor& q,
         q.const_data_ptr<uint8_t>(), pool.const_data_ptr<uint8_t>(),
         weights.const_data_ptr<float>(), seq_lens.const_data_ptr<int32_t>(),
         block_table.const_data_ptr<int32_t>(), sched.const_data_ptr<int32_t>(),
-        logits.mutable_data_ptr<float>(), max_pages, max_model_len, (int)clean_logits);
+        logits.mutable_data_ptr<float>(), max_pages, max_model_len, pool.stride(0),
+        (int)clean_logits);
   };
   STD_TORCH_CHECK(next_n == 1 || next_n == 2, "next_n must be 1 or 2");
   if (next_n == 1) {
diff --git a/csrc/libtorch_stable/attention/sm89_dsa/sm89_sparse_mla_fwd.cu b/csrc/libtorch_stable/attention/sm89_dsa/sm89_sparse_mla_fwd.cu
index c0d6dd9aaf..02c4a97213 100644
--- a/csrc/libtorch_stable/attention/sm89_dsa/sm89_sparse_mla_fwd.cu
+++ b/csrc/libtorch_stable/attention/sm89_dsa/sm89_sparse_mla_fwd.cu
@@ -54,22 +54,27 @@
 
 namespace sm89_dsa {
 
-constexpr int D = 576;   // 512 nope + 64 rope
-constexpr int DN = 512;  // value dims (MLA absorb, V == K nope)
-constexpr int ROW_BYTES = 656;
-constexpr int ROW_CHUNKS = ROW_BYTES / 16;  // 41
 constexpr int BI = 32;                      // keys per pipeline iteration
 constexpr int STAGES = 2;
 constexpr int HT = 16;                    // head tile (one m16 mma tile)
-constexpr int SKV_STRIDE = D + 8;         // in halves; 1168B rows stay 16B-aligned, 4-word bank skew
-constexpr int SQ_STRIDE = D + 8;
 constexpr int SS_STRIDE = BI + 1;         // fp32; +1 pad -> per-row bank skew
 
-constexpr int SMEM_STAGING = STAGES * BI * ROW_BYTES;  // 41984
-constexpr int SMEM_SKV = BI * SKV_STRIDE * 2;          // 37376
-constexpr int SMEM_SQ = HT * SQ_STRIDE * 2;            // 18688
-constexpr int SMEM_SS = HT * SS_STRIDE * 4;            // 2112
-constexpr int SMEM_TOTAL = SMEM_STAGING + SMEM_SKV + SMEM_SQ + SMEM_SS;  // 100160 <= 99KB opt-in
+template <bool DSV4>
+struct Layout {
+  static constexpr int D = DSV4 ? 512 : 576;
+  static constexpr int NOPE = DSV4 ? 448 : 512;
+  static constexpr int VALUE = 512;
+  // V4 stages its 576-byte data row followed by a padded 16-byte scale row.
+  static constexpr int ROW_BYTES = DSV4 ? 592 : 656;
+  static constexpr int ROW_CHUNKS = ROW_BYTES / 16;
+  static constexpr int SKV_STRIDE = D + 8;
+  static constexpr int SQ_STRIDE = D + 8;
+  static constexpr int SMEM_STAGING = STAGES * BI * ROW_BYTES;
+  static constexpr int SMEM_SKV = BI * SKV_STRIDE * 2;
+  static constexpr int SMEM_SQ = HT * SQ_STRIDE * 2;
+  static constexpr int SMEM_SS = HT * SS_STRIDE * 4;
+  static constexpr int SMEM_TOTAL = SMEM_STAGING + SMEM_SKV + SMEM_SQ + SMEM_SS;
+};
 
 DEVINL void cp_async_16(void* smem_dst, const void* gmem_src) {
   uint32_t dst = static_cast<uint32_t>(__cvta_generic_to_shared(smem_dst));
@@ -106,7 +111,7 @@ DEVINL float2 fp8x2_to_float2(uint16_t v) {
 
 // HAS_LENS == false compiles the len clamp away, leaving the full-topk loop
 // untouched (the extra kernel param is never read).
-template <bool HAS_LENS>
+template <bool HAS_LENS, bool DSV4>
 __global__ void __launch_bounds__(128, 1)
 sparse_mla_fwd_kernel(const __nv_bfloat16* __restrict__ q,   // [T, h, 576]
                       const uint8_t* __restrict__ pool,      // [S, 656]
@@ -114,12 +119,14 @@ sparse_mla_fwd_kernel(const __nv_bfloat16* __restrict__ q,   // [T, h, 576]
                       __nv_bfloat16* __restrict__ out,       // [T, h, 512]
                       float* __restrict__ lse,                // [T, h]
                       const int32_t* __restrict__ topk_lens,  // [T]; read iff HAS_LENS
-                      int h, int topk, float sm_scale) {
+                      int h, int topk, float sm_scale, int cache_block_size,
+                      int64_t cache_block_stride) {
+  using L = Layout<DSV4>;
   extern __shared__ uint8_t smem[];
   uint8_t* staging = smem;
-  __nv_bfloat16* sKV = reinterpret_cast<__nv_bfloat16*>(smem + SMEM_STAGING);
-  __nv_bfloat16* sQ = reinterpret_cast<__nv_bfloat16*>(smem + SMEM_STAGING + SMEM_SKV);
-  float* sS = reinterpret_cast<float*>(smem + SMEM_STAGING + SMEM_SKV + SMEM_SQ);
+  __nv_bfloat16* sKV = reinterpret_cast<__nv_bfloat16*>(smem + L::SMEM_STAGING);
+  __nv_bfloat16* sQ = reinterpret_cast<__nv_bfloat16*>(smem + L::SMEM_STAGING + L::SMEM_SKV);
+  float* sS = reinterpret_cast<float*>(smem + L::SMEM_STAGING + L::SMEM_SKV + L::SMEM_SQ);
 
   const int t = blockIdx.x;
   const int h_base = blockIdx.y * HT;
@@ -135,22 +142,38 @@ sparse_mla_fwd_kernel(const __nv_bfloat16* __restrict__ q,   // [T, h, 576]
   auto key_valid = [&](int k) { return k < topk && idx_row[k] >= 0; };
 
   // ---- q tile -> sQ (rows past h zeroed; pad cols never read)
-  for (int c = tid; c < HT * (D / 8); c += 128) {
-    const int r = c / (D / 8), cc = c % (D / 8);
+  for (int c = tid; c < HT * (L::D / 8); c += 128) {
+    const int r = c / (L::D / 8), cc = c % (L::D / 8);
     uint4 v = make_uint4(0u, 0u, 0u, 0u);
     if (h_base + r < h)
-      v = *reinterpret_cast<const uint4*>(q + ((int64_t)t * h + h_base + r) * D + cc * 8);
-    *reinterpret_cast<uint4*>(sQ + r * SQ_STRIDE + cc * 8) = v;
+      v = *reinterpret_cast<const uint4*>(q + ((int64_t)t * h + h_base + r) * L::D + cc * 8);
+    *reinterpret_cast<uint4*>(sQ + r * L::SQ_STRIDE + cc * 8) = v;
   }
 
   auto issue_iter = [&](int it, int stage) {
-    uint8_t* dst = staging + stage * (BI * ROW_BYTES);
+    uint8_t* dst = staging + stage * (BI * L::ROW_BYTES);
     const int kbase = it * BI;
-    for (int c = tid; c < BI * ROW_CHUNKS; c += 128) {
-      const int j = c / ROW_CHUNKS, cc = c % ROW_CHUNKS;
+    for (int c = tid; c < BI * L::ROW_CHUNKS; c += 128) {
+      const int j = c / L::ROW_CHUNKS, cc = c % L::ROW_CHUNKS;
       const int k = kbase + j;
       const int32_t slot = (k < topk) ? max(idx_row[k], 0) : 0;  // -1/tail -> slot 0 (masked)
-      cp_async_16(dst + j * ROW_BYTES + cc * 16, pool + (int64_t)slot * ROW_BYTES + cc * 16);
+      if constexpr (DSV4) {
+        const int block = slot / cache_block_size;
+        const int pos = slot % cache_block_size;
+        const uint8_t* block_base = pool + block * cache_block_stride;
+        if (cc < 36) {
+          cp_async_16(dst + j * L::ROW_BYTES + cc * 16,
+                      block_base + (int64_t)pos * 576 + cc * 16);
+        } else {
+          const uint8_t* scale_src =
+              block_base + (int64_t)cache_block_size * 576 + (int64_t)pos * 8;
+          *reinterpret_cast<uint64_t*>(dst + j * L::ROW_BYTES + 576) =
+              *reinterpret_cast<const uint64_t*>(scale_src);
+        }
+      } else {
+        cp_async_16(dst + j * L::ROW_BYTES + cc * 16,
+                    pool + (int64_t)slot * L::ROW_BYTES + cc * 16);
+      }
     }
     cp_async_commit();
   };
@@ -185,10 +208,35 @@ sparse_mla_fwd_kernel(const __nv_bfloat16* __restrict__ q,   // [T, h, 576]
 
     // ---- dequant staging[stage] -> sKV bf16 [32][584]; thread (row j=tid/4, tile p=tid%4)
     {
-      const uint8_t* srow = staging + stage * (BI * ROW_BYTES) + (tid >> 2) * ROW_BYTES;
+      const uint8_t* srow = staging + stage * (BI * L::ROW_BYTES) + (tid >> 2) * L::ROW_BYTES;
       const int p = tid & 3;
-      const float scale = reinterpret_cast<const float*>(srow + DN)[p];
-      __nv_bfloat16* drow = sKV + (tid >> 2) * SKV_STRIDE;
+      __nv_bfloat16* drow = sKV + (tid >> 2) * L::SKV_STRIDE;
+      if constexpr (DSV4) {
+#pragma unroll
+        for (int block = p; block < 7; block += 4) {
+          const int exponent = static_cast<int>(srow[576 + block]) - 127;
+          const float scale = ldexpf(1.0f, exponent);
+#pragma unroll
+          for (int d = 0; d < 64; d += 8) {
+            const uint2 raw = *reinterpret_cast<const uint2*>(srow + block * 64 + d);
+            const uint16_t* b2 = reinterpret_cast<const uint16_t*>(&raw);
+            uint32_t packed[4];
+#pragma unroll
+            for (int i = 0; i < 4; ++i) {
+              const float2 f = fp8x2_to_float2(b2[i]);
+              packed[i] = bf16x2(f.x * scale, f.y * scale);
+            }
+            *reinterpret_cast<uint4*>(drow + block * 64 + d) =
+                make_uint4(packed[0], packed[1], packed[2], packed[3]);
+          }
+        }
+        // The 64 rope bf16 values are part of the 576-byte data row.
+        const uint4* rs = reinterpret_cast<const uint4*>(srow + L::NOPE);
+        uint4* rd = reinterpret_cast<uint4*>(drow + L::NOPE);
+        rd[2 * p] = rs[2 * p];
+        rd[2 * p + 1] = rs[2 * p + 1];
+      } else {
+      const float scale = reinterpret_cast<const float*>(srow + L::NOPE)[p];
 #pragma unroll
       for (int d = 0; d < 128; d += 8) {
         const uint2 raw = *reinterpret_cast<const uint2*>(srow + p * 128 + d);
@@ -202,10 +250,11 @@ sparse_mla_fwd_kernel(const __nv_bfloat16* __restrict__ q,   // [T, h, 576]
         *reinterpret_cast<uint4*>(drow + p * 128 + d) = make_uint4(o[0], o[1], o[2], o[3]);
       }
       // rope passthrough (64 bf16 = 8 uint4); thread p copies 2
-      const uint4* rs = reinterpret_cast<const uint4*>(srow + DN + 16);
-      uint4* rd = reinterpret_cast<uint4*>(drow + DN);
+      const uint4* rs = reinterpret_cast<const uint4*>(srow + L::NOPE + 16);
+      uint4* rd = reinterpret_cast<uint4*>(drow + L::NOPE);
       rd[2 * p] = rs[2 * p];
       rd[2 * p + 1] = rs[2 * p + 1];
+      }
     }
     __syncthreads();  // sKV ready; staging[stage] consumed -> safe to refill
 
@@ -215,11 +264,11 @@ sparse_mla_fwd_kernel(const __nv_bfloat16* __restrict__ q,   // [T, h, 576]
     // ---- QK mma; warp owns keys [8*warp, 8*warp+8), 36 k-steps over 576 dims
     float c[4] = {0.f, 0.f, 0.f, 0.f};
     {
-      const __nv_bfloat16* bk = sKV + (8 * warp + group) * SKV_STRIDE + 2 * quad;
-      const __nv_bfloat16* qa0 = sQ + r0 * SQ_STRIDE + 2 * quad;
-      const __nv_bfloat16* qa1 = sQ + r1 * SQ_STRIDE + 2 * quad;
+      const __nv_bfloat16* bk = sKV + (8 * warp + group) * L::SKV_STRIDE + 2 * quad;
+      const __nv_bfloat16* qa0 = sQ + r0 * L::SQ_STRIDE + 2 * quad;
+      const __nv_bfloat16* qa1 = sQ + r1 * L::SQ_STRIDE + 2 * quad;
 #pragma unroll
-      for (int ks = 0; ks < D / 16; ++ks) {
+      for (int ks = 0; ks < L::D / 16; ++ks) {
         uint32_t a[4], b[2];
         a[0] = *reinterpret_cast<const uint32_t*>(qa0 + ks * 16);
         a[1] = *reinterpret_cast<const uint32_t*>(qa1 + ks * 16);
@@ -286,7 +335,7 @@ sparse_mla_fwd_kernel(const __nv_bfloat16* __restrict__ q,   // [T, h, 576]
 #pragma unroll
       for (int j = 0; j < 8; ++j) {  // 16-col chunks
         uint32_t r[4];
-        const __nv_bfloat16* src = sKV + (16 * s + mr) * SKV_STRIDE + 128 * warp + 16 * j + mc;
+        const __nv_bfloat16* src = sKV + (16 * s + mr) * L::SKV_STRIDE + 128 * warp + 16 * j + mc;
         ldmatrix_x4_trans(r, src);
         mma_bf16(pf[s], r, acc[2 * j]);          // cols [.. +0, +8)
         mma_bf16(pf[s], r + 2, acc[2 * j + 1]);  // cols [.. +8, +16)
@@ -302,7 +351,7 @@ sparse_mla_fwd_kernel(const __nv_bfloat16* __restrict__ q,   // [T, h, 576]
     if (hh >= h) continue;
     const float l = l_st[rr];
     const float inv = (l > 0.f) ? 1.f / l : 0.f;
-    __nv_bfloat16* orow = out + ((int64_t)t * h + hh) * DN;
+    __nv_bfloat16* orow = out + ((int64_t)t * h + hh) * L::VALUE;
 #pragma unroll
     for (int nt = 0; nt < 16; ++nt) {
       const int col = 128 * warp + 8 * nt + 2 * quad;
@@ -328,12 +377,16 @@ void sm89_sparse_mla_fwd(const torch::stable::Tensor& q,
   STD_TORCH_CHECK(q.is_cuda() && pool.is_cuda() && indices.is_cuda() &&
                       out.is_cuda() && lse.is_cuda(),
                   "all tensors must be CUDA");
-  STD_TORCH_CHECK(q.dim() == 3 && q.size(2) == 576, "q must be [T, h, 576]");
+  STD_TORCH_CHECK(q.dim() == 3 && (q.size(2) == 576 || q.size(2) == 512),
+                  "q must be [T, h, 576] or DeepSeek-V4 [T, h, 512]");
+  const bool dsv4 = q.size(2) == 512;
   STD_TORCH_CHECK(q.scalar_type() == torch::headeronly::ScalarType::BFloat16,
                   "q must be bf16");
-  STD_TORCH_CHECK(pool.dim() == 2 && pool.size(1) == 656 &&
-                      pool.scalar_type() == torch::headeronly::ScalarType::Byte,
-                  "pool must be [S, 656] u8");
+  const bool valid_pool = dsv4
+      ? pool.dim() == 3 && pool.size(2) == 584
+      : pool.dim() == 2 && pool.size(1) == 656;
+  STD_TORCH_CHECK(valid_pool && pool.scalar_type() == torch::headeronly::ScalarType::Byte,
+                  "pool must be [S, 656] u8 or DeepSeek-V4 [B, block_size, 584] u8");
   STD_TORCH_CHECK(indices.dim() == 2 &&
                       indices.scalar_type() == torch::headeronly::ScalarType::Int,
                   "indices must be [T, topk] i32");
@@ -346,9 +399,11 @@ void sm89_sparse_mla_fwd(const torch::stable::Tensor& q,
   STD_TORCH_CHECK(lse.dim() == 2 && lse.size(0) == T && lse.size(1) == h &&
                       lse.scalar_type() == torch::headeronly::ScalarType::Float,
                   "lse must be [T, h] f32");
-  STD_TORCH_CHECK(q.is_contiguous() && pool.is_contiguous() && indices.is_contiguous() &&
-                      out.is_contiguous() && lse.is_contiguous(),
-                  "all tensors must be contiguous");
+  STD_TORCH_CHECK(q.is_contiguous() && indices.is_contiguous() && out.is_contiguous() &&
+                      lse.is_contiguous(),
+                  "q, indices, out, and lse must be contiguous");
+  STD_TORCH_CHECK(!dsv4 || (pool.stride(1) == 584 && pool.stride(2) == 1),
+                  "DeepSeek-V4 pool rows must have dense 584-byte storage");
   const int32_t* lens_ptr = nullptr;
   if (topk_lens.has_value()) {
     const torch::stable::Tensor& lens = topk_lens.value();
@@ -361,20 +416,42 @@ void sm89_sparse_mla_fwd(const torch::stable::Tensor& q,
 
   const cudaStream_t stream = get_current_cuda_stream();
   dim3 grid(T, (h + sm89_dsa::HT - 1) / sm89_dsa::HT);
-  if (lens_ptr != nullptr) {
-    cudaFuncSetAttribute(sm89_dsa::sparse_mla_fwd_kernel<true>,
-                         cudaFuncAttributeMaxDynamicSharedMemorySize, sm89_dsa::SMEM_TOTAL);
-    sm89_dsa::sparse_mla_fwd_kernel<true><<<grid, 128, sm89_dsa::SMEM_TOTAL, stream>>>(
+  const int cache_block_size = dsv4 ? pool.size(1) : 0;
+  const int64_t cache_block_stride = dsv4 ? pool.stride(0) : 0;
+  if (dsv4 && lens_ptr != nullptr) {
+    constexpr int smem = sm89_dsa::Layout<true>::SMEM_TOTAL;
+    cudaFuncSetAttribute(sm89_dsa::sparse_mla_fwd_kernel<true, true>,
+                         cudaFuncAttributeMaxDynamicSharedMemorySize, smem);
+    sm89_dsa::sparse_mla_fwd_kernel<true, true><<<grid, 128, smem, stream>>>(
+        static_cast<const __nv_bfloat16*>(q.const_data_ptr()), pool.const_data_ptr<uint8_t>(),
+        indices.const_data_ptr<int32_t>(), static_cast<__nv_bfloat16*>(out.mutable_data_ptr()),
+        lse.mutable_data_ptr<float>(), lens_ptr, h, topk, (float)sm_scale, cache_block_size,
+        cache_block_stride);
+  } else if (dsv4) {
+    constexpr int smem = sm89_dsa::Layout<true>::SMEM_TOTAL;
+    cudaFuncSetAttribute(sm89_dsa::sparse_mla_fwd_kernel<false, true>,
+                         cudaFuncAttributeMaxDynamicSharedMemorySize, smem);
+    sm89_dsa::sparse_mla_fwd_kernel<false, true><<<grid, 128, smem, stream>>>(
+        static_cast<const __nv_bfloat16*>(q.const_data_ptr()), pool.const_data_ptr<uint8_t>(),
+        indices.const_data_ptr<int32_t>(), static_cast<__nv_bfloat16*>(out.mutable_data_ptr()),
+        lse.mutable_data_ptr<float>(), nullptr, h, topk, (float)sm_scale, cache_block_size,
+        cache_block_stride);
+  } else if (lens_ptr != nullptr) {
+    constexpr int smem = sm89_dsa::Layout<false>::SMEM_TOTAL;
+    cudaFuncSetAttribute(sm89_dsa::sparse_mla_fwd_kernel<true, false>,
+                         cudaFuncAttributeMaxDynamicSharedMemorySize, smem);
+    sm89_dsa::sparse_mla_fwd_kernel<true, false><<<grid, 128, smem, stream>>>(
         static_cast<const __nv_bfloat16*>(q.const_data_ptr()), pool.const_data_ptr<uint8_t>(),
         indices.const_data_ptr<int32_t>(), static_cast<__nv_bfloat16*>(out.mutable_data_ptr()),
-        lse.mutable_data_ptr<float>(), lens_ptr, h, topk, (float)sm_scale);
+        lse.mutable_data_ptr<float>(), lens_ptr, h, topk, (float)sm_scale, 0, 0);
   } else {
-    cudaFuncSetAttribute(sm89_dsa::sparse_mla_fwd_kernel<false>,
-                         cudaFuncAttributeMaxDynamicSharedMemorySize, sm89_dsa::SMEM_TOTAL);
-    sm89_dsa::sparse_mla_fwd_kernel<false><<<grid, 128, sm89_dsa::SMEM_TOTAL, stream>>>(
+    constexpr int smem = sm89_dsa::Layout<false>::SMEM_TOTAL;
+    cudaFuncSetAttribute(sm89_dsa::sparse_mla_fwd_kernel<false, false>,
+                         cudaFuncAttributeMaxDynamicSharedMemorySize, smem);
+    sm89_dsa::sparse_mla_fwd_kernel<false, false><<<grid, 128, smem, stream>>>(
         static_cast<const __nv_bfloat16*>(q.const_data_ptr()), pool.const_data_ptr<uint8_t>(),
         indices.const_data_ptr<int32_t>(), static_cast<__nv_bfloat16*>(out.mutable_data_ptr()),
-        lse.mutable_data_ptr<float>(), nullptr, h, topk, (float)sm_scale);
+        lse.mutable_data_ptr<float>(), nullptr, h, topk, (float)sm_scale, 0, 0);
   }
   STD_CUDA_KERNEL_LAUNCH_CHECK();
 #else