Xid 31 and the reserved top of VRAM

Xid 31 is the NVIDIA driver's GPU memory page fault: the MMU refused a memory access. On unlocked CMP 170HX cards it has two main sources. The first is the unlocker: early cmpunlocker builds handed the protected top of VRAM, which backs the GSP firmware's WPR2 region, to the driver's memory allocator, so ordinary allocations could land there and crash the card when VRAM was nearly full or under heavy allocation churn. That bug is fixed in cmpunlocker v0.3 (2026-08-30) and later. The second is the inference engine: int32 address overflows and oversized transient buffers in vLLM kernels raise Xid 31 at repeatable context lengths whatever unlocker you run. A short summary is on Memory errors and ECC; every Xid number is on Xid errors.

First checks after an Xid 31, each explained below:

Run cmpunlocker v0.3 or later, built from a fresh checkout. A failed git pull leaves the old late-PMA code in place without any error.
Read the kernel log from the last driver load. A build with the fix logs that the late-PMA extension was skipped; a late-PMA range that overlaps the firmware WPR range means the fix is missing.
Note the context length at the crash. An Xid 31 that comes back at the same depth or KV pool size is an engine bug: look for the engine's int64 patch.
Count the Xid lines with dmesg, then reboot the host before restarting the engine. Do not reload nvidia_uvm.
Test the card with cuda_memtest over 100% of memory. memtest_vulkan, and gpu-burn at its default 90% fill, do not reach the top blocks.

First signs

Allocation churn, 141 MB wrongly freed, and the last blocks cuda_memtest cannot write

Before the late-PMA fix was understood, several reports pointed at the top of the unlocked memory. SKYY (2026-07-29) ran a python3 ctypes loop that loads libcudart.so and calls cudaMalloc of 100 MiB (104857600 bytes) and cudaFree 100,000 times, printing every 1000 iterations; on one unlocked 8 GB card it consistently triggered Xid 31, and he later edited the post to say he had figured out the issue, without saying what it was. bayley_w told Amogh on 2026-08-02 that the late-PMA extension in the unlocker frees 141 MB of VRAM that should not be freed, and that under heavy allocation churn this reliably makes the GPUs fall off the bus. gym (2026-08-10) found that cuda_memtest at 100% of memory fails while memtest_vulkan does not, and that limiting cuda_memtest to about 63.75 GB works. It was the same on all his cards, so he concluded that the last few blocks are not writable rather than the VRAM being defective, and that other workloads leave about 200 MB free and never touch them. ju5t5omedude noted that gpu-burn fills only 90% of memory by default, which may be why it passes.

Reported 8 GB
Source: @skyy2333, 2026-07-29
Source: @bayley_w, 2026-08-02
Source: @ju5t5omedude, 2026-08-10
Source: @jimtjames, 2026-08-10

Cause: the late-PMA extension registered WPR2

PMA is the driver's physical memory allocator. WPR2 is a write-protected region at the top of VRAM that holds the GSP firmware's own heap. After the unlock grows the card's memory geometry, a "late PMA" step in early cmpunlocker builds took the highest reserved framebuffer region and registered it with the allocator, to win back about 141 MB. That region overlapped WPR2, so when VRAM was nearly full, or when allocation churn returned the wrong pages, ordinary CUDA allocations received protected pages and the first write to them raised Xid 31. The entries below are the evidence, in the order it came in.

The 65535 MB limit over-allocated the tail on 8 GB cards

On 2026-08-10 Jack reported that mainline cmpunlocker set the unlock limit to 65535 MB on the 8 GB card (64 GB unlocked), but a portion of the addresses at the tail are private to the GPU, so usable space is somewhat smaller; when a program hits OOM and writes into the tail, the card crashes. Far said this may explain why his 64 GB card crashed when running cuda_memtest over all memory more than once, while models never touched that memory. bayley_w confirmed a variant of the bug when vLLM thrashed pages during model loading, and gym saw it at 100% in cuda_memtest, "Consistent across all my cards, so it's not defective VRAM" (not in the Vulkan memtest). Amogh acknowledged it and said he would make another branch to test a fix. He added that 80 GB on the 10 GB card is a separate problem that needs a VBIOS patch, not only a memory geometry change.

Proven 8 GB
Source: @ludashi666, 2026-08-10
Source: @_j3ns, 2026-08-10
Source: @farawaythrower, 2026-08-10
Source: @bayley_w, 2026-08-10
Source: @jimtjames, 2026-08-10

The late-PMA extension handed protected pages to the allocator

After the 64 GiB unlock, the unlocker's late PMA extension selected the highest reserved framebuffer region and registered it with the physical memory allocator, but that region overlapped WPR (write-protected memory) and the GSP firmware heap, so ordinary CUDA allocations could receive protected pages. Copying model weights into those pages raised Xid 31 (ENGINE CE2 HUBCLIENT_HSCE2, FAULT_INFO_TYPE_REGION_VIOLATION ACCESS_TYPE_VIRT_WRITE) about 11 seconds into a vLLM weight load, reproduced on all four tested cards with both the amoghmunikote and the asm64-hooligan unlock implementations. The failure depended on which physical pages the allocator returned, not on allocating all 64 GiB. Fix: disable the late PMA extension and leave the region reserved. It only added 141 MiB (about 0.22%), and the main 64 GiB unlock keeps working; the driver log then shows CMPUNLOCK_PMA: extension SKIPPED (region overlaps WPR + GSP heap). This account is Miséricorde's (2026-09-18). bayley_w had found and fixed (with Claude) the same 141 MB read-only top region in August and traced it to the late PMA extension, which asm64's fork also carried, and on 2026-08-06 Fats recalled a write-protected chunk near the top.

Reported 8 GB
Source: @bayley_w, 2026-08-06
Source: @c.03937, 2026-08-06
Source: @wakeupmbrc1, 2026-09-18

A longer write-up by @wakeupmbrc1 (Miséricorde) exists, with his analysis of the PMA bug and of the separate int32 Xid 31 class in DFlash2; see Discord. (Pointer: the wiki does not present it as the canonical analysis until the author agrees.)

Apply the late-PMA fix or get guaranteed Xid 31

bayley_w: make sure the late PMA fix is applied to your cmpunlocker build; without it some workloads produce guaranteed Xid 31 errors. bassRAWR and Phoenix17 confirmed that cmpunlocker PR #32 (the late-pma branch) is required, or cuda_memtest errors on the reserved region; Amogh said he would verify it and push it to master. romangitz independently found that the unlock using memory that should stay reserved throws CUDA errors that resemble memory errors, and after patching, his 10 GB card is rock solid (10+ hours a day). Waka324 linked the upstream fix, commit f51da03, "Leave the high reserved region reserved: it backs WPR2". The PMA fix also lets NVIDIA DCGM diag level 4 pass. Far: the bug explains a 64 GB card crashing when cuda_memtest ran over all memory more than once, while models never used all of it; vLLM triggered it during model loading.

Proven 8 GB and 10 GB
Source: @farawaythrower, 2026-08-10
Source: @bayley_w, 2026-08-11
Source: @bassrawr, 2026-08-26
Source: @romangitz, 2026-09-02
Source: @cmdrwaka324, 2026-09-02

Keep the high reserved region reserved: how the server found the fix

Warning An unlocker build that registers the GSP's WPR2 region with the allocator gives Xid 31 under allocation churn and when VRAM fills, and the engine that happens to be running gets the blame. What the server saw:

2026-08-13: Merom Green's vLLM workload on a 64 GB-unlocked card crashed with Xid 31, ENGINE CE2 HUBCLIENT_HSCE2 REGION_VIOLATION ACCESS_TYPE_PHYS_WRITE, at an address inside the top reserved region that cmpunlocker had exposed. vLLM, first blamed together with a newer Gemma4 MoE model, was only the trigger that allocated into the bad region. He cherry-picked armand0e's late-PMA pull request and the identical workload then ran without Xid 31; armand0e, the PR author, said Merom hitting it on a 64 GB card confirmed the bug (Impactic also took part).
2026-08-19: clee's build of mudler/vllm.cpp for sm80 crashed on startup. Merom said the cause is probably the late-PMA issue and the fix is a PR to cmpunlocker, not to vLLM. clee had seen it on asm-hooligan's fork, which Humvee55 confirmed is the fork he pushes PRs to; ST suggested (untested) having the driver report the card as an A100 so engines load the right kernels.
2026-08-20: Mitakihara, with constant Xid 31 in vLLM, asked whether the late PMA patch is amoghmunikote/cmpunlocker PR #32 ("Leave the high reserved region reserved: it backs WPR2", by armand0e). Toffles said it was not officially merged yet and needed fixing on both card types, and that allocating into reserved space is bad. Merom confirmed that cherry-picking it fixed his vLLM MoE crashes. The workaround until patched was not to use the full VRAM; OB88594 suggested the bayley/cmpunlocker fork.
2026-08-23: PR #32 was merged (GitHub sweep, lab notes).
2026-08-30: MishaNya hit Xid 31 allocating memory with PyTorch for pre-training and fixed it (with Claude) at the cost of an extra 227 MiB reserved; Thad needed the same fix with 167 MiB, and Merom pointed to the late-pma patch. The same day the fix shipped in cmpunlocker v0.3 (next section).

8 GB and 10 GB
Source: @meromgreen, 2026-08-13
Source: @impactic, 2026-08-18
Source: @armand0e, 2026-08-19
Source: @c133, 2026-08-19
Source: @stcord, 2026-08-19
Source: @meromgreen, 2026-08-19
Source: @scallion_9883, 2026-08-19
Source: @mitakihara_., 2026-08-20
Source: @toffles3179, 2026-08-20
Source: @meromgreen, 2026-08-20
Source: @ob88594, 2026-08-20
Source: @mishanya13, 2026-08-30
Source: @thad1337, 2026-08-30
Source: @meromgreen, 2026-08-30
Source: @darkmaniac7, 2026-10-19 (lab notes)
Source: repo github.com/amoghmunikote/cmpunlocker by amoghmunikote on GitHub

The same overlap on the 10 GB card's 80 GB geometry: 141 MB of WPR2 in a 143 MB region

ggualerz's research repo on taking the 10 GB card from 40 GB to 80 GB found and fixed the same bug in cmpunlocker's late-pma.patch (branch fix/late-pma-wpr2-overlap, commit e093717). On the coherent 80 GB geometry the topmost reserved framebuffer region was about 143 MB, and about 141 MB of it was WPR2, the GSP's own protected heap; the patch registered all 143 MB as free, though only about 2 MB was really free. The effects: the allocator reported about 141 MB of capacity that did not exist, and once the main region was used up, client allocations received pages inside WPR2, where writes are blocked by hardware protection (copy-engine errors) or corrupt GSP state. Because the allocator uses the late region only when the main one is exhausted, this is a near-full hazard, not the cause of the 10 GB card's separate crash above about 40 GiB. The fix clamps the registered range below the start of WPR2 and leaves the region reserved in the region table. With it, a 79 GiB SM-kernel fill and verify ran clean, and an allocation that overshot the card got a clean out-of-memory error instead of phantom pages. The repo is documentation only (no license, last commit 2026-08-31); its 80 GB verdict is on The 10 GB card: provenance and the 80 GB question.

Reported 10 GB
Source: repo github.com/ggualerz/cmp170hx10g-to-80g by ggualerz on GitHub

A memtest crash in a 141 MB region: still on v0.2

On 2026-09-15 Justinjja's memtests crashed his GPU, and Claude traced it to a 141 MB region of VRAM, matching the earlier report by bayley_w. peter643, quoting ChatGPT, said a maintained cmpunlocker derivative credits bayley with a "PMA WPR overlap fix" and upstream v0.3 (2026-08-30) lists a "Late PMA fix" in its release notes; Justinjja then found he was still on 0.2. armand0e said bayley took his open PR and merged it into a fork before it was merged into the main repo, and that many people hit this bug. Fix: upgrade to v0.3 or later.

Reported
Source: @justinjja, 2026-09-15
Source: @peter643, 2026-09-15
Source: @armand0e, 2026-09-16

Fixes by build

Mainline cmpunlocker has carried the fix since v0.3. The forks below fixed the same tail problem earlier by reserving more of the top; they are frozen and lack later mainline features, so they matter mostly for reading old reports. Fork overview: Cmpunlocker forks.

PR #32 and cmpunlocker v0.3: which builds have the fix

The fix is armand0e's pull request to amoghmunikote/cmpunlocker, PR #32 ("Leave the high reserved region reserved: it backs WPR2"; upstream commit f51da03, linked by Waka324). A GitHub sweep (lab notes, 2026-10-19) records it merged by high117 on 2026-08-23, and Amogh shipped the Late PMA fix in release v0.3 on 2026-08-30, after many memory-issue reports; the v0.3 release notes list "Late PMA fix". bayley took armand0e's open PR into his own fork before it reached the main repo, and a maintained derivative credits bayley with a "PMA WPR overlap fix". Lycorascence called the patch PR #34; which number is the merged fix is an open question. Other builds in the sweep:

tlswotj febe720 (2026-08-07): a late-pma clamp that caps the allocator at 62 GiB and excludes the unbacked top HBM sliver, which fixed a full-VRAM Xid 31; lesj0610 carries it (bb8c6f0), and lesj0610 53b7927 notes that upstream now leaves the high region reserved.
ZombieLurker 18fbb66 (2026-09-14): a working build that dropped late-pma.patch because it registers the GSP WPR2 region into the allocator, causing Xid 31 wedges under alloc/free churn.
duggasco a2b4663: gates the late-pma extension and rejects fully reserved framebuffer regions.

The upstream v0.5 core still carries a late-PMA patch; the lab's ESXi backport kept it with lifecycle hardening and passed 60 GiB probes, so the trap depends on the build and version. What to do: run v0.3 or later built from a fresh checkout, and check any full-VRAM Xid 31 against the WPR2 and top-sliver reservation first.

Reported 8 GB and 10 GB
Source: @_j3ns, 2026-08-30
Source: @cmdrwaka324, 2026-09-02
Source: @peter643, 2026-09-15
Source: @armand0e, 2026-09-16
Source: @lycorascence, 2026-09-30
Source: @darkmaniac7, 2026-10-19 (lab notes)
Source: repo github.com/amoghmunikote/cmpunlocker by amoghmunikote on GitHub
Source: repo github.com/ZombieLurker/cmpunlocker by ZombieLurker on GitHub

Tail-reserve forks for the 8 GB card: buliaoyin at 65344 MiB and qg19932GH at 64768 MiB

On 2026-08-10 Jack pointed Amogh to two community forks that already fixed the tail over-allocation bug: github.com/buliaoyin/cmpunlocker, and commit aaddfd4 in github.com/qg19932GH/cmpunlocker, "fix: reserve VRAM tail in PMA". He said one fork "changed too many things", so he did not want to use its patch and preferred the fix to land in the main branch; as of that day it was not in amoghmunikote/cmpunlocker main. Studebaker added that some people who had issues with the main repo got asm64's version working. On 2026-09-09 nightdays pointed Ikyo, who was fighting a Gen1-only link with older code and forks, to the buliaoyin fork.

buliaoyin/cmpunlocker (default branch fix-tail-region-hang, last commit 9c96c7f on 2026-08-10) fixes the 8 GB tail-region crash (OOM, full-memory memtests, vLLM page thrash writing GSP-private tail addresses). It keeps the 64 GiB geometry but exposes 65344 MiB, leaving 192 MiB of margin at the top; the author tried 64000, 64768 and 65024 MiB before settling on 65344. The 10 GB card stays at 40960 MiB. Frozen: it lacks later mainline features.
qg19932GH/cmpunlocker is a fork of asm64-hooligan/cmpunlocker (branched 2026-08-07, last commit aaddfd4 on 2026-08-11) that reserves the top 768 MB of the 8 GB card's 64 GiB, so it shows 64768 MiB, to keep the CUDA allocator away from the GSP firmware, ECC and hardware-reserved tail (against Xid 79 and CUDA_ERROR_UNKNOWN 999 near 100% VRAM). Everything else is asm64's code. It ships a prebuilt binary and unsigned modules: do not run them.

Related symptom: cuda_memtest errors in the last few hundred MB of a 64 GB card on builds without the late-PMA fix (see How to test).

Reported 8 GB
Source: @_studebaker_, 2026-08-10
Source: repo github.com/buliaoyin/cmpunlocker by buliaoyin on GitHub, posted by @ludashi666, 2026-08-10
Source: @_j3ns, 2026-08-10
Source: @.nightdays, 2026-09-09
Source: @ikyo, 2026-09-09
Source: repo github.com/qg19932GH/cmpunlocker by qg19932GH on GitHub

344303947/cmpunlocker2: a carve-out fork that stabilized Xid-prone 8 GB cards

no-melitas reported on 2026-08-13 that his 8 GB card (VBIOS 92.00.6d.00.0a) threw Xid 31 after some activity, plus an occasional Xid 119, on the then-current amoghmunikote/cmpunlocker; he blames the GSP reacting badly to memory-management churn. Switching to github.com/344303947/cmpunlocker2 (first linked by Zorg33 on 2026-08-11) fixed it. The fork replaces the original rather than running beside it, and as he described it, its patches set a 1536 MB framebuffer carve-out the same way the earlier 512 MB carve-out was set; the link stayed at Gen2 x4 and nvidia-smi still reports 64000MiB total. The repo tagline translates as "limit problem-region VRAM, 62906MiB". Read from the repository itself (last commit 9a71162 on 2026-08-09, commits signed buliaoyin), the fork keeps the full 64 GiB geometry but exposes 65024 MiB, splits the framebuffer region at that boundary and stops late-PMA registration from crossing it; the 10 GB card stays at 40960 MiB. cuddylac997, quoting an AI review, said the fork gives 10 GB cards nothing new beyond not needing a reboot when a model goes over the about 37 GB cap, and MrManny had no issues on Amogh's master. On 2026-08-18 Waka324 asked whether commit 29ce0e1 ("Fix tail framebuffer region handling") is recommended; nobody answered in this history. The fork is frozen and behind mainline (no ECC, BAR1 resize, P2P or 615.x support) and carries the late-PMA logic that PR #32 replaced upstream.

Reported 8 GB
Source: @seruditus, 2026-08-13
Source: @mrmannyatx, 2026-08-13
Source: @cuddylac997, 2026-08-13
Source: @cmdrwaka324, 2026-08-18
Source: repo github.com/344303947/cmpunlocker2 by 344303947 on GitHub

Stale checkouts that silently lack the fix

NFA (2026-08-27): Amogh's cmpunlocker did not yet include the patch that prevents release of the reserved memory, so users had to apply that PR by hand or use asm64's unlocker; Xid 31 is the typical error. The fix shipped in v0.3 three days later. On 2026-09-04 machinezero's GLM-5.3 Flash on 4 cards loaded but crashed after 3 to 10 tool calls, after which the driver was broken and the cards needed a reboot to show in nvidia-smi again; he linked it to running out of memory with cmpunlocker. Schaka pointed to the late PMA fix in v0.3 (then about 5 days old), which as he understood it fixes those OOM issues (no guarantee), and mentioned a separate Chinese repo that makes all 64 GB addressable, unsure whether it is the same fix. On 2026-09-22 Pollos Naranja posted an incident: a vLLM Qwen service on four unlocked cards died with Xid 31 (MMU REGION_VIOLATION, ACCESS_TYPE_VIRT_WRITE) on every GPU and CUBLAS_STATUS_EXECUTION_FAILED in a BF16 cublasGemmEx, followed by physical-write region violations at 0xff7300000 through the copy engine. The kernel log showed firmware WPR at 0xff7300000 to 0xffff00000 while the late PMA extension registered 0xff7200000 to 0xfffffffff as allocatable: the ranges overlap, and the fault sat exactly at the WPR start, implicating the late-pma patch in his installed tree. Before the fault the cards ran at core 62-67 C and HBM 69-73 C with about 64278 of 65536 MiB in use and the KV cache at 28.1%, and no Xid 119 or 154 followed. Amogh replied that the fixed late-pma branch had been merged long before and the git pull had failed silently: update to current master and rebuild. Harmacist hit a similar error on a 10 GB card unlocked to 40 GB and has had no issues since fixing it. Checking which module is really loaded: Verifying an unlock.

Reported 8 GB and 10 GB
Source: @never_fly_away, 2026-08-27
Source: @machinezer0, 2026-09-04
Source: @schakaqt, 2026-09-04
Source: @thegreatredear, 2026-09-22
Source: @_j3ns, 2026-09-22
Source: @romangitz, 2026-09-22

Lock-ups on the 40 GB geometry

A card locks up when VRAM fills into the protected tail (40 GB and 64 GB geometries)

Lycorascence (2026-08-20): on the unlocked card, once VRAM fills and spills into the writable buffer at the end of memory, the card locks up until the computer restarts, so he keeps VRAM use lower to avoid writing into the protected region, which also slows some workloads. The same day he noted that exhausting the whole VRAM stalls his 10 GB card at 40 GB until a full machine reboot. aaron reported that re-running the unlock with the asm64-hooligan/cmpunlocker fork seemed to fix it; Lycorascence waited for that fork's pull request to reach main cmpunlocker, or for a VBIOS solution, before trying again. Petri Krohn (2026-08-27) said the memory fault reported by ahnguyen17 (maintainer of the cmp-170hx-vllm repo) is in the WPR2 region and not a defect of the 10 GB card. In a later thread (2026-09-30) Lycorascence explained that without the Late PMA patch there is usually a crash in the 40 GB geometry when writes overwrite critical memory that was relocated there, and that the 64 GB geometry had a similar issue needing a similar patch; Amogh shipped the Late PMA fix in release v0.3 on 2026-08-30 after many memory-issue reports. Whether it changes behaviour on the 80 GB geometry was unknown, and tests run with it still crashed above about 40 GiB. Usable memory at the 40 GB profile: Memory unlock on the 10 GB card.

Reported 8 GB and 10 GB
Source: @lycorascence, 2026-08-20
Source: @lycorascence, 2026-08-20
Source: @aaron0010, 2026-08-20
Source: @petrikrohn_26497, 2026-08-27
Source: @_j3ns, 2026-08-30
Source: @lycorascence, 2026-09-30

Xid 31 that is not the unlocker

An Xid 31 on a build that already has the late-PMA fix is usually the engine. In the reports below the crash comes back at the same context depth or KV pool size, often lands on one pipeline rank (the last rank holds the drafter and has the least free memory), and goes away with an engine patch, not an unlocker change. Engine patch sets in general: vLLM on sm_80: known bugs and patches.

Xid 31 at about 85K context with DFlash2: an int32 address overflow in vLLM spec decode

Running Qwen3.8 27B with the DFlash2 drafter on vLLM, Miséricorde got Xid 31 on the first cached decode step at high physical KV block IDs, repeatably at about 85K context (2026-08-30). Updating the cmpunlocker patch did not help; he states it is a vLLM runtime bug, not related to cmpunlocker (the unlocker Xid issue others reported is a different one and was patched). The fix, found by his agent, is in vllm/v1/attention/ops/spec_decode_attn.py and casts the loaded block id to int64: blk = tl.load(bt_ptr + req * stride_bt + pos // BLOCK_SIZE, mask=k_ok, other=0).to(tl.int64). Mamba safeguards (vLLM PR #50021 state-copy bounds, PR #50729 overlap-safe state movement) are worth keeping but did not fix the 85K crash; without the int64 change DFlash2 drops with Xid 31 on cache reads. Because of this bug he had tested Qwen3.8 27B only up to 200K context; his agent's attention calls at 400K, 524K and 700K tokens returned output, linearly slower, which does not prove semantic quality. His 2026-09-18 write-up separates the two crash classes: the PMA bug exposed protected physical pages, while the DFlash bug computed addresses in int32, overflowed into invalid pointers and needed the int64 cast. Drafters on these cards: Speculative decoding on the 170HX.

Reported
Source: @wakeupmbrc1, 2026-08-30
Source: @wakeupmbrc1, 2026-08-30
Source: @wakeupmbrc1, 2026-09-18

Load-time Xid 31 from GPU Marlin repack: a CPU repack fallback

Miséricorde's vLLM patch set includes a bit-exact CPU Marlin repack fallback for W4 and W8 weights on sm_80 (2026-09-01). It addresses a separate load-time Xid class caused by GPU repack and VMM churn (his issue #27). It raises W4 startup to roughly 206 seconds but avoids the dangerous GPU repack path. The same patch set carries Mamba state-copy bounds (vLLM PR #50021), overlap-safe Mamba state movement (vLLM PR #50729) and a num_accepted_tokens race fix based on commit c2881ce60. He says a CMP sm_80 build may need both this and the int64 spec-decode cast above.

Reported
Source: @wakeupmbrc1, 2026-09-01

Xid 31 at load on the "Samsung 40GB" card: memory vendor or the WPR2 reservation?

ahnguyen17/cmp-170hx-vllm, a vLLM recipe repo for Qwen3.8 27B, runs a "Samsung 40GB" card and a "Hynix 64GB" card. Going by the vendor split members reported (all 10 GB cards Samsung, all 8 GB cards Hynix), the Samsung card is the 10 GB card unlocked to 40 GB; the repo itself never names a card variant. On the Samsung card the author needed expandable_segments:False (the opposite of the Hynix setting) plus patches/sm80-int8-repack-cpu-fallback.patch, or vLLM hit an Xid 31 wedge at model load, each of the two being an independent trigger. The Xid is asynchronous, so the traceback lands in a bystander frame near gptq_marlin_repack. The fix runs the Marlin repack on the CPU (bit-exact, one host-to-GPU copy per layer, about 3 extra minutes of boot), with no Xid since. A 2026-10-07 commit adds notes (not a diff) for a MoE variant saying compiled repack bursts wedge the card's GSP past about 40 GB committed.

Contested 10 GB
Source: repo github.com/ahnguyen17/cmp-170hx-vllm by ahnguyen17 on GitHub
Source: @mrpec, 2026-07-19
Source: @giaki3003, 2026-07-25

Whether the trigger is the memory vendor or the reservation at the top of a 40 GB card was not settled. In the order posted:

Petri Krohn could not see why a kernel would work on one variant and crash on the other when both share GA100 silicon. Source: @petrikrohn_26497, 2026-08-25
giaki3003 called the Samsung explanation slop. Source: @giaki3003, 2026-08-26
chonsworth guessed vendor-specific binning. Source: @chonsworth, 2026-08-26
m3gagluk suggested the real difference was 40 GB versus 64 GB tuning with a driver lacking the memory reservation fix. Source: @m3gagluk, 2026-08-26
Petri thought that likely: Qwen3.8 27B hits the WPR2 reservation at the top of memory. Source: @petrikrohn_26497, 2026-08-26
On 2026-08-27 Petri said he had told the author that the fault is in the WPR2 region, not in the 10 GB card. Source: @petrikrohn_26497, 2026-08-27
On 2026-09-29 Lycorascence, fighting vLLM wedges on his own 10 GB card, pointed to the repo's Samsung notes but found the repo a month old and aimed at a much older vLLM; it was updated again on 2026-10-07. Source: @lycorascence, 2026-09-29

GLM-5.3-Flash at long context: int32 KDA offsets above 262,144 tokens and an unpatched second wall

dkpoulsen's GLM-5.3-Flash stack (NVFP4, PP5 on five 8 GB cards, stock glm53-flash image plus bind-mounted overrides) found two context walls, both Xid 31. (1) The flash-linear-attention KDA prefill kernels compute element offsets as token times 8192 in int32, which overflows above 262,144 tokens; the result is an Xid 31 (MMU fault, SM write) that wedges the driver. The repo's override files move that index math to int64 (the sequence-boundary loads, the length cast, the gate-row load and the per-chunk state indexing); they were bit-identical on short sequences and ran a clean 270k-token probe. (2) "Wall 2": rank-agnostic Xid 31 faults (GRAPHICS FAULT_PDE/REGION_VIOLATION) at the sparse-MLA layers during long chunked prefills, unpatched. 55.5k was verified clean and about 61k and above crashed; raising --max-num-batched-tokens to 8192 moved the wall from 60.7k to 69.6k (it died at exactly num_computed_tokens=69632), so the trigger scales with absolute position. CUDA graphs, prefix caching, the DSA indexer and the MLA prefill pipeline tested standalone were all ruled out, and the repo caps --max-model-len at 49152 for that reason. On the same rig the vowstar/vllm-sm80 fork (branch glm53-sm80, sparse MLA through a Triton sm_80 kernel) reached a 1,048,576-token context with zero Xids; wall 2 does not exist in that stack.

Reported 8 GB
Source: repo github.com/dkpoulsen/vllm-170hx-glm5.3-flash by dkpoulsen on GitHub, posted by @artistosu, 2026-08-29

GLM-5.3-Flash with PP8: int64 sparse-MLA KV offsets once the KV pool passes 2^31 bytes per layer

promisezackr's 24-patch set serves GLM-5.3-Flash NVFP4 on eight 8 GB cards with PP8 and DFlash2 over the official vllm/vllm-openai:glm53-flash image. One of its fixes moves the sparse-MLA KV offsets to int64: with int32 offsets the card hit an Xid 31 MMU fault once the KV pool exceeded 2^31 bytes per layer. The patches target that one image build; Zanooda runs them on six cards. Recipe: GLM 5.3 Flash on 8x 8 GB (promisezackr PP8 NVFP4 DFlash2).

Reported 8 GB
Source: repo github.com/promisezackr/glm53-flash-170hx-pp8 by promisezackr on GitHub, posted by @zanooda, 2026-08-31

DeepSeek V4 Flash: a transient logits buffer on a thin last rank (about 134K to 174K tokens)

tim-odonnell's DeepSeek V4 Flash repo traces Xid 31 crashes to a transient float32 logits buffer in the sparse-attention indexer. During prefill the Triton fallback that sm_80 takes (DeepGEMM is unavailable) allocates rows times context length in one piece, where rows is the prefill chunk, so the buffer grows with context; the last pipeline rank, which also holds the DSpark drafter, has only about 1 to 2 GiB free after weights and KV cache. With DSV4_LOGITS_ROW_CHUNK unset (the code default 0 means off), leinasi2014's Vision-Exp fork crashed the first time a single request passed about 174,000 tokens, in every one of 5+ controlled rounds, with CUDA graphs on or off; at 2048 rows the buffer is about 1.33 GiB there. DSV4_LOGITS_ROW_CHUNK=64 processes 64 rows at a time with the same results and a tiny buffer. The same kernel (fp8_mqa_logits_triton) killed the worker above about 134k on 0731, and an earlier DSpark config with a 1M context crashed at about 97k on the last rank with about 1.4 GiB free. Two lessons from the repo: the surviving ranks then print a misleading gloo ... Connection closed by peer, so chase the Xid, not the gloo message; and CUDA_LAUNCH_BLOCKING=1 located the kernel in one run. leinasi's first incident note blamed a batch-size change, but their relaunch script had silently dropped the row-chunk setting: when a crash follows a config change, diff the full container environment. Recipe: DeepSeek V4 Flash 0731 on 3x 8 GB (PP3 DSpark).

Reported 8 GB
Source: repo github.com/tim-odonnell/cmp170hx-deepseek-v4-flash by tim-odonnell on GitHub, posted by @irythros, 2026-08-19

Over-allocation: the 0.965 memory utilisation ceiling on three cards

In tim-odonnell's three-card DeepSeek-V4-Flash-0731 setup (PP3 with a 15,15,13 layer split, DSpark 5, Gen2 x4 without P2P, 100 W per card), a binary search on vLLM's --gpu-memory-utilization gave 338,432 tokens of context at 0.95, 574,464 at 0.96 and 691,968 at 0.965, while 0.97 and 0.98 crashed with an illegal memory access on the PP2 worker. His troubleshooting table reads an Xid 31 during a workload as something trying to use more memory than the card really has: check the workload's memory request against the real unlocked size. For Xid 45 and 154 crashes his general advice is 0.90 or below on unlocked cards, without assuming every last megabyte of the reported VRAM is usable; his tuned production setting is the measured 0.965.

Reported 8 GB
Source: repo github.com/tim-odonnell/cmp170hx-deepseek-v4-flash by tim-odonnell on GitHub, posted by @irythros, 2026-08-19

Do not reload nvidia_uvm after an Xid 31 on a multi-card job

Warning On 2026-09-29, in tim-odonnell's DeepSeek V4 Flash Vision-Exp setup, running rmmod and modprobe nvidia_uvm after one Xid 31 turned it into a "UVM global fatal error" and Xid 154 (OS reboot required) on all 4 cards, and nvidia-smi still looked healthy afterwards. His guide's order: capture the evidence first, then reboot if sudo dmesg -T | grep -i xid shows anything new; a reboot is the only reliable recovery on these unlocked cards. dkpoulsen's GLM-5.3-Flash notes agree: a crash can wedge the driver while nvidia-smi looks healthy, the next start fails cudaSetDevice on the wedged GPU, and only a host reboot recovers, the tell being a nonzero dmesg | grep -ci xid. kaka86mm needs a cold power cycle after an Xid 31 (about 1 in 10 engine starts) and says nvidia-smi -r is not supported. The other practice exists: allover326's DeepSeek recipe recovers a wedged card with nvidia-smi -r, then rmmod and modprobe nvidia_uvm, then re-applying the power limit, and Schaka's launchers also reload nvidia_uvm to clear an Xid 154 wedge. After an Xid 31 on a multi-card pipeline job, reboot instead. The full recovery ladder is on Recovery.

8 GB
Source: repo github.com/allover326/deepseek-v4-cmp170hx by allover326 on GitHub, posted by @deejay9983, 2026-08-05
Source: repo github.com/tim-odonnell/cmp170hx-deepseek-v4-flash by tim-odonnell on GitHub, posted by @irythros, 2026-08-19
Source: repo github.com/dkpoulsen/vllm-170hx-glm5.3-flash by dkpoulsen on GitHub, posted by @artistosu, 2026-08-29
Source: repo github.com/kaka86mm/glm53-flash-170hx by kaka86mm on GitHub, posted by @dkpoulsen, 2026-09-19
Source: repo github.com/Schaka/deepseek-v4-cmp170hx by Schaka on GitHub

"vram address overflow" on 2 of 4 cards

aidoc reported that real workload testing failed on 2 cards with a "vram address overflow" error in a 4-card system, while others reported no such issue (2026-08-01). No cause or fix was posted. The report predates the late-PMA fix in v0.3 (2026-08-30).

Reported
Source: @sekkit_75019, 2026-08-01

How to test

Test for the reserved-top bug: cuda_memtest at 100%, the churn loop and the kernel log

Run cuda_memtest over 100% of the card's memory. On a build without the late-PMA fix it fails in the last blocks on every card, while capping it at about 63.75 GB on a 64 GB card passes. memtest_vulkan does not show it, and gpu-burn fills only 90% of memory by default.
For allocation churn, run a loop that calls cudaMalloc of 100 MiB and cudaFree 100,000 times; it triggered Xid 31 on an unfixed 8 GB card.
Read the kernel log from the driver load. A build with the fix logs CMPUNLOCK_PMA: extension SKIPPED (region overlaps WPR + GSP heap). If the log shows a late-PMA range that overlaps the firmware WPR range (on one 64 GB card: WPR at 0xff7300000 to 0xffff00000, late-PMA range from 0xff7200000), the fix is missing: update to current master and rebuild.
With the fix, NVIDIA DCGM diag level 4 passes, and an allocation that overshoots the card gets a clean out-of-memory error instead of phantom pages.
After any crash, count the Xid lines with sudo dmesg | grep -ci xid before restarting anything. A count that went up means a card still carries the fault, even when nvidia-smi looks healthy.

The wider acceptance routine (gpu-burn, memtest_vulkan, cuda_memtest, then a real model) is on Testing and burn-in.

Reported 8 GB and 10 GB
Source: @skyy2333, 2026-07-29
Source: @ju5t5omedude, 2026-08-10
Source: @jimtjames, 2026-08-10
Source: repo github.com/dkpoulsen/vllm-170hx-glm5.3-flash by dkpoulsen on GitHub, posted by @artistosu, 2026-08-29
Source: @cmdrwaka324, 2026-09-02
Source: @wakeupmbrc1, 2026-09-18
Source: @thegreatredear, 2026-09-22
Source: @_j3ns, 2026-09-22
Source: repo github.com/ggualerz/cmp170hx10g-to-80g by ggualerz on GitHub

Crashes with less than 1 GB free: watch memory use when it happens

Fats reported occasional crashes when a card has less than 1 GB of its 64 GB free, recalled someone mentioning a write-protected chunk near the top, and suggested watching nvidia-smi -l 1 for cards above about 97% memory use when a crash happens (2026-08-06). For a check that the unlock itself came up clean, eka2513 posted a readout checklist from cmpunlocker master (2026-08-24); the unlock checks are on Verifying an unlock.

Reported 8 GB
Source: @c.03937, 2026-08-06
Source: @eka2513, 2026-08-24