[None][feat] cache transceiver nixl bounce buffer

[chuangz0]

Dev Engineer Review

• Adds NIXL bounce-buffer v2 for disaggregated KV-cache transfers.
• Packs descriptors into pre-registered arena regions and uses one RDMA write per chunk.
• Adds planning, buddy allocation, credit scheduling, CUDA gather/scatter, execution pools, wire codecs, control channels, and NVTX tracing.
• Adds ZMQ and NIXL notification channels.
• Adds capability handshakes, timeout handling, cancellation, peer recovery, and safe shutdown.
• Routes eligible transfers through Bounce v2 and falls back to standard NIXL when initialization fails.
• Extends AgentDesc serialization with bounce handshake metadata.
• Adds Python and C++ observability for bounce activation and submission counts.
• Review focus includes API consistency, CUDA resource ownership, thread safety, RDMA completion handling, shutdown ordering, fallback behavior, and configuration scope.
• CI failures require follow-up before merge.

QA Engineer Review

• Adds unit, CUDA, ZMQ, NIXL transport, RDMA, failure-recovery, production-agent, and Python integration tests.
• Covers arenas, configuration parsing, message codecs, transfer planning, allocation, scheduling, execution pools, gather/scatter, control channels, transport recovery, fairness, backpressure, handshakes, and data integrity.
• Adds production-path tests for one-way, bidirectional, concurrent, and multi-sender transfers.
• Adds test_python_nixl_cache_transceiver_uses_cpp_bounce with representative V1/V2, MLA/MHA, TP/PP, and DSA configurations.
• Adds the Python integration test to tests/integration/test_lists/test-db/l0_h100.yml.
• Coverage mapping for the C++ tests is not provided.
• Verdict: needs follow-up.

Description

NIXL Bounce-Buffer v2 — Design

Describes the current (as-built) implementation. Code lives in
cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/, integrated through
NixlTransferAgent (transferAgent.cpp) and gated at runtime by TRTLLM_NIXL_BOUNCE_ENABLE.

1. Motivation

A single disaggregated-KV submitTransferRequests often carries thousands to tens of thousands of
small (~4 KiB) scattered descriptors. Submitting them to NIXL/NIC one-by-one is dominated by
per-descriptor overhead, leaving the link far below line rate.

The data-path idea: the sender gathers the scattered small descriptors into one pre-registered
buffer → does a single RDMA write to a peer buffer → the receiver scatters back to the final
destinations:

src×N  ──gather──▶  region(sender)  ──RDMA write──▶  region(receiver)  ──scatter──▶  dst×N

The hard part is the control plane: how buffers are allocated across senders, flow-controlled,
recycled, and how errors are handled. Requirements:

• R1 Stream gracefully even when a single request exceeds the total buffer on both sides
  (bounded buffers carry an unbounded transfer).
• R2 Pipeline gather / RDMA write / scatter across GPU and NIC to saturate the NIC.
• R3 One receiver written concurrently by many senders: fair, starvation-free, no cross-talk.
• R4 Thread-safe under multi-threaded submit.
• R5 Any error / peer loss / teardown resolves wait() (SUCCESS/FAILURE) — never hangs.
• R6/R7 Modular, testable, readable; control and data planes are pluggable.
• R8 Pack many small concurrent requests (per-request bytes < buffer size) tightly into the shared
  buffer — no request reserves a whole buffer, which would waste memory and cap concurrency.

R1 and R8 pull in opposite directions — carry one over-sized request and pack many small ones into
the same fixed buffer — and together motivate the variable-size region arena (§2).

2. Overview

• One shared BounceArena (a device buffer registered once) is sliced into variable-size
  regions: each chunk gets a region of exactly its byte size. Many small requests pack tightly
  (high concurrency, no waste); a request larger than the arena streams chunk-by-chunk with recycling.
• Credit flow control: the receiver owns the buffer; a sender must obtain a credit
  (exclusive write permission to a region) before writing. Credits are granted incrementally per chunk
  (WANT/GRANT) and recycled on scatter completion (ACK) — bounded end to end.
• Eager gather (default on): submit() launches a chunk's gather before the GRANT arrives,
  overlapping the WANT→GRANT control round-trip with the gather kernel. Eager (credit-less) staging is
  capped at half the arena, so on a bidirectional deployment both sides can always still grant
  incoming regions (no mutual eager-starvation); credit-backed allocations are not capped.
• Pluggable control channel: ZMQ/TCP by default; optionally the control messages ride NIXL
  notifications (UCX active messages on the RDMA fabric, TRTLLM_NIXL_BOUNCE_USE_NIXL_NOTIFICATIONS),
  dropping a control hop from tens of microseconds to a few. Peers must use the same control kind —
  enforced by the capability handshake (§5).
• Single-IO-thread reactor: one IO thread per agent owns the credit / request state (nearly
  lock-free), plus M scatter workers.
• No notifMsg, no extra flush on the data plane: data-landed is decided purely by the sender
  polling poll==SUCCESS (the NIXL UCX backend already appends ucp_ep_flush_nbx per transfer, so
  SUCCESS ⇒ the data is visible at the remote target).
• Decoupled execution resources: the stream/event/scratch needed to run one gather/scatter kernel
  comes from a small ExecPool (copyStreamCount contexts), borrowed/returned per kernel — separate
  from a region's long lifetime (which lasts until ACK).

flowchart LR
    subgraph SND["Sender agent (IO thread)"]
      REQ["Request{ numChunks,<br/>nextPost, acked }"]
      OUT["shared BounceArena + ExecPool<br/>acquireLocal(bytes) → gather staging<br/>(eager: before GRANT, ≤ ½ arena)"]
    end
    subgraph RCV["Receiver agent (IO thread)"]
      SCH["CreditScheduler (single allocator)<br/>BuddyAllocator(arena)<br/>flows{ pending, held } + localHeld<br/>round-robin ring + drain mode"]
      IN["shared BounceArena<br/>variable-size regions (RDMA targets)"]
      SW["scatter workers ×M (borrow ExecPool)"]
    end
    REQ -- "① WANT(rid, chunk_bytes[])" --> SCH
    SCH -- "② GRANT(regionHandle, addr, len)" --> REQ
    REQ -- "③ gather→RDMA write→DATA(regionHandle, scatter runs)" --> IN
    IN --> SW
    SW -- "④ scatter_done" --> SCH
    SCH -- "⑤ ACK(regionHandle)" --> REQ
    SCH -. "owns/allocates arena regions" .-> IN
    REQ -. "③ occupy / ⑤ release" .-> OUT

Loading

One agent can be both sender and receiver, and both roles share the same arena (most disagg agents
only send or only receive, so one buffer halves memory; a dual-role agent still shares a single arena,
deadlock-free — the eager cap above is what makes the bidirectional case safe).

3. Terminology

Term	Meaning
region	A variable-size slice of BounceArena; its handle = byte offset within the arena (addr = baseAddr + offset).
chunk	A unit produced by BounceTransferPlan by bin-packing scattered (src,dst) descriptors; ≤ maxChunkSizeBytes, moved by one RDMA write.
credit	Exclusive write permission to a region, handed to exactly one sender as GRANT{regionHandle, addr, len}.
flow	One independent request stream, key = "peer\x1f rid" (concurrent requests from the same peer are distinct flows).
in-flight cap W	Per-flow cap on in-flight regions (= pipeline depth) — maxInflightChunksPerRequest.
arena	The single shared, registered-once device buffer (BounceArena).
ExecCtx	{stream, event, scratch, hostPinned} needed to run one gather/scatter, borrowed/returned from ExecPool.
gather / scatter	Batched D2D copy between many small descriptors and a contiguous region (GatherScatterKernel).
scatter run	A coalesced DATA plan entry (BounceScatterRun): adjacent descriptors merged when contiguous or uniformly strided.
handshake	Per-agent bounce capability blob in AgentDesc (§5); peers engage bounce only when compatible.

4. Modules

Pure logic (no GPU / threads / IO — unit-testable):

Module	Responsibility
BounceConfig	env → POD config snapshot (fromEnv); byte-valued vars accept binary K/M/G suffixes ("256MB", "1gb", "512KiB"); garbage values fall back to defaults instead of parsing to 0.
BounceTransferPlan	Bin-pack (src,dst) descriptors → chunks (each ≤ maxChunkSizeBytes, 32 B aligned, zero-length skipped); compute each descriptor's in-region offset and packedBytes; coalesce the scatter view into scatterRuns (contiguous or uniform-stride merge) so the DATA message shrinks from per-desc entries (hundreds of KB) to a handful of runs.
BounceMessage	Control-plane wire codec (fixed header + fixed-size entries); little-endian; includes the handshake codec (encodeHandshake/decodeHandshake) and the cancel codec (encodeCancel/isCancelWant).
BuddyAllocator	Pure power-of-two buddy allocator over byte offsets; alloc(bytes)/free(offset), coalesces buddies, no external fragmentation, internal ≤ 2×.
CreditScheduler	Receiver-side credit allocation + round-robin fairness + drain-mode anti-starvation (§5); embeds a BuddyAllocator over the arena; also serves local-sender acquireLocal(bytes) with the eager half-arena cap. Owned by the IO thread; the only cross-thread caller is acquireLocal() from submit() app threads (eager gather staging).

Device / IO / integration:

Module	Responsibility
BounceArena	An arenaSizeBytes device buffer (MNNVL via common::FabricMemory, else cudaMalloc), registered once; base()/baseAddr()/at(offset).
ExecPool	copyStreamCount ExecCtxs; tryAcquire() (non-blocking, nullptr when full) / release(), thread-safe.
GatherScatterKernel	Batched memcpy. Default: custom kernel (uint4-vectorized; byte path for misaligned) reading its `[srcs
TransferEngine (abstract)	The only data-plane ops: registerRegion / postWrite / poll / release. NixlTransferEngine (production, wraps one nixlAgent) is the sole implementation; transport tests run over real NIXL loopback.
ControlChannel (abstract)	addPeer (returns success) / removePeer / sendTo / recv / localEndpoint. Two implementations: ZmqControlChannel (default): a ROUTER for receive + one DEALER per peer for send; sendTo is non-blocking (drops on full kSendHwm, never blocks the IO thread). NixlNotifControlChannel: control messages as NIXL notifications on the RDMA fabric (no TCP sockets; "endpoint" is serialized NIXL metadata).
BounceTransport	Thin reactor: 1 IO thread + M scatter workers; holds BounceContext + BounceSender + BounceReceiver, routes control messages to the right role, drains both roles each tick; submit() / addPeer() / forgetPeer() / shutdown(); owns the capability handshake (localHandshakeBlob / registerPeerHandshake / hasPeerHandshake).
BounceContext	Dependencies shared by both roles and owned by the single IO thread: injected channel/engine/arena/exec, the single CreditScheduler (one arena serves both directions), sendGrants().
BounceSender	[S] role: submit→WANT (+ eager gather), GRANT→attach credits / gather+write, ACK→resolve; holds the request table + send-side deferred-cleanup state (mOrphanLocal / mPendingCancel).
BounceReceiver	[R] role: WANT→grant region, DATA→scatter, reply ACK; holds scatter workers + job/done queues + mScattering (orphaned flag for in-flight scatters).
BounceNvtx	NVTX ranges over the transfer pipeline (submit / gather / write / scatter / control hops) for nsys analysis.
NixlTransferAgent integration	maybeInitBounce (build arena+exec+transport, register arena; any init failure warns and falls back to the standard NIXL path — never fails agent construction), shouldUseBounce (routing decision, gated on the peer handshake), AgentDesc carries the local handshake blob, invalidateRemoteAgent→forgetPeer. Built only when NIXL + zmq are available (TLLM_BOUNCE_V2); decoupled from ENABLE_UCX.

5. Control Plane: Credit Flow Control + Fair Scheduling (R1/R3)

Messages (all over ControlChannel — zmq by default, NIXL notifications opt-in):

Message	Direction	Payload
WANT	sender → receiver	per-chunk byte sizes (empty = cancel) + the sender's own control endpoint.
GRANT	receiver → sender	credits {addr, len, devId(receiver), regionHandle}.
DATA	sender → receiver	regionHandle + coalesced scatter runs; sent only after poll==SUCCESS (the scatter trigger).
ACK	receiver → sender	regionHandle; scatter done, region recyclable.

Capability handshake (bootstrap + compatibility, key). Each agent's AgentDesc carries a bounce
handshake blob: {wireVersion, controlKind (ZMQ | NIXL_NOTIF), arenaUsableCapacityBytes, maxChunkSizeBytes, endpoint}. loadRemoteAgent → registerPeerHandshake validates it — version,
control kind, and maxChunkSizeBytes must match the local config, and the peer's endpoint must be
registrable — and only then marks the peer bounce-capable. shouldUseBounce requires
hasPeerHandshake(peer), so a peer with bounce disabled, a different control transport, or mismatched
chunking silently stays on the standard NIXL path (no WANT ever stalls to requestTimeoutMs).
An agent that cannot produce a usable local endpoint advertises no handshake (never breaks
metadata exchange).

Reverse-path bootstrap. Bounce needs a two-way control channel (sender sends WANT/DATA, receiver
replies GRANT/ACK), but the disagg metadata exchange is one-directional — the KV sender
loadRemoteAgents the receiver, the receiver never loads the sender. So WANT also carries the
sender's control endpoint, and the receiver addPeer(sender)s in onWant to bootstrap the
reverse path. A malformed endpoint in a WANT is rejected on the reactor thread (warn, no grant, no
exception escapes); a cancel is still honored even when endpoint registration fails, so it can reclaim
flow state left by an earlier valid WANT. Cancel/abort uses an empty WANT (still carrying the
endpoint) — no separate handshake / RETURN message.

Endpoints are routable IPs (multi-node, ZMQ). ZmqControlChannel must not bind 127.0.0.1
(unreachable cross-node). maybeInitBounce resolves the local routable IP via the shared
common::getLocalIp(getEnvNixlInterface(), rank) (TRTLLM_NIXL_INTERFACE picks the NIC, else
auto-detect by egress route / hostname — identical to UCX/NIXL addressing) and binds tcp://<ip>:*;
localEndpoint() reads the actual tcp://<ip>:<port> from zmq last_endpoint and advertises it via
the handshake / WANT. Unit tests that construct ZmqControlChannel directly keep the
tcp://127.0.0.1:* default. IPv6: zmq disables IPv6 by default, so an IPv6 bind address is
bracketed (tcp://[<ip>]:*) and the ROUTER sets ZMQ_IPV6; the DEALER (addPeer) sets ZMQ_IPV6
unconditionally (harmless for IPv4) — aligned with ucx_utils. The ROUTER additionally sets
ZMQ_ROUTER_HANDOVER so a peer that is forgotten (removePeer drops its DEALER) and later reconnects
with the same routing id is accepted, rather than having its messages silently dropped while the stale
connection is reaped. addPeer validates the endpoint before connecting and reports failure to the
caller. With NixlNotifControlChannel the "endpoint" is serialized NIXL metadata; no TCP is involved.

Receiver state (lives on the IO thread, lock-free): BuddyAllocator arena +
flows{ pending: per-chunk bytes, held: region offsets, blockedAtGrantSequence } + a round-robin
ring of active flow keys + a cursor + a grant sequence counter. Fixed per-flow cap
W = maxInflightChunksPerRequest.

schedule() — on-demand, round-robin fair, never poisons the queue, never deadlocks:

flowchart TD
    Start([event triggers schedule:<br/>onWant / onScatterDone / reclaimFlow / reclaimByPrefix / releaseLocal]) --> D0{drain mode active?<br/>a flow bypassed ≥2 full rounds}
    D0 -- yes --> DAlloc{"arena.alloc(drain flow's head)<br/>fits?"}
    DAlloc -- no --> Hold([no NEW remote grants until it fits<br/>existing regions keep freeing])
    DAlloc -- yes --> DGrant[grant it, exit drain mode] --> C1
    D0 -- no --> C1{ring non-empty?}
    C1 -- no --> Done([return accumulated GRANTs])
    C1 -- yes --> Sweep[one round-robin sweep from the cursor]
    Sweep --> Find{current flow:<br/>pending non-empty AND held.size &lt; W?}
    Find -- "none in the whole sweep" --> Done
    Find -- yes --> Alloc{"arena.alloc(pending.front())<br/>fits?"}
    Alloc -- no --> Mark[mark flow blocked at current grant sequence] --> Sweep
    Alloc -- yes --> Grant["off = the allocation<br/>pending.pop_front(), held += off<br/>accumulate GRANT{off, base+off, len}<br/>advance cursor past this flow"]
    Grant --> C1

Loading

Intuition. schedule() answers: many remote senders want to write into the same shared arena —
how to hand out arena space fairly and bounded. Three constraints:

1. Fair: flows take turns (round-robin); no early flow monopolizes.
2. Per-flow rate limit (W): a flow may hold at most W in-flight regions (held.size() < W);
   more must wait for an ACK to free one — this is the pipeline depth.
3. Aggregate rate limit (arena): all flows' in-flight regions share one arena; if the next chunk
   doesn't fit, skip it for now (backpressure, not an error).

Think of it as taking turns at a ticket counter: ring is the queue of flows, cursor is "who's
next". Each inner sweep hands out exactly one region — to the flow that is next, still wants more
(pending non-empty), is under its cap (held < W), and whose front chunk fits right now — then
advances cursor past it and starts a fresh sweep. So grants alternate across flows rather
than filling one flow's window first. The outer loop repeats until a whole sweep grants nothing
(all pending-empty / at-cap / can't-fit).

Example (W=2, arena currently fits 4 regions; flow A has 3 chunks c1/c2/c3, flow B has 2 chunks
d1/d2, cursor starts at A):

Step	cursor	Action	A.held	B.held
1	A	grant A/c1	1	0
2	B	grant B/d1	1	1
3	A	grant A/c2	2	1
4	B	grant B/d2	2	2
5	A	A at cap (2≥W), B at cap → no progress → stop	2	2

Grant order A,B,A,B (strict alternation); A's c3 stays in pending until one of A's regions is
ACKed (onScatterDone→schedule()) and the next schedule() grants it. Each schedule() returns
the new batch of GRANTs, which sendGrants splits by flow key and sends to the right peer.

Notes:

• Two-level rate limiting: per-flow W bounds a single flow's pipeline depth; arena capacity
  bounds aggregate concurrency (alloc failure ⇒ backpressure, not deadlock). There is no "divide cap
  by active flow count" logic.
• Large-chunk anti-starvation (drain mode): a flow whose head chunk fails to fit is stamped with
  the current grant sequence. When other grants have bypassed it for ≥2 full rounds
  (kBypassRounds), the receiver enters drain mode for the oldest such flow: no new remote grants
  are issued until that head chunk fits (existing regions keep progressing and freeing space).
  acquireLocal() is deliberately unaffected — this is a receiver-only admission barrier and cannot
  introduce a bidirectional circular wait. Config guarantees maxChunkSizeBytes ≤ arena usable capacity (clamped at init), so a drained arena can always fit any chunk.
• Flow lifecycle reclaim: the flow key contains a monotonic non-reused rid; when both pending
  and held are empty, eraseIfDone drops the flow immediately — otherwise flows/ring grow
  unbounded on a long-running server and schedule() degrades to O(historical requests).
• Local sender shares the arena: acquireLocal(bytes) takes a gather-staging region from the same
  BuddyAllocator; eager (credit-less) staging is capped at half the arena (§2); the conservation
  invariant holds across {free bytes, each flow's held, localHeld}.

6. Data Plane & Pipeline (R2)

• The receiver fills a flow's held up to W (when space allows, one GRANT message batches multiple
  credits) → the sender holds W credits at once → W chunks in flight; each ACK frees a region and
  the receiver refills to keep held = W.
• Eager gather hides the control RTT: with enableEagerGather (default), submit() immediately
  stages and launches gathers for the first chunks (up to the in-flight cap and the eager half-arena
  budget) before any GRANT arrives; attachCredits later binds arriving credits to
  already-gathered chunks in strict chunk order, promoting their staging regions out of the eager
  budget. The WANT→GRANT round-trip and the gather kernel overlap instead of serializing.
• W is sized by the round-trip: W ≥ ⌈(write + getXferStatus + DATA + scatter + ACK/GRANT return) / single-chunk write⌉ (bandwidth-delay product); gather/scatter (D2D ~TB/s) hide in the shadow of the
  previous chunk's RDMA (IB ~25 GB/s), so the NIC stays the bottleneck.
• Gather and scatter each borrow a separate ExecCtx stream, so different chunks' gather/write/scatter
  overlap in wall-clock.
• Small control messages: the DATA scatter plan is run-merged — adjacent descriptors whose
  (bounceOffset, dstPtr) advance contiguously or by a uniform stride collapse into one
  BounceScatterRun (the fully-dense case, e.g. ctx tp1 → gen tp4, collapses thousands of descs to a
  handful of runs). TRTLLM_NIXL_BOUNCE_DISABLE_SCATTER_RUN_MERGING restores per-desc entries for
  control-plane A/B debugging only.
• No notifMsg / no GPUDirect flush: getXferStatus==SUCCESS already includes NIXL's per-transfer
  ucp_ep_flush_nbx (complete at both origin and target). The sender sends DATA as soon as it polls
  SUCCESS; the receiver scatters on receipt.

Single-chunk timing across both planes (control = solid, data = dashed; eager gather runs ① before
②'s GRANT when enabled):

sequenceDiagram
    autonumber
    participant App as Sender app thread
    participant SIO as Sender IO thread
    participant NIC as Data plane RDMA/NIC
    participant RIO as Receiver IO thread
    participant SW as Receiver scatter worker
    App->>RIO: WANT(rid, chunk_bytes[])
    App-->>App: ① eager GATHER (D2D): N small src → arena.at(o) (launch + eventRecord, no sync)
    Note over RIO: schedule(): arena.alloc carves region s, held[flow]+=s
    RIO->>SIO: GRANT(rid, regionHandle=s, addr, len)
    Note over SIO: onGrant: attachCredits to the eager-gathered chunk (or gather now if not eager)
    Note over SIO: drainGatherReady: cudaEventQuery==success (gather done, no block)
    SIO-->>NIC: ② postWrite(arena.at(o) → addr) async RDMA, no notif, return ExecCtx on done
    NIC-->>RIO: data written into receiver region s
    Note over SIO: pollSenderHandles: poll==SUCCESS (incl. ucp_ep_flush_nbx ⇒ landed at remote target)
    SIO->>RIO: ③ DATA(rid, chunk, regionHandle=s, scatter runs)
    Note over RIO: onData: validate runs against the flow's region, enqueue ScatterJob(s, runs)
    RIO->>SW: ScatterJob(s)
    SW-->>SW: ④ SCATTER (D2D): arena.at(s) → final dst×N (borrow ExecCtx, launch + streamSync)
    SW->>RIO: scatter_done(s)
    Note over RIO: onScatterDone: arena.free(s), schedule() may GRANT the next chunk
    RIO->>SIO: ⑤ ACK(rid, chunk, regionHandle=s)
    Note over SIO: onAck: releaseLocal(region o), acked++, all ACKed ⇒ promise=SUCCESS
    SIO->>App: future ready (SUCCESS)

Loading

For a large request (K > in-flight cap / arena capacity), the loop above repeats with both sides'
memory staying at O(W) regions while the transfer size is unbounded (R1).

7. Threading (R4/R5)

• One IO thread per agent, the sole owner of the CreditScheduler + sender request table (the key
  to being lock-free; the one cross-thread entry is eager acquireLocal from submit()). Each tick:
  1. recv one control message and dispatch;
  2. drainGatherReady (gather event ready → postWrite + return ExecCtx);
  3. pollSenderHandles (poll==SUCCESS → send DATA);
  4. drainScatterDone (worker report → send ACK + free region + reschedule);
  5. drainForgets / drainPendingPosts (retry parked credits) / checkTimeouts.
• M scatter workers: take a ScatterJob → borrow ExecCtx → scatter kernel + sync → report to the IO
  thread.
• submit() never blocks: registers a Request + sends WANT (+ launches eager gathers,
  fire-and-forget), returns a shared_future; safe to call from multiple threads.
• Adaptive poll: when gather/scatter is in flight, recv uses a 0 ms timeout (low latency); fully
  idle uses 1 ms; after long 0 ms spinning it backs off ~50 µs so a long model-kernel gather delay
  can't busy-spin a core.

8. State Machines

Unified region lifecycle (one arena serves both roles; a region is held by exactly one kind of
owner at any time — the conservation invariant):

stateDiagram-v2
    [*] --> FREE: arena init
    FREE --> INCOMING_HELD: GRANT(region→remote flow) 【schedule()】
    FREE --> OUTGOING_HELD: acquireLocal(bytes) 【local gather staging — eager ≤ ½ arena】
    INCOMING_HELD --> QUEUED: DATA(regionHandle, runs) received
    QUEUED --> SCATTERING: scatter worker picks it up (borrow ExecPool ctx)
    SCATTERING --> FREE: scatter_done → onScatterDone (reply ACK + reschedule)
    INCOMING_HELD --> FREE: forgetPeer / reclaimByPrefix (no in-flight DATA)
    OUTGOING_HELD --> FREE: ACK / failure → releaseLocal (reschedule — freed bytes go to a waiting flow)
    FREE --> [*]: shutdown
    note right of OUTGOING_HELD
      Invariant: every region is in exactly one of
      { free arena bytes } ∪ { some remote flow's held } ∪ { localHeld }.
      INCOMING_HELD = remote RDMA-write target (QUEUED/SCATTERING are its sub-states).
      OUTGOING_HELD = local gather source. Never held by both.
    end note

Loading

Sender chunk/request lifecycle (never hangs). With eager gather a chunk may reach Gathered
(gather event signalled, ExecCtx returned) while its GRANT is still in flight; attachCredits then
promotes it straight to the write:

stateDiagram-v2
    [*] --> WANT_SENT: submit() registers Request + sends WANT(chunk_bytes[]) + eager gathers
    WANT_SENT --> GATHERING: eager acquireLocal + launch gather (no credit yet)
    GATHERING --> GATHERED: drainGatherReady sees cudaEventQuery==success (waiting for credit)
    WANT_SENT --> POSTING: GRANT received (onGrant, credits queued)
    GATHERED --> IN_FLIGHT: attachCredits → postWrite
    POSTING --> POSTING: arena/exec full → credit parked, retried by drainPendingPosts (no block)
    POSTING --> GATHERING: borrow ExecCtx + acquireLocal(bytes) → launch gather + eventRecord (no sync)
    GATHERING --> IN_FLIGHT: gather done + credit attached → postWrite + return ExecCtx
    IN_FLIGHT --> DATA_SENT: pollSenderHandles sees poll==SUCCESS → send DATA
    DATA_SENT --> POSTING: ACK received and chunks remain (releaseLocal that region)
    DATA_SENT --> SUCCESS: acked == numChunks
    WANT_SENT --> FAILURE: checkTimeouts, no progress beyond requestTimeoutMs
    POSTING --> FAILURE: forgetPeer / shutdown
    GATHERING --> FAILURE: gather launch/record/stream error / forgetPeer / shutdown
    IN_FLIGHT --> FAILURE: poll==kFailed / forgetPeer / shutdown
    SUCCESS --> [*]: promise=SUCCESS
    FAILURE --> [*]: promise=FAILURE (empty WANT retracts credits)
    note right of WANT_SENT
      Every terminal state resolves the promise →
      the caller's wait() always returns, never hangs (R5)
    end note

Loading

9. Error Handling & Lifecycle (R5)

Every request reaches a terminal state; wait() never hangs. Bounce failures degrade, never break:
maybeInitBounce catches any construction error (fabric alloc, zmq bind, NIXL registration) with a
warning and leaves the agent on the standard per-desc NIXL path.

Case	Trigger	Handling
Peer incompatible (bounce off, other control kind, chunk-size mismatch, bad handshake)	registerPeerHandshake at loadRemoteAgent	Peer not marked bounce-capable → shouldUseBounce routes to standard NIXL (no WANT is ever sent, no timeout burned).
Peer never GRANTs (unreachable/not ready)	checkTimeouts exceeds requestTimeoutMs	Request → FAILURE (empty WANT retracts credits).
Malformed sender endpoint in WANT	onWant addPeer fails/throws	WANT rejected with a warning on the reactor thread (no grant); a cancel is still honored to reclaim earlier flow state.
Transfer engine error	poll==kFailed	Request → FAILURE, release in-flight handle.
gather launch/record error	pumpRequest flags it → drainGatherReady	Request → FAILURE (an unrecorded event would be misread as complete, so fail explicitly).
scatter kernel/sync error	scatter worker	No ACK → sender times out; never falsely reports landed (which would silently corrupt KV). Region is still freed (no leak).
Peer loss	invalidateRemoteAgent → forgetPeer	forgetPeer drops the peer's DEALER + its handshake registration synchronously; the IO thread reclaimByPrefix("peer\x1f") reclaims all of the peer's flows + fails its in-flight requests. A fresh loadRemoteAgent must re-validate a new handshake.
Control message can't be sent (peer stalled, queue full kSendHwm)	non-blocking sendTo returns EAGAIN	Drop the message + WARNING; never block the IO thread. Affected request → FAILURE via requestTimeoutMs (no hang, no corruption).
shutdown	—	join threads → cudaDeviceSynchronize → fail all in-flight requests.

Concurrency-safety points:

• In-flight scatter vs. re-grant race: the receiver keeps a mScattering map (region offset → is
  orphaned) of all scattering regions. If forgetPeer reclaims a region still in mScattering (a
  worker is reading it), reclaimByPrefix defers the free (marks it orphaned) and only
  freeOrphanRegions it on scatter completion — preventing "new sender's RDMA write ⊥ worker's read".
• Gather on the failure path: before reclaiming a region still GATHERING, cudaStreamSynchronize
  its stream (so an abandoned gather can't write a re-granted region); a sync error WARNs and clears the
  sticky error.
• In-flight write on the failure path: a region in state Writing may still be read by the NIC as
  a source, so it cannot be freed immediately. It's recorded in mOrphanLocal;
  drainOrphanLocal() polls its xfer to a terminal state before release + releaseLocal (the
  send-side orphan mechanism, symmetric to the receive-side mScattering orphaned flag).
• Cancel (empty WANT): precise reclaim + late-DATA guard: on sender failure/abort it sends an empty
  WANT; the receiver reclaimFlow immediately frees the flow's granted-but-unwritten regions
  (otherwise they stay held until peer loss, leaking on a long-running receiver), deferring any
  scattering ones. Correspondingly onData validates with heldByFlow: if an empty WANT raced ahead of
  a DATA so the region was freed/re-granted, that late DATA is dropped (never scatter a region now
  owned by someone else).
• Scatter input validation (defense in depth): scatter runs come from the peer's DATA; before
  launch the receiver checks every run's source range lies within this flow's granted region, and
  that the expanded plan doesn't exceed local scratch capacity. Any out-of-bounds → no launch, no ACK.
• CUDA errors: gather/scatter/sync return codes are not swallowed; errors WARN (with
  cudaGetErrorString).

Entry routing (transparent to callers; disabling is byte-equivalent to the original NIXL path):

flowchart TD
    S[submitTransferRequests] --> E{shouldUseBounce?<br/>WRITE + both-VRAM + no syncMsg<br/>+ peer handshake OK + descCount/avg<br/>+ per-side uniform deviceId}
    E -- no --> N[standard NIXL path]
    E -- yes --> SUB[submit: register Request + send WANT + eager gathers + return future]
    SUB --> POST[IO thread: GRANT→attach/gather+postWrite→DATA, pipelined]
    POST --> ERR{poll==kFailed / scatter fail / peer gone / stalled beyond requestTimeoutMs?}
    ERR -- yes --> F[Request → FAILURE]
    ERR -- no --> OK{all chunks poll SUCCESS AND all ACKed?}
    OK -- yes --> SU[wait = SUCCESS]
    F --> W[wait = FAILURE → caller task.fail]

Loading

Bootstrap: the bounce handshake blob is serialized with the AgentDesc (getLocalAgentDesc /
loadRemoteAgent(AgentDesc) → registerPeerHandshake), i.e. the path production disagg already uses;
the first WANT starts directly with no separate handshake round-trip.

10. Configuration (env)

All prefixed TRTLLM_NIXL_BOUNCE_. Byte-valued variables accept case-insensitive binary suffixes
(K/KB/KiB, M/MB/MiB, G/GB/GiB; all powers of two), e.g. ARENA_SIZE_BYTES=512MB. Unparsable
values fall back to the default (never silently become 0).

Field	env (suffix after prefix)	Default	Meaning
enabled	ENABLE	off	Master switch.
arenaSizeBytes	ARENA_SIZE_BYTES	256 MiB	Shared region arena size.
arenaAllocationGranularityBytes	ARENA_ALLOCATION_GRANULARITY_BYTES	1 MiB	Buddy minimum block (allocation granularity).
maxChunkSizeBytes	MAX_CHUNK_SIZE_BYTES	32 MiB	Per-chunk byte cap (plan bin-pack cap); clamped to the arena's usable capacity if larger; must match the peer's (handshake-checked).
maxInflightChunksPerRequest	MAX_INFLIGHT_CHUNKS_PER_REQUEST	8	Per-flow in-flight cap W (pipeline depth).
copyStreamCount	COPY_STREAM_COUNT	8	ExecPool context count (GPU kernel concurrency cap).
scatterWorkerCount	SCATTER_WORKER_COUNT	4	Scatter worker thread count.
minDescriptorCount	MIN_DESCRIPTOR_COUNT	1024	Routing gate: minimum descriptor count.
maxAverageDescriptorSizeBytes	MAX_AVERAGE_DESCRIPTOR_SIZE_BYTES	16 KiB	Routing gate: maximum average descriptor bytes.
requestTimeoutMs	REQUEST_TIMEOUT_MS	30000	No-progress timeout.
disableFabricMemory	DISABLE_FABRIC_MEMORY	off	Use cudaMalloc instead of MNNVL fabric memory (CI/x86).
enableEagerGather	ENABLE_EAGER_GATHER	on	Launch gathers at submit() before GRANT (overlap control RTT); eager staging capped at ½ arena.
useNixlNotifications	USE_NIXL_NOTIFICATIONS	off	Control plane over NIXL notifications (UCX active messages) instead of ZMQ/TCP; must match the peer (handshake-checked).
useZeroCopyArguments	USE_ZERO_COPY_ARGUMENTS	on	Copy kernel reads its plan directly from pinned host memory (no H2D staging).
useCubCopy	USE_CUB_COPY	off	Use cub::DeviceMemcpy::Batched instead of the custom copy kernel (experimental).
disableScatterRunMerging	DISABLE_SCATTER_RUN_MERGING	off	DEBUG ONLY: per-desc DATA plan instead of coalesced runs.

shouldUseBounce fires when: op is WRITE, src/dst are both VRAM, no syncMessage, the peer passed
the capability handshake, descCount ≥ minDescriptorCount, all srcs are on this agent's device and
all dsts on one device, and average desc bytes ≤ maxAverageDescriptorSizeBytes; otherwise the
standard NIXL path is used.

11. Test Coverage

Tests live under cpp/tests/unit_tests/executor/bounce/.

• Pure logic (no GPU): buddyAllocatorTest (split/coalesce/fragmentation/boundaries/overflow),
  creditSchedulerTest (in-flight cap/fairness/drain-mode anti-starvation/reclaim/conservation/
  reclaim-defer/orphan/eager budget), bounceMessageCodecTest (round-trip/truncation/magic/
  cross-type-reject/large-count/handshake codec), bounceTransferPlanTest (bin-pack boundaries +
  scatter-run merging), bounceConfigTest (env parsing, byte suffixes, garbage fallback).
• GPU unit: bounceArenaTest, execPoolTest, gatherScatterKernelTest (custom kernel /
  zero-copy args / cub backend).
• Transport (real NIXL loopback + zmq): zmqControlChannelTest (incl. endpoint validation),
  bounceTransportTest (end-to-end, byte-exact; handshake compatibility; malformed-WANT rejection),
  bounceTransportFailureTest (no-GRANT timeout / engine failure / shutdown in-flight / forgetPeer
  in-flight / multi-peer shared-arena over-subscription no-deadlock / multi-threaded submit).
• Real NIXL RDMA e2e: nixlTransferEngineTest; bounceNixlE2ETest (RealRdmaLoopback single
  transfer / ConcurrentBidirectionalRealRdma 8-thread bidirectional /
  MultiAgentManySendersToOneReceiver / ForgetPeerInFlightRecovers); bounceAgentE2ETest
  (production submitTransferRequests path: single transfer + ConcurrentSubmitUsesBounce).
  All e2e tests verify byte-exactly (seed-distinct pattern per transfer, ruling out cross-talk).
• Python integration: test_cache_transceiver_single_process.py drives the NIXL bounce path
  through the Python cache transceiver (added to l0_h100.yml).

12. perf compare

https://docs.google.com/document/d/1J8ROqb1D-TQryIEyqLoYtP4Z3Hqk_7AreluczjdtH5w/edit?usp=sharing
gptoss gb200_gpt-oss-120b-fp4_8k1k_con128_ctx1_tp1_gen1_tp4_eplb0_mtp0_ccb-NIXL

con128

	gen_side p50/p90/p99/mean (ms)	max (ms)
py_bounce	10.1 / 108.6 / 116.7 / 25.1	2660
py_bounce_t1	9.8 / 113.6 / 119.3 / 40.5	2675
bounce_v2	4.2 / 5.0 / 6.0 / 4.3	201
bounce_v2_zmq	4.0 / 4.9 / 5.7 / 4.1	105

con1024

	gen_side p50/p90/p99/mean (ms)	max (ms)
py_bounce	10.2 / 113.6 / 119.1 / 38.0	2721
py_bounce_t1	10.1 / 116.8 / 120.3 / 45.0	2700
bounce_v2	4.5 / 5.5 / 6.3 / 4.6	509
bounce_v2_zmq	4.6 / 5.5 / 6.4 / 4.6	197

PR Checklist

Please review the following before submitting your PR:

• PR description clearly explains what and why. If using CodeRabbit's summary, please make sure it makes sense.

• PR Follows TRT-LLM CODING GUIDELINES to the best of your knowledge.

• Test cases are provided for new code paths (see test instructions)

• If PR introduces API changes, an appropriate PR label is added - either api-compatible or api-breaking. For api-breaking, include BREAKING in the PR title.

• Any new dependencies have been scanned for license and vulnerabilities

• CODEOWNERS updated if ownership changes

• Documentation updated as needed

• Update tava architecture diagram if there is a significant design change in PR.

• The reviewers assigned automatically/manually are appropriate for the PR.

• Please check this after reviewing the above items as appropriate for this PR.

GitHub Bot Help

To see a list of available CI bot commands, please comment /bot help.

[chuangz0]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #56919 [ run ] triggered by Bot. Commit: 006072a Link to invocation

[tensorrt-cicd]

PR_Github #56919 [ run ] completed with state SUCCESS. Commit: 006072a
/LLM/main/L0_MergeRequest_PR pipeline #45725 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

• Please check the failed tests and fix your PR
• If you cannot view the failures, ask the CI triggerer to share details
• Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

[chuangz0]

/bot run --disable-fail-fast

[chuangz0]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #61555 [ run ] triggered by Bot. Commit: b7c648d Link to invocation

[tensorrt-cicd]

PR_Github #61555 [ run ] completed with state FAILURE. Commit: b7c648d
/LLM/main/L0_MergeRequest_PR pipeline #49766 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

• Please check the failed tests and fix your PR
• If you cannot view the failures, ask the CI triggerer to share details
• Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

[chuangz0]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #61779 [ run ] triggered by Bot. Commit: 2aef331 Link to invocation

[tensorrt-cicd]

PR_Github #61779 [ run ] completed with state FAILURE. Commit: 2aef331
/LLM/main/L0_MergeRequest_PR pipeline #49979 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

• Please check the failed tests and fix your PR
• If you cannot view the failures, ask the CI triggerer to share details
• Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

[chuangz0]

/bot run --disable-fail-fast

[coderabbitai]

Walkthrough

Adds an opt-in Bounce v2 transport for NIXL GPU transfers. The change defines wire protocols, memory and scheduling components, control channels, asynchronous transfer handling, agent handshake integration, build configuration, and C++, CUDA, NIXL, ZeroMQ, and Python tests.

Changes

Bounce v2 transport

Layer / File(s)	Summary
Protocol, configuration, and memory components cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/*	Adds environment configuration, wire codecs, capability handshakes, CUDA arenas, buddy allocation, transfer planning, execution pools, batched copy kernels, and NVTX instrumentation.
Control channels and transfer engine cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/*	Adds abstract and NIXL/ZeroMQ control channels plus asynchronous NIXL write registration, posting, polling, and cleanup.
Transport reactor and scheduling cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/BounceTransport.*, cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/CreditScheduler.*	Adds sender and receiver state machines, credit scheduling, CUDA gather/scatter execution, acknowledgements, timeouts, peer reclamation, shutdown, and failure handling.
Agent and descriptor integration cpp/include/tensorrt_llm/executor/transferAgent.h, cpp/tensorrt_llm/executor/cache_transmission/transferAgent.cpp, cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/transferAgent.*, cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/agentBindings.cpp, tensorrt_llm/_torch/disaggregation/nixl/_agent_cpp.py	Adds serialized bounce handshakes, optional transport initialization, eligibility checks, peer registration, bounce submission status, fallback to standard NIXL, and Python observability properties.
Build and validation coverage cpp/CMakeLists.txt, cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/CMakeLists.txt, cpp/tests/unit_tests/executor/bounce/*, tests/integration/test_lists/test-db/l0_h100.yml, tests/unittest/disaggregated/test_cache_transceiver_single_process.py	Adds conditional build wiring and pure-logic, CUDA, ZeroMQ, NIXL, failure, concurrency, end-to-end, and Python integration tests.

Estimated code review effort: 5 (Critical) | ~120 minutes

Possibly related PRs

• NVIDIA/TensorRT-LLM#15727: Adds Python NIXL transfer orchestration related to the C++ Bounce v2 path.
• NVIDIA/TensorRT-LLM#17334: Extends bounce-buffer infrastructure for disaggregated serving.

Suggested reviewers: brnguyen2

Sequence Diagram(s)

sequenceDiagram
  participant NixlTransferAgent
  participant BounceTransport
  participant ControlChannel
  participant NixlTransferEngine
  NixlTransferAgent->>BounceTransport: Submit eligible VRAM descriptors
  BounceTransport->>ControlChannel: Exchange WANT and GRANT messages
  BounceTransport->>NixlTransferEngine: Post remote writes
  NixlTransferEngine-->>BounceTransport: Report write completion
  BounceTransport->>ControlChannel: Exchange DATA and ACK messages
  BounceTransport-->>NixlTransferAgent: Resolve transfer status

Loading

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name	Status	Explanation	Resolution
Docstring Coverage	⚠️ Warning	Docstring coverage is 28.14% which is insufficient. The required threshold is 80.00%.	Write docstrings for the functions missing them to satisfy the coverage threshold.

✅ Passed checks (4 passed)

Check name	Status	Explanation
Title check	✅ Passed	The title clearly identifies a feature that adds NIXL bounce-buffer support for the cache transceiver.
Description check	✅ Passed	The description fully explains the design, motivation, implementation, configuration, test coverage, performance results, and checklist status.
Linked Issues check	✅ Passed	Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check	✅ Passed	Check skipped because no linked issues were found for this pull request.

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create PR with unit tests

---

_{Comment @coderabbitai help to get the list of available commands.}

[coderabbitai]

Actionable comments posted: 16

🧹 Nitpick comments (14)

tests/unittest/disaggregated/test_cache_transceiver_single_process.py (2)

1564-1567: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low value

Remove the redundant env.pop call.

The loop at lines 1549-1551 already deletes every variable whose name starts with TRTLLM_NIXL_BOUNCE_. TRTLLM_NIXL_BOUNCE_DISABLE_FABRIC_MEMORY matches that prefix, so line 1567 never removes anything. Keep the comment, which explains the intent, and drop the call.

♻️ Proposed cleanup

     # Do not override TRTLLM_NIXL_BOUNCE_DISABLE_FABRIC_MEMORY: the test must
     # exercise the production default, including automatic cudaMalloc fallback
-    # on devices without fabric-memory support.
-    env.pop("TRTLLM_NIXL_BOUNCE_DISABLE_FABRIC_MEMORY", None)
+    # on devices without fabric-memory support. The prefix loop above already
+    # cleared it from the inherited environment.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@tests/unittest/disaggregated/test_cache_transceiver_single_process.py` around
lines 1564 - 1567, Remove the redundant env.pop call for
TRTLLM_NIXL_BOUNCE_DISABLE_FABRIC_MEMORY while preserving the existing comment
and prefix-based environment cleanup loop.

---

1501-1582: 📐 Maintainability & Code Quality | 🔵 Trivial

Test coverage summary (required for tests/**).

1. Test functions added, modified, or removed

Added, Python:

• test_python_nixl_cache_transceiver_uses_cpp_bounce with four parameters: v2_mha, v2_mla_tp2, v1_dsa_indexer, v2_ctx_tp2_gen_pp2.

Modified, Python:

• run_transfer_test now pins kv_cache_bounce_size_mb=0. This is a shared helper, so the change affects every existing caller in the file.

Added, C++ gtest (files supplied in this review):

• bounceMessageCodecTest.cpp: 16 tests covering WANT, cancel, GRANT, DATA, ACK, empty payloads, >65535 entry counts, short blobs, bad magic, inflated payload length, cross-type decode, and handshake round-trip, garbage, and truncation.
• bounceTransferPlanTest.cpp: 16 tests covering empty plans, 32-byte alignment, chunk-size and descriptor-count splitting, oversize rejection, device-id mismatch, zero-length descriptors, descriptor merging, and scatter-run coalescing.
• buddyAllocatorTest.cpp: 12 tests covering rounding, no-overlap, exhaustion, coalescing, recycling, overflow rejection, double free, unknown free, and fragmentation.
• creditSchedulerTest.cpp: 20 tests covering inflight caps, fairness, reclamation, busy-region deferral, tombstone removal, cursor reset, variable-size packing, and large-chunk starvation.
• zmqControlChannelTest.cpp: 5 tests covering bidirectional TCP, IPv6, FIFO ordering, idle timeout, endpoint validation, and non-blocking send.

Added, C++ gtest targets declared in bounce/CMakeLists.txt but not supplied in this review: bounceConfigTest, bounceArenaTest, execPoolTest, gatherScatterKernelTest, bounceTransportTest, bounceTransportFailureTest, bounceNixlE2ETest, bounceAgentE2ETest, and the opt-in bounceTransportTsanTest.

2. Test list membership

• test_python_nixl_cache_transceiver_uses_cpp_bounce is added to tests/integration/test_lists/test-db/l0_h100.yml at line 101. The entry has no parameter filter, so all four parameters run.
• The C++ gtest targets are registered through add_gtest in cpp/tests/unit_tests/executor/bounce/CMakeLists.txt and run through the C++ test harness. They do not belong in the integration test lists.

3. Coverage verdict: needs follow-up

Reasons:

• Nine of the declared C++ test files in this layer were not supplied in this prompt. Their assertions and gating are unverified.
• The v2_mla_tp2 and v2_ctx_tp2_gen_pp2 parameters need two GPUs. Confirm that the l0_h100 tier at line 101 provides them, and confirm the tier's wall-time budget for four nested pytest subprocesses.
• No test asserts the fallback behavior that the PR objectives describe, where bounce initialization failure falls back to the standard NIXL path. Add a case that forces bounce initialization to fail and confirms the transfer still completes.
• The effect of kv_cache_bounce_size_mb=0 on the pre-existing tests in test_cache_transceiver_single_process.py is unconfirmed. See the separate comment on lines 998-1000.

As per path instructions: "Always produce a test coverage summary, even if no issues are found."

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@tests/unittest/disaggregated/test_cache_transceiver_single_process.py` around
lines 1501 - 1582, Extend test coverage around
test_python_nixl_cache_transceiver_uses_cpp_bounce by adding a scenario that
forces bounce initialization failure and verifies the standard NIXL path still
completes. Validate that the v2_mla_tp2 and v2_ctx_tp2_gen_pp2 parameters are
gated to a two-GPU-capable tier and fit its timeout budget, and verify the
shared run_transfer_test kv_cache_bounce_size_mb=0 change does not regress
existing callers. Supply or explicitly gate the declared bounce C++ test targets
so their assertions and execution coverage are verifiable.

Source: Path instructions

cpp/tests/unit_tests/executor/bounce/bounceTransferPlanTest.cpp (1)

94-131: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Assert TllmException instead of accepting any exception.

Replace the five EXPECT_ANY_THROW assertions with EXPECT_THROW(..., tensorrt_llm::common::TllmException). Add a direct include for tensorrt_llm/common/tllmException.h.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/tests/unit_tests/executor/bounce/bounceTransferPlanTest.cpp` around lines
94 - 131, In the BounceTransferPlan tests, replace all five EXPECT_ANY_THROW
assertions with EXPECT_THROW assertions specifying
tensorrt_llm::common::TllmException, including the cases in
DescLargerThanChunkThrows, MaxChunkSizeBytesAboveU32Throws,
MixedSourceDeviceIdsThrow, and MixedDestinationDeviceIdsThrow. Add a direct
include for tensorrt_llm/common/tllmException.h.

cpp/tests/unit_tests/executor/bounce/CMakeLists.txt (1)

57-69: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Use PkgConfig::ZMQ for both ZMQ consumers.

Change pkg_check_modules(ZMQ libzmq) to pkg_check_modules(ZMQ IMPORTED_TARGET libzmq). Link PkgConfig::ZMQ and remove the separate ${ZMQ_INCLUDE_DIRS} and ${ZMQ_LIBRARIES} entries. This preserves pkg-config include paths, library paths, and linker flags.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/tests/unit_tests/executor/bounce/CMakeLists.txt` around lines 57 - 69,
Update the ZMQ discovery in the CMake configuration around pkg_check_modules to
request the imported target, then link zmqControlChannelTest against
PkgConfig::ZMQ. Remove the separate ${ZMQ_INCLUDE_DIRS} include path and
${ZMQ_LIBRARIES} link entry while retaining the cppzmq include directory.

cpp/tests/unit_tests/executor/bounce/bounceTransportTest.cpp (1)

100-105: 🎯 Functional Correctness | 🔵 Trivial | ⚡ Quick win

ConcurrentRequestsToSameReceiver does not test concurrency or a shared receiver.

The test calls runTransfer twice in sequence. Each call builds a new agent pair (btConc1A/btConc1B, then btConc2A/btConc2B), submits one transfer, waits for it, and shuts both nodes down. No two requests are ever in flight at the same time, and the two transfers do not share a receiver.

Concurrent submission is covered by BounceTransportFailure.ConcurrentMultiThreadedSubmit in cpp/tests/unit_tests/executor/bounce/bounceTransportFailureTest.cpp. Either submit two futures on one node pair before waiting, or rename the test to describe the sequential behavior it actually checks.

♻️ Sketch of a real concurrent-submit variant

TEST(BounceTransport, ConcurrentRequestsToSameReceiver)
{
    // ... build one node pair A/B, wirePair(*A, *B) ...
    auto b1 = bounce_test::makeXferBufs(8, 1024, /*seed=*/1);
    auto b2 = bounce_test::makeXferBufs(8, 1024, /*seed=*/2);
    auto f1 = A->tx->submit(b1.srcDescs, b1.dstDescs, B->name); // both in flight
    auto f2 = A->tx->submit(b2.srcDescs, b2.dstDescs, B->name);
    // ... then wait on f1 and f2 and verify both ...
}

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/tests/unit_tests/executor/bounce/bounceTransportTest.cpp` around lines
100 - 105, Update the ConcurrentRequestsToSameReceiver test to create one shared
sender/receiver pair, submit two distinct transfers through the same receiver
before waiting on either result, then await and verify both futures before
cleanup. Do not use separate sequential runTransfer calls; reuse the existing
node-pair setup and transfer verification helpers.

cpp/tests/unit_tests/executor/bounce/bounceArenaTest.cpp (2)

43-44: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low value

Brace the skip guards.

The repository C++ style requires braces on every if body. Lines 43-44 and 59-60 use unbraced bodies. The same pattern exists in other bounce test files.

As per coding guidelines: "Use Allman brace style; always brace if/else, loop, and switch bodies".

Also applies to: 59-60

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/tests/unit_tests/executor/bounce/bounceArenaTest.cpp` around lines 43 -
44, Add Allman-style braces around the skip guards controlled by hasCuda() at
both affected locations, and apply the same bracing style to matching unbraced
skip guards in the other bounce test files.

Source: Coding guidelines

---

43-60: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low value

Unbraced single-statement bodies across the bounce tests. Several new test files write single-statement if and for bodies without braces. The repository C++ style requires braces on all such bodies. One consistent edit fixes every site.

• cpp/tests/unit_tests/executor/bounce/bounceArenaTest.cpp#L43-L60: brace the two if (!hasCuda()) GTEST_SKIP(); guards at lines 43-44 and 59-60.
• cpp/tests/unit_tests/executor/bounce/bounceTransportFailureTest.cpp#L107-L112: brace the hasCuda() and node-null guards at lines 107-108, 111-112, 127-128, 154-155, 158-159, 175-176, 179-180, 202-203, 216-217, and the for bodies at lines 242-243 and 293-294.
• cpp/tests/unit_tests/executor/bounce/execPoolTest.cpp#L41-L42: brace the skip guards at lines 41-42 and 82-83.
• cpp/tests/unit_tests/executor/bounce/gatherScatterKernelTest.cpp#L180-L182: brace the nested for bodies at lines 180-182 and 255-257, and the loop and if bodies at lines 262-267.

As per coding guidelines: "Use Allman brace style; always brace if/else, loop, and switch bodies".

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/tests/unit_tests/executor/bounce/bounceArenaTest.cpp` around lines 43 -
60, Brace every single-statement if and for body using Allman style across the
bounce tests: in cpp/tests/unit_tests/executor/bounce/bounceArenaTest.cpp lines
43-60, cpp/tests/unit_tests/executor/bounce/bounceTransportFailureTest.cpp lines
107-112, 111-112, 127-128, 154-155, 158-159, 175-176, 179-180, 202-203, 216-217,
242-243, and 293-294, cpp/tests/unit_tests/executor/bounce/execPoolTest.cpp
lines 41-42 and 82-83, and
cpp/tests/unit_tests/executor/bounce/gatherScatterKernelTest.cpp lines 180-182,
255-257, and 262-267. Preserve the existing guard and loop behavior while adding
braces to all listed bodies.

Source: Coding guidelines

cpp/tests/unit_tests/executor/bounce/gatherScatterKernelTest.cpp (1)

93-101: 🎯 Functional Correctness | 🔵 Trivial | ⚡ Quick win

The gather step does not compact the buffers.

gatherDst[i] uses the same off[i] as gatherSrc[i], so the "packed" region has the identical sparse, 256-aligned layout as the source. The kernel therefore never copies into a compacted destination layout, which is the layout the transfer planner produces for a real chunk.

Use a second, tightly-packed offset array for the arena region so the test covers unequal src and dst layouts.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/tests/unit_tests/executor/bounce/gatherScatterKernelTest.cpp` around
lines 93 - 101, The gather/scatter test currently uses the sparse source offsets
for both source and packed-region addresses. Introduce a separate tightly packed
offset sequence for the arena, and use it for gatherDst and scatterSrc while
retaining off[i] for gatherSrc and scatterDst, so the test exercises unequal
source and destination layouts.

cpp/tests/unit_tests/executor/bounce/bounceAgentE2ETest.cpp (1)

140-148: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Restore the bounce environment variables with an RAII guard.

setenv calls in this file are undone only by the explicit clearBounceEnv() at the end of each test. If the agent construction fails, the test at Line 162 calls GTEST_SKIP() at Line 188 without clearing the variables. The same happens if any ASSERT_* in the test body aborts the test early. The variables then stay set for every later test in the same process, and bounce becomes enabled where the test did not request it.

bounceConfigTest.cpp already defines a ScopedEnv class for this purpose. Extract it into a shared test header and use it here, or add a small local RAII guard whose destructor calls clearBounceEnv().

Also note that setBounceEnv() never sets TRTLLM_NIXL_BOUNCE_USE_NIXL_NOTIFICATIONS, but clearBounceEnv() unsets it. That asymmetry is safe today, but it hides the dependency on the notification test at Line 275.

Also applies to: 162-189

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/tests/unit_tests/executor/bounce/bounceAgentE2ETest.cpp` around lines 140
- 148, Replace the manual bounce environment cleanup around setBounceEnv and
clearBounceEnv with an RAII guard, preferably by reusing a shared ScopedEnv from
bounceConfigTest.cpp or adding a local guard whose destructor calls
clearBounceEnv. Ensure the guard covers the test setup and body so GTEST_SKIP or
ASSERT_* still restores all bounce variables; include
TRTLLM_NIXL_BOUNCE_USE_NIXL_NOTIFICATIONS in the environment setup or otherwise
keep the managed variable set symmetric.

cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/ControlChannel.h (1)

72-76: 🩺 Stability & Availability | 🔵 Trivial | ⚡ Quick win

Mark addPeer as [[nodiscard]].

addPeer reports a failed endpoint load through its return value. Every other value-returning method in this interface is [[nodiscard]]. If a caller drops the result, all later sendTo(peer, ...) calls are dropped silently and the request stalls until requestTimeoutMs. Add the attribute so the compiler catches an ignored failure.

♻️ Proposed change

-    virtual bool addPeer(std::string const& peer, std::string const& endpoint) = 0;
+    [[nodiscard]] virtual bool addPeer(std::string const& peer, std::string const& endpoint) = 0;

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In
`@cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/ControlChannel.h`
around lines 72 - 76, Mark the virtual addPeer method in the ControlChannel
interface as [[nodiscard]] so callers must handle endpoint-registration
failures. Preserve its existing signature, behavior, and documentation.

cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/BounceNvtx.h (1)

95-103: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Add a printf format attribute to the variadic helpers.

All call sites pass literal format strings, so the static analysis warning about a non-literal format string is a false positive here. The compiler still cannot check the arguments against the format. Add __attribute__((format(printf, N, N+1))) so a format/argument mismatch becomes a compile-time diagnostic instead of undefined behavior at run time. For the constructor the implicit this is argument 1, so the format index is 3.

♻️ Proposed fix

 `#ifndef` NVTX_DISABLE
+    __attribute__((format(printf, 3, 4)))
     BounceNvtxScope(std::uint32_t argb, char const* fmt, ...)

 `#ifndef` NVTX_DISABLE
+__attribute__((format(printf, 2, 3)))
 inline std::uint64_t bounceRangeStart(std::uint32_t argb, char const* fmt, ...)

Apply the same attribute to the NVTX_DISABLE variants so both builds check identically.

Also applies to: 120-128

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/BounceNvtx.h`
around lines 95 - 103, Add the printf format attribute to the variadic
BounceNvtxScope constructor, using format index 3 and argument index 4 because
of the implicit this parameter. Apply the equivalent attribute to the
NVTX_DISABLE variant as well, so both enabled and disabled builds validate
format strings and arguments consistently.

Source: Linters/SAST tools

cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/ZmqControlChannel.cpp (1)

158-189: 🩺 Stability & Availability | 🔵 Trivial | ⚡ Quick win

Handle a partially read multipart message.

If the second recv at line 168 returns no value, the function returns false after it consumed the identity frame. The body frame stays in the socket queue. The next recv call then reads that body frame as an identity frame, which is the desync the comment at line 175 describes. Drain the remaining frames of the current message before returning false.

♻️ Proposed fix

     zmq::message_t bodyFrame;
     auto r2 = mRouter.recv(bodyFrame, zmq::recv_flags::none);
     if (!r2.has_value())
     {
+        // The identity frame is already consumed; drop the rest of this message so the next
+        // recv() starts on a frame boundary.
+        while (idFrame.more())
+        {
+            zmq::message_t drop;
+            if (!mRouter.recv(drop, zmq::recv_flags::none).has_value())
+            {
+                break;
+            }
+            idFrame.swap(drop);
+        }
         return false;
     }

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In
`@cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/ZmqControlChannel.cpp`
around lines 158 - 189, Update the receive path in the surrounding function so
that when the second mRouter.recv call for bodyFrame fails, it drains the
remaining multipart frames of the current message before returning false. Reuse
the existing trailing-frame drain pattern based on the message’s more flag,
while preserving the current successful [identity, body] handling and failure
return behavior.

cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/ExecPool.cpp (1)

80-92: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Add braces to the if bodies.

The repository C++ style requires braces on all if bodies. Lines 82-91 use unbraced single statements.

♻️ Proposed fix

     for (auto& c : mCtxs)
     {
         if (c.scratch != nullptr)
+        {
             TLLM_CUDA_CHECK_WARN(cudaFree(c.scratch));
+        }
         if (c.cubTemp != nullptr)
+        {
             TLLM_CUDA_CHECK_WARN(cudaFree(c.cubTemp));
+        }
         if (c.hostPinned != nullptr)
+        {
             TLLM_CUDA_CHECK_WARN(cudaFreeHost(c.hostPinned)); // also frees its hostPinnedDev alias
+        }
         if (c.stream != nullptr)
+        {
             TLLM_CUDA_CHECK_WARN(cudaStreamDestroy(c.stream));
+        }
         if (c.event != nullptr)
+        {
             TLLM_CUDA_CHECK_WARN(cudaEventDestroy(c.event));
+        }
     }

Based on coding guidelines: "Use Allman brace style; always brace if/else, loop, and switch bodies".

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/ExecPool.cpp`
around lines 80 - 92, Add Allman-style braces around every single-statement if
body in the mCtxs cleanup loop, covering scratch, cubTemp, hostPinned, stream,
and event resource cleanup while preserving the existing conditions and calls.

Source: Coding guidelines

cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/CMakeLists.txt (1)

53-74: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Consider the imported pkg-config target for libzmq.

pkg_check_modules(ZMQ libzmq) without IMPORTED_TARGET gives bare library names in ${ZMQ_LIBRARIES} and does not apply ${ZMQ_LIBRARY_DIRS}. If libzmq lives outside the default linker search path, the link fails. Add IMPORTED_TARGET and link PkgConfig::ZMQ, which carries include dirs, link dirs, and libraries.

The cppzmq include path also hardcodes the FetchContent layout ${CMAKE_BINARY_DIR}/_deps/cppzmq-src. Prefer the cppzmq_SOURCE_DIR variable that FetchContent populates.

♻️ Proposed CMake adjustment

-    pkg_check_modules(ZMQ libzmq) # optional: bounce is skipped if libzmq is
-                                  # absent
+    pkg_check_modules(ZMQ IMPORTED_TARGET libzmq) # optional: bounce is skipped
+                                                  # if libzmq is absent

     target_include_directories(
-      ${NIXL_WRAPPER_TARGET} PRIVATE ${CMAKE_BINARY_DIR}/_deps/cppzmq-src
-                                     ${ZMQ_INCLUDE_DIRS})
-    target_link_libraries(${NIXL_WRAPPER_TARGET} PRIVATE ${ZMQ_LIBRARIES}
-                                                         ${NIXL_BUILD_LIBRARY})
+      ${NIXL_WRAPPER_TARGET} PRIVATE ${cppzmq_SOURCE_DIR})
+    target_link_libraries(${NIXL_WRAPPER_TARGET} PRIVATE PkgConfig::ZMQ
+                                                         NIXL::nixl)

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/CMakeLists.txt`
around lines 53 - 74, Update the libzmq discovery and linking in the ZMQ_FOUND
block: call pkg_check_modules with IMPORTED_TARGET and link the
NIXL_WRAPPER_TARGET against PkgConfig::ZMQ so its include paths, library
directories, and libraries are propagated. Replace the hardcoded cppzmq path in
target_include_directories with the FetchContent-provided cppzmq_SOURCE_DIR
variable.

🤖 Prompt for all review comments with AI agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In
`@cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/BounceConfig.h`:
- Around line 192-200: Update the request-timeout configuration call using
envInt to pass false for allowZero, matching the other operational knobs and
ensuring TRTLLM_NIXL_BOUNCE_REQUEST_TIMEOUT_MS=0 falls back to the default.
Apply the same change to the corresponding occurrence identified around the
second location.

In
`@cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/BounceMessage.cpp`:
- Around line 209-218: Update decodeHandshake to check blob.size() against the
combined sizes of kHandshakeMagic and kHandshakeFormatVersion before calling
serialize_utils::deserialize. Return false for undersized blobs, then preserve
the existing deserialization and validation flow for sufficiently large
handshakes.

In
`@cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/BounceTransport.cpp`:
- Around line 816-828: Add the same nTotal <= maxEntries capacity guard used by
the scatter path to the gather path before planBufs(ctx, nTotal) and
appendSplitInto write entries. Use the existing guard behavior and return or
propagate the established failure result when the planned entry count exceeds
maxEntries, ensuring no buffer planning or writes occur beyond capacity.
- Around line 1498-1528: Add an exception boundary in BounceTransport::ioLoop
around recv(), dispatch(), and all drain operations so zmq::error_t or other
failures are caught, logged, and converted into failure handling for affected
transfers without terminating the IO thread or process. Preserve
status-code-based transfer behavior. NixlTransferEngine.cpp lines 110-166
requires no direct change; do not add exception wrappers there.

In `@cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/ExecPool.cpp`:
- Around line 50-73: Update the constructor’s resource-allocation loop to catch
exceptions from TLLM_CUDA_CHECK, invoke a private cleanup helper for all
already-initialized entries in mCtxs, then rethrow; ensure the helper safely
releases each allocated stream, event, device buffer, pinned buffer, and
optional cubTemp before construction fails.

In
`@cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/NixlNotifControlChannel.cpp`:
- Around line 110-129: Update NixlNotifControlChannel::sendTo to check mPeers
membership while holding mMu, then release the mutex before invoking
mAgent->genNotif. Preserve the unknown-peer warning and early return, while
ensuring the fabric operation and its failure logging occur outside the lock.
- Around line 88-95: In the addPeer flow around loadRemoteMD, track whether this
call inserted the metadata and only invalidate or remove it on a
loadedName-versus-peer mismatch when newly inserted. Preserve existing metadata
for agents already used by the data plane, while retaining the warning and false
return behavior.

In `@cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/transferAgent.cpp`:
- Around line 254-276: Update shouldUseBounce in
cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/transferAgent.cpp at
lines 254-276 to reject requests when any source/destination descriptor length
pair differs or any descriptor exceeds cfg.maxChunkSizeBytes, while retaining
the existing eligibility checks. In BounceTransport.cpp at lines 595-597, keep
the gate as the primary defense and catch TllmException from
BounceTransferPlan::build, resolving the returned future with kFAILURE instead
of propagating the exception.

In `@cpp/tensorrt_llm/executor/cache_transmission/transferAgent.cpp`:
- Around line 128-132: Update AgentDesc deserialization around bounceHandshake
to accept EOF immediately after the VRAM regions by using an empty handshake,
while continuing to fail when a present handshake is truncated. Preserve
validation for region data and add a regression test covering blobs produced by
the legacy serializer layout.

In `@cpp/tests/unit_tests/executor/bounce/bounceNixlE2ETest.cpp`:
- Around line 50-203: Remove the duplicated helper implementations from
bounceNixlE2ETest and include bounceTestNixlNode.h instead. Update references to
hasCuda, XferBufs, makeXferBufs, verifyXferBufs, Node, makeNode, and wirePair to
use the bounce_test:: namespace, and use the shared freeXferBufs helper where
cleanup is needed.

In `@cpp/tests/unit_tests/executor/bounce/buddyAllocatorTest.cpp`:
- Around line 298-304: Apply Allman-style braces to every specified conditional
body: in cpp/tests/unit_tests/executor/bounce/buddyAllocatorTest.cpp:298-304,
brace both the if and else branches inside the for loop; in
cpp/tests/unit_tests/executor/bounce/creditSchedulerTest.cpp:282-286, brace the
s.heldCount(p) > 0 body, and also brace the repeated statement at lines 296-297.

In `@cpp/tests/unit_tests/executor/bounce/CMakeLists.txt`:
- Around line 76-94: Update the NIXL reactor-test conditional around
BOUNCE_REACTOR_SRCS to require both ZMQ_FOUND and the existence of
${CMAKE_BINARY_DIR}/_deps/cppzmq-src, matching the earlier test gate. Keep the
existing sources, include directories, and libraries unchanged.

In `@cpp/tests/unit_tests/executor/bounce/creditSchedulerTest.cpp`:
- Around line 288-300: In the loop around the served-count guard, assert that
m.owner is non-empty before dereferencing m.owner.begin(). Keep the existing
flow that reads the owner entry, frees it, and re-grants it unchanged after the
assertion.

In `@cpp/tests/unit_tests/executor/bounce/gatherScatterKernelTest.cpp`:
- Around line 146-153: Add the existing CUDA-device availability guard used by
the other bounce tests at the start of `GatherScatterKernel.ZeroBuffersIsNoop`,
before calling `cudaStreamCreate`, so the test skips when no device is present
while preserving its current assertions otherwise.

In `@tests/integration/test_lists/test-db/l0_h100.yml`:
- Around line 100-101: Reduce the pre_merge runtime of the test entry for
test_python_nixl_cache_transceiver_uses_cpp_bounce in l0_h100.yml by limiting it
to representative parameter cases or moving the full parameter coverage to a
separate suite. Preserve coverage of the C++ bounce v2 integration while
preventing all four potentially 180-second nested pytest cases from running
serially in pre_merge.

In `@tests/unittest/disaggregated/test_cache_transceiver_single_process.py`:
- Around line 1548-1563: Update the subprocess environment setup before the
child launch to remove UCX_NET_DEVICES and explicitly set UCX_TLS to
"^ib,gdr_copy" and TRTLLM_NIXL_NUM_THREADS to "1". Keep the existing
TRTLLM_NIXL_BOUNCE_* and logging configuration unchanged.

---

Nitpick comments:
In `@cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/BounceNvtx.h`:
- Around line 95-103: Add the printf format attribute to the variadic
BounceNvtxScope constructor, using format index 3 and argument index 4 because
of the implicit this parameter. Apply the equivalent attribute to the
NVTX_DISABLE variant as well, so both enabled and disabled builds validate
format strings and arguments consistently.

In
`@cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/ControlChannel.h`:
- Around line 72-76: Mark the virtual addPeer method in the ControlChannel
interface as [[nodiscard]] so callers must handle endpoint-registration
failures. Preserve its existing signature, behavior, and documentation.

In `@cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/ExecPool.cpp`:
- Around line 80-92: Add Allman-style braces around every single-statement if
body in the mCtxs cleanup loop, covering scratch, cubTemp, hostPinned, stream,
and event resource cleanup while preserving the existing conditions and calls.

In
`@cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/ZmqControlChannel.cpp`:
- Around line 158-189: Update the receive path in the surrounding function so
that when the second mRouter.recv call for bodyFrame fails, it drains the
remaining multipart frames of the current message before returning false. Reuse
the existing trailing-frame drain pattern based on the message’s more flag,
while preserving the current successful [identity, body] handling and failure
return behavior.

In `@cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/CMakeLists.txt`:
- Around line 53-74: Update the libzmq discovery and linking in the ZMQ_FOUND
block: call pkg_check_modules with IMPORTED_TARGET and link the
NIXL_WRAPPER_TARGET against PkgConfig::ZMQ so its include paths, library
directories, and libraries are propagated. Replace the hardcoded cppzmq path in
target_include_directories with the FetchContent-provided cppzmq_SOURCE_DIR
variable.

In `@cpp/tests/unit_tests/executor/bounce/bounceAgentE2ETest.cpp`:
- Around line 140-148: Replace the manual bounce environment cleanup around
setBounceEnv and clearBounceEnv with an RAII guard, preferably by reusing a
shared ScopedEnv from bounceConfigTest.cpp or adding a local guard whose
destructor calls clearBounceEnv. Ensure the guard covers the test setup and body
so GTEST_SKIP or ASSERT_* still restores all bounce variables; include
TRTLLM_NIXL_BOUNCE_USE_NIXL_NOTIFICATIONS in the environment setup or otherwise
keep the managed variable set symmetric.

In `@cpp/tests/unit_tests/executor/bounce/bounceArenaTest.cpp`:
- Around line 43-44: Add Allman-style braces around the skip guards controlled
by hasCuda() at both affected locations, and apply the same bracing style to
matching unbraced skip guards in the other bounce test files.
- Around line 43-60: Brace every single-statement if and for body using Allman
style across the bounce tests: in
cpp/tests/unit_tests/executor/bounce/bounceArenaTest.cpp lines 43-60,
cpp/tests/unit_tests/executor/bounce/bounceTransportFailureTest.cpp lines
107-112, 111-112, 127-128, 154-155, 158-159, 175-176, 179-180, 202-203, 216-217,
242-243, and 293-294, cpp/tests/unit_tests/executor/bounce/execPoolTest.cpp
lines 41-42 and 82-83, and
cpp/tests/unit_tests/executor/bounce/gatherScatterKernelTest.cpp lines 180-182,
255-257, and 262-267. Preserve the existing guard and loop behavior while adding
braces to all listed bodies.

In `@cpp/tests/unit_tests/executor/bounce/bounceTransferPlanTest.cpp`:
- Around line 94-131: In the BounceTransferPlan tests, replace all five
EXPECT_ANY_THROW assertions with EXPECT_THROW assertions specifying
tensorrt_llm::common::TllmException, including the cases in
DescLargerThanChunkThrows, MaxChunkSizeBytesAboveU32Throws,
MixedSourceDeviceIdsThrow, and MixedDestinationDeviceIdsThrow. Add a direct
include for tensorrt_llm/common/tllmException.h.

In `@cpp/tests/unit_tests/executor/bounce/bounceTransportTest.cpp`:
- Around line 100-105: Update the ConcurrentRequestsToSameReceiver test to
create one shared sender/receiver pair, submit two distinct transfers through
the same receiver before waiting on either result, then await and verify both
futures before cleanup. Do not use separate sequential runTransfer calls; reuse
the existing node-pair setup and transfer verification helpers.

In `@cpp/tests/unit_tests/executor/bounce/CMakeLists.txt`:
- Around line 57-69: Update the ZMQ discovery in the CMake configuration around
pkg_check_modules to request the imported target, then link
zmqControlChannelTest against PkgConfig::ZMQ. Remove the separate
${ZMQ_INCLUDE_DIRS} include path and ${ZMQ_LIBRARIES} link entry while retaining
the cppzmq include directory.

In `@cpp/tests/unit_tests/executor/bounce/gatherScatterKernelTest.cpp`:
- Around line 93-101: The gather/scatter test currently uses the sparse source
offsets for both source and packed-region addresses. Introduce a separate
tightly packed offset sequence for the arena, and use it for gatherDst and
scatterSrc while retaining off[i] for gatherSrc and scatterDst, so the test
exercises unequal source and destination layouts.

In `@tests/unittest/disaggregated/test_cache_transceiver_single_process.py`:
- Around line 1564-1567: Remove the redundant env.pop call for
TRTLLM_NIXL_BOUNCE_DISABLE_FABRIC_MEMORY while preserving the existing comment
and prefix-based environment cleanup loop.
- Around line 1501-1582: Extend test coverage around
test_python_nixl_cache_transceiver_uses_cpp_bounce by adding a scenario that
forces bounce initialization failure and verifies the standard NIXL path still
completes. Validate that the v2_mla_tp2 and v2_ctx_tp2_gen_pp2 parameters are
gated to a two-GPU-capable tier and fit its timeout budget, and verify the
shared run_transfer_test kv_cache_bounce_size_mb=0 change does not regress
existing callers. Supply or explicitly gate the declared bounce C++ test targets
so their assertions and execution coverage are verifiable.

🪄 Autofix

Fix all unresolved CodeRabbit comments on this PR:

• Push a commit to this branch (recommended)
• Create a new PR with the fixes

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 7f5ace51-4d35-4a57-aa6e-1c4f9c1027f0

📥 Commits

Reviewing files that changed from the base of the PR and between 07a5591 and be9dd02.

📒 Files selected for processing (50)

• cpp/CMakeLists.txt
• cpp/include/tensorrt_llm/executor/transferAgent.h
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/CMakeLists.txt
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/BounceArena.cpp
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/BounceArena.h
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/BounceConfig.h
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/BounceMessage.cpp
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/BounceMessage.h
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/BounceNvtx.h
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/BounceTransferPlan.cpp
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/BounceTransferPlan.h
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/BounceTransport.cpp
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/BounceTransport.h
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/BuddyAllocator.cpp
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/BuddyAllocator.h
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/ControlChannel.h
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/CreditScheduler.cpp
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/CreditScheduler.h
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/ExecPool.cpp
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/ExecPool.h
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/GatherScatterKernel.cu
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/GatherScatterKernel.h
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/NixlNotifControlChannel.cpp
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/NixlNotifControlChannel.h
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/NixlTransferEngine.cpp
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/NixlTransferEngine.h
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/TransferEngine.h
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/ZmqControlChannel.cpp
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/bounce/ZmqControlChannel.h
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/transferAgent.cpp
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/transferAgent.h
• cpp/tensorrt_llm/executor/cache_transmission/transferAgent.cpp
• cpp/tests/unit_tests/executor/CMakeLists.txt
• cpp/tests/unit_tests/executor/bounce/CMakeLists.txt
• cpp/tests/unit_tests/executor/bounce/bounceAgentE2ETest.cpp
• cpp/tests/unit_tests/executor/bounce/bounceArenaTest.cpp
• cpp/tests/unit_tests/executor/bounce/bounceConfigTest.cpp
• cpp/tests/unit_tests/executor/bounce/bounceMessageCodecTest.cpp
• cpp/tests/unit_tests/executor/bounce/bounceNixlE2ETest.cpp
• cpp/tests/unit_tests/executor/bounce/bounceTestNixlNode.h
• cpp/tests/unit_tests/executor/bounce/bounceTransferPlanTest.cpp
• cpp/tests/unit_tests/executor/bounce/bounceTransportFailureTest.cpp
• cpp/tests/unit_tests/executor/bounce/bounceTransportTest.cpp
• cpp/tests/unit_tests/executor/bounce/buddyAllocatorTest.cpp
• cpp/tests/unit_tests/executor/bounce/creditSchedulerTest.cpp
• cpp/tests/unit_tests/executor/bounce/execPoolTest.cpp
• cpp/tests/unit_tests/executor/bounce/gatherScatterKernelTest.cpp
• cpp/tests/unit_tests/executor/bounce/zmqControlChannelTest.cpp
• tests/integration/test_lists/test-db/l0_h100.yml
• tests/unittest/disaggregated/test_cache_transceiver_single_process.py

[coderabbitai]

🧹 Nitpick comments (3)

tests/unittest/disaggregated/test_cache_transceiver_single_process.py (3)

957-957: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Add the required return annotation to run_transfer_test.

This helper returns no value. Add -> None to the function signature.

Proposed fix

 def run_transfer_test(
     ...
     expect_cpp_bounce: bool = False,
-):
+) -> None:

As per coding guidelines: annotate every function and use None for procedures.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@tests/unittest/disaggregated/test_cache_transceiver_single_process.py` at
line 957, Update the run_transfer_test function signature to include the
required -> None return annotation, preserving its existing parameters and
behavior.

Source: Coding guidelines

---

1139-1142: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Use the bound property directly.

These agents are expected to expose bounce_enabled. Replace getattr(agent, "bounce_enabled", False) with agent.bounce_enabled so the test checks the declared binding contract directly.

Proposed fix

-            assert all(getattr(agent, "bounce_enabled", False) for agent in agents), (
+            assert all(agent.bounce_enabled for agent in agents), (

As per coding guidelines: avoid reflection when ordinary explicit code is sufficient.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@tests/unittest/disaggregated/test_cache_transceiver_single_process.py` around
lines 1139 - 1142, Update the agent validation in the test’s list comprehension
to access the declared bounce_enabled property directly as agent.bounce_enabled
instead of using getattr with a fallback, while preserving the existing
all-agents assertion and error message.

Source: Coding guidelines

---

1534-1559: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low value

Add the missing type annotation and remove unnecessary reflection.

• Annotate run_transfer_test with -> None.
• Access the required agent.bounce_enabled property directly.
• Coverage is sufficient. test_python_nixl_cache_transceiver_uses_cpp_bounce is listed in test-db/l0_h100.yml; no applicable QA entry exists.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@tests/unittest/disaggregated/test_cache_transceiver_single_process.py` around
lines 1534 - 1559, Update run_transfer_test to explicitly return None, and
replace any reflection-based access to the agent’s bounce state with direct
agent.bounce_enabled property access. Keep the existing test coverage and
behavior unchanged, including
test_python_nixl_cache_transceiver_uses_cpp_bounce.

Source: Path instructions

🤖 Prompt for all review comments with AI agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Nitpick comments:
In `@tests/unittest/disaggregated/test_cache_transceiver_single_process.py`:
- Line 957: Update the run_transfer_test function signature to include the
required -> None return annotation, preserving its existing parameters and
behavior.
- Around line 1139-1142: Update the agent validation in the test’s list
comprehension to access the declared bounce_enabled property directly as
agent.bounce_enabled instead of using getattr with a fallback, while preserving
the existing all-agents assertion and error message.
- Around line 1534-1559: Update run_transfer_test to explicitly return None, and
replace any reflection-based access to the agent’s bounce state with direct
agent.bounce_enabled property access. Keep the existing test coverage and
behavior unchanged, including
test_python_nixl_cache_transceiver_uses_cpp_bounce.

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 4dac4641-6c7e-476f-939f-8aab96a5d30d

📥 Commits

Reviewing files that changed from the base of the PR and between 99d1ef1 and 3dba868.

📒 Files selected for processing (5)

• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/agentBindings.cpp
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/transferAgent.cpp
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/transferAgent.h
• tensorrt_llm/_torch/disaggregation/nixl/_agent_cpp.py
• tests/unittest/disaggregated/test_cache_transceiver_single_process.py

🚧 Files skipped from review as they are similar to previous changes (2)

• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/transferAgent.h
• cpp/tensorrt_llm/executor/cache_transmission/nixl_utils/transferAgent.cpp