diff --git a/python/cudnn/AGENTS.md b/python/cudnn/AGENTS.md
index 82573f1a1..83ec7420e 100644
--- a/python/cudnn/AGENTS.md
+++ b/python/cudnn/AGENTS.md
@@ -96,22 +96,23 @@ neither names the thing that breaks it most directly: a device-to-host read.
 Known violations, all pre-existing and each needing a kernel-side change, so
 none is precedent:
 
-- THD `cu_seqlens` host cumsum (`sdpa/fwd/api_dsl.py`, `_execute_thd` on both
-  SM100 and SM120). RESOLVED on all SM100 f16 families
-  (issue #552) — the reference for the SM120 port: the kernels compile
-  with dynamic token extents (Rule 4), the setup kernel builds the metadata
-  buffer device-side from the caller's length tensors, every ragged view
-  binds its buffer's capacity, and the grid is the plan-time envelope
-  (`b * ceil(s_q_declared / CGA_TILE_M) * qh`) with dead units exiting via
-  `_thd_decode`'s `batch == n_batch` sentinel — zero host reads, pinned by
-  the sync-debug and CUDA-graph capture tests. Still open: SM120 (exact
-  host-computed grid + host-built metadata; its own engine class and grid
-  mechanism, so its own port), and
-  the envelope's dead-tile tax under far-oversized declarations (measured
-  ~145 ns/dead unit exposed, <0.2% at realistic declarations; a capped
-  persistent grid reading a device-side live-unit count would bound it by
-  resident clusters). `sdpa_fwd_wrapper_sm80` shows the caller-provided
-  alternative — it requires `max_s_q` rather than deriving it.
+- THD `cu_seqlens` host cumsum (`sdpa/fwd/api_dsl.py`). RESOLVED
+  (issue #552) on every forward THD engine — SM100 f16 (all families) and
+  SM120 f16/fp8 — and the reference for future ports: the kernels compile
+  with dynamic token extents (Rule 4), a setup kernel builds the metadata
+  buffer device-side from the caller's length tensors (cu prefixes
+  normalized), every ragged view binds its buffer's capacity, and the grid
+  is the plan-time declared-S_q envelope — SM100 with dead units exiting
+  via `_thd_decode`'s `batch == n_batch` sentinel, SM120 with its
+  per-sequence rectangular grid whose past-the-length tiles always drained
+  without stores. Zero host reads, pinned per arch by the sync-debug and
+  CUDA-graph capture tests; `_thd_host_lens` is gone from the adapter.
+  Remaining follow-up: the envelope's dead-tile tax under far-oversized
+  declarations (measured ~145 ns/dead unit exposed on SM100, <0.2% at
+  realistic declarations; a capped persistent grid reading a device-side
+  live-unit count would bound it by resident clusters).
+  `sdpa_fwd_wrapper_sm80` shows the caller-provided alternative — it
+  requires `max_s_q` rather than deriving it.
 - Per-tensor FP8 descale readback (`_scalar` in the same file): fold on device,
   passing the pointers, as the backend FP8 sdpa does.
 - The FP8/MXFP8 `seq_len_q` guard in `sdpa/fwd/engines.py`. This one cannot be
diff --git a/python/cudnn/sdpa/fwd/api_dsl.py b/python/cudnn/sdpa/fwd/api_dsl.py
index fd9bbc1f5..21b96c562 100644
--- a/python/cudnn/sdpa/fwd/api_dsl.py
+++ b/python/cudnn/sdpa/fwd/api_dsl.py
@@ -337,8 +337,8 @@ def __init__(
         self.seq_q_lens_present = bool(seq_q_lens_present)
         # cu_seq_len form (cuDNN 9.24+): the corresponding seq-lens execute
         # argument arrives as a (B+1,)-int32 PREFIX-SUM tensor instead of
-        # (B,) per-batch lengths. THD-only today: the ragged lowering derives
-        # both forms host-side from its inherent tolist round-trip; the dense
+        # (B,) per-batch lengths. THD-only today: the setup kernels consume
+        # either form on device (issue #552); the dense
         # kernels have no CU read mode yet (check_support rejects).
         self.cu_seq_q_lens = bool(cu_seq_q_lens)
         self.cu_seq_kv_lens = bool(cu_seq_kv_lens)
@@ -566,9 +566,9 @@ def _checked_cu_seq_lens(self, cu_seq_lens: torch.Tensor, name: str) -> torch.Te
         """Validate a caller-provided (B+1,)-int32 prefix-sum tensor (cu_seq_len form).
 
         Strictly a view, like :meth:`_checked_seq_lens`. The prefix-sum
-        INVARIANTS (starts at 0, non-decreasing) are runtime values — they are
-        validated host-side by the THD lowering's inherent tolist round-trip,
-        not here.
+        INVARIANTS (non-decreasing; any base — the setup kernel normalizes by
+        subtracting element 0) are runtime values and caller contract: a
+        validation that needs a device read is not a validation (Rule 3).
         """
         self._value_error_if(
             cu_seq_lens.dtype != torch.int32,
@@ -584,27 +584,6 @@ def _checked_cu_seq_lens(self, cu_seq_lens: torch.Tensor, name: str) -> torch.Te
         )
         return cu_seq_lens.reshape(-1)
 
-    def _thd_host_lens(self, seq_lens, name: str, cu_form: bool) -> tuple[list, list]:
-        """One inherent D2H round-trip -> (per-batch lens, prefix sums) host lists.
-
-        Consumes EITHER length form: per-batch ``(B,)`` lengths (prefix sums
-        built by a Python scan) or the ``(B+1,)`` cu_seq_len prefix-sum form
-        (lengths are adjacent differences; the prefix-sum invariants are
-        validated here, where they are free to check).
-        """
-        if cu_form:
-            cu_host = [int(x) for x in self._checked_cu_seq_lens(seq_lens, name).tolist()]
-            self._value_error_if(
-                cu_host[0] != 0 or any(cu_host[i] > cu_host[i + 1] for i in range(len(cu_host) - 1)),
-                f"{name} must be a non-decreasing prefix sum starting at 0; got {cu_host}",
-            )
-            return [cu_host[i + 1] - cu_host[i] for i in range(len(cu_host) - 1)], cu_host
-        lens_host = [int(x) for x in self._checked_seq_lens(seq_lens, name).tolist()]
-        cu_host = [0]
-        for n in lens_host:
-            cu_host.append(cu_host[-1] + n)
-        return lens_host, cu_host
-
     def _check_seq_lens_contract(self, seq_q_lens, seq_kv_lens) -> None:
         """Reject seq-length tensors inconsistent with the compiled specialization.
 
@@ -1011,8 +990,8 @@ def scratch_workspace_bytes(self) -> int:
             # caller's ragged Stats buffer directly (token-major (T, H) or
             # head-major (H, head_stride)); without one it compiles with
             # has_lse=False and no LSE buffer exists at all. No slq/slk
-            # copies either: the metadata is built host-side from the tolist
-            # round-trip and uploaded in one H2D copy. o_desc: 16 int64 per
+            # copies either: the metadata is built DEVICE-side by the setup
+            # kernel (issue #552). o_desc: 16 int64 per
             # sequence + 16 spare, the per-sequence O TMA descriptors the
             # builder kernel fills.
             return ws_align((3 * b + 2) * 4) + ws_align((b * 16 + 16) * 8) + (0 if self.has_sink else ws_align(qh * 4))
@@ -2266,7 +2245,10 @@ def execute(
             seq_q_lens,
             seq_kv_lens,
             cutlass.Float32(scale_softmax_log2),
-            cutlass.Int32(0),  # thd_max_sq: THD-only runtime grid extent
+            cutlass.Int32(0),  # thd_max_sq: THD-only plan-time envelope grid extent
+            None,  # thd_q_lens / thd_kv_lens / thd_lens_form: THD-only, folded out
+            None,
+            None,
             current_stream,
         )
         if o_needs_copy_back:
@@ -2350,13 +2332,11 @@ def _scalar(t, default=1.0):
             # This kernel's ragged LSE store is head-major (H, head_stride):
             # tokens contiguous within a head row. (The f16 cell specializes on
             # either layout; check_support declines token-major here.)
+            # head_stride covering the packed total is caller contract — the
+            # total is a device value (Rule 3).
             lse = None
             if lse_tensor is not None:
                 head_stride = self.thd_stats_head_stride
-                self._value_error_if(
-                    head_stride < pack.t_q,
-                    f"head-major THD LSE head_stride ({head_stride}) must cover the packed Q token total ({pack.t_q})",
-                )
                 lse = lse_tensor.as_strided((self.h_q, head_stride), (head_stride, 1), lse_tensor.storage_offset())
         else:
             lse = self._checked_lse_view(lse_tensor) if lse_tensor is not None else None
@@ -2387,6 +2367,9 @@ def _scalar(t, default=1.0):
             cutlass.Float32(o_scale_fused),
             cutlass.Float32(ss),
             cutlass.Int32(pack.max_sq if pack is not None else 0),
+            pack.q_lens_dev if pack is not None else None,
+            pack.kv_lens_dev if pack is not None else None,
+            cutlass.Int32(pack.lens_form) if pack is not None else None,
             current_stream,
         )
         # Both of these consume what the kernel just wrote, so they belong on
@@ -2434,22 +2417,24 @@ def _key(desc):
         )
         return kwargs
 
-    def _thd_pack(self, q_buf, k_buf, v_buf, o_buf, seq_q_lens, seq_kv_lens, workspace, label, declared_views=False, current_stream=None):
-        """Shared THD (ragged) packing: cu_seqlens metadata + ``(1, T, H, D)`` views.
+    def _thd_pack(self, q_buf, k_buf, v_buf, o_buf, seq_q_lens, seq_kv_lens, workspace, label, declared_views=False, current_stream=None, lse_tokens_cap=None):
+        """Shared THD (ragged) packing: metadata buffer + ``(1, T, H, D)`` views.
 
         Serves the same fully-packed contract as the SM100 THD path
-        (``ragged_offset == cumsum(seq_len) * H * D`` from 0, multiplier 1);
-        the offsets are re-derived from ``seq_len_q``/``seq_len_kv``. The two
-        ``.tolist()`` D2H syncs feed the metadata upload, the ragged views'
-        extents and the launch grid (they no longer key any compile — the
-        kernels compile with DYNAMIC token extents, issue #552; removing the
-        sync itself needs the plan-time-max grid redesign tracked there).
-        The torch work (allocation, length
-        reads, the H2D upload) runs on ``current_stream`` — the LAUNCH
-        stream — so it is ordered against the kernel that consumes it.
-
-        Returns ``None`` when the packed Q total is zero (nothing to launch).
-        """
+        (``ragged_offset == cumsum(seq_len) * H * D`` from 0, multiplier 1).
+        Host round-trips (issue #552): NONE — the metadata buffer is built
+        DEVICE-side by the kernels' setup launch from the caller's length
+        tensors (returned as ``q_lens_dev``/``kv_lens_dev`` + the
+        ``lens_form`` bitmask), every ragged view binds its buffer's
+        CAPACITY (the kernels read the real per-sequence lengths from the
+        device metadata; ``lse_tokens_cap`` joins the Q/O floor when the LSE
+        shares their dynamic token symbol), and ``max_sq`` is the PLAN-TIME
+        declared S_q envelope that sizes the per-sequence grid — tiles past
+        a sequence's real length drain without loads or stores. The torch
+        allocation runs on ``current_stream`` — the LAUNCH stream — so it is
+        ordered against the kernels that consume it.
+
+        Returns ``None`` when no Q token is addressable (zero capacity)."""
         b, qh, kh = self.batch_size, self.h_q, self.h_kv
         d_qk, d_v = self.head_dim_qk, self.head_dim_v
         dev = q_buf.device
@@ -2458,19 +2443,28 @@ def _thd_pack(self, q_buf, k_buf, v_buf, o_buf, seq_q_lens, seq_kv_lens, workspa
         # [seq_kv(B) | cu_q(B+1) | cu_k(B+1)] — bound as the kernel's
         # seq_kv_lens tensor; the leading B words alias the per-sequence KV
         # lengths so the kernel's existing padded-mask read works unchanged.
-        # Built HOST-side from the (inherent) tolist round-trip and uploaded
-        # in ONE H2D copy: a device-side cumsum would allocate its scan-temp
-        # storage and launch kernels on the execute hot path. Either length
-        # form feeds it — per-batch (B,) lengths or the (B+1,) cu_seq_len
-        # prefix sums — at identical cost.
+        # WRITTEN by the setup kernel; the prefix-sum invariants are caller
+        # contract (a validation that needs a device read is not a
+        # validation — AGENTS.md Rule 3; cu prefixes are normalized by the
+        # setup kernel).
         with _torch_stream_context(current_stream, dev):
             meta = carver.take(3 * b + 2, torch.int32) if carver is not None else torch.empty(3 * b + 2, dtype=torch.int32, device=dev)
-            slq_host, cu_q_host = self._thd_host_lens(seq_q_lens, "cu_seq_len_q" if self.cu_seq_q_lens else "seq_q_lens", self.cu_seq_q_lens)
-            slk_host, cu_k_host = self._thd_host_lens(seq_kv_lens, "cu_seq_len_kv" if self.cu_seq_kv_lens else "seq_kv_lens", self.cu_seq_kv_lens)
-            meta.copy_(torch.tensor(slk_host + cu_q_host + cu_k_host, dtype=torch.int32))
-        t_q = cu_q_host[-1]
-        t_kv = cu_k_host[-1]
-        max_sq = max(slq_host) if slq_host else 0
+        q_lens_dev = self._checked_cu_seq_lens(seq_q_lens, "cu_seq_len_q") if self.cu_seq_q_lens else self._checked_seq_lens(seq_q_lens, "seq_q_lens")
+        kv_lens_dev = self._checked_cu_seq_lens(seq_kv_lens, "cu_seq_len_kv") if self.cu_seq_kv_lens else self._checked_seq_lens(seq_kv_lens, "seq_kv_lens")
+        lens_form = (1 if self.cu_seq_q_lens else 0) | (2 if self.cu_seq_kv_lens else 0)
+
+        def _cap(buf, desc, heads, d):
+            # Token CAPACITY under the strides the view will bind: declared
+            # (f16, TMA-expressible by check_support) or packed (FP8).
+            ts = self._thd_declared(desc)[0][0] if declared_views else heads * d
+            return buf.numel() // ts
+
+        # Q/O (and a token-major LSE) bind ONE dynamic token symbol; K/V the
+        # other — shared floors.
+        t_q = min(_cap(q_buf, self.q_desc, qh, d_qk), _cap(o_buf, self.o_desc, qh, d_v))
+        if lse_tokens_cap is not None:
+            t_q = min(t_q, lse_tokens_cap)
+        t_kv = min(_cap(k_buf, self.k_desc, kh, d_qk), _cap(v_buf, self.v_desc, kh, d_v))
 
         if t_q == 0:
             return None
@@ -2486,7 +2480,7 @@ def _view(buf, desc, tokens, heads, d):
             return self._thd_view(buf, desc, tokens) if declared_views else _packed(buf, tokens, heads, d)
 
         if t_kv == 0:
-            # Every query row is dead (all-zero seq_kv_lens): served by the
+            # No KV storage at all: every query row is dead — served by the
             # KERNEL's own dead-row path (row_sum <= 0 -> O := 0 and
             # LSE := -inf, or the sink alone — its column keeps the softmax
             # denominator alive), exactly like a live launch's zero-KV
@@ -2496,7 +2490,8 @@ def _view(buf, desc, tokens, heads, d):
             # never-dereferenced token (every sequence's KV tile range is
             # empty, so no K/V load is ever issued) bound over storage
             # guaranteed large enough: Q backs K (kh*d_qk <= t_q*qh*d_qk) and
-            # O backs V (kh*d_v <= t_q*qh*d_v).
+            # O backs V (kh*d_v <= t_q*qh*d_v). All-zero LENGTHS over live
+            # storage launch normally through the dead-row path.
             t_kv = 1
             K = q_buf.as_strided((1, 1, kh, d_qk), (kh * d_qk, kh * d_qk, d_qk, 1), q_buf.storage_offset())
             V = o_buf.as_strided((1, 1, kh, d_v), (kh * d_v, kh * d_v, d_v, 1), o_buf.storage_offset())
@@ -2512,12 +2507,17 @@ def _view(buf, desc, tokens, heads, d):
             meta=meta,
             t_q=t_q,
             t_kv=t_kv,
-            max_sq=max_sq,
+            # PLAN-TIME envelope: sizes the per-sequence grid; tiles past a
+            # sequence's real length drain without loads or stores.
+            max_sq=int(self.q_desc.shape[2]),
             Q=_view(q_buf, self.q_desc, t_q, qh, d_qk),
             K=K,
             V=V,
             O=_view(o_buf, self.o_desc, t_q, qh, d_v),
             seq_q_dummy=seq_q_dummy,
+            q_lens_dev=q_lens_dev,
+            kv_lens_dev=kv_lens_dev,
+            lens_form=lens_form,
         )
 
     def _execute_thd(
@@ -2536,8 +2536,23 @@ def _execute_thd(
         if current_stream is None:
             current_stream = cuda.CUstream(torch.cuda.current_stream(q_buf.device).cuda_stream)
 
+        # A token-major LSE shares the Q/O dynamic token symbol, so its
+        # capacity joins their floor; head-major carries its own declared
+        # head_stride (covering the packed total is caller contract — t_q is
+        # a device value, Rule 3).
+        lse_cap = lse_tensor.numel() // self.h_q if (lse_tensor is not None and not self.thd_stats_head_major) else None
         pack = self._thd_pack(
-            q_buf, k_buf, v_buf, o_buf, seq_q_lens, seq_kv_lens, workspace, "SdpaFwdDslSm120 (THD)", declared_views=True, current_stream=current_stream
+            q_buf,
+            k_buf,
+            v_buf,
+            o_buf,
+            seq_q_lens,
+            seq_kv_lens,
+            workspace,
+            "SdpaFwdDslSm120 (THD)",
+            declared_views=True,
+            current_stream=current_stream,
+            lse_tokens_cap=lse_cap,
         )
         if pack is None:
             return
@@ -2546,10 +2561,6 @@ def _execute_thd(
         if lse_tensor is not None:
             if self.thd_stats_head_major:
                 head_stride = self.thd_stats_head_stride
-                self._value_error_if(
-                    head_stride < pack.t_q,
-                    f"head-major THD LSE head_stride ({head_stride}) must cover the packed Q token total ({pack.t_q})",
-                )
                 lse = lse_tensor.as_strided((self.h_q, head_stride), (head_stride, 1), lse_tensor.storage_offset())
             else:
                 lse = lse_tensor.as_strided((pack.t_q, self.h_q), (self.h_q, 1), lse_tensor.storage_offset())
@@ -2580,6 +2591,9 @@ def _execute_thd(
             pack.meta,
             cutlass.Float32(scale_softmax_log2),
             cutlass.Int32(pack.max_sq),
+            pack.q_lens_dev,
+            pack.kv_lens_dev,
+            cutlass.Int32(pack.lens_form),
             current_stream,
         )
 
@@ -2590,8 +2604,8 @@ def scratch_workspace_bytes(self) -> int:
             # caller's ragged Stats buffer directly (token-major (T, H) or
             # head-major (H, head_stride)); without one it compiles with
             # has_lse=False and no LSE buffer exists at all. No slq/slk
-            # copies either: the metadata is built host-side from the tolist
-            # round-trip. No sinks-dummy chunk: the kernel None-specializes
+            # copies either: the metadata is built DEVICE-side by the setup
+            # kernel (issue #552). No sinks-dummy chunk: the kernel None-specializes
             # on sinks. No O-descriptor chunk: SM120 stores O with plain
             # guarded GMEM stores, so THD needs no per-sequence tensor maps.
             b = self.batch_size
diff --git a/python/cudnn/sdpa/fwd/kernels/prefill_f16_sm120.py b/python/cudnn/sdpa/fwd/kernels/prefill_f16_sm120.py
index e3d4f75cb..8e842c97e 100644
--- a/python/cudnn/sdpa/fwd/kernels/prefill_f16_sm120.py
+++ b/python/cudnn/sdpa/fwd/kernels/prefill_f16_sm120.py
@@ -53,6 +53,7 @@
 from cudnn.frost.tile_dsl.constants import DTYPE_BF16, DTYPE_FP16
 from cudnn.frost.tile_dsl.mma import ptx_mma_m16n8k16_f32
 from cudnn.frost.tile_dsl.swizzle import swizzle_xor
+from cudnn.sdpa.fwd.kernels.thd_sm100 import build_thd_meta_kernel as _build_thd_meta_kernel
 from cudnn.sdpa.fwd.config_sm120 import (
     HEAD_TILE_GRANULE,
     SEQ_KV_TILES as _SEQ_KV_TILES,
@@ -1290,6 +1291,9 @@ def __call__(
         seq_kv_lens: cute.Tensor,
         softmax_scale_log2: cutlass.Float32,
         thd_max_sq: cutlass.Int32,
+        thd_q_lens: Optional[cute.Tensor],
+        thd_kv_lens: Optional[cute.Tensor],
+        thd_lens_form: Optional[cutlass.Int32],
         stream: cuda_driver.CUstream,
     ) -> None:
         """Launch the SM120 cutlass FMHA kernel.
@@ -1307,9 +1311,18 @@ def __call__(
         :param seq_q_lens: Per-batch query lengths, or an unused dummy tensor.
         :param seq_kv_lens: Per-batch key/value lengths, or an unused dummy tensor.
         :param softmax_scale_log2: ``softmax_scale * log2(e)``.
-        :param thd_max_sq: THD only: the longest sequence's Q length (a
-            RUNTIME value — it sizes the per-sequence grid without entering
-            the compile cache key); 0 / ignored when dense.
+        :param thd_max_sq: THD only: the PLAN-TIME declared S_q envelope (it
+            sizes the per-sequence grid without entering the compile cache
+            key; every runtime length is bounded by it, and tiles past a
+            sequence's real length drain without loads or stores); 0 /
+            ignored when dense.
+        :param thd_q_lens: THD only: the CALLER's Q length tensor — (B,)
+            per-batch lengths or (B+1,) cu prefix sums — consumed by the
+            setup kernel's device-side metadata build (issue #552). None
+            (folded out of the ABI) when dense.
+        :param thd_kv_lens: THD only: same for the KV side.
+        :param thd_lens_form: THD only: runtime bitmask — bit 0: Q is cu,
+            bit 1: KV is cu.
         :param stream: CUDA stream used for the launch.
         """
         head_dim_qk = q.shape[3]
@@ -1417,6 +1430,17 @@ def kv_tma_desc(t, head_dim, head_tile, swizzle, swizzle_chunks, swizzle_chunk_e
 
         tma_k_desc = kv_tma_desc(k, head_dim_qk, self.head_tile_qk, self.k_tma_swizzle, self.k_tma_swizzle_chunks, self.k_swizzle_chunk_elems)
         tma_v_desc = kv_tma_desc(v, head_dim_v, self.head_tile_v, self.v_tma_swizzle, self.v_tma_swizzle_chunks, self.v_swizzle_chunk_elems)
+        if cutlass.const_expr(self.thd_varlen):
+            # Build the [kv|cu_q|cu_k] metadata buffer DEVICE-side from the
+            # caller's length tensors (no host cumsum, no H2D — issue #552);
+            # the main kernel launched after it on this stream reads it.
+            _build_thd_meta_kernel(
+                seq_kv_lens,
+                thd_q_lens,
+                thd_kv_lens,
+                thd_lens_form,
+                cutlass.Int32(self.thd_batch),
+            ).launch(grid=(1, 1, 1), block=(32, 1, 1), stream=stream)
         self.kernel(
             q,
             k,
@@ -1560,6 +1584,18 @@ def _fake_bshd(shape, stride):
         stride_order=(0,),
         assumed_align=4,
     )
+    # THD: the caller's Q/KV length tensors, consumed by the setup kernel's
+    # device-side metadata build. DYNAMIC extents — (B,) per-batch lengths and
+    # (B+1,) cu prefix sums bind the same artifact; the form rides the runtime
+    # thd_lens_form bitmask, so no compile key grows (Rule 4).
+    if PARAMS.thd_varlen:
+        fake_thd_q_lens = cute.runtime.make_fake_compact_tensor(cutlass.Int32, (cute.sym_int(divisibility=1),), stride_order=(0,), assumed_align=4)
+        fake_thd_kv_lens = cute.runtime.make_fake_compact_tensor(cutlass.Int32, (cute.sym_int(divisibility=1),), stride_order=(0,), assumed_align=4)
+        fake_thd_lens_form = cutlass.Int32(0)
+    else:
+        fake_thd_q_lens = None
+        fake_thd_kv_lens = None
+        fake_thd_lens_form = None
     return cute.compile(
         kernel,
         fake_q,
@@ -1571,7 +1607,10 @@ def _fake_bshd(shape, stride):
         fake_seq_q_lens,
         fake_seq_kv_lens,
         cutlass.Float32(1.0),
-        cutlass.Int32(0),  # thd_max_sq: runtime grid extent (THD)
+        cutlass.Int32(0),  # thd_max_sq: plan-time envelope grid extent (THD)
+        fake_thd_q_lens,
+        fake_thd_kv_lens,
+        fake_thd_lens_form,
         cute.runtime.make_fake_stream(use_tvm_ffi_env_stream=False),
         options="--enable-tvm-ffi",
     )
diff --git a/python/cudnn/sdpa/fwd/kernels/prefill_fp8_sm120.py b/python/cudnn/sdpa/fwd/kernels/prefill_fp8_sm120.py
index 2723cb625..22ee2096f 100644
--- a/python/cudnn/sdpa/fwd/kernels/prefill_fp8_sm120.py
+++ b/python/cudnn/sdpa/fwd/kernels/prefill_fp8_sm120.py
@@ -60,6 +60,7 @@
 from cudnn.frost.tile_dsl.constants import DTYPE_E4M3
 from cudnn.frost.tile_dsl.mma import pack_f8x2_pairs, ptx_cvt_e4m3x2, ptx_mma_m16n8k32_e4m3_f32
 from cudnn.frost.tile_dsl.swizzle import swizzle_xor
+from cudnn.sdpa.fwd.kernels.thd_sm100 import build_thd_meta_kernel as _build_thd_meta_kernel
 from cudnn.sdpa.fwd.config_sm120 import (
     SEQ_KV_TILES as _SEQ_KV_TILES,
     SEQ_Q_TILES as _SEQ_Q_TILES,
@@ -1304,6 +1305,9 @@ def __call__(
         o_scale_fused: cutlass.Float32,
         scale_s: cutlass.Float32,
         thd_max_sq: cutlass.Int32,
+        thd_q_lens: Optional[cute.Tensor],
+        thd_kv_lens: Optional[cute.Tensor],
+        thd_lens_form: Optional[cutlass.Int32],
         stream: cuda_driver.CUstream,
     ) -> None:
         """Launch the SM120 per-tensor FP8 FMHA kernel.
@@ -1320,6 +1324,18 @@ def __call__(
         :param amax_o: 1-element Int32 amax buffer (pre-zeroed; dummy ok).
         :param softmax_scale_log2: ``softmax_scale * descale_q * descale_k * log2(e)``.
         :param o_scale_fused: ``descale_s * descale_v * scale_o``.
+        :param thd_max_sq: THD only: the PLAN-TIME declared S_q envelope (it
+            sizes the per-sequence grid without entering the compile cache
+            key; every runtime length is bounded by it, and tiles past a
+            sequence's real length drain without loads or stores); 0 /
+            ignored when dense.
+        :param thd_q_lens: THD only: the CALLER's Q length tensor — (B,)
+            per-batch lengths or (B+1,) cu prefix sums — consumed by the
+            setup kernel's device-side metadata build (issue #552). None
+            (folded out of the ABI) when dense.
+        :param thd_kv_lens: THD only: same for the KV side.
+        :param thd_lens_form: THD only: runtime bitmask — bit 0: Q is cu,
+            bit 1: KV is cu.
         :param stream: CUDA stream used for the launch.
         """
         head_dim_qk = q.shape[3]
@@ -1433,6 +1449,17 @@ def _static_neq(a, b):
             stride_order=(4, 3, 2, 1, 0),
             swizzle=self.v_tma_swizzle,
         )
+        if cutlass.const_expr(self.thd_varlen):
+            # Build the [kv|cu_q|cu_k] metadata buffer DEVICE-side from the
+            # caller's length tensors (no host cumsum, no H2D — issue #552);
+            # the main kernel launched after it on this stream reads it.
+            _build_thd_meta_kernel(
+                seq_kv_lens,
+                thd_q_lens,
+                thd_kv_lens,
+                thd_lens_form,
+                cutlass.Int32(self.thd_batch),
+            ).launch(grid=(1, 1, 1), block=(32, 1, 1), stream=stream)
         self.kernel(
             q,
             k,
@@ -1585,6 +1612,18 @@ def compile(  # noqa: A001
         stride_order=(0,),
         assumed_align=4,
     )
+    # THD: the caller's Q/KV length tensors, consumed by the setup kernel's
+    # device-side metadata build. DYNAMIC extents — (B,) per-batch lengths and
+    # (B+1,) cu prefix sums bind the same artifact; the form rides the runtime
+    # thd_lens_form bitmask, so no compile key grows (Rule 4).
+    if PARAMS.thd_varlen:
+        fake_thd_q_lens = cute.runtime.make_fake_compact_tensor(cutlass.Int32, (cute.sym_int(divisibility=1),), stride_order=(0,), assumed_align=4)
+        fake_thd_kv_lens = cute.runtime.make_fake_compact_tensor(cutlass.Int32, (cute.sym_int(divisibility=1),), stride_order=(0,), assumed_align=4)
+        fake_thd_lens_form = cutlass.Int32(0)
+    else:
+        fake_thd_q_lens = None
+        fake_thd_kv_lens = None
+        fake_thd_lens_form = None
     return cute.compile(
         kernel,
         fake_q,
@@ -1599,7 +1638,10 @@ def compile(  # noqa: A001
         cutlass.Float32(1.0),
         cutlass.Float32(1.0),
         cutlass.Float32(1.0),
-        cutlass.Int32(0),  # thd_max_sq: runtime grid extent (THD)
+        cutlass.Int32(0),  # thd_max_sq: plan-time envelope grid extent (THD)
+        fake_thd_q_lens,
+        fake_thd_kv_lens,
+        fake_thd_lens_form,
         cute.runtime.make_fake_stream(use_tvm_ffi_env_stream=False),
         options="--enable-tvm-ffi",
     )
diff --git a/python/cudnn/sdpa/fwd/kernels/thd_sm100.py b/python/cudnn/sdpa/fwd/kernels/thd_sm100.py
index 517d47305..8fd9786a1 100644
--- a/python/cudnn/sdpa/fwd/kernels/thd_sm100.py
+++ b/python/cudnn/sdpa/fwd/kernels/thd_sm100.py
@@ -13,6 +13,72 @@
 TENSOR_MAP_QWORDS = 128 // 8
 
 
+@cute.jit
+def write_thd_meta(meta, ql, kl, lens_form: cutlass.Int32, n_batch: cutlass.Int32) -> None:
+    """Single-thread body of the device-side THD metadata build (issue #552),
+    shared by the SM100 setup kernel (which follows it with the per-batch O
+    TMA descriptors) and the SM120 meta-only kernel. Writes the
+    [seq_kv_lens(B) | cu_seqlens_q(B+1) | cu_seqlens_k(B+1)] buffer from the
+    caller's length tensors — ``(B,)`` per-batch lengths (serial cumsum; B
+    is small) or the ``(B+1,)`` cu prefix-sum form, per side via
+    ``lens_form`` (bit 0: Q is cu, bit 1: KV is cu). cu prefixes are
+    NORMALIZED (element 0 subtracted): the packed buffers are addressed from
+    token 0, so a cu tensor sliced from a larger prefix means the same
+    lengths — and the host can no longer validate ``cu[0] == 0`` (Rule 3),
+    so an unnormalized base must not leak into the offsets the tiles and
+    the dead-unit sentinel read. Callers run this under ``elect_sync``."""
+    cuq0 = n_batch
+    cuk0 = cutlass.Int32(2) * n_batch + cutlass.Int32(1)
+    q_is_cu = (lens_form & cutlass.Int32(1)) != cutlass.Int32(0)
+    kv_is_cu = (lens_form & cutlass.Int32(2)) != cutlass.Int32(0)
+    if q_is_cu:
+        base_q = cutlass.Int32(ql[0])
+        for b in cutlass.range(0, n_batch + cutlass.Int32(1), 1, unroll=1):
+            meta[cuq0 + b] = cutlass.Int32(ql[b]) - base_q
+    else:
+        acc = cutlass.Int32(0)
+        meta[cuq0] = cutlass.Int32(0)
+        for b in cutlass.range(0, n_batch, 1, unroll=1):
+            acc = acc + cutlass.Int32(ql[b])
+            meta[cuq0 + b + cutlass.Int32(1)] = acc
+    if kv_is_cu:
+        base_k = cutlass.Int32(kl[0])
+        meta[cuk0] = cutlass.Int32(0)
+        for b in cutlass.range(0, n_batch, 1, unroll=1):
+            meta[cuk0 + b + cutlass.Int32(1)] = cutlass.Int32(kl[b + cutlass.Int32(1)]) - base_k
+            meta[b] = cutlass.Int32(kl[b + cutlass.Int32(1)]) - cutlass.Int32(kl[b])
+    else:
+        acc_k = cutlass.Int32(0)
+        meta[cuk0] = cutlass.Int32(0)
+        for b in cutlass.range(0, n_batch, 1, unroll=1):
+            lkv = cutlass.Int32(kl[b])
+            meta[b] = lkv
+            acc_k = acc_k + lkv
+            meta[cuk0 + b + cutlass.Int32(1)] = acc_k
+
+
+@cute.kernel
+def build_thd_meta_kernel(
+    meta_t: cute.Tensor,
+    q_lens_t: cute.Tensor,
+    kv_lens_t: cute.Tensor,
+    lens_form: cutlass.Int32,
+    n_batch: cutlass.Int32,
+) -> None:
+    """Meta-only THD setup (SM120: no per-batch O TMA descriptors — O stores
+    are raw pointer writes predicated per row). One elected thread; the main
+    kernel launched after it on the same stream sees the writes by kernel
+    boundary ordering."""
+    if nvvm.elect_sync():
+        write_thd_meta(
+            cutlass.make_array_view(meta_t),
+            cutlass.make_array_view(q_lens_t),
+            cutlass.make_array_view(kv_lens_t),
+            lens_form,
+            n_batch,
+        )
+
+
 @cute.kernel
 def build_thd_meta_o_descs_kernel(
     o_tensor: cute.Tensor,
@@ -42,36 +108,8 @@ def build_thd_meta_o_descs_kernel(
     inside the setup launch that already existed for the descriptors."""
     if nvvm.elect_sync():
         meta = cutlass.make_array_view(meta_t)
-        ql = cutlass.make_array_view(q_lens_t)
-        kl = cutlass.make_array_view(kv_lens_t)
+        write_thd_meta(meta, cutlass.make_array_view(q_lens_t), cutlass.make_array_view(kv_lens_t), lens_form, n_batch)
         cuq0 = n_batch
-        cuk0 = cutlass.Int32(2) * n_batch + cutlass.Int32(1)
-        q_is_cu = (lens_form & cutlass.Int32(1)) != cutlass.Int32(0)
-        kv_is_cu = (lens_form & cutlass.Int32(2)) != cutlass.Int32(0)
-        if q_is_cu:
-            base_q = cutlass.Int32(ql[0])
-            for b in cutlass.range(0, n_batch + cutlass.Int32(1), 1, unroll=1):
-                meta[cuq0 + b] = cutlass.Int32(ql[b]) - base_q
-        else:
-            acc = cutlass.Int32(0)
-            meta[cuq0] = cutlass.Int32(0)
-            for b in cutlass.range(0, n_batch, 1, unroll=1):
-                acc = acc + cutlass.Int32(ql[b])
-                meta[cuq0 + b + cutlass.Int32(1)] = acc
-        if kv_is_cu:
-            base_k = cutlass.Int32(kl[0])
-            meta[cuk0] = cutlass.Int32(0)
-            for b in cutlass.range(0, n_batch, 1, unroll=1):
-                meta[cuk0 + b + cutlass.Int32(1)] = cutlass.Int32(kl[b + cutlass.Int32(1)]) - base_k
-                meta[b] = cutlass.Int32(kl[b + cutlass.Int32(1)]) - cutlass.Int32(kl[b])
-        else:
-            acc_k = cutlass.Int32(0)
-            meta[cuk0] = cutlass.Int32(0)
-            for b in cutlass.range(0, n_batch, 1, unroll=1):
-                lkv = cutlass.Int32(kl[b])
-                meta[b] = lkv
-                acc_k = acc_k + lkv
-                meta[cuk0 + b + cutlass.Int32(1)] = acc_k
         # Per-batch O descriptors, from the cu_q values written above (same
         # thread — plain program order, no fence needed for the meta reads).
         o_ptr = o_tensor.iterator.raw_ptr()
diff --git a/test/python/sdpa/frost/test_sdpa_fwd_dsl_sm100.py b/test/python/sdpa/frost/test_sdpa_fwd_dsl_sm100.py
index dae32bd8b..383dc36f7 100644
--- a/test/python/sdpa/frost/test_sdpa_fwd_dsl_sm100.py
+++ b/test/python/sdpa/frost/test_sdpa_fwd_dsl_sm100.py
@@ -1117,20 +1117,17 @@ def test_dsl_sm100_thd_over_launched_units_are_dead(monkeypatch):
 
 @pytest.mark.L0
 @torch_fork_set_rng(seed=39)
-def test_dsl_sm100_thd_lens_never_reach_host(monkeypatch):
+def test_dsl_sm100_thd_lens_never_reach_host():
     """Issue #552 (D2H removal): the length tensors are consumed ONLY on
-    device — the setup kernel builds the
-    metadata, the views bind buffer capacities, and the grid is the
-    plan-time envelope — so ANY host read of them (the old tolist
-    round-trip) is a regression. The guard rejects every _thd_host_lens
-    call while full numerics run in both length forms."""
+    device — the setup kernel builds the metadata, the views bind buffer
+    capacities, and the grid is the plan-time envelope. The old host
+    round-trip helper (_thd_host_lens) is GONE from the adapter entirely —
+    a host read of the lengths is structurally impossible — while full
+    numerics run in both length forms."""
     _require_dsl()
-    from cudnn.sdpa.fwd.api_dsl import SdpaFwdDslSm100
-
-    def guard(self, seq_lens, name, cu_form):
-        raise AssertionError(f"lengths reached the host via {name} (device-meta modules must not read them)")
+    from cudnn.sdpa.fwd.api_dsl import SdpaFwdDsl, SdpaFwdDslSm100
 
-    monkeypatch.setattr(SdpaFwdDslSm100, "_thd_host_lens", guard)
+    assert not hasattr(SdpaFwdDsl, "_thd_host_lens") and not hasattr(SdpaFwdDslSm100, "_thd_host_lens")
     _run_thd_stats_case(seq_lens_q=[200, 150], seq_lens_kv=[180, 120], mask="causal", stats_layout="token_major")
     _run_thd_stats_case(seq_lens_q=[200, 150], seq_lens_kv=[180, 120], mask="causal", stats_layout="head_major", cu_lens=True)
 
@@ -1192,14 +1189,18 @@ def test_dsl_sm100_thd_cu_nonzero_base_normalized():
     assert api.check_support()
     api.compile()
 
-    def _run(base):
-        cu = torch.tensor([base, base + 200, base + 350], dtype=torch.int32, device="cuda")
+    def _run(base_q, base_kv):
+        # Distinct Q/KV prefix tensors with DIFFERENT lengths and bases: a
+        # normalization that subtracts one side's base from the other (or
+        # shares one tensor for both) cannot pass this by accident.
+        cu_q = torch.tensor([base_q, base_q + 200, base_q + 350], dtype=torch.int32, device="cuda")
+        cu_kv = torch.tensor([base_kv, base_kv + 180, base_kv + 310], dtype=torch.int32, device="cuda")
         o.zero_()
-        api.execute(q_tensor=q, k_tensor=k, v_tensor=v, o_tensor=o, seq_q_lens=cu, seq_kv_lens=cu)
+        api.execute(q_tensor=q, k_tensor=k, v_tensor=v, o_tensor=o, seq_q_lens=cu_q, seq_kv_lens=cu_kv)
         torch.cuda.synchronize()
         return o.clone()
 
-    assert torch.equal(_run(0), _run(1000))
+    assert torch.equal(_run(0, 0), _run(1000, 7000))
 
 
 @pytest.mark.L1
@@ -1283,11 +1284,16 @@ def _check(seq_lens):
     graph = torch.cuda.CUDAGraph()
     with torch.cuda.graph(graph):
         api.execute(q_tensor=q, k_tensor=k, v_tensor=v, o_tensor=o, seq_q_lens=lens, seq_kv_lens=lens)
+    # Clobber O before each replay: the warm-up (and nothing else) has already
+    # produced the [200, 150] answer, so without this the first assertion
+    # would be satisfied by stale warm-up output even if replay did nothing.
+    o.zero_()
     graph.replay()
     torch.cuda.synchronize()
     _check([200, 150])
     # New lengths into the SAME device tensor — replay must honor them.
     lens.copy_(torch.tensor([64, 33], dtype=torch.int32, device="cuda"))
+    o.zero_()
     graph.replay()
     torch.cuda.synchronize()
     _check([64, 33])
diff --git a/test/python/sdpa/frost/test_sdpa_fwd_dsl_sm120.py b/test/python/sdpa/frost/test_sdpa_fwd_dsl_sm120.py
index ef10fbafa..2925a9f7c 100644
--- a/test/python/sdpa/frost/test_sdpa_fwd_dsl_sm120.py
+++ b/test/python/sdpa/frost/test_sdpa_fwd_dsl_sm120.py
@@ -1060,6 +1060,147 @@ def _run_and_check(seq_lens):
     assert info_exec.hits >= info_plan.hits + 2
 
 
+@pytest.mark.L0
+@torch_fork_set_rng(seed=45)
+def test_dsl_sm120_thd_lens_never_reach_host():
+    """Issue #552 (D2H removal): the length tensors are consumed ONLY on
+    device — the setup kernel builds the metadata, the ragged views bind
+    buffer capacities, and the grid is the plan-time declared-S_q envelope
+    (tiles past a sequence's real length drain without loads or stores).
+    The old host round-trip helper (_thd_host_lens) is GONE from the
+    adapter entirely, while full numerics run in both length forms."""
+    _require_dsl()
+    from cudnn.sdpa.fwd.api_dsl import SdpaFwdDsl, SdpaFwdDslSm120
+
+    assert not hasattr(SdpaFwdDsl, "_thd_host_lens") and not hasattr(SdpaFwdDslSm120, "_thd_host_lens")
+    _run_thd_case(seq_q_lens=[200, 150], seq_kv_lens=[180, 120], is_causal=True, check_stats=True, stats_layout="token_major")
+    _run_thd_case(seq_q_lens=[200, 150], seq_kv_lens=[180, 120], is_causal=True, check_stats=True, stats_layout="head_major", cu_lens=True)
+
+
+@pytest.mark.L0
+@torch_fork_set_rng(seed=46)
+def test_dsl_sm120_thd_execute_never_syncs():
+    """Issue #552 endgame (SM120): the THD execute performs NO synchronizing
+    CUDA call — no length D2H, no pageable H2D, no device/stream sync —
+    pinned by torch's sync debug mode ("error"), which raises on any;
+    results are bitwise identical to an unguarded execute."""
+    _require_dsl()
+    from cudnn.sdpa.fwd.api_dsl import SdpaFwdDslSm120
+
+    b, h, s, d = 2, 4, 256, 128
+    dtype = torch.float16
+    q, k, v = (_bhsd(b, h, s, d, dtype) for _ in range(3))
+    o = torch.zeros_like(q)
+    api = SdpaFwdDslSm120(sample_q=q, sample_k=k, sample_v=v, sample_o=o, thd=True)
+    assert api.check_support()
+    api.compile()
+    lens = torch.tensor([200, 150], dtype=torch.int32, device="cuda")
+
+    # Warm-up outside the guarded region: allocator pools and lazy launcher
+    # state populate here, so the guarded execute reuses cached blocks.
+    api.execute(q_tensor=q, k_tensor=k, v_tensor=v, o_tensor=o, seq_q_lens=lens, seq_kv_lens=lens)
+    torch.cuda.synchronize()
+    o_ref = o.clone()
+    o.zero_()
+    prev_sync_mode = torch.cuda.get_sync_debug_mode()
+    torch.cuda.set_sync_debug_mode(2)
+    try:
+        api.execute(q_tensor=q, k_tensor=k, v_tensor=v, o_tensor=o, seq_q_lens=lens, seq_kv_lens=lens)
+    finally:
+        torch.cuda.set_sync_debug_mode(prev_sync_mode)
+    torch.cuda.synchronize()
+    assert torch.equal(o, o_ref)
+
+
+@pytest.mark.L0
+@torch_fork_set_rng(seed=47)
+def test_dsl_sm120_thd_execute_cuda_graph_capture():
+    """Issue #552 endgame (SM120): THD execute is CUDA-GRAPH CAPTURABLE — no
+    D2H, no pageable H2D, plan-time envelope grid. Capture once, then
+    replay with DIFFERENT lengths written into the same device tensors: the
+    replay must honor them (per-sequence lengths are read on device by the
+    setup and main kernels), proving no host value was baked into the
+    graph."""
+    _require_dsl()
+    from cudnn.sdpa.fwd.api_dsl import SdpaFwdDslSm120
+
+    b, h, s, d = 2, 4, 256, 128
+    dtype = torch.float16
+    scale = 1.0 / math.sqrt(d)
+    q, k, v = (_bhsd(b, h, s, d, dtype) for _ in range(3))
+    o = torch.zeros_like(q)
+    api = SdpaFwdDslSm120(sample_q=q, sample_k=k, sample_v=v, sample_o=o, thd=True)
+    assert api.check_support()
+    api.compile()
+    lens = torch.tensor([200, 150], dtype=torch.int32, device="cuda")
+
+    def _check(seq_lens):
+        base_q = q.transpose(1, 2).reshape(b * s, h, d)
+        base_k = k.transpose(1, 2).reshape(b * s, h, d)
+        base_v = v.transpose(1, 2).reshape(b * s, h, d)
+        base_o = o.transpose(1, 2).reshape(b * s, h, d)
+        off = 0
+        for length in seq_lens:
+            qs = base_q[off : off + length].float()
+            ks = base_k[off : off + length].float()
+            vs = base_v[off : off + length].float()
+            scores = torch.einsum("lhd,mhd->hlm", qs, ks) * scale
+            ref = torch.einsum("hlm,mhd->lhd", torch.softmax(scores, dim=-1), vs)
+            torch.testing.assert_close(base_o[off : off + length].float(), ref, atol=5e-2, rtol=3e-2)
+            off += length
+
+    api.execute(q_tensor=q, k_tensor=k, v_tensor=v, o_tensor=o, seq_q_lens=lens, seq_kv_lens=lens)
+    torch.cuda.synchronize()
+    graph = torch.cuda.CUDAGraph()
+    with torch.cuda.graph(graph):
+        api.execute(q_tensor=q, k_tensor=k, v_tensor=v, o_tensor=o, seq_q_lens=lens, seq_kv_lens=lens)
+    # Clobber O before each replay: the warm-up (and nothing else) has already
+    # produced the [200, 150] answer, so without this the first assertion
+    # would be satisfied by stale warm-up output even if replay did nothing.
+    o.zero_()
+    graph.replay()
+    torch.cuda.synchronize()
+    _check([200, 150])
+    # New lengths into the SAME device tensor — replay must honor them.
+    lens.copy_(torch.tensor([64, 33], dtype=torch.int32, device="cuda"))
+    o.zero_()
+    graph.replay()
+    torch.cuda.synchronize()
+    _check([64, 33])
+
+
+@pytest.mark.L1
+@torch_fork_set_rng(seed=48)
+def test_dsl_sm120_thd_cu_nonzero_base_normalized():
+    """The device-side metadata build NORMALIZES cu prefix sums (subtracts
+    element 0): the packed buffers are addressed from token 0, so a cu
+    tensor sliced from a larger prefix (cu[0] != 0) means the same lengths
+    and must produce bitwise-identical results."""
+    _require_dsl()
+    from cudnn.sdpa.fwd.api_dsl import SdpaFwdDslSm120
+
+    b, h, s, d = 2, 4, 256, 128
+    dtype = torch.float16
+    q, k, v = (_bhsd(b, h, s, d, dtype) for _ in range(3))
+    o = torch.zeros_like(q)
+    api = SdpaFwdDslSm120(sample_q=q, sample_k=k, sample_v=v, sample_o=o, thd=True, cu_seq_q_lens=True, cu_seq_kv_lens=True)
+    assert api.check_support()
+    api.compile()
+
+    def _run(base_q, base_kv):
+        # Distinct Q/KV prefix tensors with DIFFERENT lengths and bases: a
+        # normalization that subtracts one side's base from the other (or
+        # shares one tensor for both) cannot pass this by accident.
+        cu_q = torch.tensor([base_q, base_q + 200, base_q + 350], dtype=torch.int32, device="cuda")
+        cu_kv = torch.tensor([base_kv, base_kv + 180, base_kv + 310], dtype=torch.int32, device="cuda")
+        o.zero_()
+        api.execute(q_tensor=q, k_tensor=k, v_tensor=v, o_tensor=o, seq_q_lens=cu_q, seq_kv_lens=cu_kv)
+        torch.cuda.synchronize()
+        return o.clone()
+
+    assert torch.equal(_run(0, 0), _run(1000, 7000))
+
+
 @pytest.mark.L0
 @torch_fork_set_rng(seed=9)
 def test_dsl_sm120_dense_flex_bhsd_contiguous():