@@ -111,29 +111,39 @@ def _validate_trailing(Sources, Destinations, Operations, Expected, Count, Error
                        WIDTH: tl.constexpr, MemoryCapacity, PAGE_TOKENS: tl.constexpr,
                        LIVE: tl.constexpr):
     row = tl.program_id(0)
-    source, destination = tl.load(Sources + row), tl.load(Destinations + row)
+    source = tl.load(Sources + row)
     PAGES = tl.load(MemoryCapacity + 1)
-    endpoint = tl.maximum(0, tl.load(Expected + row) - 1)
+    endpoint = tl.load(Expected + row)
     slot = tl.load(RequestSlots + row)
     state = tl.load(StateRows + slot * 8, (slot >= 0) & (slot < LIVE), other=0)
     snapshot = source == state
     active = (row < tl.load(Count)) & (tl.load(Errors + row) == 0) & (tl.load(Operations + row) == 0)
-    tail = active & (endpoint % PAGE_TOKENS != 0)
-    page = tl.load(Tables + slot * WIDTH + endpoint // PAGE_TOKENS, tail, other=-1)
     others = tl.arange(0, LIVE)
-    other_source, other_dest = tl.load(Sources + others), tl.load(Destinations + others)
+    other_source = tl.load(Sources + others)
     other_slot = tl.load(RequestSlots + others)
     other_state = tl.load(StateRows + other_slot * 8, (other_slot >= 0) & (other_slot < LIVE), other=0)
     other_snapshot = other_source == other_state
-    other_endpoint = tl.maximum(0, tl.load(Expected + others) - 1)
+    other_endpoint = tl.load(Expected + others)
     other_active = (others < tl.load(Count)) & (others != row) & (tl.load(Operations + others) == 0)
-    other_active &= (other_slot >= 0) & (other_slot < LIVE) & (other_endpoint // PAGE_TOKENS < WIDTH)
-    other_active &= other_endpoint % PAGE_TOKENS != 0
-    other_page = tl.load(Tables + other_slot * WIDTH + other_endpoint // PAGE_TOKENS, other_active, other=-1)
-    conflict = tl.sum((other_active & (page == other_page) & (~snapshot | ~other_snapshot)).to(tl.int32), 0) > 0
+    other_active &= (other_slot >= 0) & (other_slot < LIVE)
     error = tl.load(Errors + row)
-    error = tl.where(tail & ((page < 0) | (page >= PAGES)), 6, error)
-    error = tl.where(tail & conflict, 4, error)
+    for part in tl.static_range(2):
+        logical = endpoint // PAGE_TOKENS + part - 1
+        tail = active & (endpoint > 1) & (logical >= 0)
+        if part == 1:
+            tail &= endpoint % PAGE_TOKENS >= 2
+        bounded = (logical >= 0) & (logical < WIDTH)
+        page = tl.load(Tables + slot * WIDTH + logical, tail & bounded, other=-1)
+        conflict = tl.full((), False, tl.int1)
+        for other_part in tl.static_range(2):
+            other_logical = other_endpoint // PAGE_TOKENS + other_part - 1
+            used = other_active & (other_endpoint > 1) & (other_logical >= 0) & (other_logical < WIDTH)
+            if other_part == 1:
+                used &= other_endpoint % PAGE_TOKENS >= 2
+            other_page = tl.load(Tables + other_slot * WIDTH + other_logical, used, other=-1)
+            conflict |= tl.sum((used & (page == other_page) & (~snapshot | ~other_snapshot)).to(tl.int32), 0) > 0
+        error = tl.where(tail & ((page < 0) | (page >= PAGES)), 6, error)
+        error = tl.where(tail & conflict, 4, error)
     tl.store(Errors + row, error)
 
 
@@ -156,18 +166,28 @@ def _copy_rows(Backing, Sources, Destinations, Operations, Count, Errors,
 def _copy_pages(Pointers, Backing, Tables, Sources, Destinations, Operations, Expected, Count, Errors, RequestSlots, StateRows,
                 ROW_BYTES: tl.constexpr, PAGE_BYTES: tl.constexpr, WIDTH: tl.constexpr,
                 PAGE_TOKENS: tl.constexpr, BLOCK: tl.constexpr, LIVE: tl.constexpr = 4,
-                ENDPOINT_SHIFT: tl.constexpr = 0):
+                DRAFT_PART: tl.constexpr = -1):
     row, region, tile = tl.program_id(0), tl.program_id(1), tl.program_id(2)
     source, destination = tl.load(Sources + row), tl.load(Destinations + row)
-    endpoint = tl.maximum(0, tl.load(Expected + row) + ENDPOINT_SHIFT)
+    endpoint = tl.load(Expected + row)
     slot = tl.load(RequestSlots + row)
     state = tl.load(StateRows + slot * 8, (slot >= 0) & (slot < LIVE), other=0)
     snapshot = source == state
     active = (row < tl.load(Count)) & (tl.load(Errors + row) == 0)
-    active &= (tl.load(Operations + row) == 0) & (endpoint % PAGE_TOKENS != 0)
-    page = tl.load(Tables + slot * WIDTH + endpoint // PAGE_TOKENS, active, other=0)
+    active &= tl.load(Operations + row) == 0
+    logical = endpoint // PAGE_TOKENS
+    if DRAFT_PART < 0:
+        active &= endpoint % PAGE_TOKENS != 0
+    else:
+        logical += DRAFT_PART - 1
+        active &= (endpoint > 1) & (logical >= 0)
+        if DRAFT_PART == 1:
+            active &= endpoint % PAGE_TOKENS >= 2
+    page = tl.load(Tables + slot * WIDTH + logical, active, other=0)
     cache = tl.load(Pointers + region * 2).to(tl.pointer_type(tl.uint8))
     offset = tl.load(Pointers + region * 2 + 1)
+    if DRAFT_PART >= 0:
+        offset += DRAFT_PART * PAGE_BYTES
     saved_slot = tl.where(snapshot, destination, source)
     index = tile * BLOCK + tl.arange(0, BLOCK)
     saved = Backing + saved_slot * ROW_BYTES + offset + index
@@ -185,6 +205,9 @@ class PreparedStateTransfer:
 
     ``pages`` contains (state region name, contiguous uint8 [pages, page_bytes]
     cache) pairs. Cache owners and pointer tables remain alive until graph drain.
+    ``trailing_pages`` uses a one-page-shifted table with a one-token lag. Each
+    saved region holds two pages: the preceding boundary page, then the current
+    one. Both may be private even when the preceding draft page is complete.
     These maintenance operations are captured directly. Tensor-parallel callers
     must combine validation failures across ranks between validate() and apply().
     A nonzero command error suppresses both row and external-page writes.
@@ -233,11 +256,11 @@ def __init__(self, pool: StateStorage, commands: StateCommands, tables, *, pages
             _check_tensor(trailing_tables, tables.shape, torch.int32, device, "trailing_page_tables")
             for name, cache in trailing_pages:
                 region = pool.layout.region(name)
-                if name in names or region.dtype != torch.uint8 or cache.shape != (page_count, region.nbytes):
-                    raise ValueError("trailing partial pages must match the model page capacity and state regions")
+                if name in names or region.dtype != torch.uint8 or cache.shape != (page_count, region.nbytes // 2) or region.nbytes % 2:
+                    raise ValueError("shifted draft pages require two saved pages per cache plane")
                 _check_tensor(cache, cache.shape, torch.uint8, device, name)
                 names.add(name)
-                trailing_groups[region.nbytes].append((cache.data_ptr(), region.offset))
+                trailing_groups[cache.shape[1]].append((cache.data_ptr(), region.offset))
         self.trailing_pointers = {size: torch.tensor(values, dtype=torch.int64, device=device)
                                  for size, values in trailing_groups.items()}
         self.pointers = {size: torch.tensor(values, dtype=torch.int64, device=device) for size, values in groups.items()}
@@ -253,7 +276,7 @@ def __init__(self, pool: StateStorage, commands: StateCommands, tables, *, pages
             for size, pointer in self.pointers.items()]
         self.page_args.extend((pointer, pool.backing, trailing_tables, c.source_slots, c.destination_slots, c.operation,
             c.expected_consumed, c.num_ops, c.errors, c.request_slots, state_rows, pool.layout.row_bytes, size, tables.shape[1], page_tokens,
-            1024, self.capacity, -1) for size, pointer in self.trailing_pointers.items())
+            1024, self.capacity, part) for size, pointer in self.trailing_pointers.items() for part in (0, 1))
         self.trailing_validation_args = None if not trailing_pages else (
             c.source_slots, c.destination_slots, c.operation, c.expected_consumed, c.num_ops, c.errors,
             trailing_tables, c.request_slots, state_rows, tables.shape[1], pool.memory_capacity, page_tokens, self.capacity)
 
@@ -40,8 +40,8 @@ def continuation_layout(layer_types, world_size, *, mtp=False):
             raise ValueError(f"unsupported canonical GLM attention type: {kind}")
     if mtp:
         regions.extend((("draft.index_tail", (2, 12, 128), torch.bfloat16),
-                        ("draft.latent_page", (4, 64, 528), torch.uint8),
-                        ("draft.index_page", (64, 132), torch.uint8)))
+                        ("draft.latent_page", (2, 4, 64, 528), torch.uint8),
+                        ("draft.index_page", (2, 64, 132), torch.uint8)))
     regions.extend((("last_hidden", (4096,), torch.bfloat16), ("consumed_tokens", (1,), torch.int64)))
     return StateLayout.from_specs(regions)
 
 
@@ -88,9 +88,9 @@ def usage(prepared):
     ranks=MappingProxyType({
         # DGX Station can hold GLM on one GPU; TP1 must remain a supported rank count.
         # One rank holds every head: twice TP2's recurrent rows less the replicated part.
-        1: _costs(161_123_584, _GRAPH_BYTES, ranks=1, speculating=True),
-        2: _costs(81_464_576, _GRAPH_BYTES, ranks=2, speculating=True),
-        4: _costs(41_635_072, _GRAPH_BYTES, ranks=4, speculating=True),
+        1: _costs(161_267_200, _GRAPH_BYTES, ranks=1, speculating=True),
+        2: _costs(81_608_192, _GRAPH_BYTES, ranks=2, speculating=True),
+        4: _costs(41_778_688, _GRAPH_BYTES, ranks=4, speculating=True),
     }),
     plain_ranks=MappingProxyType({
         1: _costs(149_230_592, _GRAPH_BYTES, ranks=1, speculating=False),
 
@@ -64,7 +64,7 @@ def _run(rank, world, rendezvous=None):
                         pages=pool.partial_pages(), page_tokens=256, state_rows=inputs.request_state_rows, requests=requests)
                     slots, epochs, count = (inputs.row_slot_ids, inputs.row_generations, inputs.live_rows) if mode == "snapshot" else (
                         admission.admission_slots, admission.admission_generations, admission.admission_count)
-                    local = LocalMaintenance(mode, requests, transfer, slots, epochs, count, records, frame, sequence,
+                    local = LocalMaintenance(mode, requests, transfer, slots, epochs, count, completion_outputs(records), frame, sequence,
                                               world_size=1, admission_buffers=admission)
                     local.prepare()
                     def call(local=local):
 
@@ -148,12 +148,16 @@ def test_rejected_commands_preserve_every_state_and_cache_byte(transfers, fault)
         assert torch.equal(cache, caches[name])
 
 
-@pytest.mark.parametrize("endpoint", [1, 4, 255, 256, 257, 512])
+@pytest.mark.parametrize("endpoint", [1, 2, 4, 255, 256, 257, 258, 511, 512, 513])
 def test_mtp_trailing_partial_page_restores_at_target_page_boundaries(endpoint):
     pool = ContinuationPool.allocate(("linear_attention", "indexed_attention"), 4, 8, 6, "cuda:0", mtp=True)
     commands = StateCommands.allocate("cuda:0")
     tables = torch.tensor([[0, 1, 2], [3, 4, 5], [-1, -1, -1], [-1, -1, -1]], device="cuda:0", dtype=torch.int32)
-    draft_tables = tables.clone()
+    # Helios shares complete target pages and shifts the draft view by one.
+    shared = endpoint // 256
+    tables[1, :shared] = tables[0, :shared]
+    draft_tables = torch.full_like(tables, -1)
+    draft_tables[:, :-1] = tables[:, 1:]
     requests = DeviceRequestState.allocate("cuda:0", vocab_size=128, consumed_tokens=pool.consumed_tokens[1:5, 0])
     state_rows = torch.arange(1, 5, device="cuda:0")[:, None].expand(4, 8).contiguous()
     operation = PreparedStateTransfer(pool, commands, tables, pages=pool.partial_pages(), page_tokens=256, state_rows=state_rows, requests=requests,
@@ -163,7 +167,11 @@ def test_mtp_trailing_partial_page_restores_at_target_page_boundaries(endpoint):
     pool.consumed_tokens[1] = endpoint
     for _, cache in (*pool.partial_pages(), *pool.draft_pages()):
         cache.random_(0, 256)
-    expected = {name: cache[(endpoint - 1) // 256].clone() for name, cache in pool.draft_pages()}
+    logical = [endpoint // 256 - 1] if endpoint > 1 and endpoint // 256 > 0 else []
+    if endpoint > 1 and endpoint % 256 >= 2:
+        logical.append(endpoint // 256)
+    expected = {(name, page): cache[draft_tables[0, page]].clone()
+                for name, cache in pool.draft_pages() for page in logical}
     graph = torch.cuda.CUDAGraph()
     _commands(commands, [(1, 5, 0, endpoint)])
     with torch.cuda.graph(graph):
@@ -177,12 +185,50 @@ def test_mtp_trailing_partial_page_restores_at_target_page_boundaries(endpoint):
     torch.cuda.synchronize()
     assert not commands.errors.count_nonzero()
     for name, cache in pool.draft_pages():
-        if (endpoint - 1) % 256:
-            torch.testing.assert_close(cache[3 + (endpoint - 1) // 256], expected[name], rtol=0, atol=0)
-    if (endpoint - 1) % 256:
+        for page in logical:
+            torch.testing.assert_close(cache[draft_tables[1, page]], expected[name, page], rtol=0, atol=0)
+    if logical:
         before = pool.backing.clone()
-        draft_tables[1, (endpoint - 1) // 256] = -1
+        draft_tables[1, logical[-1]] = -1
         graph.replay()
         torch.cuda.synchronize()
         assert commands.errors[0] == 6
         torch.testing.assert_close(pool.backing, before, rtol=0, atol=0)
+
+
+@pytest.mark.parametrize("snapshot", [False, True])
+def test_shifted_draft_boundary_aliases_are_validated_across_both_saved_pages(snapshot):
+    pool = ContinuationPool.allocate(("indexed_attention",), 4, 8, 6, "cuda:0", mtp=True)
+    commands = StateCommands.allocate("cuda:0")
+    tables = torch.tensor([[0, 1, 2], [3, 4, 5], [-1, -1, -1], [-1, -1, -1]],
+                          device="cuda", dtype=torch.int32)
+    # A preceding full page in lane zero aliases lane one's current partial page.
+    draft_tables = torch.tensor([[4, 2, -1], [4, 5, -1], [-1, -1, -1], [-1, -1, -1]],
+                                device="cuda", dtype=torch.int32)
+    requests = DeviceRequestState.allocate("cuda:0", vocab_size=128,
+                                          consumed_tokens=pool.consumed_tokens[1:5, 0])
+    state_rows = torch.arange(1, 5, device="cuda")[:, None].expand(4, 8).contiguous()
+    operation = PreparedStateTransfer(pool, commands, tables, pages=pool.partial_pages(),
+        page_tokens=256, state_rows=state_rows, requests=requests,
+        trailing_pages=pool.draft_pages(), trailing_tables=draft_tables)
+    operation.prepare()
+    pool.backing.random_(0, 256)
+    for _, cache in (*pool.partial_pages(), *pool.draft_pages()):
+        cache.random_(0, 256)
+    for live, saved, endpoint in ((1, 5, 257), (2, 6, 2)):
+        pool.consumed_tokens[live if snapshot else saved] = endpoint
+    rows = [(1, 5, 0, 257), (2, 6, 0, 2)] if snapshot else [(5, 1, 0, 257), (6, 2, 0, 2)]
+    _commands(commands, rows)
+    before = pool.backing.clone()
+    caches = {name: cache.clone() for name, cache in (*pool.partial_pages(), *pool.draft_pages())}
+    graph = torch.cuda.CUDAGraph()
+    with torch.cuda.graph(graph):
+        operation.validate()
+        operation.apply()
+    graph.replay()
+    torch.cuda.synchronize()
+    assert commands.errors[:2].tolist() == ([0, 0] if snapshot else [4, 4])
+    if not snapshot:
+        torch.testing.assert_close(pool.backing, before, rtol=0, atol=0)
+    for name, cache in (*pool.partial_pages(), *pool.draft_pages()):
+        torch.testing.assert_close(cache, caches[name], rtol=0, atol=0)