From 8e727a9c8e770267d74a9ee61c353b2868e31177 Mon Sep 17 00:00:00 2001
From: Shinto C V <cshintov@gmail.com>
Date: Thu, 27 Aug 2026 15:03:48 +0530
Subject: [PATCH] chore: make the branch pre-commit clean

The qwen4-main-squashed branch fails CI lint because it was never run
through the repo's pre-commit hooks. Formatting only, applied by
`pre-commit run --all-files`; no behaviour change.
---
 .../jit/csrc/attention/qsa_indexer.cuh        | 168 ++++++------------
 .../jit/csrc/elementwise/fast_topk.cuh        |  20 +--
 .../elementwise/grouped_gemma_rmsnorm.cuh     |  17 +-
 .../jit/csrc/elementwise/hc_combine.cuh       |  75 +++-----
 .../kernels/ops/attention/qsa_indexer.py      |   8 +-
 .../ops/attention/triton_gdn_fused_proj.py    |   5 +-
 .../kernels/ops/elementwise/fast_topk.py      |   2 +-
 .../kernels/ops/elementwise/hc_combine.py     |  10 +-
 .../sglang/kernels/ops/elementwise/hc_mix.py  |   4 +-
 ...fer_pr4266_dense_bf16_gemm_sm100_splitk.py |  24 +--
 .../ops/layernorm/grouped_gemma_rmsnorm.py    |   6 +-
 python/sglang/srt/disaggregation/decode.py    |   1 -
 .../attention/hybrid_linear_attn_backend.py   |   4 +-
 .../sglang/srt/layers/attention/qsa/config.py |   7 +-
 .../srt/layers/attention/qsa/dsa_indexer.py   |  32 +---
 .../layers/attention/qsa/graph_metadata.py    |   8 +-
 .../sglang/srt/layers/attention/qsa/kernel.py |   8 +-
 .../srt/layers/attention/qsa/metadata.py      |  13 +-
 python/sglang/srt/layers/attention/qsa/mqa.py |   8 +-
 .../srt/layers/attention/qsa/qsa_indexer.py   |  39 ++--
 .../attention/qwen_sparse_attn_backend.py     |  80 +++------
 python/sglang/srt/layers/hc_mix_triton.py     |   4 +-
 python/sglang/srt/layers/hyperconnection.py   |   4 +-
 .../sglang/srt/layers/quantization/unquant.py |   8 +-
 python/sglang/srt/mem_cache/common.py         |   3 +-
 python/sglang/srt/mem_cache/qsa_kv_pool.py    |   1 -
 python/sglang/srt/models/qwen4_exp_mtp.py     |  12 +-
 .../srt/multimodal/processors/qwen_vl.py      |   2 +-
 python/sglang/srt/speculative/draft_utils.py  |   4 +-
 .../benchmark/attention/bench_qsa_indexer.py  |   1 +
 .../benchmark/elementwise/bench_fast_topk.py  |   4 +-
 .../benchmark/elementwise/bench_hc_combine.py |   4 +-
 .../layernorm/bench_grouped_gemma_rmsnorm.py  |   4 +-
 .../layernorm/bench_qwen4_ple_norm.py         |   8 +-
 .../ops/attention/test_qkvzba_split_ratio3.py |   1 +
 .../kernels/ops/attention/test_qsa_indexer.py |  17 +-
 .../kernels/ops/elementwise/test_fast_topk.py |   4 +-
 .../ops/elementwise/test_hc_combine.py        |   8 +-
 .../ops/layernorm/test_qwen4_ple_norm.py      |   4 +-
 test/registered/kernels/test_hc_mix_triton.py |  11 +-
 test/registered/kernels/test_qsa.py           | 115 ++++++------
 .../kernels/test_verify_commit_triton.py      |   1 +
 .../unit/mem_cache/test_paged_free_segment.py |   1 -
 43 files changed, 287 insertions(+), 473 deletions(-)

diff --git a/python/sglang/kernels/jit/csrc/attention/qsa_indexer.cuh b/python/sglang/kernels/jit/csrc/attention/qsa_indexer.cuh
index 1e9e3b2ceec0..82dc7f6eb9ca 100644
--- a/python/sglang/kernels/jit/csrc/attention/qsa_indexer.cuh
+++ b/python/sglang/kernels/jit/csrc/attention/qsa_indexer.cuh
@@ -155,10 +155,7 @@ SGL_DEVICE void qsa_mrope_apply(
  */
 template <typename T, int kHeadDim>
 SGL_DEVICE void qsa_gemma_norm_row(
-    const T* __restrict__ x_row,
-    const T* __restrict__ weight,
-    const float eps,
-    T* __restrict__ smem_row) {
+    const T* __restrict__ x_row, const T* __restrict__ weight, const float eps, T* __restrict__ smem_row) {
   using namespace device;
   constexpr int kPerLane = kHeadDim / kWarpThreads;
   using vec_t = AlignedVector<T, kPerLane>;
@@ -204,15 +201,15 @@ SGL_DEVICE void qsa_gemma_norm_row(
 }
 
 struct QsaIndexQPrepParams {
-  const void* qk;                  // [tokens, (num_q_heads + 1) * kHeadDim]
-  void* q_out;                     // [tokens, q_heads_padded, kHeadDim]
-  const void* weight;              // [kHeadDim]
-  const float* cos_sin_cache;      // [positions_capacity, rotary_dim]
-  const int32_t* axis_map;         // [rotary_dim / 2]
-  const int64_t* positions;        // [num_axes, tokens] (row stride may exceed tokens)
-  const int64_t* cache_loc;        // [tokens]
-  void* key_state_buffer;          // [slots, kHeadDim]
-  int64_t* rope_position_buffer;   // [slots, 3]
+  const void* qk;                 // [tokens, (num_q_heads + 1) * kHeadDim]
+  void* q_out;                    // [tokens, q_heads_padded, kHeadDim]
+  const void* weight;             // [kHeadDim]
+  const float* cos_sin_cache;     // [positions_capacity, rotary_dim]
+  const int32_t* axis_map;        // [rotary_dim / 2]
+  const int64_t* positions;       // [num_axes, tokens] (row stride may exceed tokens)
+  const int64_t* cache_loc;       // [tokens]
+  void* key_state_buffer;         // [slots, kHeadDim]
+  int64_t* rope_position_buffer;  // [slots, 3]
   int64_t positions_stride;
   int32_t num_axes;
   int32_t num_q_heads;
@@ -227,8 +224,7 @@ struct QsaIndexQPrepParams {
  * One CTA (4 warps) per token; one warp per query head.
  */
 template <typename T, int kHeadDim, bool kIsNeox, bool kUsePDL>
-__global__ __launch_bounds__(128) void qsa_index_q_prep_kernel(
-    const QsaIndexQPrepParams __grid_constant__ params) {
+__global__ __launch_bounds__(128) void qsa_index_q_prep_kernel(const QsaIndexQPrepParams __grid_constant__ params) {
   using namespace device;
   constexpr int kPerLane = kHeadDim / kWarpThreads;
   using vec_t = AlignedVector<T, kPerLane>;
@@ -239,8 +235,7 @@ __global__ __launch_bounds__(128) void qsa_index_q_prep_kernel(
 
   device::PDLWaitPrimary<kUsePDL>();
 
-  const int64_t qk_row =
-      static_cast<int64_t>(token) * (params.num_q_heads + 1) * kHeadDim;
+  const int64_t qk_row = static_cast<int64_t>(token) * (params.num_q_heads + 1) * kHeadDim;
   const int64_t loc = params.cache_loc[token];
   int64_t pos[3];
 #pragma unroll
@@ -249,19 +244,13 @@ __global__ __launch_bounds__(128) void qsa_index_q_prep_kernel(
     pos[a] = params.positions[ax * params.positions_stride + token];
   }
 
-  for (int32_t h = static_cast<int32_t>(warp); h < params.q_heads_padded;
-       h += 4) {
-    T* out_row =
-        static_cast<T*>(params.q_out) +
-        (static_cast<int64_t>(token) * params.q_heads_padded + h) * kHeadDim;
+  for (int32_t h = static_cast<int32_t>(warp); h < params.q_heads_padded; h += 4) {
+    T* out_row = static_cast<T*>(params.q_out) + (static_cast<int64_t>(token) * params.q_heads_padded + h) * kHeadDim;
     if (h < params.num_q_heads) {
       const T* x_row = static_cast<const T*>(params.qk) + qk_row + h * kHeadDim;
-      qsa_gemma_norm_row<T, kHeadDim>(
-          x_row, static_cast<const T*>(params.weight), params.eps,
-          smem_rows[warp]);
+      qsa_gemma_norm_row<T, kHeadDim>(x_row, static_cast<const T*>(params.weight), params.eps, smem_rows[warp]);
       qsa_mrope_apply<T, kHeadDim, kIsNeox>(
-          smem_rows[warp], out_row, params.cos_sin_cache, params.axis_map, pos,
-          params.rotary_dim);
+          smem_rows[warp], out_row, params.cos_sin_cache, params.axis_map, pos, params.rotary_dim);
     } else {
       vec_t zv;
       zv.fill(DTypeTrait<T>::from(0.0f));
@@ -276,8 +265,7 @@ __global__ __launch_bounds__(128) void qsa_index_q_prep_kernel(
     kv.load(
         static_cast<const T*>(params.qk) + qk_row + params.num_q_heads * kHeadDim,
         lane);  // offset is in vector units
-    kv.store(static_cast<T*>(params.key_state_buffer) + loc * kHeadDim,
-             lane);
+    kv.store(static_cast<T*>(params.key_state_buffer) + loc * kHeadDim, lane);
   }
   if (warp == 1 && lane < 3) {
     params.rope_position_buffer[loc * 3 + lane] = pos[lane];
@@ -287,14 +275,14 @@ __global__ __launch_bounds__(128) void qsa_index_q_prep_kernel(
 }
 
 struct QsaIndexKCompressParams {
-  const void* key_state_buffer;      // [slots, kHeadDim]
-  const int32_t* group_locs;         // [groups, compress_ratio]
+  const void* key_state_buffer;         // [slots, kHeadDim]
+  const int32_t* group_locs;            // [groups, compress_ratio]
   const int64_t* rope_position_buffer;  // [slots, 3]
-  const float* cos_sin_cache;        // [positions_capacity, rotary_dim]
-  const int32_t* axis_map;           // [rotary_dim / 2]
-  const void* weight;                // [kHeadDim]
-  const int32_t* write_locs;         // [groups]
-  void* compressed_k_buffer;         // [compressed_slots, kHeadDim]
+  const float* cos_sin_cache;           // [positions_capacity, rotary_dim]
+  const int32_t* axis_map;              // [rotary_dim / 2]
+  const void* weight;                   // [kHeadDim]
+  const int32_t* write_locs;            // [groups]
+  void* compressed_k_buffer;            // [compressed_slots, kHeadDim]
   int32_t compress_ratio;
   int32_t rotary_dim;
   int32_t num_groups;
@@ -307,8 +295,8 @@ struct QsaIndexKCompressParams {
  * One warp per group.
  */
 template <typename T, int kHeadDim, bool kIsNeox, bool kUsePDL>
-__global__ __launch_bounds__(128) void qsa_index_k_compress_kernel(
-    const QsaIndexKCompressParams __grid_constant__ params) {
+__global__
+__launch_bounds__(128) void qsa_index_k_compress_kernel(const QsaIndexKCompressParams __grid_constant__ params) {
   using namespace device;
   constexpr int kPerLane = kHeadDim / kWarpThreads;
   using vec_t = AlignedVector<T, kPerLane>;
@@ -333,8 +321,7 @@ __global__ __launch_bounds__(128) void qsa_index_k_compress_kernel(
     for (int32_t r = 0; r < params.compress_ratio; ++r) {
       vec_t v;
       v.load(
-          static_cast<const T*>(params.key_state_buffer) +
-              static_cast<int64_t>(locs[r]) * kHeadDim,
+          static_cast<const T*>(params.key_state_buffer) + static_cast<int64_t>(locs[r]) * kHeadDim,
           lane);  // offset is in vector units
 #pragma unroll
       for (int i = 0; i < kPerLane; ++i) {
@@ -366,10 +353,8 @@ __global__ __launch_bounds__(128) void qsa_index_k_compress_kernel(
     const float nf = math::rsqrt(ss / kHeadDim + params.eps);
 #pragma unroll
     for (int i = 0; i < kPerLane; ++i) {
-      const float wf = static_cast<float>(
-          static_cast<const T*>(params.weight)[lane * kPerLane + i]);
-      smem_rows[warp][lane * kPerLane + i] =
-          DTypeTrait<T>::from(mf[i] * nf * (1.0f + wf));
+      const float wf = static_cast<float>(static_cast<const T*>(params.weight)[lane * kPerLane + i]);
+      smem_rows[warp][lane * kPerLane + i] = DTypeTrait<T>::from(mf[i] * nf * (1.0f + wf));
     }
     __syncwarp();
   }
@@ -380,11 +365,9 @@ __global__ __launch_bounds__(128) void qsa_index_k_compress_kernel(
     pos[a] = params.rope_position_buffer[static_cast<int64_t>(loc0) * 3 + a];
   }
 
-  T* out_row = static_cast<T*>(params.compressed_k_buffer) +
-               static_cast<int64_t>(params.write_locs[group]) * kHeadDim;
+  T* out_row = static_cast<T*>(params.compressed_k_buffer) + static_cast<int64_t>(params.write_locs[group]) * kHeadDim;
   qsa_mrope_apply<T, kHeadDim, kIsNeox>(
-      smem_rows[warp], out_row, params.cos_sin_cache, params.axis_map, pos,
-      params.rotary_dim);
+      smem_rows[warp], out_row, params.cos_sin_cache, params.axis_map, pos, params.rotary_dim);
 
   device::PDLTriggerSecondary<kUsePDL>();
 }
@@ -418,48 +401,25 @@ void qsa_index_q_prep(
   device.set_options<kDLCUDA>();
   constexpr int64_t D = kHeadDim;
 
-  TensorMatcher({tokens, (num_q_heads + 1) * D})
-      .with_dtype<T>()
-      .with_device(device)
-      .verify(qk);
+  TensorMatcher({tokens, (num_q_heads + 1) * D}).with_dtype<T>().with_device(device).verify(qk);
   auto heads_padded = SymbolicSize{"heads_padded"};
-  TensorMatcher({tokens, heads_padded, D})
-      .with_dtype<T>()
-      .with_device(device)
-      .verify(q_out);
+  TensorMatcher({tokens, heads_padded, D}).with_dtype<T>().with_device(device).verify(q_out);
   TensorMatcher({D}).with_dtype<T>().with_device(device).verify(weight);
   auto cache_rows = SymbolicSize{"cos_sin_cache_rows"};
-  TensorMatcher({cache_rows, rotary_dim})
-      .with_dtype<fp32_t>()
-      .with_device(device)
-      .verify(cos_sin_cache);
-  TensorMatcher({rotary_dim / 2})
-      .with_dtype<int32_t>()
-      .with_device(device)
-      .verify(axis_map);
-  TensorMatcher({num_axes, tokens})
-      .with_dtype<int64_t>()
-      .with_device(device)
-      .with_strides({-1, 1})
-      .verify(positions);
-  TensorMatcher({tokens}).with_dtype<int64_t>().with_device(device).verify(
-      cache_loc);
+  TensorMatcher({cache_rows, rotary_dim}).with_dtype<fp32_t>().with_device(device).verify(cos_sin_cache);
+  TensorMatcher({rotary_dim / 2}).with_dtype<int32_t>().with_device(device).verify(axis_map);
+  TensorMatcher({num_axes, tokens}).with_dtype<int64_t>().with_device(device).with_strides({-1, 1}).verify(positions);
+  TensorMatcher({tokens}).with_dtype<int64_t>().with_device(device).verify(cache_loc);
   auto slots = SymbolicSize{"state_slots"};
-  TensorMatcher({slots, D}).with_dtype<T>().with_device(device).verify(
-      key_state_buffer);
-  TensorMatcher({slots, 3})
-      .with_dtype<int64_t>()
-      .with_device(device)
-      .verify(rope_position_buffer);
+  TensorMatcher({slots, D}).with_dtype<T>().with_device(device).verify(key_state_buffer);
+  TensorMatcher({slots, 3}).with_dtype<int64_t>().with_device(device).verify(rope_position_buffer);
 
   const int64_t num_tokens = tokens.unwrap();
   const int64_t q_heads_padded = heads_padded.unwrap();
   CHECK_HOST(num_tokens > 0) << "qsa_index_q_prep: no tokens";
-  CHECK_HOST(num_axes == 1 || num_axes == 3)
-      << "qsa_index_q_prep: positions must have 1 or 3 axes, got " << num_axes;
+  CHECK_HOST(num_axes == 1 || num_axes == 3) << "qsa_index_q_prep: positions must have 1 or 3 axes, got " << num_axes;
   CHECK_HOST(q_heads_padded >= num_q_heads)
-      << "qsa_index_q_prep: padded heads " << q_heads_padded
-      << " < num_q_heads " << num_q_heads;
+      << "qsa_index_q_prep: padded heads " << q_heads_padded << " < num_q_heads " << num_q_heads;
   CHECK_HOST(rotary_dim > 0 && rotary_dim % 2 == 0 && rotary_dim <= D)
       << "qsa_index_q_prep: invalid rotary_dim " << rotary_dim;
 
@@ -472,8 +432,7 @@ void qsa_index_q_prep(
       .positions = static_cast<const int64_t*>(positions.data_ptr()),
       .cache_loc = static_cast<const int64_t*>(cache_loc.data_ptr()),
       .key_state_buffer = key_state_buffer.data_ptr(),
-      .rope_position_buffer =
-          static_cast<int64_t*>(rope_position_buffer.data_ptr()),
+      .rope_position_buffer = static_cast<int64_t*>(rope_position_buffer.data_ptr()),
       .positions_stride = positions.stride(0),
       .num_axes = static_cast<int32_t>(num_axes),
       .num_q_heads = static_cast<int32_t>(num_q_heads),
@@ -482,8 +441,7 @@ void qsa_index_q_prep(
       .eps = eps,
   };
   LaunchKernel(static_cast<uint32_t>(num_tokens), 128, device.unwrap())
-      .enable_pdl(kUsePDL)(
-          qsa_index_q_prep_kernel<T, kHeadDim, kIsNeox, kUsePDL>, params);
+      .enable_pdl(kUsePDL)(qsa_index_q_prep_kernel<T, kHeadDim, kIsNeox, kUsePDL>, params);
 }
 
 /**
@@ -508,34 +466,17 @@ void qsa_index_k_compress(
   constexpr int64_t D = kHeadDim;
 
   auto slots = SymbolicSize{"state_slots"};
-  TensorMatcher({slots, D}).with_dtype<T>().with_device(device).verify(
-      key_state_buffer);
+  TensorMatcher({slots, D}).with_dtype<T>().with_device(device).verify(key_state_buffer);
   auto groups = SymbolicSize{"groups"};
-  TensorMatcher({groups, compress_ratio})
-      .with_dtype<int32_t>()
-      .with_device(device)
-      .verify(group_locs);
-  TensorMatcher({slots, 3})
-      .with_dtype<int64_t>()
-      .with_device(device)
-      .verify(rope_position_buffer);
+  TensorMatcher({groups, compress_ratio}).with_dtype<int32_t>().with_device(device).verify(group_locs);
+  TensorMatcher({slots, 3}).with_dtype<int64_t>().with_device(device).verify(rope_position_buffer);
   auto cache_rows = SymbolicSize{"cos_sin_cache_rows"};
-  TensorMatcher({cache_rows, rotary_dim})
-      .with_dtype<fp32_t>()
-      .with_device(device)
-      .verify(cos_sin_cache);
-  TensorMatcher({rotary_dim / 2})
-      .with_dtype<int32_t>()
-      .with_device(device)
-      .verify(axis_map);
+  TensorMatcher({cache_rows, rotary_dim}).with_dtype<fp32_t>().with_device(device).verify(cos_sin_cache);
+  TensorMatcher({rotary_dim / 2}).with_dtype<int32_t>().with_device(device).verify(axis_map);
   TensorMatcher({D}).with_dtype<T>().with_device(device).verify(weight);
-  TensorMatcher({groups}).with_dtype<int32_t>().with_device(device).verify(
-      write_locs);
+  TensorMatcher({groups}).with_dtype<int32_t>().with_device(device).verify(write_locs);
   auto compressed_slots = SymbolicSize{"compressed_slots"};
-  TensorMatcher({compressed_slots, D})
-      .with_dtype<T>()
-      .with_device(device)
-      .verify(compressed_k_buffer);
+  TensorMatcher({compressed_slots, D}).with_dtype<T>().with_device(device).verify(compressed_k_buffer);
 
   const int64_t num_groups = groups.unwrap();
   CHECK_HOST(num_groups > 0) << "qsa_index_k_compress: no groups";
@@ -547,8 +488,7 @@ void qsa_index_k_compress(
   const auto params = QsaIndexKCompressParams{
       .key_state_buffer = key_state_buffer.data_ptr(),
       .group_locs = static_cast<const int32_t*>(group_locs.data_ptr()),
-      .rope_position_buffer =
-          static_cast<const int64_t*>(rope_position_buffer.data_ptr()),
+      .rope_position_buffer = static_cast<const int64_t*>(rope_position_buffer.data_ptr()),
       .cos_sin_cache = static_cast<const float*>(cos_sin_cache.data_ptr()),
       .axis_map = static_cast<const int32_t*>(axis_map.data_ptr()),
       .weight = weight.data_ptr(),
@@ -559,10 +499,8 @@ void qsa_index_k_compress(
       .num_groups = static_cast<int32_t>(num_groups),
       .eps = eps,
   };
-  LaunchKernel(
-      static_cast<uint32_t>(div_ceil(num_groups, 4)), 128, device.unwrap())
-      .enable_pdl(kUsePDL)(
-          qsa_index_k_compress_kernel<T, kHeadDim, kIsNeox, kUsePDL>, params);
+  LaunchKernel(static_cast<uint32_t>(div_ceil(num_groups, 4)), 128, device.unwrap())
+      .enable_pdl(kUsePDL)(qsa_index_k_compress_kernel<T, kHeadDim, kIsNeox, kUsePDL>, params);
 }
 
 }  // namespace sglang
diff --git a/python/sglang/kernels/jit/csrc/elementwise/fast_topk.cuh b/python/sglang/kernels/jit/csrc/elementwise/fast_topk.cuh
index d2741a66f453..ca5e05712158 100644
--- a/python/sglang/kernels/jit/csrc/elementwise/fast_topk.cuh
+++ b/python/sglang/kernels/jit/csrc/elementwise/fast_topk.cuh
@@ -9,6 +9,7 @@
 // is unspecified (atomic collection order), matching the AOT kernel.
 #include <sgl_kernel/tensor.h>
 #include <sgl_kernel/utils.h>
+
 #include <sgl_kernel/utils.cuh>
 
 #include <tvm/ffi/container/tensor.h>
@@ -33,8 +34,7 @@ struct FastTopKParams {
 SGL_DEVICE auto convert_to_uint8(float x) -> uint8_t {
   const __half h = __float2half_rn(x);
   const uint16_t bits = __half_as_ushort(h);
-  const uint16_t key = (bits & 0x8000) ? static_cast<uint16_t>(~bits)
-                                       : static_cast<uint16_t>(bits | 0x8000);
+  const uint16_t key = (bits & 0x8000) ? static_cast<uint16_t>(~bits) : static_cast<uint16_t>(bits | 0x8000);
   return static_cast<uint8_t>(key >> 8);
 }
 
@@ -45,8 +45,7 @@ SGL_DEVICE auto convert_to_uint32(float x) -> uint32_t {
 
 // When length <= kTopK, write the indices directly.
 template <int kTopK>
-SGL_DEVICE void naive_topk(
-    const float* __restrict__ score, int32_t* __restrict__ indice, int32_t length) {
+SGL_DEVICE void naive_topk(const float* __restrict__ score, int32_t* __restrict__ indice, int32_t length) {
   const auto tid = threadIdx.x;
   for (int i = tid; i < kTopK; i += kThreadsPerBlock) {
     indice[i] = (i < length) ? i : -1;
@@ -55,8 +54,7 @@ SGL_DEVICE void naive_topk(
 
 // Radix-select top-k. Assumes length > kTopK (checked by the caller).
 template <int kTopK>
-SGL_DEVICE void radix_select_topk(
-    const float* __restrict__ input, int* __restrict__ index, int row_start, int length) {
+SGL_DEVICE void radix_select_topk(const float* __restrict__ input, int* __restrict__ index, int row_start, int length) {
   int topk = kTopK;
   constexpr auto BLOCK_SIZE = kThreadsPerBlock;
   constexpr auto RADIX = 256;
@@ -156,9 +154,7 @@ SGL_DEVICE void radix_select_topk(
 
     // clip here to prevent overflow
     const auto _raw_num_input = s_num_input[r_idx];
-    const auto num_input = (_raw_num_input < int(SMEM_INPUT_SIZE))
-                               ? _raw_num_input
-                               : int(SMEM_INPUT_SIZE);
+    const auto num_input = (_raw_num_input < int(SMEM_INPUT_SIZE)) ? _raw_num_input : int(SMEM_INPUT_SIZE);
 
     run_cumsum();
     if (tx < RADIX && s_histogram[tx] > topk && s_histogram[tx + 1] <= topk) {
@@ -293,11 +289,7 @@ struct FastTopKKernel {
     };
 
     const auto num_rows = static_cast<uint32_t>(B.unwrap());
-    LaunchKernel(
-        num_rows,
-        fast_topk_detail::kThreadsPerBlock,
-        device.unwrap(),
-        fast_topk_detail::kSmemBytes)
+    LaunchKernel(num_rows, fast_topk_detail::kThreadsPerBlock, device.unwrap(), fast_topk_detail::kSmemBytes)
         .enable_pdl(kUsePDL)(kernel, params);
   }
 };
diff --git a/python/sglang/kernels/jit/csrc/elementwise/grouped_gemma_rmsnorm.cuh b/python/sglang/kernels/jit/csrc/elementwise/grouped_gemma_rmsnorm.cuh
index c5fa5f8aacaa..abdde55a6c44 100644
--- a/python/sglang/kernels/jit/csrc/elementwise/grouped_gemma_rmsnorm.cuh
+++ b/python/sglang/kernels/jit/csrc/elementwise/grouped_gemma_rmsnorm.cuh
@@ -63,12 +63,9 @@ __global__ __launch_bounds__(kGroupSize / 16) void grouped_gemma_rmsnorm_kernel(
 
   PDLWaitPrimary<kUsePDL>();
 
-  const auto input_ptr =
-      pointer::offset<Float>(params.input, static_cast<int64_t>(bid) * kGroupSize);
-  const auto output_ptr =
-      pointer::offset<Float>(params.output, static_cast<int64_t>(bid) * kGroupSize);
-  const auto weight_ptr =
-      pointer::offset<Float>(params.weight, static_cast<int64_t>(group) * kGroupSize);
+  const auto input_ptr = pointer::offset<Float>(params.input, static_cast<int64_t>(bid) * kGroupSize);
+  const auto output_ptr = pointer::offset<Float>(params.output, static_cast<int64_t>(bid) * kGroupSize);
+  const auto weight_ptr = pointer::offset<Float>(params.weight, static_cast<int64_t>(group) * kGroupSize);
 
   Storage input_vec[kNumLoads];
   Storage weight_vec[kNumLoads];
@@ -113,8 +110,7 @@ __global__ __launch_bounds__(kGroupSize / 16) void grouped_gemma_rmsnorm_kernel(
     for (uint32_t i = 0; i < kVecLen; ++i) {
       const auto [ix, iy] = cast<fp32x2_t>(input_vec[j][i]);
       const auto [wx, wy] = cast<fp32x2_t>(weight_vec[j][i]);
-      output_vec[i] = cast<Float2>(
-          fp32x2_t{ix * norm_factor * (1.0f + wx), iy * norm_factor * (1.0f + wy)});
+      output_vec[i] = cast<Float2>(fp32x2_t{ix * norm_factor * (1.0f + wx), iy * norm_factor * (1.0f + wy)});
     }
     gmem.store(output_ptr, output_vec, j);
   }
@@ -161,9 +157,8 @@ struct GroupedGemmaRMSNormKernel {
         .verify(output);
 
     const int64_t hidden_size = H.unwrap();
-    CHECK_HOST(hidden_size % kGroupSize == 0)
-        << "grouped_gemma_rmsnorm: hidden_size (" << hidden_size
-        << ") must be divisible by group_size (" << kGroupSize << ")";
+    CHECK_HOST(hidden_size % kGroupSize == 0) << "grouped_gemma_rmsnorm: hidden_size (" << hidden_size
+                                              << ") must be divisible by group_size (" << kGroupSize << ")";
 
     const auto params = GroupedGemmaRMSNormParams{
         .input = input.data_ptr(),
diff --git a/python/sglang/kernels/jit/csrc/elementwise/hc_combine.cuh b/python/sglang/kernels/jit/csrc/elementwise/hc_combine.cuh
index edcf86a665b4..5892d8826e35 100644
--- a/python/sglang/kernels/jit/csrc/elementwise/hc_combine.cuh
+++ b/python/sglang/kernels/jit/csrc/elementwise/hc_combine.cuh
@@ -38,8 +38,7 @@ struct HcCombineParams {
  * \tparam Float       Element type: bf16_t | fp16_t.
  */
 template <int64_t kHcCount, int64_t kHiddenSize, bool kUsePDL, typename Float>
-__global__ __launch_bounds__(256) void hc_combine_kernel(
-    const HcCombineParams __grid_constant__ params) {
+__global__ __launch_bounds__(256) void hc_combine_kernel(const HcCombineParams __grid_constant__ params) {
   using namespace device;
   using Float2 = packed_t<Float>;
   using Storage = AlignedVector<Float2, 4>;  // 8 elements, 16 bytes
@@ -54,15 +53,11 @@ __global__ __launch_bounds__(256) void hc_combine_kernel(
   const auto gmem = tile::Memory<Storage>::cta(kNumThreads);
   const uint32_t m = blockIdx.x;
 
-  const auto y_ptr =
-      pointer::offset<Float>(params.block_output, static_cast<int64_t>(m) * kHiddenSize);
-  const auto r_ptr =
-      pointer::offset<Float>(params.residual, static_cast<int64_t>(m) * kRowSize);
-  const auto n_ptr =
-      pointer::offset<Float>(params.normed_residual, static_cast<int64_t>(m) * kRowSize);
+  const auto y_ptr = pointer::offset<Float>(params.block_output, static_cast<int64_t>(m) * kHiddenSize);
+  const auto r_ptr = pointer::offset<Float>(params.residual, static_cast<int64_t>(m) * kRowSize);
+  const auto n_ptr = pointer::offset<Float>(params.normed_residual, static_cast<int64_t>(m) * kRowSize);
   const auto w_ptr = static_cast<const Float*>(params.inject_weight);
-  const auto out_ptr =
-      pointer::offset<Float>(params.output, static_cast<int64_t>(m) * kRowSize);
+  const auto out_ptr = pointer::offset<Float>(params.output, static_cast<int64_t>(m) * kRowSize);
 
   PDLWaitPrimary<kUsePDL>();
 
@@ -143,10 +138,8 @@ template <int64_t kHcCount, int64_t kHiddenSize, bool kUsePDL, typename DType>
 struct HcCombineKernel {
   static_assert(sizeof(DType) == 2, "HcCombine only supports 2-byte dtypes");
   static_assert(kHcCount > 0, "kHcCount must be positive");
-  static_assert(kHiddenSize > 0 && kHiddenSize % 8 == 0,
-                "kHiddenSize must be a multiple of 8");
-  static_assert((kHcCount * kHiddenSize) % (256 * 8) == 0,
-                "kHcCount * kHiddenSize must be a multiple of 2048");
+  static_assert(kHiddenSize > 0 && kHiddenSize % 8 == 0, "kHiddenSize must be a multiple of 8");
+  static_assert((kHcCount * kHiddenSize) % (256 * 8) == 0, "kHcCount * kHiddenSize must be a multiple of 2048");
   static constexpr auto kernel = hc_combine_kernel<kHcCount, kHiddenSize, kUsePDL, DType>;
   static constexpr uint32_t kBlockSize = 256;
 
@@ -198,7 +191,6 @@ struct HcCombineKernel {
   }
 };
 
-
 struct HcCombineSplitParams {
   const void* block_output;
   const void* residual;
@@ -230,8 +222,8 @@ constexpr uint32_t kVecLen = 8;
  * buffer clearing are needed.
  */
 template <int64_t kHcCount, int64_t kHiddenSize, bool kUsePDL, typename Float>
-__global__ __launch_bounds__(hc_combine_split_detail::kGateThreads)
-    void hc_combine_gate_kernel(const HcCombineSplitParams __grid_constant__ params) {
+__global__ __launch_bounds__(hc_combine_split_detail::kGateThreads) void hc_combine_gate_kernel(
+    const HcCombineSplitParams __grid_constant__ params) {
   using namespace device;
   using namespace hc_combine_split_detail;
   using Float2 = packed_t<Float>;
@@ -248,8 +240,7 @@ __global__ __launch_bounds__(hc_combine_split_detail::kGateThreads)
   const uint32_t c = blockIdx.y % kHcCount;
   const uint32_t ref_tid = split * kGateThreads + threadIdx.x;
 
-  const auto n_ptr =
-      pointer::offset<Float>(params.normed_residual, static_cast<int64_t>(m) * kRowSize);
+  const auto n_ptr = pointer::offset<Float>(params.normed_residual, static_cast<int64_t>(m) * kRowSize);
   const auto w_ptr = static_cast<const Float*>(params.inject_weight);
 
   PDLWaitPrimary<kUsePDL>();
@@ -276,8 +267,7 @@ __global__ __launch_bounds__(hc_combine_split_detail::kGateThreads)
     }
     sum = warp::reduce_sum(sum);
     if (threadIdx.x == 0) {
-      params.partials[(static_cast<int64_t>(m) * kSplit + split) * kHcCount + c] =
-          sum;
+      params.partials[(static_cast<int64_t>(m) * kSplit + split) * kHcCount + c] = sum;
     }
   }
 
@@ -292,8 +282,8 @@ __global__ __launch_bounds__(hc_combine_split_detail::kGateThreads)
  * per-CTA scalar.
  */
 template <int64_t kHcCount, int64_t kHiddenSize, bool kUsePDL, typename Float>
-__global__ __launch_bounds__(hc_combine_split_detail::kApplyThreads)
-    void hc_combine_apply_kernel(const HcCombineSplitParams __grid_constant__ params) {
+__global__ __launch_bounds__(hc_combine_split_detail::kApplyThreads) void hc_combine_apply_kernel(
+    const HcCombineSplitParams __grid_constant__ params) {
   using namespace device;
   using namespace hc_combine_split_detail;
   using Float2 = packed_t<Float>;
@@ -310,20 +300,16 @@ __global__ __launch_bounds__(hc_combine_split_detail::kApplyThreads)
   const uint32_t vec_base = split * kVecsPerSplit;
   const uint32_t branch = vec_base / kVecsPerBranch;
 
-  const auto y_ptr =
-      pointer::offset<Float>(params.block_output, static_cast<int64_t>(m) * kHiddenSize);
-  const auto r_ptr =
-      pointer::offset<Float>(params.residual, static_cast<int64_t>(m) * kRowSize);
-  const auto out_ptr =
-      pointer::offset<Float>(params.output, static_cast<int64_t>(m) * kRowSize);
+  const auto y_ptr = pointer::offset<Float>(params.block_output, static_cast<int64_t>(m) * kHiddenSize);
+  const auto r_ptr = pointer::offset<Float>(params.residual, static_cast<int64_t>(m) * kRowSize);
+  const auto out_ptr = pointer::offset<Float>(params.output, static_cast<int64_t>(m) * kRowSize);
 
   PDLWaitPrimary<kUsePDL>();
 
   float total = 0.0f;
 #pragma unroll
   for (uint32_t s = 0; s < kSplit; ++s) {
-    total += params.partials[(static_cast<int64_t>(m) * kSplit + s) * kHcCount +
-                             branch];
+    total += params.partials[(static_cast<int64_t>(m) * kSplit + s) * kHcCount + branch];
   }
   const float a = 2.0f / (1.0f + math::exp(-total / kHcCount));
 
@@ -351,10 +337,8 @@ __global__ __launch_bounds__(hc_combine_split_detail::kApplyThreads)
 template <int64_t kHcCount, int64_t kHiddenSize, bool kUsePDL, typename DType>
 struct HcCombineSplitKernel {
   static_assert(sizeof(DType) == 2, "HcCombine only supports 2-byte dtypes");
-  static constexpr auto gate_kernel =
-      hc_combine_gate_kernel<kHcCount, kHiddenSize, kUsePDL, DType>;
-  static constexpr auto apply_kernel =
-      hc_combine_apply_kernel<kHcCount, kHiddenSize, kUsePDL, DType>;
+  static constexpr auto gate_kernel = hc_combine_gate_kernel<kHcCount, kHiddenSize, kUsePDL, DType>;
+  static constexpr auto apply_kernel = hc_combine_apply_kernel<kHcCount, kHiddenSize, kUsePDL, DType>;
 
   static void
   run(const tvm::ffi::TensorView block_output,
@@ -369,25 +353,16 @@ struct HcCombineSplitKernel {
     auto device = SymbolicDevice{};
     device.set_options<kDLCUDA>();
 
-    TensorMatcher({M, kHiddenSize})
-        .with_dtype<DType>()
-        .with_device(device)
-        .verify(block_output);
+    TensorMatcher({M, kHiddenSize}).with_dtype<DType>().with_device(device).verify(block_output);
     TensorMatcher({M, kHcCount * kHiddenSize})
         .with_dtype<DType>()
         .with_device(device)
         .verify(residual)
         .verify(normed_residual)
         .verify(output);
-    TensorMatcher({kHcCount, kHcCount * kHiddenSize})
-        .with_dtype<DType>()
-        .with_device(device)
-        .verify(inject_weight);
+    TensorMatcher({kHcCount, kHcCount * kHiddenSize}).with_dtype<DType>().with_device(device).verify(inject_weight);
     auto part_rows = SymbolicSize{"partial_rows"};
-    TensorMatcher({part_rows, kSplit, kHcCount})
-        .with_dtype<fp32_t>()
-        .with_device(device)
-        .verify(partials);
+    TensorMatcher({part_rows, kSplit, kHcCount}).with_dtype<fp32_t>().with_device(device).verify(partials);
 
     const auto params = HcCombineSplitParams{
         .block_output = block_output.data_ptr(),
@@ -399,11 +374,9 @@ struct HcCombineSplitKernel {
     };
 
     const auto num_tokens = static_cast<uint32_t>(M.unwrap());
-    LaunchKernel(dim3(num_tokens, kSplit * kHcCount, 1), kGateThreads,
-                 device.unwrap())
+    LaunchKernel(dim3(num_tokens, kSplit * kHcCount, 1), kGateThreads, device.unwrap())
         .enable_pdl(kUsePDL)(gate_kernel, params);
-    LaunchKernel(dim3(num_tokens, kSplit, 1), kApplyThreads, device.unwrap())
-        .enable_pdl(kUsePDL)(apply_kernel, params);
+    LaunchKernel(dim3(num_tokens, kSplit, 1), kApplyThreads, device.unwrap()).enable_pdl(kUsePDL)(apply_kernel, params);
   }
 };
 
diff --git a/python/sglang/kernels/ops/attention/qsa_indexer.py b/python/sglang/kernels/ops/attention/qsa_indexer.py
index 59e38518a44f..054fba6eac7a 100644
--- a/python/sglang/kernels/ops/attention/qsa_indexer.py
+++ b/python/sglang/kernels/ops/attention/qsa_indexer.py
@@ -39,9 +39,7 @@ def _jit_qsa_indexer_module(
 ) -> Module:
     """Compile and cache the JIT QSA indexer module for one specialisation."""
     if dtype not in (torch.bfloat16, torch.float16):
-        raise RuntimeError(
-            f"Unsupported dtype {dtype}. Supported: bfloat16, float16"
-        )
+        raise RuntimeError(f"Unsupported dtype {dtype}. Supported: bfloat16, float16")
     if head_dim not in (64, 128, 256):
         raise RuntimeError(
             f"Unsupported index head_dim {head_dim}. Supported: 64, 128, 256"
@@ -162,9 +160,7 @@ def qsa_index_k_compress_store(
     is_neox_style    : NeoX (True) or GPT-J (False) RoPE pairing
     """
     head_dim = weight.shape[0]
-    module = _jit_qsa_indexer_module(
-        key_state_buffer.dtype, head_dim, is_neox_style
-    )
+    module = _jit_qsa_indexer_module(key_state_buffer.dtype, head_dim, is_neox_style)
     module.k_compress(
         key_state_buffer,
         group_locs,
diff --git a/python/sglang/kernels/ops/attention/triton_gdn_fused_proj.py b/python/sglang/kernels/ops/attention/triton_gdn_fused_proj.py
index 74633be7813d..dbcf8a086e9e 100644
--- a/python/sglang/kernels/ops/attention/triton_gdn_fused_proj.py
+++ b/python/sglang/kernels/ops/attention/triton_gdn_fused_proj.py
@@ -243,10 +243,7 @@ def fused_qkvzba_split_reshape_cat_contiguous_kernel(
         + offs_v
     )
     blk_z_st_ptr = (
-        z
-        + i_bs * NUM_HEADS_V * HEAD_V
-        + i_qk * V_PER_GROUP * HEAD_V
-        + offs_v
+        z + i_bs * NUM_HEADS_V * HEAD_V + i_qk * V_PER_GROUP * HEAD_V + offs_v
     )
 
     tl.store(blk_q_st_ptr, tl.load(blk_q_ptr))
diff --git a/python/sglang/kernels/ops/elementwise/fast_topk.py b/python/sglang/kernels/ops/elementwise/fast_topk.py
index c242ba74ae8b..7e8757fc75ae 100644
--- a/python/sglang/kernels/ops/elementwise/fast_topk.py
+++ b/python/sglang/kernels/ops/elementwise/fast_topk.py
@@ -1,6 +1,6 @@
 from __future__ import annotations
 
-from typing import Optional, TYPE_CHECKING
+from typing import TYPE_CHECKING, Optional
 
 import torch
 
diff --git a/python/sglang/kernels/ops/elementwise/hc_combine.py b/python/sglang/kernels/ops/elementwise/hc_combine.py
index c4a9e1794207..4245cd07c30c 100644
--- a/python/sglang/kernels/ops/elementwise/hc_combine.py
+++ b/python/sglang/kernels/ops/elementwise/hc_combine.py
@@ -1,6 +1,6 @@
 from __future__ import annotations
 
-from typing import Optional, TYPE_CHECKING
+from typing import TYPE_CHECKING, Optional
 
 import torch
 
@@ -16,15 +16,15 @@
 
 
 @cache_once
-def _jit_hc_combine_module(hc_count: int, hidden_size: int, dtype: torch.dtype) -> Module:
+def _jit_hc_combine_module(
+    hc_count: int, hidden_size: int, dtype: torch.dtype
+) -> Module:
     """Compile and cache the JIT HC combine module for a given shape/dtype."""
     # Checks on the compile key live here, not in `hc_combine`: `cache_once`
     # keys on (hc_count, hidden_size, dtype), so this runs once per
     # specialisation instead of once per call.
     if dtype not in (torch.bfloat16, torch.float16):
-        raise RuntimeError(
-            f"Unsupported dtype {dtype}. Supported: bfloat16, float16"
-        )
+        raise RuntimeError(f"Unsupported dtype {dtype}. Supported: bfloat16, float16")
     if hidden_size <= 0 or hidden_size % 8 != 0:
         raise RuntimeError(
             f"Unsupported hidden_size {hidden_size}. Must be a multiple of 8."
diff --git a/python/sglang/kernels/ops/elementwise/hc_mix.py b/python/sglang/kernels/ops/elementwise/hc_mix.py
index 36a43d227840..d1c911ea6227 100644
--- a/python/sglang/kernels/ops/elementwise/hc_mix.py
+++ b/python/sglang/kernels/ops/elementwise/hc_mix.py
@@ -92,9 +92,7 @@ def hc_mix(
     if rows == 0:
         return out
     inv_hc = 1.0 / hc_count
-    t_pad = _get_scratch(
-        lowrank, hyper_input_normed.dtype, hyper_input_normed.device
-    )
+    t_pad = _get_scratch(lowrank, hyper_input_normed.dtype, hyper_input_normed.device)
     for row_start in range(0, rows, _MAX_ROWS):
         row_end = min(row_start + _MAX_ROWS, rows)
         m = row_end - row_start
diff --git a/python/sglang/kernels/ops/gemm/flashinfer_pr4266_dense_bf16_gemm_sm100_splitk.py b/python/sglang/kernels/ops/gemm/flashinfer_pr4266_dense_bf16_gemm_sm100_splitk.py
index 1b3cf2528e38..fe2d03229c9d 100644
--- a/python/sglang/kernels/ops/gemm/flashinfer_pr4266_dense_bf16_gemm_sm100_splitk.py
+++ b/python/sglang/kernels/ops/gemm/flashinfer_pr4266_dense_bf16_gemm_sm100_splitk.py
@@ -316,6 +316,7 @@ def _store_shared_remote_v4(
 def _sigmoid_f32(v):
     return cute_math.rcp(cute_math.exp(v * -1.0) + 1.0)
 
+
 #: Epilogue modes; "none" preserves the vendored store path byte-for-byte.
 _EPILOGUE_MODES = ("none", "silu", "gate")
 
@@ -370,9 +371,7 @@ def __init__(
                     f"{gate_out_elems} outputs; must be a multiple of 128"
                 )
             gate_smem = (
-                _align_up(
-                    _smem_bytes(tactic, tactic.ab_stages), _GATE_TILE_ALIGN_BYTES
-                )
+                _align_up(_smem_bytes(tactic, tactic.ab_stages), _GATE_TILE_ALIGN_BYTES)
                 + tactic.mma_m * tactic.mma_n * _FP32_BYTES
             )
             if gate_smem > _SMEM_CAPACITY_BYTES:
@@ -888,9 +887,7 @@ def epilog_warp(
                 )
                 rSig.store(_sigmoid_f32(rAcc.load()))
                 sGate_epi = cute.flat_divide(gate_tile, epi_tile)
-                cute_ext.partition_and_copy(
-                    thr_t2r, rSig, sGate_epi[None, None, 0, 0]
-                )
+                cute_ext.partition_and_copy(thr_t2r, rSig, sGate_epi[None, None, 0, 0])
                 cute.arch.barrier(
                     barrier_id=_GATE_BARRIER_ID,
                     number_of_threads=self.epilog_threads,
@@ -912,9 +909,9 @@ def epilog_warp(
                             sig = gate_tile[j_local * group + g, m_local]
                             xv = mX[m_global, g * hs + j_global].to(cutlass.Float32)
                             gated = gated + sig * xv
-                        mOut[m_global, j_global] = (
-                            gated * self.epilogue_scale
-                        ).to(c_dtype)
+                        mOut[m_global, j_global] = (gated * self.epilogue_scale).to(
+                            c_dtype
+                        )
             else:
                 rD.store(rAcc.load().to(c_dtype))
                 # Preserve TMEM coordinates; the copy predicates output tails.
@@ -1038,12 +1035,8 @@ def _make_compile_repr_tensors(
     if epilogue_mode == "gate":
         return (
             *tensors,
-            _from_dlpack_dynamic(
-                _torch.empty((n, m), dtype=dtype, device="cuda"), 1
-            ),
-            _from_dlpack_dynamic(
-                _torch.empty((n, m), dtype=dtype, device="cuda"), 1
-            ),
+            _from_dlpack_dynamic(_torch.empty((n, m), dtype=dtype, device="cuda"), 1),
+            _from_dlpack_dynamic(_torch.empty((n, m), dtype=dtype, device="cuda"), 1),
         )
     if not has_bias:
         return (*tensors, None)
@@ -1154,6 +1147,7 @@ def _validate_runtime_tensors(a, b, bias, out) -> tuple[int, int, int]:
         tensor.dtype != a.dtype for tensor in tensors
     ):
         raise ValueError("a, b, out, and bias must share BF16 or FP16 dtype")
+
     def _is_dense_2d(tensor: _torch.Tensor) -> bool:
         rows, cols = tensor.shape
         return (tensor.stride(1) == 1 and tensor.stride(0) >= cols) or (
diff --git a/python/sglang/kernels/ops/layernorm/grouped_gemma_rmsnorm.py b/python/sglang/kernels/ops/layernorm/grouped_gemma_rmsnorm.py
index febfeff2ecb0..8fd57a6378cd 100644
--- a/python/sglang/kernels/ops/layernorm/grouped_gemma_rmsnorm.py
+++ b/python/sglang/kernels/ops/layernorm/grouped_gemma_rmsnorm.py
@@ -1,6 +1,6 @@
 from __future__ import annotations
 
-from typing import Optional, TYPE_CHECKING
+from typing import TYPE_CHECKING, Optional
 
 import torch
 
@@ -22,9 +22,7 @@ def _jit_grouped_gemma_rmsnorm_module(group_size: int, dtype: torch.dtype) -> Mo
     # `cache_once` keys on (group_size, dtype), so this runs once per
     # specialisation instead of once per call.
     if dtype not in (torch.bfloat16, torch.float16):
-        raise RuntimeError(
-            f"Unsupported dtype {dtype}. Supported: bfloat16, float16"
-        )
+        raise RuntimeError(f"Unsupported dtype {dtype}. Supported: bfloat16, float16")
     if group_size <= 0 or group_size % 512 != 0:
         raise RuntimeError(
             f"Unsupported group_size {group_size}. Must be a multiple of 512."
diff --git a/python/sglang/srt/disaggregation/decode.py b/python/sglang/srt/disaggregation/decode.py
index 495c4c397cf8..97eaa0625aa4 100644
--- a/python/sglang/srt/disaggregation/decode.py
+++ b/python/sglang/srt/disaggregation/decode.py
@@ -268,7 +268,6 @@ def clear(self):
         self.mamba_allocator.clear()
 
 
-
 @dataclass
 class DecodeRequest:
     req: Req
diff --git a/python/sglang/srt/layers/attention/hybrid_linear_attn_backend.py b/python/sglang/srt/layers/attention/hybrid_linear_attn_backend.py
index beea4dcf3825..fbc5c1d5d198 100644
--- a/python/sglang/srt/layers/attention/hybrid_linear_attn_backend.py
+++ b/python/sglang/srt/layers/attention/hybrid_linear_attn_backend.py
@@ -1031,9 +1031,7 @@ def init_forward_metadata_in_graph(self, forward_batch: ForwardBatch):
 
     def get_indexer_metadata(self, layer_id: int, forward_batch: ForwardBatch):
         if layer_id in self.full_attn_layers:
-            return self.full_attn_backend.get_indexer_metadata(
-                layer_id, forward_batch
-            )
+            return self.full_attn_backend.get_indexer_metadata(layer_id, forward_batch)
         return None
 
     def on_after_cuda_graph_warmup(self):
diff --git a/python/sglang/srt/layers/attention/qsa/config.py b/python/sglang/srt/layers/attention/qsa/config.py
index 7c7cc2ad84fa..fef9b8dc6003 100644
--- a/python/sglang/srt/layers/attention/qsa/config.py
+++ b/python/sglang/srt/layers/attention/qsa/config.py
@@ -169,10 +169,9 @@ def parse_qsa_profile(config) -> Optional[QSAProfile]:
     if text_config is None:
         return None
     has_compressed = getattr(text_config, "indexer_n_heads", None) is not None
-    has_tokenwise = (
-        getattr(text_config, "index_topk", None) is not None
-        and _is_qwen_family(text_config)
-    )
+    has_tokenwise = getattr(
+        text_config, "index_topk", None
+    ) is not None and _is_qwen_family(text_config)
     if has_compressed and has_tokenwise:
         raise ValueError(
             "Ambiguous QSA config: both compressed (indexer_*) and tokenwise "
diff --git a/python/sglang/srt/layers/attention/qsa/dsa_indexer.py b/python/sglang/srt/layers/attention/qsa/dsa_indexer.py
index c8aff144d954..408cd2e53def 100644
--- a/python/sglang/srt/layers/attention/qsa/dsa_indexer.py
+++ b/python/sglang/srt/layers/attention/qsa/dsa_indexer.py
@@ -95,9 +95,7 @@ def __init__(
         if page_size != 64:
             # The paged index-K layout and every fast path assume 64-token
             # pages, matching qsa_0511.
-            raise ValueError(
-                f"tokenwise QSA requires page_size = 64, got {page_size}"
-            )
+            raise ValueError(f"tokenwise QSA requires page_size = 64, got {page_size}")
         if envs.SGLANG_QWEN_DSA_USE_FP8_INDEXER.get():
             raise NotImplementedError(
                 "The FP8 tokenwise indexer path (qsa_0511 deep_gemm) is not "
@@ -174,9 +172,7 @@ def __init__(
             dtype=torch.get_default_dtype(),
         )
 
-    def project_qkw(
-        self, hidden_states: torch.Tensor, positions: torch.Tensor
-    ):
+    def project_qkw(self, hidden_states: torch.Tensor, positions: torch.Tensor):
         """Fused Q/K/W projection, per-head RMS norm and indexer RoPE."""
 
         qkw, _ = self.index_qkw_proj(hidden_states)
@@ -200,15 +196,11 @@ def forward_cuda(
         indexer_metadata,
     ) -> torch.Tensor:
         forward_mode = forward_batch.forward_mode
-        is_target_verify = getattr(
-            forward_mode, "is_target_verify", lambda: False
-        )()
-        is_draft_extend = getattr(
-            forward_mode, "is_draft_extend", lambda **_: False
-        )(include_v2=True)
-        is_paged = (
-            forward_mode.is_decode() or is_target_verify or is_draft_extend
+        is_target_verify = getattr(forward_mode, "is_target_verify", lambda: False)()
+        is_draft_extend = getattr(forward_mode, "is_draft_extend", lambda **_: False)(
+            include_v2=True
         )
+        is_paged = forward_mode.is_decode() or is_target_verify or is_draft_extend
         if is_paged:
             # See the compressed QSAIndexer: speculative/decode rows derive
             # their physical causal length from the paged metadata, not from
@@ -217,9 +209,7 @@ def forward_cuda(
         else:
             logical_positions = getattr(forward_batch, "positions", None)
             if logical_positions is None:
-                logical_positions = (
-                    positions[0] if positions.ndim == 2 else positions
-                )
+                logical_positions = positions[0] if positions.ndim == 2 else positions
             logical_positions = logical_positions.flatten()
 
         # DP padding adds token rows without assigning them to a request;
@@ -268,9 +258,7 @@ def forward_cuda(
         out_cache_loc = getattr(indexer_metadata, "out_cache_loc", None)
         if out_cache_loc is None:
             out_cache_loc = forward_batch.out_cache_loc
-        pool.set_dsa_index_k_buffer(
-            self.layer_id, out_cache_loc[:num_valid_tokens], k
-        )
+        pool.set_dsa_index_k_buffer(self.layer_id, out_cache_loc[:num_valid_tokens], k)
 
         if is_paged:
             return self._select_paged(q, w, indexer_metadata)
@@ -298,9 +286,7 @@ def _select_paged(
         output = torch.full(
             (rows, self.token_topk), -1, dtype=torch.int32, device=q.device
         )
-        row_chunk = _qsa_prefill_row_chunk_size(
-            rows, max_len, self.index_n_heads
-        )
+        row_chunk = _qsa_prefill_row_chunk_size(rows, max_len, self.index_n_heads)
         table_long = table.long()
         for row_start in range(0, rows, row_chunk):
             row_end = min(row_start + row_chunk, rows)
diff --git a/python/sglang/srt/layers/attention/qsa/graph_metadata.py b/python/sglang/srt/layers/attention/qsa/graph_metadata.py
index 16ed0bb2d9b0..2c8964ecd2ff 100644
--- a/python/sglang/srt/layers/attention/qsa/graph_metadata.py
+++ b/python/sglang/srt/layers/attention/qsa/graph_metadata.py
@@ -164,7 +164,9 @@ def _qsa_graph_row_metadata_kernel(
     for p0 in range(0, max_pages, PAGE_BLOCK):
         idx = p0 + offs
         valid = idx < tl.minimum(max_pages, row_width_pages)
-        loc = tl.load(req_to_token_ptr + token_row + idx * FULL_PAGE, mask=valid, other=0)
+        loc = tl.load(
+            req_to_token_ptr + token_row + idx * FULL_PAGE, mask=valid, other=0
+        )
         tl.store(table_row + idx, tl.maximum(loc // FULL_PAGE, 0), mask=valid)
 
 
@@ -173,9 +175,7 @@ def supports_graph_metadata_kernels(pool, device) -> bool:
 
     from sglang.srt.mem_cache.qsa_kv_pool import QSATokenToKVPool
 
-    return torch.device(device).type == "cuda" and isinstance(
-        pool, QSATokenToKVPool
-    )
+    return torch.device(device).type == "cuda" and isinstance(pool, QSATokenToKVPool)
 
 
 def launch_graph_metadata(
diff --git a/python/sglang/srt/layers/attention/qsa/kernel.py b/python/sglang/srt/layers/attention/qsa/kernel.py
index c9500f1bf785..1d8e9ec2d9f7 100644
--- a/python/sglang/srt/layers/attention/qsa/kernel.py
+++ b/python/sglang/srt/layers/attention/qsa/kernel.py
@@ -40,9 +40,7 @@ def qsa_fast_topk(
 
         from sgl_kernel import top_k as top_k_module
 
-        supported_topk = getattr(
-            top_k_module, "_FAST_TOPK_SUPPORTED_K", (2048,)
-        )
+        supported_topk = getattr(top_k_module, "_FAST_TOPK_SUPPORTED_K", (2048,))
         if topk in supported_topk:
             return top_k_module.fast_topk_v2(
                 logits, lengths, topk=topk, row_starts=starts
@@ -196,9 +194,7 @@ def _expand_qsa_block_indices_kernel(
         ).to(tl.int32),
         axis=0,
     )
-    valid_token_count = tl.minimum(
-        valid_block_count * COMPRESS_RATIO, TOKEN_TOPK
-    )
+    valid_token_count = tl.minimum(valid_block_count * COMPRESS_RATIO, TOKEN_TOPK)
 
     query_position = tl.load(query_positions + row)
     visible_tokens = query_position + 1
diff --git a/python/sglang/srt/layers/attention/qsa/metadata.py b/python/sglang/srt/layers/attention/qsa/metadata.py
index b7a41afd9c29..261333b4f5c8 100644
--- a/python/sglang/srt/layers/attention/qsa/metadata.py
+++ b/python/sglang/srt/layers/attention/qsa/metadata.py
@@ -140,9 +140,7 @@ def get_prefill_mqa_inputs(
         sequence_lengths = self.sequence_lengths.to(torch.int32)
         sequence_lengths_list = sequence_lengths.tolist()
         for sequence_id in range(len(sequence_lengths_list)):
-            complete_blocks = (
-                int(sequence_lengths_list[sequence_id]) // ratio
-            )
+            complete_blocks = int(sequence_lengths_list[sequence_id]) // ratio
             if complete_blocks == 0:
                 continue
             # DSV4-style addressing: a group's compressed slot is its first
@@ -150,8 +148,9 @@ def get_prefill_mqa_inputs(
             # contiguous in one page), read straight off the request's
             # token-slot row.
             compressed_locs = (
-                self.token_slot_table[sequence_id, : complete_blocks * ratio : ratio]
-                .long()
+                self.token_slot_table[
+                    sequence_id, : complete_blocks * ratio : ratio
+                ].long()
                 // ratio
             )
             parts.append(compressed_buffer.index_select(0, compressed_locs))
@@ -274,7 +273,9 @@ def build_group_ring_slots(
     """Ring slots of a planned group's members, oldest first."""
     requests = req_pool_indices.long()[sequence_ids]
     offsets = torch.arange(
-        compress_ratio - 1, -1, -1,
+        compress_ratio - 1,
+        -1,
+        -1,
         device=group_end_positions.device,
         dtype=torch.long,
     )
diff --git a/python/sglang/srt/layers/attention/qsa/mqa.py b/python/sglang/srt/layers/attention/qsa/mqa.py
index a2bfed445dac..02719934dd79 100644
--- a/python/sglang/srt/layers/attention/qsa/mqa.py
+++ b/python/sglang/srt/layers/attention/qsa/mqa.py
@@ -258,9 +258,7 @@ def kernel(
                                         0,
                                         :,
                                     ],
-                                    k_shared[
-                                        sp * page_size : (sp + 1) * page_size, :
-                                    ],
+                                    k_shared[sp * page_size : (sp + 1) * page_size, :],
                                 )
                         T.gemm(
                             k_shared,
@@ -307,9 +305,7 @@ def tilelang_qsa_mqa_prefill(
     # Allocate the padded output once. Appending even a few padding rows with
     # torch.cat would allocate and copy the entire [rows, keys] FP32 matrix,
     # temporarily doubling the dominant prefill buffer for long contexts.
-    logits = torch.zeros(
-        (padded_rows, keys), dtype=torch.float32, device=q.device
-    )
+    logits = torch.zeros((padded_rows, keys), dtype=torch.float32, device=q.device)
     q_padded = q.to(torch.bfloat16).contiguous()
     starts = row_starts.to(device=q.device, dtype=torch.int32).contiguous()
     ends = row_ends.to(device=q.device, dtype=torch.int32).contiguous()
diff --git a/python/sglang/srt/layers/attention/qsa/qsa_indexer.py b/python/sglang/srt/layers/attention/qsa/qsa_indexer.py
index ddbbdaf2a48a..a2e47b542247 100644
--- a/python/sglang/srt/layers/attention/qsa/qsa_indexer.py
+++ b/python/sglang/srt/layers/attention/qsa/qsa_indexer.py
@@ -23,19 +23,18 @@
 from sglang.srt.layers.utils import MultiPlatformOp
 from sglang.srt.model_executor.runner import get_is_capture_mode
 
-
 # Bound the dominant FP32 [query_rows, compressed_keys] prefill workspace.
 # Top-k is row-independent, so large scheduler chunks can be scored in smaller
 # row tiles without changing the selected blocks.
 _QSA_PREFILL_LOGITS_BUDGET_BYTES = 128 * 1024 * 1024
+
+
 def _qsa_prefill_row_chunk_size(rows: int, keys: int, heads: int) -> int:
     if rows <= 0 or keys <= 0:
         return max(rows, 1)
     block_q = max(1, 128 // heads)
     bytes_per_row = keys * torch.float32.itemsize
-    max_padded_rows = max(
-        block_q, _QSA_PREFILL_LOGITS_BUDGET_BYTES // bytes_per_row
-    )
+    max_padded_rows = max(block_q, _QSA_PREFILL_LOGITS_BUDGET_BYTES // bytes_per_row)
     max_padded_rows = max(block_q, max_padded_rows // block_q * block_q)
     return min(rows, max_padded_rows)
 
@@ -218,9 +217,10 @@ def normalize_compressed_keys(
         return self.apply_rope(block_positions, normalized)
 
     def _use_fused_compress(self, pool) -> bool:
-        return (
-            getattr(pool, "qsa_rope_position_buffer", None) is not None
-            and self._use_fused_prep(pool.get_qsa_key_state_buffer(self.layer_id))
+        return getattr(
+            pool, "qsa_rope_position_buffer", None
+        ) is not None and self._use_fused_prep(
+            pool.get_qsa_key_state_buffer(self.layer_id)
         )
 
     def _fused_compress_store(
@@ -284,7 +284,6 @@ def _group_ring_slots(
             compress_ratio=self.compress_ratio,
         )
 
-
     def update_key_state_and_compress(
         self,
         token_k: torch.Tensor,
@@ -408,10 +407,14 @@ def apply_rope(self, positions: torch.Tensor, tensor: torch.Tensor) -> torch.Ten
         if tensor.numel() == 0:
             return tensor
         positions = positions.long()
-        num_positions = positions.shape[-1] if positions.ndim == 2 else positions.numel()
+        num_positions = (
+            positions.shape[-1] if positions.ndim == 2 else positions.numel()
+        )
         if num_positions != tensor.shape[0]:
             raise ValueError("QSA RoPE positions must match the token dimension")
-        if not get_is_capture_mode() and hasattr(self.rotary_emb, "_ensure_cos_sin_cache_length"):
+        if not get_is_capture_mode() and hasattr(
+            self.rotary_emb, "_ensure_cos_sin_cache_length"
+        ):
             self.rotary_emb._ensure_cos_sin_cache_length(int(positions.max().item()))
 
         # Let the exact Qwen4-Exp RoPE instance compose regular or three-axis
@@ -540,12 +543,8 @@ def forward_cuda(
         indexer_metadata,
     ) -> torch.Tensor:
         forward_mode = forward_batch.forward_mode
-        is_target_verify = getattr(
-            forward_mode, "is_target_verify", lambda: False
-        )()
-        is_draft_extend = getattr(
-            forward_mode, "is_draft_extend_v2", lambda: False
-        )()
+        is_target_verify = getattr(forward_mode, "is_target_verify", lambda: False)()
+        is_draft_extend = getattr(forward_mode, "is_draft_extend_v2", lambda: False)()
         if forward_mode.is_decode() or is_target_verify or is_draft_extend:
             # EAGLE/MTP may advance the model's RoPE coordinate independently
             # from the physical paged-KV position.  Compression and sparse
@@ -557,9 +556,7 @@ def forward_cuda(
         else:
             logical_positions = getattr(forward_batch, "positions", None)
             if logical_positions is None:
-                logical_positions = (
-                    positions[0] if positions.ndim == 2 else positions
-                )
+                logical_positions = positions[0] if positions.ndim == 2 else positions
             logical_positions = logical_positions.flatten()
         # DP MAX_LEN padding adds token rows without assigning them to a
         # request. token_to_batch_idx is the source of truth for semantic rows.
@@ -634,9 +631,7 @@ def forward_cuda(
             )
 
         compressed_keys, row_starts, row_ends, sequence_lengths = (
-            indexer_metadata.get_prefill_mqa_inputs(
-                self.layer_id, logical_positions
-            )
+            indexer_metadata.get_prefill_mqa_inputs(self.layer_id, logical_positions)
         )
         query_sequence_ids = indexer_metadata.get_token_to_batch_idx()
         row_sequence_lengths = sequence_lengths.index_select(
diff --git a/python/sglang/srt/layers/attention/qwen_sparse_attn_backend.py b/python/sglang/srt/layers/attention/qwen_sparse_attn_backend.py
index 7758a88feaea..98dbc8cef7ec 100644
--- a/python/sglang/srt/layers/attention/qwen_sparse_attn_backend.py
+++ b/python/sglang/srt/layers/attention/qwen_sparse_attn_backend.py
@@ -78,9 +78,7 @@ def _resolve_flash_attn_varlen_func():
     except ImportError:
         pass
     try:
-        from flash_attn.cute.interface import (
-            flash_attn_varlen_func as cute_varlen_func,
-        )
+        from flash_attn.cute.interface import flash_attn_varlen_func as cute_varlen_func
 
         def flash_attn_varlen_func(*args, **kwargs):
             output = cute_varlen_func(*args, **kwargs)
@@ -95,7 +93,6 @@ def flash_attn_varlen_func(*args, **kwargs):
         ) from exc
 
 
-
 class QwenSparseAttnMetadata(msgspec.Struct, frozen=True):
     """Per-forward metadata consumed by core sparse attention."""
 
@@ -257,8 +254,7 @@ def _require_chain_speculation(self, forward_mode, spec_info) -> None:
             return
         if int(getattr(spec_info, "topk", 1) or 1) != 1:
             raise NotImplementedError(
-                "Qwen QSA target verification supports only "
-                "speculative_eagle_topk=1"
+                "Qwen QSA target verification supports only " "speculative_eagle_topk=1"
             )
         draft_tokens = int(getattr(spec_info, "draft_token_num", 0) or 0)
         if draft_tokens > self.compress_ratio:
@@ -288,9 +284,7 @@ def _speculative_max_row_length(forward_batch, sequence_lengths) -> int:
             )
             return max(1, int(sequence_lengths.max()))
         spec_info = forward_batch.spec_info
-        draft_window = (
-            int(spec_info.draft_token_num) if spec_info is not None else 0
-        )
+        draft_window = int(spec_info.draft_token_num) if spec_info is not None else 0
         return max(1, int(seq_lens_cpu.max()) + draft_window)
 
     @staticmethod
@@ -378,9 +372,7 @@ def _graph_speculative_layout(
                     extend_lengths = torch.cat(
                         [
                             extend_lengths,
-                            torch.zeros(
-                                bs - extend_lengths.numel(), dtype=torch.int32
-                            ),
+                            torch.zeros(bs - extend_lengths.numel(), dtype=torch.int32),
                         ]
                     )
             else:
@@ -414,9 +406,7 @@ def _graph_speculative_layout(
                 torch.full((int(extend_len),), prefix_len, dtype=torch.int32)
             )
         row_lengths = (
-            torch.cat(row_lengths)
-            if row_lengths
-            else torch.empty(0, dtype=torch.int32)
+            torch.cat(row_lengths) if row_lengths else torch.empty(0, dtype=torch.int32)
         )
         row_prefix_lengths = (
             torch.cat(row_prefix_lengths)
@@ -429,12 +419,8 @@ def _graph_speculative_layout(
                 "QSA CUDA graph speculative layout has inconsistent token count: "
                 f"capacity={num_tokens}, actual={actual_rows}"
             )
-        repeats = extend_lengths.to(
-            device=req_pool_indices.device, dtype=torch.long
-        )
-        row_req_pool_indices = torch.repeat_interleave(
-            req_pool_indices[:bs], repeats
-        )
+        repeats = extend_lengths.to(device=req_pool_indices.device, dtype=torch.long)
+        row_req_pool_indices = torch.repeat_interleave(req_pool_indices[:bs], repeats)
         # Draft-extend graphs always execute the captured static token shape,
         # while a replay can contain fewer accepted tokens.  The runner packs
         # real rows first and zero-fills the tail, so give those tail rows safe
@@ -542,9 +528,7 @@ def _qsa_write_plan(
             # member sits chunk-locally at (block * ratio - prefix).
             member_rows = torch.where(
                 valid,
-                row_token_starts[rows]
-                + blocks * compress_ratio
-                - prefix_lens[rows],
+                row_token_starts[rows] + blocks * compress_ratio - prefix_lens[rows],
                 torch.zeros_like(blocks),
             )
         return write_locs, group_end_positions, rows, member_rows
@@ -609,10 +593,7 @@ def _metadata_from_forward_batch(self, forward_batch) -> QwenSparseAttnMetadata:
             self.device = forward_batch.seq_lens.device
         if not self.max_context_len:
             self.max_context_len = self.req_to_token.shape[1]
-        if (
-            forward_batch.forward_mode.is_idle()
-            or forward_batch.seq_lens.numel() == 0
-        ):
+        if forward_batch.forward_mode.is_idle() or forward_batch.seq_lens.numel() == 0:
             # DP attention runs IDLE dummy forwards on ranks without work, and
             # the MTP multi-step wrapper forwards them as zero-row DECODE
             # steps.  Model layers skip attention for these batches, but
@@ -620,9 +601,7 @@ def _metadata_from_forward_batch(self, forward_batch) -> QwenSparseAttnMetadata:
             # of falling into the extend/decode paths on empty tensors.
             return self._empty_metadata(forward_batch)
         original_mode = getattr(forward_batch, "_original_forward_mode", None)
-        if original_mode is not None and self._is_speculative_paged_mode(
-            original_mode
-        ):
+        if original_mode is not None and self._is_speculative_paged_mode(original_mode):
             # DP MAX_LEN pseudo-extend rewrites the mode to EXTEND with
             # extend_seq_lens == 1 per request, which loses the per-request
             # draft fan-out of target_verify/draft_extend.  Refuse to guess
@@ -632,9 +611,7 @@ def _metadata_from_forward_batch(self, forward_batch) -> QwenSparseAttnMetadata:
                 f"speculative mode {original_mode}: token rows would be "
                 "mis-mapped to requests"
             )
-        speculative_paged = self._is_speculative_paged_mode(
-            forward_batch.forward_mode
-        )
+        speculative_paged = self._is_speculative_paged_mode(forward_batch.forward_mode)
         if speculative_paged:
             logical_positions = forward_batch.positions
             if logical_positions.ndim == 2:
@@ -690,9 +667,7 @@ def _metadata_from_forward_batch(self, forward_batch) -> QwenSparseAttnMetadata:
             else:
                 extend_seq_lens = forward_batch.extend_seq_lens
                 if extend_seq_lens is None:
-                    raise ValueError(
-                        "QSA extend metadata requires extend_seq_lens"
-                    )
+                    raise ValueError("QSA extend metadata requires extend_seq_lens")
                 token_to_batch_idx = torch.repeat_interleave(
                     torch.arange(
                         batch_size,
@@ -912,8 +887,7 @@ def init_cuda_graph_state(self, max_bs: int, max_num_tokens: int) -> None:
             max_bs, dtype=torch.int32, device=self.device
         )
         self._graph_extend_lens_pin = [
-            torch.zeros(max_bs, dtype=torch.int32, pin_memory=True)
-            for _ in range(2)
+            torch.zeros(max_bs, dtype=torch.int32, pin_memory=True) for _ in range(2)
         ]
         self._extend_lens_pin_idx = 0
 
@@ -1064,9 +1038,7 @@ def _replay_cuda_graph_metadata(
             )
         else:
             metadata.sequence_lengths.copy_(seq_lens[:bs].to(torch.int32))
-            metadata.row_req_pool_indices.copy_(
-                req_pool_indices[:bs].to(torch.int32)
-            )
+            metadata.row_req_pool_indices.copy_(req_pool_indices[:bs].to(torch.int32))
             metadata.indexer_metadata.graph_prefix_lengths.copy_(
                 (seq_lens[:bs] - 1).clamp_min(0).to(torch.int32)
             )
@@ -1228,8 +1200,7 @@ def _update_qsa_cuda_graph_metadata(
         row_width_pages = self.req_to_token.shape[1] // full_page
         num_pages = min(max_pages, row_width_pages)
         table = (
-            self.req_to_token[req_indices, : num_pages * full_page : full_page]
-            .long()
+            self.req_to_token[req_indices, : num_pages * full_page : full_page].long()
             // full_page
         ).clamp_min(0)
         page_table[:, :num_pages].copy_(table.to(torch.int32))
@@ -1531,7 +1502,6 @@ def _get_fa2_scratch(
             self._fa2_scratch[key] = buffers
         return buffers[0][:capacity], buffers[1][:capacity]
 
-
     def _get_trtllm_sparse_tables(self, batch, pages_per_row, page, device):
         key = (batch, pages_per_row, device)
         cached = self._trtllm_sparse_tables.get(key)
@@ -1541,9 +1511,7 @@ def _get_trtllm_sparse_tables(self, batch, pages_per_row, page, device):
             block_tables = (
                 torch.arange(batch, dtype=torch.int32, device=device)[:, None]
                 * pages_per_row
-                + torch.arange(pages_per_row, dtype=torch.int32, device=device)[
-                    None, :
-                ]
+                + torch.arange(pages_per_row, dtype=torch.int32, device=device)[None, :]
             ).contiguous()
             cached = (cu, block_tables)
             self._trtllm_sparse_tables[key] = cached
@@ -1582,9 +1550,7 @@ def _forward_trtllm_sparse(
         cu_strided, block_tables = self._get_trtllm_sparse_tables(
             batch, pages_per_row, page, device
         )
-        capacity_rows = (
-            self._cuda_graph_max_tokens if metadata.is_cuda_graph else batch
-        )
+        capacity_rows = self._cuda_graph_max_tokens if metadata.is_cuda_graph else batch
         packed_k, packed_v = self._get_fa2_scratch(
             max(capacity_rows, batch) * stride,
             k_buffer.shape[1],
@@ -1812,9 +1778,7 @@ def _step_out_cache_loc(self, forward_batch, step: int):
             .reshape(steps, -1)[step]
         )
 
-    def _make_step_forward_batch(
-        self, forward_batch, step: int, num_padding: int = 0
-    ):
+    def _make_step_forward_batch(self, forward_batch, step: int, num_padding: int = 0):
         step_forward_batch = copy(forward_batch)
         step_forward_batch.forward_mode = ForwardMode.DECODE
         step_forward_batch.seq_lens = (forward_batch.seq_lens + step + 1).to(
@@ -1841,9 +1805,7 @@ def _make_step_forward_batch(
                 )
                 step_forward_batch.seq_lens_cpu[-num_padding:] = 1
         step_forward_batch.batch_size = int(step_forward_batch.seq_lens.numel())
-        step_forward_batch.out_cache_loc = self._step_out_cache_loc(
-            forward_batch, step
-        )
+        step_forward_batch.out_cache_loc = self._step_out_cache_loc(forward_batch, step)
         return step_forward_batch
 
     def set_mtp_shared_sparse_indices(self, state) -> None:
@@ -1860,9 +1822,7 @@ def init_cuda_graph_state(self, max_bs: int, max_num_tokens: int):
         for backend in self.attn_backends:
             backend.init_cuda_graph_state(max_bs, max_num_tokens)
 
-    def init_forward_metadata_out_graph(
-        self, forward_batch, in_capture: bool = False
-    ):
+    def init_forward_metadata_out_graph(self, forward_batch, in_capture: bool = False):
         if in_capture:
             for step, backend in enumerate(self.attn_backends):
                 # Every capture row is synthetic.  Keep its sequence length
diff --git a/python/sglang/srt/layers/hc_mix_triton.py b/python/sglang/srt/layers/hc_mix_triton.py
index eae575f143a5..734ad2f342c5 100644
--- a/python/sglang/srt/layers/hc_mix_triton.py
+++ b/python/sglang/srt/layers/hc_mix_triton.py
@@ -216,9 +216,7 @@ def fused_hc_mix(
     device = hyper_input_normed.device
     num_ctas = torch.cuda.get_device_properties(device).multi_processor_count
     t_raw = torch.empty((rows_pad, lowrank), dtype=torch.float32, device=device)
-    out = torch.empty(
-        (rows, hs), dtype=hyper_input_normed.dtype, device=device
-    )
+    out = torch.empty((rows, hs), dtype=hyper_input_normed.dtype, device=device)
     if rows == 0:
         return out
     _hc_mix_persistent_kernel[(num_ctas,)](
diff --git a/python/sglang/srt/layers/hyperconnection.py b/python/sglang/srt/layers/hyperconnection.py
index 30f7e83008d0..654a0a66e56a 100644
--- a/python/sglang/srt/layers/hyperconnection.py
+++ b/python/sglang/srt/layers/hyperconnection.py
@@ -33,9 +33,7 @@ def __init__(
         self.weight.weight_loader = self._weight_loader
         # The JIT kernel requires group_size to be a multiple of 512; this is
         # init-static, so resolve it once here (device/dtype stay per-call).
-        effective_group_size = (
-            group_size if group_size is not None else hidden_size
-        )
+        effective_group_size = group_size if group_size is not None else hidden_size
         self._jit_group_size = (
             effective_group_size if effective_group_size % 512 == 0 else None
         )
diff --git a/python/sglang/srt/layers/quantization/unquant.py b/python/sglang/srt/layers/quantization/unquant.py
index 5d72e46f7b76..c0cf4fc82cf5 100644
--- a/python/sglang/srt/layers/quantization/unquant.py
+++ b/python/sglang/srt/layers/quantization/unquant.py
@@ -211,13 +211,9 @@ def _precompile_splitk_tactics() -> None:
     torch.cuda.synchronize()
 
 
-def _flashinfer_pr4266_bf16_gemm(
-    x: torch.Tensor, weight: torch.Tensor
-) -> torch.Tensor:
+def _flashinfer_pr4266_bf16_gemm(x: torch.Tensor, weight: torch.Tensor) -> torch.Tensor:
     x_2d = x.view(-1, x.shape[-1])
-    out = torch.empty(
-        (x_2d.shape[0], weight.shape[0]), dtype=x.dtype, device=x.device
-    )
+    out = torch.empty((x_2d.shape[0], weight.shape[0]), dtype=x.dtype, device=x.device)
     tactic = _flashinfer_pr4266_splitk_tactic(
         *_FLASHINFER_PR4266_TUNED_TACTICS[
             (x_2d.shape[0], weight.shape[0], weight.shape[1])
diff --git a/python/sglang/srt/mem_cache/common.py b/python/sglang/srt/mem_cache/common.py
index b8183e5aa7b3..4bc61811b882 100644
--- a/python/sglang/srt/mem_cache/common.py
+++ b/python/sglang/srt/mem_cache/common.py
@@ -15,7 +15,7 @@
 from sglang.srt.mem_cache.hicache_storage import PoolTransfer
 from sglang.srt.mem_cache.memory_pool import HybridReqToTokenPool, ReqToTokenPool
 from sglang.srt.runtime_context import get_serving, get_spec
-from sglang.srt.utils.common import ceil_align, ceil_div
+from sglang.srt.utils.common import ceil_align
 
 if TYPE_CHECKING:
     from sglang.srt.managers.schedule_batch import Req
@@ -138,7 +138,6 @@ def evict_from_tree_cache(tree_cache: BasePrefixCache | None, num_tokens: int):
             tree_cache.evict(EvictParams(num_tokens=num_tokens - available_size))
 
 
-
 def retraction_backup(
     req: Req,
     tree_cache: BasePrefixCache,
diff --git a/python/sglang/srt/mem_cache/qsa_kv_pool.py b/python/sglang/srt/mem_cache/qsa_kv_pool.py
index 8e84e12aa57b..e2b63daae571 100644
--- a/python/sglang/srt/mem_cache/qsa_kv_pool.py
+++ b/python/sglang/srt/mem_cache/qsa_kv_pool.py
@@ -52,7 +52,6 @@ def qsa_bytes_per_token(
         )
         return index_k_bytes // compress_ratio * num_layers
 
-
     def __init__(
         self,
         *,
diff --git a/python/sglang/srt/models/qwen4_exp_mtp.py b/python/sglang/srt/models/qwen4_exp_mtp.py
index 115fd4504c4d..5f0becb2bdf0 100644
--- a/python/sglang/srt/models/qwen4_exp_mtp.py
+++ b/python/sglang/srt/models/qwen4_exp_mtp.py
@@ -79,10 +79,14 @@ def _init_pre_fc_norms(self, config: PretrainedConfig) -> None:
             if self.hc_count > 1
             else config.hidden_size
         )
-        self.pre_fc_norm_hidden = GemmaRMSNorm(hidden_norm_size, eps=config.rms_norm_eps)
+        self.pre_fc_norm_hidden = GemmaRMSNorm(
+            hidden_norm_size, eps=config.rms_norm_eps
+        )
 
     def _init_linear_projections(self, config: PretrainedConfig) -> None:
-        self.fc_embedding = nn.Linear(config.hidden_size, config.hidden_size, bias=False)
+        self.fc_embedding = nn.Linear(
+            config.hidden_size, config.hidden_size, bias=False
+        )
         self.fc_hidden = nn.Linear(config.hidden_size, config.hidden_size, bias=False)
 
     def _init_standard_fusion(self, config: PretrainedConfig):
@@ -213,7 +217,9 @@ def forward(
         logits_output = self.logits_processor(
             input_ids, hidden_states, self.lm_head, forward_batch
         )
-        self._set_hc_logits_hidden_states(logits_output, hc_hidden_states, forward_batch)
+        self._set_hc_logits_hidden_states(
+            logits_output, hc_hidden_states, forward_batch
+        )
         return logits_output
 
 
diff --git a/python/sglang/srt/multimodal/processors/qwen_vl.py b/python/sglang/srt/multimodal/processors/qwen_vl.py
index cec691583cbd..c4e2a33c726a 100644
--- a/python/sglang/srt/multimodal/processors/qwen_vl.py
+++ b/python/sglang/srt/multimodal/processors/qwen_vl.py
@@ -29,9 +29,9 @@
 )
 from sglang.srt.models.qwen3_5_mtp import Qwen3_5ForCausalLMMTP
 from sglang.srt.models.qwen3_omni_moe import Qwen3OmniMoeForConditionalGeneration
-from sglang.srt.models.qwen4_exp import Qwen4ExpForConditionalGeneration
 from sglang.srt.models.qwen3_vl import Qwen3VLForConditionalGeneration
 from sglang.srt.models.qwen3_vl_moe import Qwen3VLMoeForConditionalGeneration
+from sglang.srt.models.qwen4_exp import Qwen4ExpForConditionalGeneration
 from sglang.srt.multimodal.processors.base_processor import (
     BaseMultimodalProcessor as SGLangBaseProcessor,
 )
diff --git a/python/sglang/srt/speculative/draft_utils.py b/python/sglang/srt/speculative/draft_utils.py
index cc2c98dec906..300588bef5bd 100644
--- a/python/sglang/srt/speculative/draft_utils.py
+++ b/python/sglang/srt/speculative/draft_utils.py
@@ -113,9 +113,7 @@ def create_draft_extend_backend(self):
                 parse_qsa_profile,
             )
 
-            profile = parse_qsa_profile(
-                self.draft_model_runner.model_config.hf_config
-            )
+            profile = parse_qsa_profile(self.draft_model_runner.model_config.hf_config)
             if profile is not None and profile.variant != QSA_VARIANT_COMPRESSED:
                 # Tokenwise QSA has no graph-stable indexer metadata; keep
                 # the intentional eager draft-extend path and never fall
diff --git a/test/registered/kernels/benchmark/attention/bench_qsa_indexer.py b/test/registered/kernels/benchmark/attention/bench_qsa_indexer.py
index 4afb9401f10e..25727ae8eb7f 100644
--- a/test/registered/kernels/benchmark/attention/bench_qsa_indexer.py
+++ b/test/registered/kernels/benchmark/attention/bench_qsa_indexer.py
@@ -9,6 +9,7 @@
 from types import SimpleNamespace
 
 import torch
+
 from sglang.kernels.jit.benchmark import marker
 from sglang.test.ci.ci_register import register_cuda_ci
 
diff --git a/test/registered/kernels/benchmark/elementwise/bench_fast_topk.py b/test/registered/kernels/benchmark/elementwise/bench_fast_topk.py
index e6c674fa0a80..56e2f3e9a373 100644
--- a/test/registered/kernels/benchmark/elementwise/bench_fast_topk.py
+++ b/test/registered/kernels/benchmark/elementwise/bench_fast_topk.py
@@ -5,7 +5,9 @@
 from sglang.kernels.ops.elementwise.fast_topk import fast_topk as jit_fast_topk
 from sglang.test.ci.ci_register import register_cuda_ci
 
-register_cuda_ci(est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-large")
+register_cuda_ci(
+    est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-large"
+)
 
 
 def torch_topk(score, lengths, topk):
diff --git a/test/registered/kernels/benchmark/elementwise/bench_hc_combine.py b/test/registered/kernels/benchmark/elementwise/bench_hc_combine.py
index 74a915c6ab9a..367043ad7b5f 100644
--- a/test/registered/kernels/benchmark/elementwise/bench_hc_combine.py
+++ b/test/registered/kernels/benchmark/elementwise/bench_hc_combine.py
@@ -6,7 +6,9 @@
 from sglang.kernels.ops.elementwise.hc_combine import hc_combine as jit_hc_combine
 from sglang.test.ci.ci_register import register_cuda_ci
 
-register_cuda_ci(est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-large")
+register_cuda_ci(
+    est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-large"
+)
 
 HC_COUNT = 4
 HIDDEN_SIZE = 2560
diff --git a/test/registered/kernels/benchmark/layernorm/bench_grouped_gemma_rmsnorm.py b/test/registered/kernels/benchmark/layernorm/bench_grouped_gemma_rmsnorm.py
index 16bbc4fa490c..b5355d2b14b3 100644
--- a/test/registered/kernels/benchmark/layernorm/bench_grouped_gemma_rmsnorm.py
+++ b/test/registered/kernels/benchmark/layernorm/bench_grouped_gemma_rmsnorm.py
@@ -7,7 +7,9 @@
 )
 from sglang.test.ci.ci_register import register_cuda_ci
 
-register_cuda_ci(est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-large")
+register_cuda_ci(
+    est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-large"
+)
 
 
 def torch_impl_grouped_gemma_rmsnorm(
diff --git a/test/registered/kernels/benchmark/layernorm/bench_qwen4_ple_norm.py b/test/registered/kernels/benchmark/layernorm/bench_qwen4_ple_norm.py
index fa9d0a79b24f..42ef0e88ec69 100644
--- a/test/registered/kernels/benchmark/layernorm/bench_qwen4_ple_norm.py
+++ b/test/registered/kernels/benchmark/layernorm/bench_qwen4_ple_norm.py
@@ -5,12 +5,12 @@
 from sglang.srt.models.qwen4_exp import Qwen4ExpPLEGroupedNorm
 from sglang.test.ci.ci_register import register_cuda_ci
 
-register_cuda_ci(est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-large")
+register_cuda_ci(
+    est_time=30, stage="base-b-kernel-benchmark", runner_config="1-gpu-large"
+)
 
 
-def torch_impl_ple_norm(
-    norm: Qwen4ExpPLEGroupedNorm, x: torch.Tensor
-) -> torch.Tensor:
+def torch_impl_ple_norm(norm: Qwen4ExpPLEGroupedNorm, x: torch.Tensor) -> torch.Tensor:
     """Eager baseline (the production unfused fp32 chain)."""
     x_float = x.float()
     group_shape = x_float.shape[:-1] + (-1, norm.group_size)
diff --git a/test/registered/kernels/ops/attention/test_qkvzba_split_ratio3.py b/test/registered/kernels/ops/attention/test_qkvzba_split_ratio3.py
index f4513b915492..627d605ae07b 100644
--- a/test/registered/kernels/ops/attention/test_qkvzba_split_ratio3.py
+++ b/test/registered/kernels/ops/attention/test_qkvzba_split_ratio3.py
@@ -1,4 +1,5 @@
 import sys
+
 import pytest
 import torch
 
diff --git a/test/registered/kernels/ops/attention/test_qsa_indexer.py b/test/registered/kernels/ops/attention/test_qsa_indexer.py
index 7c50b1ca16e4..41d263204681 100644
--- a/test/registered/kernels/ops/attention/test_qsa_indexer.py
+++ b/test/registered/kernels/ops/attention/test_qsa_indexer.py
@@ -11,6 +11,7 @@
 
 import pytest
 import torch
+
 from sglang.test.ci.ci_register import register_cuda_ci
 
 register_cuda_ci(est_time=120, stage="base-b-kernel-unit", runner_config="1-gpu-large")
@@ -89,9 +90,7 @@ class FakePool:
     """Minimal stand-in for the QSA KV pool buffers used by the indexer."""
 
     def __init__(self, num_slots, num_compressed, device, dtype=torch.bfloat16):
-        self.key_state = torch.zeros(
-            num_slots, 1, HEAD_DIM, dtype=dtype, device=device
-        )
+        self.key_state = torch.zeros(num_slots, 1, HEAD_DIM, dtype=dtype, device=device)
         self.qsa_rope_position_buffer = torch.zeros(
             num_slots, 3, dtype=torch.int64, device=device
         )
@@ -186,9 +185,9 @@ def _run_case(
         positions = logical_positions
 
     # Distinct state slots per token; slot 0 is deliberately unused.
-    cache_loc = torch.randperm(max(num_tokens + 8, 4096), device=device)[
-        :num_tokens
-    ].long() + 1
+    cache_loc = (
+        torch.randperm(max(num_tokens + 8, 4096), device=device)[:num_tokens].long() + 1
+    )
     token_slot_table = torch.zeros(1, 65536, dtype=torch.int32, device=device)
     token_slot_table[0, logical_positions.long()] = cache_loc.to(torch.int32)
 
@@ -202,9 +201,7 @@ def _run_case(
     pool_new = FakePool(8192, 4096, device, dtype)
 
     # Reference: pre-fusion eager path (toggle the real switches off).
-    q_ref, token_k_ref, stored_ref = _force_eager(indexer).project_qk(
-        hidden, positions
-    )
+    q_ref, token_k_ref, stored_ref = _force_eager(indexer).project_qk(hidden, positions)
     assert not stored_ref
     indexer.update_key_state_and_compress(
         token_k_ref,
@@ -363,8 +360,8 @@ def test_decode_selection_equivalent():
     scores are fp32 sums of 128-dim dot products, so a 1-ulp change in one
     component only matters on exact ties.
     """
-    from sglang.srt.layers.attention.qsa.mqa import torch_qsa_mqa_decode
     from sglang.srt.layers.attention.qsa.kernel import qsa_fast_topk
+    from sglang.srt.layers.attention.qsa.mqa import torch_qsa_mqa_decode
 
     device = torch.device("cuda")
     dtype = torch.bfloat16
diff --git a/test/registered/kernels/ops/elementwise/test_fast_topk.py b/test/registered/kernels/ops/elementwise/test_fast_topk.py
index 9814b984101e..360ad6f50e99 100644
--- a/test/registered/kernels/ops/elementwise/test_fast_topk.py
+++ b/test/registered/kernels/ops/elementwise/test_fast_topk.py
@@ -21,7 +21,9 @@ def _check_topk_values(score, lengths, indices, topk, row_starts):
         row = indices[b]
         if length <= topk:
             # naive path: identity indices, then -1 fill
-            assert torch.equal(row[:length].cpu(), torch.arange(length, dtype=torch.int32))
+            assert torch.equal(
+                row[:length].cpu(), torch.arange(length, dtype=torch.int32)
+            )
             assert (row[length:] == -1).all()
             continue
         assert (row >= 0).all(), "long rows must fill every slot"
diff --git a/test/registered/kernels/ops/elementwise/test_hc_combine.py b/test/registered/kernels/ops/elementwise/test_hc_combine.py
index e9a4f735aaec..7cf54125e7ea 100644
--- a/test/registered/kernels/ops/elementwise/test_hc_combine.py
+++ b/test/registered/kernels/ops/elementwise/test_hc_combine.py
@@ -36,7 +36,9 @@ def _reference_hc_combine(
     return (R + injection).flatten(-2)
 
 
-def _make_inputs(num_tokens: int, dtype: torch.dtype, hc: int = HC_COUNT, hs: int = HIDDEN_SIZE):
+def _make_inputs(
+    num_tokens: int, dtype: torch.dtype, hc: int = HC_COUNT, hs: int = HIDDEN_SIZE
+):
     torch.manual_seed(0)
     block_output = torch.randn(num_tokens, hs, dtype=dtype, device="cuda")
     residual = torch.randn(num_tokens, hc * hs, dtype=dtype, device="cuda")
@@ -191,7 +193,9 @@ def test_hc_combine_dtype_mismatch():
 def test_hc_combine_shape_mismatch():
     dtype = torch.bfloat16
     block_output, residual, normed_residual, inject_weight = _make_inputs(4, dtype)
-    bad_weight = torch.randn(HC_COUNT, HC_COUNT * HIDDEN_SIZE + 2048, dtype=dtype, device="cuda")
+    bad_weight = torch.randn(
+        HC_COUNT, HC_COUNT * HIDDEN_SIZE + 2048, dtype=dtype, device="cuda"
+    )
     with pytest.raises(RuntimeError):
         hc_combine(
             block_output,
diff --git a/test/registered/kernels/ops/layernorm/test_qwen4_ple_norm.py b/test/registered/kernels/ops/layernorm/test_qwen4_ple_norm.py
index cf84e0d1cd90..71d4bfbc800c 100644
--- a/test/registered/kernels/ops/layernorm/test_qwen4_ple_norm.py
+++ b/test/registered/kernels/ops/layernorm/test_qwen4_ple_norm.py
@@ -106,7 +106,9 @@ def test_qwen4_ple_norm_eager_fallback():
 
 def test_qwen4_ple_norm_non_cuda_fallback():
     torch.manual_seed(0)
-    norm = Qwen4ExpPLEGroupedNorm(_PLE_HIDDEN_SIZE, eps=1e-6, group_size=_PLE_GROUP_SIZE)
+    norm = Qwen4ExpPLEGroupedNorm(
+        _PLE_HIDDEN_SIZE, eps=1e-6, group_size=_PLE_GROUP_SIZE
+    )
     x = torch.randn(4, _PLE_HIDDEN_SIZE, dtype=torch.bfloat16)
 
     out = norm(x)
diff --git a/test/registered/kernels/test_hc_mix_triton.py b/test/registered/kernels/test_hc_mix_triton.py
index b0c9a676585d..601d0e259da2 100644
--- a/test/registered/kernels/test_hc_mix_triton.py
+++ b/test/registered/kernels/test_hc_mix_triton.py
@@ -1,4 +1,5 @@
 import sys
+
 import pytest
 import torch
 import torch.nn.functional as F
@@ -37,12 +38,10 @@ def _make_inputs(num_tokens: int, dtype: torch.dtype):
     torch.manual_seed(0)
     x = torch.randn(num_tokens, HC_COUNT * HIDDEN_SIZE, dtype=dtype, device="cuda")
     w_down = (
-        torch.randn(LOWRANK, HC_COUNT * HIDDEN_SIZE, dtype=dtype, device="cuda")
-        * 0.02
+        torch.randn(LOWRANK, HC_COUNT * HIDDEN_SIZE, dtype=dtype, device="cuda") * 0.02
     )
     w_up = (
-        torch.randn(HC_COUNT * HIDDEN_SIZE, LOWRANK, dtype=dtype, device="cuda")
-        * 0.02
+        torch.randn(HC_COUNT * HIDDEN_SIZE, LOWRANK, dtype=dtype, device="cuda") * 0.02
     )
     return x, w_down, w_up
 
@@ -60,9 +59,7 @@ def test_fused_hc_mix_matches_reference(dtype, num_tokens):
     assert fused_hc_mix_supported(x, w_down, w_up)
     out = fused_hc_mix(x, w_down, w_up, HC_COUNT, HIDDEN_SIZE)
     ref = _reference_mix(x, w_down, w_up, HC_COUNT, HIDDEN_SIZE)
-    torch.testing.assert_close(
-        out.to(torch.float64), ref, **_TOLERANCES[dtype]
-    )
+    torch.testing.assert_close(out.to(torch.float64), ref, **_TOLERANCES[dtype])
 
 
 def test_fused_hc_mix_no_less_accurate_than_eager():
diff --git a/test/registered/kernels/test_qsa.py b/test/registered/kernels/test_qsa.py
index b9cbed29bc97..eab9ec3d2878 100644
--- a/test/registered/kernels/test_qsa.py
+++ b/test/registered/kernels/test_qsa.py
@@ -5,34 +5,32 @@
 import torch
 
 from sglang.srt.configs.qwen4_exp import Qwen4ExpConfig
+from sglang.srt.layers.attention import qwen_sparse_attn_backend as qsa_backend_module
 from sglang.srt.layers.attention.attention_registry import ATTENTION_BACKENDS
-from sglang.srt.layers.attention.qsa import qsa_indexer as qsa_indexer_module
 from sglang.srt.layers.attention.qsa import dsa_indexer as dsa_indexer_module
-from sglang.srt.layers.attention.qsa.metadata import QSAIndexerMetadata
+from sglang.srt.layers.attention.qsa import qsa_indexer as qsa_indexer_module
 from sglang.srt.layers.attention.qsa.kernel import (
     average_pool_qsa_keys,
     expand_qsa_block_indices,
-    torch_expand_qsa_block_indices,
-    triton_expand_qsa_block_indices,
     qsa_fast_topk,
     qsa_sparse_attention,
+    torch_expand_qsa_block_indices,
+    triton_expand_qsa_block_indices,
+)
+from sglang.srt.layers.attention.qsa.metadata import (
+    QSAIndexerMetadata,
+    build_qsa_row_ranges,
 )
 from sglang.srt.layers.attention.qsa.mqa import (
     qsa_mqa_decode,
     qsa_mqa_prefill,
 )
-from sglang.srt.layers.attention.qsa.metadata import build_qsa_row_ranges
 from sglang.srt.layers.attention.qsa.qsa_indexer import QSAIndexer
-from sglang.srt.layers.attention import qwen_sparse_attn_backend as qsa_backend_module
 from sglang.srt.layers.attention.qwen_sparse_attn_backend import (
     QwenSparseAttnBackend,
     QwenSparseMultiStepDraftBackend,
 )
-from sglang.srt.mem_cache.allocator.token import TokenToKVPoolAllocator
-from sglang.srt.mem_cache.memory_pool import ReqToTokenPool
-from sglang.srt.mem_cache.qsa_kv_pool import QSATokenToKVPool
 from sglang.srt.model_executor.forward_batch_info import ForwardMode
-
 from sglang.test.ci.ci_register import register_cuda_ci
 
 register_cuda_ci(est_time=60, stage="base-b-kernel-unit", runner_config="1-gpu-large")
@@ -210,7 +208,9 @@ def test_qsa_glue_builds_indexer_per_variant(monkeypatch):
     recorded = {}
 
     class _FakeIndexer:
-        def __init__(self, config, layer_id, quant_config=None, prefix="", rotary_emb=None):
+        def __init__(
+            self, config, layer_id, quant_config=None, prefix="", rotary_emb=None
+        ):
             recorded.update(
                 config=config,
                 layer_id=layer_id,
@@ -240,9 +240,7 @@ def __init__(self, config, layer_id, quant_config=None, prefix="", **kw):
                 dsa_prefix=prefix,
             )
 
-    monkeypatch.setattr(
-        dsa_indexer_module, "QwenDSAIndexer", _FakeDSAIndexer
-    )
+    monkeypatch.setattr(dsa_indexer_module, "QwenDSAIndexer", _FakeDSAIndexer)
     dsa_indexer = build_qsa_indexer(
         _tokenwise_config_namespace(), layer_id=2, prefix="q"
     )
@@ -553,9 +551,7 @@ def test_qwen_dsa_select_prefill_respects_causal_windows():
     metadata = SimpleNamespace(
         sequence_lengths=torch.tensor([4, 5], dtype=torch.int32),
         token_slot_table=token_slot_table,
-        token_to_batch_idx=torch.tensor(
-            [0, 0, 0, 1, 1, 1, 1], dtype=torch.int32
-        ),
+        token_to_batch_idx=torch.tensor([0, 0, 0, 1, 1, 1, 1], dtype=torch.int32),
         token_to_kv_pool=pool,
     )
     indexer = _make_dsa_indexer_stub(topk=4)
@@ -563,12 +559,8 @@ def test_qwen_dsa_select_prefill_respects_causal_windows():
     q[..., 1:] = 0
     w = torch.ones(7, 2, dtype=torch.bfloat16)
     # seq 0 extends positions 1..3, seq 1 extends positions 1..4.
-    logical_positions = torch.tensor(
-        [1, 2, 3, 1, 2, 3, 4], dtype=torch.int64
-    )
-    out = QwenDSAIndexer._select_prefill(
-        indexer, q, w, logical_positions, metadata
-    )
+    logical_positions = torch.tensor([1, 2, 3, 1, 2, 3, 4], dtype=torch.int64)
+    out = QwenDSAIndexer._select_prefill(indexer, q, w, logical_positions, metadata)
     assert out.shape == (7, 4)
     # A row may never select a column beyond its causal end — in particular
     # seq 1 rows before position 4 must not see the hot col 4.
@@ -659,9 +651,7 @@ def test_qsa_idle_skips_metadata_construction():
     backend = QwenSparseAttnBackend.__new__(QwenSparseAttnBackend)
     backend.forward_metadata = object()
 
-    backend.init_forward_metadata(
-        SimpleNamespace(forward_mode=ForwardMode.IDLE)
-    )
+    backend.init_forward_metadata(SimpleNamespace(forward_mode=ForwardMode.IDLE))
 
     assert backend.forward_metadata is None
 
@@ -759,9 +749,7 @@ def test_qsa_target_verify_rejects_branching_speculation():
 
 
 def test_qsa_mtp_cuda_graph_padding_stays_below_compression_boundary():
-    backend, forward_batch, _ = _make_mtp_draft_batch(
-        steps=4, seq_lens=(3, 1)
-    )
+    backend, forward_batch, _ = _make_mtp_draft_batch(steps=4, seq_lens=(3, 1))
 
     class Recorder:
         def __init__(self):
@@ -1117,9 +1105,7 @@ def test_qsa_idle_metadata_builds_empty_rows():
     assert step_metadata.row_req_pool_indices.numel() == 0
     # Per-step out_cache_loc slicing must stay empty without allocating rows.
     for attn_backend in draft.attn_backends:
-        assert (
-            attn_backend.forward_metadata.indexer_metadata.out_cache_loc.numel() == 0
-        )
+        assert attn_backend.forward_metadata.indexer_metadata.out_cache_loc.numel() == 0
 
 
 def test_qsa_decode_requires_one_query_row_per_request():
@@ -1240,9 +1226,7 @@ def test_qsa_speculative_pseudo_extend_is_rejected():
         except ValueError as exc:
             assert "pseudo-extend" in str(exc)
         else:
-            raise AssertionError(
-                f"QSA must reject pseudo-extend of {original_mode}"
-            )
+            raise AssertionError(f"QSA must reject pseudo-extend of {original_mode}")
 
 
 def test_qsa_indexer_requires_compress_ratio_above_one():
@@ -1565,9 +1549,7 @@ def test_qsa_prefill_selection_microchunks_rows(monkeypatch):
     starts = torch.zeros(rows, dtype=torch.int32)
     ends = torch.full((rows,), keys, dtype=torch.int32)
     positions = torch.full((rows,), keys * compress_ratio - 1, dtype=torch.long)
-    sequence_lengths = torch.full(
-        (rows,), keys * compress_ratio, dtype=torch.int32
-    )
+    sequence_lengths = torch.full((rows,), keys * compress_ratio, dtype=torch.int32)
 
     monkeypatch.setattr(
         qsa_indexer_module,
@@ -1645,9 +1627,7 @@ def test_qsa_reranks_wider_candidate_set():
         _rerank_qsa_topk_candidates,
     )
 
-    logits = torch.tensor(
-        [[0.0, 9.0, 8.0, 7.0, 6.0, 10.0]], dtype=torch.float32
-    )
+    logits = torch.tensor([[0.0, 9.0, 8.0, 7.0, 6.0, 10.0]], dtype=torch.float32)
     starts = torch.tensor([1], dtype=torch.int32)
     candidates = torch.tensor([[0, 1, 2, 3, 4]], dtype=torch.int32)
 
@@ -1768,9 +1748,7 @@ def test_qsa_mtp_step_out_cache_loc_matches_draft_forward_layout():
     backend.topk, backend.speculative_num_steps = 1, 3
     bs, topk, steps = 4, 1, 3
     flat = torch.arange(bs * topk * steps, dtype=torch.int64)
-    fb = SimpleNamespace(
-        out_cache_loc=flat, batch_size=bs, seq_lens=torch.ones(bs)
-    )
+    fb = SimpleNamespace(out_cache_loc=flat, batch_size=bs, seq_lens=torch.ones(bs))
     # Reference: the exact draft_forward expression chain.
     reference = flat.reshape(bs, topk, steps).permute(2, 0, 1).reshape(steps, -1)
     for step in range(steps):
@@ -1892,9 +1870,7 @@ def make_metadata():
         # Path 2: legacy host refresh over the layout the kernels produced.
         host_metadata = make_metadata()
         host_metadata.sequence_lengths.copy_(kernel_metadata.sequence_lengths)
-        host_metadata.row_req_pool_indices.copy_(
-            kernel_metadata.row_req_pool_indices
-        )
+        host_metadata.row_req_pool_indices.copy_(kernel_metadata.row_req_pool_indices)
         backend._update_qsa_cuda_graph_metadata(
             host_metadata.indexer_metadata, host_metadata.row_req_pool_indices
         )
@@ -1940,8 +1916,10 @@ def _qsa_expected_graph_layout(
             row_prefix.append(max(base - 1, 0))
             row_reqs.append(req_pool[pid])
             continue
-        eff = (extend_len if pid < real_reqs else 0) if mode == 1 else (
-            extend_lens[pid] if pid < real_reqs else 0
+        eff = (
+            (extend_len if pid < real_reqs else 0)
+            if mode == 1
+            else (extend_lens[pid] if pid < real_reqs else 0)
         )
         prefix, limit = (base, base + eff) if mode == 1 else (max(base - eff, 0), base)
         for j in range(eff):
@@ -1988,7 +1966,9 @@ def run_case(mode, bs, num_rows, seq_lens, extend_lens, extend_len, num_padding)
         indexer = QSAIndexerMetadata(
             sequence_lengths=torch.zeros(num_rows, dtype=torch.int32, device=device),
             token_to_batch_idx=torch.arange(num_rows, dtype=torch.int32, device=device),
-            token_slot_table=torch.zeros((num_rows, 1), dtype=torch.int32, device=device),
+            token_slot_table=torch.zeros(
+                (num_rows, 1), dtype=torch.int32, device=device
+            ),
             out_cache_loc=torch.zeros(num_rows, dtype=torch.int64, device=device),
             token_to_kv_pool=pool,
             compress_ratio=ratio,
@@ -2002,7 +1982,9 @@ def run_case(mode, bs, num_rows, seq_lens, extend_lens, extend_len, num_padding)
             graph_compressed_lengths=torch.zeros(
                 num_rows, dtype=torch.int32, device=device
             ),
-            graph_prefix_lengths=torch.zeros(num_rows, dtype=torch.int32, device=device),
+            graph_prefix_lengths=torch.zeros(
+                num_rows, dtype=torch.int32, device=device
+            ),
             decode_logical_positions=torch.zeros(
                 num_rows, dtype=torch.int32, device=device
             ),
@@ -2016,7 +1998,9 @@ def run_case(mode, bs, num_rows, seq_lens, extend_lens, extend_len, num_padding)
             token_to_batch_idx=indexer.token_to_batch_idx,
             token_slot_table=indexer.token_slot_table,
             indexer_metadata=indexer,
-            row_req_pool_indices=torch.zeros(num_rows, dtype=torch.int32, device=device),
+            row_req_pool_indices=torch.zeros(
+                num_rows, dtype=torch.int32, device=device
+            ),
             is_cuda_graph=True,
         )
         launch_graph_metadata(
@@ -2063,18 +2047,33 @@ def run_case(mode, bs, num_rows, seq_lens, extend_lens, extend_len, num_padding)
 
     # Target verify: uniform 4-token window, 2 padded request slots.
     run_case(
-        mode=1, bs=6, num_rows=32, seq_lens=[254, 255, 256, 300, 7, 7],
-        extend_lens=None, extend_len=4, num_padding=2,
+        mode=1,
+        bs=6,
+        num_rows=32,
+        seq_lens=[254, 255, 256, 300, 7, 7],
+        extend_lens=None,
+        extend_len=4,
+        num_padding=2,
     )
     # Draft extend: per-request extend lengths (accept-count dependent), padded.
     run_case(
-        mode=2, bs=5, num_rows=24, seq_lens=[260, 512, 257, 9, 9],
-        extend_lens=[3, 1, 4, 0, 0], extend_len=0, num_padding=2,
+        mode=2,
+        bs=5,
+        num_rows=24,
+        seq_lens=[260, 512, 257, 9, 9],
+        extend_lens=[3, 1, 4, 0, 0],
+        extend_len=0,
+        num_padding=2,
     )
     # Decode with a padded tail (dummy rows alias request slot 0).
     run_case(
-        mode=0, bs=4, num_rows=4, seq_lens=[256, 1024, 1025, 4096],
-        extend_lens=None, extend_len=0, num_padding=0,
+        mode=0,
+        bs=4,
+        num_rows=4,
+        seq_lens=[256, 1024, 1025, 4096],
+        extend_lens=None,
+        extend_len=0,
+        num_padding=0,
     )
 
 
diff --git a/test/registered/kernels/test_verify_commit_triton.py b/test/registered/kernels/test_verify_commit_triton.py
index 2545b308e945..41c61a07ca2c 100644
--- a/test/registered/kernels/test_verify_commit_triton.py
+++ b/test/registered/kernels/test_verify_commit_triton.py
@@ -1,4 +1,5 @@
 import sys
+
 import pytest
 import torch
 
diff --git a/test/registered/unit/mem_cache/test_paged_free_segment.py b/test/registered/unit/mem_cache/test_paged_free_segment.py
index b81faf83030d..b2630cb8e134 100644
--- a/test/registered/unit/mem_cache/test_paged_free_segment.py
+++ b/test/registered/unit/mem_cache/test_paged_free_segment.py
@@ -217,7 +217,6 @@ def free(self, free_index: torch.Tensor):
         self.freed.append(free_index)
 
 
-
 class TestBaseFallbackFreeSegments(unittest.TestCase):
     def test_trim_dedups_boundary_page_before_fallback_free(self):
         # fallback allocators (UnifiedMamba/SWA) dedup per free() call at best;