[None][refactor] BREAKING: Remove the Python backend of KVCacheManagerV2

.gitignore

@@ -59,8 +59,6 @@ tensorrt_llm/flash_mla_cpp_tllm.*.so
 tensorrt_llm/flash_mla_cpp_tllm.pyi
 /3rdparty/fmha_sm100/
 /3rdparty/nccl_extensions/
-tensorrt_llm/runtime/kv_cache_manager_v2/**/*.so
-**/*__mypyc*.so
 tensorrt_llm/scripts
 *docs/cpp_docs*
 *docs/source/_cpp_gen*

cpp/tensorrt_llm/batch_manager/kv_cache_manager_v2/AGENTS.md

@@ -468,11 +468,12 @@ Use extra review and tests for changes involving:
   `radixBlockTreeTest.cpp`, `kvCacheManagerTest.cpp`,
   `kvCacheManagerV2DigestPoolTest.cpp`, `kvCacheManagerV2HostMemTest.cpp`,
   `kvCacheManagerV2StatsTest.cpp`, and `kvCacheManagerV2TypedIndexTest.cpp`.
-- Python behavior and backend-parity tests are in
-  `tests/unittest/kv_cache_manager_v2_tests/`. During development, prefer the
-  fast path below: set `PYTHONPATH` to `tensorrt_llm/runtime/` and execute the
-  test file directly with `python`. Do not use `pytest` for this fast path; the
-  file's test runner avoids importing the full `tensorrt_llm` package.
+- Python behavior tests are in `tests/unittest/kv_cache_manager_v2_tests/`, and
+  drive this C++ implementation through the nanobind bindings. During
+  development, prefer the fast path below: set `PYTHONPATH` to
+  `tensorrt_llm/runtime/` and execute the test file directly with `python`. Do
+  not use `pytest` for this fast path; the file's test runner avoids importing
+  the full `tensorrt_llm` package.
 
   ```bash
   REPO_ROOT="$(git rev-parse --show-toplevel)"

cpp/tensorrt_llm/batch_manager/kv_cache_manager_v2/exceptions.h

@@ -111,8 +111,8 @@ class LogicError : public std::logic_error
 };
 
 // Mirrors a Python `assert` failure: the binding layer translates this to a
-// Python AssertionError so shared tests observe the same exception type as the
-// pure-Python backend.
+// Python AssertionError, so a configuration mistake surfaces to Python callers
+// as the exception type they would expect.
 class AssertionError : public std::logic_error
 {
 public:

cpp/tensorrt_llm/nanobind/batch_manager/kvCacheManagerV2.cpp

@@ -21,6 +21,7 @@
 #include "kv_cache_manager_v2/coldPageCodec.h"
 #include "kv_cache_manager_v2/common.h"
 #include "kv_cache_manager_v2/config.h"
+#include "kv_cache_manager_v2/cudaVirtMem.h"
 #include "kv_cache_manager_v2/eventManager.h"
 #include "kv_cache_manager_v2/exceptions.h"
 #include "kv_cache_manager_v2/introspection.h"
@@ -2227,6 +2228,20 @@ void KvCacheManagerV2Bindings::initBindings(nb::module_& m)
         nb::arg("quota"), nb::arg("slot_size_lists"), nb::arg("ratio_list"), nb::arg("granularity"),
         nb::arg("min_slots"), nb::call_guard<nb::gil_scoped_release>());
 
+    // CUDA virtual-memory primitives, reached through _introspection because they carry no
+    // stability promise: the native disaggregated bounce buffer reserves one contiguous fabric
+    // region with them and maps physical chunks into it up front.
+    nb::class_<kv::PooledPhysMemAllocator>(mIntrospection, "PooledPhysMemAllocator")
+        .def(nb::init<size_t>(), nb::arg("phys_mem_size"))
+        .def_prop_ro("device_id", &kv::PooledPhysMemAllocator::deviceId);
+    nb::class_<kv::VirtMem>(mIntrospection, "VirtMem")
+        // keep_alive<1, 3>: VirtMem holds PooledPhysMemAllocator by reference, so the allocator
+        // must outlive it. Argument 3 is the allocator (1 is self, 2 is vm_size).
+        .def(nb::init<size_t, kv::PooledPhysMemAllocator&, size_t>(), nb::arg("vm_size"), nb::arg("phys_mem_allocator"),
+            nb::arg("init_num_phys_mem") = 0, nb::keep_alive<1, 3>())
+        .def("destroy", &kv::VirtMem::destroy)
+        .def_prop_ro("address", &kv::VirtMem::address);
+
     // ---- Cold-page codec --------------------------------------------------
     nb::class_<kv::IKvCacheColdPageCodec>(m, "IKvCacheColdPageCodec");
     m.def("create_default_kv_cache_cold_page_codec", &kv::createDefaultKvCacheColdPageCodec,

cpp/tests/unit_tests/batch_manager/kvCacheManagerV2StatsTest.cpp

@@ -39,6 +39,7 @@ namespace
 
 using namespace tensorrt_llm::batch_manager::kv_cache_manager_v2;
 using tensorrt_llm::batch_manager::kv_cache_manager_v2::test::makeConfig;
+using tensorrt_llm::batch_manager::kv_cache_manager_v2::test::makeHybridTieredConfig;
 using tensorrt_llm::batch_manager::kv_cache_manager_v2::test::makeTieredConfig;
 
 TEST(KvCacheManagerV2StatsTest, StatsDeltaArithmetic)
@@ -454,4 +455,167 @@ TEST(KvCacheManagerV2StatsTest, DisabledStatsSuppressSuspendResumeCounters)
     EXPECT_EQ(cudaStreamDestroy(stream), cudaSuccess);
 }
 
+// Hybrid attention + SSM page-movement statistics.
+//
+// Iteration statistics are keyed by life cycle and must report recurrent (SSM) page
+// movement alongside attention movement, otherwise KDA recurrent-state offload, onboard
+// and drop are invisible to callers. The global cache-hit counters are the deliberate
+// exception: they stay attention-only.
+//
+// The tiers below give the attention life cycle 4 GPU and 2 host slots of 1 MiB, and the
+// SSM life cycle 1 GPU and 1 host slot of 2 MiB, so a second sequence evicts the first and
+// a second eviction round overflows the host pool.
+namespace
+{
+constexpr int kHybridBlocks = 3;
+constexpr CacheLevel kHybridHostLevel{1};
+
+int64_t slotBytesFor(StorageManager const& storage, CacheLevel level, LifeCycleId lifeCycle)
+{
+    int64_t bytes = 0;
+    for (size_t const size : storage.slotSize(storage.getPoolGroupIndex(level, lifeCycle)))
+    {
+        bytes += static_cast<int64_t>(size);
+    }
+    return bytes;
+}
+
+std::vector<TokenIdExt> makeTokens(KvCacheManager const& manager, int firstToken)
+{
+    std::vector<TokenIdExt> tokens;
+    for (int offset = 0; offset < kHybridBlocks * manager.tokensPerBlock(); ++offset)
+    {
+        tokens.emplace_back(TokenId{firstToken + offset});
+    }
+    return tokens;
+}
+
+// Fill a sequence, park it, then start a second one that needs the same slots. Returns the
+// still-open second sequence so the caller can close it to trigger the onboard.
+std::pair<std::shared_ptr<KvCache>, std::shared_ptr<KvCache>> evictFirstSequence(
+    KvCacheManager& manager, cudaStream_t stream, int firstToken)
+{
+    auto const tokens = makeTokens(manager, firstToken);
+    auto first = manager.createKvCache();
+    EXPECT_TRUE(first->resume(reinterpret_cast<CUstream>(stream)));
+    EXPECT_TRUE(first->resize(static_cast<int>(tokens.size())));
+    first->commit(toSpan(tokens));
+    first->suspend();
+
+    auto second = manager.createKvCache();
+    EXPECT_TRUE(second->resume(reinterpret_cast<CUstream>(stream)));
+    EXPECT_TRUE(second->resize(static_cast<int>(tokens.size())));
+    return {std::move(first), std::move(second)};
+}
+} // namespace
+
+TEST(KvCacheManagerV2StatsTest, OffloadAndOnboardAreRecordedForAttentionAndSsmLifeCycles)
+{
+    ASSERT_EQ(cudaSetDevice(0), cudaSuccess);
+    cudaStream_t stream{};
+    ASSERT_EQ(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking), cudaSuccess);
+    auto manager = std::make_shared<KvCacheManager>(makeHybridTieredConfig());
+    auto& storage = manager->storage();
+    LifeCycleId const attention{0};
+    LifeCycleId const ssm{1};
+    ASSERT_TRUE(std::holds_alternative<AttnLifeCycle>(manager->lifeCycles().getLifeCycle(attention)));
+    ASSERT_FALSE(std::holds_alternative<AttnLifeCycle>(manager->lifeCycles().getLifeCycle(ssm)));
+
+    manager->getAndResetIterationStats();
+    auto [first, second] = evictFirstSequence(*manager, stream, 0);
+
+    auto const offload = manager->getAndResetIterationStats();
+    ASSERT_EQ(offload.size(), 2) << "both life cycles must report offload";
+    for (LifeCycleId const lifeCycle : {attention, ssm})
+    {
+        auto const& stats = offload.at(lifeCycle);
+        EXPECT_GT(stats.iterOffloadBlocks, 0) << "life cycle " << lifeCycle.value();
+        EXPECT_EQ(stats.iterOffloadBytes, stats.iterOffloadBlocks * slotBytesFor(storage, kHotLevel, lifeCycle));
+    }
+
+    second->close();
+    ASSERT_TRUE(first->resume(reinterpret_cast<CUstream>(stream)));
+
+    auto const onboard = manager->getAndResetIterationStats();
+    ASSERT_EQ(onboard.size(), 2) << "both life cycles must report onboard";
+    for (LifeCycleId const lifeCycle : {attention, ssm})
+    {
+        auto const& stats = onboard.at(lifeCycle);
+        EXPECT_GT(stats.iterOnboardBlocks, 0) << "life cycle " << lifeCycle.value();
+        EXPECT_EQ(stats.iterOnboardBytes, stats.iterOnboardBlocks * slotBytesFor(storage, kHotLevel, lifeCycle));
+    }
+
+    first->close();
+    EXPECT_EQ(cudaStreamDestroy(stream), cudaSuccess);
+}
+
+TEST(KvCacheManagerV2StatsTest, SsmOnboardLeavesGlobalAllocCountersToAttention)
+{
+    ASSERT_EQ(cudaSetDevice(0), cudaSuccess);
+    cudaStream_t stream{};
+    ASSERT_EQ(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking), cudaSuccess);
+    auto manager = std::make_shared<KvCacheManager>(makeHybridTieredConfig());
+    LifeCycleId const attention{0};
+    LifeCycleId const ssm{1};
+
+    manager->getAndResetIterationStats();
+    auto [first, second] = evictFirstSequence(*manager, stream, 0);
+    manager->getAndResetIterationStats();
+
+    second->close();
+    auto const allocTotalBefore = manager->getCommittedStats().allocTotalBlocks;
+    auto const allocNewBefore = manager->getCommittedStats().allocNewBlocks;
+    ASSERT_TRUE(first->resume(reinterpret_cast<CUstream>(stream)));
+    auto const allocTotalDelta = manager->getCommittedStats().allocTotalBlocks - allocTotalBefore;
+    auto const allocNewDelta = manager->getCommittedStats().allocNewBlocks - allocNewBefore;
+
+    auto const onboard = manager->getAndResetIterationStats();
+    ASSERT_EQ(onboard.size(), 2);
+    auto const attentionOnboard = onboard.at(attention).iterAllocTotalBlocks;
+    auto const ssmOnboard = onboard.at(ssm).iterAllocTotalBlocks;
+    // Both life cycles onboard, so a global delta equal to the attention share alone is
+    // only possible if the SSM share was excluded.
+    ASSERT_GT(attentionOnboard, 0);
+    ASSERT_GT(ssmOnboard, 0);
+    EXPECT_EQ(allocTotalDelta, attentionOnboard);
+    EXPECT_EQ(allocNewDelta, attentionOnboard);
+
+    first->close();
+    EXPECT_EQ(cudaStreamDestroy(stream), cudaSuccess);
+}
+
+TEST(KvCacheManagerV2StatsTest, HostDropIsRecordedForAttentionAndSsmLifeCycles)
+{
+    ASSERT_EQ(cudaSetDevice(0), cudaSuccess);
+    cudaStream_t stream{};
+    ASSERT_EQ(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking), cudaSuccess);
+    auto manager = std::make_shared<KvCacheManager>(makeHybridTieredConfig());
+    auto& storage = manager->storage();
+    LifeCycleId const attention{0};
+    LifeCycleId const ssm{1};
+
+    // First round fills the host pools.
+    auto [first, second] = evictFirstSequence(*manager, stream, 0);
+    second->close();
+    first->close();
+
+    // Second round uses disjoint tokens, so nothing is reused and the host pools overflow.
+    manager->getAndResetIterationStats();
+    auto [third, fourth] = evictFirstSequence(*manager, stream, 1000);
+
+    auto const dropped = manager->getAndResetIterationStats();
+    ASSERT_EQ(dropped.size(), 2) << "both life cycles must report host drops";
+    for (LifeCycleId const lifeCycle : {attention, ssm})
+    {
+        auto const& stats = dropped.at(lifeCycle);
+        EXPECT_GT(stats.iterHostDroppedBlocks, 0) << "life cycle " << lifeCycle.value();
+        EXPECT_EQ(stats.iterHostDroppedBytes,
+            stats.iterHostDroppedBlocks * slotBytesFor(storage, kHybridHostLevel, lifeCycle));
+    }
+
+    fourth->close();
+    third->close();
+    EXPECT_EQ(cudaStreamDestroy(stream), cudaSuccess);
+}
+
 } // namespace

cpp/tests/unit_tests/batch_manager/kvCacheManagerV2TestUtils.h

@@ -51,4 +51,35 @@ inline KVCacheManagerConfig makeTieredConfig()
     return config;
 }
 
+//! Attention and SSM life cycles side by side, over a GPU and a host tier.
+//!
+//! Life cycles are registered in layer order, so the attention layer is LifeCycleId{0} and
+//! the SSM layer is LifeCycleId{1}. The buffer sizes differ so per-life-cycle byte counters
+//! identify which life cycle they came from.
+//!
+//! The quotas give attention 4 GPU and 2 host slots of 1 MiB, and SSM 1 GPU and 1 host slot
+//! of 2 MiB. A three-block sequence therefore fits on the GPU but a second sequence evicts
+//! it, and a second eviction round overflows the host pools.
+inline KVCacheManagerConfig makeHybridTieredConfig()
+{
+    KVCacheManagerConfig config;
+    config.tokensPerBlock = 4;
+    config.cacheTiers.emplace_back(GpuCacheTierConfig{6UL << 20});
+    config.cacheTiers.emplace_back(HostCacheTierConfig{4UL << 20});
+
+    AttentionLayerConfig attention;
+    attention.layerId = 0;
+    attention.buffers.push_back(BufferConfig{"key", 1UL << 20, std::nullopt});
+    config.layers.emplace_back(std::move(attention));
+
+    SsmLayerConfig ssm;
+    ssm.layerId = 1;
+    ssm.buffers.push_back(BufferConfig{"ssm_state", 2UL << 20, std::nullopt});
+    config.layers.emplace_back(std::move(ssm));
+
+    // KVCacheManagerConfig::validate() rejects an SSM layer without this.
+    config.commitMinSnapshot = true;
+    return config;
+}
+
 } // namespace tensorrt_llm::batch_manager::kv_cache_manager_v2::test

docs/source/features/kvcache.md

@@ -278,7 +278,14 @@ Events are buffered per rank, gathered onto rank 0 under attention data parallel
 pulled per iteration through `LLM.get_kv_cache_events()` / `LLM.get_kv_cache_events_async()`,
 or over the `/kv_cache_events` endpoint of `trtllm-serve`.
 
-#### Streaming path (prototype)
+#### Streaming path (unsupported)
+
+```{note}
+The streaming path has no implementation: `kv_cache_config.kv_events_config` is rejected
+at startup. Use the buffered path via `kv_cache_config.event_buffer_max_size` instead. The
+wire format and endpoint convention below describe the contract a future native event sink
+must satisfy.
+```
 
 Configured with ```kv_cache_config.kv_events_config```. Each rank encodes its own events and
 publishes them directly over a ZeroMQ `PUB` socket from a background thread, so there is no
@@ -297,12 +304,10 @@ kv_cache_config = KvCacheConfig(
 )
 ```
 
-**Constraints.** The streaming path requires KV cache manager V2 running on its Python
-backend (`TLLM_KV_CACHE_MANAGER_V2_BACKEND=python`); the default `cpp` backend cannot
-consume the Python event sink and raises an error naming this variable. Pipeline
-parallelism and context parallelism are rejected. Events are not published for draft
-models or during KV-cache-size estimation. When streaming is enabled the buffered pull API
-returns an empty list rather than raising.
+**Constraints.** Enabling the streaming path raises at startup. A Python event sink cannot
+serve it, because the KV cache manager V2 radix tree calls its sink natively rather than
+through Python; re-enabling it needs a native sink. Pipeline parallelism and context
+parallelism are rejected independently.
 
 **Endpoint convention.** Every attention-DP rank binds `base_port + rank` using its
 **global** rank, so `N` ranks occupy `[base_port, base_port + N - 1]` cluster-wide and

docs/source/installation/build-from-source.md

@@ -104,7 +104,6 @@ With `--build_root <dir>` set, the following default under `<dir>` instead of th
 | Build virtual environment | `venv-<python version>` | run inside an activated venv, or `--no-venv` |
 | Wheel staging tree and `*.egg-info` | `wheel-staging` | `TRTLLM_WHEEL_STAGING_DIR` |
 | ccache directory (with `--use_ccache`) | `ccache` | `CCACHE_DIR` |
-| Intermediate extension-module objects | `kv_cache_manager_v2-temp` | — |
 
 Conan's `cpp/CMakeUserPresets.json` convenience file is also skipped in this mode, since it would reference the (possibly ephemeral) out-of-tree build directory.
 

examples/kv_cache_compression/nvfp4_cold_page.md

@@ -71,9 +71,8 @@ and KV-cache block reuse remains supported because Page and token identity are
 unchanged. One-model MTP-EAGLE and EAGLE3 are supported; HELIX context
 parallelism is not currently supported.
 
-Set `kv_cache_config.use_kv_cache_manager_v2: true` explicitly, and do not set
-`TLLM_KV_CACHE_MANAGER_V2_BACKEND=python`. A nonzero Host or Disk cache is also
-required for Pages to cross a compression boundary.
+Set `kv_cache_config.use_kv_cache_manager_v2: true` explicitly. A nonzero Host
+or Disk cache is also required for Pages to cross a compression boundary.
 
 On Linux 6.11 through 6.13, mixed models that need both NVFP4 Attention
 lifecycles and lossless SSM/GDN fallback lifecycles are not supported. See the

pyproject.toml

@@ -924,10 +924,6 @@ disallow_incomplete_defs = false
 disallow_untyped_defs = false
 warn_return_any = false
 
-[[tool.mypy.overrides]]
-module = ["tensorrt_llm.runtime.kv_cache_manager_v2.*"]
-disallow_any_generics = false
-
 [[tool.mypy.overrides]]
 module = ["tensorrt_llm.bindings.*"]
 ignore_errors = true