kvcache-ai/Mooncake

Fork 0

mirror of https://github.com/kvcache-ai/Mooncake.git synced 2026-11-05 01:17:55 +00:00

[GH-ISSUE #4592] [Bug]: TENT MnnvlTransport reuses an imported mapping after the peer re-registers the same address #1170

New issue

Closed

opened 2026-10-29 21:23:33 +00:00 by bsakaguchi · 2 comments

bsakaguchi commented

2026-10-29 21:23:33 +00:00

Owner

Originally created by @SongOf on GitHub (Oct 12, 2026).
Original GitHub issue: https://github.com/kvcache-ai/Mooncake/issues/4592

Bug Report

Version: origin/main 559f83cb73eb870d208a8d4f333b1a24ee992f26
Environment: TENT with the MNNVL transport, Fabric or POSIX-FD handle type. Reproducible without a second node; requires two GPUs with P2P or one GPU with two processes.

Description

MnnvlTransport caches imported mappings keyed only by (target_id, remote base address). When the peer publishes a new allocation at the same virtual address (process restart, or free + re-allocate inside the same process) the consumer keeps using the mapping imported from the old mnnvl_handle. Transfers complete successfully against the old allocation, which the consumer's own mapping keeps alive.

Three contributing points in the code:

The cache hit never reads the current descriptor. relocateSharedMemoryAddress resolves from the map before consulting SegmentManager, and the map entry stores only mnnvl_addr, length, cuda_id — no handle identity:
github.com/kvcache-ai/Mooncake@559f83cb73/mooncake-transfer-engine/tent/include/tent/transport/mnnvl/mnnvl_transport.h (L113-L117)
github.com/kvcache-ai/Mooncake@559f83cb73/mooncake-transfer-engine/tent/src/transport/mnnvl/mnnvl_transport.cpp (L799)
Hits are served from a thread_local copy of the map that is only refilled when empty or after a miss, so invalidating the global map does not reach threads that already hold a copy:
github.com/kvcache-ai/Mooncake@559f83cb73/mooncake-transfer-engine/tent/src/transport/mnnvl/mnnvl_transport.cpp (L769-L778)
github.com/kvcache-ai/Mooncake@559f83cb73/mooncake-transfer-engine/tent/src/transport/mnnvl/mnnvl_transport.cpp (L857)
Nothing ever retires a mapping. uninstall() clears the map without cuMemUnmap / cuMemRelease / cuMemAddressFree, and there is no path for a desc update to retire entries:
github.com/kvcache-ai/Mooncake@559f83cb73/mooncake-transfer-engine/tent/src/transport/mnnvl/mnnvl_transport.cpp (L242-L247)

SegmentManager's push invalidation (NotifySegmentUpdated) bumps version_ for the descriptor cache only; transport mapping caches are not notified. The subscriber list is held in the provider process, so a provider restart sends no notification and consumers rely on TENT_SEGMENT_DESC_TTL_MS (1 hour).

Steps to reproduce

Process P0: allocateLocalMemory + registerLocalMemory a buffer B0 at address A, write pattern 0xA0, publish segment S.
Process C: open S, submitTransfer READ of A → observes 0xA0. (Mapping for (S, A) is now cached in C.)
P0 frees B0 and allocates again; with the same allocation sequence the new buffer B1 lands at A. Register B1, write pattern 0xB1, publish S again. (Equivalently: kill P0, start P1 with the same segment name; it gets A again.)
Process C: submitTransfer READ of A.

Expected: C reads 0xB1, or the transfer fails with a clear error.
Actual: C reads 0xA0 and the transfer reports success. With the push notification delivered, the descriptor cache in C is refreshed but the result is unchanged because the mapping cache is keyed by address only and the thread_local copy is never invalidated. Without a notification (restart case) the same result persists until the 1h TTL.

Side effect: P0's B0 allocation is not released while C holds the mapping, so after a few restart cycles the provider can fail cuMemCreate with OOM.

Correction after reproduction (2026-10-12)

Reproduced on b198a191 with POSIX-FD handles, see the comment below. Two things in the description above need adjusting:

The deterministic trigger is in-process re-allocation, not provider restart. When the provider frees and re-allocates inside the same process, the VMM address reservation lands on the same VA every time, the push notification refreshes the consumer's descriptor cache, and the consumer still reads the old data because the mapping cache is keyed by address only. This isolates the mapping-cache defect from descriptor freshness.
Provider restart corrupts data only when the new process gets the same VA. With ASLR enabled, the new provider usually reserves a different address, so the consumer takes the NeedsRefreshCache path and imports the new buffer correctly; the old mapping is then only a leak that keeps the old allocation alive. The restart case in the repro needs setarch -R (or randomize_va_space=0) to hit; with ASLR off it is deterministic.

The reproduction also exposed an independent blocker: POSIX-FD import passes &fd instead of the fd value to cuMemImportFromShareableHandle, so FD mode cannot import anything on main. Fixed in #4597, which is a prerequisite for FD-mode tests of this issue.

Proposed fix (no wire-format change)

Store mnnvl_handle (and length) in OpenedMnnvlEntry; a hit requires base address and handle to match.
Remove the thread_local copy, or tag it with SegmentManager::version_ and drop it on change.
On mismatch, return NeedsRefreshCache so withCachedSegment refreshes once and retries; still mismatched → InvalidEntry. Never fall back to the old mapping.
Add an in-flight-safe retire primitive (mark stale, release after the stream drains, then cuMemUnmap → cuMemRelease → cuMemAddressFree; close the imported fd in POSIX-FD mode) and use it from uninstall() and from a SegmentManager observer hook on desc replacement.
Optionally carry a per-registration id in transport_attrs[MNNVL], which needs no BufferDesc field change.

Related: #2832 (RFC covering this and the legacy NvlinkTransport; this issue scopes the TENT part that does not need the wire-format decision), #3581 (same pattern for TENT SHM), #4108.

Before submitting...

Ensure you searched for relevant issues and read the documentation

Originally created by @SongOf on GitHub (Oct 12, 2026). Original GitHub issue: https://github.com/kvcache-ai/Mooncake/issues/4592 ## Bug Report **Version:** origin/main `559f83cb73eb870d208a8d4f333b1a24ee992f26` **Environment:** TENT with the MNNVL transport, Fabric or POSIX-FD handle type. Reproducible without a second node; requires two GPUs with P2P or one GPU with two processes. ### Description `MnnvlTransport` caches imported mappings keyed only by `(target_id, remote base address)`. When the peer publishes a new allocation at the same virtual address (process restart, or free + re-allocate inside the same process) the consumer keeps using the mapping imported from the old `mnnvl_handle`. Transfers complete successfully against the old allocation, which the consumer's own mapping keeps alive. Three contributing points in the code: 1. The cache hit never reads the current descriptor. `relocateSharedMemoryAddress` resolves from the map before consulting `SegmentManager`, and the map entry stores only `mnnvl_addr`, `length`, `cuda_id` — no handle identity: https://github.com/kvcache-ai/Mooncake/blob/559f83cb73eb870d208a8d4f333b1a24ee992f26/mooncake-transfer-engine/tent/include/tent/transport/mnnvl/mnnvl_transport.h#L113-L117 https://github.com/kvcache-ai/Mooncake/blob/559f83cb73eb870d208a8d4f333b1a24ee992f26/mooncake-transfer-engine/tent/src/transport/mnnvl/mnnvl_transport.cpp#L799 2. Hits are served from a `thread_local` copy of the map that is only refilled when empty or after a miss, so invalidating the global map does not reach threads that already hold a copy: https://github.com/kvcache-ai/Mooncake/blob/559f83cb73eb870d208a8d4f333b1a24ee992f26/mooncake-transfer-engine/tent/src/transport/mnnvl/mnnvl_transport.cpp#L769-L778 https://github.com/kvcache-ai/Mooncake/blob/559f83cb73eb870d208a8d4f333b1a24ee992f26/mooncake-transfer-engine/tent/src/transport/mnnvl/mnnvl_transport.cpp#L857 3. Nothing ever retires a mapping. `uninstall()` clears the map without `cuMemUnmap` / `cuMemRelease` / `cuMemAddressFree`, and there is no path for a desc update to retire entries: https://github.com/kvcache-ai/Mooncake/blob/559f83cb73eb870d208a8d4f333b1a24ee992f26/mooncake-transfer-engine/tent/src/transport/mnnvl/mnnvl_transport.cpp#L242-L247 `SegmentManager`'s push invalidation (`NotifySegmentUpdated`) bumps `version_` for the descriptor cache only; transport mapping caches are not notified. The subscriber list is held in the provider process, so a provider restart sends no notification and consumers rely on `TENT_SEGMENT_DESC_TTL_MS` (1 hour). ### Steps to reproduce 1. Process P0: `allocateLocalMemory` + `registerLocalMemory` a buffer `B0` at address `A`, write pattern `0xA0`, publish segment `S`. 2. Process C: open `S`, `submitTransfer` READ of `A` → observes `0xA0`. (Mapping for `(S, A)` is now cached in C.) 3. P0 frees `B0` and allocates again; with the same allocation sequence the new buffer `B1` lands at `A`. Register `B1`, write pattern `0xB1`, publish `S` again. (Equivalently: kill P0, start P1 with the same segment name; it gets `A` again.) 4. Process C: `submitTransfer` READ of `A`. **Expected:** C reads `0xB1`, or the transfer fails with a clear error. **Actual:** C reads `0xA0` and the transfer reports success. With the push notification delivered, the descriptor cache in C is refreshed but the result is unchanged because the mapping cache is keyed by address only and the `thread_local` copy is never invalidated. Without a notification (restart case) the same result persists until the 1h TTL. Side effect: P0's `B0` allocation is not released while C holds the mapping, so after a few restart cycles the provider can fail `cuMemCreate` with OOM. ### Correction after reproduction (2026-10-12) Reproduced on `b198a191` with POSIX-FD handles, see the comment below. Two things in the description above need adjusting: - **The deterministic trigger is in-process re-allocation**, not provider restart. When the provider frees and re-allocates inside the same process, the VMM address reservation lands on the same VA every time, the push notification refreshes the consumer's descriptor cache, and the consumer still reads the old data because the mapping cache is keyed by address only. This isolates the mapping-cache defect from descriptor freshness. - **Provider restart corrupts data only when the new process gets the same VA.** With ASLR enabled, the new provider usually reserves a different address, so the consumer takes the `NeedsRefreshCache` path and imports the new buffer correctly; the old mapping is then only a leak that keeps the old allocation alive. The restart case in the repro needs `setarch -R` (or `randomize_va_space=0`) to hit; with ASLR off it is deterministic. The reproduction also exposed an independent blocker: POSIX-FD import passes `&fd` instead of the fd value to `cuMemImportFromShareableHandle`, so FD mode cannot import anything on main. Fixed in #4597, which is a prerequisite for FD-mode tests of this issue. ### Proposed fix (no wire-format change) - Store `mnnvl_handle` (and length) in `OpenedMnnvlEntry`; a hit requires base address and handle to match. - Remove the `thread_local` copy, or tag it with `SegmentManager::version_` and drop it on change. - On mismatch, return `NeedsRefreshCache` so `withCachedSegment` refreshes once and retries; still mismatched → `InvalidEntry`. Never fall back to the old mapping. - Add an in-flight-safe retire primitive (mark stale, release after the stream drains, then `cuMemUnmap` → `cuMemRelease` → `cuMemAddressFree`; close the imported fd in POSIX-FD mode) and use it from `uninstall()` and from a `SegmentManager` observer hook on desc replacement. - Optionally carry a per-registration id in `transport_attrs[MNNVL]`, which needs no `BufferDesc` field change. Related: #2832 (RFC covering this and the legacy `NvlinkTransport`; this issue scopes the TENT part that does not need the wire-format decision), #3581 (same pattern for TENT SHM), #4108. ## Before submitting... - [x] Ensure you searched for relevant issues and read the documentation

bsakaguchi closed this issue

2026-10-29 21:23:33 +00:00

bsakaguchi commented

2026-10-29 21:23:34 +00:00

Author

Owner

@github-actions[bot] commented on GitHub (Oct 12, 2026):

Thanks for opening this issue, @SongOf!

Field	Value
Issue	#4592
GitHub user ID	`46475785`
Reporter	@SongOf

A maintainer will triage this when possible. To help us respond faster, please include:

Mooncake version or commit SHA
Environment (OS, CUDA/driver, RDMA stack if relevant)
Steps to reproduce and expected vs. actual behavior

Useful links: Documentation · Contributing guide

This message was posted automatically by the issue bot.

@github-actions[bot] commented on GitHub (Oct 12, 2026): Thanks for opening this issue, @SongOf! | Field | Value | |-------|-------| | **Issue** | #4592 | | **GitHub user ID** | `46475785` | | **Reporter** | @SongOf | A maintainer will triage this when possible. To help us respond faster, please include: - Mooncake version or commit SHA - Environment (OS, CUDA/driver, RDMA stack if relevant) - Steps to reproduce and expected vs. actual behavior Useful links: [Documentation](https://kvcache-ai.github.io/Mooncake/) · [Contributing guide](https://github.com/kvcache-ai/Mooncake/blob/main/CONTRIBUTING.md) > This message was posted automatically by the issue bot.

bsakaguchi commented

2026-10-29 21:23:34 +00:00

Author

Owner

@SongOf commented on GitHub (Oct 12, 2026):

Reproduced on current main

Reproduced on b198a191 (main) with POSIX-FD handles, two processes sharing one GPU. Both the in-process re-allocation case and the provider-restart case return the old allocation's data, and the transfer reports COMPLETED.

Environment

2× NVIDIA L20 (PCIe, no P2P between them), driver 580.126.09, CUDA 12.8, Ubuntu 22.04. This is not MNNVL hardware; it is used as a single-node repro box with POSIX-FD handles, and only GPU 0 is used.
Build: -DUSE_CUDA=ON -DUSE_TENT=ON
Run: MC_ENABLE_MNNVL=1 CUDA_VISIBLE_DEVICES=0, config metadata_type=p2p, transports/nvlink/handle_type=posix_fd, TCP/RDMA/io_uring disabled so MNNVL is used
Only the public TransferEngine API is used. The provider is fork+exec'd, so neither process forks after CUDA init.

Results

Scenario	Peer address after change	2nd READ by consumer	Status
realloc: P0 `unregister`+`free` B0 (0xA0), `allocate`+`register` B1 (0xB1)	same VA	0xA0	`COMPLETED`
restart: P0 exits, P1 starts on the same `ip:port`, fills 0xB1 (`setarch -R` so the VA repeats)	same VA	0xA0	`COMPLETED` (2/2 runs)

=== realloc ===
[provider 11507] B0 at 0x7fef6f800000 filled 0xa0, segment 127.0.0.1:17506
[provider 11507] freed B0, B1 at 0x7fef6f800000 (same VA: yes) filled 0xb1
step 2: READ 0x7fef6f800000 -> 0xa0 status=COMPLETED (expect 0xa0)
step 3: provider re-published 0x7fef6f800000 (same address: yes)
step 4: READ 0x7fef6f800000 -> 0xa0 status=COMPLETED (expect 0xb1 or a failure)
RESULT[realloc]: BUG REPRODUCED - stale mapping served old allocation

=== restart (setarch -R) ===
[provider 11660] B0 at 0x7fffd5800000 filled 0xa0, segment 127.0.0.1:17659
[provider 11706] B0 at 0x7fffd5800000 filled 0xb1, segment 127.0.0.1:17659
step 2: READ 0x7fffd5800000 -> 0xa0 status=COMPLETED (expect 0xa0)
step 3: provider 11660 exited; starting a new one on 127.0.0.1:17659
step 3: provider re-published 0x7fffd5800000 (same address: yes)
step 4: READ 0x7fffd5800000 -> 0xa0 status=COMPLETED (expect 0xb1 or a failure)
RESULT[restart]: BUG REPRODUCED - stale mapping served old allocation

In the restart case the consumer still reads 0xA0 after P0 has exited. This confirms the side effect described above: the consumer's imported mapping keeps the old physical allocation alive. With address-space randomization enabled, P1 usually gets a different VA and the restart case does not trigger. The realloc case hits the same VA every time.

Separate blocker: POSIX-FD import is broken on main

To reach this code path I had to locally patch an unrelated bug. relocateSharedMemoryAddress passes the fd by address:

github.com/kvcache-ai/Mooncake@b198a191/mooncake-transfer-engine/tent/src/transport/mnnvl/mnnvl_transport.cpp (L818)

CHECK_CU(cuMemImportFromShareableHandle(&handle, &local_fd, handle_type_));

For CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR the driver expects the fd value itself cast to void *. Passing &local_fd fails with CUDA_ERROR_UNKNOWN (999), and the LOG(FATAL) in submitTransferTasks then aborts the process. A standalone CUDA program confirms it: (void *)(uintptr_t)fd imports fine, &fd returns 999. So on main, MNNVL in POSIX-FD mode cannot import any peer buffer. The results above were obtained with only this line changed to (void *)(uintptr_t)local_fd. Fixed in #4597.

Notes on Fabric mode

Fabric mode could not be tested on this box. CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_FABRIC_SUPPORTED reports 1, so supportFabricMem() passes and MNNVL installs. However, cuMemCreate with requestedHandleTypes = FABRIC returns CUDA_ERROR_NOT_PERMITTED (800): there is no IMEX channel and no NVLink fabric on L20. Because the cache logic does not depend on the handle type, the same behaviour is expected with Fabric handles on an IMEX-enabled system, but that still needs confirming on GB200-class hardware.

The repro pins CUDA_VISIBLE_DEVICES=0 because the two L20s have no P2P path and MNNVL grants every visible device access to a mapping. That only matters on PCIe boxes like this one, not on NVLink-domain hardware, and is unrelated to this issue.

Repro program (mnnvl_reregister_repro.cpp)

Build it as an extra target next to the TENT tests, linked against tent_link_group, then run:

CUDA_VISIBLE_DEVICES=0 MC_ENABLE_MNNVL=1 ./tent_mnnvl_reregister_repro consumer realloc
CUDA_VISIBLE_DEVICES=0 MC_ENABLE_MNNVL=1 setarch $(uname -m) -R ./tent_mnnvl_reregister_repro consumer restart

// Repro for kvcache-ai/Mooncake#4592: TENT MnnvlTransport keeps using an
// imported mapping after the peer re-registers a new allocation at the same
// virtual address.
//
//   repro consumer <scenario>      scenario = realloc | restart
//   repro provider <port> <ctl_rd> <ack_wr> <mode> <pattern1> [pattern2]
//
// Run with MC_ENABLE_MNNVL=1. Providers are fork+exec'd so no process forks
// after CUDA initialization.

#include <cuda_runtime.h>
#include <signal.h>
#include <sys/wait.h>
#include <unistd.h>

#include <chrono>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <memory>
#include <string>
#include <thread>
#include <vector>

#include "tent/common/config.h"
#include "tent/transfer_engine.h"

using namespace mooncake::tent;

static constexpr size_t kLen = 4 << 20;

#define DIE(...)                               \
    do {                                       \
        fprintf(stderr, "[%d] FATAL: ", getpid()); \
        fprintf(stderr, __VA_ARGS__);          \
        fprintf(stderr, "\n");                 \
        _exit(10);                             \
    } while (0)

#define CHK(expr)                                                     \
    do {                                                              \
        auto _s = (expr);                                             \
        if (!_s.ok()) DIE("%s -> %s", #expr, _s.ToString().c_str()); \
    } while (0)

static std::shared_ptr<Config> makeConfig(int port) {
    auto conf = std::make_shared<Config>();
    conf->set("metadata_type", std::string("p2p"));
    conf->set("rpc_server_hostname", std::string("127.0.0.1"));
    if (port) conf->set("rpc_server_port", port);
    conf->set("transports/tcp/enable", false);
    conf->set("transports/rdma/enable", false);
    conf->set("transports/io_uring/enable", false);
    // Fabric export needs IMEX; POSIX FD works on any VMM-capable GPU.
    conf->set("transports/nvlink/handle_type", std::string("posix_fd"));
    conf->set("log_level", std::string("warning"));
    return conf;
}

static void writeLine(int fd, const std::string &s) {
    std::string l = s + "\n";
    if (write(fd, l.data(), l.size()) != (ssize_t)l.size()) DIE("write");
}

static std::string readLine(int fd) {
    std::string s;
    char c;
    while (true) {
        ssize_t n = read(fd, &c, 1);
        if (n <= 0) DIE("peer closed pipe");
        if (c == '\n') return s;
        s += c;
    }
}

// ---------------------------------------------------------------- provider
struct ProviderBuf {
    void *addr = nullptr;
};

static void allocFill(TransferEngine &te, ProviderBuf &b, int pattern) {
    CHK(te.allocateLocalMemory(&b.addr, kLen, "cuda:0"));
    CHK(te.registerLocalMemory(b.addr, kLen));
    if (cudaMemset(b.addr, pattern, kLen) != cudaSuccess) DIE("cudaMemset");
    cudaDeviceSynchronize();
}

static int providerMain(int argc, char **argv) {
    int port = atoi(argv[2]);
    int ctl = atoi(argv[3]), ack = atoi(argv[4]);
    std::string mode = argv[5];
    int p1 = strtol(argv[6], nullptr, 0);
    cudaSetDevice(0);
    TransferEngine te(makeConfig(port));
    if (!te.available()) DIE("provider engine not available");

    ProviderBuf b;
    allocFill(te, b, p1);
    fprintf(stderr, "[provider %d] B0 at %p filled 0x%02x, segment %s\n",
            getpid(), b.addr, p1, te.getSegmentName().c_str());
    writeLine(ack, te.getSegmentName() + " " +
                       std::to_string((uint64_t)b.addr));

    if (mode == "realloc") {
        int p2 = strtol(argv[7], nullptr, 0);
        readLine(ctl);  // "realloc"
        CHK(te.unregisterLocalMemory(b.addr, kLen));
        CHK(te.freeLocalMemory(b.addr));
        void *old = b.addr;
        allocFill(te, b, p2);
        fprintf(stderr,
                "[provider %d] freed B0, B1 at %p (same VA: %s) filled 0x%02x\n",
                getpid(), b.addr, old == b.addr ? "yes" : "no", p2);
        writeLine(ack, std::to_string((uint64_t)b.addr));
    }
    readLine(ctl);  // "exit"
    return 0;
}

// ---------------------------------------------------------------- consumer
struct Provider {
    pid_t pid;
    int ctl_wr, ack_rd;
};

static Provider spawnProvider(const char *self, int port,
                              const std::vector<std::string> &extra) {
    int ctl[2], ack[2];
    if (pipe(ctl) || pipe(ack)) DIE("pipe");
    pid_t pid = fork();
    if (pid == 0) {
        close(ctl[1]);
        close(ack[0]);
        std::vector<std::string> args = {self, "provider", std::to_string(port),
                                         std::to_string(ctl[0]),
                                         std::to_string(ack[1])};
        args.insert(args.end(), extra.begin(), extra.end());
        std::vector<char *> av;
        for (auto &a : args) av.push_back(const_cast<char *>(a.c_str()));
        av.push_back(nullptr);
        execv(self, av.data());
        _exit(127);
    }
    close(ctl[0]);
    close(ack[1]);
    return {pid, ctl[1], ack[0]};
}

static void stopProvider(Provider &p) {
    writeLine(p.ctl_wr, "exit");
    int st = 0;
    waitpid(p.pid, &st, 0);
    close(p.ctl_wr);
    close(p.ack_rd);
}

// READ kLen bytes from (seg, remote) into local device buffer; returns the
// first byte observed and the transfer status.
static int readRemote(TransferEngine &te, SegmentID seg, uint64_t remote,
                      void *local, TransferStatusEnum &result) {
    cudaMemset(local, 0, kLen);
    cudaDeviceSynchronize();
    auto batch = te.allocateBatch(1);
    Request r{};
    r.opcode = Request::READ;
    r.source = local;
    r.target_id = seg;
    r.target_offset = remote;
    r.length = kLen;
    CHK(te.submitTransfer(batch, {r}));
    TransferStatus st{};
    auto deadline = std::chrono::steady_clock::now() + std::chrono::seconds(10);
    do {
        CHK(te.getTransferStatus(batch, 0, st));
    } while (st.s == PENDING &&
             std::chrono::steady_clock::now() < deadline);
    te.freeBatch(batch);
    result = st.s;
    std::vector<unsigned char> host(kLen);
    cudaMemcpy(host.data(), local, kLen, cudaMemcpyDeviceToHost);
    // Make sure the whole buffer is uniform.
    for (size_t i = 1; i < kLen; ++i)
        if (host[i] != host[0]) return -1;
    return host[0];
}

static const char *statusName(TransferStatusEnum s) {
    switch (s) {
        case COMPLETED: return "COMPLETED";
        case FAILED: return "FAILED";
        case PENDING: return "PENDING(timeout)";
        default: return "OTHER";
    }
}

static int consumerMain(const char *self, const std::string &scenario) {
    const int port = 17000 + (getpid() % 1000);
    Provider p;
    std::vector<std::string> ack;
    if (scenario == "realloc")
        p = spawnProvider(self, port, {"realloc", "0xA0", "0xB1"});
    else
        p = spawnProvider(self, port, {"single", "0xA0"});

    auto line = readLine(p.ack_rd);
    std::string seg_name = line.substr(0, line.find(' '));
    uint64_t A = std::stoull(line.substr(line.find(' ') + 1));

    // CUDA is first touched here, after all forks except `restart`'s second
    // spawn, which exec()s immediately.
    cudaSetDevice(0);
    TransferEngine te(makeConfig(0));
    if (!te.available()) DIE("consumer engine not available");
    void *local = nullptr;
    CHK(te.allocateLocalMemory(&local, kLen, "cuda:0"));
    CHK(te.registerLocalMemory(local, kLen));

    SegmentID seg;
    CHK(te.openSegment(seg, seg_name));

    TransferStatusEnum s;
    int v = readRemote(te, seg, A, local, s);
    printf("step 2: READ %#lx -> 0x%02x status=%s (expect 0xa0)\n", A, v,
           statusName(s));

    uint64_t A2;
    if (scenario == "realloc") {
        writeLine(p.ctl_wr, "realloc");
        A2 = std::stoull(readLine(p.ack_rd));
    } else {
        stopProvider(p);
        printf("step 3: provider %d exited; starting a new one on %s\n",
               p.pid, seg_name.c_str());
        p = spawnProvider(self, port, {"single", "0xB1"});
        line = readLine(p.ack_rd);
        A2 = std::stoull(line.substr(line.find(' ') + 1));
        if (line.substr(0, line.find(' ')) != seg_name)
            DIE("new provider has a different segment name");
    }
    printf("step 3: provider re-published %#lx (same address: %s)\n", A2,
           A2 == A ? "yes" : "NO - repro precondition not met");

    // Let any push notification land.
    std::this_thread::sleep_for(std::chrono::milliseconds(500));
    v = readRemote(te, seg, A2, local, s);
    printf("step 4: READ %#lx -> 0x%02x status=%s (expect 0xb1 or a failure)\n",
           A2, v, statusName(s));

    bool bug = (s == COMPLETED && v == 0xA0);
    printf("RESULT[%s]: %s\n", scenario.c_str(),
           bug ? "BUG REPRODUCED - stale mapping served old allocation"
               : (v == 0xB1 ? "not reproduced (read new data)"
                            : "inconclusive"));
    stopProvider(p);
    te.unregisterLocalMemory(local, kLen);
    te.freeLocalMemory(local);
    return bug ? 1 : 0;
}

int main(int argc, char **argv) {
    signal(SIGPIPE, SIG_IGN);
    if (argc >= 7 && std::string(argv[1]) == "provider")
        return providerMain(argc, argv);
    if (argc == 3 && std::string(argv[1]) == "consumer")
        return consumerMain("/proc/self/exe", argv[2]);
    fprintf(stderr, "usage: %s consumer realloc|restart\n", argv[0]);
    return 2;
}

@SongOf commented on GitHub (Oct 12, 2026): ## Reproduced on current main Reproduced on `b198a191` (main) with **POSIX-FD** handles, two processes sharing one GPU. Both the in-process re-allocation case and the provider-restart case return the old allocation's data, and the transfer reports `COMPLETED`. ### Environment - 2× NVIDIA L20 (PCIe, no P2P between them), driver 580.126.09, CUDA 12.8, Ubuntu 22.04. This is not MNNVL hardware; it is used as a single-node repro box with POSIX-FD handles, and only GPU 0 is used. - Build: `-DUSE_CUDA=ON -DUSE_TENT=ON` - Run: `MC_ENABLE_MNNVL=1 CUDA_VISIBLE_DEVICES=0`, config `metadata_type=p2p`, `transports/nvlink/handle_type=posix_fd`, TCP/RDMA/io_uring disabled so MNNVL is used - Only the public `TransferEngine` API is used. The provider is fork+exec'd, so neither process forks after CUDA init. ### Results | Scenario | Peer address after change | 2nd READ by consumer | Status | |---|---|---|---| | **realloc**: P0 `unregister`+`free` B0 (0xA0), `allocate`+`register` B1 (0xB1) | same VA | **0xA0** | `COMPLETED` | | **restart**: P0 exits, P1 starts on the same `ip:port`, fills 0xB1 (`setarch -R` so the VA repeats) | same VA | **0xA0** | `COMPLETED` (2/2 runs) | ``` === realloc === [provider 11507] B0 at 0x7fef6f800000 filled 0xa0, segment 127.0.0.1:17506 [provider 11507] freed B0, B1 at 0x7fef6f800000 (same VA: yes) filled 0xb1 step 2: READ 0x7fef6f800000 -> 0xa0 status=COMPLETED (expect 0xa0) step 3: provider re-published 0x7fef6f800000 (same address: yes) step 4: READ 0x7fef6f800000 -> 0xa0 status=COMPLETED (expect 0xb1 or a failure) RESULT[realloc]: BUG REPRODUCED - stale mapping served old allocation === restart (setarch -R) === [provider 11660] B0 at 0x7fffd5800000 filled 0xa0, segment 127.0.0.1:17659 [provider 11706] B0 at 0x7fffd5800000 filled 0xb1, segment 127.0.0.1:17659 step 2: READ 0x7fffd5800000 -> 0xa0 status=COMPLETED (expect 0xa0) step 3: provider 11660 exited; starting a new one on 127.0.0.1:17659 step 3: provider re-published 0x7fffd5800000 (same address: yes) step 4: READ 0x7fffd5800000 -> 0xa0 status=COMPLETED (expect 0xb1 or a failure) RESULT[restart]: BUG REPRODUCED - stale mapping served old allocation ``` In the restart case the consumer still reads 0xA0 **after P0 has exited**. This confirms the side effect described above: the consumer's imported mapping keeps the old physical allocation alive. With address-space randomization enabled, P1 usually gets a different VA and the restart case does not trigger. The realloc case hits the same VA every time. ### Separate blocker: POSIX-FD import is broken on main To reach this code path I had to locally patch an unrelated bug. `relocateSharedMemoryAddress` passes the fd **by address**: https://github.com/kvcache-ai/Mooncake/blob/b198a191/mooncake-transfer-engine/tent/src/transport/mnnvl/mnnvl_transport.cpp#L818 ```cpp CHECK_CU(cuMemImportFromShareableHandle(&handle, &local_fd, handle_type_)); ``` For `CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR` the driver expects the fd value itself cast to `void *`. Passing `&local_fd` fails with `CUDA_ERROR_UNKNOWN` (999), and the `LOG(FATAL)` in `submitTransferTasks` then aborts the process. A standalone CUDA program confirms it: `(void *)(uintptr_t)fd` imports fine, `&fd` returns 999. **So on main, MNNVL in POSIX-FD mode cannot import any peer buffer.** The results above were obtained with only this line changed to `(void *)(uintptr_t)local_fd`. Fixed in #4597. ### Notes on Fabric mode Fabric mode could not be tested on this box. `CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_FABRIC_SUPPORTED` reports 1, so `supportFabricMem()` passes and MNNVL installs. However, `cuMemCreate` with `requestedHandleTypes = FABRIC` returns `CUDA_ERROR_NOT_PERMITTED` (800): there is no IMEX channel and no NVLink fabric on L20. Because the cache logic does not depend on the handle type, the same behaviour is expected with Fabric handles on an IMEX-enabled system, but that still needs confirming on GB200-class hardware. The repro pins `CUDA_VISIBLE_DEVICES=0` because the two L20s have no P2P path and MNNVL grants every visible device access to a mapping. That only matters on PCIe boxes like this one, not on NVLink-domain hardware, and is unrelated to this issue. <details> <summary>Repro program (<code>mnnvl_reregister_repro.cpp</code>)</summary> Build it as an extra target next to the TENT tests, linked against `tent_link_group`, then run: ``` CUDA_VISIBLE_DEVICES=0 MC_ENABLE_MNNVL=1 ./tent_mnnvl_reregister_repro consumer realloc CUDA_VISIBLE_DEVICES=0 MC_ENABLE_MNNVL=1 setarch $(uname -m) -R ./tent_mnnvl_reregister_repro consumer restart ``` ```cpp // Repro for kvcache-ai/Mooncake#4592: TENT MnnvlTransport keeps using an // imported mapping after the peer re-registers a new allocation at the same // virtual address. // // repro consumer <scenario> scenario = realloc | restart // repro provider <port> <ctl_rd> <ack_wr> <mode> <pattern1> [pattern2] // // Run with MC_ENABLE_MNNVL=1. Providers are fork+exec'd so no process forks // after CUDA initialization. #include <cuda_runtime.h> #include <signal.h> #include <sys/wait.h> #include <unistd.h> #include <chrono> #include <cstdio> #include <cstdlib> #include <cstring> #include <memory> #include <string> #include <thread> #include <vector> #include "tent/common/config.h" #include "tent/transfer_engine.h" using namespace mooncake::tent; static constexpr size_t kLen = 4 << 20; #define DIE(...) \ do { \ fprintf(stderr, "[%d] FATAL: ", getpid()); \ fprintf(stderr, __VA_ARGS__); \ fprintf(stderr, "\n"); \ _exit(10); \ } while (0) #define CHK(expr) \ do { \ auto _s = (expr); \ if (!_s.ok()) DIE("%s -> %s", #expr, _s.ToString().c_str()); \ } while (0) static std::shared_ptr<Config> makeConfig(int port) { auto conf = std::make_shared<Config>(); conf->set("metadata_type", std::string("p2p")); conf->set("rpc_server_hostname", std::string("127.0.0.1")); if (port) conf->set("rpc_server_port", port); conf->set("transports/tcp/enable", false); conf->set("transports/rdma/enable", false); conf->set("transports/io_uring/enable", false); // Fabric export needs IMEX; POSIX FD works on any VMM-capable GPU. conf->set("transports/nvlink/handle_type", std::string("posix_fd")); conf->set("log_level", std::string("warning")); return conf; } static void writeLine(int fd, const std::string &s) { std::string l = s + "\n"; if (write(fd, l.data(), l.size()) != (ssize_t)l.size()) DIE("write"); } static std::string readLine(int fd) { std::string s; char c; while (true) { ssize_t n = read(fd, &c, 1); if (n <= 0) DIE("peer closed pipe"); if (c == '\n') return s; s += c; } } // ---------------------------------------------------------------- provider struct ProviderBuf { void *addr = nullptr; }; static void allocFill(TransferEngine &te, ProviderBuf &b, int pattern) { CHK(te.allocateLocalMemory(&b.addr, kLen, "cuda:0")); CHK(te.registerLocalMemory(b.addr, kLen)); if (cudaMemset(b.addr, pattern, kLen) != cudaSuccess) DIE("cudaMemset"); cudaDeviceSynchronize(); } static int providerMain(int argc, char **argv) { int port = atoi(argv[2]); int ctl = atoi(argv[3]), ack = atoi(argv[4]); std::string mode = argv[5]; int p1 = strtol(argv[6], nullptr, 0); cudaSetDevice(0); TransferEngine te(makeConfig(port)); if (!te.available()) DIE("provider engine not available"); ProviderBuf b; allocFill(te, b, p1); fprintf(stderr, "[provider %d] B0 at %p filled 0x%02x, segment %s\n", getpid(), b.addr, p1, te.getSegmentName().c_str()); writeLine(ack, te.getSegmentName() + " " + std::to_string((uint64_t)b.addr)); if (mode == "realloc") { int p2 = strtol(argv[7], nullptr, 0); readLine(ctl); // "realloc" CHK(te.unregisterLocalMemory(b.addr, kLen)); CHK(te.freeLocalMemory(b.addr)); void *old = b.addr; allocFill(te, b, p2); fprintf(stderr, "[provider %d] freed B0, B1 at %p (same VA: %s) filled 0x%02x\n", getpid(), b.addr, old == b.addr ? "yes" : "no", p2); writeLine(ack, std::to_string((uint64_t)b.addr)); } readLine(ctl); // "exit" return 0; } // ---------------------------------------------------------------- consumer struct Provider { pid_t pid; int ctl_wr, ack_rd; }; static Provider spawnProvider(const char *self, int port, const std::vector<std::string> &extra) { int ctl[2], ack[2]; if (pipe(ctl) || pipe(ack)) DIE("pipe"); pid_t pid = fork(); if (pid == 0) { close(ctl[1]); close(ack[0]); std::vector<std::string> args = {self, "provider", std::to_string(port), std::to_string(ctl[0]), std::to_string(ack[1])}; args.insert(args.end(), extra.begin(), extra.end()); std::vector<char *> av; for (auto &a : args) av.push_back(const_cast<char *>(a.c_str())); av.push_back(nullptr); execv(self, av.data()); _exit(127); } close(ctl[0]); close(ack[1]); return {pid, ctl[1], ack[0]}; } static void stopProvider(Provider &p) { writeLine(p.ctl_wr, "exit"); int st = 0; waitpid(p.pid, &st, 0); close(p.ctl_wr); close(p.ack_rd); } // READ kLen bytes from (seg, remote) into local device buffer; returns the // first byte observed and the transfer status. static int readRemote(TransferEngine &te, SegmentID seg, uint64_t remote, void *local, TransferStatusEnum &result) { cudaMemset(local, 0, kLen); cudaDeviceSynchronize(); auto batch = te.allocateBatch(1); Request r{}; r.opcode = Request::READ; r.source = local; r.target_id = seg; r.target_offset = remote; r.length = kLen; CHK(te.submitTransfer(batch, {r})); TransferStatus st{}; auto deadline = std::chrono::steady_clock::now() + std::chrono::seconds(10); do { CHK(te.getTransferStatus(batch, 0, st)); } while (st.s == PENDING && std::chrono::steady_clock::now() < deadline); te.freeBatch(batch); result = st.s; std::vector<unsigned char> host(kLen); cudaMemcpy(host.data(), local, kLen, cudaMemcpyDeviceToHost); // Make sure the whole buffer is uniform. for (size_t i = 1; i < kLen; ++i) if (host[i] != host[0]) return -1; return host[0]; } static const char *statusName(TransferStatusEnum s) { switch (s) { case COMPLETED: return "COMPLETED"; case FAILED: return "FAILED"; case PENDING: return "PENDING(timeout)"; default: return "OTHER"; } } static int consumerMain(const char *self, const std::string &scenario) { const int port = 17000 + (getpid() % 1000); Provider p; std::vector<std::string> ack; if (scenario == "realloc") p = spawnProvider(self, port, {"realloc", "0xA0", "0xB1"}); else p = spawnProvider(self, port, {"single", "0xA0"}); auto line = readLine(p.ack_rd); std::string seg_name = line.substr(0, line.find(' ')); uint64_t A = std::stoull(line.substr(line.find(' ') + 1)); // CUDA is first touched here, after all forks except `restart`'s second // spawn, which exec()s immediately. cudaSetDevice(0); TransferEngine te(makeConfig(0)); if (!te.available()) DIE("consumer engine not available"); void *local = nullptr; CHK(te.allocateLocalMemory(&local, kLen, "cuda:0")); CHK(te.registerLocalMemory(local, kLen)); SegmentID seg; CHK(te.openSegment(seg, seg_name)); TransferStatusEnum s; int v = readRemote(te, seg, A, local, s); printf("step 2: READ %#lx -> 0x%02x status=%s (expect 0xa0)\n", A, v, statusName(s)); uint64_t A2; if (scenario == "realloc") { writeLine(p.ctl_wr, "realloc"); A2 = std::stoull(readLine(p.ack_rd)); } else { stopProvider(p); printf("step 3: provider %d exited; starting a new one on %s\n", p.pid, seg_name.c_str()); p = spawnProvider(self, port, {"single", "0xB1"}); line = readLine(p.ack_rd); A2 = std::stoull(line.substr(line.find(' ') + 1)); if (line.substr(0, line.find(' ')) != seg_name) DIE("new provider has a different segment name"); } printf("step 3: provider re-published %#lx (same address: %s)\n", A2, A2 == A ? "yes" : "NO - repro precondition not met"); // Let any push notification land. std::this_thread::sleep_for(std::chrono::milliseconds(500)); v = readRemote(te, seg, A2, local, s); printf("step 4: READ %#lx -> 0x%02x status=%s (expect 0xb1 or a failure)\n", A2, v, statusName(s)); bool bug = (s == COMPLETED && v == 0xA0); printf("RESULT[%s]: %s\n", scenario.c_str(), bug ? "BUG REPRODUCED - stale mapping served old allocation" : (v == 0xB1 ? "not reproduced (read new data)" : "inconclusive")); stopProvider(p); te.unregisterLocalMemory(local, kLen); te.freeLocalMemory(local); return bug ? 1 : 0; } int main(int argc, char **argv) { signal(SIGPIPE, SIG_IGN); if (argc >= 7 && std::string(argv[1]) == "provider") return providerMain(argc, argv); if (argc == 3 && std::string(argv[1]) == "consumer") return consumerMain("/proc/self/exe", argv[2]); fprintf(stderr, "usage: %s consumer realloc|restart\n", argv[0]); return 2; } ``` </details>

bsakaguchi referenced this issue

2026-10-29 21:32:22 +00:00

[PR #1170] [MERGED] [Store] feat:mooncake store enable ascend fabric mem mode #2048