On aarch64 with `CONFIG_ARM64_64K_PAGES=y` and driver 590.44.01 or newer, a mapp…ing request of 4 GiB or
more that goes through the driver's DMA submap path comes back with an unusable 64 KiB page just below
each 4 GiB boundary. `cudaMalloc` and `cuMemCreate`/`cuMemMap` succeed; the first read or write to such a page faults:
```
NVRM: Xid (PCI:000f:01:00): 31, ... MMU Fault: ENGINE GRAPHICS GPC0 GPCCLIENT_T1_0
faulted @ 0x4_3fff0000. Fault is of type FAULT_PTE ACCESS_TYPE_VIRT_WRITE
```
The attached `repro-64k-4gib-fault.cu` shows it in about ten seconds: a 5 GiB `cudaMalloc`, then one-byte
writes at `4GiB-128KiB` (passes), `4GiB` (passes), `4GiB-64KiB` (faults). It exits 1 on reproduce and
refuses to run on a 4 KiB host. Practical impact: vLLM dies at engine init allocating a ~90 GiB KV cache.
Not affected: 4 KiB-page kernels (tested to 90 GiB), requests below 4 GiB, `cudaMallocManaged` (which
appears not to go through this path), and driver 580.x.
Same machine, same kernel binary, same firmware throughout:
| clean | faulting |
|---|---|
| 580.142, 580.159.03, 580.173.02 | 590.44.01, 595.45.04, 595.84, 610.43.02, 610.43.03 |
### Cause
`kernel-open/common/inc/nv-linux.h:1078` guards submap sizing with `CONFIG_ARM64_4K_PAGES`. The comment
above it opens *"On 4K ARM kernels, use max submap size a multiple of 2M…"* and states the requirement as
*"the mapped IOVA range must be aligned at 2M boundary"*. The 4K qualifier looks deliberate, but the
pre-590 version of that same comment required only 64 KiB alignment — which every 64 KiB page count
satisfies for free, so the `#else` was adequate then. When the requirement tightened to 2 MiB in
590.44.01 (comment rewrite plus the 4K arm's `- 16` → `- 512`, both first appearing there; 580.173.02
still carries `- 16`), only the arm with an explicit constant was updated. `CONFIG_ARM64_4K_PAGES` occurs
once in the tree and the `#else` is the generic x86/RISC-V fallback, so a 64 KiB-page ARM64 kernel now
gets `NV_U32_MAX >> PAGE_SHIFT` = 65535 pages = `0xFFFF0000`, which is not 2 MiB-aligned.
Sweeping the constant on one machine shows the behaviour follows alignment, not size:
| pages | submap bytes | 2 MiB-aligned | result | faulting offset observed |
|---|---|---|---|---|
| 65472 | `0xFFC00000` | yes | pass | — |
| 65488 | `0xFFD00000` | no | FAULT | `0xFFD00000` (probed directly) |
| 65504 | `0xFFE00000` | yes | pass | — |
| 65520 | `0xFFF00000` | no | FAULT | `0xFFF00000` |
| 65535 (current) | `0xFFFF0000` | no | FAULT | `0xFFFF0000` |
65488 faults while both a smaller (65472) and a larger (65504) value pass, so a size threshold does not
fit. Writing `S` for the submap size and `δ = roundup(S, 2MiB) − S`, the unmapped range is the top `δ`
bytes of each window — `[roundup(S,2MiB) − δ, roundup(S,2MiB))` — and is empty when `δ = 0`. That
accounts for every offset tested, including three that are individually counter-intuitive: at the current
value `δ` is one page, giving exactly one unusable page per 4 GiB; at `0xFFF00000` the range is 1 MiB
wide, so `0xFFFE0000` also faults; and at `0xFFD00000` the offsets `0xFFFE0000` and `0xFFFF0000` both
**pass**, even though `0xFFFF0000` is the faulting offset under the shipping value. In other words the
unmapped range moves with the constant rather than being anchored to 4 GiB.
### Fix (attached, `git am`-clean against 610.43.03)
Scope the arm to `CONFIG_ARM64` and express the margin as 2 MiB in OS pages. `0x200000 >> 12` is the
existing `512`, so ARM64-4K is byte-identical and unchanged; x86_64 and RISC-V keep the untouched `#else`.
The patch also widens the comment's "On 4K ARM kernels" qualifier to match. (`NVCPU_AARCH64` is the more
common arch idiom in this tree; `CONFIG_ARM64` is used here because the expression depends on the
kernel's `PAGE_SHIFT` and the site already tests a `CONFIG_` symbol.)
```c
#if defined(CONFIG_ARM64)
/* 2MiB IOVA alignment (see above), expressed in OS pages: 512 @4K, 128 @16K, 32 @64K */
#define NV_DMA_SUBMAP_2M_PAGES ((NvU32)(0x200000 >> PAGE_SHIFT))
#define NV_DMA_SUBMAP_MAX_PAGES ((NvU32)(((NV_U32_MAX >> PAGE_SHIFT) + 1) - NV_DMA_SUBMAP_2M_PAGES))
#else
#define NV_DMA_SUBMAP_MAX_PAGES ((NvU32)(NV_U32_MAX >> PAGE_SHIFT))
#endif
```
Built against a stock kernel.org 6.18.38 aarch64 64 KiB-page kernel: the reproducer passes, a 90 GiB
allocation passes, and vLLM serves with no Xid events — where the unpatched driver on that same kernel
faults at engine init. Rebuilt against a 4 KiB-page kernel the value is byte-identical to today's and
re-verified there (90 GiB allocation and a vLLM serve both pass). On a vendor-built 6.18.39 the cycle was
run in both directions: unpatched faults, patched serves, reverting restores the fault.
Scoping to `CONFIG_ARM64` is the minimal change, but it may not be the right axis: the L4T tree carries a
TODO on this same macro saying it should subtract the **GPU MMU** page size rather than a constant
("iGPU supports MAX 2MB page size"), which would be arch-independent. The alignment sensitivity itself is
established behaviourally by the sweep above rather than by code inspection — I did not identify the
consumer that depends on it. A 2 MiB alignment requirement for huge mappings does exist elsewhere in the
RM (`iovaspaceApplyDefaultAlignment_IMPL` in `io_vaspace.c`, with `RM_PAGE_SIZE_HUGE` = 2 MiB), which is
corroborating but not the same code path.
### Environment
GB10 (DGX Spark). Faulting: 595.84, 595.45.04, 610.43.03, 610.43.02, 590.44.01. Clean: 580.173.02,
580.159.03, 580.142. Reproduced on both a stock kernel.org 6.18.38 build and a vendor-built 6.18.39,
each with `CONFIG_ARM64_64K_PAGES=y`; CUDA 13.0; Rocky Linux 10.2. Several kernels are installed on this
machine, so the attached bug-report log lists more than one — it was captured on the stock kernel.org
6.18.38 boot, reproducing.
This configuration is documented rather than exotic: NVIDIA's Grace Performance Tuning Guide gives
*"The recommended default value for the page size is 64K"*, and Red Hat's DGX Spark kernel article
(2026-06-23) installs a 64k kernel alongside driver 595.58.03.
One template item I could not complete: the closed kernel module is not offered for this GPU, so I have
no way to run that comparison. The nearest control is driver 580.x on the identical kernel, which is clean.
---
<details>
<summary>Reproducer — <code>repro-64k-4gib-fault.cu</code> (build: <code>nvcc -O2 -o repro repro-64k-4gib-fault.cu</code>)</summary>
```cuda
// repro-64k-4gib-fault.cu — reproducer for the NVIDIA 64K-page DMA-submap IOVA misalignment.
//
// On an aarch64 host kernel built with CONFIG_ARM64_64K_PAGES=y, driver 595.x / 610.43.x:
// GPU mapping requests of >= 4 GiB contain unusable pages. cudaMalloc SUCCEEDS; the first
// access to such a page raises Xid 31 / FAULT_PTE and the process dies with
// "an illegal memory access was encountered".
//
// With the shipping NV_DMA_SUBMAP_MAX_PAGES (65535 pages = 0xFFFF0000) the affected page is
// at offset 4GiB - 64KiB, and recurs every 4 GiB.
//
// Not affected: driver 580.x, 4 KiB-page kernels, requests < 4 GiB, cudaMallocManaged.
//
// Requires: a 64 KiB-page host kernel, and ~5 GiB of free GPU memory.
// Build: nvcc -O2 -o repro repro-64k-4gib-fault.cu
// Exit: 0 = not reproduced 1 = REPRODUCED 2 = could not test (wrong host page size, no GPU,
// insufficient memory, CUDA error)
#include <cstdio>
#include <unistd.h>
#include <cuda_runtime.h>
#define ALLOC_GIB 5
#define NEED_BYTES ((size_t)ALLOC_GIB << 30)
__global__ void poke(unsigned char *p) { *p = 0x5a; }
enum Outcome { OK, FAULTED, ENVFAIL };
// allocate ALLOC_GIB GiB, write one byte at `off`, report
static Outcome probe(const char *label, size_t off) {
unsigned char *p = nullptr;
cudaError_t e = cudaMalloc(&p, NEED_BYTES);
if (e != cudaSuccess) {
printf(" %-24s CANNOT TEST: cudaMalloc(%d GiB) failed: %s\n", label, ALLOC_GIB, cudaGetErrorString(e));
return ENVFAIL;
}
poke<<<1, 1>>>(p + off);
e = cudaGetLastError();
if (e == cudaSuccess) e = cudaDeviceSynchronize();
if (e != cudaSuccess) {
// the CUDA context is unusable after an MMU fault; caller must stop
printf(" %-24s FAULT at base+0x%zx (%s)\n", label, off, cudaGetErrorString(e));
return FAULTED;
}
if (cudaFree(p) != cudaSuccess) { printf(" %-24s CANNOT TEST: cudaFree failed\n", label); return ENVFAIL; }
printf(" %-24s pass\n", label);
return OK;
}
int main() {
long ps = sysconf(_SC_PAGESIZE);
if (ps != 65536) {
printf("CANNOT TEST: host page size is %ld bytes; this reproducer requires a 64 KiB-page\n"
"kernel (CONFIG_ARM64_64K_PAGES=y). A 4 KiB-page host is not affected by this bug.\n", ps);
return 2;
}
int dev = 0;
if (cudaGetDevice(&dev) != cudaSuccess) { printf("CANNOT TEST: no usable CUDA device\n"); return 2; }
cudaDeviceProp prop;
if (cudaGetDeviceProperties(&prop, dev) != cudaSuccess) { printf("CANNOT TEST: cannot query device\n"); return 2; }
// No free-memory pre-check: on unified-memory parts cudaMemGetInfo reports host MemFree, so a warm
// page cache would make this refuse to run on an affected box. The cudaMalloc failure path below
// already reports insufficient memory.
printf("GPU: %s host page size: %ld bytes (needs ~%d GiB of GPU memory)\n", prop.name, ps, ALLOC_GIB);
printf("A %d GiB cudaMalloc, then a one-byte write at three offsets around the 4 GiB mark:\n", ALLOC_GIB);
// two controls that bracket the subject page, then the subject
Outcome r = probe("4GiB - 128KiB (control)", 0xFFFE0000ULL);
if (r == ENVFAIL) return 2;
if (r == FAULTED) { printf("\nREPRODUCED (at a control offset — the affected extent is wider than one page).\n"); return 1; }
r = probe("4GiB + 0 (control)", 0x100000000ULL);
if (r == ENVFAIL) return 2;
if (r == FAULTED) { printf("\nREPRODUCED (at a control offset — the affected extent is wider than one page).\n"); return 1; }
r = probe("4GiB - 64KiB (SUBJECT)", 0xFFFF0000ULL);
if (r == ENVFAIL) return 2;
if (r == FAULTED) {
printf("\nREPRODUCED: the page at 4 GiB - 64 KiB is mapped but unusable.\n");
return 1;
}
printf("\nNOT REPRODUCED: all three offsets are usable on this driver.\n");
return 0;
}
```
</details>
<details>
<summary>Patch — <code>git am</code>-clean against tag <code>610.43.03</code></summary>
```
From 7e4d19d63acfcddf6a7ebcdc5962dc3f180bc01e Mon Sep 17 00:00:00 2001
From: Max Spevack <max.spevack@gmail.com>
Date: Tue, 28 Jul 2026 15:34:50 -0700
Subject: [PATCH] nv-linux.h: apply the 2MiB IOVA submap alignment on all ARM64
page sizes
The comment above NV_DMA_SUBMAP_MAX_PAGES states the invariant the macro
exists to satisfy: "use max submap size as (4G-2M) since the mapped IOVA
range must be aligned at 2M boundary". The guard tests
CONFIG_ARM64_4K_PAGES, and the #else is the generic x86/RISC-V fallback,
so a CONFIG_ARM64_64K_PAGES kernel gets NV_U32_MAX >> PAGE_SHIFT = 65535
pages = 0xFFFF0000, which is not 2MiB-aligned. Before 590.44.01 the stated
requirement was 64KiB alignment, which every 64KiB page count satisfies for
free; when it tightened to 2MiB only the arm with an explicit constant was
updated. The comment's "On 4K ARM kernels" qualifier is widened to match.
Expressing the margin as 2MiB in OS pages is byte-identical at 4K
(0x200000 >> 12 == 512), so ARM64-4K does not change, and it makes 16K and
64K correct too. x86_64 and RISC-V keep the untouched else-branch.
Observed on aarch64 CONFIG_ARM64_64K_PAGES=y with drivers 590.44.01 and
newer (590.44.01, 595.45.04, 595.84, 610.43.02, 610.43.03 all reproduce;
580.142, 580.159.03, 580.173.02 do not): mappings of >= 4 GiB through the
DMA submap path come back with unusable pages. cudaMalloc and
cuMemCreate/cuMemMap succeed, and the first access faults with Xid 31
FAULT_PTE. vLLM dies at engine init allocating a large KV cache.
Sweeping the constant on GB10 / 610.43.03 shows the behaviour follows 2MiB
alignment rather than size. Writing S for the submap size and
delta = roundup(S, 2MiB) - S, the unmapped range in each window is the top
delta bytes, i.e. [roundup(S,2MiB) - delta, roundup(S,2MiB)):
65472 (0xFFC00000, aligned, delta=0) no unmapped range pass
65488 (0xFFD00000, delta=1MiB) 0xFFD00000..0xFFE00000 faults
65504 (0xFFE00000, aligned, delta=0) no unmapped range pass
65520 (0xFFF00000, delta=1MiB) 0xFFF00000..0x100000000 faults
65535 (0xFFFF0000, delta=64KiB) 0xFFFF0000..0x100000000 faults
65488 faults while both a smaller (65472) and a larger (65504) value pass,
so a size threshold does not fit. At 65488 the offsets 0xFFFE0000 and
0xFFFF0000 both pass even though 0xFFFF0000 is the faulting offset under
the shipping value -- the unmapped range moves with the constant.
With this change, built against a stock kernel.org 6.18.38 aarch64 64KiB-page
kernel: the reproducer passes, a 90 GiB allocation passes, and vLLM serves,
where the unpatched driver on the same kernel faults at engine init. Built
against a 4KiB-page kernel the value is byte-identical to today's and
re-verified there. On a vendor-built 6.18.39 the cycle was run both ways:
unpatched faults, patched serves, reverting restores the fault.
---
kernel-open/common/inc/nv-linux.h | 9 +++++----
1 file changed, 5 insertions(+), 4 deletions(-)
diff --git a/kernel-open/common/inc/nv-linux.h b/kernel-open/common/inc/nv-linux.h
index 1a70b9f..857d586 100644
--- a/kernel-open/common/inc/nv-linux.h
+++ b/kernel-open/common/inc/nv-linux.h
@@ -1066,7 +1066,7 @@ typedef struct nv_dma_map_s {
i++, sm = &dm->mapping.discontig.submaps[i])
/*
- * On 4K ARM kernels, use max submap size a multiple of 2M to avoid breaking up 2M page size
+ * On ARM kernels, use max submap size a multiple of 2M to avoid breaking up 2M page size
* sysmem allocations.
*
* Instead of using (4G-4K), use max submap size as (4G-2M) since the mapped IOVA range
@@ -1075,9 +1075,10 @@ typedef struct nv_dma_map_s {
* Bug 5401803: Tracks migrating away from making IOMMU mappings using submaps in favor of
* using sg_chain() to chain a single large scatterlist.
*/
-#if defined(CONFIG_ARM64_4K_PAGES)
-#define NV_DMA_U32_MAX_4K_PAGES ((NvU32)((NV_U32_MAX >> PAGE_SHIFT) + 1))
-#define NV_DMA_SUBMAP_MAX_PAGES ((NvU32)(NV_DMA_U32_MAX_4K_PAGES - 512))
+#if defined(CONFIG_ARM64)
+/* 2MiB IOVA alignment (see above), expressed in OS pages: 512 @4K, 128 @16K, 32 @64K */
+#define NV_DMA_SUBMAP_2M_PAGES ((NvU32)(0x200000 >> PAGE_SHIFT))
+#define NV_DMA_SUBMAP_MAX_PAGES ((NvU32)(((NV_U32_MAX >> PAGE_SHIFT) + 1) - NV_DMA_SUBMAP_2M_PAGES))
#else
#define NV_DMA_SUBMAP_MAX_PAGES ((NvU32)(NV_U32_MAX >> PAGE_SHIFT))
#endif
--
2.55.0
```
</details>
`nvidia-bug-report.log.gz` is attached below, captured on the stock kernel.org 6.18.38 / 64 KiB boot while
reproducing.
I could not tick the template's "does not happen with the proprietary driver package" box: the closed
kernel module is not offered for this GPU, so that comparison is not available to me. The nearest control
is driver 580.x on the identical kernel, which is clean.