diff --git a/.agents/NOW.md b/.agents/NOW.md
index bf7ca04de..b23e75c6a 100644
--- a/.agents/NOW.md
+++ b/.agents/NOW.md
@@ -28,7 +28,7 @@ Working head: `row/backend-rocm-w0` (#41). Prior: benchmark checkpoint
 | Supported-models list | **LANDED**: FEATURES arch table CI-bound (33 archs) | — |
 | `/v1/videos` OpenAI shape | **MERGED** (#71): Sora `model`/`size`/`seconds` + `GET /{id}/content` | `row/SERVE-VIDEOS-REFS` PR open: reference conditioning |
 | `BACKEND-ROCM` | **(b) fix in; #140 gfx1201 hipBLAS + Gemma-4 MoE landed (contributor, authorship-preserved); W0 green 4 archs** | compile + M2 ([spec](specs/rocm-unified-memory-b.md)) |
-| TP spike #287 (PR #143) | **TP-W1 LANDED**: rank-group table + `LoadedModel` TP handle (6/6) | dispatch TP-W2 (linears + loader) |
+| TP spike #287 (PR #143) | **TP-W1 LANDED**: rank-group table + TP handle (6/6); DSR leak FIXED (unblocks #127/#154/#155) | TP-W2 (linears + loader) |
 | Release | SPIKE; 30/30 | #129 |
 | Surface coverage (`ARCH-ONE-SURFACE`) | ROW 8 + #139 IN; **ROW 6 IN REVIEW (#137): embeddings LIVE — `LlamaModel` arch, PoolingRunner in the step, `vllm_embed` v15, `/v1/embeddings`, fold gate 4/4-231, 9 kills** | Merge #137; real-ckpt oracle cosine residual |
 
diff --git a/.agents/state.md b/.agents/state.md
index 0e50fef63..a7ae2d214 100644
--- a/.agents/state.md
+++ b/.agents/state.md
@@ -43289,4 +43289,34 @@ inertness and thread_local isolation. No regression: existing `test_tp_forward`
 2/2·60 and `test_communicator` 8/8·50 unchanged (the LoadedModel field is
 additive/byte-neutral). Header-only, no CMake wiring beyond the test. Next: TP-W2
 (row/input-dim shard + rank-0 bias + per-rank Hq/Hkv + QKV kv-replication + vocab
-embed/lm_head + logits all-gather).
\ No newline at end of file
+embed/lm_head + logits all-gather).
+
+## 2026-08-08 — FIX: device-leakage regression on main (async readback becomes a Backend capability) — unblocks 3 contributor PRs
+<!-- state: 2026-08-09T00:05 -->
+
+`check-device-leakage` was RED on main (`1a021b1b`, and back through `c05cee1d`):
+`src/vllm/v1/worker/gpu/runner.cpp:107` named `vt::DeviceType::kCUDA` in the
+device-agnostic shared layer (DSR bucket `kcuda` 1 > baseline 0). It came in with
+the `QueueSupportsAsyncInputCombine` rescope during the PR #140 fix round — our
+regression, and it was failing CI on all three open contributor PRs (#127, #154,
+#155), exactly as richiejp reported in #127's "Honest gaps". Premise verified in
+the tree before acting (not taken on report).
+
+FIX mirrors the `SupportsAuxStream` precedent the guard's own message prescribes:
+move the question onto the backend. New
+`vt::Backend::SupportsAsyncSampledTokenReadback()` (backend.h, base **false**),
+overridden true by CPU (host and device memory are one allocation) and CUDA (the
+sampled id is device-mirrored). The runner now asks
+`vt::TryGetBackend(queue.device.type)` — nullptr (device not built) yields false,
+which also subsumes the old `#ifdef VLLM_CPP_CUDA` guard. SEMANTICS UNCHANGED:
+CPU async-ON, CUDA async-ON, discrete non-CUDA (ROCm gfx1201) async-OFF — the
+"!"-token hazard that motivated the original rescope stays closed, and the ROCm
+TODO now points at the override rather than at a device branch.
+
+Gates: `check-device-leakage` RED→GREEN (`kcuda=0`, DSR 32 == baseline 32);
+`test_async_llm` 8/8·347, `test_engine_core` 6/6·44, `test_llm_engine` 11/11·204
+(CPU still resolves async-ON, the regression this could have caused); clean
+-Werror CPU build. SEPARATE pre-existing main breakage recorded, not fixed here:
+`check-doc-checkpoint` is RED on `1a021b1b` itself (that policy commit touched
+70+ `.agents/` files without docs/BENCHMARKS.md + docs/FEATURES.md); it validates
+HEAD, so it self-heals as compliant commits land.
\ No newline at end of file
diff --git a/.githooks/pre-push b/.githooks/pre-push
index 2eecf8802..472bd15f8 100755
--- a/.githooks/pre-push
+++ b/.githooks/pre-push
@@ -25,7 +25,13 @@ ZERO='0000000000000000000000000000000000000000'
 # The checkers and the paths they read. Keep both lists in sync with
 # scripts/agent-preflight.sh's CHECKERS entries for the same two gates.
 CHECKERS=(check-policy.py check-prompt-contract.py check-readme-structure.py check-public-doc-tables.py)
-EXPORT_PATHS=(README.md docs scripts .agents)
+# AGENTS.md + .env.example are exported too. check-policy.py asserts AGENTS.md is
+# a non-symlink regular file (policy_contract.py:428) and resolves its Markdown
+# links against this sandbox, so omitting either made EVERY push fail here while
+# passing in the real tree ("AGENTS.md must be a non-symlink regular file", then
+# "broken Markdown link: .env.example"). Keep this list a SUPERSET of what the
+# CHECKERS read, or the hook fails closed on content that is actually fine.
+EXPORT_PATHS=(README.md AGENTS.md .env.example docs scripts .agents)
 
 shas=()
 while read -r _local_ref local_sha _remote_ref _remote_sha; do
diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md
index 145c78a44..f6f568047 100644
--- a/docs/BENCHMARKS.md
+++ b/docs/BENCHMARKS.md
@@ -7,6 +7,7 @@
 | **Binary release matrix (spiked)** | `ENG-RELEASE-BINARIES`: primary host-ABI fat-CUDA + adaptive-CPU static-core bundles; optional per-SM diagnostics; experimental literal-static musl CPU | **PENDING:** pins 10-SM fat CUDA, adaptive no-AVX2 CPU, W1-W13/W10-W12 policy, public pending states; 30 tests GREEN. No archive, staged smoke, runtime, correctness, or performance evidence | n/a |
 | **Developer agent entry point (implemented)** | `DOCS-AGENT-PROTOCOL-ENTRYPOINT`: public contribution guide + synchronized, mutation-gated pre-claim intake rule | Rebased documentation/protocol only; benchmark void | n/a |
 | **ARCH audit: ABI is text-only** | 4 capabilities (H3 video, Laguna, Kimi-Linear, DeepSeek-V4) reachable only from `examples/`, none registry-backed. No gate asks whether a CONSUMER can reach a capability. Documentation only |
+| **DSR fix: async readback capability (2026-08-08)** | **No number owed**: behavior-neutral (CPU/CUDA async-ON, discrete non-CUDA async-OFF, unchanged); moves a `kCUDA` check onto `Backend`, unblocking red CI on #127/#154/#155 |
 | **`ROAD-V1-MEM` M1+M2 (2026-08-08)** | KV auto-sizing CPU brick: `--kv-cache-memory` sizes the pool from a byte budget via the group-aware `KVBytesPerBlock` divisor (ABI v16, CPU-gated). M3 profile run dgx-gated |
 | **Record repair 2026-08-07** | `main` was red on `check-agent-record` + `check-env-doc`, blocking every PR. Dangling `kda-chunk-aot/` link and two undocumented env vars. No behaviour change |
 | **vLLM** | Qwen3.6-27B NVFP4, GB10 | ahead 4.5% at c1, **tie** at c2 to c32 | identical |
diff --git a/docs/STATUS.md b/docs/STATUS.md
index 6ccd1857d..e7e0dd6ef 100644
--- a/docs/STATUS.md
+++ b/docs/STATUS.md
@@ -431,10 +431,11 @@ tp=1). The **real TP-2 GPU run + NCCL build-verify remain HW-blocked** (no
 parallelism mode is enumerated and ranked in
 [parallelism-modes](../.agents/specs/parallelism-modes.md), noting vLLM's
 "sequence parallel" is a TP compilation pass, not an axis).
+**Shared layer names no device**: async readback is a `Backend` capability.
+
 **Tensor parallelism is scoped end-to-end at the pin** (#287,
 [tensor-parallelism-spike](../.agents/specs/tensor-parallelism-spike.md)):
-~40% of the TP surface already landed/reusable; the landed TP wiring
-dead-ends at the layer boundary (no loader shards yet); **TP-W1 LANDED** (group table); W2..W4+W7
+~40% of the TP surface landed/reusable; **TP-W1 LANDED** (group table); W2..W4+W7
 (the engine-level TP2-on-CPU token-exact gate) are CPU-completable
 NOW; only NCCL + gate-model perf wait on hardware. The DSpark speculator
 (DFlash-derived block drafter for our Qwen3 + Gemma4 families) is grounded
diff --git a/include/vt/backend.h b/include/vt/backend.h
index 21d884a82..cf4ff60ee 100644
--- a/include/vt/backend.h
+++ b/include/vt/backend.h
@@ -121,6 +121,25 @@ class Backend {
   // the DSR scan) so the model file stops naming a device at the aux-stream gate.
   virtual bool SupportsAuxStream() const { return false; }
 
+  // Whether the host may validly read the SAMPLED TOKEN ID back between steps,
+  // which is what the depth-2 async input-combine path requires
+  // (gpu/runner.cpp: QueueSupportsAsyncInputCombine). Base false — a DISCRETE
+  // non-CUDA GPU (e.g. ROCm gfx1201) is the hazard: the non-CUDA leg of
+  // sample_tokens_async Synchronizes and then host-dereferences `dev_ids`, a
+  // device Alloc that is garbage off-device (the "!"-token corruption on the lab
+  // R9700, 2026-08-07), so those queues MUST stay synchronous. Overridden true
+  // by CPU (host and device memory are the same allocation, so the read is
+  // always valid) and by CUDA (the sampled id is device-mirrored,
+  // async_device_mirror()). This is the capability the runner's
+  // `device == kCUDA` gate actually asked; it lives on Backend (src/vt, off the
+  // DSR scan) so the device-agnostic shared layer stops naming a device — the
+  // same move SupportsAuxStream made for the aux-stream gate.
+  // TODO(rocm): an INTEGRATED non-CUDA GPU reports UnifiedMemory()==true (see
+  // row/ROCM-UNIFIED-MEMORY-B), where the alias IS valid; such a backend may
+  // override this true once a HIP sampled-token mirror or a D2H copy of dev_ids
+  // lands.
+  virtual bool SupportsAsyncSampledTokenReadback() const { return false; }
+
   // Optional graph/command capture (CUDA Graphs / Metal ICB / Vulkan CB).
   virtual bool SupportsGraphCapture() const { return false; }
   virtual void BeginCapture(Queue& q);
diff --git a/src/vllm/v1/worker/gpu/runner.cpp b/src/vllm/v1/worker/gpu/runner.cpp
index 1a2524ff6..4bf531bdb 100644
--- a/src/vllm/v1/worker/gpu/runner.cpp
+++ b/src/vllm/v1/worker/gpu/runner.cpp
@@ -86,28 +86,19 @@ static bool AsyncRunnerEnvDefault() {
 }
 
 // Async input-combine reads the sampled token id back on the host between
-// steps. Where that read is valid the default-ON async path stays on:
-//   - kCPU: host and device memory are the same allocation, so the read is
-//     always valid. This path was correct, default-ON, and contract-tested
-//     before the ROCm work; it MUST stay true (else the CPU backend silently
-//     regresses to synchronous depth-1).
-//   - kCUDA: the sampled id is device-mirrored (async_device_mirror()).
-// A DISCRETE non-CUDA GPU (e.g. ROCm gfx1201) is the real hazard: the non-CUDA
-// leg of sample_tokens_async Synchronizes and then host-dereferences dev_ids,
-// which is a device Alloc — valid on CPU/UMA, garbage off-device. That is the
-// root cause of the "!" tokens on the lab R9700 (2026-08-07), not an embed
-// race. Keep those queues synchronous until a HIP sampled-token mirror or a D2H
-// copy of dev_ids lands.
-// TODO(rocm): an INTEGRATED non-CUDA GPU reports UnifiedMemory()==true (see
-// row/ROCM-UNIFIED-MEMORY-B), where the alias is valid and async would be safe;
-// route it through the backend UnifiedMemory() seam once reachable here.
+// steps. Whether that read is VALID is a backend CAPABILITY, not a device name:
+// ask the backend (vt::Backend::SupportsAsyncSampledTokenReadback, backend.h),
+// which answers true for CPU (host and device memory are one allocation) and
+// CUDA (the sampled id is device-mirrored, async_device_mirror()), and false for
+// a DISCRETE non-CUDA GPU (e.g. ROCm gfx1201) whose sample_tokens_async leg
+// host-dereferences a device Alloc — the root cause of the "!" tokens on the lab
+// R9700 (2026-08-07). An absent backend (device not built into this binary)
+// yields nullptr and therefore false, which also subsumes the old
+// #ifdef VLLM_CPP_CUDA guard. Keeping the question on the backend is what stops
+// this device-agnostic shared layer from naming a device (check-device-leakage).
 static bool QueueSupportsAsyncInputCombine(const vt::Queue& queue) {
-  if (queue.device.type == vt::DeviceType::kCPU) return true;
-#ifdef VLLM_CPP_CUDA
-  if (queue.device.type == vt::DeviceType::kCUDA) return true;
-#endif
-  (void)queue;
-  return false;
+  const vt::Backend* backend = vt::TryGetBackend(queue.device.type);
+  return backend != nullptr && backend->SupportsAsyncSampledTokenReadback();
 }
 
 // GDN step-geometry diagnostic (default OFF). When VT_GDN_DIAG_STEP_LOG=1, each
diff --git a/src/vt/cpu/cpu_backend.cpp b/src/vt/cpu/cpu_backend.cpp
index 96274a1d9..4a1591b2d 100644
--- a/src/vt/cpu/cpu_backend.cpp
+++ b/src/vt/cpu/cpu_backend.cpp
@@ -9,7 +9,13 @@ namespace vt::cpu {
 namespace {
 
 class CpuBackend final : public Backend {
+  // Host and device memory are the SAME allocation here, so reading the sampled
+  // token id back between steps is always valid: the depth-2 async input-combine
+  // path stays default-ON for CPU (it was correct and contract-tested before the
+  // ROCm work; regressing it would silently drop the CPU backend to depth-1).
  public:
+  bool SupportsAsyncSampledTokenReadback() const override { return true; }
+
   void* Alloc(size_t bytes) override {
     VT_CHECK(bytes <= SIZE_MAX - 63, "cpu alloc size overflow");
     void* p = std::aligned_alloc(64, ((bytes + 63) / 64) * 64);  // 64B-aligned, padded size
diff --git a/src/vt/cuda/cuda_backend.cu b/src/vt/cuda/cuda_backend.cu
index aeb8fd342..06a0f7b78 100644
--- a/src/vt/cuda/cuda_backend.cu
+++ b/src/vt/cuda/cuda_backend.cu
@@ -188,6 +188,9 @@ class CudaBackend final : public Backend {
   bool SupportsGraphCapture() const override { return true; }
   // S7: CUDA has a secondary stream for the MoE shared-expert overlap fork.
   bool SupportsAuxStream() const override { return true; }
+  // The sampled token id is device-mirrored (async_device_mirror()), so the
+  // between-steps host readback the depth-2 async input-combine needs is valid.
+  bool SupportsAsyncSampledTokenReadback() const override { return true; }
   void BeginCapture(Queue& q) override {
     Check(cudaStreamBeginCapture(AsStream(q), cudaStreamCaptureModeThreadLocal),
           "cudaStreamBeginCapture");