diff --git a/python/cudnn/deepseek_sparse_attention/sparse_attention_backward/_interface_sm100.py b/python/cudnn/deepseek_sparse_attention/sparse_attention_backward/_interface_sm100.py
index 3030f8198..ed7a3274c 100644
--- a/python/cudnn/deepseek_sparse_attention/sparse_attention_backward/_interface_sm100.py
+++ b/python/cudnn/deepseek_sparse_attention/sparse_attention_backward/_interface_sm100.py
@@ -8,7 +8,7 @@
 import cutlass.cute as cute
 
 from cudnn.deepseek_sparse_attention.utils.compiler import compile_options
-from cudnn.deepseek_sparse_attention.utils.runtime import resolve_stream
+from cudnn.deepseek_sparse_attention.utils.runtime import resolve_stream, torch_stream_context
 from cudnn.deepseek_sparse_attention.utils.tensor_conversion import to_cute_tensor
 from .dsa_bwd_sm100 import FlashAttentionDSABackwardSm100
 
@@ -39,10 +39,10 @@ def flash_attn_bwd_sm100(
     Internally wraps as batch=1 for the CuTe DSL kernel.
 
     Args:
-        q: (total_S_q, nheads, headdim) bfloat16
-        kv: (total_S_kv, headdim) bfloat16  (K=V, MQA h_kv=1)
-        out: (total_S_q, nheads, headdim_v) bfloat16
-        dout: (total_S_q, nheads, headdim_v) bfloat16
+        q: (total_S_q, nheads, headdim) float16 or bfloat16
+        kv: (total_S_kv, headdim) float16 or bfloat16  (K=V, MQA h_kv=1)
+        out: (total_S_q, nheads, headdim_v) float16 or bfloat16
+        dout: (total_S_q, nheads, headdim_v) float16 or bfloat16
         lse: (total_S_q, nheads) float32, FlashMLA KV-only LSE excluding sink
         attn_sink: (nheads,) float32
         topk_idxs: (total_S_q, topk_max) int32, global indices
@@ -56,6 +56,10 @@ def flash_attn_bwd_sm100(
     """
     total_S_q, num_head, head_dim = q.shape
     total_S_kv = kv.shape[0]
+    # Mirror the check_support gate: the SM100 kernel is tiled only for
+    # head_dim in {512, 576}; any other value indexes shared memory out of
+    # bounds and crashes inside the kernel.
+    assert head_dim in (512, 576), f"head_dim must be 512 or 576, got {head_dim}"
     head_dim_v = 512 if head_dim == 576 else head_dim
     device = q.device
 
@@ -67,7 +71,21 @@ def flash_attn_bwd_sm100(
     tensors_to_check = [q, kv, out, dout, lse, attn_sink, topk_idxs]
     if topk_length is not None:
         tensors_to_check.append(topk_length)
-    assert all(t.is_cuda for t in tensors_to_check)
+    assert all(t.is_cuda and t.device == device for t in tensors_to_check), f"all inputs must be CUDA tensors on {device}"
+
+    # Cross-tensor shape validation: every tensor below is indexed with
+    # coordinates derived from q, so a mismatched shape silently reads or
+    # writes out of place instead of failing.
+    assert kv.ndim == 2 and kv.shape[1] == head_dim, f"kv shape mismatch: expected (total_S_kv, {head_dim}), got {tuple(kv.shape)}"
+    expected_o_shape = (total_S_q, num_head, head_dim_v)
+    assert out.shape == expected_o_shape, f"out shape mismatch: expected {expected_o_shape}, got {tuple(out.shape)}"
+    assert dout.shape == expected_o_shape, f"dout shape mismatch: expected {expected_o_shape}, got {tuple(dout.shape)}"
+    assert lse.shape == (total_S_q, num_head), f"lse shape mismatch: expected {(total_S_q, num_head)}, got {tuple(lse.shape)}"
+    assert attn_sink.shape == (num_head,), f"attn_sink shape mismatch: expected {(num_head,)}, got {tuple(attn_sink.shape)}"
+    assert topk_idxs.ndim == 2 and topk_idxs.shape[0] == total_S_q, f"topk_idxs shape mismatch: expected ({total_S_q}, topk_max), got {tuple(topk_idxs.shape)}"
+    if topk_length is not None:
+        assert topk_length.dtype == torch.int32, f"topk_length dtype mismatch: expected torch.int32, got {topk_length.dtype}"
+        assert topk_length.shape == (total_S_q,), f"topk_length shape mismatch: expected {(total_S_q,)}, got {tuple(topk_length.shape)}"
 
     if softmax_scale is None:
         softmax_scale = 1.0 / math.sqrt(head_dim)
@@ -76,58 +94,75 @@ def flash_attn_bwd_sm100(
     num_head_blocks = (num_head + block_tile - 1) // block_tile
     batch_size = 1
 
-    # Ensure contiguous
-    q, kv, out, dout = [t.contiguous() for t in (q, kv, out, dout)]
-    lse = lse.contiguous()
-
-    # Allocate output tensors
-    if dq is None:
-        dq = torch.empty_like(q)
-    else:
-        assert dq.shape == q.shape, f"dq shape mismatch: expected {q.shape}, got {dq.shape}"
-        assert dq.dtype == q.dtype, f"dq dtype mismatch: expected {q.dtype}, got {dq.dtype}"
-        assert dq.device == device, f"dq device mismatch: expected {device}, got {dq.device}"
-    if dkv is None:
-        dkv = torch.zeros(total_S_kv, head_dim, dtype=kv.dtype, device=device)
-    else:
-        expected_dkv_shape = (total_S_kv, head_dim)
-        assert dkv.shape == expected_dkv_shape, f"dkv shape mismatch: expected {expected_dkv_shape}, got {dkv.shape}"
-        assert dkv.dtype == kv.dtype, f"dkv dtype mismatch: expected {kv.dtype}, got {dkv.dtype}"
-        assert dkv.device == device, f"dkv device mismatch: expected {device}, got {dkv.device}"
-        dkv.fill_(0)
-    d_sink = torch.zeros_like(attn_sink)
-
-    # Allocate workspace tensors
-    acc_dtype = cutlass.Float32
-    ws_lse_odo_shape = FlashAttentionDSABackwardSm100._get_workspace_size_LSE_OdO(
-        total_S_q,
-        head_dim,
-        num_head,
-        batch_size,
-        acc_dtype,
-    )
-    workspace_LSE_OdO = torch.zeros(
-        *ws_lse_odo_shape,
-        dtype=torch.uint8,
-        device=device,
-    )
-
-    ws_dkv_shape = FlashAttentionDSABackwardSm100._get_workspace_size_dKV(
-        total_S_kv,
-        head_dim,
-        batch_size,
-        acc_dtype,
-    )
-    workspace_dKV = torch.zeros(
-        *ws_dkv_shape,
-        dtype=torch.uint8,
-        device=device,
-    )
+    current_stream = resolve_stream(current_stream)
+
+    # Normalize inputs and allocate outputs/workspaces on the execution stream:
+    # the kernel below launches on `current_stream`, so the semantically
+    # required zero-initialization of dkv/d_sink and both workspaces (and any
+    # contiguity copies) must be stream-ordered with it, not with the ambient
+    # torch stream the caller happens to be on.
+    with torch_stream_context(current_stream):
+        # Ensure contiguous
+        q, kv, out, dout = [t.contiguous() for t in (q, kv, out, dout)]
+        lse = lse.contiguous()
+        attn_sink = attn_sink.contiguous()
+        topk_idxs = topk_idxs.contiguous()
+        if topk_length is not None:
+            topk_length = topk_length.contiguous()
+
+        # Allocate output tensors
+        if dq is None:
+            dq = torch.empty_like(q)
+        else:
+            assert dq.shape == q.shape, f"dq shape mismatch: expected {q.shape}, got {dq.shape}"
+            assert dq.dtype == q.dtype, f"dq dtype mismatch: expected {q.dtype}, got {dq.dtype}"
+            assert dq.device == device, f"dq device mismatch: expected {device}, got {dq.device}"
+            # The compile cache is keyed without output strides, so a caller
+            # provided output must match the contiguous layout the kernel was
+            # compiled for (it is not copied: that would break out-parameter
+            # identity).
+            assert dq.is_contiguous(), "dq must be contiguous"
+        if dkv is None:
+            dkv = torch.zeros(total_S_kv, head_dim, dtype=kv.dtype, device=device)
+        else:
+            expected_dkv_shape = (total_S_kv, head_dim)
+            assert dkv.shape == expected_dkv_shape, f"dkv shape mismatch: expected {expected_dkv_shape}, got {dkv.shape}"
+            assert dkv.dtype == kv.dtype, f"dkv dtype mismatch: expected {kv.dtype}, got {dkv.dtype}"
+            assert dkv.device == device, f"dkv device mismatch: expected {device}, got {dkv.device}"
+            assert dkv.is_contiguous(), "dkv must be contiguous"
+            dkv.fill_(0)
+        d_sink = torch.zeros_like(attn_sink)
+
+        # Allocate workspace tensors
+        acc_dtype = cutlass.Float32
+        ws_lse_odo_shape = FlashAttentionDSABackwardSm100._get_workspace_size_LSE_OdO(
+            total_S_q,
+            head_dim,
+            num_head,
+            batch_size,
+            acc_dtype,
+        )
+        workspace_LSE_OdO = torch.zeros(
+            *ws_lse_odo_shape,
+            dtype=torch.uint8,
+            device=device,
+        )
+
+        ws_dkv_shape = FlashAttentionDSABackwardSm100._get_workspace_size_dKV(
+            total_S_kv,
+            head_dim,
+            batch_size,
+            acc_dtype,
+        )
+        workspace_dKV = torch.zeros(
+            *ws_dkv_shape,
+            dtype=torch.uint8,
+            device=device,
+        )
 
     problem_shape = (total_S_q, total_S_kv, head_dim, (num_head, batch_size))
 
     dtype = torch2cute_dtype_map[q.dtype]
-    current_stream = resolve_stream(current_stream)
 
     has_topk_length = topk_length is not None
     max_topk = topk_idxs.shape[1]
@@ -149,6 +184,7 @@ def flash_attn_bwd_sm100(
         workspace_dKV_tensor = to_cute_tensor(workspace_dKV)
 
         kernel_obj = FlashAttentionDSABackwardSm100(
+            element_dtype=dtype,
             head_dim=head_dim,
             head_dim_v=head_dim_v,
             block_tile=block_tile,
diff --git a/python/cudnn/deepseek_sparse_attention/sparse_attention_backward/api.py b/python/cudnn/deepseek_sparse_attention/sparse_attention_backward/api.py
index 7ce340394..5fb2dab2e 100644
--- a/python/cudnn/deepseek_sparse_attention/sparse_attention_backward/api.py
+++ b/python/cudnn/deepseek_sparse_attention/sparse_attention_backward/api.py
@@ -20,8 +20,8 @@
 class SparseAttentionBackward(APIBase):
     def __init__(
         self,
-        sample_q: torch.Tensor,  # (total_S_q, H, D) BF16
-        sample_kv: torch.Tensor,  # (total_S_kv, D) BF16 (K=V)
+        sample_q: torch.Tensor,  # (total_S_q, H, D) FP16/BF16
+        sample_kv: torch.Tensor,  # (total_S_kv, D) FP16/BF16 (K=V)
         sample_out: torch.Tensor,  # (total_S_q, H, D_v)
         sample_dout: torch.Tensor,  # (total_S_q, H, D_v)
         sample_lse: torch.Tensor,  # (total_S_q, H) FP32, KV-only LSE
@@ -69,6 +69,79 @@ def check_support(self) -> bool:
         self._check_dtype(self.lse_desc, torch.float32, name="LSE")
         self._check_dtype(self.attn_sink_desc, torch.float32, name="attn_sink")
         self._check_dtype(self.topk_idxs_desc, torch.int32, name="topk_idxs")
+        self._check_dtype(self.out_desc, self.q_desc.dtype, name="out", extra_error_msg="out must have same dtype as Q")
+        self._check_dtype(self.dout_desc, self.q_desc.dtype, name="dout", extra_error_msg="dout must have same dtype as Q")
+        if self.topk_length_desc is not None:
+            self._check_dtype(self.topk_length_desc, torch.int32, name="topk_length")
+
+        # Device placement + cross-tensor device consistency. The SM90/SM100
+        # kernels are CUDA-only and reject CPU or cross-device inputs at
+        # execution time (see the is_cuda / same-device assert in
+        # ``_interface_sm100.flash_attn_bwd_sm100``), so a placement mismatch
+        # must fail the support gate here rather than compile/launch and crash.
+        ref_device = self.q_desc.device
+        descriptors = [
+            self.q_desc,
+            self.kv_desc,
+            self.out_desc,
+            self.dout_desc,
+            self.lse_desc,
+            self.attn_sink_desc,
+            self.topk_idxs_desc,
+        ]
+        if self.topk_length_desc is not None:
+            descriptors.append(self.topk_length_desc)
+        self._value_error_if(
+            ref_device.type != "cuda",
+            f"Q must live on CUDA, got {ref_device}",
+        )
+        self._value_error_if(
+            any(desc.device != ref_device for desc in descriptors),
+            f"All inputs must share Q's device {ref_device}, got {[desc.device for desc in descriptors]}",
+        )
+
+        # Cross-tensor shape contract: every companion tensor is indexed with
+        # coordinates derived from Q, so a mismatched shape silently reads or
+        # writes out of place at execution time instead of failing.
+        total_s_q, num_heads, head_dim = self.q_desc.shape
+        # The SM100 kernel is tiled only for head_dim in {512, 576} (the 576
+        # MLA case splits QK=576 / V=512); any other head_dim compiles to a
+        # layout that indexes shared memory out of bounds and crashes.
+        self._value_error_if(
+            head_dim not in (512, 576),
+            f"head_dim must be 512 or 576, got {head_dim}",
+        )
+        head_dim_v = 512 if head_dim == 576 else head_dim
+        expected_o_shape = (total_s_q, num_heads, head_dim_v)
+        self._value_error_if(
+            self.kv_desc.shape[1] != head_dim,
+            f"KV must have shape (total_S_kv, {head_dim}), got {self.kv_desc.shape}",
+        )
+        self._value_error_if(
+            self.out_desc.shape != expected_o_shape,
+            f"out must have shape {expected_o_shape}, got {self.out_desc.shape}",
+        )
+        self._value_error_if(
+            self.dout_desc.shape != expected_o_shape,
+            f"dout must have shape {expected_o_shape}, got {self.dout_desc.shape}",
+        )
+        self._value_error_if(
+            self.lse_desc.shape != (total_s_q, num_heads),
+            f"LSE must have shape {(total_s_q, num_heads)}, got {self.lse_desc.shape}",
+        )
+        self._value_error_if(
+            self.attn_sink_desc.shape != (num_heads,),
+            f"attn_sink must have shape {(num_heads,)}, got {self.attn_sink_desc.shape}",
+        )
+        self._value_error_if(
+            self.topk_idxs_desc.ndim != 2 or self.topk_idxs_desc.shape[0] != total_s_q,
+            f"topk_idxs must have shape ({total_s_q}, topk_max), got {self.topk_idxs_desc.shape}",
+        )
+        if self.topk_length_desc is not None:
+            self._value_error_if(
+                self.topk_length_desc.shape != (total_s_q,),
+                f"topk_length must have shape {(total_s_q,)}, got {self.topk_length_desc.shape}",
+            )
 
         self._is_supported = True
         return True
diff --git a/python/cudnn/deepseek_sparse_attention/sparse_attention_backward/dsa_bwd_sm100.py b/python/cudnn/deepseek_sparse_attention/sparse_attention_backward/dsa_bwd_sm100.py
index 5d78f31f0..3cb138714 100644
--- a/python/cudnn/deepseek_sparse_attention/sparse_attention_backward/dsa_bwd_sm100.py
+++ b/python/cudnn/deepseek_sparse_attention/sparse_attention_backward/dsa_bwd_sm100.py
@@ -4,7 +4,7 @@
 
 import cutlass
 import cutlass.cute as cute
-from cutlass.cute.typing import Float32, Int32, BFloat16, Int64
+from cutlass.cute.typing import Float32, Int32, Int64
 import cutlass.pipeline as pipeline
 from cutlass.cute.nvgpu import OperandMajorMode, cpasync, tcgen05
 import cutlass.utils.blackwell_helpers as sm100_utils
@@ -17,6 +17,7 @@ class FlashAttentionDSABackwardSm100:
 
     def __init__(
         self,
+        element_dtype: Type[cutlass.Numeric],
         head_dim: int,
         head_dim_v: int,
         block_tile: int,
@@ -40,9 +41,11 @@ def __init__(
         self.QdS_cta_tiler = (head_dim_main, block_tile, block_tile)
         self.cluster_shape_mn = (1, 1)
 
-        self.element_dtype = BFloat16
-        # User constraint: dKV accumulation must stay FP32. BFloat16 is only
-        # used for element/output storage.
+        if element_dtype not in [cutlass.Float16, cutlass.BFloat16]:
+            raise ValueError(f"Unsupported element dtype: {element_dtype}")
+        self.element_dtype = element_dtype
+        # dKV accumulation stays FP32; element_dtype controls element/output
+        # storage.
         self.acc_dtype = Float32
 
         # =============== Sum OdO ================
diff --git a/test/python/fe_api/dsa/dsa_utils.py b/test/python/fe_api/dsa/dsa_utils.py
index 269f83c13..134b35d9e 100644
--- a/test/python/fe_api/dsa/dsa_utils.py
+++ b/test/python/fe_api/dsa/dsa_utils.py
@@ -16,6 +16,11 @@ def _require_sm90():
         pytest.skip("SM90 GPU required")
 
 
+def _require_sm100():
+    if not torch.cuda.is_available() or torch.cuda.get_device_capability()[0] != 10:
+        pytest.skip("SM100 GPU required")
+
+
 # Parameterization marks shared by every DSA test
 DSA_PARAM_MARKS = [
     pytest.mark.parametrize("dtype", [torch.bfloat16]),
diff --git a/test/python/fe_api/dsa/test_DSA_sparse_attention_backward.py b/test/python/fe_api/dsa/test_DSA_sparse_attention_backward.py
index d0e6548e4..66e5d6704 100644
--- a/test/python/fe_api/dsa/test_DSA_sparse_attention_backward.py
+++ b/test/python/fe_api/dsa/test_DSA_sparse_attention_backward.py
@@ -14,6 +14,7 @@
 
 from fe_api.dsa.dsa_utils import (
     _require_sm90,
+    _require_sm100,
     dsa_init,
     with_dsa_sparse_attention_backward_params,
 )
@@ -302,3 +303,450 @@ def rel_l2(a, b):
 
     assert rel_l2(dkv, dkv_ref) < 1e-4, "dkv parity vs stage-1 baseline"
     assert rel_l2(d_sink, d_sink_ref) < 1e-4, "d_sink parity vs stage-1 baseline"
+
+
+@pytest.mark.L0
+@pytest.mark.gpu_exclusive
+@pytest.mark.xdist_group(name="gpu_exclusive")
+@torch_fork_set_rng(seed=7)
+def test_DSA_sparse_attention_backward_nondefault_stream_zero_init_ordering():
+    """The SM100 interface allocates and zero-initializes dq/dkv/d_sink and the
+    two workspaces with plain torch calls, which enqueue on the ambient torch
+    stream, while the kernel launches on the caller-provided ``current_stream``.
+    Without explicit stream scoping the two are unordered: with a busy ambient
+    stream, the semantically required zero-fills land *after* the kernel and
+    wipe the dkv/d_sink accumulation (or, in the other interleaving, the kernel
+    accumulates on top of uninitialized memory).
+
+    The ambient default stream is parked on ``torch.cuda._sleep`` so the
+    unordered interleaving is reached reliably (the zero-fills cannot start
+    until the sleep retires, while the side-stream kernel is free to run);
+    the test needs the GPU to itself for that reason."""
+    try:
+        from cudnn import DSA
+        from cuda.bindings import driver as cuda
+    except ImportError:
+        pytest.skip("Environment not supported: cudnn[cutedsl] not installed")
+
+    _require_sm100()
+    device = torch.device("cuda")
+    s_q, s_kv, num_heads = 256, 1024, 64
+    head_dim, topk = 512, 64
+    softmax_scale = 1.0 / math.sqrt(head_dim)
+
+    q = torch.randn(s_q, num_heads, head_dim, dtype=torch.bfloat16, device=device) / 10
+    kv = torch.randn(s_kv, head_dim, dtype=torch.bfloat16, device=device) / 10
+    attn_sink = torch.randn(num_heads, dtype=torch.float32, device=device)
+    topk_idxs = torch.stack([torch.randperm(s_kv, device=device)[:topk] for _ in range(s_q)]).to(torch.int32)
+
+    out, lse = ref_sparse_attention_forward(
+        q,
+        kv,
+        attn_sink,
+        topk_idxs,
+        softmax_scale=softmax_scale,
+    )
+    dout = torch.randn_like(out)
+
+    def run(stream):
+        result = DSA.sparse_attention_backward_wrapper(
+            q,
+            kv,
+            out,
+            dout,
+            lse,
+            attn_sink,
+            topk_idxs,
+            softmax_scale=softmax_scale,
+            stream=stream,
+        )
+        torch.cuda.synchronize()
+        return result["dq"], result["dkv"], result["d_sink"]
+
+    # Control on the ambient (default) stream; also primes the compile cache
+    # so the raced call below is a pure execute.
+    default_stream = cuda.CUstream(torch.cuda.current_stream().cuda_stream)
+    dq_ref, dkv_ref, d_sink_ref = run(default_stream)
+    assert (dkv_ref != 0).any(), "control dkv must have nonzero content"
+    assert (d_sink_ref != 0).any(), "control d_sink must have nonzero content"
+
+    # Park the ambient default stream, then launch on a side stream. The
+    # interface's zero-fills must be ordered with the side-stream kernel, not
+    # queued behind the sleep on the default stream.
+    side_stream = torch.cuda.Stream()
+    torch.cuda._sleep(2_000_000_000)
+    dq, dkv, d_sink = run(cuda.CUstream(side_stream.cuda_stream))
+
+    assert (dkv != 0).any(), "dkv accumulation was wiped by a zero-init racing on another stream"
+    assert (d_sink != 0).any(), "d_sink accumulation was wiped by a zero-init racing on another stream"
+
+    def rel_l2(a, b):
+        return ((a.float() - b.float()).norm() / b.float().norm().clamp_min(1e-30)).item()
+
+    assert torch.equal(dq, dq_ref), "dq must not depend on the launch stream"
+    assert rel_l2(dkv, dkv_ref) < 1e-4, "dkv parity vs default-stream control"
+    assert rel_l2(d_sink, d_sink_ref) < 1e-4, "d_sink parity vs default-stream control"
+
+
+@pytest.mark.L0
+@torch_fork_set_rng(seed=16)
+def test_DSA_sparse_attention_backward_fp16_sm100_numerics():
+    """SM100 must compile FP16 inputs with FP16 MMA/storage semantics."""
+    try:
+        from cudnn import DSA
+    except ImportError:
+        pytest.skip("Environment not supported: cudnn[cutedsl] not installed")
+
+    _require_sm100()
+    device = torch.device("cuda")
+    s_q, s_kv, num_heads = 4, 128, 64
+    head_dim, topk = 512, 64
+    softmax_scale = 1.0 / math.sqrt(head_dim)
+
+    q = torch.randn(s_q, num_heads, head_dim, dtype=torch.float16, device=device)
+    kv = torch.randn(s_kv, head_dim, dtype=torch.float16, device=device)
+    attn_sink = torch.linspace(-2.0, 2.0, num_heads, dtype=torch.float32, device=device)
+    topk_idxs = torch.stack([torch.randperm(s_kv, device=device)[:topk] for _ in range(s_q)]).to(torch.int32)
+    topk_length = torch.tensor([16, 32, 48, 64], dtype=torch.int32, device=device)
+
+    out, lse = ref_sparse_attention_forward(
+        q,
+        kv,
+        attn_sink,
+        topk_idxs,
+        topk_length=topk_length,
+        softmax_scale=softmax_scale,
+    )
+    dout = torch.randn_like(out)
+    assert q.dtype == kv.dtype == out.dtype == dout.dtype == torch.float16
+    assert lse.dtype == torch.float32
+
+    result = DSA.sparse_attention_backward_wrapper(
+        q,
+        kv,
+        out,
+        dout,
+        lse,
+        attn_sink,
+        topk_idxs,
+        softmax_scale=softmax_scale,
+        topk_length=topk_length,
+    )
+    torch.cuda.synchronize()
+    dq, dkv, d_sink = result["dq"], result["dkv"], result["d_sink"]
+
+    assert dq.dtype == dkv.dtype == torch.float16
+    assert d_sink.dtype == torch.float32
+    assert torch.isfinite(dq).all()
+    assert torch.isfinite(dkv).all()
+    assert torch.isfinite(d_sink).all()
+    check_ref_dsa_sparse_attention_backward(
+        q,
+        kv,
+        attn_sink,
+        topk_idxs,
+        out,
+        dout,
+        lse,
+        dq,
+        dkv,
+        d_sink,
+        softmax_scale=softmax_scale,
+        topk_length=topk_length,
+        atol=5e-2,
+        rtol=5e-2,
+    )
+
+
+@pytest.mark.L0
+@torch_fork_set_rng(seed=23)
+def test_DSA_sparse_attention_backward_noncontiguous_aux_inputs():
+    """The interface normalizes q/kv/out/dout/lse to contiguous but not
+    attn_sink/topk_idxs/topk_length. Non-contiguous aux tensors previously
+    escaped down to the CuTe DSL layer and failed there with low-level stride
+    errors (a signature mismatch against the shared compile-cache entry on the
+    warm path, a leading-stride assert on the cold path). They must be
+    normalized like every other input; both cache paths are covered here."""
+    try:
+        from cudnn.deepseek_sparse_attention.sparse_attention_backward import _interface_sm100
+    except ImportError:
+        pytest.skip("Environment not supported: cudnn[cutedsl] not installed")
+
+    _require_sm100()
+    device = torch.device("cuda")
+    s_q, s_kv, num_heads = 256, 1024, 64
+    head_dim, topk = 576, 64
+    softmax_scale = 1.0 / math.sqrt(head_dim)
+
+    q = torch.randn(s_q, num_heads, head_dim, dtype=torch.bfloat16, device=device) / 10
+    kv = torch.randn(s_kv, head_dim, dtype=torch.bfloat16, device=device) / 10
+    attn_sink = torch.randn(num_heads, dtype=torch.float32, device=device)
+    topk_idxs = torch.stack([torch.randperm(s_kv, device=device)[:topk] for _ in range(s_q)]).to(torch.int32)
+    topk_length = torch.randint(1, topk + 1, (s_q,), dtype=torch.int32, device=device)
+
+    out, lse = ref_sparse_attention_forward(
+        q,
+        kv,
+        attn_sink,
+        topk_idxs,
+        topk_length=topk_length,
+        softmax_scale=softmax_scale,
+    )
+    dout = torch.randn_like(out)
+
+    def run(attn_sink_, topk_idxs_, topk_length_):
+        dq, dkv, d_sink = _interface_sm100.flash_attn_bwd_sm100(
+            q,
+            kv,
+            out,
+            dout,
+            lse,
+            attn_sink_,
+            topk_idxs_,
+            softmax_scale=softmax_scale,
+            topk_length=topk_length_,
+        )
+        torch.cuda.synchronize()
+        return dq.clone(), dkv.clone(), d_sink.clone()
+
+    def strided_copy_1d(t):
+        base = torch.zeros(2 * t.shape[0], dtype=t.dtype, device=t.device)
+        base[::2] = t
+        view = base[::2]
+        assert not view.is_contiguous() and torch.equal(view, t)
+        return view
+
+    def strided_copy_2d(t):
+        base = torch.zeros(t.shape[0], 2 * t.shape[1], dtype=t.dtype, device=t.device)
+        base[:, ::2] = t
+        view = base[:, ::2]
+        assert not view.is_contiguous() and torch.equal(view, t)
+        return view
+
+    def rel_l2(a, b):
+        return ((a.float() - b.float()).norm() / b.float().norm().clamp_min(1e-30)).item()
+
+    # Cold path: nothing cached for this compile key, first call is
+    # non-contiguous (previously a leading-stride error inside the DSL).
+    _interface_sm100.flash_attn_bwd_sm100.compile_cache.clear()
+    dq_cold, dkv_cold, d_sink_cold = run(strided_copy_1d(attn_sink), strided_copy_2d(topk_idxs), strided_copy_1d(topk_length))
+
+    # Contiguous control (same compile key, warm cache).
+    dq_ref, dkv_ref, d_sink_ref = run(attn_sink, topk_idxs, topk_length)
+
+    # Warm path: non-contiguous call against the cached contiguous signature
+    # (previously a signature stride mismatch inside the DSL).
+    dq, dkv, d_sink = run(strided_copy_1d(attn_sink), strided_copy_2d(topk_idxs), strided_copy_1d(topk_length))
+
+    for tag, (dq_t, dkv_t, d_sink_t) in {"cold": (dq_cold, dkv_cold, d_sink_cold), "warm": (dq, dkv, d_sink)}.items():
+        assert torch.equal(dq_t, dq_ref), f"{tag}: dq must not depend on aux input contiguity"
+        assert rel_l2(dkv_t, dkv_ref) < 1e-4, f"{tag}: dkv parity vs contiguous control"
+        assert rel_l2(d_sink_t, d_sink_ref) < 1e-4, f"{tag}: d_sink parity vs contiguous control"
+
+
+@pytest.mark.L0
+@torch_fork_set_rng(seed=42)
+def test_DSA_sparse_attention_backward_cross_shape_validation():
+    """The compiled kernel takes every dimension as a dynamic value derived
+    from q, so mis-shaped companion tensors do not fail: a transposed dout or
+    lse runs without any error and returns silently corrupted gradients
+    (measured rel-L2 vs the correct result: ~1.1 and ~45 respectively).
+    The interface must validate the cross-tensor contract up front. Caller
+    provided dq/dkv must additionally be contiguous: the compile cache is
+    keyed without output strides, so a strided out-parameter would be written
+    through the wrong layout."""
+    try:
+        from cudnn.deepseek_sparse_attention.sparse_attention_backward import _interface_sm100
+    except ImportError:
+        pytest.skip("Environment not supported: cudnn[cutedsl] not installed")
+
+    _require_sm100()
+    device = torch.device("cuda")
+    s_q, s_kv, num_heads = 4, 128, 64
+    head_dim, head_dim_v, topk = 576, 512, 64
+    softmax_scale = 1.0 / math.sqrt(head_dim)
+
+    q = torch.randn(s_q, num_heads, head_dim, dtype=torch.bfloat16, device=device)
+    kv = torch.randn(s_kv, head_dim, dtype=torch.bfloat16, device=device)
+    out = torch.randn(s_q, num_heads, head_dim_v, dtype=torch.bfloat16, device=device)
+    dout = torch.randn_like(out)
+    lse = torch.randn(s_q, num_heads, dtype=torch.float32, device=device)
+    attn_sink = torch.randn(num_heads, dtype=torch.float32, device=device)
+    topk_idxs = torch.stack([torch.randperm(s_kv, device=device)[:topk] for _ in range(s_q)]).to(torch.int32)
+    topk_length = torch.full((s_q,), topk, dtype=torch.int32, device=device)
+
+    good = dict(
+        q=q,
+        kv=kv,
+        out=out,
+        dout=dout,
+        lse=lse,
+        attn_sink=attn_sink,
+        topk_idxs=topk_idxs,
+        topk_length=topk_length,
+        dq=None,
+        dkv=None,
+    )
+
+    def call(args):
+        _interface_sm100.flash_attn_bwd_sm100(
+            args["q"],
+            args["kv"],
+            args["out"],
+            args["dout"],
+            args["lse"],
+            args["attn_sink"],
+            args["topk_idxs"],
+            softmax_scale=softmax_scale,
+            topk_length=args["topk_length"],
+            dq=args["dq"],
+            dkv=args["dkv"],
+        )
+
+    shape_cases = {
+        "kv": kv[:, : head_dim - 64].contiguous(),
+        "out": out.transpose(1, 2).contiguous(),
+        "dout": dout.transpose(1, 2).contiguous(),
+        "lse": lse.transpose(0, 1).contiguous(),
+        "attn_sink": attn_sink[: num_heads // 2].contiguous(),
+        "topk_idxs": topk_idxs[: s_q - 1].contiguous(),
+        "topk_length": topk_length[: s_q - 1].contiguous(),
+    }
+    for name, bad_tensor in shape_cases.items():
+        args = dict(good)
+        args[name] = bad_tensor
+        with pytest.raises(AssertionError, match=f"{name} shape mismatch"):
+            call(args)
+
+    args = dict(good)
+    args["topk_length"] = topk_length.to(torch.int64)
+    with pytest.raises(AssertionError, match="topk_length dtype mismatch"):
+        call(args)
+
+    # Caller-provided out-params must be contiguous (they are not copied).
+    dq_strided = torch.empty(s_q, num_heads, 2 * head_dim, dtype=q.dtype, device=device)[..., ::2]
+    args = dict(good)
+    args["dq"] = dq_strided
+    with pytest.raises(AssertionError, match="dq must be contiguous"):
+        call(args)
+    dkv_strided = torch.empty(s_kv, 2 * head_dim, dtype=kv.dtype, device=device)[:, ::2]
+    args = dict(good)
+    args["dkv"] = dkv_strided
+    with pytest.raises(AssertionError, match="dkv must be contiguous"):
+        call(args)
+
+
+@pytest.mark.L0
+@torch_fork_set_rng(seed=51)
+def test_DSA_sparse_attention_backward_check_support_validates_contract():
+    """check_support works on metadata-only descriptors and is the advertised
+    support gate, so the cross-tensor contract must be enforced there as well,
+    not only by the runtime asserts in the execution interface."""
+    try:
+        from cudnn.deepseek_sparse_attention.sparse_attention_backward import SparseAttentionBackward
+    except ImportError:
+        pytest.skip("Environment not supported: cudnn[cutedsl] not installed")
+
+    _require_sm100()
+    device = torch.device("cuda")
+    s_q, s_kv, num_heads = 4, 128, 64
+    head_dim, head_dim_v, topk = 576, 512, 64
+
+    q = torch.randn(s_q, num_heads, head_dim, dtype=torch.bfloat16, device=device)
+    kv = torch.randn(s_kv, head_dim, dtype=torch.bfloat16, device=device)
+    out = torch.randn(s_q, num_heads, head_dim_v, dtype=torch.bfloat16, device=device)
+    dout = torch.randn_like(out)
+    lse = torch.randn(s_q, num_heads, dtype=torch.float32, device=device)
+    attn_sink = torch.randn(num_heads, dtype=torch.float32, device=device)
+    topk_idxs = torch.stack([torch.randperm(s_kv, device=device)[:topk] for _ in range(s_q)]).to(torch.int32)
+    topk_length = torch.full((s_q,), topk, dtype=torch.int32, device=device)
+
+    good = dict(
+        sample_q=q,
+        sample_kv=kv,
+        sample_out=out,
+        sample_dout=dout,
+        sample_lse=lse,
+        sample_attn_sink=attn_sink,
+        sample_topk_idxs=topk_idxs,
+        sample_topk_length=topk_length,
+    )
+    assert SparseAttentionBackward(**good).check_support()
+
+    fp16_good = dict(good)
+    for name in ("sample_q", "sample_kv", "sample_out", "sample_dout"):
+        fp16_good[name] = fp16_good[name].to(torch.float16)
+    assert SparseAttentionBackward(**fp16_good).check_support()
+
+    bad_cases = {
+        "sample_kv": kv[:, : head_dim - 64].contiguous(),
+        "sample_out": out.transpose(1, 2).contiguous(),
+        "sample_dout": dout.transpose(1, 2).contiguous(),
+        "sample_lse": lse.transpose(0, 1).contiguous(),
+        "sample_attn_sink": attn_sink[: num_heads // 2].contiguous(),
+        "sample_topk_idxs": topk_idxs[: s_q - 1].contiguous(),
+        "sample_topk_length": topk_length[: s_q - 1].contiguous(),
+    }
+    for name, bad_tensor in bad_cases.items():
+        kwargs = dict(good)
+        kwargs[name] = bad_tensor
+        with pytest.raises(ValueError):
+            SparseAttentionBackward(**kwargs).check_support()
+
+    # Device placement: check_support must reject CPU inputs that the
+    # SM90/SM100 runtime would otherwise reject at launch time.
+    cpu_kwargs = {name: tensor.to("cpu") for name, tensor in good.items()}
+    with pytest.raises(ValueError, match="Q must live on CUDA"):
+        SparseAttentionBackward(**cpu_kwargs).check_support()
+
+    # Cross-device: Q stays on CUDA, KV moved to CPU (no new CUDA allocation;
+    # reuses the good tensors) -> device-consistency failure.
+    cross_kwargs = dict(good)
+    cross_kwargs["sample_kv"] = good["sample_kv"].to("cpu")
+    with pytest.raises(ValueError, match="must share Q's device"):
+        SparseAttentionBackward(**cross_kwargs).check_support()
+
+    # head_dim must be one of {512, 576}: the kernel is tiled only for those.
+    bad_head_dim = dict(good)
+    bad_head_dim["sample_q"] = torch.randn(s_q, num_heads, 128, dtype=torch.bfloat16, device=device)
+    bad_head_dim["sample_kv"] = torch.randn(s_kv, 128, dtype=torch.bfloat16, device=device)
+    with pytest.raises(ValueError, match="head_dim must be 512 or 576"):
+        SparseAttentionBackward(**bad_head_dim).check_support()
+
+
+@pytest.mark.L0
+@torch_fork_set_rng(seed=7)
+def test_DSA_sparse_attention_backward_rejects_unsupported_head_dim_runtime():
+    """The SM100 kernel is tiled only for head_dim in {512, 576}; any other
+    head_dim indexes shared memory out of bounds and crashes inside the kernel.
+    The interface must reject it before any compile/launch."""
+    try:
+        from cudnn.deepseek_sparse_attention.sparse_attention_backward import _interface_sm100
+    except ImportError:
+        pytest.skip("Environment not supported: cudnn[cutedsl] not installed")
+
+    _require_sm100()
+    device = torch.device("cuda")
+    s_q, s_kv, num_heads = 4, 128, 64
+    head_dim, head_dim_v, topk = 128, 128, 64
+    softmax_scale = 1.0 / math.sqrt(head_dim)
+
+    q = torch.randn(s_q, num_heads, head_dim, dtype=torch.bfloat16, device=device)
+    kv = torch.randn(s_kv, head_dim, dtype=torch.bfloat16, device=device)
+    out = torch.randn(s_q, num_heads, head_dim_v, dtype=torch.bfloat16, device=device)
+    dout = torch.randn_like(out)
+    lse = torch.randn(s_q, num_heads, dtype=torch.float32, device=device)
+    attn_sink = torch.randn(num_heads, dtype=torch.float32, device=device)
+    topk_idxs = torch.stack([torch.randperm(s_kv, device=device)[:topk] for _ in range(s_q)]).to(torch.int32)
+
+    with pytest.raises(AssertionError, match="head_dim must be 512 or 576"):
+        _interface_sm100.flash_attn_bwd_sm100(
+            q,
+            kv,
+            out,
+            dout,
+            lse,
+            attn_sink,
+            topk_idxs,
+            softmax_scale=softmax_scale,
+        )