You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
{{ message }}
Repository navigation
⚠️ WSL2 GPU OOM leads to Hyper-V kernel panic and Windows BLUE SCREEN OF DEATH — total system crash, not a recoverable error #40732
WSL2 GPU OOM (Out of Memory) due to repeated CUDA allocation failures leads to a Hyper-V VM kernel panic followed by a Windows host bugcheck (BSOD) — the entire system crashes and reboots, not just a recoverable CUDA error.
Steps to Reproduce
Load a large model (>18GB VRAM) in WSL2 via PyTorch or llama.cpp — e.g., Qwen3.6-35B on an RTX 5090 (24GB)
While it is loaded, attempt to load a second model (e.g., LocateAnything-3B, ~14GB)
PyTorch CUDA caching allocator / CUDA driver API returns OOM
Retry without cleaning up VRAM fragments 3-5 times
⚠️Instead of returning a clean CUDA error, the WSL2 VM kernel panics, and Windows host bugchecks (BSOD)
Full system restart required
Expected Behavior
CUDA OOM should return a clean error code (cudaErrorMemoryAllocation). The process should fail gracefully. The WSL VM should never crash the Windows host.
Data loss from unsaved work, productivity disruption, and erodes trust in WSL2 for GPU workloads. This makes WSL2 unusable for any development involving large GPU models where OOM is a risk.
If this a feature request, please reply with '/feature'. If this is a question, reply with '/question'. Otherwise, please attach logs by following the instructions below, your issue will not be reviewed unless they are added. These logs will help us understand what is going on in your machine.
How to collect WSL logs
Download and execute collect-wsl-logs.ps1 in an administrative powershell prompt:
The script will output the path of the log file once done.
If this is a networking issue, please use .\collect-wsl-logs.ps1 -LogProfile networking instead, following the instructions in Collect WSL logs for networking issues
Once completed please upload the output files to this GitHub issue.
If you choose to email these logs instead of attaching them to the bug, please send them to wsl-gh-logs@microsoft.com with the GitHub issue number in the subject, and include a link to your GitHub issue comment in the message body, and reply with '/emailed-logs'.
If this a feature request, please reply with '/feature'. If this is a question, reply with '/question'. Otherwise, please attach logs by following the instructions below, your issue will not be reviewed unless they are added. These logs will help us understand what is going on in your machine.
Same symptom, with a ~250-line pure-CUDA reproducer: a guest CUDA process bugchecks the Windows host, and one flag flips it
This issue describes the symptom I am seeing. Adding what I could pin down, in case it helps —
the important part is that it is now reducible to a single .cu file with no PyTorch and no
libraries beyond libcudart, and that one command-line flag decides whether the host
survives.
What happens
A CUDA process inside WSL2 oversubscribes VRAM and keeps allocating. Free VRAM reaches 0, WDDM
starts paging, and a few steps later the Windows host bugchecks and reboots. From inside the
guest there is nothing: no cudaErrorMemoryAllocation, no exception, no kernel message, no dmesg entry. The log simply stops after the last line it managed to flush — which names the last
call entered, not necessarily the one that hung (the line is printed before the launch, and a cudaDeviceSynchronize() follows two statements later).
For a GPU-P guest this is a containment problem regardless of whose driver stalls: an unprivileged
process in the VM can take the host down, and gets no error it could handle.
Environment
Windows kernel version 26100, OS build 26200.9168; WSL2 with GPU-P (/dev/dxg)
⚠ Five of these eleven have no known workload and I am not attributing them. The reproducer
below accounts for six more crashes deliberately produced afterwards.
There is not a single Display (EventID 4101, TDR) or nvlddmkm event in the System log — so
this is not a display-driver reset; the host dies before TDR reports anything. That query and its
positive control are in EVIDENCE_no_display_driver_events.txt.
The reproducer, and the flag that flips it
Five allocations per step, nothing ever freed: 4 sources of 64 MiB ((239674,70) fp32) plus one
destination of 256 MiB ((239674,280) fp32), one kernel reading all four and writing the
destination with a concat-along-dim-1 pattern, one cudaDeviceSynchronize(). 512 MiB per step.
Free VRAM reaches 0 at step 14; the host goes down at step 16-18, holding 7.0-8.5 GiB.
The kernel is a faithful transcription of the one torch.cat actually launches — identified by profiling, not by reading the source:
and a --verify mode checks every output cell against what cat(dim=1) must produce: 1 147 720 cells, 0 unwritten, 0 wrong.
flag
load instruction
result
--vec-bytes 16
LDG.E.128
host bugchecks, 5 of 5
--vec-bytes 8
LDG.E.64
host bugchecks, 1 of 1
--vec-bytes 4
LDG.E.32
survives 40 steps / 20 GiB, 2 of 2
Everything else is identical, including — measured — the throughput: the crashing and the surviving
arm both do the first 15 steps (7.5 GiB) in 0.4 s. So it is not pressure, not parallelism (52
blocks × 128 threads in both), and not bandwidth. It is the width of the load.
Also eliminated, one variable at a time: PyTorch in the process, cuRAND Philox in the fill, block
count (832 vs 52), and whether the vector straddles an output row (70 % 4 = 2 vs 72 % 4 = 0,
byte volume held constant). All irrelevant — both sides crash.
The transcription compiles to the same code, checked in the SASS.cuobjdump -sass shows
exactly one wide load per instantiation — LDG.E.128 for --vec-bytes 16, LDG.E.64 for 8,
plain 32-bit LDG.E for 4 — so the flag really does change the instruction and not just a
constant. And the 16-byte instantiation compiles to 31 registers/thread, which is the same
figure nsys reports for PyTorch's own CatArrayBatchedCopy_alignedK_contig on the real
workload.
Optional: reproducing it through PyTorch
This is how I found it, and it is worth running once because it shows the defect hits a real
framework workload, not only a hand-written kernel. It is not the path to bisect on: PyTorch
gives you no way to switch the load width, so you cannot flip the behaviour the way the .cu above
can.
torch.cat on four contiguous, 16-byte-aligned fp32 tensors dispatches to exactly the kernel
transcribed above — CatArrayBatchedCopy_alignedK_contig<…, align=16> — which is what the profile
shows. Same five allocations per step, same 512 MiB, same crash point.
#!/usr/bin/env python3# WSL2 with GPU-P, Windows kernel 26100 / build 26200.9168, NVIDIA driver 597.06,# RTX PRO 2000 Blackwell Laptop GPU 8151 MiB, torch 2.11.0+cu128 (CUDA 12.8 runtime).## Expected: torch.cuda.OutOfMemoryError, the process exits, the host is fine.# Observed: once mem_get_info() reports 0 free, the WINDOWS HOST bugchecks and reboots. The guest# gets no error and logs nothing.## Measured: 6 crashes in 6 PyTorch runs whose workload contained torch.cat. They died on step 14,# 15, 15 and 17; THIS exact minimal form (randn x4 + cat, nothing else) accounts for one of them --# it died on step 17, after step 16 took 19.9 s and RECOVERED.## NOTE: this is the OPTIONAL path. The primary reproducer is repro_cat_alignedk.cu -- pure CUDA,# no PyTorch, and one flag (--vec-bytes) decides whether the host survives. Use this file to show# that the defect hits a real framework workload, not to bisect it.importos, torchlog=open('repro.log', 'w')
defL(msg): # fsync every line: a host bugcheck loses the page cacheprint(msg, flush=True); log.write(msg+'\n'); log.flush(); os.fsync(log.fileno())
rows=239674# 4 x (rows,70) fp32 = 256 MiB; cat's output is another 256 MiBheld= [] # nothing is ever freedforstepinrange(1, 41):
free, total=torch.cuda.mem_get_info()
st=torch.cuda.memory_stats()
L(f'step {step}: free {free>>20} MiB, allocated {torch.cuda.memory_allocated() >>20} MiB, 'f'cudaMalloc={st.get("num_device_alloc", 0)} cudaFree={st.get("num_device_free", 0)} 'f'retries={st.get("num_alloc_retries", 0)}')
try:
parts= [torch.randn(rows, 70, device='cuda') for_inrange(4)]
# If this is the last line in the log, it names the last call ENTERED -- not necessarily# the one that hung: it is printed before cat, and synchronize() follows two lines later.L(f'step {step}: cat <- LAST LINE = LAST CALL ENTERED')
X=torch.cat(parts, 1) # <-- ablation shows this is the necessary ingredientheld.append((parts, X))
torch.cuda.synchronize()
excepttorch.cuda.OutOfMemoryErrorase:
L(f'# CORRECT BEHAVIOUR: OutOfMemoryError at step {step}: {str(e)[:120]}')
breakL('# host survived')
6 crashes in 6 PyTorch runs whose workload contained torch.cat, dying on step 14, 15, 15
and 17. This exact minimal form accounts for one of them: it died on step 17, after step 16
took 19.9 s and recovered.
Allocator counters right up to the crash: num_device_alloc +5 per step, num_device_free = 0, num_alloc_retries = 0, num_sync_all_streams = 0, num_ooms = 0 —
the caching allocator only ever grows; there is no free/realloc churn involved.
⚠ Two honest limits on this path:
The last log line names the last call entered, not necessarily the one that hung — it is
printed before cat, and synchronize() follows two statements later.
Removing torch.cat and keeping everything else (same allocations, same bytes, same live-buffer
count) does not crash; but the arm that allocated the fifth block without writing it also
did not crash, so "the copy is what matters" is supported by the pure-CUDA ablation above rather
than by PyTorch alone.
★ One untested prediction you can falsify cheaply
parallel_cat picks the vectorised kernel only when every input passes is_aligned_vec4. So
feeding torch.cat inputs that are not 16-byte aligned — e.g. t[:, 1:] views rather than
freshly allocated tensors — should dispatch to the non-vectorised CatArrayBatchedCopy and,
if my ablation generalises, should not take the host down at the same point.
I have not run this. I am stating it because it is a cheap, falsifiable check on my own
conclusion, and if it fails, my conclusion is wrong.
Host side, and the dxgkrnl part
In 6 of 7 readable dumps — including two from the pure-CUDA reproducer — a thread inside nvlddmkm holds a heavily contended ERESOURCE while itself blocked, for 8.4 to 17.3 s
(measured: 8.375 · 8.625 · 8.890 · 11.109 · 11.250 · 17.296), on contention counts
1940 · 2344 · 2713 · 2943 · 8474 · 21995:
The other outer path, in two PyTorch dumps, is dxgkrnl!DxgkEscape ← win32u!NtGdiDdDDIEscape ← nvapi64+…, called from NVDisplay.Container.exe.
⚠ The part that would matter most for WSL is the weakest. In one dump the guest's own dxgkrnl!DXG_HOST_VIRTUALGPU_VMBUS::VmBusCreateAllocation thread — the host-side end of the
guest's CREATEALLOCATION — is WrResource-blocked in nt!ExAcquireResourceSharedLite behind
exactly that lock, for 2.421 s (thread Cid 0004.17f4, attached process vmwp.exe). In a second dump that thread is running, not queued. In four others there is
no such thread in the dump at all. So "the guest's allocation queues behind an NVIDIA lock" is n = 1 and I am flagging it rather than asserting it.
Up to eight of 16 processors end up in nt!KiIpiStallOnPacketTargetsPrcb; the bugcheck arrives by
NMI. PROCESS_NAME is a different innocent bystander in each dump (five distinct values), and FAILURE_BUCKET_ID names intelppm!HvRequestIdle, which is just the idle loop of the processors
that happen to be idle at NMI time.
The stall is measurable without crashing anything, and from outside the process
In one run, steps at zero free VRAM took 32.4 s and then 40.0 s and recovered before a
later step never returned. A separate process doing nothing but nvidia-smi --query-gpu=… once
per second logged gaps of 33 s and 45 s at exactly those moments — so the stall is a
property of the system, not of the submitting process's queue. The surviving --vec-bytes 4 arm
shows none of it (longest step 0.2 s, 40 steps in 3.0 s, sampler never misses a beat).
What I am not claiming
Not a mechanism — I have the correlation and the invariant, not the reason.
The guest CREATEALLOCATION chain is n = 1 (see above).
One dump contradicts the invariant: the 0x101 dump from the 52-block run reports 1 lock currently held with contention 2 and nonvlddmkm+0xfe1ab frame. So the two
bugcheck codes may not be two detectors of one state; the invariant is anchored on the 0x20001 class, where it holds 6 of 6.
A dxgmms2!VidSchiControlVSyncThread → nt!PpmAcquireLock chain appears in two dumps, but that
thread owns a different resource — it is not a waiter on the contended one.
The surviving arms are n ≤ 2; "scalar loads are safe" is not established.
The three BugcheckCode = 0 shutdowns wrote no dump. For one, the on-screen code was read as 0x101; for the others nothing was read — so that class is a write failure, not a separate
failure mode.
Five of the eleven shutdowns have no known workload.
Two notes for anyone reproducing this
nsys hangs on WSL2 on any CUDA program, even a 30-line one — it is CPU sampling, not CUDA
(nsys status --environment: CPU Profiling Environment (system-wide): Fail). Use --sample=none --cpuctxsw=none. Without it you cannot see which kernel actually runs, and I
spent nine transcriptions reproducing the wrong one.
Log with fsync per line — a host bugcheck takes the WSL2 VM down with its page cache, so a
buffered log loses exactly the lines that matter.
Happy to supply the reproducer, the run logs, and the kernel dumps.
Another machine reproducing the same hypervisor IOMMU bugcheck (0x20001, Arg1=0x28) under WSL2 GPU-PV, with a kernel memory dump
I have five Windows host crashes/hangs in four days, all while running CUDA (PyTorch) inside WSL2. Two of them are HYPERVISOR_ERROR 0x20001 with bit-identical arguments(0x28, 0x1, 0x29b92701, 0xfc810000), i.e. the same internal assertion in the hypervisor's I/O MMU module; two are CLOCK_WATCHDOG_TIMEOUT 0x101; one is a hang with no bugcheck recorded. This matches the pattern in this issue (0x20001 with Arg1=0x28, 0x101, and no-dump hangs). The fifth crash has a kernel memory dump; its stack shows the hypervisor delivering an NMI to an idle root-partition processor, with no Windows driver on the stack.
Environment
HP OMEN MAX 16-ah0xxx laptop, Intel Core Ultra 9 275HX (24 logical processors), 32 GB RAM, NVIDIA GeForce RTX 5080 Laptop GPU (16 GB), BIOS F.23 (2026-07-28)
Windows 11 Pro 26H2, wsl --version reports Windows version 10.0.26300.9550; VBS running with HVCI (HalpHvciEnabled: 1 in the dump); Hyper-V scheduler type 0x4 (event log)
Likelihood evaluation of small OLMo-architecture checkpoints (allenai/DataDecide-*, loaded with hf_olmo), fp32, batches of 16 sequences of about 300 tokens, computing [16, ~300, 50304] fp32 logits. With a 300M model the CUDA caching allocator takes the card to about 15.9 of 16.3 GB. Models are downloaded and evaluated one at a time; the GPU is otherwise idle.
Crashes
#
Time
Bugcheck
Driver
What was running
Dump
1
2026-10-11 23:45
0x20001 (0x28, 0x1, 0x29b92701, 0xfc810000)
610.62
~45 GB of HF downloads plus loading a 150M model for a 50-pair smoke test
minidump
2
2026-10-14 ~00:45
none recorded (host hung; forced power-off)
610.62
evaluating 4M–10M models in sequence (37th model)
—
3
2026-10-14 09:35
0x101 (no dump written; Kernel-Power 41 has BugcheckCode=0)
610.62
a 10M model, ~70 s after GPU work started on a fresh boot
—
4
2026-10-14 17:14
0x101 (0x8, 0x0, PRCB, 0x17) — processor 23 hung
617.42
first 300M model
minidump
5
2026-10-15 09:18
0x20001 (0x28, 0x1, 0x29b92701, 0xfc810000)
617.42
third 300M model; VRAM 15,980 / 16,303 MiB
kernel memory dump (4.2 GB) + minidump
Correlation with dxgkio_make_resident: Ioctl failed: -12 in the WSL kernel log
Under 610.62 the -12 line appeared on every model load including 4M ones, and crashes #1 and #3 followed such lines. In crash #5 the sequence was: -12 at 09:18:19, last VM log line 09:18:28, Windows EventLog 6008 "previous shutdown at 09:18" — i.e. the hypervisor error hit within ten seconds of the residency failure. VRAM sampled every 15 s before crash #5: 9,908 → 11,902 → 15,980 MiB (100 % utilisation, 68 °C, host MemAvailable 29 GB). Under 617.42 the VM log also shows a constant stream of dxgkio_escape: Ioctl failed: -22 (about 80 per minute from VM boot, independent of load); I do not know whether that is related.
HYPERVISOR_ERROR (20001)
Arg1: 0000000000000028, An internal error was detected in the I/O MMU module.
Arg2: 0000000000000001 Arg3: 0000000029b92701 Arg4: 00000000fc810000
Key : Stack.Pointer Value: NMI
Key : SecureKernel.HalpHvciEnabled Value: 1
PROCESS_NAME: System FAULTING processor: 15 (idle)
nt!KeBugCheckEx
nt!HvlSkCrashdumpCallbackRoutine+0x8a
nt!KiProcessNMI+0xb1
nt!KxNmiInterrupt+0x82
nt!KiNmiInterrupt+0x26e
nt!HvcallInitiateHypercall+0x68
nt!HvlRequestProcessorHalt+0x26
intelppm!HvRequestIdle+0x2d
intelppm!PepIdleExecute+0x2d
nt!PpmIdleExecuteTransition+0x518
nt!PoIdle+0x190
nt!KiIdleLoop+0x54
FAILURE_BUCKET_ID: 0x20001_28_1_intelppm!HvRequestIdle
The root partition was idle when the hypervisor raised the NMI; intelppm!HvRequestIdle is just where the idle loop was interrupted. Crash #4's minidump is CLOCK_WATCHDOG_TIMEOUT_nt!KeAccumulateTicks with the hung processor's stack unavailable (triage dump), which is why I switched to kernel dumps. Both dumps (and the 10-11 minidump with the same 0x20001 arguments) are available on request; I would rather not attach a 4.2 GB memory dump publicly.
Ruled out so far
NVIDIA driver update 610.62 → 617.42: -12 disappeared for small models but not for 300M models; crashes continued.
Not thermal (GPU 68 °C), not host RAM pressure (29 GB available), no WHEA errors except corrected PCIe AER on the Wi-Fi NIC at every boot, no OOM in the VM.
Hardware: no machine-check errors logged; the 0x20001 arguments are identical across two crashes four days apart, which looks like a software assertion rather than a failing core.
Reproduction attempts
The real workload reproduced it on 2 of 4 runs at the 300M model size (it is my evaluation pipeline: public allenai/DataDecide-*-300M checkpoints loaded with hf_olmo, fp32, batches of 16 padded texts from public sources, three tests in a row per model in a fresh process each time — about 20,000 limerick-length texts (~60 tokens), then 4,000 passages of ~170 tokens, then 4,000 of ~300 tokens; the -12 line and the crash came about 30 s into the third test both times). I can share the exact scripts and data privately.
A standalone script that mimics the workload (same model, random tokens, same batch shape) did not reproduce it in four 15-minute runs: fixed 300-token batches (reserved 5.9 GB); lengths uniform in 50–400 (reserved 15.0 GB, PyTorch reporting 0 MiB free for over ten minutes); alternating 60/400-token phases (7.8 GB); and a replay of the real test order with jittered lengths (11.8 GB). No make_resident -12 was logged in any of them. So on this machine "free VRAM reaches zero" by itself is not sufficient; whatever the real pipeline does differently (process-per-model with the previous CUDA context just torn down, real text length distribution, tokenizer path) matters. The script is below in case it helps someone else; the -12 kernel line was the earliest warning in every crash here.
wzgrx As described in the thread, repeated CUDA allocation failures in WSL are followed by a Windows 11 host BSOD, rather than a recoverable out-of-memory error. Your guest log contains a warning in dxgvmb_send_wait_sync_object_gpu , but its connection to the crash is unclear. The root cause is not established by the supplied attachments.
Please share the Windows kernel crash dump and corresponding BugCheck/WER event records, together with the reproduction command, timestamped application logs, exact crash time, and Windows/WSL/NVIDIA driver versions. Use a private upload channel for dumps and traces.
For additional host graphics tracing, run these checks in elevated Windows PowerShell:
wpr.exe -profiles
wpr.exe -status
If GPU is listed and no WPR recording is active, start before reproduction:
wpr.exe -start GPU -filemode
If Windows remains responsive, save the trace:
wpr.exe -stop "$env:USERPROFILE\Desktop\host-gpu.etl"
Please share the resulting ETL. A BSOD may prevent WPR from finalizing the trace, so the crash dump remains the priority. File-mode tracing can grow large; keep the recording as short as practical and monitor free disk space.
ksingara I'm not the OP, but I have everything in your list from a second machine (my comment above, Oct 15), plus six more crashes since then that change the picture. Summary first, details below.
What changed since Oct 15. Eleven host crashes in total on this machine between Oct 11 and Oct 20, all under sustained CUDA load from a PyTorch evaluation pipeline. Three of them happened with no GPU work inside WSL at all (native Windows CUDA, WSL idle), one happened with VT-x disabled in firmware (no hypervisor: the bugcheck then becomes CLOCK\_WATCHDOG\_TIMEOUT 0x101 instead of HYPERVISOR\_ERROR 0x20001), and one happened after opting the GPU and NIC drivers out of DMA remapping (DmaRemappingCompatible=0 for nvlddmkm and rt25cx21), which only changed the IOMMU internal error code (Arg1 0x28 → 0x13). So on this machine WSL2 GPU-PV is the heaviest trigger but not the cause; the common factor is the hypervisor's IOMMU path (DMA remapping and interrupt remapping) under heavy device DMA, and in both kernel dumps I have analysed (crashes 6 and 11) USB xHCI DMA map/unmap hypercalls were in flight at the moment of the crash, whatever the GPU was doing. The dxgkio\_make\_resident -12 correlation in my first comment is a symptom of the same pressure, not the root cause. I stopped running GPU work on this machine on Oct 20 and moved it to a cloud GPU.
Environment (HP OMEN MAX 16-ah0xxx, Core Ultra 9 275HX, RTX 5080 Laptop 16 GB, BIOS Insyde F.23 2026-07-28; VBS + HVCI on, Kernel DMA Protection On, Secure Boot on; VT-d cannot be disabled in this firmware)
Windows 11 Pro 26H2, build 26300.9550 (unchanged across all crashes)
WSL 3.0.1.0, kernel 6.18.40.1-1 at the time (3.0.3.0 / 6.18.54.1-1 now); WSLg 1.0.79; Direct3D 1.611.1-81528511; DXCore 10.0.26100.1
NVIDIA 610.62 (crashes 1–3), 617.42 (crashes 4–11), 621.01 now; CUDA 13.0 in WSL, torch 2.14.1+cu130
Third-party kernel drivers in the DMA path: VMware Workstation 25.0.0 is installed and its USB arbitration driver hcmon.sys 8.11.16.0 (Broadcom) is running; it appears in the crash crash 11 stack (see below). USB xHCI carries Bluetooth (keyboard and mouse), the internal and an external webcam, and a dozen HIDs; no USB storage.
WER BugCheck (Event 1001) records for crashes 7–11 are in my System log (it was cleared on 10-20 01:17, so crashes 1–6 are gone from the log; their dumps are not). I can paste the records or attach them.
What the kernel dumps show
crash 5 (!analyze -v): 0x20001\_28\_1, "An internal error was detected in the I/O MMU module"; stack is the NMI-driven HvlSkCrashdumpCallbackRoutine, the interrupted thread is an idle processor in intelppm!HvRequestIdle; HalpHvciEnabled: 1. No Windows driver on the stack.
crash 6 (!running -it): 20 of 24 processors idle; the two busy ones are both inside an IOMMU unmap hypercall: CPU 0 USBXHCI!Bulk\_Stage\_FreeScatterGatherList → nt!HalPutScatterGatherListDmarThin → nt!IommuUnmapLogicalRange → nt!HvlDmaUnmapDeviceLogicalRange → nt!HvcallpExtendedFastHypercall; CPU 18 netadaptercx!NxDmaAdapter::PutScatterGatherList → … → nt!HvlDmaUnmapDeviceLogicalRange → nt!HvcallpExtendedFastHypercall (Realtek 2.5GbE, rt25cx21x64.sys 1125.26.710.2025). No nvlddmkm / dxgkrnl frames anywhere.
crash 9 (no hypervisor): clock watchdog with the hung processor in ExecutionContext!PollWorkerEntry; i.e. with Hyper-V out of the way the same hang surfaces as a plain CPU watchdog.
crash 11 (!analyze -v; !running -it, kernel dump): Arg1 0x13 = "The hypervisor has been spinning excessively" (the other seven are 0x28, "internal error in the I/O MMU module"). The NMI landed on CPU 11 inside a device-interrupt retarget hypercall: nt!HvlRetargetDeviceInterrupt → HalpHvRetargetDeviceMsiInterrupt ← HalpInterruptSetRemappedDestinationHv ← KiIntSteerDistributeInterrupts ← PpmParkSteerInterrupts (core parking re-steering MSIs), a DPC that interrupted an nvlddmkm thread of NVDisplay.Container.exe. At the same instant CPU 12 was in an IOMMU map hypercall (HvlDmaMapDeviceLogicalRange ← HalGetScatterGatherListDmarThin ← USBXHCI!Control\_Transfer\_Map ← ucx01000 ← hcmon.sys) and CPU 13 in an IOMMU unmap hypercall (HvlDmaUnmapDeviceLogicalRange ← HalPutScatterGatherListDmarThin ← USBXHCI!Bulk\_Stage\_FreeScatterGatherList); CPU 20 was switching to VTL1 (HVCI); four more CPUs were stalled in KeFlushProcessWriteBuffers IPIs waiting for those. So with the GPU and NIC opted out of DMA remapping, the hypercalls in flight were USB xHCI DMA map/unmap plus interrupt remapping, and the hypervisor reported itself spinning instead of an I/O MMU internal error. hcmon.sys in the CPU 12 stack is VMware Workstation's USB arbitration driver (8.11.16.0), which sits in the USB control-transfer path on this machine; I have not yet tested with it disabled.
Reproduction. The real workload reproduces it within minutes to hours: python -m litq.eval.batch over public allenai/DataDecide-\* checkpoints (300M and 1B; hf\_olmo, fp32, batches of 16 texts of 60–300 tokens, one fresh process per model). A standalone script with the same model and batch shape did not reproduce it in four 15-minute runs (details in my Oct 15 comment). I can give private access to the exact repository, job list and logs; every crash is bracketed by timestamped application logs and nvidia-smi samples at 15 s.
WPR. I have not captured an ETL: with a BSOD the trace never finalises, as you note, and each reproduction risks the disk. If a GPU trace is actually useful to you, tell me and I will run wpr.exe -start GPU -filemode once before a deliberate reproduction on this machine.
What I need from you: a private upload channel for the three full kernel dumps (4.5 GB, 3.0 GB, 5.0 GB) and the five minidumps. The Feedback Hub report https://aka.ms/AA13us97 already has the Oct 11–15 minidumps.
Activity
github-actions commentedon Jun 6, 2026
Logs are required for review from WSL team
If this a feature request, please reply with '/feature'. If this is a question, reply with '/question'.
Otherwise, please attach logs by following the instructions below, your issue will not be reviewed unless they are added. These logs will help us understand what is going on in your machine.
How to collect WSL logs
Download and execute collect-wsl-logs.ps1 in an administrative powershell prompt:
The script will output the path of the log file once done.
If this is a networking issue, please use
.\collect-wsl-logs.ps1 -LogProfile networkinginstead, following the instructions in Collect WSL logs for networking issuesOnce completed please upload the output files to this GitHub issue.
See Collect WSL logs (recommended method).
If you choose to email these logs instead of attaching them to the bug, please send them to wsl-gh-logs@microsoft.com with the GitHub issue number in the subject, and include a link to your GitHub issue comment in the message body, and reply with '/emailed-logs'.
wzgrx commentedon Jun 6, 2026
01-wsl-dmesg.log
02-system-info.log
03-gateway-errors.log
04-vision-errors.log
05-crash-timeline.txt
WSL VM-death fix: never set expandable_segments; buffer hygiene + VRA…
stalek71 commentedon Aug 31, 2026
Same symptom, with a ~250-line pure-CUDA reproducer: a guest CUDA process bugchecks the Windows host, and one flag flips it
This issue describes the symptom I am seeing. Adding what I could pin down, in case it helps —
the important part is that it is now reducible to a single
.cufile with no PyTorch and nolibraries beyond
libcudart, and that one command-line flag decides whether the hostsurvives.
What happens
A CUDA process inside WSL2 oversubscribes VRAM and keeps allocating. Free VRAM reaches 0, WDDM
starts paging, and a few steps later the Windows host bugchecks and reboots. From inside the
guest there is nothing: no
cudaErrorMemoryAllocation, no exception, no kernel message, nodmesgentry. The log simply stops after the last line it managed to flush — which names the lastcall entered, not necessarily the one that hung (the line is printed before the launch, and a
cudaDeviceSynchronize()follows two statements later).For a GPU-P guest this is a containment problem regardless of whose driver stalls: an unprivileged
process in the VM can take the host down, and gets no error it could handle.
Environment
/dev/dxg)CUDA 12.8 runtime (
torch 2.11.0+cu128)CrashDumpEnabled = 2(kernel dump)11 unexpected shutdowns
0x20001HYPERVISOR_ERRORArg1=0x28(I/O MMU module),Arg2=1,Arg3=0x29b92701,Arg4=0xfc801000(×5) /0xfc800000(×1)0x101CLOCK_WATCHDOG_TIMEOUTArg1=0xc;Arg4(hung processor) = 9, then 5BugcheckCode = 0, no dump written⚠ Five of these eleven have no known workload and I am not attributing them. The reproducer
below accounts for six more crashes deliberately produced afterwards.
There is not a single
Display(EventID 4101, TDR) ornvlddmkmevent in the System log — sothis is not a display-driver reset; the host dies before TDR reports anything. That query and its
positive control are in
EVIDENCE_no_display_driver_events.txt.The reproducer, and the flag that flips it
Five allocations per step, nothing ever freed: 4 sources of 64 MiB (
(239674,70)fp32) plus onedestination of 256 MiB (
(239674,280)fp32), one kernel reading all four and writing thedestination with a concat-along-dim-1 pattern, one
cudaDeviceSynchronize(). 512 MiB per step.Free VRAM reaches 0 at step 14; the host goes down at step 16-18, holding 7.0-8.5 GiB.
The kernel is a faithful transcription of the one
torch.catactually launches — identified byprofiling, not by reading the source:
and a
--verifymode checks every output cell against whatcat(dim=1)must produce:1 147 720 cells, 0 unwritten, 0 wrong.
--vec-bytes 16LDG.E.128--vec-bytes 8LDG.E.64--vec-bytes 4LDG.E.32Everything else is identical, including — measured — the throughput: the crashing and the surviving
arm both do the first 15 steps (7.5 GiB) in 0.4 s. So it is not pressure, not parallelism (52
blocks × 128 threads in both), and not bandwidth. It is the width of the load.
Also eliminated, one variable at a time: PyTorch in the process, cuRAND Philox in the fill, block
count (832 vs 52), and whether the vector straddles an output row (
70 % 4 = 2vs72 % 4 = 0,byte volume held constant). All irrelevant — both sides crash.
The transcription compiles to the same code, checked in the SASS.
cuobjdump -sassshowsexactly one wide load per instantiation —
LDG.E.128for--vec-bytes 16,LDG.E.64for8,plain 32-bit
LDG.Efor4— so the flag really does change the instruction and not just aconstant. And the 16-byte instantiation compiles to 31 registers/thread, which is the same
figure
nsysreports for PyTorch's ownCatArrayBatchedCopy_alignedK_contigon the realworkload.
Optional: reproducing it through PyTorch
This is how I found it, and it is worth running once because it shows the defect hits a real
framework workload, not only a hand-written kernel. It is not the path to bisect on: PyTorch
gives you no way to switch the load width, so you cannot flip the behaviour the way the
.cuabovecan.
torch.caton four contiguous, 16-byte-aligned fp32 tensors dispatches to exactly the kerneltranscribed above —
CatArrayBatchedCopy_alignedK_contig<…, align=16>— which is what the profileshows. Same five allocations per step, same 512 MiB, same crash point.
Measured behaviour, from the run logs:
torch 2.11.0+cu128(CUDA 12.8 runtime), 8151 MiB card, 7033 MiB free at starttorch.cat, dying on step 14, 15, 15and 17. This exact minimal form accounts for one of them: it died on step 17, after step 16
took 19.9 s and recovered.
num_device_alloc+5 per step,num_device_free = 0,num_alloc_retries = 0,num_sync_all_streams = 0,num_ooms = 0—the caching allocator only ever grows; there is no free/realloc churn involved.
⚠ Two honest limits on this path:
printed before
cat, andsynchronize()follows two statements later.torch.catand keeping everything else (same allocations, same bytes, same live-buffercount) does not crash; but the arm that allocated the fifth block without writing it also
did not crash, so "the copy is what matters" is supported by the pure-CUDA ablation above rather
than by PyTorch alone.
★ One untested prediction you can falsify cheaply
parallel_catpicks the vectorised kernel only when every input passesis_aligned_vec4. Sofeeding
torch.catinputs that are not 16-byte aligned — e.g.t[:, 1:]views rather thanfreshly allocated tensors — should dispatch to the non-vectorised
CatArrayBatchedCopyand,if my ablation generalises, should not take the host down at the same point.
I have not run this. I am stating it because it is a cheap, falsifiable check on my own
conclusion, and if it fails, my conclusion is wrong.
Host side, and the
dxgkrnlpartIn 6 of 7 readable dumps — including two from the pure-CUDA reproducer — a thread inside
nvlddmkmholds a heavily contendedERESOURCEwhile itself blocked, for 8.4 to 17.3 s(measured: 8.375 · 8.625 · 8.890 · 11.109 · 11.250 · 17.296), on contention counts
1940 · 2344 · 2713 · 2943 · 8474 · 21995:
The other outer path, in two PyTorch dumps, is
dxgkrnl!DxgkEscape←win32u!NtGdiDdDDIEscape←nvapi64+…, called fromNVDisplay.Container.exe.⚠ The part that would matter most for WSL is the weakest. In one dump the guest's own
dxgkrnl!DXG_HOST_VIRTUALGPU_VMBUS::VmBusCreateAllocationthread — the host-side end of theguest's CREATEALLOCATION — is
WrResource-blocked innt!ExAcquireResourceSharedLitebehindexactly that lock, for 2.421 s (thread
Cid 0004.17f4, attached processvmwp.exe). In a second dump that thread is running, not queued. In four others there isno such thread in the dump at all. So "the guest's allocation queues behind an NVIDIA lock" is
n = 1 and I am flagging it rather than asserting it.
Up to eight of 16 processors end up in
nt!KiIpiStallOnPacketTargetsPrcb; the bugcheck arrives byNMI.
PROCESS_NAMEis a different innocent bystander in each dump (five distinct values), andFAILURE_BUCKET_IDnamesintelppm!HvRequestIdle, which is just the idle loop of the processorsthat happen to be idle at NMI time.
The stall is measurable without crashing anything, and from outside the process
In one run, steps at zero free VRAM took 32.4 s and then 40.0 s and recovered before a
later step never returned. A separate process doing nothing but
nvidia-smi --query-gpu=…onceper second logged gaps of 33 s and 45 s at exactly those moments — so the stall is a
property of the system, not of the submitting process's queue. The surviving
--vec-bytes 4armshows none of it (longest step 0.2 s, 40 steps in 3.0 s, sampler never misses a beat).
What I am not claiming
0x101dump from the 52-block run reports1 lock currently heldwith contention 2 and nonvlddmkm+0xfe1abframe. So the twobugcheck codes may not be two detectors of one state; the invariant is anchored on the
0x20001class, where it holds 6 of 6.dxgmms2!VidSchiControlVSyncThread→nt!PpmAcquireLockchain appears in two dumps, but thatthread owns a different resource — it is not a waiter on the contended one.
BugcheckCode = 0shutdowns wrote no dump. For one, the on-screen code was read as0x101; for the others nothing was read — so that class is a write failure, not a separatefailure mode.
Two notes for anyone reproducing this
nsyshangs on WSL2 on any CUDA program, even a 30-line one — it is CPU sampling, not CUDA(
nsys status --environment: CPU Profiling Environment (system-wide): Fail). Use--sample=none --cpuctxsw=none. Without it you cannot see which kernel actually runs, and Ispent nine transcriptions reproducing the wrong one.
fsyncper line — a host bugcheck takes the WSL2 VM down with its page cache, so abuffered log loses exactly the lines that matter.
Happy to supply the reproducer, the run logs, and the kernel dumps.
repro_cat_alignedk.zip
stalek71 commentedon Sep 3, 2026
Issue description on NVIDIA forum
stalek71 commentedon Sep 23, 2026
Has anyone looked into this and started working on a fix?
QinBojia commentedon Oct 15, 2026
Another machine reproducing the same hypervisor IOMMU bugcheck (0x20001, Arg1=0x28) under WSL2 GPU-PV, with a kernel memory dump
I have five Windows host crashes/hangs in four days, all while running CUDA (PyTorch) inside WSL2. Two of them are
HYPERVISOR_ERROR 0x20001with bit-identical arguments(0x28, 0x1, 0x29b92701, 0xfc810000), i.e. the same internal assertion in the hypervisor's I/O MMU module; two areCLOCK_WATCHDOG_TIMEOUT 0x101; one is a hang with no bugcheck recorded. This matches the pattern in this issue (0x20001 with Arg1=0x28, 0x101, and no-dump hangs). The fifth crash has a kernel memory dump; its stack shows the hypervisor delivering an NMI to an idle root-partition processor, with no Windows driver on the stack.Environment
wsl --versionreports Windows version 10.0.26300.9550; VBS running with HVCI (HalpHvciEnabled: 1in the dump); Hyper-V scheduler type 0x4 (event log)acornnpm package (inability to symlink on mounted drives) #3 happened on 610.62; updating did not stop them); CUDA 13.0 runtime in WSL, torch 2.14.1+cu130, transformers 4.57.6, ai2-olmo 0.6.0.wslconfig:[wsl2] networkingMode=mirrored, dnsTunneling=true, firewall=true, autoProxy=true, memory=32GB, swap=96GB(noprocessorslimit)Workload
Likelihood evaluation of small OLMo-architecture checkpoints (
allenai/DataDecide-*, loaded withhf_olmo), fp32, batches of 16 sequences of about 300 tokens, computing[16, ~300, 50304]fp32 logits. With a 300M model the CUDA caching allocator takes the card to about 15.9 of 16.3 GB. Models are downloaded and evaluated one at a time; the GPU is otherwise idle.Crashes
Correlation with
dxgkio_make_resident: Ioctl failed: -12in the WSL kernel logUnder driver 617.42:
make_resident -12loggedUnder 610.62 the
-12line appeared on every model load including 4M ones, and crashes #1 and #3 followed such lines. In crash #5 the sequence was:-12at 09:18:19, last VM log line 09:18:28, WindowsEventLog 6008"previous shutdown at 09:18" — i.e. the hypervisor error hit within ten seconds of the residency failure. VRAM sampled every 15 s before crash #5: 9,908 → 11,902 → 15,980 MiB (100 % utilisation, 68 °C, host MemAvailable 29 GB). Under 617.42 the VM log also shows a constant stream ofdxgkio_escape: Ioctl failed: -22(about 80 per minute from VM boot, independent of load); I do not know whether that is related.Kernel dump of crash #5 (
!analyze -v, abridged)The root partition was idle when the hypervisor raised the NMI;
intelppm!HvRequestIdleis just where the idle loop was interrupted. Crash #4's minidump isCLOCK_WATCHDOG_TIMEOUT_nt!KeAccumulateTickswith the hung processor's stack unavailable (triage dump), which is why I switched to kernel dumps. Both dumps (and the 10-11 minidump with the same 0x20001 arguments) are available on request; I would rather not attach a 4.2 GB memory dump publicly.Ruled out so far
-12disappeared for small models but not for 300M models; crashes continued.Reproduction attempts
The real workload reproduced it on 2 of 4 runs at the 300M model size (it is my evaluation pipeline: public
allenai/DataDecide-*-300Mcheckpoints loaded withhf_olmo, fp32, batches of 16 padded texts from public sources, three tests in a row per model in a fresh process each time — about 20,000 limerick-length texts (~60 tokens), then 4,000 passages of ~170 tokens, then 4,000 of ~300 tokens; the-12line and the crash came about 30 s into the third test both times). I can share the exact scripts and data privately.A standalone script that mimics the workload (same model, random tokens, same batch shape) did not reproduce it in four 15-minute runs: fixed 300-token batches (reserved 5.9 GB); lengths uniform in 50–400 (reserved 15.0 GB, PyTorch reporting 0 MiB free for over ten minutes); alternating 60/400-token phases (7.8 GB); and a replay of the real test order with jittered lengths (11.8 GB). No
make_resident -12was logged in any of them. So on this machine "free VRAM reaches zero" by itself is not sufficient; whatever the real pipeline does differently (process-per-model with the previous CUDA context just torn down, real text length distribution, tokenizer path) matters. The script is below in case it helps someone else; the-12kernel line was the earliest warning in every crash here.Feedback Hub report with the three minidumps attached: https://aka.ms/AA13us97 (the 4.2 GB kernel dump is available on request).
Happy to run further diagnostics (ETW traces, a specific
.wslconfig, driver versions) if that helps.ksingara commentedon Oct 21, 2026
wzgrx As described in the thread, repeated CUDA allocation failures in WSL are followed by a Windows 11 host BSOD, rather than a recoverable out-of-memory error. Your guest log contains a warning in dxgvmb_send_wait_sync_object_gpu , but its connection to the crash is unclear. The root cause is not established by the supplied attachments.
Please share the Windows kernel crash dump and corresponding BugCheck/WER event records, together with the reproduction command, timestamped application logs, exact crash time, and Windows/WSL/NVIDIA driver versions. Use a private upload channel for dumps and traces.
For additional host graphics tracing, run these checks in elevated Windows PowerShell:
wpr.exe -profiles
wpr.exe -status
If GPU is listed and no WPR recording is active, start before reproduction:
wpr.exe -start GPU -filemode
If Windows remains responsive, save the trace:
wpr.exe -stop "$env:USERPROFILE\Desktop\host-gpu.etl"
Please share the resulting ETL. A BSOD may prevent WPR from finalizing the trace, so the crash dump remains the priority. File-mode tracing can grow large; keep the recording as short as practical and monitor free disk space.
QinBojia commentedon Oct 28, 2026
ksingara I'm not the OP, but I have everything in your list from a second machine (my comment above, Oct 15), plus six more crashes since then that change the picture. Summary first, details below.
What changed since Oct 15. Eleven host crashes in total on this machine between Oct 11 and Oct 20, all under sustained CUDA load from a PyTorch evaluation pipeline. Three of them happened with no GPU work inside WSL at all (native Windows CUDA, WSL idle), one happened with VT-x disabled in firmware (no hypervisor: the bugcheck then becomes
CLOCK\_WATCHDOG\_TIMEOUT 0x101instead ofHYPERVISOR\_ERROR 0x20001), and one happened after opting the GPU and NIC drivers out of DMA remapping (DmaRemappingCompatible=0fornvlddmkmandrt25cx21), which only changed the IOMMU internal error code (Arg10x28→0x13). So on this machine WSL2 GPU-PV is the heaviest trigger but not the cause; the common factor is the hypervisor's IOMMU path (DMA remapping and interrupt remapping) under heavy device DMA, and in both kernel dumps I have analysed (crashes 6 and 11) USB xHCI DMA map/unmap hypercalls were in flight at the moment of the crash, whatever the GPU was doing. Thedxgkio\_make\_resident -12correlation in my first comment is a symptom of the same pressure, not the root cause. I stopped running GPU work on this machine on Oct 20 and moved it to a cloud GPU.Environment (HP OMEN MAX 16-ah0xxx, Core Ultra 9 275HX, RTX 5080 Laptop 16 GB, BIOS Insyde F.23 2026-07-28; VBS + HVCI on, Kernel DMA Protection On, Secure Boot on; VT-d cannot be disabled in this firmware)
hcmon.sys8.11.16.0 (Broadcom) is running; it appears in the crash crash 11 stack (see below). USB xHCI carries Bluetooth (keyboard and mouse), the internal and an external webcam, and a dozen HIDs; no USB storage.All crashes (local time, UTC+8)
CLOCK\_WATCHDOG\_TIMEOUT\_ExecutionContext!PollWorkerEntryDmaRemappingCompatible=0for nvlddmkm + rt25cx21WER
BugCheck(Event 1001) records for crashes 7–11 are in my System log (it was cleared on 10-20 01:17, so crashes 1–6 are gone from the log; their dumps are not). I can paste the records or attach them.What the kernel dumps show
!analyze -v):0x20001\_28\_1, "An internal error was detected in the I/O MMU module"; stack is the NMI-drivenHvlSkCrashdumpCallbackRoutine, the interrupted thread is an idle processor inintelppm!HvRequestIdle;HalpHvciEnabled: 1. No Windows driver on the stack.!running -it): 20 of 24 processors idle; the two busy ones are both inside an IOMMU unmap hypercall: CPU 0USBXHCI!Bulk\_Stage\_FreeScatterGatherList → nt!HalPutScatterGatherListDmarThin → nt!IommuUnmapLogicalRange → nt!HvlDmaUnmapDeviceLogicalRange → nt!HvcallpExtendedFastHypercall; CPU 18netadaptercx!NxDmaAdapter::PutScatterGatherList → … → nt!HvlDmaUnmapDeviceLogicalRange → nt!HvcallpExtendedFastHypercall(Realtek 2.5GbE, rt25cx21x64.sys 1125.26.710.2025). Nonvlddmkm/dxgkrnlframes anywhere.ExecutionContext!PollWorkerEntry; i.e. with Hyper-V out of the way the same hang surfaces as a plain CPU watchdog.!analyze -v; !running -it, kernel dump): Arg10x13= "The hypervisor has been spinning excessively" (the other seven are0x28, "internal error in the I/O MMU module"). The NMI landed on CPU 11 inside a device-interrupt retarget hypercall:nt!HvlRetargetDeviceInterrupt → HalpHvRetargetDeviceMsiInterrupt ← HalpInterruptSetRemappedDestinationHv ← KiIntSteerDistributeInterrupts ← PpmParkSteerInterrupts(core parking re-steering MSIs), a DPC that interrupted annvlddmkmthread ofNVDisplay.Container.exe. At the same instant CPU 12 was in an IOMMU map hypercall (HvlDmaMapDeviceLogicalRange ← HalGetScatterGatherListDmarThin ← USBXHCI!Control\_Transfer\_Map ← ucx01000 ← hcmon.sys) and CPU 13 in an IOMMU unmap hypercall (HvlDmaUnmapDeviceLogicalRange ← HalPutScatterGatherListDmarThin ← USBXHCI!Bulk\_Stage\_FreeScatterGatherList); CPU 20 was switching to VTL1 (HVCI); four more CPUs were stalled inKeFlushProcessWriteBuffersIPIs waiting for those. So with the GPU and NIC opted out of DMA remapping, the hypercalls in flight were USB xHCI DMA map/unmap plus interrupt remapping, and the hypervisor reported itself spinning instead of an I/O MMU internal error.hcmon.sysin the CPU 12 stack is VMware Workstation's USB arbitration driver (8.11.16.0), which sits in the USB control-transfer path on this machine; I have not yet tested with it disabled.Reproduction. The real workload reproduces it within minutes to hours:
python -m litq.eval.batchover publicallenai/DataDecide-\*checkpoints (300M and 1B;hf\_olmo, fp32, batches of 16 texts of 60–300 tokens, one fresh process per model). A standalone script with the same model and batch shape did not reproduce it in four 15-minute runs (details in my Oct 15 comment). I can give private access to the exact repository, job list and logs; every crash is bracketed by timestamped application logs andnvidia-smisamples at 15 s.WPR. I have not captured an ETL: with a BSOD the trace never finalises, as you note, and each reproduction risks the disk. If a GPU trace is actually useful to you, tell me and I will run
wpr.exe -start GPU -filemodeonce before a deliberate reproduction on this machine.What I need from you: a private upload channel for the three full kernel dumps (4.5 GB, 3.0 GB, 5.0 GB) and the five minidumps. The Feedback Hub report https://aka.ms/AA13us97 already has the Oct 11–15 minidumps.
bugcheck_events_2026-10-27.txt