Add SM100 FP8 SDPA support for d192/d128

python/cudnn/engines/manifest.py

@@ -184,6 +184,7 @@ def offered_ids(self) -> Dict[str, int]:
             "sdpa_fwd_prefill_sm100_d192_d128": EngineSlot(6, opt_in=True),
             "sdpa_fwd_prefill_sm120_fp8": EngineSlot(7, opt_in=True),
             "sdpa_fwd_prefill_sm80": EngineSlot(8, opt_in=True),
+            "sdpa_fwd_prefill_sm100_d192_d128_fp8": EngineSlot(9, opt_in=True),
         },
         analyzer=("cudnn.sdpa.graph_analyzer", "analyze"),
         heuristics=("cudnn.sdpa.fwd.heuristics", "recommend"),

python/cudnn/frost/tile_dsl/tma.py

@@ -41,7 +41,16 @@ def cp_async_bulk_shared_cluster_shared_cta(dst_mem, src_mem, mbar, size, *, pre
 
 
 @cute.jit
-def tma_load_tile(smem_tile, gmem_slice, mbar, *, cta_group: int = 1, mcast_mask=None, acquire: cutlass.Constexpr[bool] = True):
+def tma_load_tile(
+    smem_tile,
+    gmem_slice,
+    mbar,
+    *,
+    cta_group: int = 1,
+    mcast_mask=None,
+    acquire: cutlass.Constexpr[bool] = True,
+    l2_cache_hint=None,
+):
     num_iters = smem_tile.tma_loads_per_tile
     granu_elems = smem_tile.tma_granu_elems
     sub_stride = smem_tile.tma_subtile_stride_elems
@@ -70,6 +79,7 @@ def tma_load_tile(smem_tile, gmem_slice, mbar, *, cta_group: int = 1, mcast_mask
                     tma_desc_ptr,
                     coords,
                     mbar,
+                    l2_cache_hint=l2_cache_hint,
                 )
             else:
                 nvvm.cp_async_bulk_tensor_shared_cluster_global(
@@ -80,6 +90,7 @@ def tma_load_tile(smem_tile, gmem_slice, mbar, *, cta_group: int = 1, mcast_mask
                     [],
                     multicast_mask=mcast_mask,
                     group=nvvm.CTAGroup.CTA_2,
+                    l2_cache_hint=l2_cache_hint,
                 )
 
 

python/cudnn/sdpa/fwd/api_dsl.py

@@ -74,12 +74,23 @@ def dtype_name(buffer) -> str:
     torch.float16: DTYPE_FP16,
 }
 _SM100_FP8_DTYPES = (torch.float8_e4m3fn, torch.float8_e5m2)
-# d128 FP8 kernels (E4M3/E5M2 in, BF16/FP16/FP8 out). Block-scale MXFP8 (per-32-block
-# E8M0 SF) vs per-tensor FP8 (scalar descales). Selected by the graph op (sdpa_mxfp8 vs
-# sdpa_fp8); the f16/bf16 flavors use _SM100_KERNEL_FILES.
+# FP8 kernels use E4M3/E5M2 inputs and BF16/FP16/FP8 outputs. Block-scale
+# MXFP8 remains d128-only; per-tensor FP8 has exact d128/d128 and d192/d128
+# kernels.
 _SM100_MXFP8_KERNEL_FILE = "prefill_d128_mxfp8_sm100.py"
-_SM100_FP8_KERNEL_FILE = "prefill_d128_fp8_sm100.py"
 _SM107_FP8_KERNEL_FILE = "prefill_d128_fp8_sm107.py"
+_SM100_FP8_KERNEL_FILES = {
+    (128, 128): "prefill_d128_fp8_sm100.py",
+    (192, 128): "prefill_d192_d128_fp8_sm100.py",
+}
+
+
+def _sm100_fp8_shapes(pertensor: bool, device_cc: tuple[int, int]) -> frozenset[tuple[int, int]]:
+    if not pertensor or device_cc == (10, 7):
+        return frozenset({(128, 128)})
+    return frozenset({(128, 128), (192, 128)})
+
+
 # Both flavors tile KV in TILE_N=128 columns; the KV tail is only masked when
 # the padded/causal mask paths are active (see check_support).
 _SM100_TILE_N = 128
@@ -204,7 +215,7 @@ def _pick_flavor(d_qk: int, d_v: int) -> tuple[int, int]:
     the tile box stays the compile-time D, so loads past d_qk / d_v hardware
     zero-fill (adding exact zero terms to every QK^T dot product — S, softmax
     and P·V are bit-identical to the unpadded problem) and O stores past d_v
-    are OOB-clipped. FP8/MXFP8 stays exact-match d128 (gated in
+    are OOB-clipped. FP8/MXFP8 uses exact native shapes (gated in
     check_support); alignment (d % 8, the TMA 16-byte global-stride rule at
     2 bytes/elem) is also gated in check_support / engines.mismatch.
     """
@@ -231,11 +242,11 @@ def _load_sm100_kernel_module(flavor: tuple[int, int], params: Sm100TemplatePara
     dense K=64 FP8 path baked in — see prefill_d128_fp8_sm107.py)."""
 
     tag = _flavor_tag(flavor)
-    if fp8 and pertensor and rubin:
+    if fp8 and pertensor and rubin and flavor == (128, 128):
         filename = _SM107_FP8_KERNEL_FILE
         tag = f"sdpa_fwd_sm107_fp8_{tag}"
     elif fp8:
-        filename = _SM100_FP8_KERNEL_FILE if pertensor else _SM100_MXFP8_KERNEL_FILE
+        filename = _SM100_FP8_KERNEL_FILES[flavor] if pertensor else _SM100_MXFP8_KERNEL_FILE
         tag = f"sdpa_fwd_sm100_{'fp8' if pertensor else 'mxfp8'}_{tag}"
     else:
         filename = _SM100_KERNEL_FILES[flavor]
@@ -782,11 +793,14 @@ def check_support(self) -> bool:
             f"SdpaFwdDslSm100 requires {_allowed_msg}; found SM{major}{minor} on {device}",
         )
 
-        # FP8/MXFP8: exact-match d128 only — the FP8 kernels' SF plumbing and
-        # QMMA geometry are not audited for envelope zero-padding.
+        # FP8 paths use exact native shapes. Per-tensor FP8 also has the
+        # d192/d128 flavor; MXFP8 remains d128-only until its scale-factor
+        # descriptors and pipeline are extended.
+        fp8_shapes = _sm100_fp8_shapes(self._pertensor, self._device_cc)
         self._value_error_if(
-            self._fp8 and (int(d_qk), int(d_v)) != (128, 128),
-            f"FP8/MXFP8 (E4M3/E5M2 inputs) requires exact D_QK=D_V=128 (no envelope padding); got (D_QK={d_qk}, D_V={d_v})",
+            self._fp8 and (int(d_qk), int(d_v)) not in fp8_shapes,
+            f"{'FP8' if self._pertensor else 'MXFP8'} (E4M3/E5M2 inputs) requires an exact native shape in {sorted(fp8_shapes)} "
+            f"(no envelope padding); got (D_QK={d_qk}, D_V={d_v})",
         )
         # Envelope alignment gate: the TMA descriptors are built from the
         # actual tensor extents, and cuTensorMapEncodeTiled requires every
@@ -936,16 +950,37 @@ def compile(self) -> None:
                 d_v=d_v_sched,
                 elem_bytes=1 if self._fp8 else 2,
             )
+        lpt_head_group = 1
+        if self._fp8 and self._pertensor and self.flavor == (192, 128) and not self.thd and (self.batch_size * self.h_q) % 16 == 0:
+            lpt_head_group = 16
+        lpt_q_tiles = 0
+        if self._fp8 and self._pertensor and self.flavor == (192, 128) and not self.thd:
+            lpt_q_tiles = (self.s_q_max + 511) // 512
+        template_window_right = self.window_right
+        if (
+            self._fp8
+            and self._pertensor
+            and self.flavor == (192, 128)
+            and self.window_left is None
+            and self.window_right is None
+            and not self.seq_kv_lens_present
+        ):
+            # CUTLASS DSL 4.7 does not finish lowering the large-shape FP8
+            # MASK_NONE x32 path. A right bound of S_kv removes no valid K but
+            # selects the equivalent masked-interior lowering.
+            template_window_right = self.s_k_max
         params = Sm100TemplateParams(
             dtype_qkv=_SM100_DTYPE_QKV_CODE[self.dtype],
             dtype_o=_SM100_DTYPE_QKV_CODE[self.dtype_o],
             window_left=self.window_left,
-            window_right=self.window_right,
+            window_right=template_window_right,
             bottom_right=self.causal_bottom_right,
             has_sink=self.has_sink,
             seq_kv_lens_present=self.seq_kv_lens_present,
             seq_q_lens_present=self.seq_q_lens_present,
             sched_policy=sched_policy,
+            lpt_head_group=lpt_head_group,
+            lpt_q_tiles=lpt_q_tiles,
             thd_varlen=self.thd,
             fused_ldtm_stat=fused_ldtm_stat,
         )
@@ -960,7 +995,7 @@ def compile(self) -> None:
             # f16-only.
             self._compiled_kernel = self._k_mod.compile(**self._thd_compile_kwargs())
         elif self._fp8:
-            # FP8/MXFP8 kernels are exact-match d128 (gated in check_support);
+            # FP8/MXFP8 kernels use exact native shapes (gated in check_support);
             # their compile() has no envelope head-dim parameters. has_lse=False
             # (no Stats output) compiles the LSE store out — no dummy buffer at
             # any level (the amax_o atomicMax write is independent).

python/cudnn/sdpa/fwd/config_sm100.py

@@ -87,6 +87,12 @@ class TemplateParams:
     # cu_seqlens instead.
     seq_q_lens_present: bool = False
     sched_policy: int = SCHED_NATURAL
+    # Compile-time LPT head/batch grouping. Keep 1 unless the selected kernel
+    # and concrete graph shape opt into a divisor of B*Hq.
+    lpt_head_group: int = 1
+    # Dense D192 FP8 may specialize the reverse-row LPT decoder to its exact
+    # number of query tiles. Zero keeps the existing runtime derivation.
+    lpt_q_tiles: int = 0
     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)
@@ -125,8 +131,8 @@ 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}")
     fp8 = k.dtype_qkv in (DTYPE_E4M3, DTYPE_E5M2)
-    if fp8 and flavor != "d128":
-        raise ValueError(f"{flavor}: FP8/MXFP8 inputs (DTYPE_QKV 0/1) are only supported on d128")
+    if fp8 and flavor not in ("d128", "d192"):
+        raise ValueError(f"{flavor}: FP8/MXFP8 inputs (DTYPE_QKV 0/1) are only supported on d128 and d192")
     dtype_o = k.dtype_qkv if k.dtype_o < 0 else k.dtype_o
     if dtype_o not in (DTYPE_E4M3, DTYPE_E5M2, DTYPE_BF16, DTYPE_FP16):
         raise ValueError(f"{flavor}: DTYPE_O must be 0..3; got {dtype_o}")
@@ -173,6 +179,10 @@ def _validate_params(flavor: str, k: TemplateParams) -> None:
             # 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)")
+    if k.lpt_head_group not in (1, 16):
+        raise ValueError(f"{flavor}: LPT_HEAD_GROUP must be 1 or 16; got {k.lpt_head_group}")
+    if fp8 and flavor == "d192" and k.split_kv != 1:
+        raise ValueError("d192: split_kv is not implemented by the per-tensor FP8 kernel")
 
 
 def _mask_flags_from(params: TemplateParams) -> int:
@@ -773,7 +783,7 @@ class CfgD192(CfgD128):
 
 
 def _d192_smem_bytes(cfg) -> int:
-    """Data-buffer SMEM for the d192 pipeline (Q/O always aliased).
+    """Data-buffer SMEM for the d192 pipeline.
 
     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:
@@ -792,32 +802,43 @@ def _d192_smem_bytes(cfg) -> int:
 
 def _validate_cfg_d192(cfg: CfgD192) -> None:
     """Consistency checks on the native DSv3 d192/d128 geometry."""
+    fp8 = cfg.DTYPE_QKV in (DTYPE_E4M3, DTYPE_E5M2)
     checks = (
-        (cfg.DTYPE_QKV in (DTYPE_BF16, DTYPE_FP16), "d192: only BF16/FP16 inputs are supported"),
-        (cfg.DTYPE_O == cfg.DTYPE_QKV, "d192: DTYPE_O must equal DTYPE_QKV"),
+        (
+            cfg.DTYPE_O in (DTYPE_E4M3, DTYPE_E5M2, DTYPE_BF16, DTYPE_FP16) if fp8 else cfg.DTYPE_O == cfg.DTYPE_QKV,
+            "d192: DTYPE_O must equal DTYPE_QKV for half input; FP8 allows an independent output dtype",
+        ),
         (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 == 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",
+            cfg.STAGES_KV == (2 if fp8 else 1) * cfg.CTA_MMA,
+            "d192: STAGES_KV must scale with input dtype and cluster width " "(FP8: 2/4 at cga1/cga2; half: 1/2)",
         ),
         (
             _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.QO_ALIAS == (0 if fp8 else 1), "d192: Q/O SMEM alias must be disabled for FP8 and enabled for half input"),
         (cfg.TILES_Q == 2, "d192: TILES_Q must be 2"),
         (cfg.SOFTMAX_WARPGROUPS == 2, "d192: SOFTMAX_WARPGROUPS must be 2"),
         (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.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"),
+        (
+            cfg.TILE_K_HW_BMM1 == (32 if fp8 else 16) and cfg.TILE_K_HW_BMM2 == (32 if fp8 else 16),
+            "d192: TILE_K_HW must be 32 for FP8 and 16 for BF16/FP16",
+        ),
+        (
+            cfg.Q_SWZ_BYTES == (64 if fp8 else 128) and cfg.K_SWZ_BYTES == (64 if fp8 else 128),
+            "d192: Q/K swizzle must be 64B for FP8 and 128B for BF16/FP16",
+        ),
+        (
+            cfg.V_SWZ_BYTES == v_swz_bytes(128, cfg.CTA_MMA, cfg.BPE) and cfg.O_SWZ_BYTES in ((64, 128) if fp8 else (128,)),
+            "d192: V/O swizzle is inconsistent with the input/output dtype",
+        ),
     )
     for ok, msg in checks:
         if not ok:
@@ -827,32 +848,34 @@ def _validate_cfg_d192(cfg: CfgD192) -> None:
 def make_cfg_d192(params: TemplateParams) -> Tuple[CfgD192, TmaIters]:
     _validate_params("d192", params)
     b = bpe(params.dtype_qkv)
+    fp8 = params.dtype_qkv in (DTYPE_E4M3, DTYPE_E5M2)
+    dtype_o = params.dtype_qkv if params.dtype_o < 0 else params.dtype_o
+    b_o = bpe(dtype_o)
     cfg = CfgD192(
         DTYPE_QKV=params.dtype_qkv,
-        DTYPE_O=params.dtype_qkv,
+        DTYPE_O=dtype_o,
         BPE=b,
-        BPE_O=b,
+        BPE_O=b_o,
         SPLIT_KV=int(params.split_kv),
+        QO_ALIAS=0 if fp8 else 1,
         Q_SWZ_BYTES=q_swz_bytes(192, b),
         K_SWZ_BYTES=q_swz_bytes(192, 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),
+        O_SWZ_BYTES=o_swz_bytes(128, b_o),
         RESCALE_THRESHOLD=rescale_threshold(params.dtype_qkv),
-        TILE_K_HW_BMM1=tile_k_hw(params.dtype_qkv),
-        TILE_K_HW_BMM2=tile_k_hw(params.dtype_qkv),
+        TILE_K_HW_BMM1=32 if fp8 else tile_k_hw(params.dtype_qkv),
+        TILE_K_HW_BMM2=32 if fp8 else tile_k_hw(params.dtype_qkv),
+        STAGES_KV=(2 if fp8 else 1) * params.cta_mma,
         MASK_FLAGS=_mask_flags_from(params),
         WINDOW_LEFT=params.window_left or 0,
         WINDOW_RIGHT=params.window_right or 0,
         HAS_SINK=int(params.has_sink),
         BOTTOM_RIGHT=int(params.bottom_right),
         SCHEDULER_POLICY=1,
-        SOFTMAX_REGS=216 if _mask_flags_from(params) == MASK_NONE else 192,
-        CORRECTION_REGS=40 if _mask_flags_from(params) == MASK_NONE else 88,
+        SOFTMAX_REGS=184 if fp8 else 216 if _mask_flags_from(params) == MASK_NONE else 192,
+        CORRECTION_REGS=104 if fp8 else 40 if _mask_flags_from(params) == MASK_NONE else 88,
         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),

python/cudnn/sdpa/fwd/engines.py

@@ -150,6 +150,10 @@ class Capabilities:
     swa: bool = False
     padded: bool = False
     sink: bool = False
+    # Optional dtype subset for sink support. None means every dtype served by
+    # the engine; a subset lets one exact-shape flavor decline an unsupported
+    # low-precision sink path without affecting its non-sink coverage.
+    sink_dtypes: Optional[frozenset] = None
     stats: bool = False
     # The adapter accepts lse_tensor=None (its kernel None-specializes the LSE
     # store), so a stats-less graph needs no dummy-LSE workspace chunk. Every
@@ -339,6 +343,9 @@ def mismatch(capabilities: Capabilities, facts: "ga.SdpaGraphFacts", knobs: Opti
         if fact and not cap:
             return f"graph uses {label}, which this engine does not support"
 
+    if facts.has_sink and capabilities.sink_dtypes is not None and facts.dtype not in capabilities.sink_dtypes:
+        return f"sink token with dtype {facts.dtype} not in {sorted(str(d) for d in capabilities.sink_dtypes)}"
+
     if facts.right_band_widening and facts.right_bound is not None and facts.right_bound < 0:
         return f"negative diagonal_band_right_bound ({facts.right_bound}) is not supported"
 
@@ -481,8 +488,14 @@ def _sm100_mxfp8_spec(d: int) -> EngineSpec:
     )
 
 
-def _sm100_fp8_spec(d: int) -> EngineSpec:
-    """d128 per-tensor FP8 engine (E4M3/E5M2 in + scalar descales, half/FP8 out).
+def _sm100_fp8_spec(
+    d: int,
+    d_v: Optional[int] = None,
+    *,
+    dtypes: Optional[frozenset] = None,
+    sink_dtypes: Optional[frozenset] = None,
+) -> EngineSpec:
+    """Exact-shape per-tensor FP8 engine with scalar descales.
 
     Padding mask (per-batch ``seq_len_kv`` → KV-side masking) is supported: KV-only
     padding leaves every query row real, so each row's total_sum > 0 and the
@@ -491,15 +504,19 @@ def _sm100_fp8_spec(d: int) -> EngineSpec:
     (dense execute only for v1), so thd=False.
     """
 
+    d_v = d if d_v is None else d_v
+    suffix = f"d{d}" if d_v == d else f"d{d}_d{d_v}"
+    if dtypes is None:
+        dtypes = frozenset({cudnn.data_type.FP8_E4M3, cudnn.data_type.FP8_E5M2})
     return EngineSpec(
-        name=f"sdpa_fwd_prefill_sm100_d{d}_fp8",
+        name=f"sdpa_fwd_prefill_sm100_{suffix}_fp8",
         capabilities=Capabilities(
             sm_lo=_BLACKWELL[0],
             sm_hi=_BLACKWELL[1],
             phase="prefill",
             d_qk=frozenset({d}),
-            d_v=frozenset({d}),
-            dtypes=frozenset({cudnn.data_type.FP8_E4M3, cudnn.data_type.FP8_E5M2}),
+            d_v=frozenset({d_v}),
+            dtypes=dtypes,
             out_dtypes=frozenset({cudnn.data_type.HALF, cudnn.data_type.BFLOAT16, cudnn.data_type.FP8_E4M3, cudnn.data_type.FP8_E5M2}),
             is_fp8=True,
             causal=True,
@@ -508,6 +525,7 @@ def _sm100_fp8_spec(d: int) -> EngineSpec:
             swa=True,
             padded=True,
             sink=True,
+            sink_dtypes=sink_dtypes,
             stats=True,
             lse_optional=True,
             # The fp8 kernel lacks the SEQ_Q_LENS_PRESENT epilogue trim, but its
@@ -1026,6 +1044,14 @@ def _sm120_fp8_spec() -> EngineSpec:
     _sm100_spec(512),
     _sm100_mxfp8_spec(128),
     _sm100_fp8_spec(128),
+    _sm100_fp8_spec(
+        192,
+        d_v=128,
+        dtypes=frozenset({cudnn.data_type.FP8_E4M3, cudnn.data_type.FP8_E5M2}),
+        # The D192 E5M2 sink path has a distinct FP8 online-softmax rounding
+        # trajectory that exceeds the frontend tolerance on sparse CI seeds.
+        sink_dtypes=frozenset({cudnn.data_type.FP8_E4M3}),
+    ),
     _sm120_spec(),
     _sm120_fp8_spec(),
     _sm80_spec(),