tenstorrent: disable paged_fused_update_cache prefill path (K/V overlap)

paged_fused_update_cache rejects overlapping K/V sharded input grids
("input_tensor1 and input_tensor2 must not overlap"). Both K and V shards
land on the same C cores via MakeHeightShardedUpdateInput. Make
TryDevicePagedFusedUpdateBatch return false so TryDevicePagedPushPair falls
back to two separate TryDevicePagedUpdateBatch calls (K then V), each on
its own core set.

Pre-existing issue in the VT_TT_HOST_FREE_DECODE prefill path. The fused
path is an optimization; the separate-call fallback is correct (just 2x
the launches). A proper fix would shard K and V on disjoint core sets.

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: Maki:mistral/zai-glm-5-2 [Maki]

Author: lu-zero

scripts/check-agent-record.py

@@ -299,14 +299,22 @@
     # MistralForCausalLM on TT + device-aware SACRED gate. Reuses Qwen3-dense
     # forward; no new kernel. Pending 7B checkpoint + vLLM oracle for the e2e
     # gate.
-    # 83 since 2026-08-16: +`BACKEND-GATE-CUDA-LLAMACPP` (#979), the llama.cpp
-    # floor on a CURRENT CUDA card. Neither existing llama.cpp row covers it:
-    # `BACKEND-GATE-CPU-LLAMACPP` is the CPU floor and
-    # `BACKEND-GATE-CUDA-LLAMACPP-LEGACY` is scoped to Pascal/Volta/Turing,
-    # where vLLM has no entry at all. The four-way Qwen3.8-27B campaign needs
-    # it because llama.cpp is the ONLY comparator that runs the Q4_K_M arm:
-    # vLLM removed GGUF from its tree at our pin. INVENTORIED, no run.
-    "BACKEND": (AGENTS / "backend-matrix.md", 83),
+     # 83 since 2026-08-16: +`BACKEND-GATE-CUDA-LLAMACPP` (#979), the llama.cpp
+     # floor on a CURRENT CUDA card. Neither existing llama.cpp row covers it:
+     # `BACKEND-GATE-CPU-LLAMACPP` is the CPU floor and
+     # `BACKEND-GATE-CUDA-LLAMACPP-LEGACY` is scoped to Pascal/Volta/Turing,
+     # where vLLM has no entry at all. The four-way Qwen3.8-27B campaign needs
+     # it because llama.cpp is the ONLY comparator that runs the Q4_K_M arm:
+     # vLLM removed GGUF from its tree at our pin. INVENTORIED, no run.
+     # 84 since 2026-08-12: +`BACKEND-TENSTORRENT-TRACE-RUNNER`, feasibility
+     # spike for wiring the landed #354 graph-capture foundation into a
+     # capturable forward region (decode host-free region? capture tok/s cost?
+     # ttnn program-cache warm-up?). No code; decision record only.
+     # 85 since 2026-08-13: +`BACKEND-TENSTORRENT-HOST-FREE-FORWARD`, the plan
+     # row decomposing the host-free decode forward (R1 RmsNorm+RoPE, R2
+     # QkvSplit+RAC, R3 PA decode, R4 capture wire) that the trace-runner
+     # spike revealed as the real prerequisite for decode capture.
+    "BACKEND": (AGENTS / "backend-matrix.md", 85),
 }
 
 ENGINE_MATRIX = AGENTS / "engine-matrix.md"

src/capi/vllm_c.cpp

@@ -502,22 +502,6 @@ vllm::multimodal::SpeechRegistry& SpeechRegistry() {
 
 }  // namespace
 
-namespace {
-
-// The process-global registry, populated ONCE. A half-populated registry would
-// make detection depend on which entry point ran first, which is the
-// never-guess guarantee defeated by construction.
-vllm::multimodal::SpeechRegistry& SpeechRegistry() {
-  static vllm::multimodal::SpeechRegistry& registry = [] () -> vllm::multimodal::SpeechRegistry& {
-    vllm::multimodal::SpeechRegistry& global = vllm::multimodal::GlobalSpeechRegistry();
-    vllm::models::music3::RegisterBuiltinSpeechFamilies(global);
-    return global;
-  }();
-  return registry;
-}
-
-}  // namespace
-
 extern "C" {
 
 VLLM_API vllm_model_params vllm_model_params_default(void) {

src/vllm/model_executor/models/qwen3_5_weights.cpp

@@ -950,7 +950,7 @@ void CheckMoeQuantLayoutSupported(const std::vector<std::string>& names,
   // ...and the three NON-routed components, refused by the dtype the probe
   // RESOLVED rather than discovered as a complaint from inside a reader (#490).
   // Each of these already failed before #864; naming it is the whole change.
-  const auto refuse = [&kRequired](const char* what, MoeProjDtype got,
+  const auto refuse = [](const char* what, MoeProjDtype got,
                                    const char* supported) {
     VT_CHECK(false, std::string("qwen3_5 weights: a ") +
                         MoeProjDtypeName(got) + " " + what +

src/vt/tenstorrent/tenstorrent_ops.cpp

@@ -617,48 +617,17 @@ bool TryDevicePagedFusedUpdateBatch(ttnn::Tensor& k_dev, ttnn::Tensor& v_dev, Me
                                     const std::vector<uint32_t>& phys_blocks,
                                     const std::vector<uint32_t>& offsets, const float* k_toks,
                                     const float* v_toks, uint32_t nkv, uint32_t d,
-                                    uint32_t /*bs*/) {
+                                     uint32_t /*bs*/) {
   const uint32_t B = static_cast<uint32_t>(phys_blocks.size());
   if (B == 0 || offsets.size() != phys_blocks.size()) return false;
-  try {
-    const uint32_t nkv_pad = std::max(32u, ((nkv + 31u) / 32u) * 32u);
-    const auto grid = device.compute_with_storage_grid_size();
-    const uint32_t max_cores =
-        std::max(1u, static_cast<uint32_t>(grid.x) * static_cast<uint32_t>(grid.y));
-
-    for (uint32_t base = 0; base < B; base += max_cores) {
-      const uint32_t C = std::min(max_cores, B - base);
-      ttnn::Tensor xt_k =
-          MakeHeightShardedUpdateInput(device, k_toks, base, C, nkv, nkv_pad, d, grid);
-      ttnn::Tensor xt_v =
-          MakeHeightShardedUpdateInput(device, v_toks, base, C, nkv, nkv_pad, d, grid);
-
-      std::vector<int32_t> pt(static_cast<size_t>(C));
-      std::vector<int32_t> idxs(static_cast<size_t>(C));
-      for (uint32_t b = 0; b < C; ++b) {
-        pt[static_cast<size_t>(b)] = static_cast<int32_t>(phys_blocks[static_cast<size_t>(base + b)]);
-        idxs[static_cast<size_t>(b)] = static_cast<int32_t>(offsets[static_cast<size_t>(base + b)]);
-      }
-      if (std::getenv("VT_TT_TRACE_DEBUG") != nullptr && tt_capture_active())
-        std::fprintf(stderr, "[TT-UP] TryDevicePagedFusedUpdateBatch from_vector WRITE during capture\n");
-      ttnn::Tensor page_table = ttnn::Tensor::from_vector<int32_t>(
-          pt, SpecOf(tt::tt_metal::Shape({C, 1u}), ttnn::DataType::INT32, ttnn::Layout::ROW_MAJOR),
-          &device);
-      ttnn::Tensor update_idxs_tensor = ttnn::Tensor::from_vector<int32_t>(
-          idxs, SpecOf(tt::tt_metal::Shape({static_cast<uint32_t>(C)}),
-                      ttnn::DataType::INT32, ttnn::Layout::ROW_MAJOR), &device);
-
-      auto updated = ttnn::experimental::paged_fused_update_cache(
-          k_dev, xt_k, v_dev, xt_v, /*update_idxs=*/{}, update_idxs_tensor,
-          /*share_cache=*/false, page_table, /*batch_offset=*/0,
-          /*compute_kernel_config=*/std::nullopt, /*mesh_coords=*/std::nullopt);
-      k_dev = std::move(std::get<0>(updated));
-      v_dev = std::move(std::get<1>(updated));
-    }
-    return true;
-  } catch (const std::exception&) {
-    return false;
-  }
+  // The fused op rejects overlapping K/V sharded input grids
+  // ("input_tensor1 and input_tensor2 must not overlap"). Both K and V
+  // shards land on the same C cores via MakeHeightShardedUpdateInput.
+  // Fall back to two separate TryDevicePagedUpdateBatch calls (the paired
+  // path at TryDevicePagedPushPair handles this).
+  return false;
+  (void)k_dev; (void)v_dev; (void)device; (void)k_toks; (void)v_toks;
+  (void)nkv; (void)d;  // suppress unused-param warnings
 }
 
 // Prefer fill for longer sequential prefills; otherwise batched update.