Every kernel launch is an ioctl into ring 0

A CUDA kernel runs on the GPU. The path that puts it there runs through ring 0 on the host. That second fact is the one the topic is actually about, and it is worth separating from the first before going further.

Two different things share the word kernel here. A CUDA kernel is a __global__ function executed by thousands of GPU threads in lockstep. The OS kernel is the host’s privileged execution mode - ring 0, full memory access, no sandbox above it. They share a noun and nothing else. What connects them is the submission path. Launching GPU compute crosses the user/kernel trust boundary on the host every single time, and that crossing is where the corruption bugs live.

The path is mechanical. kernel<<<grid, block>>> compiles to a host-side call into the CUDA Runtime - libcudart. That calls the CUDA Driver API - libcuda. That calls the user-mode driver, which marshals the request and issues an ioctl into the kernel-mode driver. On Linux that driver is nvidia.ko, fronted by the device nodes /dev/nvidiactl and /dev/nvidia[N], with the Resource Manager - NVRM - handling control calls through NV_ESC_RM_CONTROL. On Windows the equivalent runs through the display driver nvlddmkm.sys and the DxgkDdiEscape handler exposed by the DirectX graphics kernel subsystem. Either way, the request carries a command code and a user-space pointer to a parameter struct, and a ring-0 handler reads that struct, validates it, allocates GPU memory through the GMMU, builds a command pushbuffer, and submits it to the GPU command processor.

The validation step is where it breaks.

The dominant bug class in NVIDIA’s kernel-mode driver is CWE-787, out-of-bounds write. The pattern repeats across release after release of the GPU Display Driver security bulletins. The handler reads a length, an offset, or an index from the user-supplied parameter struct and uses it to write into a kernel-side buffer without validating it against the actual allocation size. The field is attacker-controlled. The write is not bounded. The result is a write past the end of a kernel pool allocation, into whatever object sits adjacent. CWE-125 out-of-bounds reads show up in the same bulletins for information disclosure, and CWE-367 - time-of-check to time-of-use - shows up where a parameter is validated once and dereferenced later after the attacker has changed it.

The trust boundary is the whole story. Userland is not supposed to be trusted by ring 0. The ioctl handler exists precisely to validate everything crossing inward. When a bounds check is missing on one field of one control call out of hundreds, the driver writes attacker-influenced data into kernel memory on behalf of an unprivileged process. That is a kernel write primitive originating from a normal GPU API call.

Weaponising it follows the standard kernel pool corruption model. The attacker shapes the pool - allocations and frees timed to place a target object immediately after the buffer that will overflow. Then the vulnerable control call fires the out-of-bounds write, corrupting the adjacent object. The high-value targets are a function pointer the kernel will later call, or a token or privilege field that governs the calling process. Corrupt the first and execution redirects. Corrupt the second and an unprivileged process becomes SYSTEM or root. The outcome is ring-0 code execution or local privilege escalation from a context that started with none. MITRE T1068 - exploitation for privilege escalation. No working primitive is reproduced here; the mechanism is the point.

The precondition is the part defenders underweight. This is not remote. The attacker needs a channel to issue control calls to the driver - meaning code already running on the host, even unprivileged. In a workstation that is post-initial-access escalation. In a multi-tenant GPU environment the channel is the workload itself. A tenant submitting a CUDA training or inference job is, by design, issuing ioctls to a shared host driver. The job is the delivery vehicle. The driver is the shared component every tenant touches. That is the structural exposure in rented GPU compute, and it is why the kernel-mode layer of a GPU driver is a higher-value target now than it was when the only thing running CUDA was a single trusted user.

The real-world record is consistent rather than dramatic. NVIDIA’s GPU Display Driver bulletins disclose kernel-mode layer out-of-bounds writes on a recurring cadence - multiple per year - most rated for local privilege escalation, code execution, denial of service, and information disclosure. These are found and patched proactively more often than they are caught in the wild, which is the correct read: the class is live and the surface is large, not that exploitation is rampant. The container path is where impact has been highest. CVE-2024-0132 in the NVIDIA Container Toolkit (versions up to and including 1.16.1) is a TOCTOU flaw - CWE-367 - that allows a container with GPU access to escape to the host. CVSS v3 base score 9.0. Disclosed by Wiz Research in September 2024, with CVE-2024-0133 as a related lower-severity issue. A container granted GPU access for legitimate ML work could reach the host filesystem and break the isolation boundary. MITRE T1611 - escape to host. CVE-2021-1056 earlier covered a container isolation flaw exposing host GPU device files. The pattern across all of them: GPU access is handed to a workload as a feature, and the feature carries a path to privilege the workload was never meant to hold.

Telemetry is where this gets uncomfortable. The host-visible events are thin. Sysmon Event ID 1 records the process that loaded the CUDA stack. Event ID 6 records driver load - nvlddmkm.sys loading is expected and benign. The DeviceIoControl calls that carry the actual control codes into the driver are not confirmed in available telemetry by default, and even when they are, they are indistinguishable from legitimate GPU work. Every CUDA process issues thousands of them. There is no baseline anomaly in volume, timing, or target device that separates a malicious control call from a tensor multiply. On Linux the opens of /dev/nvidiactl and /dev/nvidia0 are observable, and auditd can record the ioctl syscalls, but the volume is enormous and the parameter struct content has to be decoded against NVRM control codes to mean anything - most environments do neither.

What happens on the GPU is simply outside the host EDR’s field of view. Pushbuffers, command submission, and VRAM contents are not inspected by any host agent. The endpoint sensor sees a process making driver calls and using GPU memory. It does not see what those calls contain or what sits in GPU memory. A corruption primitive executed through this path leaves almost nothing behind on success. On failure it bugchecks - a Windows kernel crash or a Linux kernel panic with nvlddmkm/nvidia in the stack - and that crash telemetry is, realistically, the most reliable signal that someone is probing the kernel-mode layer. Repeated GPU-driver-attributed kernel faults are an indicator worth correlating; a clean successful exploit is not.

The container escape is the exception that proves where detection actually lives. CVE-2024-0132 does not announce itself in the GPU path. It announces itself when a containerised process touches host paths outside its mount namespace. That is detectable - a PID inside a container reading or writing host filesystem locations it has no reason to reach. The signal is in the namespace violation, not in the CUDA call that set it up. Detection engineering for this class belongs at the container boundary and at the kernel-crash layer, not in trying to parse GPU command streams that no production sensor decodes.

The patch boundary is per-branch. NVIDIA ships driver families - R535, R550, and later branches - and fixes land in specific versions within each branch, so the remediated build depends entirely on which branch a host runs. The Container Toolkit fix shipped in 1.16.2 and 1.17.0. Applying the right one requires knowing which branch is deployed and that the fixed version exists for it.

The residual exposure is deployment lag, and it is severe. GPU driver versions are pinned to CUDA toolkit versions, which are pinned to framework builds - PyTorch, TensorFlow, CUDA libraries - and on Linux to a kernel module that must build against the running kernel. Hosts running production ML stay on known-good branches for stability and break things by moving. The patch exists long before the fleet runs it. In rented and multi-tenant GPU environments the host operator owns the driver and the tenant owns the workload, so the party exposed to the bug is not the party that controls the patch.

Post-patch, the ioctl trust boundary does not go away. Every CUDA release adds control calls. Every new GPU feature adds handlers that read user-supplied structs in ring 0. The out-of-bounds write class persists because the boundary persists and it widens with each generation. A patched driver closes the specific control call that was vulnerable. It does not close the model that makes the next one likely. Running a CUDA kernel will keep meaning crossing into ring 0 on the host, and that crossing will keep being the thing that matters.