[TENT] fix: TCP transport implicitly creates CUDA contexts

[gogongxt]

Description

TENT's TCP transport implicitly creates CUDA primary contexts in processes that never touched CUDA, or on GPUs the transfer never uses — several hundred MiB per process (~520 MiB on H200, ~414 MiB on A800). In per-rank deployments (CUDA_VISIBLE_DEVICES=<rank>), every rank's TCP traffic burns ~0.5 GB of GPU memory on its own card.

Three independent mechanisms, each verified on current main (2×A800, CUDA 12.9, TENT-enabled wheel, per-process nvidia-fd tracking at staged checkpoints):

1. Host-to-host copies go through the CUDA stream pool. CudaPlatform::copy() classifies both pointers as "current device", CUDAStreamPool::acquire() then calls cudaGetDevice() (device 0 on a fresh RPC worker thread) and cudaStreamCreateWithFlags() creates a primary context on GPU 0. This is reached on every TCP WRITE via the server-side ControlService::onSendData() (control_plane.cpp:516 — no MTYPE_CPU fast path, unlike sendData()/onRecvData()) and on every TCP READ via ControlClient::recvData() (control_plane.cpp:148).
2. Runtime-API pointer classification. cudaPointerGetAttributes() on a thread with no prior cudaSetDevice() implicitly initializes the calling thread's current device (default GPU 0).
3. cudaMemcpyAsync() validates the stream against the current device's context. When a thread with no context (current device 0) copies through a stream belonging to another GPU, the runtime creates a primary context on GPU 0 — even though both buffer and stream live elsewhere. Confirmed with a minimal C-level repro (all other calls in the sequence are innocent).

Fix

• New tent/platform/cuda_utils.h: ensureCudaDriverInit() (idempotent cuInit, creates no context — the driver API does not initialize itself lazily) and getCudaDeviceForPtr(), which classifies pointers via cuPointerGetAttribute and reads the driver's global pointer table without creating any context or implicitly initializing a device.
• CudaPlatform::getMemoryType() / getPointerDeviceId() / getLocation() / free() and the CUDA device plugin now classify pointers through the driver API.
• CudaPlatform::copy(): host-to-host copies use plain ::memcpy; device-involved copies pin the buffer's device for the whole acquire/copy/synchronize sequence (restoring the caller's binding only when the thread already had a context), so a context is only ever created on the GPU that actually owns the buffer.

This supersedes #2330 (authored against a June tree; copy() has since been reworked by #3476 and the probe guards by #3261/#3945, so that diff no longer applies — and it did not cover mechanism 3 or the missing cuInit). It is the TENT counterpart of #2307, which fixed the classic (non-TENT) TcpTransport.

Module

• Transfer Engine (mooncake-transfer-engine)

Type of Change

• Bug fix

How Has This Been Tested?

Test commands: two-process transfer over loopback TCP (MC_USE_TENT=1, tcp-only MC_TENT_CONF, GPU buffers allocated via ctypes libcudart so the test itself never pulls in torch), observing nvidia-smi --query-compute-apps plus per-process /proc/<pid>/fd nvidia-fd counts (driver init vs. context creation) at staged checkpoints; round-trip data-integrity checks in both directions. Also ran the exact repro script shared by @chestnut-Q in the #2307 discussion.

Test results: (2×A800-SXM4-80GB, CUDA 12.9, wheel built from this branch)

Scenario (pure TCP, no torch)	Before	After
DRAM→DRAM WRITE	CPU-only target gets a 414 MiB context on GPU 0	no context
DRAM→DRAM READ	CPU-only initiator gets a 414 MiB context on GPU 0	no context
same + CUDA_VISIBLE_DEVICES=1	414 MiB context on the rank's own GPU (the reported ~0.5 GB/card symptom)	no context; round-trip data verified
cuda:1 → DRAM WRITE	initiator: contexts on GPU 0 and GPU 1; target polluted on GPU 0	initiator: GPU 1 only (its own allocation); target clean; round-trip data verified
@chestnut-Q's repro script (#2307)	extra ~320 MiB on GPU 0 in the initiator (per their comment on #2330)	GPU 0 completely clean; only the initiator's own GPU 1 allocation remains; 50× WRITE completes

• Manual testing done (described above)
• Unit tests pass (deferred to CI)

Checklist

• I have performed a self-review of my own code
• I have formatted my code using ./scripts/code_format.sh (clang-format clean on all changed files)
• I have run pre-commit on the files changed in this PR and all hooks pass
• I have updated the documentation (if applicable) — N/A
• I have added tests to prove my changes are effective (manual repro matrix above; happy to add a unit test if maintainers prefer)
• For changes >500 LOC: I have filed an RFC issue — N/A (~100 LOC)

AI Assistance Disclosure

• AI tools were used (specify below)

Claude Code assisted with root-cause analysis (fd-tracking repro, staged isolation, minimal C-level repro), implementing the fix, and running the verification matrix. The human submitter reviewed every changed line and can defend the change end-to-end.

[alogfans]

Could you please check the CI failure since it seems related to your changes?