Port three host fixes from the official TensorFold: RDMA GID auto-detect, NCCL cuMem off, page cache release

NOTICE

@@ -17,6 +17,7 @@ nodes. TensorFold's text above, LICENSE, LICENSES/ and the upstream sections of
 Added:
   src/tensorfold/families/deepseek_v41/   the model family (engine, kernels, server integration)
   src/tensorfold/cuda/rdma.py, src/tensorfold/cuda/rdma_gather.cu   RDMA-write all-gather for TP decode
+  src/tensorfold/cuda/roce_gid.py          RoCE v2 IPv4 GID index scan of an RDMA port (sysfs)
   tools/dsv41/                             tests, checks, benchmark clients, the reference model, DESIGN.md,
                                            REPORT.md and ATTRIBUTION.md
   A "DeepSeek-V4.1-Flash on CUDA (TP2)" section in THIRD_PARTY_NOTICES.md, and a fork note at the top of README.md.
@@ -49,4 +50,10 @@ Modified upstream files:
   src/tensorfold/server/authentication.py      the no-key startup line names the CUDA server's TF_API_KEY gate
   .gitignore                                   ignores macOS ._* files
 
+Later work in the official TensorFold (https://github.com/ashhart/TensorFold, main after v0.6.6, Apache-2.0, Copyright
+2026 TensorFold contributors) gave three ideas that this fork wrote again in Python, without copying its Zig code:
+NCCL's cuMem allocations off by default (its pull request #995), the GID index scan of an RDMA port (`gidIndex` in
+`zig/src/cuda/roce.zig`) and the release of a loaded file's page cache (its pull requests #766 and #916). They are in
+src/tensorfold/families/deepseek_v41/cuda/hostenv.py, src/tensorfold/cuda/roce_gid.py and the callers of those files.
+
 The "git log" of this branch and tools/dsv41/ATTRIBUTION.md describe the changes and their sources.

THIRD_PARTY_NOTICES.md

@@ -94,6 +94,12 @@ This section was added by the fork that adds the `deepseek_v41` family; see `NOT
 - The grouped EXL3 expert kernel additions (`src/tensorfold/cuda/exl3/experts_grouped.cuh`, `experts.cu`) follow
   [ExLlamaV3](https://github.com/turboderp-org/exllamav3)'s EXL3 format, MIT License, Copyright (c) 2025 Turboderp
   (text below), as the rest of the shared EXL3 module does.
+- Three host settings follow ideas of the official [TensorFold](https://github.com/ashhart/TensorFold) main branch
+  (Apache-2.0, Copyright 2026 TensorFold contributors), written again in Python without taking its Zig code: NCCL's
+  cuMem allocations off by default (`hostenv.py`; its pull request #995), the RoCE v2 IPv4 GID scan of a port
+  (`src/tensorfold/cuda/roce_gid.py`; the idea of `gidIndex` in `zig/src/cuda/roce.zig`) and the page cache release
+  after a weight load (`hostenv.py`, `weights.py`; its pull requests #766 and #916). None of its DeepSeek-V4.1 files
+  (`families/deepseek_v41/`, `kernels/cuda/deepseek_v41/`) was opened.
 - The baseline used in `tools/dsv41/REPORT.md` (a vLLM kit under AGPL-3.0) was only run as a black box over HTTP. None
   of its code was read or included. `tools/dsv41/ATTRIBUTION.md` lists every source in full.
 

src/tensorfold/cuda/rdma.py

@@ -12,6 +12,10 @@
 TF_RDMA_STAGE_FENCE / TF_RDMA_OWN_EARLY / TF_RDMA_COLLECT_UNROLL = 0 each restore c100's way. TF_RDMA_SLOTS (default 4)
 sets the ring slots (2 halves the pinned rings; any count >= 2 is safe: a rank stages gather q + 2 only after its own
 collect of q + 1, which needs the peer's data of q + 1, sent after the peer's collect of q).
+
+The GID index of each port is the lowest RoCE v2 IPv4 entry of its GID table, read at start (roce_gid.py): the tables of
+two nodes can differ, and a link flap can move an entry. TF_RDMA_GID_INDEX or NCCL_IB_GID_INDEX with a value of 0 or more
+fixes the index for every port; unset, empty or -1 means "detect".
 """
 
 from __future__ import annotations
@@ -24,6 +28,8 @@
 
 import torch
 
+from .roce_gid import active_roce_devices, resolve_gid_indices
+
 
 @lru_cache(maxsize=1)
 def _ext():
@@ -42,16 +48,7 @@ def device_names() -> list[str]:
     if named:
         return named
     first = (os.environ.get("NCCL_IB_HCA") or "").strip().lstrip("=^").split(",")[0].split(":")[0].strip()
-    root = Path("/sys/class/infiniband")
-    active = []
-    for d in sorted(p.name for p in root.iterdir()) if root.exists() else []:
-        try:
-            state = (root / d / "ports" / "1" / "state").read_text()
-            layer = (root / d / "ports" / "1" / "link_layer").read_text().strip()
-        except OSError:
-            continue
-        if "ACTIVE" in state and layer == "Ethernet":
-            active.append(d)
+    active = active_roce_devices()
     if first in active:
         active.remove(first)
         active.insert(0, first)
@@ -65,15 +62,17 @@ def __init__(self, store, rank: int, world: int, *, max_bytes: int, slots: int |
                  prefix: str = "tf_rdma", devices: list[str] | None = None, gid_index: int | None = None) -> None:
         ext = _ext()
         slots = int(os.environ.get("TF_RDMA_SLOTS") or 4) if slots is None else int(slots)
-        gid = int(os.environ.get("NCCL_IB_GID_INDEX", "5")) if gid_index is None else int(gid_index)
         cpu = int(os.environ.get("TF_RDMA_CPU", "-1"))
         self.ext, self.rank, self.world, self.max_bytes = ext, rank, world, int(max_bytes)
         self.devices = device_names() if devices is None else list(devices)
         error = ""
         try:
             if slots < 2:
                 raise ValueError(f"TF_RDMA_SLOTS must be at least 2 (got {slots})")
-            self.h = ext.create(rank, world, self.max_bytes, slots, self.devices, gid, cpu)
+            gids = resolve_gid_indices(self.devices, fixed=gid_index)
+            for name, (index, how) in zip(self.devices, gids):
+                print(f"[tensorfold] rank {rank}: RDMA GID index of {name} port 1: {index} ({how})", flush=True)
+            self.h = ext.create(rank, world, self.max_bytes, slots, self.devices, [g[0] for g in gids], cpu)
             info = pickle.dumps((ext.local_info(self.h), slots, self.max_bytes))
         except Exception as exc:           # noqa: BLE001  (published, so every rank refuses together)
             error, info = f"{type(exc).__name__}: {exc}", b""

src/tensorfold/cuda/rdma_gather.cu

@@ -86,6 +86,7 @@ struct Nic {
     ibv_mr* mr = nullptr;                 // the send ring's (the whole block's when the rings are cudaHostAlloc memory)
     ibv_mr* mr_recv = nullptr;            // the receive ring's (== mr in that case)
     ibv_gid gid{};
+    int gid_index = 0;                    // this port's GID table entry (RoCE v2, IPv4), chosen by rdma.py
     ibv_mtu mtu = IBV_MTU_4096;           // the port's active MTU (RoCE: the netdev MTU less headers, rounded down)
     std::vector<ibv_qp*> qps;             // per peer
     std::vector<uint32_t> psns;
@@ -98,7 +99,7 @@ struct Peer {
 };
 
 struct Gather {
-    int rank = 0, world = 1, slots = 4, port = 1, gid_index = 5, cpu = -1;
+    int rank = 0, world = 1, slots = 4, port = 1, cpu = -1;
     size_t max_bytes = 0;
     char* host = nullptr;                 // the metadata's block (host_reg 0: rings + metadata in one cudaHostAlloc)
     size_t host_bytes = 0;
@@ -375,7 +376,7 @@ void proxy_loop(Gather* gp) {
     }
 }
 
-Nic open_nic(const std::string& name, Gather& g, int world, int rank) {
+Nic open_nic(const std::string& name, Gather& g, int world, int rank, int gid_index) {
     int count = 0;
     ibv_device** list = ibv_get_device_list(&count);
     TORCH_CHECK(list && count > 0, "rdma gather: no RDMA devices");
@@ -402,7 +403,9 @@ Nic open_nic(const std::string& name, Gather& g, int world, int rank) {
         TORCH_CHECK(nic.mr_recv, "rdma gather: ibv_reg_mr of the ", g.recv_bytes, "-byte receive ring on ", name,
                     " failed");
     }
-    TORCH_CHECK(ibv_query_gid(nic.ctx, g.port, g.gid_index, &nic.gid) == 0, "rdma gather: ibv_query_gid failed");
+    nic.gid_index = gid_index;
+    TORCH_CHECK(gid_index >= 0 && ibv_query_gid(nic.ctx, g.port, gid_index, &nic.gid) == 0,
+                "rdma gather: ibv_query_gid failed for GID index ", gid_index, " of ", name, " port ", g.port);
     // The QPs take the port's active MTU, not a fixed 4096: at netdev MTU 1500 RoCE runs at 1024, and a QP set to 4096
     // sends frames the link drops (every write then fails with "transport retry counter exceeded").
     ibv_port_attr port{};
@@ -444,16 +447,16 @@ Nic open_nic(const std::string& name, Gather& g, int world, int rank) {
 }
 
 int64_t create(int64_t rank, int64_t world, int64_t max_bytes, int64_t slots, const std::vector<std::string>& devices,
-               int64_t gid_index, int64_t cpu) {
+               const std::vector<int64_t>& gid_indices, int64_t cpu) {
     TORCH_CHECK(world >= 2 && rank >= 0 && rank < world, "rdma gather: rank ", rank, " of ", world);
     TORCH_CHECK(max_bytes > 0 && max_bytes % 64 == 0, "rdma gather: max_bytes must be a positive multiple of 64");
     TORCH_CHECK(!devices.empty(), "rdma gather: name at least one RDMA device");
+    TORCH_CHECK(gid_indices.size() == devices.size(), "rdma gather: one GID index a device");
     auto g = std::make_unique<Gather>();
     g->rank = (int)rank;
     g->world = (int)world;
     g->slots = (int)slots;
     g->max_bytes = (size_t)max_bytes;
-    g->gid_index = (int)gid_index;
     g->cpu = (int)cpu;
     const size_t send_b = (size_t)slots * max_bytes, recv_b = (size_t)slots * world * max_bytes;
     const size_t meta_b = ((size_t)slots * world + 1 + 2 * (size_t)slots + 1) * 8;
@@ -506,7 +509,8 @@ int64_t create(int64_t rank, int64_t world, int64_t max_bytes, int64_t slots, co
     C10_CUDA_CHECK(cudaMemset(g->staged, 0, 8));
     C10_CUDA_CHECK(cudaMalloc(reinterpret_cast<void**>(&g->probe), 8 * 8));
     C10_CUDA_CHECK(cudaMemset(g->probe, 0, 8 * 8));
-    for (const auto& name : devices) g->nics.push_back(open_nic(name, *g, (int)world, (int)rank));
+    for (size_t i = 0; i < devices.size(); ++i)
+        g->nics.push_back(open_nic(devices[i], *g, (int)world, (int)rank, (int)gid_indices[i]));
     g->peers.assign(world, Peer{});
     gathers.push_back(std::move(g));
     return (int64_t)gathers.size() - 1;
@@ -561,7 +565,7 @@ void link_peers(int64_t h, const std::vector<py::tuple>& remote) {
             rtr.min_rnr_timer = 12;
             rtr.ah_attr.is_global = 1;
             std::memcpy(rtr.ah_attr.grh.dgid.raw, peer.gid[d].data(), 16);
-            rtr.ah_attr.grh.sgid_index = g.gid_index;
+            rtr.ah_attr.grh.sgid_index = g.nics[d].gid_index;
             rtr.ah_attr.grh.hop_limit = 64;
             rtr.ah_attr.port_num = g.port;
             TORCH_CHECK(ibv_modify_qp(g.nics[d].qps[p], &rtr,

src/tensorfold/cuda/roce_gid.py

@@ -0,0 +1,150 @@
+"""Find the RoCE v2 IPv4 GID index of an RDMA port (no torch needed).
+
+The GID table of a CX7 port can differ from node to node, and it can change after a link flap. A fixed index
+then points at the wrong entry, or at an empty one. The kernel shows the table in sysfs:
+``<root>/<device>/ports/<port>/gids/<n>`` holds the GID and ``.../gid_attrs/types/<n>`` holds its type.
+An entry is the one we want when its type is "RoCE v2" and its GID is an IPv4-mapped address.
+
+The scan idea follows ``gidIndex`` in ``zig/src/cuda/roce.zig`` of the official TensorFold
+(https://github.com/ashhart/TensorFold, Apache-2.0). The code here is written for this fork.
+"""
+
+from __future__ import annotations
+
+import os
+from collections.abc import Mapping, Sequence
+from pathlib import Path
+
+SYSFS = Path("/sys/class/infiniband")
+# The first 96 bits of an IPv4-mapped IPv6 address (::ffff:a.b.c.d), as the kernel prints a GID.
+IPV4_MAPPED = "0000:0000:0000:0000:0000:ffff:"
+# The variables that can fix the index. The first one that has a value decides. A value of -1 means "detect".
+OVERRIDE_VARS = ("TF_RDMA_GID_INDEX", "NCCL_IB_GID_INDEX")
+
+
+def _read(path: Path) -> str | None:
+    try:
+        return path.read_text().strip()
+    except OSError:                         # an empty table entry has no type file that can be read
+        return None
+
+
+def _entries(directory: Path) -> list[int]:
+    try:
+        return sorted(int(p.name) for p in directory.iterdir() if p.name.isdigit())
+    except OSError:
+        return []
+
+
+def gid_entries(device: str, port: int = 1, *, root: Path | str = SYSFS) -> list[tuple[int, str, str]]:
+    """Every readable entry of the GID table as ``(index, type, gid)``, in index order."""
+
+    base = Path(root) / device / "ports" / str(port)
+    out = []
+    for n in _entries(base / "gids"):
+        kind = _read(base / "gid_attrs" / "types" / str(n))
+        gid = _read(base / "gids" / str(n))
+        if kind is not None and gid is not None:
+            out.append((n, kind, gid))
+    return out
+
+
+def detect_gid_index(device: str, port: int = 1, *, root: Path | str = SYSFS) -> int | None:
+    """The lowest index of a RoCE v2 entry with an IPv4 address, or None when the table has none."""
+
+    for n, kind, gid in gid_entries(device, port, root=root):
+        if not kind.startswith("RoCE v2"):
+            continue
+        gid = gid.lower()
+        if gid.startswith(IPV4_MAPPED) and gid[len(IPV4_MAPPED):] != "0000:0000":
+            return n
+    return None
+
+
+def active_roce_devices(*, root: Path | str = SYSFS) -> list[str]:
+    """The RDMA devices whose port 1 is ACTIVE with an Ethernet link layer (RoCE), in name order."""
+
+    root = Path(root)
+    out = []
+    try:
+        names = sorted(p.name for p in root.iterdir())
+    except OSError:
+        return out
+    for d in names:
+        state, layer = _read(root / d / "ports" / "1" / "state"), _read(root / d / "ports" / "1" / "link_layer")
+        if state is not None and "ACTIVE" in state and layer == "Ethernet":
+            out.append(d)
+    return out
+
+
+def nccl_hca_ports(env: Mapping[str, str] | None = None, *, root: Path | str = SYSFS) -> list[tuple[str, int]] | None:
+    """The ``(device, port)`` pairs that NCCL_IB_HCA lets NCCL use, among the active RoCE devices.
+
+    NCCL_IB_HCA is ``[^][=]<name>[:<port>],...``: a name matches the start of a device name, or the whole name after
+    ``=``. Unset or empty: every active RoCE device (port 1). A list with ``^`` excludes devices: returns None, as the
+    scan does not follow NCCL's rules for it. A pair without a port uses port 1.
+    """
+
+    env = os.environ if env is None else env
+    raw = (env.get("NCCL_IB_HCA") or "").strip()
+    active = active_roce_devices(root=root)
+    if not raw:
+        return [(d, 1) for d in active]
+    if raw.startswith("^"):
+        return None
+    exact = raw.startswith("=")
+    out = []
+    for item in raw.lstrip("=").split(","):
+        name, _, port = item.strip().partition(":")
+        if not name:
+            continue
+        try:
+            number = int(port) if port.strip() else 1
+        except ValueError:
+            return None
+        for d in active:
+            if (d == name if exact else d.startswith(name)) and (d, number) not in out:
+                out.append((d, number))
+    return out
+
+
+def gid_override(env: Mapping[str, str] | None = None) -> int | None:
+    """The index that the user fixed, or None when the user did not fix one (variable unset, empty or -1)."""
+
+    env = os.environ if env is None else env
+    for name in OVERRIDE_VARS:
+        raw = (env.get(name) or "").strip()
+        if not raw:
+            continue
+        try:
+            value = int(raw)
+        except ValueError:
+            raise ValueError(f"{name} must be a whole number (got {raw!r})") from None
+        return value if value >= 0 else None
+    return None
+
+
+def resolve_gid_indices(devices: Sequence[str], port: int = 1, *, fixed: int | None = None,
+                        env: Mapping[str, str] | None = None,
+                        root: Path | str = SYSFS) -> list[tuple[int, str]]:
+    """One ``(index, how)`` for each device. ``how`` is "fixed" (an argument or a variable) or "detected".
+
+    ``fixed`` (an argument) wins over the variables. A negative ``fixed`` means "detect", and the variables are not read.
+    Raises RuntimeError when a device has no RoCE v2 IPv4 entry and nothing fixes the index.
+    """
+
+    if fixed is None:
+        fixed = gid_override(env)
+    elif fixed < 0:
+        fixed = None
+    out = []
+    for device in devices:
+        if fixed is not None:
+            out.append((fixed, "fixed"))
+            continue
+        found = detect_gid_index(device, port, root=root)
+        if found is None:
+            raise RuntimeError(f"{device} port {port} has no RoCE v2 IPv4 GID in {Path(root) / device}; "
+                               f"set {OVERRIDE_VARS[0]} to the index to use")
+        out.append((found, "detected"))
+    return out

src/tensorfold/families/deepseek_v41/cuda/engine.py

@@ -82,6 +82,8 @@ def __init__(self, model_dir: Path, *, rank: int, world: int, master: str, port:
         # gather on the CX7 link against 7.1 with NCCL's choice (decode windows go over the RDMA gather)
         os.environ.setdefault("NCCL_PROTO", "Simple")
         os.environ.setdefault("NCCL_MIN_NCHANNELS", "4")
+        if world > 1:
+            self._nccl_host_env()
         nccl = NCCL(rank, world, master, port) if world > 1 else None
         self.nccl = nccl
         self.w = load(model_dir, rank, world, dspark=drafts > 0)
@@ -189,6 +191,16 @@ def __init__(self, model_dir: Path, *, rank: int, world: int, master: str, port:
             print(f"[tensorfold] DeepSeek-V4.1 engine ready: {world} rank(s), context {self.limit}, "
                   f"{'DSpark ' + str(drafts) + ' drafts' if self.drafter else 'serial decode'}", flush=True)
 
+    def _nccl_host_env(self) -> None:
+        """NCCL's cuMem allocations off and its GID index set (hostenv.py); a value that the user set stays."""
+
+        from .hostenv import apply_nccl_defaults
+
+        done = apply_nccl_defaults()
+        if done:
+            print(f"[tensorfold] rank {self.rank}: NCCL settings by default: "
+                  + " ".join(f"{k}={v}" for k, v in done.items()), flush=True)
+
     def _memory_ceiling(self) -> None:
         """After the warm-up, torch may grow only while the host keeps TF_DS_MEM_FLOOR_GIB (default 2; 0: no ceiling)
         available: a step that would need more fails with CUDA's out-of-memory error (that request fails, the lane's