Summary
On WSL2, after heavy in-process pinned-memory churn (repeated alloc/free of large pinned buffers), H2D copies from freshly pinned buffers intermittently fail silently: cudaMemcpyBatchAsync (and plain cudaMemcpyAsync via torch.Tensor.copy_) return cudaSuccess, but the device buffer is left unchanged (zeros). Host-side reads of the same source buffer remain correct, and is_pinned() reports True.
This is not a FreeToken code bug — the copies are enqueued and synchronized correctly. It looks like a WSL2/Windows-driver pinned-mapping issue (stale or never-established device mapping of the pinned region).
Environment
- Windows 11 25H2 (build 26200), WSL2 Ubuntu 22.04.5
- 2× RTX 3090, Windows NVIDIA driver 591.86 (CUDA 13.1-capable)
- CUDA 13.0 toolkit in WSL (nvcc 13.0.48), torch 2.11.0+cu130
- FreeToken 0.1.2 built from source (
uv pip install -e ".[accel]")
Observed behavior
- After the churn below, a single freshly pinned 16-byte source buffer: the first device read sometimes succeeds, and all later device reads of the same buffer silently return nothing (dst stays zero), while
src.tolist() on the host is still correct.
- A
time.sleep(0.005) between pin_memory() and the copy makes it dramatically worse (19/20 copies failed vs 1/20 with no sleep) — consistent with a deferred mapping/unmap race in the driver.
CUDA_LAUNCH_BLOCKING=1 avoids the failure entirely (timing change).
- Pure pinned churn (no FreeToken code, no executor) also triggers it at a lower rate.
Impact on FreeToken
tests/moe/test_prefill_hit_d2d.py::test_batch_memcpy_roundtrip fails in ~50% of full-suite runs on this box (the MoE offload tests alloc/free ~192 MB pinned banks per test).
- The
batch_memcpy probe (python/freetoken/kernel/batch_memcpy.py::_probe) catches this and OffloadMoeCache falls back to full-layer copies, so the batch path is protected.
- However the plain-copy path has no equivalent guard: the full-layer copy fallback (and any H2D from pinned memory, e.g. offload prefetch) could hit the same silent no-op in a long-running process that has churned pinned allocations. That would be silent data corruption, not an error.
Minimal repro
# run inside WSL2 with the freetoken venv: python pinned_dma_repro.py
import time
from types import SimpleNamespace
import torch
from freetoken.kernel.pinned import alloc_pinned_tensor
from freetoken.moe.cpu_executor import CpuMoeExecutor
torch.cuda.init()
dev = torch.device("cuda")
L, E, H, I, top_k = 4, 16, 1024, 512, 4
# 1) churn: ~192 MB pinned banks alloc+fill+free, with a CpuMoeExecutor per round
for i in range(20):
gu = alloc_pinned_tensor(L * E, 2 * I, H, dtype=torch.bfloat16)
dn = alloc_pinned_tensor(L * E, H, I, dtype=torch.bfloat16)
gu.copy_(torch.randn(L * E, 2 * I, H) * 0.1)
dn.copy_(torch.randn(L * E, H, I) * 0.1)
cache = SimpleNamespace(
quant_format="bf16",
bank_sources={"gate_up": list(gu.split(E)), "down": list(dn.split(E))},
num_layers=L, num_experts=E, decode_target="cpu", cpu_executor=None,
)
ex = CpuMoeExecutor(cache, top_k=top_k, activation="silu",
apply_router_weight_on_input=False, num_threads=1,
max_tokens=1, device=dev)
del ex, cache, gu, dn
# 2) one pinned source, repeated device reads with a little idle time
src = torch.arange(16, dtype=torch.uint8).pin_memory()
for i in range(8):
time.sleep(0.005 * (i + 1))
dst = torch.zeros(16, dtype=torch.uint8, device=dev)
s = torch.cuda.Stream()
with torch.cuda.stream(s):
dst.copy_(src)
s.synchronize()
print(f"copy{i}: gpu_ok={torch.equal(dst.cpu(), src)} "
f"host_readable={src.tolist() == list(range(16))}")
Typical output on the affected box (fails in most runs):
copy0: gpu_ok=True host_readable=True
copy1: gpu_ok=False host_readable=True
copy2: gpu_ok=False host_readable=True
...
Without the executor (pure alloc_pinned_tensor churn ×100), the same probe fails at a lower rate (~1/20).
Suggested next steps / questions
- Is this known WSL2 behavior with the current driver generation? Any WSL2 setting that changes pinned-memory backing?
- Should the plain-copy offload paths get a cheap integrity guard (e.g. the batch-memcpy probe pattern) for WSL hosts?
- Happy to test a newer Windows driver (591.86 → latest) and report back.
Summary
On WSL2, after heavy in-process pinned-memory churn (repeated alloc/free of large pinned buffers), H2D copies from freshly pinned buffers intermittently fail silently:
cudaMemcpyBatchAsync(and plaincudaMemcpyAsyncviatorch.Tensor.copy_) returncudaSuccess, but the device buffer is left unchanged (zeros). Host-side reads of the same source buffer remain correct, andis_pinned()reportsTrue.This is not a FreeToken code bug — the copies are enqueued and synchronized correctly. It looks like a WSL2/Windows-driver pinned-mapping issue (stale or never-established device mapping of the pinned region).
Environment
uv pip install -e ".[accel]")Observed behavior
src.tolist()on the host is still correct.time.sleep(0.005)betweenpin_memory()and the copy makes it dramatically worse (19/20 copies failed vs 1/20 with no sleep) — consistent with a deferred mapping/unmap race in the driver.CUDA_LAUNCH_BLOCKING=1avoids the failure entirely (timing change).Impact on FreeToken
tests/moe/test_prefill_hit_d2d.py::test_batch_memcpy_roundtripfails in ~50% of full-suite runs on this box (the MoE offload tests alloc/free ~192 MB pinned banks per test).batch_memcpyprobe (python/freetoken/kernel/batch_memcpy.py::_probe) catches this andOffloadMoeCachefalls back to full-layer copies, so the batch path is protected.Minimal repro
Typical output on the affected box (fails in most runs):
Without the executor (pure
alloc_pinned_tensorchurn ×100), the same probe fails at a lower rate (~1/20).Suggested next steps / questions