diff --git a/python/cudnn/sdpa/fwd/config_sm100.py b/python/cudnn/sdpa/fwd/config_sm100.py
index f69331962..a01a4fa28 100644
--- a/python/cudnn/sdpa/fwd/config_sm100.py
+++ b/python/cudnn/sdpa/fwd/config_sm100.py
@@ -88,12 +88,39 @@ class TemplateParams:
     seq_q_lens_present: bool = False
     sched_policy: int = SCHED_NATURAL
     thd_varlen: bool = False
+    # KV split: each Q tile's KV loop range is cut into ``split_kv`` contiguous
+    # chunks, each run as its own persistent tile writing a partial (O, LSE)
+    # that kernels/split_combine_sm100.py reduces.  1 = off (byte-identical
+    # codegen to the single-pass kernel).
+    split_kv: int = 1
+    # MMA cluster width: 2 = cga2 collective tcgen05.mma.cta_group::2 (a CTA
+    # pair share one MMA, each holding half of every K/V tile); 1 = cga1, one
+    # independent CTA per tile.
+    #
+    # cga1 has no collective MMA to halve per-CTA K/V, so at a fixed STAGES_KV
+    # it doubles that footprint.  d128 buys the 64 KiB back by aliasing Q and O
+    # into one slab (make_cfg_d128 turns QO_ALIAS on for cga1; see
+    # _validate_cfg_d128's SMEM check); the fp8 family instead scales the stage
+    # count with the width so stages x per-CTA-buffer stays constant.
+    cta_mma: int = 2
     # cc10.3+ fuses the S_acc row-max into the LDTM (tcgen05.ld.red.f32.max); cc10.0
     # lacks it and uses the manual load + software reduction. Auto-set from the device
     # capability at compile time (MXFP8 only; the f16/fp8 kernels do not read it).
     fused_ldtm_stat: bool = False
 
 
+# split_kv / cta_mma live on the TemplateParams shared by every SM100 flavor, but
+# a flavor only honours them once its make_cfg_* threads them into a Cfg AND its
+# kernel reads them.  Accepting them elsewhere would silently ignore them — and
+# for split_kv that is not merely surprising but WRONG: the caller sizes an
+# (S*B)-batch partial workspace and runs the combine, while the kernel keeps
+# writing only slots [0, B).  The untouched slots keep lse_partial = 0 rather
+# than -inf, so they carry weight exp(0 - M) != 0 through the log-sum-exp and
+# corrupt the result instead of dropping out.  Grow these sets as flavors land.
+_SPLIT_KV_FLAVORS = frozenset({"d128", "d192", "d256", "d512"})
+_CTA_MMA_FLAVORS = frozenset({"d128", "d192"})
+
+
 def _validate_params(flavor: str, k: TemplateParams) -> None:
     if k.dtype_qkv not in (DTYPE_E4M3, DTYPE_E5M2, DTYPE_BF16, DTYPE_FP16):
         raise ValueError(f"{flavor}: DTYPE_QKV must be E4M3/E5M2/BF16/FP16 (0..3); got {k.dtype_qkv}")
@@ -121,6 +148,31 @@ def _validate_params(flavor: str, k: TemplateParams) -> None:
             raise ValueError(f"{flavor}: SEQ_Q_LENS_PRESENT requires SEQ_KV_LENS_PRESENT (padding mask)")
     if k.sched_policy not in (SCHED_NATURAL, SCHED_LPT, SCHED_LPT_L2):
         raise ValueError(f"{flavor}: only SCHED_NATURAL (0) / SCHED_LPT (1) / SCHED_LPT_L2 (2) are wired up; got {k.sched_policy}")
+    if k.cta_mma not in (1, 2):
+        raise ValueError(f"{flavor}: cta_mma must be 1 (cga1) or 2 (cga2); got {k.cta_mma}")
+    # split_kv / cta_mma live on the TemplateParams shared by every SM100 flavor,
+    # but only make_cfg_d128 threads them into a Cfg and only the d128 kernel
+    # reads them.  Accepting them elsewhere would silently ignore them — and for
+    # split_kv that is not merely surprising but WRONG: the caller sizes an
+    # (S*B)-batch partial workspace and runs the combine, while the kernel keeps
+    # writing only slots [0, B).  The untouched slots keep lse_partial = 0 rather
+    # than -inf, so they carry weight exp(0 - M) != 0 through the log-sum-exp and
+    # corrupt the result instead of dropping out.  Reject at the door.
+    if k.split_kv != 1 and flavor not in _SPLIT_KV_FLAVORS:
+        raise ValueError(f"{flavor}: split_kv is not implemented on this flavor (got {k.split_kv}); supported: {sorted(_SPLIT_KV_FLAVORS)}")
+    if k.cta_mma != 2 and flavor not in _CTA_MMA_FLAVORS:
+        raise ValueError(f"{flavor}: cta_mma is not selectable on this flavor (got {k.cta_mma}); supported: {sorted(_CTA_MMA_FLAVORS)}")
+    if k.split_kv < 1:
+        raise ValueError(f"{flavor}: split_kv must be >= 1 (1 = KV-split off); got {k.split_kv}")
+    if k.split_kv > 1:
+        # Each of these would need extra machinery in the combine pass, so the
+        # backstop rejects them rather than silently producing a wrong answer.
+        if k.thd_varlen:
+            raise ValueError(f"{flavor}: split_kv > 1 is dense-only (THD packs its own flat grid)")
+        if k.has_sink:
+            # The sink logit is folded into the softmax denominator in the
+            # per-tile epilogue, so every split would add its own copy of it.
+            raise ValueError(f"{flavor}: split_kv > 1 with attention sink is not supported (the sink would be counted once per split)")
 
 
 def _mask_flags_from(params: TemplateParams) -> int:
@@ -290,6 +342,9 @@ class CfgD256:
 
     THD_VARLEN: int = 0
 
+    # KV split; 1 = off.  See TemplateParams.split_kv.
+    SPLIT_KV: int = 1
+
 
 def _validate_cfg_d256(cfg: CfgD256) -> None:
     """Consistency checks on the (mostly hardcoded) d256 geometry."""
@@ -335,6 +390,7 @@ def make_cfg_d256(params: TemplateParams) -> Tuple[CfgD256, TmaIters]:
         SEQ_KV_LENS_PRESENT=1 if (params.thd_varlen or params.seq_kv_lens_present) else 0,
         SEQ_Q_LENS_PRESENT=int(params.seq_q_lens_present),
         THD_VARLEN=int(params.thd_varlen),
+        SPLIT_KV=int(params.split_kv),
     )
     _validate_cfg_d256(cfg)
     return cfg, _tma_iters(cfg)
@@ -429,6 +485,9 @@ class CfgD512:
 
     THD_VARLEN: int = 0
 
+    # KV split; 1 = off.  See TemplateParams.split_kv.
+    SPLIT_KV: int = 1
+
 
 def _validate_cfg_d512(cfg: CfgD512) -> None:
     """Consistency checks on the (mostly hardcoded) d512 geometry."""
@@ -478,6 +537,7 @@ def make_cfg_d512(params: TemplateParams) -> Tuple[CfgD512, TmaIters]:
         SEQ_KV_LENS_PRESENT=1 if (params.thd_varlen or params.seq_kv_lens_present) else 0,
         SEQ_Q_LENS_PRESENT=int(params.seq_q_lens_present),
         THD_VARLEN=int(params.thd_varlen),
+        SPLIT_KV=int(params.split_kv),
     )
     _validate_cfg_d512(cfg)
     return cfg, _tma_iters(cfg)
@@ -575,6 +635,34 @@ class CfgD128:
 
     THD_VARLEN: int = 0
 
+    # KV split; 1 = off.  See TemplateParams.split_kv.
+    SPLIT_KV: int = 1
+
+
+# Blackwell SM100 per-CTA dynamic SMEM cap (228 KiB physical, 227 KiB usable).
+_SM100_MAX_DYN_SMEM = 227 * 1024
+
+
+def _d128_smem_bytes(cfg) -> int:
+    """Data-buffer SMEM for the d128 pipeline (barriers/TMEM ptr are noise).
+
+    Q and O are TILES_Q slabs each; under QO_ALIAS they share one slab sized to
+    the larger.  K/V are STAGES_KV buffers each, and their PER-CTA size is
+    divided by CTA_MMA because the cga2 collective MMA lets a CTA pair hold half
+    of every tile.  That divisor is exactly what cga1 gives up, which is why
+    cga1 needs the Q/O alias to break even:
+
+        cga2, no alias : 64(Q) + 64(O) + 32(K) + 32(V) = 192 KiB
+        cga1, no alias : 64(Q) + 64(O) + 64(K) + 64(V) = 256 KiB  (over cap)
+        cga1, alias    : 64(Q u O)     + 64(K) + 64(V) = 192 KiB
+    """
+    q_slab = cfg.TILE_M * cfg.TILE_K * cfg.BPE
+    o_slab = cfg.TILE_M * cfg.TILE_O * cfg.BPE_O
+    qo = cfg.TILES_Q * (max(q_slab, o_slab) if cfg.QO_ALIAS else q_slab + o_slab)
+    k = cfg.STAGES_KV * (cfg.TILE_N * cfg.TILE_K * cfg.BPE // cfg.CTA_MMA)
+    v = cfg.STAGES_KV * (cfg.TILE_O * cfg.TILE_N * cfg.BPE // cfg.CTA_MMA)
+    return qo + k + v
+
 
 def _validate_cfg_d128(cfg: CfgD128) -> None:
     """Consistency checks on the (mostly hardcoded) d128 (llama) geometry."""
@@ -583,14 +671,27 @@ def _validate_cfg_d128(cfg: CfgD128) -> None:
         (cfg.MMA_REGS == cfg.TMALDG_REGS == cfg.TMASTG_REGS == cfg.SCHEDULER_REGS, "d128: MMA/TMALDG/TMASTG/SCHEDULER regs must match"),
         (cfg.MMA_REGS + cfg.CORRECTION_REGS + cfg.SOFTMAX_WARPGROUPS * cfg.SOFTMAX_REGS <= 512, "d128: register budget over 512"),
         (cfg.MMA_REGS % 8 == 0 and cfg.CORRECTION_REGS % 8 == 0 and cfg.SOFTMAX_REGS % 8 == 0, "d128: per-role regs must be multiples of 8"),
-        (cfg.CGA_M == 2 and cfg.CTA_MMA == 2, "d128 SM100 is cga2-only (CGA_M == CTA_MMA == 2)"),
+        (cfg.CGA_M == cfg.CTA_MMA and cfg.CTA_MMA in (1, 2), "d128 SM100: CGA_M must equal CTA_MMA, and CTA_MMA must be 1 (cga1) or 2 (cga2)"),
+        (
+            cfg.QO_ALIAS == 1 if cfg.CTA_MMA == 1 else True,
+            "d128 cga1: QO_ALIAS is mandatory — cga1 doubles per-CTA K/V (no collective MMA to halve it), "
+            "so Q and O must share one slab to stay inside the SMEM cap",
+        ),
+        (
+            _d128_smem_bytes(cfg) <= _SM100_MAX_DYN_SMEM,
+            f"d128: SMEM {_d128_smem_bytes(cfg) // 1024} KiB over the SM100 {_SM100_MAX_DYN_SMEM // 1024} KiB per-CTA cap",
+        ),
         (cfg.TILE_K == 128 and cfg.TILE_O == 128, "d128: d_qk = d_v = 128"),
         (cfg.TILES_Q == 2, "d128 (llama): TILES_Q must be 2"),
         (cfg.SOFTMAX_WARPGROUPS == 2, "d128 (llama): SOFTMAX_WARPGROUPS must be 2"),
         (cfg.CORRECTION_WARPS == 4, "d128 (llama): CORRECTION_WARPS must be 4"),
         (cfg.TOTAL_WARPS == 16 and cfg.THREADS_PER_CTA == 512, "d128 (llama): 16 warps / 512 threads"),
         (cfg.READ_TILE_ARRIVERS == 15, f"d128 llama: expected READ_TILE_ARRIVERS=15, got {cfg.READ_TILE_ARRIVERS}"),
-        (cfg.STAGES_KV == (4 if _fp8 else 2), "d128 SM100: STAGES_KV must be 4 (fp8/mxfp8) / 2 (f16, 192 KiB SMEM budget)"),
+        (
+            cfg.STAGES_KV == ((2 if cfg.CTA_MMA == 1 else 4) if _fp8 else 2),
+            "d128 SM100: STAGES_KV must be 2 (f16/bf16) or, for fp8/mxfp8, 4 at cga2 and 2 at cga1 — "
+            "the stage depth scales with the cluster width so stages x per-CTA-buffer stays constant",
+        ),
         (
             cfg.TILE_K_HW_BMM1 == (32 if _fp8 else 16) and cfg.TILE_K_HW_BMM2 == (32 if _fp8 else 16),
             "d128: TILE_K_HW must be 32 (fp8/mxfp8 K=32 QMMA) / 16 (f16, 1-chunk on SM10x)",
@@ -621,14 +722,27 @@ def make_cfg_d128(params: TemplateParams) -> Tuple[CfgD128, TmaIters]:
         DTYPE_O=dtype_o,
         BPE=b,
         BPE_O=b_o,
+        CGA_M=params.cta_mma,
+        CTA_MMA=params.cta_mma,
+        # cga1 has no collective MMA to halve per-CTA K/V, so Q and O must share
+        # one slab to stay under the SMEM cap (_validate_cfg_d128 enforces it).
+        QO_ALIAS=1 if params.cta_mma == 1 else 0,
         Q_SWZ_BYTES=q_swz_bytes(128, b),
         K_SWZ_BYTES=q_swz_bytes(128, b),
-        V_SWZ_BYTES=v_swz_bytes(128, 2, b),
+        V_SWZ_BYTES=v_swz_bytes(128, params.cta_mma, b),
         O_SWZ_BYTES=o_swz_bytes(128, b_o),
         RESCALE_THRESHOLD=rescale_threshold(params.dtype_qkv),
         TILE_K_HW_BMM1=tile_k_hw_fp8,
         TILE_K_HW_BMM2=tile_k_hw_fp8,
-        STAGES_KV=4 if fp8 else 2,
+        # KV stage depth scales with the cluster width, as in cuDNN's own
+        # kernels (stages_kv = N * CTA_MMA): cga1 has no collective MMA to halve
+        # per-CTA K/V, so the stage count halves instead to keep the product --
+        # and hence the SMEM -- constant.  Only the fp8 family needs this here:
+        # f16/bf16 already fit at cga1 by aliasing Q and O, and their verified
+        # cga1 configuration keeps STAGES_KV=2.  mxfp8 additionally stages E8M0
+        # scale factors that the SMEM model cannot see, and at STAGES_KV=4 that
+        # pushed a cga1 CTA to 237024 B against the 232448 B cap.
+        STAGES_KV=(2 if params.cta_mma == 1 else 4) if fp8 else 2,
         MASK_FLAGS=_mask_flags_from(params),
         WINDOW_LEFT=params.window_left or 0,
         WINDOW_RIGHT=params.window_right or 0,
@@ -638,6 +752,7 @@ def make_cfg_d128(params: TemplateParams) -> Tuple[CfgD128, TmaIters]:
         SEQ_KV_LENS_PRESENT=1 if (params.thd_varlen or params.seq_kv_lens_present) else 0,
         SEQ_Q_LENS_PRESENT=int(params.seq_q_lens_present),
         THD_VARLEN=int(params.thd_varlen),
+        SPLIT_KV=int(params.split_kv),
     )
     _validate_cfg_d128(cfg)
     return cfg, _tma_iters(cfg)
@@ -657,6 +772,24 @@ class CfgD192(CfgD128):
     CORRECTION_REGS: int = 88
 
 
+def _d192_smem_bytes(cfg) -> int:
+    """Data-buffer SMEM for the d192 pipeline (Q/O always aliased).
+
+    Same shape as _d128_smem_bytes, but d_qk = 192 makes the Q and K slabs 1.5x
+    the d128 ones, which is why this flavor needs a shallower KV pipeline:
+
+        cga2, STAGES_KV=2 : 96(Q u O) + 48(K) + 32(V) = 176 KiB
+        cga1, STAGES_KV=2 : 96        + 96    + 64    = 256 KiB  (over cap)
+        cga1, STAGES_KV=1 : 96        + 48    + 32    = 176 KiB
+    """
+    q_slab = cfg.TILE_M * cfg.TILE_K * cfg.BPE
+    o_slab = cfg.TILE_M * cfg.TILE_O * cfg.BPE_O
+    qo = cfg.TILES_Q * (max(q_slab, o_slab) if cfg.QO_ALIAS else q_slab + o_slab)
+    k = cfg.STAGES_KV * (cfg.TILE_N * cfg.TILE_K * cfg.BPE // cfg.CTA_MMA)
+    v = cfg.STAGES_KV * (cfg.TILE_O * cfg.TILE_N * cfg.BPE // cfg.CTA_MMA)
+    return qo + k + v
+
+
 def _validate_cfg_d192(cfg: CfgD192) -> None:
     """Consistency checks on the native DSv3 d192/d128 geometry."""
     checks = (
@@ -665,7 +798,16 @@ def _validate_cfg_d192(cfg: CfgD192) -> None:
         (cfg.MMA_REGS == cfg.TMALDG_REGS == cfg.TMASTG_REGS == cfg.SCHEDULER_REGS, "d192: MMA/TMALDG/TMASTG/SCHEDULER regs must match"),
         (cfg.MMA_REGS + cfg.CORRECTION_REGS + cfg.SOFTMAX_WARPGROUPS * cfg.SOFTMAX_REGS <= 512, "d192: register budget over 512"),
         (cfg.MMA_REGS % 8 == 0 and cfg.CORRECTION_REGS % 8 == 0 and cfg.SOFTMAX_REGS % 8 == 0, "d192: per-role regs must be multiples of 8"),
-        (cfg.CGA_M == 2 and cfg.CTA_MMA == 2, "d192 SM100 is cga2-only (CGA_M == CTA_MMA == 2)"),
+        (cfg.CGA_M == cfg.CTA_MMA and cfg.CTA_MMA in (1, 2), "d192 SM100: CGA_M must equal CTA_MMA, and CTA_MMA must be 1 (cga1) or 2 (cga2)"),
+        (
+            cfg.STAGES_KV == (1 if cfg.CTA_MMA == 1 else 2),
+            "d192: STAGES_KV must scale with the cluster width (2 at cga2, 1 at cga1) — cga1 doubles per-CTA K/V, "
+            "and cuDNN's own d192 kernel uses stages_kv = 1 * CTA_MMA for exactly this reason",
+        ),
+        (
+            _d192_smem_bytes(cfg) <= _SM100_MAX_DYN_SMEM,
+            f"d192: SMEM {_d192_smem_bytes(cfg) // 1024} KiB over the SM100 {_SM100_MAX_DYN_SMEM // 1024} KiB per-CTA cap",
+        ),
         (cfg.TILE_K == 192 and cfg.TILE_O == 128, "d192: expected D_QK tile 192 and D_V tile 128"),
         (cfg.QO_ALIAS == 1, "d192: Q/O SMEM alias is required to stay within SM100 SMEM budget"),
         (cfg.TILES_Q == 2, "d192: TILES_Q must be 2"),
@@ -673,7 +815,6 @@ def _validate_cfg_d192(cfg: CfgD192) -> None:
         (cfg.CORRECTION_WARPS == 4, "d192: CORRECTION_WARPS must be 4"),
         (cfg.TOTAL_WARPS == 16 and cfg.THREADS_PER_CTA == 512, "d192: 16 warps / 512 threads"),
         (cfg.READ_TILE_ARRIVERS == 15, f"d192: expected READ_TILE_ARRIVERS=15, got {cfg.READ_TILE_ARRIVERS}"),
-        (cfg.STAGES_KV == 2, "d192 SM100: STAGES_KV must be 2 for BF16/FP16"),
         (cfg.TILE_K_HW_BMM1 == 16 and cfg.TILE_K_HW_BMM2 == 16, "d192: TILE_K_HW must be 16 for BF16/FP16 on SM10x"),
         (cfg.Q_SWZ_BYTES == 128 and cfg.K_SWZ_BYTES == 128, "d192: Q/K swizzle must be 128B"),
         (cfg.V_SWZ_BYTES == 128 and cfg.O_SWZ_BYTES == 128, "d192: V/O swizzle must be 128B"),
@@ -691,9 +832,15 @@ def make_cfg_d192(params: TemplateParams) -> Tuple[CfgD192, TmaIters]:
         DTYPE_O=params.dtype_qkv,
         BPE=b,
         BPE_O=b,
+        SPLIT_KV=int(params.split_kv),
         Q_SWZ_BYTES=q_swz_bytes(192, b),
         K_SWZ_BYTES=q_swz_bytes(192, b),
-        V_SWZ_BYTES=v_swz_bytes(128, 2, b),
+        CGA_M=params.cta_mma,
+        CTA_MMA=params.cta_mma,
+        # cuDNN's d192 kernel uses stages_kv = 1 * CTA_MMA; cga1 doubles per-CTA
+        # K/V, so the stage count has to halve to keep the CTA inside the cap.
+        STAGES_KV=1 if params.cta_mma == 1 else 2,
+        V_SWZ_BYTES=v_swz_bytes(128, params.cta_mma, b),
         O_SWZ_BYTES=o_swz_bytes(128, b),
         RESCALE_THRESHOLD=rescale_threshold(params.dtype_qkv),
         TILE_K_HW_BMM1=tile_k_hw(params.dtype_qkv),
diff --git a/python/cudnn/sdpa/fwd/config_sm120.py b/python/cudnn/sdpa/fwd/config_sm120.py
index 031ecb0ec..9726e04ec 100644
--- a/python/cudnn/sdpa/fwd/config_sm120.py
+++ b/python/cudnn/sdpa/fwd/config_sm120.py
@@ -71,6 +71,11 @@ class TemplateParams:
     sched_policy: int = SCHED_NATURAL
     q_tile: int = SEQ_Q_TILES[0]
     kv_tile: int = SEQ_KV_TILES[0]
+    # KV split: each Q tile's KV-tile range [min_kv_tile, num_kv_tiles) is cut
+    # into ``split_kv`` contiguous chunks, each run by its own CTA writing a
+    # partial (O, LSE) that kernels/split_combine_sm100.py reduces.  1 = off.
+    # Opt-in only -- the graph front door never selects it.
+    split_kv: int = 1
 
 
 def validate_params(
@@ -94,6 +99,22 @@ def validate_params(
         raise ValueError(f"SM120 SDPA: dtype_qkv must be one of {allowed_dtypes}; got {params.dtype_qkv}")
     if params.dtype_o not in allowed_o_dtypes:
         raise ValueError(f"SM120 SDPA: dtype_o must be one of {allowed_o_dtypes}; got {params.dtype_o}")
+    if params.split_kv < 1:
+        raise ValueError(f"SM120 SDPA: split_kv must be >= 1 (1 = KV-split off); got {params.split_kv}")
+    if params.split_kv > 1:
+        # Each of these would need extra machinery in the combine pass.
+        if params.thd_varlen:
+            raise ValueError("SM120 SDPA: split_kv > 1 is dense-only (THD packs its own flat grid)")
+        if params.has_sink:
+            # The sink logit is folded into the softmax denominator per tile, so
+            # every split would add its own copy of it.
+            raise ValueError("SM120 SDPA: split_kv > 1 with an attention sink is not supported")
+        if params.sched_policy != SCHED_NATURAL:
+            # The split index rides the NATURAL grid's batch axis (y = batch +
+            # split*B); LPT / LPT_L2 flatten the grid to 1-D and derive the batch
+            # from the linear tile id, so there is no axis left to fold it into.
+            raise ValueError(f"SM120 SDPA: split_kv > 1 currently requires SCHED_NATURAL; got sched_policy={params.sched_policy}")
+
     if params.window_right is not None and params.window_right < 0:
         raise ValueError(f"SM120 SDPA: window_right must be None (unbounded) or >= 0 (0 = plain causal); got {params.window_right}")
     if not allow_right_band and params.window_right not in (None, 0):
diff --git a/python/cudnn/sdpa/fwd/kernels/_common_sm100.py b/python/cudnn/sdpa/fwd/kernels/_common_sm100.py
index 336fca559..9023cbf53 100644
--- a/python/cudnn/sdpa/fwd/kernels/_common_sm100.py
+++ b/python/cudnn/sdpa/fwd/kernels/_common_sm100.py
@@ -302,6 +302,204 @@ def compute_kv_loop_bounds(
     )
 
 
+class SplitHelpers(NamedTuple):
+    """Split-aware decode / bounds closures, plus the two flags kernels fold on."""
+
+    SPLIT_KV: int
+    # True when a tile's KV range can come out empty (right <= left).  See
+    # make_split_helpers for why KV split makes this reachable without a mask.
+    MAY_BE_EMPTY: bool
+    split_chunk: object
+    decode_initial_split: object
+    decode_payload_split: object
+    bounds_for_tile_split: object
+    nomask_range_split: object
+    partial_batch: object
+
+
+def make_split_helpers(CFG, *, bounds_for_tile, dispatch_decode_initial, dispatch_decode_payload) -> SplitHelpers:
+    """Split-aware decode / bounds closures shared by the SM100 prefill flavors.
+
+    ``bounds_for_tile`` is the caller's own bounds closure, taking
+    ``(q_super_idx, seqlen_q, seqlen_kv, cta_in_pair, seq_q_lens_tensor,
+    batch_idx)`` — flavors differ in whether they apply the dead-Q-tile trim, so
+    the split narrowing composes on top of whatever they already do.
+
+    At SPLIT_KV == 1 every closure below folds away and the traced code is the
+    classic single-pass kernel.
+    """
+    SPLIT_KV = int(getattr(CFG, "SPLIT_KV", 1))
+
+    # The split index rides the BATCH axis (grid.z), not grid.x: the decode
+    # already recovers the batch coordinate on BOTH the blockIdx and the
+    # scheduler-handout paths -- it is the high half of the packed head|batch
+    # word -- so a composite z = batch + split*B travels with it for free, with
+    # no in-place mutation of the shared tile id and no dependence on the grid's
+    # x extent (which is q_clusters * CGA_M, NOT the n_q_supers the kernel is
+    # handed, on any flavor where CGA_M != CTA_MMA -- d512).
+
+    # Can a tile's KV range come out EMPTY (right <= left)?
+    #
+    # Before KV split the answer was "only under a mask" — a SWA/causal/padded
+    # tile can fall entirely outside the band — so the empty-tile handshake
+    # (mb_empty_mainloop: correction arrives, MMA waits, TMA-LDG skips its
+    # loads) was gated on MASK_FLAGS != 0 and folded away at MASK_NONE, where
+    # [0, S_kv/TILE_N) is never empty.
+    #
+    # KV split breaks that WITHOUT a mask: a split past the end of a short range
+    # legitimately gets zero tiles.  Correction detects empties with a RUNTIME
+    # test, so if the gate const-folds to False the warp groups disagree —
+    # correction jumps to its epilogue while MMA waits on mb_q_full and TMA-LDG
+    # issues loads nobody consumes — and the kernel deadlocks.
+    MAY_BE_EMPTY = (CFG.MASK_FLAGS != 0) or (SPLIT_KV > 1)
+
+    # Which grid does this flavor launch?  SCHED_NATURAL uses a 3-D
+    # (q_super, head, batch) grid; the LPT policy flattens everything into x.
+    # NOTE this is the flavor's EFFECTIVE policy, not the requested one:
+    # make_cfg_d192 hardcodes SCHEDULER_POLICY=1 regardless of params, so a
+    # params-level check would miss it (and did -- d192 silently launched the
+    # unsplit grid because the split multiplier was only on the NATURAL branch).
+    IS_LPT = CFG.SCHEDULER_POLICY != SCHED_NATURAL
+
+    @cute.jit
+    def _lpt_split_of(raw, n_q_supers, n_qh, n_batch):
+        """(within-split raw x, split) for the flattened LPT grid.
+
+        The LPT tile space is q_tiles * n_qh * n_batch clusters; KV split
+        appends SPLIT_KV copies of it, split-major, so the split is the high
+        digit of the cluster index.  The CGA lane (raw % CGA_M) is preserved so
+        the caller's decode still sees a well-formed x coordinate.
+        """
+        cga = cutlass.Int32(CFG.CGA_M)
+        linear = raw // cga
+        q_tiles = n_q_supers // cutlass.Int32(CFG.CTA_MMA)
+        per_split = q_tiles * n_qh * n_batch
+        split = linear // per_split
+        rest = (linear % per_split) * cga + (raw % cga)
+        return rest, split
+
+    @cute.jit
+    def _split_chunk(left, right, split_idx):
+        """Cut ``[left, right)`` into SPLIT_KV near-equal chunks.
+
+        The FIRST ``rem`` splits get one extra tile, so chunk sizes differ by at
+        most 1 however the mask has already narrowed the range — a balanced cut
+        matters because the slowest split sets the critical path.  A split past
+        the end collapses to ``lo == hi``; the existing empty-mainloop path then
+        writes O := 0 / LSE := -inf, exactly the identity of the combine's
+        log-sum-exp, so no special case is needed downstream.
+        """
+        n_tiles = right - left
+        per = n_tiles // cutlass.Int32(SPLIT_KV)
+        rem = n_tiles % cutlass.Int32(SPLIT_KV)
+        lo = left + split_idx * per + cute.math.min(split_idx, rem)
+        extra = cutlass.Int32(
+            arith.select(
+                (split_idx < rem).ir_value(),
+                cutlass.Int32(1).ir_value(),
+                cutlass.Int32(0).ir_value(),
+            )
+        )
+        return lo, lo + per + extra
+
+    @cute.jit
+    def _decode_initial_split(bidx, bidy, bidz, cta_in_pair, n_q_supers, n_qh, n_batch, seq_kv_lens_t, qh_per_kh=None, seqlen_kv=None):
+        """decode_initial + this tile's split index.
+
+        NATURAL: the split rides the BATCH axis (see the note above on why not
+        grid.x).  The host launches z = B * SPLIT_KV, so the split falls out of
+        the DECODED batch coordinate as ``b // n_batch``, leaving the real batch
+        as ``b % n_batch``.
+
+        LPT / LPT_L2: the grid is flat, so there is no batch axis to ride and the
+        split is folded into the linear tile id instead; ``_lpt_split_of`` peels
+        it back off before the flavor's dispatcher sees the id.
+
+        ``qh_per_kh`` / ``seqlen_kv`` are the LPT_L2 cost-model inputs; they are
+        opaque here and forwarded to the flavor's dispatcher unchanged.
+        """
+        if cutlass.const_expr(SPLIT_KV > 1 and IS_LPT):
+            raw, split = _lpt_split_of(bidx, n_q_supers, n_qh, n_batch)
+            q, h, b = dispatch_decode_initial(raw, bidy, bidz, cta_in_pair, n_q_supers, n_qh, n_batch, seq_kv_lens_t, qh_per_kh, seqlen_kv)
+            return q, h, b, split
+        q, h, b = dispatch_decode_initial(bidx, bidy, bidz, cta_in_pair, n_q_supers, n_qh, n_batch, seq_kv_lens_t, qh_per_kh, seqlen_kv)
+        if cutlass.const_expr(SPLIT_KV == 1):
+            return q, h, b, cutlass.Int32(0)
+        return q, h, b % n_batch, b // n_batch
+
+    @cute.jit
+    def _decode_payload_split(t0, t1, cta_in_pair, n_q_supers, n_qh, n_batch, seq_kv_lens_t, qh_per_kh=None, seqlen_kv=None):
+        """decode_payload + split index; ``t0`` is the try_cancel cluster-base id."""
+        if cutlass.const_expr(SPLIT_KV > 1 and IS_LPT):
+            raw, split = _lpt_split_of(t0, n_q_supers, n_qh, n_batch)
+            q, h, b = dispatch_decode_payload(raw, t1, cta_in_pair, n_q_supers, n_qh, n_batch, seq_kv_lens_t, qh_per_kh, seqlen_kv)
+            return q, h, b, split
+        q, h, b = dispatch_decode_payload(t0, t1, cta_in_pair, n_q_supers, n_qh, n_batch, seq_kv_lens_t, qh_per_kh, seqlen_kv)
+        if cutlass.const_expr(SPLIT_KV == 1):
+            return q, h, b, cutlass.Int32(0)
+        return q, h, b % n_batch, b // n_batch
+
+    @cute.jit
+    def _bounds_for_tile_split(q_super_idx, seqlen_q, seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx):
+        """The flavor's bounds, narrowed to this split's slice of the KV range.
+
+        Splitting the ALREADY-masked ``[left, right)`` rather than the raw KV
+        extent is what keeps causal / SWA correct AND balanced: each split gets
+        an equal share of the tile's real work, not of the sequence.  The
+        unmasked band is clamped into the slice, which preserves the
+        ``left <= unmasked_lo <= unmasked_hi <= right`` invariant the mainloop
+        relies on, because clamping is monotone.
+        """
+        b = bounds_for_tile(q_super_idx, seqlen_q, seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+        if cutlass.const_expr(SPLIT_KV == 1):
+            return b
+        lo, hi = _split_chunk(b.left, b.right, split_idx)
+        return KvLoopBounds(
+            left=lo,
+            unmasked_lo=cute.math.min(cute.math.max(b.unmasked_lo, lo), hi),
+            unmasked_hi=cute.math.min(cute.math.max(b.unmasked_hi, lo), hi),
+            right=hi,
+        )
+
+    @cute.jit
+    def _nomask_range_split(seqlen_kv, split_idx):
+        """MASK_NONE fast-path KV range, split-aware.
+
+        MMA / TMA-LDG take this path while softmax / correction go through
+        bounds_for_tile_split; every warp group must land on the SAME chunk
+        boundaries or their mbarrier handshakes desync.  The split branch
+        therefore divides with the div-up compute_kv_loop_bounds uses, not the
+        floor of the historical fast path.
+        """
+        if cutlass.const_expr(SPLIT_KV == 1):
+            return cutlass.Int32(0), seqlen_kv // cutlass.Int32(CFG.TILE_N)
+        n_tiles = (seqlen_kv + cutlass.Int32(CFG.TILE_N - 1)) // cutlass.Int32(CFG.TILE_N)
+        return _split_chunk(cutlass.Int32(0), n_tiles, split_idx)
+
+    @cute.jit
+    def _partial_batch(batch_idx, split_idx, n_batch):
+        """Batch coord of this split's partial O / LSE slot (split-major).
+
+        Stacking the partials on the BATCH axis (extent B*SPLIT_KV) means the O
+        TMA descriptor is untouched — only the coord shifts.  Folds to batch_idx
+        at SPLIT_KV == 1.
+        """
+        if cutlass.const_expr(SPLIT_KV == 1):
+            return batch_idx
+        return batch_idx + split_idx * n_batch
+
+    return SplitHelpers(
+        SPLIT_KV=SPLIT_KV,
+        MAY_BE_EMPTY=MAY_BE_EMPTY,
+        split_chunk=_split_chunk,
+        decode_initial_split=_decode_initial_split,
+        decode_payload_split=_decode_payload_split,
+        bounds_for_tile_split=_bounds_for_tile_split,
+        nomask_range_split=_nomask_range_split,
+        partial_batch=_partial_batch,
+    )
+
+
 class SdpaHelpers(NamedTuple):
     decode_initial: object
     decode_payload: object
diff --git a/python/cudnn/sdpa/fwd/kernels/prefill_d128_f16_sm100.py b/python/cudnn/sdpa/fwd/kernels/prefill_d128_f16_sm100.py
index 0d313815a..9cfdd3321 100644
--- a/python/cudnn/sdpa/fwd/kernels/prefill_d128_f16_sm100.py
+++ b/python/cudnn/sdpa/fwd/kernels/prefill_d128_f16_sm100.py
@@ -9,9 +9,13 @@
 threads, persistent try_cancel scheduler, no Q∪O alias.
 
 SM100 resource layout:
-  1. **cga2-only, STAGES_KV=2**: per-CTA K/V halved by the collective MMA →
-     SMEM 192 KiB (sQ 64 + sK 32 + sV 32 + sO 64), under the Blackwell
-     ~228 KiB cap.
+  1. **STAGES_KV=2, both cluster widths at 192 KiB**: at cga2 (CTA_MMA=2,
+     the default) the collective MMA halves per-CTA K/V → sQ 64 + sK 32 +
+     sV 32 + sO 64.  At cga1 (``cta_mma=1``) there is no collective MMA, so
+     K/V double to 64 + 64 and the 64 KiB is bought back by aliasing Q and O
+     into one slab (``QO_ALIAS``, mandatory at cga1 and enforced by
+     ``_validate_cfg_d128``) → 64(Q∪O) + 64 + 64.  Both land under the
+     Blackwell ~228 KiB cap; ``_d128_smem_bytes`` is the checked model.
   2. **TMEM stats** (512-col Blackwell cap): S_acc 0/128 + O 256..511 fill
      all 512 cols, so stats ride the FREE HEAD of each sub-tile's S_acc
      slot — sub-tile 0 → col 0, sub-tile 1 → col 128 (P only aliases the
@@ -42,6 +46,30 @@
 
 FP8 (E4M3 / E5M2, incl. output-dtype override) lives in the sibling
 ``prefill_d128_fp8_sm100.py`` (shares this flavor's config).
+
+KV split (``CFG.SPLIT_KV > 1``): each chunk runs the UNCHANGED mainloop and
+epilogue, writing its partial O and LSE into a split-major workspace at batch
+coord ``b + s*B`` -- so the O TMA descriptor is untouched, only the batch coord
+shifts.  ``split_combine_sm100.py`` reduces over the split axis.
+
+KV split composes with the cluster width, and the pair is what closes the gap
+to cuDNN's own SM100 prefill split-K.  cga1 halves both the wasted MMA work at
+small S_q (a tile covers 256 Q rows, not 512) and the CTAs per tile, so twice
+as many splits fit in one wave.  Measured on B200 at B=1, H=16, S_q=128,
+S_kv=32K, d=128, end-to-end incl. combine, max|O-ref| 2e-5 throughout:
+
+    cga2 SPLIT_KV=1   414 us    (32 CTAs — the classic kernel)
+    cga2 SPLIT_KV=4   117 us    (128 CTAs)  3.6x
+    cga1 SPLIT_KV=8    73 us    (128 CTAs)  5.7x   <- best
+    (cuDNN 9.26/9.30's best plan for the same shape: 69 us, cga1 + 8 splits)
+
+On square shapes (S_q == S_kv, 1K/4K/8K) cga1 and cga2 measure within noise —
+same CTA count, same work per CTA — so cga1 is a small-S_q lever, not a
+regression risk elsewhere.
+Gated (config_sm100._validate_params) to SCHED_NATURAL, dense (non-THD), no
+sink; requires ``has_lse=True`` since the per-split LSE drives the combine.
+At SPLIT_KV == 1 every split helper const-folds away and the traced code is the
+classic single-pass kernel.
 """
 
 from functools import lru_cache
@@ -135,6 +163,7 @@
 
 from cudnn.sdpa.fwd.kernels._common_sm100 import (
     Bars,
+    make_split_helpers,
     KvLoopBounds,
     make_classic_bars,
     compute_kv_loop_bounds,
@@ -200,6 +229,28 @@
 _thd_tma_offsets = _sdpa_h.thd_tma_offsets
 
 
+# === KV split ===
+#
+# The mechanics live in _common_sm100.make_split_helpers (shared with the other
+# SM100 prefill flavors); this is the small-S_q lever for this kernel: a cga2
+# cluster covers TILES_Q*TILE_M*CTA_MMA = 512 Q rows, so at S_q = 128 the whole
+# problem is ceil(S_q/512) * H * B clusters (32 CTAs at H=16, B=1) on a 148-SM
+# part, each walking the entire 32K-token KV loop alone.
+_split_h = make_split_helpers(
+    CFG,
+    bounds_for_tile=_bounds_for_tile,
+    dispatch_decode_initial=_dispatch_decode_initial,
+    dispatch_decode_payload=_dispatch_decode_payload,
+)
+SPLIT_KV = _split_h.SPLIT_KV
+MAY_BE_EMPTY = _split_h.MAY_BE_EMPTY
+_decode_initial_split = _split_h.decode_initial_split
+_decode_payload_split = _split_h.decode_payload_split
+_bounds_for_tile_split = _split_h.bounds_for_tile_split
+_nomask_range_split = _split_h.nomask_range_split
+_partial_batch = _split_h.partial_batch
+
+
 # P (BMM2 operand) aliases the TAIL of each 128-col S_acc slot since BMM1
 # finishes (and softmax loads S into registers) before P is written.
 # fp16/bf16 pack 2 probs per FP32 cell → P width = TILE_N/2 ≤ 64 cols → P
@@ -617,7 +668,7 @@ def _tmaldg_warp_group(
     tma_k = GmemTileTma(tma_k_desc)
     tma_v = GmemTileTma(tma_v_desc)
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -634,13 +685,15 @@ def _tmaldg_warp_group(
     q_row_base = cute.arch.make_warp_uniform(q_super_idx * cutlass.Int32(CFG.TILES_Q * CFG.TILE_M))
     q_seq_off, kv_seq_off, tma_batch = _thd_tma_offsets(seq_kv_lens_tensor, batch_idx, n_batch)
 
-    if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+    if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV == 1):
         kv_left = cutlass.Int32(0)
         kv_right = seqlen_kv // cutlass.Int32(CFG.TILE_N)
+    elif cutlass.const_expr(CFG.MASK_FLAGS == 0):
+        kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
     else:
         eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
         eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds_init = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+        bounds_init = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
         kv_left = bounds_init.left
         kv_right = bounds_init.right
 
@@ -654,7 +707,7 @@ def _tmaldg_warp_group(
     while is_valid_tile > cutlass.Int32(0):
         read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
 
-        if cutlass.const_expr(CFG.MASK_FLAGS != 0) and (kv_right <= kv_left):
+        if cutlass.const_expr(MAY_BE_EMPTY) and (kv_right <= kv_left):
             pass
         else:
             # Prologue interleave Q[0] -> K[first] -> Q[1] -> V[first] -> mainloop —
@@ -760,7 +813,7 @@ def _tmaldg_warp_group(
         nxt_q = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load())
         nxt_hb = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load())
         nxt_v = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load())
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -777,10 +830,12 @@ def _tmaldg_warp_group(
         q_seq_off, kv_seq_off, tma_batch = _thd_tma_offsets(seq_kv_lens_tensor, batch_idx, n_batch)
         is_valid_tile = nxt_v & cutlass.Int32(1)
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
-        if cutlass.const_expr(CFG.MASK_FLAGS != 0):
+        if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV > 1):
+            kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+        elif cutlass.const_expr(CFG.MASK_FLAGS != 0):
             eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
             eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-            bounds_next = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+            bounds_next = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
             kv_left = bounds_next.left
             kv_right = bounds_next.right
 
@@ -822,7 +877,7 @@ def _tmastg_warp_group(
 
     tma_o = GmemTileTma(tma_o_desc)
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -841,6 +896,10 @@ def _tmastg_warp_group(
         read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
 
         q_row_base = q_super_idx * cutlass.Int32(CFG.TILES_Q * CFG.TILE_M)
+        # KV split: partials are stacked split-major on the BATCH axis of the
+        # workspace (extent B*SPLIT_KV), so the store needs no new descriptor —
+        # only a shifted batch coord.  Folds to batch_idx at SPLIT_KV == 1.
+        o_batch = _partial_batch(batch_idx, split_idx, n_batch)
 
         for qs in cutlass.range_constexpr(CFG.TILES_Q):
             bars.mb_o_full[qs].wait(o_full_phase)
@@ -861,7 +920,7 @@ def _tmastg_warp_group(
             else:
                 tma_store_tile(
                     sO[qs],
-                    tma_o(cutlass.Int32(0), head_idx, q_row_base + cutlass.Int32(qs * CFG.TILE_M), batch_idx),
+                    tma_o(cutlass.Int32(0), head_idx, q_row_base + cutlass.Int32(qs * CFG.TILE_M), o_batch),
                 )
 
             tma_store_commit()
@@ -879,7 +938,7 @@ def _tmastg_warp_group(
         nxt_q = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load()
         nxt_hb = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load()
         nxt_v = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load()
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -1022,11 +1081,13 @@ def _mma_warp_group(
     desc_Q0 = sQ[0].desc()
     desc_Q1 = sQ[1].desc()
 
-    if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+    if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV == 1):
         kv_left = cutlass.Int32(0)
         kv_right = seqlen_kv // cutlass.Int32(CFG.TILE_N)
     else:
-        q_super_idx, _hd, batch_idx = _dispatch_decode_initial(
+        # Under KV split the MASK_NONE path still has to decode, because the
+        # split index (and hence this tile's KV slice) lives in the tile id.
+        q_super_idx, _hd, batch_idx, split_idx = _decode_initial_split(
             sched.bidx_init,
             sched.bidy_init,
             sched.bidz_init,
@@ -1038,11 +1099,14 @@ def _mma_warp_group(
             qh_per_kh,
             seqlen_kv,
         )
-        eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
-        eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds_init = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
-        kv_left = bounds_init.left
-        kv_right = bounds_init.right
+        if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+            kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+        else:
+            eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
+            eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
+            bounds_init = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
+            kv_left = bounds_init.left
+            kv_right = bounds_init.right
 
     q_full_phase = cutlass.Int32(0)
     kv_state = PipelineState.start(phase=0)
@@ -1065,7 +1129,7 @@ def _mma_warp_group(
     while is_valid_tile > cutlass.Int32(0):
         read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
 
-        if cutlass.const_expr(CFG.MASK_FLAGS != 0) and (kv_right <= kv_left):
+        if cutlass.const_expr(MAY_BE_EMPTY) and (kv_right <= kv_left):
             # Empty mainloop: fire both bmm2_done so softmax/correction phase trackers stay in lockstep.
             bars.mb_empty_mainloop.wait(empty_mainloop_phase)
             empty_mainloop_phase = empty_mainloop_phase ^ cutlass.Int32(1)
@@ -1224,14 +1288,14 @@ def _mma_warp_group(
         nvvm.bar_warp_sync(cute.arch.FULL_MASK)
 
         wait(sched.mb_scheduler.subview(sched_state.idx), sched_state.phase)
-        if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+        if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV == 1):
             nxt_v = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load()
             is_valid_tile = nxt_v & cutlass.Int32(1)
         else:
             nxt_q = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load())
             nxt_hb = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load())
             nxt_v = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load())
-            q_super_idx, _hd, batch_idx = _dispatch_decode_payload(
+            q_super_idx, _hd, batch_idx, split_idx = _decode_payload_split(
                 nxt_q,
                 nxt_hb,
                 cta_in_pair,
@@ -1243,11 +1307,14 @@ def _mma_warp_group(
                 seqlen_kv,
             )
             is_valid_tile = nxt_v & cutlass.Int32(1)
-            eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
-            eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-            bounds_next = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
-            kv_left = bounds_next.left
-            kv_right = bounds_next.right
+            if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+                kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+            else:
+                eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
+                eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
+                bounds_next = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
+                kv_left = bounds_next.left
+                kv_right = bounds_next.right
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
 
     bars.mb_tmem_dealloc.wait(cutlass.Int32(0))
@@ -1509,7 +1576,7 @@ def _softmax_warp_group(
         cutlass.Float32,
     )
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -1527,7 +1594,7 @@ def _softmax_warp_group(
     eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
 
     eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-    bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+    bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
 
     softmax_wg_base_const = CFG.SOFTMAX_WG0_BASE if sub_tile_id == 0 else CFG.SOFTMAX_WG1_BASE
     tid_in_wg = cute.arch.thread_idx()[0] - cutlass.Int32(softmax_wg_base_const * 32)
@@ -1644,7 +1711,7 @@ def _softmax_warp_group(
         nxt_q = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load())
         nxt_hb = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load())
         nxt_v = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load())
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -1659,7 +1726,7 @@ def _softmax_warp_group(
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
         eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
         eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+        bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
 
 
 @cute.jit
@@ -1713,7 +1780,7 @@ def _correction_warp_group(
     bmm2_done_phase = cutlass.Int32(0)
     o_empty_phase = cutlass.Int32(1)  # bootstrap
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -1731,7 +1798,7 @@ def _correction_warp_group(
     eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
 
     eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-    bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+    bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
 
     while is_valid_tile > cutlass.Int32(0):
         read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
@@ -1884,7 +1951,12 @@ def _correction_warp_group(
             else:
                 if q_row_global < seqlen_q:
                     lse_arr = cutlass.make_array_view(lse_tensor)
-                    lse_row = lse_arr[batch_idx, head_idx, :]
+                    # KV split: this chunk's LSE goes to its own split-major
+                    # slot (batch extent B*SPLIT_KV), matching where TMA-STG put
+                    # the chunk's O.  The pair (O_s, lse_s) is everything the
+                    # combine needs.  Folds to batch_idx at SPLIT_KV == 1.
+                    lse_batch = _partial_batch(batch_idx, split_idx, n_batch)
+                    lse_row = lse_arr[lse_batch, head_idx, :]
                     lse_row[q_row_global] = lse_val
 
             sO_sub_base = sO[qs].base
@@ -1929,7 +2001,7 @@ def _correction_warp_group(
         nxt_q = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load()
         nxt_hb = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load()
         nxt_v = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load()
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -1944,7 +2016,7 @@ def _correction_warp_group(
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
         eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
         eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+        bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
 
     # End-of-warp tmem_dealloc: under cga2 each corr lane ALSO DSMEM-arrives
     # on the peer so the peer's local mbar accumulates the full CGA-total count.
@@ -2072,7 +2144,13 @@ def _tma_swz(byte_w: int):
         grid_shape = (n_thd_units * cutlass.Int32(CFG.CGA_M), cutlass.Int32(1), cutlass.Int32(1))
     else:
         # Grid Python-folds on Cfg constant (avoids DSL if staging).
-        grid_shape = (grid_q_supers, QH, B) if cutlass.const_expr(CFG.SCHEDULER_POLICY == SCHED_NATURAL) else (grid_q_supers * QH * B, 1, 1)
+        # KV split rides the BATCH axis: z = batch + split*B.  The decode
+        # already recovers the batch coord on both the blockIdx and the
+        # scheduler-handout paths, so the split travels with it for free.  SPLIT_KV > 1 is
+        # gated to SCHED_NATURAL by the config validator.
+        grid_shape = (
+            (grid_q_supers, QH, B * SPLIT_KV) if cutlass.const_expr(CFG.SCHEDULER_POLICY == SCHED_NATURAL) else (grid_q_supers * QH * B * SPLIT_KV, 1, 1)
+        )
     _kernel(
         tma_q_desc,
         tma_k_desc,
@@ -2147,6 +2225,10 @@ def compile(  # noqa: A001
         raise ValueError(f"d128 envelope: need 0 < d_qk <= {CFG.TILE_K} and 0 < d_v <= {CFG.TILE_O}; got ({d_qk}, {d_v})")
     if (d_qk * CFG.BPE) % 16 != 0 or (d_v * CFG.BPE_O) % 16 != 0:
         raise ValueError(f"d128 envelope: d_qk*BPE and d_v*BPE must be 16-byte multiples (TMA global-stride rule); got ({d_qk}, {d_v}) at BPE={CFG.BPE}")
+    if SPLIT_KV > 1 and not has_lse:
+        # Each split's LSE is not optional under KV split — it IS the weight the
+        # combine reduces with.  Without it the partials cannot be recombined.
+        raise ValueError("d128: split_kv > 1 requires has_lse=True (the per-split LSE drives the combine)")
     _fake_batch = 1 if CFG.THD_VARLEN else b
     if CFG.THD_VARLEN:
         # Dynamic packed token totals: one symbol per ragged group (Q/O and
@@ -2154,6 +2236,11 @@ def compile(  # noqa: A001
         # compiled artifact instead of minting a new one (issue #552).
         sq = cute.sym_int(divisibility=1)
         skv = cute.sym_int(divisibility=1)
+    # KV split: O and LSE are the PARTIAL workspaces, stacked split-major on the
+    # batch axis (B*SPLIT_KV).  Q/K/V keep the real batch — only the outputs
+    # grow.  At SPLIT_KV == 1 these are the plain output tensors.
+    _o_batch = _fake_batch * SPLIT_KV
+    _lse_batch = b * SPLIT_KV
 
     def _fake_bshd(shape, stride, dtype=STORAGE_DTYPE, bpe=CFG.BPE):
         if stride is None:
@@ -2172,7 +2259,7 @@ def _fake_bshd(shape, stride, dtype=STORAGE_DTYPE, bpe=CFG.BPE):
     fake_q = _fake_bshd((_fake_batch, sq, qh, d_qk), q_stride)
     fake_k = _fake_bshd((_fake_batch, skv, kh, d_qk), k_stride)
     fake_v = _fake_bshd((_fake_batch, skv, kh, d_v), v_stride)
-    fake_o = _fake_bshd((_fake_batch, sq, qh, d_v), o_stride, dtype=STORAGE_DTYPE)
+    fake_o = _fake_bshd((_o_batch, sq, qh, d_v), o_stride, dtype=STORAGE_DTYPE)
     if not has_lse:
         # No Stats output: the LSE argument is None-specialized and the store
         # is compiled out entirely — no dummy buffer exists at any level.
@@ -2212,7 +2299,7 @@ def _fake_bshd(shape, stride, dtype=STORAGE_DTYPE, bpe=CFG.BPE):
             raise ValueError("lse_head_major / lse_head_stride are THD-only (dense LSE is compact (B, H, Sq))")
         fake_lse = cute.runtime.make_fake_compact_tensor(
             cutlass.Float32,
-            (b, qh, sq),
+            (_lse_batch, qh, sq),
             stride_order=(2, 1, 0),
             assumed_align=16,
         )
diff --git a/python/cudnn/sdpa/fwd/kernels/prefill_d128_fp8_sm100.py b/python/cudnn/sdpa/fwd/kernels/prefill_d128_fp8_sm100.py
index 84d2e5c12..fd9118a39 100644
--- a/python/cudnn/sdpa/fwd/kernels/prefill_d128_fp8_sm100.py
+++ b/python/cudnn/sdpa/fwd/kernels/prefill_d128_fp8_sm100.py
@@ -151,6 +151,7 @@
 
 
 from cudnn.sdpa.fwd.kernels._common_sm100 import (
+    make_split_helpers,
     Bars,
     KvLoopBounds,
     make_classic_bars,
@@ -195,6 +196,41 @@
 _thd_tma_offsets = _sdpa_h.thd_tma_offsets
 
 
+# === KV split ===
+#
+# Mechanics live in _common_sm100.make_split_helpers, shared with the other
+# SM100 prefill flavors: each Q tile's KV loop range is cut into SPLIT_KV
+# contiguous chunks, each run as its own persistent tile, and each writing a
+# normalized partial O + its own LSE into a split-major workspace that
+# split_combine_sm100 folds with the exact log-sum-exp identity.  At
+# SPLIT_KV == 1 every closure folds away and this is the classic kernel.
+
+
+@cute.jit
+def _bounds_for_tile_uniform(q_super_idx, seqlen_q, seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx):
+    """Uniform 6-arg bounds signature for make_split_helpers.
+
+    This flavor has no dead-Q-tile trim, so the trailing two args are
+    accepted and ignored — callers pass None for them.
+    """
+    return _bounds_for_tile(q_super_idx, seqlen_q, seqlen_kv, cta_in_pair)
+
+
+_split_h = make_split_helpers(
+    CFG,
+    bounds_for_tile=_bounds_for_tile_uniform,
+    dispatch_decode_initial=_dispatch_decode_initial,
+    dispatch_decode_payload=_dispatch_decode_payload,
+)
+SPLIT_KV = _split_h.SPLIT_KV
+MAY_BE_EMPTY = _split_h.MAY_BE_EMPTY
+_decode_initial_split = _split_h.decode_initial_split
+_decode_payload_split = _split_h.decode_payload_split
+_bounds_for_tile_split = _split_h.bounds_for_tile_split
+_nomask_range_split = _split_h.nomask_range_split
+_partial_batch = _split_h.partial_batch
+
+
 @dataclass(frozen=True)
 class KernelTmemLayout:
     """Column offsets for the classic 2-sub-tile SDPA pipeline (FP8).
@@ -586,7 +622,7 @@ def _tmaldg_warp_group(
     tma_k = GmemTileTma(tma_k_desc)
     tma_v = GmemTileTma(tma_v_desc)
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -603,13 +639,15 @@ def _tmaldg_warp_group(
     q_row_base = cute.arch.make_warp_uniform(q_super_idx * cutlass.Int32(CFG.TILES_Q * CFG.TILE_M))
     q_seq_off, kv_seq_off, tma_batch = _thd_tma_offsets(seq_kv_lens_tensor, batch_idx, n_batch)
 
-    if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+    if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV == 1):
         kv_left = cutlass.Int32(0)
         kv_right = seqlen_kv // cutlass.Int32(CFG.TILE_N)
+    elif cutlass.const_expr(CFG.MASK_FLAGS == 0):
+        kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
     else:
         eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
         eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds_init = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair)
+        bounds_init = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, None, None, split_idx)
         kv_left = bounds_init.left
         kv_right = bounds_init.right
 
@@ -623,7 +661,7 @@ def _tmaldg_warp_group(
     while is_valid_tile > cutlass.Int32(0):
         read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
 
-        if cutlass.const_expr(CFG.MASK_FLAGS != 0) and (kv_right <= kv_left):
+        if cutlass.const_expr(MAY_BE_EMPTY) and (kv_right <= kv_left):
             pass
         else:
             # Prologue interleave: Q[0] → K[first] → Q[1] → V[first] → mainloop.
@@ -722,7 +760,7 @@ def _tmaldg_warp_group(
         nxt_q = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load())
         nxt_hb = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load())
         nxt_v = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load())
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -739,10 +777,12 @@ def _tmaldg_warp_group(
         q_seq_off, kv_seq_off, tma_batch = _thd_tma_offsets(seq_kv_lens_tensor, batch_idx, n_batch)
         is_valid_tile = nxt_v & cutlass.Int32(1)
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
-        if cutlass.const_expr(CFG.MASK_FLAGS != 0):
+        if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV > 1):
+            kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+        elif cutlass.const_expr(CFG.MASK_FLAGS != 0):
             eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
             eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-            bounds_next = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair)
+            bounds_next = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, None, None, split_idx)
             kv_left = bounds_next.left
             kv_right = bounds_next.right
 
@@ -780,7 +820,7 @@ def _tmastg_warp_group(
 
     tma_o = GmemTileTma(tma_o_desc)
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -799,6 +839,10 @@ def _tmastg_warp_group(
         read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
 
         q_row_base = q_super_idx * cutlass.Int32(CFG.TILES_Q * CFG.TILE_M)
+        # KV split: partials are stacked split-major on the workspace BATCH axis
+        # (extent B*SPLIT_KV), so the store needs no new descriptor — only a
+        # shifted batch coord.  Folds to batch_idx at SPLIT_KV == 1.
+        o_batch = _partial_batch(batch_idx, split_idx, n_batch)
 
         for qs in cutlass.range_constexpr(CFG.TILES_Q):
             bars.mb_o_full[qs].wait(o_full_phase)
@@ -806,7 +850,7 @@ def _tmastg_warp_group(
             # O TMA params follow O's swizzle, not V's (V and O swizzles may differ).
             tma_store_tile(
                 sO[qs],
-                tma_o(cutlass.Int32(0), head_idx, q_row_base + cutlass.Int32(qs * CFG.TILE_M), batch_idx),
+                tma_o(cutlass.Int32(0), head_idx, q_row_base + cutlass.Int32(qs * CFG.TILE_M), o_batch),
             )
 
             tma_store_commit()
@@ -820,7 +864,7 @@ def _tmastg_warp_group(
         nxt_q = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load()
         nxt_hb = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load()
         nxt_v = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load()
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -955,11 +999,11 @@ def _mma_warp_group(
     desc_Q0 = sQ[0].desc()
     desc_Q1 = sQ[1].desc()
 
-    if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+    if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV == 1):
         kv_left = cutlass.Int32(0)
         kv_right = seqlen_kv // cutlass.Int32(CFG.TILE_N)
     else:
-        q_super_idx, _hd, batch_idx = _dispatch_decode_initial(
+        q_super_idx, _hd, batch_idx, split_idx = _decode_initial_split(
             sched.bidx_init,
             sched.bidy_init,
             sched.bidz_init,
@@ -971,11 +1015,14 @@ def _mma_warp_group(
             qh_per_kh,
             seqlen_kv,
         )
-        eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
-        eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds_init = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair)
-        kv_left = bounds_init.left
-        kv_right = bounds_init.right
+        if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+            kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+        else:
+            eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
+            eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
+            bounds_init = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, None, None, split_idx)
+            kv_left = bounds_init.left
+            kv_right = bounds_init.right
 
     q_full_phase = cutlass.Int32(0)
     kv_state = PipelineState.start(phase=0)
@@ -999,7 +1046,7 @@ def _mma_warp_group(
     while is_valid_tile > cutlass.Int32(0):
         read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
 
-        if cutlass.const_expr(CFG.MASK_FLAGS != 0) and (kv_right <= kv_left):
+        if cutlass.const_expr(MAY_BE_EMPTY) and (kv_right <= kv_left):
             # Empty-kv tile: fire bmm2_done so softmax/corr phases stay in lockstep.
             bars.mb_empty_mainloop.wait(empty_mainloop_phase)
             empty_mainloop_phase = empty_mainloop_phase ^ cutlass.Int32(1)
@@ -1150,14 +1197,14 @@ def _mma_warp_group(
         nvvm.bar_warp_sync(cute.arch.FULL_MASK)
 
         wait(sched.mb_scheduler.subview(sched_state.idx), sched_state.phase)
-        if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+        if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV == 1):
             nxt_v = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load()
             is_valid_tile = nxt_v & cutlass.Int32(1)
         else:
             nxt_q = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load())
             nxt_hb = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load())
             nxt_v = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load())
-            q_super_idx, _hd, batch_idx = _dispatch_decode_payload(
+            q_super_idx, _hd, batch_idx, split_idx = _decode_payload_split(
                 nxt_q,
                 nxt_hb,
                 cta_in_pair,
@@ -1169,11 +1216,14 @@ def _mma_warp_group(
                 seqlen_kv,
             )
             is_valid_tile = nxt_v & cutlass.Int32(1)
-            eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
-            eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-            bounds_next = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair)
-            kv_left = bounds_next.left
-            kv_right = bounds_next.right
+            if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+                kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+            else:
+                eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
+                eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
+                bounds_next = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, None, None, split_idx)
+                kv_left = bounds_next.left
+                kv_right = bounds_next.right
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
 
     bars.mb_tmem_dealloc.wait(cutlass.Int32(0))
@@ -1404,7 +1454,7 @@ def _softmax_warp_group(
         cutlass.Float32,
     )
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -1422,7 +1472,7 @@ def _softmax_warp_group(
     eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
 
     eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-    bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair)
+    bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, None, None, split_idx)
 
     softmax_wg_base_const = CFG.SOFTMAX_WG0_BASE if sub_tile_id == 0 else CFG.SOFTMAX_WG1_BASE
     tid_in_wg = cute.arch.thread_idx()[0] - cutlass.Int32(softmax_wg_base_const * 32)
@@ -1539,7 +1589,7 @@ def _softmax_warp_group(
         nxt_q = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load())
         nxt_hb = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load())
         nxt_v = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load())
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -1554,7 +1604,7 @@ def _softmax_warp_group(
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
         eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
         eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair)
+        bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, None, None, split_idx)
 
 
 @cute.jit
@@ -1608,7 +1658,7 @@ def _correction_warp_group(
     bmm2_done_phase = cutlass.Int32(0)
     o_empty_phase = cutlass.Int32(1)  # bootstrap pre-armed at phase 1 so first wait passes
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -1626,7 +1676,7 @@ def _correction_warp_group(
     eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
 
     eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-    bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair)
+    bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, None, None, split_idx)
 
     while is_valid_tile > cutlass.Int32(0):
         read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
@@ -1748,7 +1798,11 @@ def _correction_warp_group(
             if _row_valid:
                 if cutlass.const_expr(lse_tensor is not None):
                     lse_arr = cutlass.make_array_view(lse_tensor)
-                    lse_row = lse_arr[batch_idx, head_idx, :]
+                    # This chunk's LSE goes to its own split-major slot, matching where
+                    # TMA-STG put the chunk's O.  The pair (O_s, lse_s) is everything
+                    # the combine needs.
+                    lse_batch = _partial_batch(batch_idx, split_idx, n_batch)
+                    lse_row = lse_arr[lse_batch, head_idx, :]
                     lse_row[q_row_global] = lse_val
 
             # amax_o = max over valid rows of |o_scaled| (the fp32 pre-cast output). Divided
@@ -1837,8 +1891,14 @@ def _correction_warp_group(
                         bars.mb_o_empty[qs].wait(o_empty_phase)
                     smem_ptr.store_swizzled(o_half, alignment=64, swizzle=_O_SMEM_SWIZZLE)
 
-            if _row_valid:
-                nvvm.atomicrmw(nvvm.AtomicOp.MAX, _amax_o_ptr, _amax_o_local.bitcast(cutlass.Int32))
+            # Under KV split this epilogue sees only its OWN partial, and the
+            # recombined O is a convex combination of the partials -- so a max
+            # over partials over-reports the output amax (~2.9x at 8 splits).
+            # split_combine_sm100 computes it over the recombined O instead;
+            # this write has to stay out of the way, since atomicMax only grows.
+            if cutlass.const_expr(SPLIT_KV == 1):
+                if _row_valid:
+                    nvvm.atomicrmw(nvvm.AtomicOp.MAX, _amax_o_ptr, _amax_o_local.bitcast(cutlass.Int32))
 
             # fence_proxy needed before TMA reads SMEM written by stores above.
             nvvm.fence_proxy("async.shared", space="cta")
@@ -1855,7 +1915,7 @@ def _correction_warp_group(
         nxt_q = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load()
         nxt_hb = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load()
         nxt_v = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load()
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -1870,7 +1930,7 @@ def _correction_warp_group(
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
         eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
         eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair)
+        bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, None, None, split_idx)
 
     # tmem_dealloc fan-out: fire one arrive per lane; cga2 also DSMEM-arrives
     # on the peer so peer's local mbar accumulates the full CGA count.
@@ -1956,7 +2016,10 @@ def _tma_swz(byte_w: int):
     q_clusters = (SQ + rows_per_cluster - 1) // rows_per_cluster
     grid_q_supers = q_clusters * CFG.CTA_MMA
     q_supers = grid_q_supers
-    grid_shape = (grid_q_supers, QH, B) if cutlass.const_expr(CFG.SCHEDULER_POLICY == SCHED_NATURAL) else (grid_q_supers * QH * B, 1, 1)
+    # KV split rides the BATCH axis: z = batch + split*B.  The decode
+    # already recovers the batch coord on both the blockIdx and the
+    # scheduler-handout paths, so the split travels with it for free.
+    grid_shape = (grid_q_supers, QH, B * SPLIT_KV) if cutlass.const_expr(CFG.SCHEDULER_POLICY == SCHED_NATURAL) else (grid_q_supers * QH * B * SPLIT_KV, 1, 1)
     _kernel(
         tma_q_desc,
         tma_k_desc,
@@ -1993,6 +2056,14 @@ def compile(b: int = 1, qh: int = 1, kh: int = 1, sq: int = 256, skv: int = 128,
     ``has_lse=False`` compiles the LSE store out (the kernel specializes on a
     ``None`` LSE argument) — callers without a Stats output pass no LSE buffer
     at all; the amax_o atomicMax write is independent and unchanged."""
+    if SPLIT_KV > 1 and not has_lse:
+        # Each split's LSE is not optional under KV split — it IS the weight
+        # the combine reduces with.  Without it the partials cannot be recombined.
+        raise ValueError("split_kv > 1 requires has_lse=True (the per-split LSE drives the combine)")
+    # KV split: O and LSE are the PARTIAL workspaces, stacked split-major on
+    # the batch axis (B*SPLIT_KV).  Q/K/V keep the real batch.
+    _o_batch = b * SPLIT_KV
+    _lse_batch = b * SPLIT_KV
     fake_q = cute.runtime.make_fake_compact_tensor(
         STORAGE_DTYPE,
         (b, sq, qh, CFG.TILE_K),
@@ -2013,7 +2084,7 @@ def compile(b: int = 1, qh: int = 1, kh: int = 1, sq: int = 256, skv: int = 128,
     )
     fake_o = cute.runtime.make_fake_compact_tensor(
         OUT_STORAGE_DTYPE,
-        (b, sq, qh, CFG.TILE_O),
+        (_o_batch, sq, qh, CFG.TILE_O),
         stride_order=(3, 2, 1, 0),
         assumed_align=16,
     )
@@ -2024,7 +2095,7 @@ def compile(b: int = 1, qh: int = 1, kh: int = 1, sq: int = 256, skv: int = 128,
     else:
         fake_lse = cute.runtime.make_fake_compact_tensor(
             cutlass.Float32,
-            (b, qh, sq),
+            (_lse_batch, qh, sq),
             stride_order=(2, 1, 0),
             assumed_align=16,
         )
diff --git a/python/cudnn/sdpa/fwd/kernels/prefill_d128_mxfp8_sm100.py b/python/cudnn/sdpa/fwd/kernels/prefill_d128_mxfp8_sm100.py
index e629faa0c..c6870fed0 100644
--- a/python/cudnn/sdpa/fwd/kernels/prefill_d128_mxfp8_sm100.py
+++ b/python/cudnn/sdpa/fwd/kernels/prefill_d128_mxfp8_sm100.py
@@ -196,6 +196,7 @@ def _round_up(a: int, b: int) -> int:
 
 
 from cudnn.sdpa.fwd.kernels._common_sm100 import (
+    make_split_helpers,
     Bars,
     KvLoopBounds,
     make_classic_bars,
@@ -238,6 +239,41 @@ def _round_up(a: int, b: int) -> int:
 _dispatch_decode_initial = _sdpa_h.dispatch_decode_initial
 _dispatch_decode_payload = _sdpa_h.dispatch_decode_payload
 _thd_tma_offsets = _sdpa_h.thd_tma_offsets
+
+
+# === KV split ===
+#
+# Mechanics live in _common_sm100.make_split_helpers, shared with the other
+# SM100 prefill flavors: each Q tile's KV loop range is cut into SPLIT_KV
+# contiguous chunks, each run as its own persistent tile, and each writing a
+# normalized partial O + its own LSE into a split-major workspace that
+# split_combine_sm100 folds with the exact log-sum-exp identity.  At
+# SPLIT_KV == 1 every closure folds away and this is the classic kernel.
+
+
+@cute.jit
+def _bounds_for_tile_uniform(q_super_idx, seqlen_q, seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx):
+    """Uniform 6-arg bounds signature for make_split_helpers.
+
+    This flavor has no dead-Q-tile trim, so the trailing two args are
+    accepted and ignored — callers pass None for them.
+    """
+    return _bounds_for_tile(q_super_idx, seqlen_q, seqlen_kv, cta_in_pair)
+
+
+_split_h = make_split_helpers(
+    CFG,
+    bounds_for_tile=_bounds_for_tile_uniform,
+    dispatch_decode_initial=_dispatch_decode_initial,
+    dispatch_decode_payload=_dispatch_decode_payload,
+)
+SPLIT_KV = _split_h.SPLIT_KV
+MAY_BE_EMPTY = _split_h.MAY_BE_EMPTY
+_decode_initial_split = _split_h.decode_initial_split
+_decode_payload_split = _split_h.decode_payload_split
+_bounds_for_tile_split = _split_h.bounds_for_tile_split
+_nomask_range_split = _split_h.nomask_range_split
+_partial_batch = _split_h.partial_batch
 _thd_sf_tile_bases = _sdpa_h.thd_sf_tile_bases
 
 
@@ -773,7 +809,7 @@ def _tmaldg_warp_group(
     k_sf_peer_row = cta_in_pair * cutlass.Int32(K_SF_ROWS_PER_PEER)
     v_sf_peer_row = cta_in_pair * cutlass.Int32(V_SF_ROWS_PER_PEER)
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -793,13 +829,15 @@ def _tmaldg_warp_group(
     q_seq_off, kv_seq_off, tma_batch = _thd_tma_offsets(seq_kv_lens_tensor, batch_idx, n_batch)
     cu_sf_q_base, cu_sf_k_base = _thd_sf_tile_bases(seq_kv_lens_tensor, batch_idx, n_batch)
 
-    if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+    if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV == 1):
         kv_left = cutlass.Int32(0)
         kv_right = seqlen_kv // cutlass.Int32(CFG.TILE_N)
+    elif cutlass.const_expr(CFG.MASK_FLAGS == 0):
+        kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
     else:
         eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
         eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds_init = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair)
+        bounds_init = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, None, None, split_idx)
         kv_left = bounds_init.left
         kv_right = bounds_init.right
 
@@ -818,7 +856,7 @@ def _tmaldg_warp_group(
         read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
 
         # Empty-kv tile: MMA's empty-mainloop branch handles the matching mbar phases.
-        if cutlass.const_expr(CFG.MASK_FLAGS != 0) and (kv_right <= kv_left):
+        if cutlass.const_expr(MAY_BE_EMPTY) and (kv_right <= kv_left):
             pass
         else:
             q_sf_tile_base = q_row_base // cutlass.Int32(CFG.TILE_M)
@@ -976,7 +1014,7 @@ def _tmaldg_warp_group(
         nxt_q = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load())
         nxt_hb = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load())
         nxt_v = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load())
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -993,10 +1031,12 @@ def _tmaldg_warp_group(
         cu_sf_q_base, cu_sf_k_base = _thd_sf_tile_bases(seq_kv_lens_tensor, batch_idx, n_batch)
         is_valid_tile = nxt_v & cutlass.Int32(1)
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
-        if cutlass.const_expr(CFG.MASK_FLAGS != 0):
+        if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV > 1):
+            kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+        elif cutlass.const_expr(CFG.MASK_FLAGS != 0):
             eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
             eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-            bounds_next = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair)
+            bounds_next = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, None, None, split_idx)
             kv_left = bounds_next.left
             kv_right = bounds_next.right
 
@@ -1031,7 +1071,7 @@ def _tmastg_warp_group(
 
     tma_o = GmemTileTma(tma_o_desc)
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -1050,6 +1090,10 @@ def _tmastg_warp_group(
         read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
 
         q_row_base = q_super_idx * cutlass.Int32(CFG.TILES_Q * CFG.TILE_M)
+        # KV split: partials are stacked split-major on the workspace BATCH axis
+        # (extent B*SPLIT_KV), so the store needs no new descriptor — only a
+        # shifted batch coord.  Folds to batch_idx at SPLIT_KV == 1.
+        o_batch = _partial_batch(batch_idx, split_idx, n_batch)
 
         for qs in cutlass.range_constexpr(CFG.TILES_Q):
             bars.mb_o_full[qs].wait(o_full_phase)
@@ -1057,7 +1101,7 @@ def _tmastg_warp_group(
             # O TMA params follow O's swizzle, NOT V's — required when V_SWZ_B != O_SWZ_B.
             tma_store_tile(
                 sO[qs],
-                tma_o(cutlass.Int32(0), head_idx, q_row_base + cutlass.Int32(qs * CFG.TILE_M), batch_idx),
+                tma_o(cutlass.Int32(0), head_idx, q_row_base + cutlass.Int32(qs * CFG.TILE_M), o_batch),
             )
 
             tma_store_commit()
@@ -1071,7 +1115,7 @@ def _tmastg_warp_group(
         nxt_q = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load()
         nxt_hb = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load()
         nxt_v = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load()
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -1255,11 +1299,11 @@ def _utccp_bmm2_sf(tmem_sf_p, tmem_sf_v, smem_desc_p, smem_desc_v):
         nvvm.tcgen05_cp(nvvm.Tcgen05CpShape.SHAPE_32X128B, tmem_sf_p, smem_desc_p, group=CTA_GROUP_KIND, multicast=nvvm.Tcgen05CpMulticast.WARPX4)
         nvvm.tcgen05_cp(nvvm.Tcgen05CpShape.SHAPE_32X128B, tmem_sf_v, smem_desc_v, group=CTA_GROUP_KIND, multicast=nvvm.Tcgen05CpMulticast.WARPX4)
 
-    if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+    if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV == 1):
         kv_left = cutlass.Int32(0)
         kv_right = seqlen_kv // cutlass.Int32(CFG.TILE_N)
     else:
-        q_super_idx, _hd, batch_idx = _dispatch_decode_initial(
+        q_super_idx, _hd, batch_idx, split_idx = _decode_initial_split(
             sched.bidx_init,
             sched.bidy_init,
             sched.bidz_init,
@@ -1271,11 +1315,14 @@ def _utccp_bmm2_sf(tmem_sf_p, tmem_sf_v, smem_desc_p, smem_desc_v):
             qh_per_kh,
             seqlen_kv,
         )
-        eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
-        eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds_init = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair)
-        kv_left = bounds_init.left
-        kv_right = bounds_init.right
+        if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+            kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+        else:
+            eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
+            eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
+            bounds_init = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, None, None, split_idx)
+            kv_left = bounds_init.left
+            kv_right = bounds_init.right
 
     q_full_phase = cutlass.Int32(0)
     kv_state = PipelineState.start(phase=0)
@@ -1304,7 +1351,7 @@ def _utccp_bmm2_sf(tmem_sf_p, tmem_sf_v, smem_desc_p, smem_desc_v):
     while is_valid_tile > cutlass.Int32(0):
         read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
 
-        if cutlass.const_expr(CFG.MASK_FLAGS != 0) and (kv_right <= kv_left):
+        if cutlass.const_expr(MAY_BE_EMPTY) and (kv_right <= kv_left):
             # Empty mainloop — keep softmax/correction phase trackers in lockstep with non-empty path.
             bars.mb_empty_mainloop.wait(empty_mainloop_phase)
             empty_mainloop_phase = empty_mainloop_phase ^ cutlass.Int32(1)
@@ -1552,14 +1599,14 @@ def _utccp_bmm2_sf(tmem_sf_p, tmem_sf_v, smem_desc_p, smem_desc_v):
         nvvm.bar_warp_sync(cute.arch.FULL_MASK)
 
         wait(sched.mb_scheduler.subview(sched_state.idx), sched_state.phase)
-        if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+        if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV == 1):
             nxt_v = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load()
             is_valid_tile = nxt_v & cutlass.Int32(1)
         else:
             nxt_q = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load())
             nxt_hb = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load())
             nxt_v = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load())
-            q_super_idx, _hd, batch_idx = _dispatch_decode_payload(
+            q_super_idx, _hd, batch_idx, split_idx = _decode_payload_split(
                 nxt_q,
                 nxt_hb,
                 cta_in_pair,
@@ -1571,11 +1618,14 @@ def _utccp_bmm2_sf(tmem_sf_p, tmem_sf_v, smem_desc_p, smem_desc_v):
                 seqlen_kv,
             )
             is_valid_tile = nxt_v & cutlass.Int32(1)
-            eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
-            eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-            bounds_next = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair)
-            kv_left = bounds_next.left
-            kv_right = bounds_next.right
+            if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+                kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+            else:
+                eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
+                eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
+                bounds_next = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, None, None, split_idx)
+                kv_left = bounds_next.left
+                kv_right = bounds_next.right
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
 
     bars.mb_tmem_dealloc.wait(cutlass.Int32(0))
@@ -1838,7 +1888,7 @@ def _softmax_warp_group(
         cutlass.Float32,
     )
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -1856,7 +1906,7 @@ def _softmax_warp_group(
     eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
 
     eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-    bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair)
+    bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, None, None, split_idx)
 
     softmax_wg_base_const = CFG.SOFTMAX_WG0_BASE if sub_tile_id == 0 else CFG.SOFTMAX_WG1_BASE
     tid_in_wg = cute.arch.thread_idx()[0] - cutlass.Int32(softmax_wg_base_const * 32)
@@ -1972,7 +2022,7 @@ def _softmax_warp_group(
         nxt_q = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load())
         nxt_hb = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load())
         nxt_v = cute.arch.make_warp_uniform((sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load())
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -1987,7 +2037,7 @@ def _softmax_warp_group(
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
         eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
         eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair)
+        bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, None, None, split_idx)
 
 
 @cute.jit
@@ -2035,7 +2085,7 @@ def _correction_warp_group(
     bmm2_done_phase = cutlass.Int32(0)
     o_empty_phase = cutlass.Int32(1)  # bootstrap pre-armed at phase 1 so first wait passes immediately
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -2053,7 +2103,7 @@ def _correction_warp_group(
     eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
 
     eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-    bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair)
+    bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, None, None, split_idx)
 
     while is_valid_tile > cutlass.Int32(0):
         read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
@@ -2166,7 +2216,11 @@ def _correction_warp_group(
             if cutlass.const_expr(lse_tensor is not None):
                 if _row_valid:
                     lse_arr = cutlass.make_array_view(lse_tensor)
-                    lse_row = lse_arr[batch_idx, head_idx, :]
+                    # This chunk's LSE goes to its own split-major slot, matching where
+                    # TMA-STG put the chunk's O.  The pair (O_s, lse_s) is everything
+                    # the combine needs.
+                    lse_batch = _partial_batch(batch_idx, split_idx, n_batch)
+                    lse_row = lse_arr[lse_batch, head_idx, :]
                     lse_row[q_row_global] = lse_val
 
             sO_sub_base = sO[qs].base
@@ -2200,8 +2254,15 @@ def _correction_warp_group(
                 smem_ptr.store_swizzled(o_out, alignment=64, swizzle=_O_SMEM_SWIZZLE)
 
             # One atomic per valid row (invalid/OOB rows must not poison the global amax).
-            if _row_valid:
-                nvvm.atomicrmw(nvvm.AtomicOp.MAX, _amax_o_ptr, _amax_o_local.bitcast(cutlass.Int32))
+            #
+            # Under KV split this epilogue sees only its OWN partial, and the
+            # recombined O is a convex combination of the partials -- so a max
+            # over partials over-reports the output amax.  split_combine_sm100
+            # computes it over the recombined O instead; this write has to stay
+            # out of the way, since atomicMax only grows.
+            if cutlass.const_expr(SPLIT_KV == 1):
+                if _row_valid:
+                    nvvm.atomicrmw(nvvm.AtomicOp.MAX, _amax_o_ptr, _amax_o_local.bitcast(cutlass.Int32))
 
             # fence_proxy needed before TMA reads SMEM written by tcgen05_st.
             nvvm.fence_proxy("async.shared", space="cta")
@@ -2217,7 +2278,7 @@ def _correction_warp_group(
         nxt_q = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load()
         nxt_hb = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load()
         nxt_v = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load()
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -2232,7 +2293,7 @@ def _correction_warp_group(
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
         eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
         eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair)
+        bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, None, None, split_idx)
 
     # mb_tmem_dealloc fan-out — all-lanes arrive + DSMEM-arrive on cross-pair peer under cga2.
     if cutlass.const_expr(CFG.CTA_MMA == 2):
@@ -2353,7 +2414,10 @@ def _build_sf_desc(sf_tensor, num_tiles, sf_smem_size, num_rows_box, num_heads):
     q_clusters = (SQ + rows_per_cluster - 1) // rows_per_cluster
     grid_q_supers = q_clusters * CFG.CTA_MMA
     q_supers = grid_q_supers
-    grid_shape = (grid_q_supers, QH, B) if cutlass.const_expr(CFG.SCHEDULER_POLICY == SCHED_NATURAL) else (grid_q_supers * QH * B, 1, 1)
+    # KV split rides the BATCH axis: z = batch + split*B.  The decode
+    # already recovers the batch coord on both the blockIdx and the
+    # scheduler-handout paths, so the split travels with it for free.
+    grid_shape = (grid_q_supers, QH, B * SPLIT_KV) if cutlass.const_expr(CFG.SCHEDULER_POLICY == SCHED_NATURAL) else (grid_q_supers * QH * B * SPLIT_KV, 1, 1)
     _kernel(
         tma_q_desc,
         tma_k_desc,
@@ -2395,6 +2459,15 @@ def compile(  # noqa: A001
     ``has_lse=False`` compiles the LSE store out (the kernel specializes on a
     ``None`` LSE argument) — callers without a Stats output pass no LSE buffer
     at all; the amax_o atomicMax write is independent and unchanged."""
+    if SPLIT_KV > 1 and not has_lse:
+        # Each split's LSE is not optional under KV split — it IS the weight
+        # the combine reduces with.  Without it the partials cannot be recombined.
+        raise ValueError("split_kv > 1 requires has_lse=True (the per-split LSE drives the combine)")
+    # KV split: O and LSE are the PARTIAL workspaces, stacked split-major on
+    # the batch axis (B*SPLIT_KV).  Q/K/V keep the real batch.
+    _o_batch = b * SPLIT_KV
+    _lse_batch = b * SPLIT_KV
+
     # Q SF tiles TILE_M-row wide → num_tiles = SQ/TILE_M; K/V SF TILE_N-row wide → num_tiles = SKV/TILE_N.
     sq_tiles = (sq + CFG.TILE_M - 1) // CFG.TILE_M
     skv_tiles = (skv + CFG.TILE_N - 1) // CFG.TILE_N
@@ -2421,7 +2494,7 @@ def compile(  # noqa: A001
     )
     fake_o = cute.runtime.make_fake_compact_tensor(
         OUT_STORAGE_DTYPE,
-        (b, sq, qh, CFG.TILE_O),
+        (_o_batch, sq, qh, CFG.TILE_O),
         stride_order=(3, 2, 1, 0),
         assumed_align=16,
     )
@@ -2452,7 +2525,7 @@ def compile(  # noqa: A001
     else:
         fake_lse = cute.runtime.make_fake_compact_tensor(
             cutlass.Float32,
-            (b, qh, sq),
+            (_lse_batch, qh, sq),
             stride_order=(2, 1, 0),
             assumed_align=16,
         )
diff --git a/python/cudnn/sdpa/fwd/kernels/prefill_d192_d128_f16_sm100.py b/python/cudnn/sdpa/fwd/kernels/prefill_d192_d128_f16_sm100.py
index 6eac98386..d0b4fb547 100644
--- a/python/cudnn/sdpa/fwd/kernels/prefill_d192_d128_f16_sm100.py
+++ b/python/cudnn/sdpa/fwd/kernels/prefill_d192_d128_f16_sm100.py
@@ -137,6 +137,7 @@ def _wait_mbarrier(mb, phase):
 
 
 from cudnn.sdpa.fwd.kernels._common_sm100 import (
+    make_split_helpers,
     make_classic_bars,
     row_max_for_exp2,
     make_sdpa_helpers,
@@ -211,6 +212,38 @@ def _bounds_for_tile(
 _thd_tma_offsets = _sdpa_h.thd_tma_offsets
 
 
+# === KV split ===
+#
+# Mechanics live in _common_sm100.make_split_helpers, shared with the other
+# SM100 prefill flavors: each Q tile's KV loop range is cut into SPLIT_KV
+# contiguous chunks, each run as its own persistent tile, and each writing a
+# normalized partial O + its own LSE into a split-major workspace that
+# split_combine_sm100 folds with the exact log-sum-exp identity.  At
+# SPLIT_KV == 1 every closure folds away and this is the classic kernel.
+_split_h = make_split_helpers(
+    CFG,
+    bounds_for_tile=_bounds_for_tile,
+    dispatch_decode_initial=_dispatch_decode_initial,
+    dispatch_decode_payload=_dispatch_decode_payload,
+)
+SPLIT_KV = _split_h.SPLIT_KV
+MAY_BE_EMPTY = _split_h.MAY_BE_EMPTY
+_decode_initial_split = _split_h.decode_initial_split
+_decode_payload_split = _split_h.decode_payload_split
+_bounds_for_tile_split = _split_h.bounds_for_tile_split
+_nomask_range_split = _split_h.nomask_range_split
+_partial_batch = _split_h.partial_batch
+
+# This flavor carries its OWN empty-KV predicate (defined above), narrower than
+# MAY_BE_EMPTY: it counts only PADDED / SWA / bottom-right, the masks that can
+# empty a tile.  KV split makes an empty range reachable without any of them — a
+# split past the end of a short range gets zero tiles — so the nine sites gated
+# on CAN_HAVE_EMPTY_KV (the empty-tile handshake, the QO_ALIAS drain, and the
+# epilogue's O-store guards) have to compile in under split as well.  Leaving it
+# const-folded to False is what desynchronises the warp groups.
+CAN_HAVE_EMPTY_KV = CAN_HAVE_EMPTY_KV or SPLIT_KV > 1
+
+
 @cute.jit
 def _apply_top_left_causal_mask_chunk(reg_S, q_abs, kv_col_base, N: int = 64):
     neg_inf = cutlass.Float32(float("-inf"))
@@ -671,7 +704,7 @@ def _tmaldg_warp_group(
     tma_k = GmemTileTma(tma_k_desc)
     tma_v = GmemTileTma(tma_v_desc)
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -688,13 +721,15 @@ def _tmaldg_warp_group(
     q_row_base = cute.arch.make_warp_uniform(q_super_idx * cutlass.Int32(CFG.TILES_Q * CFG.TILE_M))
     q_seq_off, kv_seq_off, tma_batch = _thd_tma_offsets(seq_kv_lens_tensor, batch_idx, n_batch)
 
-    if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+    if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV == 1):
         kv_left = cutlass.Int32(0)
         kv_right = seqlen_kv // cutlass.Int32(CFG.TILE_N)
+    elif cutlass.const_expr(CFG.MASK_FLAGS == 0):
+        kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
     else:
         eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
         eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds_init = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+        bounds_init = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
         kv_left = bounds_init.left
         kv_right = bounds_init.right
 
@@ -708,7 +743,7 @@ def _tmaldg_warp_group(
     while is_valid_tile > cutlass.Int32(0):
         read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
 
-        if cutlass.const_expr(CFG.MASK_FLAGS != 0) and (kv_right <= kv_left):
+        if cutlass.const_expr(MAY_BE_EMPTY) and (kv_right <= kv_left):
             if cutlass.const_expr(CAN_HAVE_EMPTY_KV and IS_QO_ALIAS):
                 # TMA-STG advances the Q/O alias gate for every tile, including
                 # empty ones. Consume that transaction even though no Q reload
@@ -830,7 +865,7 @@ def _tmaldg_warp_group(
         nxt_q = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load()
         nxt_hb = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load()
         nxt_v = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load()
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -847,10 +882,12 @@ def _tmaldg_warp_group(
         q_seq_off, kv_seq_off, tma_batch = _thd_tma_offsets(seq_kv_lens_tensor, batch_idx, n_batch)
         is_valid_tile = nxt_v & cutlass.Int32(1)
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
-        if cutlass.const_expr(CFG.MASK_FLAGS != 0):
+        if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV > 1):
+            kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+        elif cutlass.const_expr(CFG.MASK_FLAGS != 0):
             eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
             eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-            bounds_next = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+            bounds_next = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
             kv_left = bounds_next.left
             kv_right = bounds_next.right
 
@@ -895,7 +932,7 @@ def _tmastg_warp_group(
 
     tma_o = GmemTileTma(tma_o_desc)
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -914,6 +951,10 @@ def _tmastg_warp_group(
         read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
 
         q_row_base = q_super_idx * cutlass.Int32(CFG.TILES_Q * CFG.TILE_M)
+        # KV split: partials are stacked split-major on the workspace BATCH axis
+        # (extent B*SPLIT_KV), so the store needs no new descriptor — only a
+        # shifted batch coord.  Folds to batch_idx at SPLIT_KV == 1.
+        o_batch = _partial_batch(batch_idx, split_idx, n_batch)
 
         for qs in cutlass.range_constexpr(CFG.TILES_Q):
             _wait_mbarrier(bars.mb_o_full[qs], o_full_phase)
@@ -934,7 +975,7 @@ def _tmastg_warp_group(
             else:
                 tma_store_tile(
                     sO[qs],
-                    tma_o(cutlass.Int32(0), head_idx, q_row_base + cutlass.Int32(qs * CFG.TILE_M), batch_idx),
+                    tma_o(cutlass.Int32(0), head_idx, q_row_base + cutlass.Int32(qs * CFG.TILE_M), o_batch),
                 )
 
             tma_store_commit()
@@ -963,7 +1004,7 @@ def _tmastg_warp_group(
         nxt_q = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load()
         nxt_hb = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load()
         nxt_v = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load()
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -1106,11 +1147,11 @@ def _mma_warp_group(
     desc_Q0 = sQ[0].desc()
     desc_Q1 = sQ[1].desc()
 
-    if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+    if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV == 1):
         kv_left = cutlass.Int32(0)
         kv_right = seqlen_kv // cutlass.Int32(CFG.TILE_N)
     else:
-        q_super_idx, _hd, batch_idx = _dispatch_decode_initial(
+        q_super_idx, _hd, batch_idx, split_idx = _decode_initial_split(
             sched.bidx_init,
             sched.bidy_init,
             sched.bidz_init,
@@ -1122,11 +1163,14 @@ def _mma_warp_group(
             qh_per_kh,
             seqlen_kv,
         )
-        eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
-        eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds_init = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
-        kv_left = bounds_init.left
-        kv_right = bounds_init.right
+        if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+            kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+        else:
+            eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
+            eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
+            bounds_init = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
+            kv_left = bounds_init.left
+            kv_right = bounds_init.right
 
     q_full_phase = cutlass.Int32(0)
     kv_state = PipelineState.start(phase=0)
@@ -1149,7 +1193,7 @@ def _mma_warp_group(
     while is_valid_tile > cutlass.Int32(0):
         read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
 
-        if cutlass.const_expr(CFG.MASK_FLAGS != 0) and (kv_right <= kv_left):
+        if cutlass.const_expr(MAY_BE_EMPTY) and (kv_right <= kv_left):
             if cutlass.const_expr(not CAN_HAVE_EMPTY_KV):
                 _wait_mbarrier(bars.mb_empty_mainloop, empty_mainloop_phase)
                 empty_mainloop_phase = empty_mainloop_phase ^ cutlass.Int32(1)
@@ -1307,14 +1351,14 @@ def _mma_warp_group(
         nvvm.bar_warp_sync(cute.arch.FULL_MASK)
 
         _wait_ptr(sched.mb_scheduler.subview(sched_state.idx), sched_state.phase)
-        if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+        if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV == 1):
             nxt_v = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load()
             is_valid_tile = nxt_v & cutlass.Int32(1)
         else:
             nxt_q = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load()
             nxt_hb = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load()
             nxt_v = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load()
-            q_super_idx, _hd, batch_idx = _dispatch_decode_payload(
+            q_super_idx, _hd, batch_idx, split_idx = _decode_payload_split(
                 nxt_q,
                 nxt_hb,
                 cta_in_pair,
@@ -1326,11 +1370,14 @@ def _mma_warp_group(
                 seqlen_kv,
             )
             is_valid_tile = nxt_v & cutlass.Int32(1)
-            eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
-            eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-            bounds_next = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
-            kv_left = bounds_next.left
-            kv_right = bounds_next.right
+            if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+                kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+            else:
+                eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
+                eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
+                bounds_next = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
+                kv_left = bounds_next.left
+                kv_right = bounds_next.right
 
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
 
@@ -1617,7 +1664,7 @@ def _softmax_warp_group(
         cutlass.Float32,
     )
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -1635,7 +1682,7 @@ def _softmax_warp_group(
     eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
 
     eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-    bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+    bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
 
     softmax_wg_base_const = CFG.SOFTMAX_WG0_BASE if sub_tile_id == 0 else CFG.SOFTMAX_WG1_BASE
     tid_in_wg = cute.arch.thread_idx()[0] - cutlass.Int32(softmax_wg_base_const * 32)
@@ -1755,7 +1802,7 @@ def _softmax_warp_group(
         nxt_q = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load()
         nxt_hb = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load()
         nxt_v = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load()
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -1770,7 +1817,7 @@ def _softmax_warp_group(
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
         eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
         eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+        bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
 
 
 @cute.jit
@@ -1824,7 +1871,7 @@ def _correction_warp_group(
     bmm2_done_phase = cutlass.Int32(0)
     o_empty_phase = cutlass.Int32(1)  # bootstrap
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -1842,7 +1889,7 @@ def _correction_warp_group(
     eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
 
     eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-    bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+    bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
 
     while is_valid_tile > cutlass.Int32(0):
         read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
@@ -1997,7 +2044,11 @@ def _correction_warp_group(
             else:
                 if q_row_global < seqlen_q:
                     lse_arr = cutlass.make_array_view(lse_tensor)
-                    lse_row = lse_arr[batch_idx, head_idx, :]
+                    # This chunk's LSE goes to its own split-major slot, matching
+                    # where TMA-STG put the chunk's O.  The pair (O_s, lse_s) is
+                    # everything the combine needs.
+                    lse_batch = _partial_batch(batch_idx, split_idx, n_batch)
+                    lse_row = lse_arr[lse_batch, head_idx, :]
                     lse_row[q_row_global] = lse_val
 
             sO_sub_base = sO[qs].base
@@ -2050,7 +2101,7 @@ def _correction_warp_group(
         nxt_q = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0))).load()
         nxt_hb = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1))).load()
         nxt_v = (sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2))).load()
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -2065,7 +2116,7 @@ def _correction_warp_group(
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
         eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
         eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+        bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
 
     # End-of-warp tmem_dealloc: under cga2 each corr lane ALSO DSMEM-arrives
     # on the peer so the peer's local mbar accumulates the full CGA-total count.
@@ -2191,7 +2242,12 @@ def _tma_swz(byte_w: int):
         grid_shape = (n_thd_units * cutlass.Int32(CFG.CGA_M), cutlass.Int32(1), cutlass.Int32(1))
     else:
         # Grid Python-folds on Cfg constant (avoids DSL if staging).
-        grid_shape = (grid_q_supers, QH, B) if cutlass.const_expr(CFG.SCHEDULER_POLICY == SCHED_NATURAL) else (grid_q_supers * QH * B, 1, 1)
+        # KV split rides the BATCH axis: z = batch + split*B.  The decode
+        # already recovers the batch coord on both the blockIdx and the
+        # scheduler-handout paths, so the split travels with it for free.
+        grid_shape = (
+            (grid_q_supers, QH, B * SPLIT_KV) if cutlass.const_expr(CFG.SCHEDULER_POLICY == SCHED_NATURAL) else (grid_q_supers * QH * B * SPLIT_KV, 1, 1)
+        )
     _kernel(
         tma_q_desc,
         tma_k_desc,
@@ -2258,6 +2314,10 @@ def compile(  # noqa: A001
         raise ValueError(f"d192 envelope: need 0 < d_qk <= {CFG.TILE_K} and 0 < d_v <= {CFG.TILE_O}; got ({d_qk}, {d_v})")
     if (d_qk * CFG.BPE) % 16 != 0 or (d_v * CFG.BPE_O) % 16 != 0:
         raise ValueError(f"d192 envelope: d_qk*BPE and d_v*BPE must be 16-byte multiples (TMA global-stride rule); got ({d_qk}, {d_v}) at BPE={CFG.BPE}")
+    if SPLIT_KV > 1 and not has_lse:
+        # Each split's LSE is not optional under KV split — it IS the weight
+        # the combine reduces with.  Without it the partials cannot be recombined.
+        raise ValueError("split_kv > 1 requires has_lse=True (the per-split LSE drives the combine)")
     _fake_batch = 1 if CFG.THD_VARLEN else b
     if CFG.THD_VARLEN:
         # Dynamic packed token totals: one symbol per ragged group (Q/O and
@@ -2265,6 +2325,10 @@ def compile(  # noqa: A001
         # compiled artifact instead of minting a new one (issue #552).
         sq = cute.sym_int(divisibility=1)
         skv = cute.sym_int(divisibility=1)
+    # KV split: O and LSE are the PARTIAL workspaces, stacked split-major on
+    # the batch axis (B*SPLIT_KV).  Q/K/V keep the real batch.
+    _o_batch = _fake_batch * SPLIT_KV
+    _lse_batch = b * SPLIT_KV
 
     def _fake_bshd(shape, stride, dtype=STORAGE_DTYPE, bpe=CFG.BPE):
         if stride is None:
@@ -2283,7 +2347,7 @@ def _fake_bshd(shape, stride, dtype=STORAGE_DTYPE, bpe=CFG.BPE):
     fake_q = _fake_bshd((_fake_batch, sq, qh, d_qk), q_stride)
     fake_k = _fake_bshd((_fake_batch, skv, kh, d_qk), k_stride)
     fake_v = _fake_bshd((_fake_batch, skv, kh, d_v), v_stride)
-    fake_o = _fake_bshd((_fake_batch, sq, qh, d_v), o_stride, dtype=STORAGE_DTYPE)
+    fake_o = _fake_bshd((_o_batch, sq, qh, d_v), o_stride, dtype=STORAGE_DTYPE)
     if not has_lse:
         # No Stats output: the LSE argument is None-specialized and the store
         # is compiled out entirely — no dummy buffer exists at any level.
@@ -2323,7 +2387,7 @@ def _fake_bshd(shape, stride, dtype=STORAGE_DTYPE, bpe=CFG.BPE):
             raise ValueError("lse_head_major / lse_head_stride are THD-only (dense LSE is compact (B, H, Sq))")
         fake_lse = cute.runtime.make_fake_compact_tensor(
             cutlass.Float32,
-            (b, qh, sq),
+            (_lse_batch, qh, sq),
             stride_order=(2, 1, 0),
             assumed_align=16,
         )
diff --git a/python/cudnn/sdpa/fwd/kernels/prefill_d256_f16_sm100.py b/python/cudnn/sdpa/fwd/kernels/prefill_d256_f16_sm100.py
index 936990e32..a3ddef5a3 100644
--- a/python/cudnn/sdpa/fwd/kernels/prefill_d256_f16_sm100.py
+++ b/python/cudnn/sdpa/fwd/kernels/prefill_d256_f16_sm100.py
@@ -83,6 +83,7 @@
 
 
 from cudnn.sdpa.fwd.kernels._common_sm100 import (
+    make_split_helpers,
     D256Bars as Bars,
     KvLoopBounds,
     make_d256_bars,
@@ -131,6 +132,29 @@
 _thd_tma_offsets = _sdpa_h.thd_tma_offsets
 
 
+# === KV split ===
+#
+# Mechanics live in _common_sm100.make_split_helpers, shared with the other
+# SM100 prefill flavors: each Q tile's KV loop range is cut into SPLIT_KV
+# contiguous chunks, each run as its own persistent tile, and each writing a
+# normalized partial O + its own LSE into a split-major workspace that
+# split_combine_sm100 folds with the exact log-sum-exp identity.  At
+# SPLIT_KV == 1 every closure folds away and this is the classic kernel.
+_split_h = make_split_helpers(
+    CFG,
+    bounds_for_tile=_bounds_for_tile,
+    dispatch_decode_initial=_dispatch_decode_initial,
+    dispatch_decode_payload=_dispatch_decode_payload,
+)
+SPLIT_KV = _split_h.SPLIT_KV
+MAY_BE_EMPTY = _split_h.MAY_BE_EMPTY
+_decode_initial_split = _split_h.decode_initial_split
+_decode_payload_split = _split_h.decode_payload_split
+_bounds_for_tile_split = _split_h.bounds_for_tile_split
+_nomask_range_split = _split_h.nomask_range_split
+_partial_batch = _split_h.partial_batch
+
+
 @dataclass(frozen=True)
 class KernelTmemLayout:
     TOTAL_COLS: int = 512
@@ -480,7 +504,7 @@ def _tmaldg_warp_group(
     tma_k = GmemTileTma(tma_k_desc)
     tma_v = GmemTileTma(tma_v_desc)
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -496,13 +520,15 @@ def _tmaldg_warp_group(
     q_row_base = cute.arch.make_warp_uniform(q_super_idx * cutlass.Int32(CFG.TILES_Q * CFG.TILE_M))
     q_seq_off, kv_seq_off, tma_batch = _thd_tma_offsets(seq_kv_lens_tensor, batch_idx, n_batch)
 
-    if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+    if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV == 1):
         kv_left = cutlass.Int32(0)
         kv_right = seqlen_kv // cutlass.Int32(CFG.TILE_N)
+    elif cutlass.const_expr(CFG.MASK_FLAGS == 0):
+        kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
     else:
         eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
         eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds_init = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+        bounds_init = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
         kv_left = bounds_init.left
         kv_right = bounds_init.right
 
@@ -518,7 +544,7 @@ def _tmaldg_warp_group(
         bars.mb_q_o_alias.wait(q_o_alias_phase)
         q_o_alias_phase = q_o_alias_phase ^ cutlass.Int32(1)
 
-        if cutlass.const_expr(CFG.MASK_FLAGS != 0) and (kv_right <= kv_left):
+        if cutlass.const_expr(MAY_BE_EMPTY) and (kv_right <= kv_left):
             pass
         else:
             if cutlass.const_expr(CFG.CTA_MMA == 2):
@@ -572,7 +598,7 @@ def _tmaldg_warp_group(
         nxt_q = cute.arch.make_warp_uniform(sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0)).load())
         nxt_hb = cute.arch.make_warp_uniform(sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1)).load())
         nxt_v = cute.arch.make_warp_uniform(sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2)).load())
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -588,10 +614,12 @@ def _tmaldg_warp_group(
         q_seq_off, kv_seq_off, tma_batch = _thd_tma_offsets(seq_kv_lens_tensor, batch_idx, n_batch)
         is_valid_tile = nxt_v & cutlass.Int32(1)
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
-        if cutlass.const_expr(CFG.MASK_FLAGS != 0):
+        if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV > 1):
+            kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+        elif cutlass.const_expr(CFG.MASK_FLAGS != 0):
             eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
             eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-            bounds_next = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+            bounds_next = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
             kv_left = bounds_next.left
             kv_right = bounds_next.right
 
@@ -624,7 +652,7 @@ def _tmastg_warp_group(
 
     tma_o = GmemTileTma(tma_o_desc)
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -649,6 +677,10 @@ def _tmastg_warp_group(
             bars.mb_o_full[chunk].wait(o_full_phase)
 
         q_row_coord = q_super_idx * cutlass.Int32(CFG.TILES_Q * CFG.TILE_M)
+        # KV split: partials are stacked split-major on the workspace BATCH
+        # axis (extent B*SPLIT_KV), so the store needs no new descriptor —
+        # only a shifted batch coord.  Folds to batch_idx at SPLIT_KV == 1.
+        o_batch = _partial_batch(batch_idx, split_idx, n_batch)
 
         if cutlass.const_expr(CFG.THD_VARLEN):
             # DEAD unit (batch == n_batch, envelope grid — issue #552): no O
@@ -661,7 +693,7 @@ def _tmastg_warp_group(
         else:
             tma_store_tile(
                 sO[0],
-                tma_o(cutlass.Int32(0), head_idx, q_row_coord, batch_idx),
+                tma_o(cutlass.Int32(0), head_idx, q_row_coord, o_batch),
             )
         tma_store_commit()
         tma_store_wait(0)
@@ -676,7 +708,7 @@ def _tmastg_warp_group(
         nxt_q = sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0)).load()
         nxt_hb = sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1)).load()
         nxt_v = sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2)).load()
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -771,11 +803,11 @@ def _mma_warp_group(
 
     desc_Q = sQ[0].desc()
 
-    if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+    if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV == 1):
         kv_left = cutlass.Int32(0)
         kv_right = seqlen_kv // cutlass.Int32(CFG.TILE_N)
     else:
-        q_super_idx, _hd, batch_idx = _dispatch_decode_initial(
+        q_super_idx, _hd, batch_idx, split_idx = _decode_initial_split(
             sched.bidx_init,
             sched.bidy_init,
             sched.bidz_init,
@@ -787,11 +819,14 @@ def _mma_warp_group(
             qh_per_kh,
             seqlen_kv,
         )
-        eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
-        eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds_init = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
-        kv_left = bounds_init.left
-        kv_right = bounds_init.right
+        if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+            kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+        else:
+            eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
+            eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
+            bounds_init = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
+            kv_left = bounds_init.left
+            kv_right = bounds_init.right
 
     q_full_phase = cutlass.Int32(0)
     kv_state_K = PipelineState.start(phase=0)
@@ -805,7 +840,7 @@ def _mma_warp_group(
     while is_valid_tile > cutlass.Int32(0):
         read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
 
-        if cutlass.const_expr(CFG.MASK_FLAGS != 0) and (kv_right <= kv_left):
+        if cutlass.const_expr(MAY_BE_EMPTY) and (kv_right <= kv_left):
             bars.mb_empty_mainloop.wait(empty_mainloop_phase)
             empty_mainloop_phase = empty_mainloop_phase ^ cutlass.Int32(1)
             elect_p = nvvm.elect_sync()
@@ -913,14 +948,14 @@ def _mma_warp_group(
         nvvm.bar_warp_sync(cute.arch.FULL_MASK)
 
         wait(sched.mb_scheduler.subview(sched_state.idx), sched_state.phase)
-        if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+        if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV == 1):
             nxt_v = sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2)).load()
             is_valid_tile = nxt_v & cutlass.Int32(1)
         else:
             nxt_q = cute.arch.make_warp_uniform(sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0)).load())
             nxt_hb = cute.arch.make_warp_uniform(sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1)).load())
             nxt_v = cute.arch.make_warp_uniform(sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2)).load())
-            q_super_idx, _hd, batch_idx = _dispatch_decode_payload(
+            q_super_idx, _hd, batch_idx, split_idx = _decode_payload_split(
                 nxt_q,
                 nxt_hb,
                 cta_in_pair,
@@ -932,11 +967,14 @@ def _mma_warp_group(
                 seqlen_kv,
             )
             is_valid_tile = nxt_v & cutlass.Int32(1)
-            eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
-            eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-            bounds_next = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
-            kv_left = bounds_next.left
-            kv_right = bounds_next.right
+            if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+                kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+            else:
+                eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
+                eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
+                bounds_next = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
+                kv_left = bounds_next.left
+                kv_right = bounds_next.right
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
 
     bars.mb_tmem_dealloc.wait(cutlass.Int32(0))
@@ -970,7 +1008,7 @@ def _softmax_warp_group(
 
     NEG_INF = cutlass.Float32(float("-inf"))
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -988,7 +1026,7 @@ def _softmax_warp_group(
     eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
 
     eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-    bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+    bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
 
     tid_in_wg = cute.arch.thread_idx()[0] - cutlass.Int32(CFG.SOFTMAX_WG0_BASE * 32)
 
@@ -1338,7 +1376,7 @@ def _softmax_warp_group(
         nxt_q = cute.arch.make_warp_uniform(sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0)).load())
         nxt_hb = cute.arch.make_warp_uniform(sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1)).load())
         nxt_v = cute.arch.make_warp_uniform(sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2)).load())
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -1353,7 +1391,7 @@ def _softmax_warp_group(
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
         eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
         eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+        bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
 
 
 @cute.jit
@@ -1386,7 +1424,7 @@ def _correction_warp_group(
     stat_mbar_state = cutlass.Int32(0)
     epilogue_state = cutlass.Int32(1)
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -1404,7 +1442,7 @@ def _correction_warp_group(
     eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
 
     eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-    bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+    bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
 
     O_CHUNK = 16
     N_CHUNKS_O = CFG.TILE_O // O_CHUNK
@@ -1552,7 +1590,11 @@ def _correction_warp_group(
         else:
             if q_row_global < seqlen_q:
                 lse_arr = cutlass.make_array_view(lse_tensor)
-                lse_row = lse_arr[batch_idx, head_idx, :]
+                # This chunk's LSE goes to its own split-major slot, matching where
+                # TMA-STG put the chunk's O.  The pair (O_s, lse_s) is everything
+                # the combine needs.
+                lse_batch = _partial_batch(batch_idx, split_idx, n_batch)
+                lse_row = lse_arr[lse_batch, head_idx, :]
                 lse_row[q_row_global] = lse_val
 
         parity_last_rt = cutlass.Int32(0)
@@ -1602,7 +1644,7 @@ def _correction_warp_group(
         nxt_q = sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0)).load()
         nxt_hb = sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1)).load()
         nxt_v = sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2)).load()
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -1617,7 +1659,7 @@ def _correction_warp_group(
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
         eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
         eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+        bounds = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
 
     if cutlass.const_expr(CFG.CTA_MMA == 2):
         peer_cta = cta_id_x ^ cutlass.Int32(1)
@@ -1720,7 +1762,12 @@ def _tma_swz(byte_w: int):
         ).launch(grid=(1, 1, 1), block=(32, 1, 1), stream=stream)
         grid_shape = (n_thd_units * cutlass.Int32(CFG.CGA_M), cutlass.Int32(1), cutlass.Int32(1))
     else:
-        grid_shape = (grid_q_supers, QH, B) if cutlass.const_expr(CFG.SCHEDULER_POLICY == SCHED_NATURAL) else (grid_q_supers * QH * B, 1, 1)
+        # KV split rides the BATCH axis: z = batch + split*B.  The decode
+        # already recovers the batch coord on both the blockIdx and the
+        # scheduler-handout paths, so the split travels with it for free.
+        grid_shape = (
+            (grid_q_supers, QH, B * SPLIT_KV) if cutlass.const_expr(CFG.SCHEDULER_POLICY == SCHED_NATURAL) else (grid_q_supers * QH * B * SPLIT_KV, 1, 1)
+        )
     _kernel(
         tma_q_desc,
         tma_k_desc,
@@ -1779,6 +1826,10 @@ def compile(  # noqa: A001
         raise ValueError(f"d256 envelope: need 0 < d_qk <= {CFG.TILE_K} and 0 < d_v <= {CFG.TILE_O}; got ({d_qk}, {d_v})")
     if (d_qk * CFG.BPE) % 16 != 0 or (d_v * CFG.BPE_O) % 16 != 0:
         raise ValueError(f"d256 envelope: d_qk*BPE and d_v*BPE must be 16-byte multiples (TMA global-stride rule); got ({d_qk}, {d_v}) at BPE={CFG.BPE}")
+    if SPLIT_KV > 1 and not has_lse:
+        # Each split's LSE is not optional under KV split — it IS the weight
+        # the combine reduces with.  Without it the partials cannot be recombined.
+        raise ValueError("split_kv > 1 requires has_lse=True (the per-split LSE drives the combine)")
     _fake_batch = 1 if CFG.THD_VARLEN else b
     if CFG.THD_VARLEN:
         # Dynamic packed token totals: one symbol per ragged group (Q/O and
@@ -1786,6 +1837,10 @@ def compile(  # noqa: A001
         # compiled artifact instead of minting a new one (issue #552).
         sq = cute.sym_int(divisibility=1)
         skv = cute.sym_int(divisibility=1)
+    # KV split: O and LSE are the PARTIAL workspaces, stacked split-major on
+    # the batch axis (B*SPLIT_KV).  Q/K/V keep the real batch.
+    _o_batch = _fake_batch * SPLIT_KV
+    _lse_batch = b * SPLIT_KV
 
     def _fake_bshd(shape, stride, dtype=STORAGE_DTYPE, bpe=CFG.BPE):
         if stride is None:
@@ -1804,7 +1859,7 @@ def _fake_bshd(shape, stride, dtype=STORAGE_DTYPE, bpe=CFG.BPE):
     fake_q = _fake_bshd((_fake_batch, sq, qh, d_qk), q_stride)
     fake_k = _fake_bshd((_fake_batch, skv, kh, d_qk), k_stride)
     fake_v = _fake_bshd((_fake_batch, skv, kh, d_v), v_stride)
-    fake_o = _fake_bshd((_fake_batch, sq, qh, d_v), o_stride, dtype=OUT_STORAGE_DTYPE, bpe=CFG.BPE_O)
+    fake_o = _fake_bshd((_o_batch, sq, qh, d_v), o_stride, dtype=OUT_STORAGE_DTYPE, bpe=CFG.BPE_O)
     if not has_lse:
         # No Stats output: the LSE argument is None-specialized and the store
         # is compiled out entirely — no dummy buffer exists at any level.
@@ -1844,7 +1899,7 @@ def _fake_bshd(shape, stride, dtype=STORAGE_DTYPE, bpe=CFG.BPE):
             raise ValueError("lse_head_major / lse_head_stride are THD-only (dense LSE is compact (B, H, Sq))")
         fake_lse = cute.runtime.make_fake_compact_tensor(
             cutlass.Float32,
-            (b, qh, sq),
+            (_lse_batch, qh, sq),
             stride_order=(2, 1, 0),
             assumed_align=16,
         )
diff --git a/python/cudnn/sdpa/fwd/kernels/prefill_d512_f16_sm100.py b/python/cudnn/sdpa/fwd/kernels/prefill_d512_f16_sm100.py
index 3eea02ac0..78c914d7d 100644
--- a/python/cudnn/sdpa/fwd/kernels/prefill_d512_f16_sm100.py
+++ b/python/cudnn/sdpa/fwd/kernels/prefill_d512_f16_sm100.py
@@ -91,6 +91,7 @@ def _require(cond, msg):
 )
 
 from cudnn.sdpa.fwd.kernels._common_sm100 import (
+    make_split_helpers,
     KvLoopBounds,
     compute_kv_loop_bounds,
     row_max_for_exp2,
@@ -322,6 +323,29 @@ class KernelTmemLayout:
 _thd_tma_offsets = _sdpa_h.thd_tma_offsets
 
 
+# === KV split ===
+#
+# Mechanics live in _common_sm100.make_split_helpers, shared with the other
+# SM100 prefill flavors: each Q tile's KV loop range is cut into SPLIT_KV
+# contiguous chunks, each run as its own persistent tile, and each writing a
+# normalized partial O + its own LSE into a split-major workspace that
+# split_combine_sm100 folds with the exact log-sum-exp identity.  At
+# SPLIT_KV == 1 every closure folds away and this is the classic kernel.
+_split_h = make_split_helpers(
+    CFG,
+    bounds_for_tile=_bounds_for_tile,
+    dispatch_decode_initial=_dispatch_decode_initial,
+    dispatch_decode_payload=_dispatch_decode_payload,
+)
+SPLIT_KV = _split_h.SPLIT_KV
+MAY_BE_EMPTY = _split_h.MAY_BE_EMPTY
+_decode_initial_split = _split_h.decode_initial_split
+_decode_payload_split = _split_h.decode_payload_split
+_bounds_for_tile_split = _split_h.bounds_for_tile_split
+_nomask_range_split = _split_h.nomask_range_split
+_partial_batch = _split_h.partial_batch
+
+
 @cute.kernel
 def _kernel(
     tma_q_desc: cutlass.GridConstant[tmap.TensorMap],
@@ -807,7 +831,7 @@ def _compute_warp_group(
     wid_in_wg = tid_in_wg // cutlass.Int32(32)
     is_lead_warp = wid_in_wg == cutlass.Int32(0)
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -822,16 +846,22 @@ def _compute_warp_group(
     is_valid_tile = cutlass.Int32(1)
     sched_state = PipelineState.start()
 
-    if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+    if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV == 1):
         kv_left = cutlass.Int32(0)
         kv_unmasked_lo = cutlass.Int32(0)
         kv_unmasked_hi = seqlen_kv // cutlass.Int32(CFG.TILE_N)
         kv_right = seqlen_kv // cutlass.Int32(CFG.TILE_N)
         eff_seqlen_kv = seqlen_kv
+    elif cutlass.const_expr(CFG.MASK_FLAGS == 0):
+        # Unmasked, so this split's whole slice is the unmasked band.
+        kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+        kv_unmasked_lo = kv_left
+        kv_unmasked_hi = kv_right
+        eff_seqlen_kv = seqlen_kv
     else:
         eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
         eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds_init = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+        bounds_init = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
         kv_left = bounds_init.left
         kv_unmasked_lo = bounds_init.unmasked_lo
         kv_unmasked_hi = bounds_init.unmasked_hi
@@ -875,7 +905,7 @@ def _compute_warp_group(
 
             q_abs = q_super_idx * cutlass.Int32(CFG.TILES_Q * CFG.TILE_M) + tid_in_wg
 
-            if cutlass.const_expr(CFG.MASK_FLAGS != 0) and (kv_right <= kv_left):
+            if cutlass.const_expr(MAY_BE_EMPTY) and (kv_right <= kv_left):
                 pass
             else:
                 if cutlass.const_expr(CFG.MASK_FLAGS == 0):
@@ -995,7 +1025,7 @@ def _compute_warp_group(
         else:
             tmem_O_base = tmem_base_addr + cutlass.Int32(LAYOUT.O_OFF)
 
-            if cutlass.const_expr(CFG.MASK_FLAGS != 0) and (kv_right <= kv_left):
+            if cutlass.const_expr(MAY_BE_EMPTY) and (kv_right <= kv_left):
                 bars.mb_empty_mainloop.arrive_on_peer(leader_cta_id, pred=is_lead_warp & nvvm.elect_sync())
             else:
                 cur_parity_0 = alpha_full_state.idx
@@ -1178,14 +1208,15 @@ def _compute_warp_group(
             else:
                 if q_row_global < seqlen_q:
                     lse_arr = cutlass.make_array_view(lse_tensor)
-                    lse_row = lse_arr[batch_idx, head_idx, :]
+                    lse_batch = _partial_batch(batch_idx, split_idx, n_batch)
+                    lse_row = lse_arr[lse_batch, head_idx, :]
                     lse_row[q_row_global] = lse
 
         wait(sched.mb_scheduler.subview(sched_state.idx), sched_state.phase)
         nxt_q = cute.arch.make_warp_uniform(sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0)).load())
         nxt_hb = cute.arch.make_warp_uniform(sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1)).load())
         nxt_v = cute.arch.make_warp_uniform(sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2)).load())
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -1198,10 +1229,12 @@ def _compute_warp_group(
         )
         is_valid_tile = nxt_v & cutlass.Int32(1)
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
-        if cutlass.const_expr(CFG.MASK_FLAGS != 0):
+        if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV > 1):
+            kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+        elif cutlass.const_expr(CFG.MASK_FLAGS != 0):
             eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
             eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-            bounds_next = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+            bounds_next = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
             kv_left = bounds_next.left
             kv_unmasked_lo = bounds_next.unmasked_lo
             kv_unmasked_hi = bounds_next.unmasked_hi
@@ -1293,7 +1326,7 @@ def _mma_warp_group(
         kind=MMA_KIND,
     )
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -1308,13 +1341,15 @@ def _mma_warp_group(
     is_valid_tile = cutlass.Int32(1)
     sched_state = PipelineState.start()
 
-    if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+    if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV == 1):
         kv_left = cutlass.Int32(0)
         kv_right = seqlen_kv // cutlass.Int32(CFG.TILE_N)
+    elif cutlass.const_expr(CFG.MASK_FLAGS == 0):
+        kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
     else:
         eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
         eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds_init = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+        bounds_init = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
         kv_left = bounds_init.left
         kv_right = bounds_init.right
 
@@ -1351,7 +1386,7 @@ def _mma_warp_group(
         read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
 
         if is_sg0:
-            if cutlass.const_expr(CFG.MASK_FLAGS != 0) and (kv_right <= kv_left):
+            if cutlass.const_expr(MAY_BE_EMPTY) and (kv_right <= kv_left):
                 pass
             else:
                 bars.mb_tma_q_full.wait(q_full_state.phase)
@@ -1386,7 +1421,7 @@ def _mma_warp_group(
 
                 bars.mb_tma_q_empty.arrive(cta_group=CFG.CTA_MMA, mcast_mask=mcast_mask, pred=nvvm.elect_sync())
         else:
-            if cutlass.const_expr(CFG.MASK_FLAGS != 0) and (kv_right <= kv_left):
+            if cutlass.const_expr(MAY_BE_EMPTY) and (kv_right <= kv_left):
                 bars.mb_empty_mainloop.wait(empty_mainloop_state.phase)
                 empty_mainloop_state = advance(empty_mainloop_state, 1)
                 bars.mb_bmm2_done[bmm2_done_prod_state.idx].arrive(cta_group=CFG.CTA_MMA, mcast_mask=mcast_mask, pred=nvvm.elect_sync())
@@ -1447,7 +1482,7 @@ def _mma_warp_group(
         nxt_q = sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0)).load()
         nxt_hb = sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1)).load()
         nxt_v = sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2)).load()
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -1460,10 +1495,12 @@ def _mma_warp_group(
         )
         is_valid_tile = nxt_v & cutlass.Int32(1)
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
-        if cutlass.const_expr(CFG.MASK_FLAGS != 0):
+        if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV > 1):
+            kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+        elif cutlass.const_expr(CFG.MASK_FLAGS != 0):
             eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
             eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-            bounds_next = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+            bounds_next = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
             kv_left = bounds_next.left
             kv_right = bounds_next.right
 
@@ -1503,7 +1540,7 @@ def _mma_warp_non_leader(
     _UTCCP_N_CALLS = LAYOUT.Q_TMEM_COLS // _UTCCP_TMEM_COLS_PER_CALL
 
     if is_sg1:
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
             sched.bidx_init,
             sched.bidy_init,
             sched.bidz_init,
@@ -1518,13 +1555,15 @@ def _mma_warp_non_leader(
         is_valid_tile = cutlass.Int32(1)
         sched_state = PipelineState.start()
 
-        if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+        if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV == 1):
             kv_left = cutlass.Int32(0)
             kv_right = seqlen_kv // cutlass.Int32(CFG.TILE_N)
+        elif cutlass.const_expr(CFG.MASK_FLAGS == 0):
+            kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
         else:
             eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
             eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-            bounds_init = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+            bounds_init = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
             kv_left = bounds_init.left
             kv_right = bounds_init.right
 
@@ -1533,7 +1572,7 @@ def _mma_warp_non_leader(
         while is_valid_tile > cutlass.Int32(0):
             read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
 
-            if cutlass.const_expr(CFG.MASK_FLAGS != 0) and (kv_right <= kv_left):
+            if cutlass.const_expr(MAY_BE_EMPTY) and (kv_right <= kv_left):
                 pass
             else:
                 for _kv in cutlass.range(kv_left, kv_right, 1, unroll=1):
@@ -1550,7 +1589,7 @@ def _mma_warp_non_leader(
             nxt_q = sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0)).load()
             nxt_hb = sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1)).load()
             nxt_v = sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2)).load()
-            q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+            q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
                 nxt_q,
                 nxt_hb,
                 cta_in_pair,
@@ -1563,10 +1602,12 @@ def _mma_warp_non_leader(
             )
             is_valid_tile = nxt_v & cutlass.Int32(1)
             sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
-            if cutlass.const_expr(CFG.MASK_FLAGS != 0):
+            if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV > 1):
+                kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+            elif cutlass.const_expr(CFG.MASK_FLAGS != 0):
                 eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
                 eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-                bounds_next = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+                bounds_next = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
                 kv_left = bounds_next.left
                 kv_right = bounds_next.right
     bars.mb_tmem_dealloc.wait(cutlass.Int32(0))
@@ -1606,7 +1647,7 @@ def _tmaldg_warp_group(
     tma_k = GmemTileTma(tma_k_desc)
     tma_v = GmemTileTma(tma_v_desc)
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -1622,13 +1663,15 @@ def _tmaldg_warp_group(
     q_row_base = cute.arch.make_warp_uniform(q_super_idx * cutlass.Int32(CFG.TILES_Q * CFG.TILE_M))
     q_seq_off, kv_seq_off, tma_batch = _thd_tma_offsets(seq_kv_lens_tensor, batch_idx, n_batch)
 
-    if cutlass.const_expr(CFG.MASK_FLAGS == 0):
+    if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV == 1):
         kv_left = cutlass.Int32(0)
         kv_right = seqlen_kv // cutlass.Int32(CFG.TILE_N)
+    elif cutlass.const_expr(CFG.MASK_FLAGS == 0):
+        kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
     else:
         eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
         eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-        bounds_init = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+        bounds_init = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
         kv_left = bounds_init.left
         kv_right = bounds_init.right
 
@@ -1641,7 +1684,7 @@ def _tmaldg_warp_group(
     while is_valid_tile > cutlass.Int32(0):
         read_tile_id_arrive(sched.mb_read_tile_id.subview(sched_state.idx), CGA_SIZE)
 
-        if cutlass.const_expr(CFG.MASK_FLAGS != 0) and (kv_right <= kv_left):
+        if cutlass.const_expr(MAY_BE_EMPTY) and (kv_right <= kv_left):
             # Empty KV loop (SWA window past the padded KV tail, or a dead-Q-tile
             # collapse): no loads — but the O∪V SMEM alias gate must stay in
             # phase.  The sg1 compute epilogue stores (trimmed) O and TMA-STG
@@ -1705,7 +1748,7 @@ def _tmaldg_warp_group(
         nxt_q = cute.arch.make_warp_uniform(sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0)).load())
         nxt_hb = cute.arch.make_warp_uniform(sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1)).load())
         nxt_v = cute.arch.make_warp_uniform(sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2)).load())
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -1721,10 +1764,12 @@ def _tmaldg_warp_group(
         q_seq_off, kv_seq_off, tma_batch = _thd_tma_offsets(seq_kv_lens_tensor, batch_idx, n_batch)
         is_valid_tile = nxt_v & cutlass.Int32(1)
         sched_state = advance(sched_state, CFG.SCHEDULER_STAGES)
-        if cutlass.const_expr(CFG.MASK_FLAGS != 0):
+        if cutlass.const_expr(CFG.MASK_FLAGS == 0 and SPLIT_KV > 1):
+            kv_left, kv_right = _nomask_range_split(seqlen_kv, split_idx)
+        elif cutlass.const_expr(CFG.MASK_FLAGS != 0):
             eff_seqlen_kv = _resolve_seqlen_kv(seq_kv_lens_tensor, batch_idx, seqlen_kv)
             eff_seqlen_q = _resolve_seqlen_q(seq_kv_lens_tensor, batch_idx, seqlen_q, n_batch)
-            bounds_next = _bounds_for_tile(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx)
+            bounds_next = _bounds_for_tile_split(q_super_idx, eff_seqlen_q, eff_seqlen_kv, cta_in_pair, seq_q_lens_tensor, batch_idx, split_idx)
             kv_left = bounds_next.left
             kv_right = bounds_next.right
 
@@ -1761,7 +1806,7 @@ def _tmastg_warp_group(
 
     tma_o = GmemTileTma(tma_o_desc)
 
-    q_super_idx, head_idx, batch_idx = _dispatch_decode_initial(
+    q_super_idx, head_idx, batch_idx, split_idx = _decode_initial_split(
         sched.bidx_init,
         sched.bidy_init,
         sched.bidz_init,
@@ -1784,6 +1829,8 @@ def _tmastg_warp_group(
                 bars.mb_tma_o_full[chunk].wait(o_full_state.phase)
 
             q_row_coord = q_super_idx * cutlass.Int32(CFG.TILES_Q * CFG.TILE_M)
+            # KV split: partials stack split-major on the workspace BATCH axis.
+            o_batch = _partial_batch(batch_idx, split_idx, n_batch)
             if cutlass.const_expr(CFG.THD_VARLEN):
                 # DEAD unit (batch == n_batch, envelope grid — issue #552): no O
                 # rows exist and descriptor slot n_batch is never built, so skip
@@ -1795,7 +1842,7 @@ def _tmastg_warp_group(
             else:
                 tma_store_tile(
                     sO[0],
-                    tma_o(cutlass.Int32(0), head_idx, q_row_coord, batch_idx),
+                    tma_o(cutlass.Int32(0), head_idx, q_row_coord, o_batch),
                 )
             tma_store_commit()
             tma_store_wait(0)
@@ -1809,7 +1856,7 @@ def _tmastg_warp_group(
         nxt_q = sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(0)).load()
         nxt_hb = sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(1)).load()
         nxt_v = sched.tile_id_smem.subview(sched_state.idx * cutlass.Int32(8) + cutlass.Int32(2)).load()
-        q_super_idx, head_idx, batch_idx = _dispatch_decode_payload(
+        q_super_idx, head_idx, batch_idx, split_idx = _decode_payload_split(
             nxt_q,
             nxt_hb,
             cta_in_pair,
@@ -1919,7 +1966,10 @@ def _tma_swz(byte_w: int):
         ).launch(grid=(1, 1, 1), block=(32, 1, 1), stream=stream)
         grid_shape = (n_thd_units * cutlass.Int32(CFG.CGA_M), cutlass.Int32(1), cutlass.Int32(1))
     else:
-        grid_shape = (grid_q_supers, QH, B) if cutlass.const_expr(CFG.SCHEDULER_POLICY == SCHED_NATURAL) else (grid_q_supers * QH * B, 1, 1)
+        # KV split rides the BATCH axis: z = batch + split*B.
+        grid_shape = (
+            (grid_q_supers, QH, B * SPLIT_KV) if cutlass.const_expr(CFG.SCHEDULER_POLICY == SCHED_NATURAL) else (grid_q_supers * QH * B * SPLIT_KV, 1, 1)
+        )
     _kernel(
         tma_q_desc,
         tma_k_desc,
@@ -1980,6 +2030,8 @@ def compile(  # noqa: A001
         raise ValueError(f"d512 envelope: need 0 < d_qk <= {CFG.TILE_K} and 0 < d_v <= {CFG.TILE_O}; got ({d_qk}, {d_v})")
     if (d_qk * CFG.BPE) % 16 != 0 or (d_v * CFG.BPE_O) % 16 != 0:
         raise ValueError(f"d512 envelope: d_qk*BPE and d_v*BPE must be 16-byte multiples (TMA global-stride rule); got ({d_qk}, {d_v}) at BPE={CFG.BPE}")
+    if SPLIT_KV > 1 and not has_lse:
+        raise ValueError("split_kv > 1 requires has_lse=True (the per-split LSE drives the combine)")
     _fake_batch = 1 if CFG.THD_VARLEN else b
     if CFG.THD_VARLEN:
         # Dynamic packed token totals: one symbol per ragged group (Q/O and
@@ -1987,6 +2039,8 @@ def compile(  # noqa: A001
         # compiled artifact instead of minting a new one (issue #552).
         sq = cute.sym_int(divisibility=1)
         skv = cute.sym_int(divisibility=1)
+    _o_batch = _fake_batch * SPLIT_KV
+    _lse_batch = b * SPLIT_KV
 
     def _fake_bshd(shape, stride, dtype=STORAGE_DTYPE, bpe=CFG.BPE):
         if stride is None:
@@ -2005,7 +2059,7 @@ def _fake_bshd(shape, stride, dtype=STORAGE_DTYPE, bpe=CFG.BPE):
     fake_q = _fake_bshd((_fake_batch, sq, qh, d_qk), q_stride)
     fake_k = _fake_bshd((_fake_batch, skv, kh, d_qk), k_stride)
     fake_v = _fake_bshd((_fake_batch, skv, kh, d_v), v_stride)
-    fake_o = _fake_bshd((_fake_batch, sq, qh, d_v), o_stride, dtype=OUT_STORAGE_DTYPE, bpe=CFG.BPE_O)
+    fake_o = _fake_bshd((_o_batch, sq, qh, d_v), o_stride, dtype=OUT_STORAGE_DTYPE, bpe=CFG.BPE_O)
     if not has_lse:
         # No Stats output: the LSE argument is None-specialized and the store
         # is compiled out entirely — no dummy buffer exists at any level.
@@ -2045,7 +2099,7 @@ def _fake_bshd(shape, stride, dtype=STORAGE_DTYPE, bpe=CFG.BPE):
             raise ValueError("lse_head_major / lse_head_stride are THD-only (dense LSE is compact (B, H, Sq))")
         fake_lse = cute.runtime.make_fake_compact_tensor(
             cutlass.Float32,
-            (b, qh, sq),
+            (_lse_batch, qh, sq),
             stride_order=(2, 1, 0),
             assumed_align=16,
         )
diff --git a/python/cudnn/sdpa/fwd/kernels/prefill_f16_sm120.py b/python/cudnn/sdpa/fwd/kernels/prefill_f16_sm120.py
index 545648226..4a4016bf0 100644
--- a/python/cudnn/sdpa/fwd/kernels/prefill_f16_sm120.py
+++ b/python/cudnn/sdpa/fwd/kernels/prefill_f16_sm120.py
@@ -210,6 +210,7 @@ def __init__(
         seq_kv_lens_present: bool = False,
         has_sink: bool = False,
         thd_varlen: bool = False,
+        split_kv: int = 1,
         thd_batch: int = 1,
         thd_lse_head_major: bool = False,
         head_tile_qk: int = 128,
@@ -282,6 +283,9 @@ def __init__(
         self.head_tile_v = head_tile_v
         self.q_tile = q_tile
         self.kv_tile = kv_tile
+        # KV split: SPLIT_KV CTAs per (q_tile, batch, head), each covering a
+        # contiguous slice of that tile's KV-tile range.  1 = off (folds away).
+        self.split_kv = split_kv
 
         # Warp roles
         if self.q_tile == 128:
@@ -831,6 +835,18 @@ def kernel(
         """
         tidx, _, _ = cute.arch.thread_idx()
         q_tile_idx, batch_idx, head_idx = cute.arch.block_idx()
+        # KV split rides the BATCH axis: grid.y = batch + split*B.  Q/K/V and the
+        # per-batch seqlens must use the REAL batch; only the O/LSE partial slot
+        # uses the composite.  Folds away entirely at split_kv == 1.  NATURAL-only
+        # (config_sm120 enforces it): LPT / LPT_L2 flatten the grid to 1-D and
+        # derive the batch from the linear tile id, leaving no axis to ride.
+        split_idx = cutlass.Int32(0)
+        o_batch_idx = batch_idx
+        if cutlass.const_expr(self.split_kv > 1):
+            n_batch_real = cutlass.Int32(q.shape[0])
+            split_idx = batch_idx // n_batch_real
+            batch_idx = batch_idx % n_batch_real
+            o_batch_idx = split_idx * n_batch_real + batch_idx
         if cutlass.const_expr(self.sched_policy != SCHED_NATURAL):
             _n_qh = cutlass.Int32(q.shape[2])
             _n_batch = cutlass.Int32(self.thd_batch if cutlass.const_expr(self.thd_varlen) else q.shape[0])
@@ -850,6 +866,7 @@ def kernel(
                 )
             else:
                 q_tile_idx, head_idx, batch_idx = lpt_tile_coords(q_tile_idx, _n_qh, _n_batch, _q_tiles)
+            o_batch_idx = batch_idx
         elif cutlass.const_expr(self.is_causal):
             # Diagonal-bounded work grows with the Q tile. Launch long tiles
             # first to avoid leaving a few expensive CTAs in the final
@@ -908,7 +925,7 @@ def kernel(
             o_head_off = q_row_base * o_seq_stride + head_idx * o_head_stride
         else:
             q_head_off = batch_idx * q_batch_stride + head_idx * q_head_stride
-            o_head_off = batch_idx * o_batch_stride + head_idx * o_head_stride
+            o_head_off = o_batch_idx * o_batch_stride + head_idx * o_head_stride
         kv_head_idx = head_idx // (num_heads_q // num_heads_kv)
 
         num_kv_tiles = ceil_div(seqlen_k, self.kv_tile)
@@ -934,6 +951,20 @@ def kernel(
                 first_q_position += seqlen_k - seqlen_q
             first_valid_col = cute.math.max(cutlass.Int32(0), first_q_position - self.window_size_left)
             min_kv_tile = first_valid_col // self.kv_tile
+        if cutlass.const_expr(self.split_kv > 1):
+            # Cut the ALREADY-masked [min_kv_tile, num_kv_tiles) into SPLIT_KV
+            # near-equal chunks; the first `rem` splits take one extra tile so
+            # the slowest split (which sets the critical path) is minimal.  A
+            # split past the end collapses to lo == hi, which drives has_kv_work
+            # false below -- the epilogue then produces row_sum = 0, i.e.
+            # O := 0 / LSE := -inf, the identity of the combine's log-sum-exp.
+            _span = num_kv_tiles - min_kv_tile
+            _per = _span // cutlass.Int32(self.split_kv)
+            _rem = _span % cutlass.Int32(self.split_kv)
+            _lo = min_kv_tile + split_idx * _per + cute.math.min(split_idx, _rem)
+            _extra = cutlass.Int32(1) if split_idx < _rem else cutlass.Int32(0)
+            min_kv_tile = _lo
+            num_kv_tiles = _lo + _per + _extra
         has_kv_work = num_kv_tiles > 0 and (num_kv_tiles - 1) >= min_kv_tile
 
         # Shared-memory layout:
@@ -1161,7 +1192,13 @@ def kernel(
                         )
                     kv_tile_idx -= 1
             else:
-                if kv_tile_idx >= 0:
+                # Guard against the split's START, not 0. Without KV split
+                # min_kv_tile is 0 on this branch so the two agree, but a split
+                # begins partway into the KV range -- and the load warp bounds
+                # its loop by min_kv_tile, so testing >= 0 here makes the compute
+                # warps run one tile the loader never fetches and the CTA
+                # deadlocks on the TMA barrier.
+                if kv_tile_idx >= min_kv_tile:
                     self.compute_one_kv_tile(
                         basic_params,
                         mma_params,
@@ -1236,7 +1273,7 @@ def kernel(
                             if lse_q_idx >= seqlen_q:
                                 lse_out = -cutlass.Float32.inf
                             if lse_q_idx < q.shape[1]:
-                                lse_row = lse_arr[batch_idx, head_idx, :]
+                                lse_row = lse_arr[o_batch_idx, head_idx, :]
                                 lse_row[lse_q_idx] = lse_out
 
             prims.barrier_cta_sync(self.bar_compute_sync, thread_count=self.threads_compute)
@@ -1372,10 +1409,13 @@ def __call__(
         def _static_neq(a, b):
             return isinstance(a, int) and isinstance(b, int) and a != b
 
+        # Under KV split, O is the split-major PARTIAL workspace: its batch mode
+        # is B*SPLIT_KV while Q/K/V keep the real batch, so O's batch is checked
+        # against that multiple rather than against Q's.
         if cutlass.const_expr(
             _static_neq(q.shape[0], k.shape[0])
             or any(_static_neq(a, b) for a, b in zip(k.shape[:3], v.shape[:3]))
-            or _static_neq(q.shape[0], o.shape[0])
+            or _static_neq(q.shape[0] * self.split_kv, o.shape[0])
             or _static_neq(q.shape[1], o.shape[1])
             or _static_neq(q.shape[2], o.shape[2])
             or q.shape[2] % k.shape[2] != 0
@@ -1405,8 +1445,10 @@ def _static_neq(a, b):
                     if cutlass.const_expr(lse.stride != (q.shape[2], 1)):
                         raise ValueError("THD LSE must be compact token-major")
             else:
-                if cutlass.const_expr(lse.shape != (q.shape[0], q.shape[2], q.shape[1])):
-                    raise ValueError("LSE must have shape (B, H, Sq)")
+                # Under KV split the LSE is the split-major partial workspace,
+                # batch mode B*SPLIT_KV (same as O).
+                if cutlass.const_expr(lse.shape != (q.shape[0] * self.split_kv, q.shape[2], q.shape[1])):
+                    raise ValueError("LSE must have shape (B * split_kv, H, Sq)")
                 if cutlass.const_expr(lse.stride != (q.shape[2] * q.shape[1], q.shape[1], 1)):
                     raise ValueError("LSE must be compact row-major")
         if cutlass.const_expr(self.has_sink != (sinks is not None)):
@@ -1485,7 +1527,9 @@ def kv_tma_desc(t, head_dim, head_tile, swizzle, swizzle_chunks, swizzle_chunk_e
         if cutlass.const_expr(self.sched_policy != SCHED_NATURAL):
             grid = (n_q_tiles * n_batch * n_head, 1, 1)
         else:
-            grid = (n_q_tiles, n_batch, n_head)
+            # KV split rides the batch axis: y = batch + split*B (config_sm120
+            # allows split_kv > 1 only under NATURAL, whose 3-D grid has one).
+            grid = (n_q_tiles, n_batch * self.split_kv, n_head)
         self.kernel(
             q,
             k,
@@ -1566,7 +1610,10 @@ def compile(  # noqa: A001
         head_tile_v=round_up_head_tile(d_v),
         q_tile=PARAMS.q_tile,
         kv_tile=PARAMS.kv_tile,
+        split_kv=PARAMS.split_kv,
     )
+    if PARAMS.split_kv > 1 and not has_lse:
+        raise ValueError("SM120 SDPA: split_kv > 1 requires an LSE output (the per-split LSE drives the combine)")
     fake_batch = 1 if PARAMS.thd_varlen else b
     if PARAMS.thd_varlen:
         # Dynamic packed token totals: one symbol per ragged group (Q/O and
@@ -1574,6 +1621,10 @@ def compile(  # noqa: A001
         # compiled artifact instead of minting a new one (issue #552).
         sq = cute.sym_int(divisibility=1)
         skv = cute.sym_int(divisibility=1)
+    # KV split: O and LSE are the PARTIAL workspaces, stacked split-major on the
+    # batch axis (B*SPLIT_KV).  Q/K/V keep the real batch.
+    o_fake_batch = fake_batch * PARAMS.split_kv
+    lse_fake_batch = fake_batch * PARAMS.split_kv
 
     def _fake_bshd(shape, stride):
         if stride is None:
@@ -1587,11 +1638,11 @@ def _fake_bshd(shape, stride):
     fake_q = _fake_bshd((fake_batch, sq, qh, d_qk), q_stride)
     fake_k = _fake_bshd((fake_batch, skv, kh, d_qk), k_stride)
     fake_v = _fake_bshd((fake_batch, skv, kh, d_v), v_stride)
-    fake_o = _fake_bshd((fake_batch, sq, qh, d_v), o_stride)
+    fake_o = _fake_bshd((o_fake_batch, sq, qh, d_v), o_stride)
     if PARAMS.thd_varlen:
         fake_lse_shape = (qh, lse_head_stride) if lse_head_major else (sq, qh)
     else:
-        fake_lse_shape = (fake_batch, qh, sq)
+        fake_lse_shape = (lse_fake_batch, qh, sq)
     fake_lse = (
         cute.runtime.make_fake_compact_tensor(
             cutlass.Float32,
diff --git a/python/cudnn/sdpa/fwd/kernels/split_combine_sm100.py b/python/cudnn/sdpa/fwd/kernels/split_combine_sm100.py
new file mode 100644
index 000000000..48aa41800
--- /dev/null
+++ b/python/cudnn/sdpa/fwd/kernels/split_combine_sm100.py
@@ -0,0 +1,199 @@
+# Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: MIT
+
+"""KV-split combine (reduction) pass for the SM100 DSL prefill SDPA kernels.
+
+Each split writes ``O_s`` (normalized by its own running sum) and ``lse_s`` into
+a split-major workspace at batch coord ``b + s*B``; this pass reduces over ``s``.
+
+A split whose range came out empty ends with total_sum == 0, which the epilogue
+turns into ``O := 0 / lse := -inf`` — the identity here, so empty splits need no
+special case.
+
+One block per (q_row, head, batch); the block's threads stride over d_v, and
+each thread walks the split axis in registers.
+"""
+
+from typing import Callable, Optional, Tuple
+
+from functools import lru_cache
+
+import cutlass
+import cutlass.cute as cute
+from cutlass._mlir.dialects import arith
+from cutlass.base_dsl.typing import Pointer
+from cutlass.experimental import primitives as nvvm
+import cuda.bindings.driver as _cuda_driver  # noqa: F401  (cute.compile pulls cuda)
+
+# One block per (q_row, head, batch); 128 lanes stride over d_v.  d_v <= 128 for
+# every flavor that uses this pass today, so the stride loop runs once.
+THREADS = 128
+
+NEG_INF = float("-inf")
+
+
+@cute.kernel
+def _combine_kernel(
+    o_partial: cute.Tensor,  # [S*B, S_q, H, D] — split-major partial O
+    lse_partial: cute.Tensor,  # [S*B, H, S_q]   — split-major partial LSE
+    o_out: cute.Tensor,  # [B, S_q, H, D]
+    lse_out: Optional[cute.Tensor],  # [B, H, S_q] or None (None-specialized)
+    # 1-element fp32 amax of the RECOMBINED O.  The per-split epilogues cannot
+    # compute this: each sees only its own partial, and O is a convex
+    # combination of those, so a max over partials over-reports (~2.9x at 8
+    # splits).  The FP8 kernels therefore skip their in-kernel amax write when
+    # SPLIT_KV > 1 and leave it to this pass.  None-specialized off otherwise.
+    amax_o: Optional[cute.Tensor],
+    n_batch: cutlass.Int32,
+    n_splits: cutlass.Int32,
+    d_v: cutlass.Int32,
+) -> None:
+    tidx, _, _ = cute.arch.thread_idx()
+    q_row = cute.arch.block_idx()[0]
+    head = cute.arch.block_idx()[1]
+    batch = cute.arch.block_idx()[2]
+
+    op = cutlass.make_array_view(o_partial)
+    lp = cutlass.make_array_view(lse_partial)
+    oo = cutlass.make_array_view(o_out)
+
+    # --- pass 1: M = max_s lse_s, then den = sum_s exp(lse_s - M) ---
+    # Every lane redundantly walks the (very short) split axis; the values are
+    # block-uniform and hit L1, which is cheaper than staging them through SMEM.
+    m = cutlass.Float32(NEG_INF)
+    for s in cutlass.range(0, n_splits, 1, unroll=1):
+        lse_row = lp[batch + s * n_batch, head, :]
+        m = cute.math.max(m, cutlass.Float32(lse_row[q_row]))
+
+    # All splits dead (every row fully masked): emit O := 0 / lse := -inf rather
+    # than exp(-inf - -inf) == NaN.  m_safe only feeds the exponentials.
+    all_dead = m == cutlass.Float32(NEG_INF)
+    m_safe = cutlass.Float32(arith.select(all_dead.ir_value(), cutlass.Float32(0.0).ir_value(), m.ir_value()))
+
+    # Same reasoning as pass 2: skip dead splits rather than trusting a fastmath
+    # exp(-inf) to be exactly 0.
+    den = cutlass.Float32(0.0)
+    for s in cutlass.range(0, n_splits, 1, unroll=1):
+        lse_row = lp[batch + s * n_batch, head, :]
+        lse_s = cutlass.Float32(lse_row[q_row])
+        if lse_s > cutlass.Float32(NEG_INF):
+            den = den + cute.math.exp(lse_s - m_safe, fastmath=True)
+
+    inv_den = cutlass.Float32(1.0) / cute.math.max(den, cutlass.Float32(1e-30))
+    inv_den = cutlass.Float32(arith.select(all_dead.ir_value(), cutlass.Float32(0.0).ir_value(), inv_den.ir_value()))
+
+    # --- pass 2: O = sum_s w_s O_s / den, accumulated in fp32 ---
+    #
+    # A dead split (empty KV range) carries lse_s = -inf, so its weight is
+    # exp(-inf) == 0 and it should contribute nothing.  Relying on the ARITHMETIC
+    # to erase it is not safe: 0 * x is NaN for a non-finite x, and under
+    # fastmath the weight itself is only approximately zero.  Skip such splits
+    # outright -- they are the identity element of this reduction by
+    # construction, so branching is exact where multiplying is not.  (Observed:
+    # d512 with 5 KV tiles over 8 splits produced NaN in the recombined O
+    # without this guard, even though every partial slot held a clean
+    # -inf / 0.)
+    neg_inf = cutlass.Float32(NEG_INF)
+    amax_local = cutlass.Float32(0.0)
+    for d0 in cutlass.range(tidx, d_v, THREADS, unroll=1):
+        acc = cutlass.Float32(0.0)
+        for s in cutlass.range(0, n_splits, 1, unroll=1):
+            lse_row = lp[batch + s * n_batch, head, :]
+            lse_s = cutlass.Float32(lse_row[q_row])
+            if lse_s > neg_inf:
+                w = cute.math.exp(lse_s - m_safe, fastmath=True)
+                o_row = op[batch + s * n_batch, q_row, head, :]
+                acc = acc + w * cutlass.Float32(o_row[d0])
+        out_row = oo[batch, q_row, head, :]
+        o_val = acc * inv_den
+        out_row[d0] = o_val.to(o_out.element_type)
+        if cutlass.const_expr(amax_o is not None):
+            # amax over the fp32 PRE-CAST value, matching what the single-pass
+            # epilogue reports.
+            amax_local = cute.math.max(amax_local, cute.math.max(o_val, -o_val))
+
+    # One atomic per lane.  The value is non-negative, so its fp32 bit pattern
+    # orders the same as the float and an integer atomicMax is exact -- the same
+    # trick the kernels' own epilogues use.
+    if cutlass.const_expr(amax_o is not None):
+        _amax_ptr = Pointer(amax_o.iterator.raw_ptr(), dtype=cutlass.Int32)
+        nvvm.atomicrmw(nvvm.AtomicOp.MAX, _amax_ptr, amax_local.bitcast(cutlass.Int32))
+
+    # --- the recombined LSE (only when the caller asked for Stats) ---
+    if cutlass.const_expr(lse_out is not None):
+        if tidx == cutlass.Int32(0):
+            lo = cutlass.make_array_view(lse_out)
+            lse_val = m_safe + cute.math.log(cute.math.max(den, cutlass.Float32(1e-30)), fastmath=True)
+            lse_val = cutlass.Float32(arith.select(all_dead.ir_value(), cutlass.Float32(NEG_INF).ir_value(), lse_val.ir_value()))
+            lse_row_out = lo[batch, head, :]
+            lse_row_out[q_row] = lse_val
+
+
+@cute.jit
+def _host(
+    o_partial: cute.Tensor,
+    lse_partial: cute.Tensor,
+    o_out: cute.Tensor,
+    lse_out: Optional[cute.Tensor],
+    amax_o: Optional[cute.Tensor],
+    problem_size: Tuple[int, int, int, int],
+    n_splits: cutlass.Int32,
+    stream: _cuda_driver.CUstream = None,
+) -> None:
+    B, H, SQ, D = problem_size
+    _combine_kernel(
+        o_partial,
+        lse_partial,
+        o_out,
+        lse_out,
+        amax_o,
+        cutlass.Int32(B),
+        n_splits,
+        cutlass.Int32(D),
+    ).launch(
+        grid=(SQ, H, B),
+        block=[THREADS, 1, 1],
+        stream=stream,
+    )
+
+
+@lru_cache(maxsize=None)
+def compile(  # noqa: A001
+    b: int,
+    h: int,
+    sq: int,
+    d_v: int,
+    splits: int,
+    dtype_o: str = "f16",
+    has_lse: bool = False,
+    has_amax: bool = False,
+) -> Callable:
+    """Compile the combine pass for one concrete (B, H, S_q, d_v, splits) shape.
+
+    ``splits`` is baked into the workspace EXTENTS (batch dim ``splits*b``) but
+    passed to the kernel as a runtime count, so the split axis is a dynamic loop
+    — the pass is bandwidth-bound, so unrolling it buys nothing.  ``has_lse``
+    controls whether the recombined LSE is written at all; with ``False`` the
+    store is None-specialized out of the traced code.  ``has_amax`` does the
+    same for the FP8-family amax of the recombined O.
+    """
+    elem = {"f16": cutlass.Float16, "bf16": cutlass.BFloat16}[dtype_o]
+
+    fake_o_partial = cute.runtime.make_fake_compact_tensor(elem, (splits * b, sq, h, d_v), stride_order=(3, 2, 1, 0), assumed_align=16)
+    fake_lse_partial = cute.runtime.make_fake_compact_tensor(cutlass.Float32, (splits * b, h, sq), stride_order=(2, 1, 0), assumed_align=16)
+    fake_o_out = cute.runtime.make_fake_compact_tensor(elem, (b, sq, h, d_v), stride_order=(3, 2, 1, 0), assumed_align=16)
+    fake_lse_out = cute.runtime.make_fake_compact_tensor(cutlass.Float32, (b, h, sq), stride_order=(2, 1, 0), assumed_align=16) if has_lse else None
+    fake_amax_o = cute.runtime.make_fake_compact_tensor(cutlass.Float32, (1,), stride_order=(0,), assumed_align=4) if has_amax else None
+
+    return cute.compile(
+        _host,
+        fake_o_partial,
+        fake_lse_partial,
+        fake_o_out,
+        fake_lse_out,
+        fake_amax_o,
+        (b, h, sq, d_v),
+        cutlass.Int32(0),
+        stream=cute.runtime.make_fake_stream(use_tvm_ffi_env_stream=False),
+        options="--enable-tvm-ffi",
+    )
diff --git a/test/python/sdpa/frost/test_sdpa_fwd_split_kv_sm100.py b/test/python/sdpa/frost/test_sdpa_fwd_split_kv_sm100.py
new file mode 100644
index 000000000..62fad4b62
--- /dev/null
+++ b/test/python/sdpa/frost/test_sdpa_fwd_split_kv_sm100.py
@@ -0,0 +1,814 @@
+# Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: MIT
+
+"""KV-split tests for the SM100 d128 f16/bf16 prefill kernel.
+
+The split must change NOTHING numerically: the recombined output has to match
+both a torch fp32 reference and the unsplit kernel to fp16 rounding.
+
+Coverage: split counts that divide the KV tile count evenly and ones that do not
+(including more splits than tiles, which leaves some splits empty — the case
+that exercises the empty-mainloop / LSE = -inf identity path), dense and causal
+masks, MHA and GQA.
+"""
+
+import math
+import os
+
+import pytest
+import torch
+
+from frost_test_utils import requires_blackwell, requires_dsl
+
+pytestmark = [requires_blackwell, requires_dsl]
+
+D = 128
+TILE_N = 128
+# One cga2 cluster covers TILES_Q * TILE_M * CTA_MMA Q rows.
+ROWS_PER_CLUSTER = 2 * 128 * 2
+
+
+def _ref_sdpa(q, k, v, scale, is_causal, kh):
+    """fp32 reference over BSHD inputs; returns BSHD."""
+    if kh != q.shape[2]:
+        rep = q.shape[2] // kh
+        k = k.repeat_interleave(rep, dim=2)
+        v = v.repeat_interleave(rep, dim=2)
+    qb, kb, vb = (t.float().permute(0, 2, 1, 3) for t in (q, k, v))
+    s = torch.matmul(qb, kb.transpose(-1, -2)) * scale
+    if is_causal:
+        s_q, s_kv = qb.shape[2], kb.shape[2]
+        i = torch.arange(s_q, device=q.device).view(s_q, 1)
+        j = torch.arange(s_kv, device=q.device).view(1, s_kv)
+        s = s.masked_fill(j > i, float("-inf"))  # top-left causal
+    p = torch.softmax(s, dim=-1)
+    return torch.matmul(p, vb).permute(0, 2, 1, 3)
+
+
+def _kernel_module(splits, dtype_qkv, causal, cta_mma=2):
+    from cudnn.frost.template_loader import load_template
+    from cudnn.sdpa.fwd import api_dsl
+    from cudnn.sdpa.fwd.config_sm100 import TemplateParams
+
+    path = os.path.join(os.path.dirname(os.path.abspath(api_dsl.__file__)), "kernels", "prefill_d128_f16_sm100.py")
+    kw = {"dtype_qkv": dtype_qkv, "split_kv": splits, "cta_mma": cta_mma}
+    if causal:
+        kw["window_right"] = 0
+    return load_template(path, TemplateParams(**kw), tag=f"splitkv{splits}_d{dtype_qkv}_{'caus' if causal else 'dense'}_cga{cta_mma}")
+
+
+def _run(splits, B, H, KH, SQ, SKV, dtype, causal, cta_mma=2):
+    """Launch the split kernel + DSL combine; returns the recombined O (BSHD fp32)."""
+    import cutlass
+    import cuda.bindings.driver as cuda_driver
+
+    from cudnn.sdpa.fwd.kernels import split_combine_sm100 as comb
+
+    dev = "cuda"
+    scale = 1.0 / math.sqrt(D)
+    torch.manual_seed(0)
+    q = torch.randn(B, SQ, H, D, device=dev, dtype=dtype)
+    k = torch.randn(B, SKV, KH, D, device=dev, dtype=dtype)
+    v = torch.randn(B, SKV, KH, D, device=dev, dtype=dtype)
+
+    mod = _kernel_module(splits, 3 if dtype == torch.float16 else 2, causal, cta_mma=cta_mma)
+    fn = mod.compile(b=B, qh=H, kh=KH, sq=SQ, skv=SKV, d_qk=D, d_v=D, has_lse=True)
+
+    o_p = torch.zeros(splits * B, SQ, H, D, device=dev, dtype=dtype)
+    lse_p = torch.zeros(splits * B, H, SQ, device=dev, dtype=torch.float32)
+    stream = cuda_driver.CUstream(torch.cuda.current_stream().cuda_stream)
+    fn(
+        q,
+        k,
+        v,
+        o_p,
+        lse_p,
+        torch.zeros(H, dtype=torch.float32, device=dev),  # sinks (ABI slot)
+        torch.zeros(B, dtype=torch.int32, device=dev),  # seq_kv_lens (unused)
+        torch.zeros(1, dtype=torch.int64, device=dev),  # o_desc (THD only)
+        (B, H, KH, SQ, SKV, 0),
+        cutlass.Float32(scale * math.log2(math.e)),
+        cutlass.Int32(0),
+        None,
+        stream=stream,
+    )
+    if splits == 1:
+        torch.cuda.synchronize()
+        return o_p.float(), (q, k, v, scale)
+
+    o_out = torch.zeros(B, SQ, H, D, device=dev, dtype=dtype)
+    lse_out = torch.zeros(B, H, SQ, device=dev, dtype=torch.float32)
+    cfn = comb.compile(b=B, h=H, sq=SQ, d_v=D, splits=splits, dtype_o="f16" if dtype == torch.float16 else "bf16", has_lse=True)
+    cfn(o_p, lse_p, o_out, lse_out, None, (B, H, SQ, D), cutlass.Int32(splits), stream=stream)
+    torch.cuda.synchronize()
+    assert not torch.isnan(o_out).any(), "NaN in combined O"
+    return o_out.float(), (q, k, v, scale)
+
+
+@pytest.mark.L0
+@pytest.mark.parametrize("splits", [2, 4, 8], ids=lambda s: f"split{s}")
+def test_split_kv_matches_reference_dense(splits):
+    """The target shape class: tiny S_q against a long KV run."""
+    B, H, SQ, SKV = 1, 4, 128, 2048
+    got, (q, k, v, scale) = _run(splits, B, H, H, SQ, SKV, torch.float16, causal=False)
+    ref = _ref_sdpa(q, k, v, scale, is_causal=False, kh=H)
+    assert (got - ref).abs().max().item() <= 2e-2
+
+
+@pytest.mark.L0
+@pytest.mark.parametrize("splits", [3, 4, 8], ids=lambda s: f"split{s}")
+def test_split_kv_uneven_and_empty_splits(splits):
+    """5 KV tiles over 3/4/8 splits: unequal chunks, and at 8 three EMPTY splits.
+
+    An empty split ends with total_sum == 0, so its epilogue writes O := 0 /
+    LSE := -inf — the identity of the combine's log-sum-exp.  Getting this wrong
+    shows up as NaN or as a wrong normalization, not as a small drift.
+    """
+    B, H, SQ, SKV = 1, 4, 128, 5 * TILE_N
+    assert SKV // TILE_N == 5
+    got, (q, k, v, scale) = _run(splits, B, H, H, SQ, SKV, torch.float16, causal=False)
+    ref = _ref_sdpa(q, k, v, scale, is_causal=False, kh=H)
+    assert (got - ref).abs().max().item() <= 2e-2
+
+
+@pytest.mark.L0
+@pytest.mark.parametrize("splits", [2, 4], ids=lambda s: f"split{s}")
+def test_split_kv_causal(splits):
+    """Causal: the split must cut the MASKED range, so each split gets real work."""
+    B, H, SQ, SKV = 1, 4, 1024, 1024
+    got, (q, k, v, scale) = _run(splits, B, H, H, SQ, SKV, torch.float16, causal=True)
+    ref = _ref_sdpa(q, k, v, scale, is_causal=True, kh=H)
+    assert (got - ref).abs().max().item() <= 2e-2
+
+
+@pytest.mark.L0
+@pytest.mark.parametrize("splits", [4], ids=lambda s: f"split{s}")
+def test_split_kv_gqa(splits):
+    B, H, KH, SQ, SKV = 2, 8, 2, 128, 2048
+    got, (q, k, v, scale) = _run(splits, B, H, KH, SQ, SKV, torch.float16, causal=False)
+    ref = _ref_sdpa(q, k, v, scale, is_causal=False, kh=KH)
+    assert (got - ref).abs().max().item() <= 2e-2
+
+
+@pytest.mark.L0
+def test_split_kv_matches_unsplit():
+    """Split vs unsplit on identical inputs — the split is occupancy-only."""
+    B, H, SQ, SKV = 1, 4, 128, 2048
+    base, _ = _run(1, B, H, H, SQ, SKV, torch.float16, causal=False)
+    for splits in (2, 4, 8):
+        got, _ = _run(splits, B, H, H, SQ, SKV, torch.float16, causal=False)
+        # Both round to fp16; the split reassociates the sum, so allow one ulp
+        # of drift at fp16 magnitudes rather than demanding bit-equality.
+        assert (got - base).abs().max().item() <= 2e-2, f"split{splits} diverged from unsplit"
+
+
+@pytest.mark.L0
+def test_split_kv_bf16():
+    B, H, SQ, SKV = 1, 4, 128, 2048
+    got, (q, k, v, scale) = _run(4, B, H, H, SQ, SKV, torch.bfloat16, causal=False)
+    ref = _ref_sdpa(q, k, v, scale, is_causal=False, kh=H)
+    assert (got - ref).abs().max().item() <= 1e-1  # bf16 has ~8 mantissa bits
+
+
+@pytest.mark.L0
+def test_split_kv_rejects_unsupported_combos():
+    """The config validator is the backstop for what the combine cannot express."""
+    from cudnn.frost.tile_dsl.scheduler import SCHED_LPT
+    from cudnn.sdpa.fwd.config_sm100 import TemplateParams, make_cfg_d128
+
+    with pytest.raises(ValueError, match="split_kv"):
+        make_cfg_d128(TemplateParams(split_kv=0))
+    # Both scheduler policies are supported: NATURAL carries the split on the
+    # batch axis, LPT on the flattened x axis (make_split_helpers._lpt_split_of).
+    assert make_cfg_d128(TemplateParams(split_kv=4, sched_policy=SCHED_LPT))[0].SPLIT_KV == 4
+    with pytest.raises(ValueError, match="sink"):
+        make_cfg_d128(TemplateParams(split_kv=4, has_sink=True))
+    with pytest.raises(ValueError, match="dense-only"):
+        make_cfg_d128(TemplateParams(split_kv=4, thd_varlen=True, seq_kv_lens_present=True))
+
+
+@pytest.mark.L0
+def test_split_kv_requires_lse():
+    """has_lse=False + split is rejected: the per-split LSE IS the combine weight."""
+    mod = _kernel_module(4, 3, causal=False)
+    with pytest.raises(ValueError, match="has_lse"):
+        mod.compile(b=1, qh=4, kh=4, sq=128, skv=2048, d_qk=D, d_v=D, has_lse=False)
+
+
+# --- cga1 (CTA_MMA=1) ---------------------------------------------------
+#
+# cga1 drops the collective 2-CTA MMA: one independent CTA per tile, covering
+# TILES_Q*TILE_M = 256 Q rows instead of 512.  It halves both the wasted MMA
+# work at small S_q and the CTAs per tile, so with KV split twice as many splits
+# fit in one wave.  It is SMEM-neutral only because make_cfg_d128 turns QO_ALIAS
+# on for cga1 (no collective MMA to halve per-CTA K/V), which is the part most
+# likely to break -- hence the direct cga1-vs-cga2 output comparison below.
+
+
+@pytest.mark.L0
+@pytest.mark.parametrize("splits", [1, 4, 8], ids=lambda s: f"split{s}")
+def test_cga1_matches_reference(splits):
+    B, H, SQ, SKV = 1, 4, 128, 2048
+    got, (q, k, v, scale) = _run(splits, B, H, H, SQ, SKV, torch.float16, causal=False, cta_mma=1)
+    ref = _ref_sdpa(q, k, v, scale, is_causal=False, kh=H)
+    assert (got - ref).abs().max().item() <= 2e-2
+
+
+@pytest.mark.L0
+def test_cga1_matches_cga2():
+    """Same inputs through both cluster widths — cga1 is an occupancy change only."""
+    B, H, SQ, SKV = 1, 4, 128, 2048
+    for splits in (1, 4):
+        a, _ = _run(splits, B, H, H, SQ, SKV, torch.float16, causal=False, cta_mma=2)
+        b, _ = _run(splits, B, H, H, SQ, SKV, torch.float16, causal=False, cta_mma=1)
+        assert (a - b).abs().max().item() <= 2e-2, f"cga1 diverged from cga2 at split{splits}"
+
+
+@pytest.mark.L0
+def test_cga1_causal_and_gqa():
+    B, H, KH, SQ, SKV = 1, 8, 2, 1024, 1024
+    got, (q, k, v, scale) = _run(4, B, H, KH, SQ, SKV, torch.float16, causal=True, cta_mma=1)
+    ref = _ref_sdpa(q, k, v, scale, is_causal=True, kh=KH)
+    assert (got - ref).abs().max().item() <= 2e-2
+
+
+@pytest.mark.L0
+def test_cga1_requires_qo_alias_and_smem_fits():
+    """cga1 must enable QO_ALIAS, and both widths must stay inside the SMEM cap."""
+    from dataclasses import replace
+
+    from cudnn.sdpa.fwd.config_sm100 import TemplateParams, make_cfg_d128, _validate_cfg_d128, _d128_smem_bytes, _SM100_MAX_DYN_SMEM
+
+    cfg1, _ = make_cfg_d128(TemplateParams(cta_mma=1))
+    cfg2, _ = make_cfg_d128(TemplateParams(cta_mma=2))
+    assert cfg1.QO_ALIAS == 1 and cfg2.QO_ALIAS == 0
+    assert cfg1.CGA_M == cfg1.CTA_MMA == 1 and cfg2.CGA_M == cfg2.CTA_MMA == 2
+    # Both fit, and cga1 only fits *because* of the alias.
+    assert _d128_smem_bytes(cfg1) <= _SM100_MAX_DYN_SMEM
+    assert _d128_smem_bytes(cfg2) <= _SM100_MAX_DYN_SMEM
+    assert _d128_smem_bytes(replace(cfg1, QO_ALIAS=0)) > _SM100_MAX_DYN_SMEM
+    with pytest.raises(ValueError, match="QO_ALIAS is mandatory"):
+        _validate_cfg_d128(replace(cfg1, QO_ALIAS=0))
+    with pytest.raises(ValueError, match="cta_mma"):
+        make_cfg_d128(TemplateParams(cta_mma=3))
+
+
+@pytest.mark.L0
+def test_split_kv_and_cta_mma_flavor_gating():
+    """Flavors must REJECT knobs they do not honour, not ignore them.
+
+    ``split_kv`` / ``cta_mma`` sit on the TemplateParams shared by every SM100
+    flavor, but a flavor only honours them once its make_cfg_* threads them into
+    a Cfg AND its kernel reads them.  Silently ignoring split_kv is a
+    wrong-answer bug, not a no-op: the caller sizes an (S*B)-batch partial
+    workspace and runs the combine while the kernel writes only slots [0, B),
+    leaving the rest at lse_partial = 0 instead of -inf — weight exp(0 - M) != 0,
+    so they corrupt the reduction rather than dropping out.
+    """
+    from cudnn.sdpa.fwd.config_sm100 import (
+        _CTA_MMA_FLAVORS,
+        _SPLIT_KV_FLAVORS,
+        TemplateParams,
+        make_cfg_d128,
+        make_cfg_d192,
+        make_cfg_d256,
+        make_cfg_d512,
+    )
+
+    mk = {"d128": make_cfg_d128, "d192": make_cfg_d192, "d256": make_cfg_d256, "d512": make_cfg_d512}
+    for name, f in mk.items():
+        f(TemplateParams())  # defaults must always build
+        if name in _SPLIT_KV_FLAVORS:
+            assert f(TemplateParams(split_kv=4))[0].SPLIT_KV == 4
+        else:
+            with pytest.raises(ValueError, match="split_kv is not implemented"):
+                f(TemplateParams(split_kv=4))
+        if name in _CTA_MMA_FLAVORS:
+            assert f(TemplateParams(cta_mma=1))[0].CTA_MMA == 1
+        else:
+            with pytest.raises(ValueError, match="cta_mma is not selectable"):
+                f(TemplateParams(cta_mma=1))
+
+
+# --- empty-split coverage across every f16 flavor -------------------------
+#
+# More splits than KV tiles leaves some splits with an EMPTY range. That path is
+# what deadlocked d128 during bring-up: the empty-tile handshake
+# (mb_empty_mainloop) was compile-time gated on MASK_FLAGS != 0, so at MASK_NONE
+# the producers ran a full prologue while correction took the empty path and the
+# kernel hung. d192 additionally has its own CAN_HAVE_EMPTY_KV predicate gating
+# nine more sites. Every flavor therefore needs this exercised, not just d128 --
+# a regression here shows up as a HANG, not a wrong number.
+
+_F16_FLAVORS = {
+    "d128": ("prefill_d128_f16_sm100.py", 128, 128),
+    "d192": ("prefill_d192_d128_f16_sm100.py", 192, 128),
+    "d256": ("prefill_d256_f16_sm100.py", 256, 256),
+    "d512": ("prefill_d512_f16_sm100.py", 512, 512),
+}
+
+
+@pytest.mark.L0
+@pytest.mark.parametrize("flavor", sorted(_F16_FLAVORS))
+def test_empty_splits_every_flavor(flavor):
+    """5 KV tiles over 8 splits -> 3 empty splits, on every f16 flavor."""
+    import math as _math
+    import os as _os
+
+    import cutlass
+    import cuda.bindings.driver as cuda_driver
+
+    from cudnn.frost.template_loader import load_template
+    from cudnn.sdpa.fwd import api_dsl
+    from cudnn.sdpa.fwd.config_sm100 import TemplateParams
+    from cudnn.sdpa.fwd.kernels import split_combine_sm100 as comb
+
+    kmod, d_qk, d_v = _F16_FLAVORS[flavor]
+    B, H, SQ, SKV, S = 1, 4, 128, 5 * TILE_N, 8
+    assert SKV // TILE_N < S, "this test must leave some splits empty"
+    dev = "cuda"
+    scale = 1.0 / _math.sqrt(d_qk)
+    torch.manual_seed(0)
+
+    q = torch.randn(B, SQ, H, d_qk, device=dev, dtype=torch.float16)
+    k = torch.randn(B, SKV, H, d_qk, device=dev, dtype=torch.float16)
+    v = torch.randn(B, SKV, H, d_v, device=dev, dtype=torch.float16)
+
+    path = _os.path.join(_os.path.dirname(_os.path.abspath(api_dsl.__file__)), "kernels", kmod)
+    mod = load_template(path, TemplateParams(dtype_qkv=3, split_kv=S), tag=f"empty_{flavor}_{S}")
+    fn = mod.compile(b=B, qh=H, kh=H, sq=SQ, skv=SKV, d_qk=d_qk, d_v=d_v, has_lse=True)
+
+    o_p = torch.zeros(S * B, SQ, H, d_v, device=dev, dtype=torch.float16)
+    lse_p = torch.zeros(S * B, H, SQ, device=dev, dtype=torch.float32)
+    stream = cuda_driver.CUstream(torch.cuda.current_stream().cuda_stream)
+    fn(
+        q,
+        k,
+        v,
+        o_p,
+        lse_p,
+        torch.zeros(H, dtype=torch.float32, device=dev),
+        torch.zeros(B, dtype=torch.int32, device=dev),
+        torch.zeros(1, dtype=torch.int64, device=dev),
+        (B, H, H, SQ, SKV, 0),
+        cutlass.Float32(scale * _math.log2(_math.e)),
+        cutlass.Int32(0),
+        None,
+        stream=stream,
+    )
+    o_out = torch.zeros(B, SQ, H, d_v, device=dev, dtype=torch.float16)
+    cfn = comb.compile(b=B, h=H, sq=SQ, d_v=d_v, splits=S, dtype_o="f16", has_lse=False)
+    cfn(o_p, lse_p, o_out, None, None, (B, H, SQ, d_v), cutlass.Int32(S), stream=stream)
+    torch.cuda.synchronize()
+
+    assert not torch.isnan(o_out).any(), f"{flavor}: NaN from an empty split"
+    ref = _ref_sdpa(q, k, v, scale, is_causal=False, kh=H)
+    assert (o_out.float() - ref).abs().max().item() <= 2e-2
+
+
+# --- cga1 dtype coverage --------------------------------------------------
+#
+# The d128 flavor's config gate admits cta_mma=1 for every dtype, because
+# f16/bf16, fp8 and mxfp8 all share make_cfg_d128. That is correct for
+# f16/bf16/fp8 but NOT for mxfp8: its per-32-block E8M0 scale factors are staged
+# in SMEM by the mxfp8 kernel alone, which config_sm100._d128_smem_bytes cannot
+# see, and they push a cga1 CTA to 237024 B against the 232448 B cap. Without
+# the kernel-local guard the launch is rejected by the driver at runtime.
+
+
+@pytest.mark.L0
+@pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16], ids=["fp16", "bf16"])
+def test_cga1_half_dtypes(dtype):
+    B, H, SQ, SKV = 1, 4, 128, 2048
+    got, (q, k, v, scale) = _run(4, B, H, H, SQ, SKV, dtype, causal=False, cta_mma=1)
+    ref = _ref_sdpa(q, k, v, scale, is_causal=False, kh=H)
+    tol = 2e-2 if dtype == torch.float16 else 1e-1
+    assert (got - ref).abs().max().item() <= tol
+
+
+@pytest.mark.L0
+def test_cga1_stage_depth_scales_with_cluster_width():
+    """cga1 halves the KV stage depth for the fp8 family, as cuDNN's kernels do.
+
+    cuDNN scales the stage count with the cluster width (stages_kv = N * CTA_MMA)
+    so that stages x per-CTA-buffer -- and hence SMEM -- stays constant. FROST
+    needs the same for fp8/mxfp8: at STAGES_KV=4 a cga1 CTA asked for 237024 B
+    against the 232448 B sm_100a cap, because mxfp8 also stages E8M0 scale
+    factors that _d128_smem_bytes cannot see. f16/bf16 already fit at cga1 via
+    the Q/O alias and keep STAGES_KV=2.
+    """
+    from cudnn.frost.tile_dsl.constants import DTYPE_E4M3, DTYPE_FP16
+    from cudnn.sdpa.fwd.config_sm100 import TemplateParams, make_cfg_d128
+
+    for cta_mma, want_fp8_stages in ((2, 4), (1, 2)):
+        cfg_fp8, _ = make_cfg_d128(TemplateParams(dtype_qkv=DTYPE_E4M3, dtype_o=DTYPE_FP16, cta_mma=cta_mma))
+        assert cfg_fp8.STAGES_KV == want_fp8_stages, (cta_mma, cfg_fp8.STAGES_KV)
+        cfg_f16, _ = make_cfg_d128(TemplateParams(dtype_qkv=DTYPE_FP16, cta_mma=cta_mma))
+        assert cfg_f16.STAGES_KV == 2, (cta_mma, cfg_f16.STAGES_KV)
+    # cga1 must also turn the Q/O alias on, which is what pays for the doubled
+    # per-CTA K/V in the first place.
+    assert make_cfg_d128(TemplateParams(dtype_qkv=DTYPE_FP16, cta_mma=1))[0].QO_ALIAS == 1
+
+
+# --- fp8 / mxfp8 split coverage ------------------------------------------
+#
+# These two flavors share make_cfg_d128 with f16/bf16 but have their own ABIs
+# (amax outputs for fp8; per-32-block E8M0 scale factors for mxfp8), so they are
+# not reachable through _run above. Both cluster widths are exercised: mxfp8 at
+# cga1 in particular only fits once STAGES_KV halves with the cluster width.
+
+
+def _fp8_family_split(kfile, dtype_qkv, splits, cta_mma, mx):
+    import math as _math
+    import os as _os
+
+    import cutlass
+    import cuda.bindings.driver as cuda_driver
+
+    from cudnn.frost.template_loader import load_template
+    from cudnn.frost.tile_dsl.constants import DTYPE_FP16
+    from cudnn.sdpa.fwd import api_dsl
+    from cudnn.sdpa.fwd.config_sm100 import TemplateParams
+    from cudnn.sdpa.fwd.kernels import split_combine_sm100 as comb
+
+    B, H, SQ, SKV, D = 1, 4, 128, 2048, 128
+    dev = "cuda"
+    scale = 1.0 / _math.sqrt(D)
+    torch.manual_seed(0)
+    kdir = _os.path.join(_os.path.dirname(_os.path.abspath(api_dsl.__file__)), "kernels")
+    params = TemplateParams(dtype_qkv=dtype_qkv, dtype_o=DTYPE_FP16, split_kv=splits, cta_mma=cta_mma)
+    mod = load_template(_os.path.join(kdir, kfile), params, tag=f"t_{kfile[:16]}_{splits}_{cta_mma}")
+    fn = mod.compile(b=B, qh=H, kh=H, sq=SQ, skv=SKV, has_lse=True)
+    stream = cuda_driver.CUstream(torch.cuda.current_stream().cuda_stream)
+    o_p = torch.zeros(splits * B, SQ, H, D, device=dev, dtype=torch.float16)
+    lse_p = torch.zeros(splits * B, H, SQ, device=dev, dtype=torch.float32)
+    amax_o = torch.zeros(1, dtype=torch.float32, device=dev)
+    zH = torch.zeros(H, dtype=torch.float32, device=dev)
+    zB = torch.zeros(B, dtype=torch.int32, device=dev)
+    ps = (B, H, H, SQ, SKV, 0)
+    log2e = cutlass.Float32(scale * _math.log2(_math.e))
+
+    if not mx:
+        mk = lambda *sh: (torch.randn(*sh, device=dev) * 0.5).clamp(-448, 448).to(torch.float8_e4m3fn)
+        q, k, v = mk(B, SQ, H, D), mk(B, SKV, H, D), mk(B, SKV, H, D)
+        # The FP8 entry takes four 1-element fp32 DEVICE scale tensors
+        # (descale_q/k/v, scale_o) — the scales fold in-kernel — and no Amax_S.
+        one = lambda: torch.ones(1, dtype=torch.float32, device=dev)
+        fn(q, k, v, o_p, lse_p, zH, zB, ps, log2e, cutlass.Float32(1.0), one(), one(), one(), one(), amax_o, stream=stream)
+        qf, kf, vf = (t.float().permute(0, 2, 1, 3) for t in (q, k, v))
+    else:
+        from sdpa.mxfp8_quant import quantize_to_mxfp8
+
+        qr = torch.randn(B, H, SQ, D, device=dev) * 0.5
+        kr = torch.randn(B, H, SKV, D, device=dev) * 0.5
+        vr = torch.randn(B, H, SKV, D, device=dev) * 0.5
+        a, adq, aswz, b_, bdq, bswz = quantize_to_mxfp8(qr, B, H, SQ, D)
+        q8, sfq, qf = a, aswz, adq.reshape(B, H, SQ, D).float()
+        a, adq, aswz, b_, bdq, bswz = quantize_to_mxfp8(kr, B, H, SKV, D)
+        k8, sfk, kf = a, aswz, adq.reshape(B, H, SKV, D).float()
+        a, adq, aswz, b_, bdq, bswz = quantize_to_mxfp8(vr, B, H, SKV, D)
+        v8, sfv, vf = b_, bswz, bdq.reshape(B, H, SKV, D).float()
+        sfq, sfk, sfv = (t.reshape(B, H, -1, 512).view(torch.int8).contiguous() for t in (sfq, sfk, sfv))
+        q8 = q8.reshape(B, H, SQ, D).permute(0, 2, 1, 3).contiguous()
+        k8 = k8.reshape(B, H, SKV, D).permute(0, 2, 1, 3).contiguous()
+        v8 = v8.reshape(B, H, SKV, D).permute(0, 2, 1, 3).contiguous()
+        fn(q8, k8, v8, o_p, sfq, sfk, sfv, lse_p, amax_o, zH, zB, ps, log2e, stream=stream)
+
+    ref = torch.matmul(torch.softmax(torch.matmul(qf, kf.transpose(-1, -2)) * scale, -1), vf).permute(0, 2, 1, 3)
+    if splits == 1:
+        torch.cuda.synchronize()
+        return o_p.float(), ref, amax_o
+    o_out = torch.zeros(B, SQ, H, D, device=dev, dtype=torch.float16)
+    # has_amax: at splits > 1 the per-split epilogues skip their amax write, so
+    # the combine is what reports it -- over the RECOMBINED O.
+    cfn = comb.compile(b=B, h=H, sq=SQ, d_v=D, splits=splits, dtype_o="f16", has_lse=False, has_amax=True)
+    cfn(o_p, lse_p, o_out, None, amax_o, (B, H, SQ, D), cutlass.Int32(splits), stream=stream)
+    torch.cuda.synchronize()
+    assert not torch.isnan(o_out).any(), "NaN in combined fp8-family O"
+    return o_out.float(), ref, amax_o
+
+
+def _assert_amax_is_of_the_output(amax_o, got):
+    """amax_o must describe the OUTPUT the caller receives, at any split count.
+
+    The per-split epilogues each see only their own partial, and the recombined
+    O is a convex combination of those, so |O| <= max_s |O_s|: a max taken over
+    partials silently over-reports (measured 1.5x at 2 splits, 2.9x at 8) and
+    would hand the caller a far-too-loose quantization scale.  Comparing against
+    the recombined tensor is what catches that; comparing against the partials
+    would pass either way.
+    """
+    reported = amax_o.item()
+    true_amax = got.abs().max().item()
+    # The kernel takes its amax on the fp32 value before the half store, so the
+    # readback can differ by one rounding step in that direction only.
+    assert reported >= true_amax * 0.99, f"amax_o {reported} under-reports |O| {true_amax}"
+    assert reported <= true_amax * 1.01, f"amax_o {reported} over-reports |O| {true_amax} — computed over partials?"
+
+
+@pytest.mark.L0
+@pytest.mark.parametrize("cta_mma", [2, 1], ids=["cga2", "cga1"])
+@pytest.mark.parametrize("splits", [1, 4], ids=lambda s: f"split{s}")
+def test_split_kv_fp8(splits, cta_mma):
+    from cudnn.frost.tile_dsl.constants import DTYPE_E4M3
+
+    got, ref, amax_o = _fp8_family_split("prefill_d128_fp8_sm100.py", DTYPE_E4M3, splits, cta_mma, mx=False)
+    assert (got - ref).abs().max().item() <= 5e-2
+    _assert_amax_is_of_the_output(amax_o, got)
+
+
+@pytest.mark.L0
+@pytest.mark.parametrize("cta_mma", [2, 1], ids=["cga2", "cga1"])
+@pytest.mark.parametrize("splits", [1, 4], ids=lambda s: f"split{s}")
+def test_split_kv_mxfp8(splits, cta_mma):
+    from cudnn.frost.tile_dsl.constants import DTYPE_E4M3
+
+    got, ref, amax_o = _fp8_family_split("prefill_d128_mxfp8_sm100.py", DTYPE_E4M3, splits, cta_mma, mx=True)
+    assert (got - ref).abs().max().item() <= 1.5e-1
+    _assert_amax_is_of_the_output(amax_o, got)
+
+
+# --- gaps found by auditing the suite against what the config permits -----
+
+
+@pytest.mark.L0
+@pytest.mark.parametrize("splits", [1, 4], ids=lambda s: f"split{s}")
+def test_combine_lse_matches_reference(splits):
+    """The RECOMBINED LSE is an output too, and nothing else here checks it.
+
+    split_combine_sm100 computes lse = M + log(sum_s exp(lse_s - M)); every other
+    test only compares O, so a wrong LSE would pass all of them.
+    """
+    import cutlass
+    import cuda.bindings.driver as cuda_driver
+
+    from cudnn.sdpa.fwd.kernels import split_combine_sm100 as comb
+
+    B, H, SQ, SKV = 2, 4, 128, 2048
+    dev = "cuda"
+    scale = 1.0 / math.sqrt(D)
+    torch.manual_seed(0)
+    q = torch.randn(B, SQ, H, D, device=dev, dtype=torch.float16)
+    k = torch.randn(B, SKV, H, D, device=dev, dtype=torch.float16)
+    v = torch.randn(B, SKV, H, D, device=dev, dtype=torch.float16)
+
+    mod = _kernel_module(splits, 3, causal=False)
+    fn = mod.compile(b=B, qh=H, kh=H, sq=SQ, skv=SKV, d_qk=D, d_v=D, has_lse=True)
+    o_p = torch.zeros(splits * B, SQ, H, D, device=dev, dtype=torch.float16)
+    lse_p = torch.zeros(splits * B, H, SQ, device=dev, dtype=torch.float32)
+    stream = cuda_driver.CUstream(torch.cuda.current_stream().cuda_stream)
+    fn(
+        q,
+        k,
+        v,
+        o_p,
+        lse_p,
+        torch.zeros(H, dtype=torch.float32, device=dev),
+        torch.zeros(B, dtype=torch.int32, device=dev),
+        torch.zeros(1, dtype=torch.int64, device=dev),
+        (B, H, H, SQ, SKV, 0),
+        cutlass.Float32(scale * math.log2(math.e)),
+        cutlass.Int32(0),
+        None,
+        stream=stream,
+    )
+
+    # Reference LSE = logsumexp of the scaled scores, in natural log.
+    qb, kb = (t.float().permute(0, 2, 1, 3) for t in (q, k))
+    ref_lse = torch.logsumexp(torch.matmul(qb, kb.transpose(-1, -2)) * scale, dim=-1)  # [B,H,SQ]
+
+    if splits == 1:
+        torch.cuda.synchronize()
+        got_lse = lse_p.view(B, H, SQ)
+    else:
+        o_out = torch.zeros(B, SQ, H, D, device=dev, dtype=torch.float16)
+        lse_out = torch.zeros(B, H, SQ, device=dev, dtype=torch.float32)
+        cfn = comb.compile(b=B, h=H, sq=SQ, d_v=D, splits=splits, dtype_o="f16", has_lse=True)
+        cfn(o_p, lse_p, o_out, lse_out, None, (B, H, SQ, D), cutlass.Int32(splits), stream=stream)
+        torch.cuda.synchronize()
+        got_lse = lse_out
+    assert not torch.isnan(got_lse).any(), "NaN in recombined LSE"
+    assert (got_lse - ref_lse).abs().max().item() <= 5e-3
+
+
+@pytest.mark.L0
+@pytest.mark.parametrize("flavor", sorted(_F16_FLAVORS))
+@pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16], ids=["fp16", "bf16"])
+def test_even_splits_every_flavor_batched(flavor, dtype):
+    """Even splits, B > 1, both half dtypes, on every flavor.
+
+    B > 1 matters specifically: the split-major slot is batch + split*B and the
+    grid.z decode is b % n_batch / b // n_batch, all of which are trivially
+    correct at B == 1 and so were barely exercised.
+    """
+    import os as _os
+
+    import cutlass
+    import cuda.bindings.driver as cuda_driver
+
+    from cudnn.frost.template_loader import load_template
+    from cudnn.sdpa.fwd import api_dsl
+    from cudnn.sdpa.fwd.config_sm100 import TemplateParams
+    from cudnn.sdpa.fwd.kernels import split_combine_sm100 as comb
+
+    kmod, d_qk, d_v = _F16_FLAVORS[flavor]
+    B, H, SQ, SKV, S = 3, 4, 128, 2048, 4
+    dev = "cuda"
+    scale = 1.0 / math.sqrt(d_qk)
+    torch.manual_seed(0)
+    q = torch.randn(B, SQ, H, d_qk, device=dev, dtype=dtype)
+    k = torch.randn(B, SKV, H, d_qk, device=dev, dtype=dtype)
+    v = torch.randn(B, SKV, H, d_v, device=dev, dtype=dtype)
+
+    path = _os.path.join(_os.path.dirname(_os.path.abspath(api_dsl.__file__)), "kernels", kmod)
+    params = TemplateParams(dtype_qkv=3 if dtype == torch.float16 else 2, split_kv=S)
+    mod = load_template(path, params, tag=f"even_{flavor}_{dtype}_{S}")
+    fn = mod.compile(b=B, qh=H, kh=H, sq=SQ, skv=SKV, d_qk=d_qk, d_v=d_v, has_lse=True)
+    o_p = torch.zeros(S * B, SQ, H, d_v, device=dev, dtype=dtype)
+    lse_p = torch.zeros(S * B, H, SQ, device=dev, dtype=torch.float32)
+    stream = cuda_driver.CUstream(torch.cuda.current_stream().cuda_stream)
+    fn(
+        q,
+        k,
+        v,
+        o_p,
+        lse_p,
+        torch.zeros(H, dtype=torch.float32, device=dev),
+        torch.zeros(B, dtype=torch.int32, device=dev),
+        torch.zeros(1, dtype=torch.int64, device=dev),
+        (B, H, H, SQ, SKV, 0),
+        cutlass.Float32(scale * math.log2(math.e)),
+        cutlass.Int32(0),
+        None,
+        stream=stream,
+    )
+    o_out = torch.zeros(B, SQ, H, d_v, device=dev, dtype=dtype)
+    cfn = comb.compile(b=B, h=H, sq=SQ, d_v=d_v, splits=S, dtype_o="f16" if dtype == torch.float16 else "bf16", has_lse=False)
+    cfn(o_p, lse_p, o_out, None, None, (B, H, SQ, d_v), cutlass.Int32(S), stream=stream)
+    torch.cuda.synchronize()
+    assert not torch.isnan(o_out).any(), f"{flavor}: NaN"
+    ref = _ref_sdpa(q, k, v, scale, is_causal=False, kh=H)
+    assert (o_out.float() - ref).abs().max().item() <= (2e-2 if dtype == torch.float16 else 1e-1)
+
+
+# --- masks the config permits with split, previously untested -------------
+#
+# split_kv only rejects sinks and THD, so SWA / bottom-right / padded are all
+# reachable. They are also where _bounds_for_tile_split does its real work:
+# it slices the ALREADY-masked [left, right) and clamps unmasked_lo/hi into the
+# slice, and none of that was exercised. The reference is the one the sibling
+# suite uses to model this kernel's exact mask semantics.
+
+
+def _run_masked(kfile, d_qk, d_v, splits, *, B, H, KH, SQ, SKV, tp_kwargs, seq_kv_lens=None, seq_q_lens=None, dtype=torch.float16, cta_mma=2):
+    """Launch split kernel + combine under an arbitrary mask; returns (got, q, k, v, scale)."""
+    import os as _os
+
+    import cutlass
+    import cuda.bindings.driver as cuda_driver
+
+    from cudnn.frost.template_loader import load_template
+    from cudnn.sdpa.fwd import api_dsl
+    from cudnn.sdpa.fwd.config_sm100 import TemplateParams
+    from cudnn.sdpa.fwd.kernels import split_combine_sm100 as comb
+
+    dev = "cuda"
+    scale = 1.0 / math.sqrt(d_qk)
+    torch.manual_seed(0)
+    q = torch.randn(B, SQ, H, d_qk, device=dev, dtype=dtype)
+    k = torch.randn(B, SKV, KH, d_qk, device=dev, dtype=dtype)
+    v = torch.randn(B, SKV, KH, d_v, device=dev, dtype=dtype)
+
+    path = _os.path.join(_os.path.dirname(_os.path.abspath(api_dsl.__file__)), "kernels", kfile)
+    params = TemplateParams(dtype_qkv=3 if dtype == torch.float16 else 2, split_kv=splits, cta_mma=cta_mma, **tp_kwargs)
+    mod = load_template(path, params, tag=f"mask_{kfile[:16]}_{splits}_{cta_mma}_{sorted(tp_kwargs.items())}")
+    fn = mod.compile(b=B, qh=H, kh=KH, sq=SQ, skv=SKV, d_qk=d_qk, d_v=d_v, has_lse=True)
+
+    o_p = torch.zeros(splits * B, SQ, H, d_v, device=dev, dtype=dtype)
+    lse_p = torch.zeros(splits * B, H, SQ, device=dev, dtype=torch.float32)
+    skv_t = seq_kv_lens if seq_kv_lens is not None else torch.zeros(B, dtype=torch.int32, device=dev)
+    stream = cuda_driver.CUstream(torch.cuda.current_stream().cuda_stream)
+    fn(
+        q,
+        k,
+        v,
+        o_p,
+        lse_p,
+        torch.zeros(H, dtype=torch.float32, device=dev),
+        skv_t,
+        torch.zeros(1, dtype=torch.int64, device=dev),
+        (B, H, KH, SQ, SKV, 0),
+        cutlass.Float32(scale * math.log2(math.e)),
+        cutlass.Int32(0),
+        seq_q_lens,
+        stream=stream,
+    )
+    if splits == 1:
+        torch.cuda.synchronize()
+        return o_p.float(), q, k, v, scale
+    o_out = torch.zeros(B, SQ, H, d_v, device=dev, dtype=dtype)
+    cfn = comb.compile(b=B, h=H, sq=SQ, d_v=d_v, splits=splits, dtype_o="f16" if dtype == torch.float16 else "bf16", has_lse=False)
+    cfn(o_p, lse_p, o_out, None, None, (B, H, SQ, d_v), cutlass.Int32(splits), stream=stream)
+    torch.cuda.synchronize()
+    assert not torch.isnan(o_out).any(), "NaN under mask+split"
+    return o_out.float(), q, k, v, scale
+
+
+def _bhsd(t):
+    return t.permute(0, 2, 1, 3)
+
+
+@pytest.mark.L0
+@pytest.mark.parametrize("splits", [1, 4], ids=lambda s: f"split{s}")
+def test_split_kv_swa_causal(splits):
+    """Causal + sliding window: both mask bits, so unmasked_lo AND hi are live."""
+    from test_sdpa_fwd_dsl_sm100 import _ref_sdpa_full
+
+    W = 256
+    B, H, SQ, SKV = 1, 4, 1024, 1024
+    got, q, k, v, scale = _run_masked(
+        "prefill_d128_f16_sm100.py", 128, 128, splits, B=B, H=H, KH=H, SQ=SQ, SKV=SKV, tp_kwargs=dict(window_right=0, window_left=W)
+    )
+    ref = _ref_sdpa_full(_bhsd(q), _bhsd(k), _bhsd(v), scale=scale, is_causal=True, swa_window=W)
+    assert (got - ref.float().permute(0, 2, 1, 3)).abs().max().item() <= 2e-2
+
+
+@pytest.mark.L0
+@pytest.mark.parametrize("splits", [1, 4], ids=lambda s: f"split{s}")
+def test_split_kv_bottom_right_causal(splits):
+    """Bottom-right anchored causal — the diagonal sits at (S_q, S_kv)."""
+    from test_sdpa_fwd_dsl_sm100 import _ref_sdpa_full
+
+    B, H, SQ, SKV = 1, 4, 128, 2048
+    got, q, k, v, scale = _run_masked(
+        "prefill_d128_f16_sm100.py", 128, 128, splits, B=B, H=H, KH=H, SQ=SQ, SKV=SKV, tp_kwargs=dict(window_right=0, bottom_right=True)
+    )
+    ref = _ref_sdpa_full(_bhsd(q), _bhsd(k), _bhsd(v), scale=scale, is_causal=True, bottom_right=True)
+    assert (got - ref.float().permute(0, 2, 1, 3)).abs().max().item() <= 2e-2
+
+
+@pytest.mark.L0
+@pytest.mark.parametrize("splits", [1, 4], ids=lambda s: f"split{s}")
+def test_split_kv_padded_kv(splits):
+    """Per-batch KV padding: the split must slice each batch's OWN live range."""
+    from test_sdpa_fwd_dsl_sm100 import _ref_sdpa_full
+
+    B, H, SQ, SKV = 2, 4, 128, 2048
+    lens = torch.tensor([2048, 1531], dtype=torch.int32, device="cuda")  # 2nd ends mid-tile
+    got, q, k, v, scale = _run_masked(
+        "prefill_d128_f16_sm100.py", 128, 128, splits, B=B, H=H, KH=H, SQ=SQ, SKV=SKV, tp_kwargs=dict(seq_kv_lens_present=True), seq_kv_lens=lens
+    )
+    ref = _ref_sdpa_full(_bhsd(q), _bhsd(k), _bhsd(v), scale=scale, seq_kv_lens=lens)
+    assert (got - ref.float().permute(0, 2, 1, 3)).abs().max().item() <= 2e-2
+
+
+@pytest.mark.L0
+@pytest.mark.parametrize("flavor", ["d192", "d256", "d512"])
+def test_split_kv_causal_other_flavors(flavor):
+    """Causal + split beyond d128 — _split_chunk slices the masked range here too."""
+    from test_sdpa_fwd_dsl_sm100 import _ref_sdpa_full
+
+    kmod, d_qk, d_v = _F16_FLAVORS[flavor]
+    B, H, SQ, SKV = 1, 4, 1024, 1024
+    got, q, k, v, scale = _run_masked(kmod, d_qk, d_v, 4, B=B, H=H, KH=H, SQ=SQ, SKV=SKV, tp_kwargs=dict(window_right=0))
+    ref = _ref_sdpa_full(_bhsd(q), _bhsd(k), _bhsd(v), scale=scale, is_causal=True)
+    assert (got - ref.float().permute(0, 2, 1, 3)).abs().max().item() <= 2e-2
+
+
+@pytest.mark.L0
+@pytest.mark.parametrize("splits", [1, 4], ids=lambda s: f"split{s}")
+def test_split_kv_padded_q_trim(splits):
+    """Dense padded-Q trim (seq_q_lens) + split.
+
+    Q rows at or past the batch's actual length must come back O := 0 (cuDNN
+    >= 9.14 convention). Under split that has to hold for EVERY split's partial,
+    or the combine mixes live and dead rows -- the trim is applied per split in
+    the epilogue, after which the dead row's lse = -inf makes it drop out.
+    """
+    from test_sdpa_fwd_dsl_sm100 import _ref_sdpa_full
+
+    B, H, SQ, SKV = 2, 4, 256, 2048
+    kv_lens = torch.tensor([2048, 2048], dtype=torch.int32, device="cuda")
+    q_lens = torch.tensor([256, 137], dtype=torch.int32, device="cuda")  # 2nd trims mid-tile
+    got, q, k, v, scale = _run_masked(
+        "prefill_d128_f16_sm100.py",
+        128,
+        128,
+        splits,
+        B=B,
+        H=H,
+        KH=H,
+        SQ=SQ,
+        SKV=SKV,
+        tp_kwargs=dict(seq_kv_lens_present=True, seq_q_lens_present=True),
+        seq_kv_lens=kv_lens,
+        seq_q_lens=q_lens,
+    )
+    ref = _ref_sdpa_full(_bhsd(q), _bhsd(k), _bhsd(v), scale=scale, seq_kv_lens=kv_lens)
+    ref = ref.float().permute(0, 2, 1, 3)  # -> BSHD
+    # Rows past the per-batch Q length must be exactly zero.
+    rows = torch.arange(SQ, device=got.device).view(1, SQ, 1, 1)
+    dead = rows >= q_lens.view(B, 1, 1, 1)
+    assert got[dead.expand_as(got)].abs().max().item() == 0.0, "trimmed Q rows are not zero"
+    live = ~dead
+    assert (got - ref)[live.expand_as(got)].abs().max().item() <= 2e-2