Inside amdgpu GART and IOMMU — How Zero-Copy GPU Access Works at the Kernel Level

You call malloc(4194304). You get back a virtual address — 0x7f8a4c000000. The kernel hasn’t even allocated physical pages yet. They’ll come on first touch, scattered across DRAM in 4 KB chunks. Your “contiguous” buffer is actually a random spray of physical addresses: 0x78F3000, 0x3A21B000, 0x112078000…

Now you want your GPU to read from this buffer. Not copy it to VRAM. Not bounce it through a staging buffer. Read it directly, in place, zero-copy.

Here is how the Linux kernel makes that happen, step by step, from malloc to the GPU’s first DMA read.

Part 1: The GPU Memory Model — Three Domains

The amdgpu driver manages memory in three distinct domains:

• VRAM is the GPU’s private high-bandwidth memory. The GPU accesses it directly at ~500 GB/s. The CPU can access it through the PCIe BAR, but at PCIe speeds (~32 GB/s). It is the fast path for GPU-only data.

• GTT (GART Table) is system RAM that has been mapped into the GPU’s address space through the GART page table. To the GPU shader, GTT memory looks just like VRAM — a flat, contiguous address space. Under the hood, the GART hardware translates each access through a page table that maps GTT offsets to physical system pages.

• System is ordinary system RAM that the GPU cannot see at all without GART mapping.

The critical insight: GTT is not a copy. It is a page table entry. When the GPU reads from GTT address 0x1000, the GART hardware walks its table, finds that 0x1000 maps to physical page 0x78F3000, and the GPU’s memory controller issues a DMA read to that physical address. The data was always in system RAM. The GPU just couldn’t see it until you gave it a map.

Part 2: GART — The Graphics Address Remapping Table

GART is the GPU’s equivalent of the CPU’s page table. It translates GPU virtual addresses (or GTT offsets, depending on the generation) into physical addresses that the GPU’s memory controllers can access over PCIe or the on-die interconnect.

Here is what happens on a GTT memory access, step by step:

1. Shader issues load to GPU virtual address 0x1000 (within the GTT aperture)
2. GPU TLB lookup — miss (first access)
3. GPUVM page table walk — the GPU has its own per-VMID 4-level page table (PML4 → PDP → PDE → PTE). The PTE contains a GART index
4. GART table lookup — the GART table entry contains a bus address: either a physical address (no IOMMU) or an I/O virtual address (with IOMMU)
5. IOMMU remapping (if active) — translates the IOVA into a physical address through the AMD-Vi page tables
6. DMA read — the GPU memory controller issues a PCIe or on-die read to the physical address
7. TLB fill — the translation is cached in the GPU TLB for subsequent accesses

GART vs GPUVM — Two Levels of Translation

On modern AMD GPUs (Vega and later), there are two translation layers:

Layer	What it translates	Unit	Managed by
GPUVM	GPU virtual address → GART index	Per-VMID, 4-level page table	amdgpu_vm.c
GART	GART index → bus address (IOVA or PA)	Flat table in VRAM	amdgpu_gart.c

The GPUVM provides per-process isolation. Each application gets its own VMID (Virtual Machine ID) with its own page table root. Two processes can both use GPU virtual address 0x1000 and map it to different GART entries — just like two CPU processes can both use 0x7f000000 and have different physical pages.

The GART provides the physical mapping. It tells the GPU hardware which bus addresses to DMA to. The GART table itself lives in a reserved region of VRAM (or, on APUs without discrete VRAM, in a carved-out region of system RAM) and is programmed by the kernel driver.

GART on APUs vs Discrete GPUs

On an AMD APU like the Ryzen 7 7730U, the GART model is simplified. The GPU sits on the same Infinity Fabric as the CPU. There is no PCIe bus between them. The GART entries point directly to system physical addresses (no IOMMU remapping needed because the GPU is a first-class citizen on the coherent fabric). This is why APU zero-copy is so fast — the GPU skips both the PCIe round-trip and the IOMMU page walk.

On a discrete GPU with an active IOMMU, every GART entry goes through an additional AMD-Vi translation. This adds latency but provides DMA protection: a buggy or malicious GPU shader cannot read arbitrary system memory; the IOMMU blocks accesses outside the mapped IOVA ranges.

Part 3: GEM and TTM — Buffer Object Lifecycle

Before the GPU can access anything, the kernel needs a buffer object. amdgpu uses the GEM (Graphics Execution Manager) API with TTM (Translation Table Manager) as the backing memory manager.

The lifecycle is:

1. Create: Userspace calls DRM_IOCTL_GEM_CREATE or imports a DMA-BUF via DRM_IOCTL_PRIME_FD_TO_HANDLE. The kernel creates a struct amdgpu_bo wrapping a struct ttm_buffer_object.

2. Place: TTM decides where the buffer lives. For GTT (system RAM mapped to GPU), it uses TTM_PL_TT. TTM allocates system pages, builds a scatter-gather table, and calls the driver’s bind callback.

3. Bind: amdgpu_ttm_backend_bind() allocates GART table entries, maps each page through the DMA API (dma_map_sg()), and writes the bus addresses into the GART table in VRAM.

4. Map: amdgpu_vm_bo_update() updates the GPU’s per-VMID page table with the GART entries. The GPU can now access the buffer at its GPU virtual address.

5. Evict: If TTM needs space, it can evict the buffer from GTT to SYSTEM domain. The GART entries are freed. The GPU page table PTE is marked invalid. On next GPU access, a VM fault occurs and the buffer must be re-bound.

Why TTM? The Eviction Problem

The GPU can have hundreds of buffer objects mapped through GART simultaneously — textures, command buffers, descriptor tables, staging buffers. GART entries are a finite resource (typically a few GB of addressable space on older GPUs, much larger on modern ones). TTM handles eviction: when GART space runs low, TTM picks a buffer (LRU policy), unmaps it from GART, and moves it to SYSTEM domain. When the GPU next touches it, TTM re-binds it — possibly to a different GART range.

This is the same problem the CPU’s virtual memory solves with swap. TTM is the GPU’s swap manager.

Part 4: The Full Zero-Copy Pipeline

Here is the complete flow from malloc to GPU DMA, with every kernel function involved:

The Kernel Call Chain

Here is the actual call chain through the amdgpu driver, annotated:

Userspace ioctl (DRM_IOCTL_GEM_CREATE or PRIME import)
  → amdgpu_gem_object_create()
    → amdgpu_bo_create()                  // allocate struct amdgpu_bo
      → ttm_bo_init()                      // TTM: create ttm_buffer_object
        → amdgpu_ttm_tt_create()           // allocate ttm_tt (TTM TT object)

GPU first access triggers binding:
  → ttm_bo_validate()                      // TTM: validate placement
    → amdgpu_ttm_backend_bind()             // bind to GTT
      → amdgpu_gart_bind()                  // write GART entries
        → amdgpu_gart_map()                 // per-page: GART[idx] = dma_addr
      → amdgpu_vm_bo_update()               // update GPUVM page table
        → amdgpu_vm_update_range()          // SDMA or CPU write to PT
          → amdgpu_vm_ptes_update()         // write PTE: GPU_VA → GART idx

The critical function is amdgpu_gart_bind(). Each call writes an array of dma_addr_t values into the GART table at a specific offset. For our 4 MB buffer with 4 KB pages, it writes 1024 entries. Each entry is 8 bytes — the GART table entry format is:

Bits [63:12]: Page frame number (physical or IOVA, 4 KB aligned)
Bits [11:0]:  Flags (valid, system, coherent, snooped, etc.)

The GPU hardware reads these entries during page table walks. The GPU TLB caches the full translation chain: GPU_VA → GART index → bus address. A TLB hit costs just a few GPU cycles. A TLB miss triggers a multi-level page table walk that can take hundreds of cycles and requires reading the GART table from VRAM.

Part 5: The DMA-BUF Path — Cross-Process Zero-Copy

The GEM/TTM path above works within a single process. What if Process A wants to share a buffer with Process B, or with a different device (a video encoder, a display controller, an FPGA)?

That is what DMA-BUF solves. It is the Linux kernel’s universal buffer sharing mechanism.

The key insight: the backing pages exist once. The dma_buf holds a single sg_table describing the physical pages. Each importer calls dma_buf_map_attachment() to get its own DMA mapping — a per-device view of the same physical pages. The IOMMU gives each device its own IOVA space, so Process A’s GPU might see the buffer at IOVA 0xAAAA0000 while Process B’s encoder sees it at 0xBBBB0000. Both map to the same physical pages. Zero copy, cross-device, cross-process.

Synchronization: dma_fence and dma_resv

DMA-BUF also provides synchronization. The dma_resv object attached to each dma_buf tracks which devices have outstanding reads or writes. When the GPU finishes writing to the buffer, it signals a dma_fence. Before the encoder reads, it waits on that fence. This is implicit synchronization — the kernel enforces ordering without userspace needing to manage it explicitly. Vulkan and modern OpenGL use explicit sync (timeline semaphores) instead, which DMA-BUF also supports through sync_file.

Part 6: HMM and SVM — The Future

The zero-copy paths above require pinning pages with pin_user_pages(). This locks pages in physical RAM — they cannot be swapped out, migrated, or reclaimed. For a 4 MB buffer, fine. For a 64 GB LLM, not fine.

HMM (Heterogeneous Memory Management) and SVM (Shared Virtual Memory) solve this:

In the SVM model with AMD IOMMUv2:

1. CPU and GPU share the same page tables. No separate GART. No pinning. The GPU uses the CPU’s mm_struct page tables via the IOMMU.

2. ATS (Address Translation Service) lets the GPU cache CPU page table entries in its Device TLB. A TLB hit means the GPU can translate virtual addresses without going through the IOMMU.

3. PRI (Page Request Interface) handles page faults. If the GPU accesses a virtual address that is valid but not present in RAM (swapped out, not yet allocated), the IOMMU reports a PRI fault. The Linux MM fault handler allocates the page, updates the page table, and the GPU retries.

4. mmu_notifier keeps the GPU’s Device TLB coherent with CPU page table changes. When the kernel swaps out a page, it notifies the IOMMU via mmu_notifier, which invalidates the corresponding Device TLB entry on the GPU. The GPU’s next access triggers an ATS miss → PRI fault → page-in.

This is the model that AMD’s ROCm stack uses for unified memory via hipMallocManaged(). The GPU and CPU share virtual addresses directly. No GART. No GUP pinning. The hardware handles faults, migrations, and coherence transparently.

The catch: ATS and PRI require modern GPU hardware (Vega 20 and later for full support) and add latency on first access. A GPU TLB miss that triggers an ATS request → IOMMU → PRI fault → MM page-in → ATS retry can take tens of microseconds — fine for bulk data access, catastrophic for fine-grained shared data structures. This is the fundamental trade-off of SVM.

Part 7: What This Looks Like on My Ryzen APU

On the Ryzen 7 7730U APU running CachyOS (kernel 7.1.6), the GART and IOMMU configuration is visible in sysfs and debugfs:

$ cat /sys/kernel/debug/dri/0/amdgpu_gtt_mm
GTT domain: 0x00000000-0xFFFFFFFF (4 GB aperture)
  used: 147 MB, free: 3.85 GB

$ cat /sys/kernel/debug/dri/0/amdgpu_vram_mm  
VRAM domain: 0x00000000-0x7FFFFFFF (2 GB aperture, carved from system RAM)
  used: 0 MB, free: 2 GB

$ dmesg | grep -i iommu
[    0.804390] iommu: Default domain type: Translated
[    0.804390] iommu: DMA domain TLB invalidation policy: lazy mode
[    0.917308] amdgpu 0000:04:00.0: amdgpu: IOMMU is active

The APU has no discrete VRAM, so VRAM is carved from system RAM (UMA frame buffer, typically 1-2 GB). The GART aperture maps the rest of system RAM into the GPU’s address space. The IOMMU is active in “Translated” mode (IOMMUv1 with DMA remapping, not SVM/PASID mode). Every GPU DMA access goes through an IOMMU page walk — the security-isolation overhead on an APU where the GPU is already on the coherent fabric. You can disable this overhead with iommu=pt (passthrough), but the kernel defaults to remapping for safety.

The actual GART table lives in the stolen VRAM region. Each 4 KB page of system RAM accessible to the GPU occupies one 8-byte GART entry. For a 16 GB system with all RAM mapped to the GPU, that’s 4,194,304 GART entries consuming 32 MB of VRAM — acceptable overhead.

The Bottom Line

Zero-copy GPU access is not magic. It is a carefully orchestrated dance between five kernel subsystems:

Subsystem	Role
GUP (get_user_pages)	Pin userspace pages in RAM
DMA Mapping	Convert physical pages to bus addresses (IOVA or PA)
GART	Write bus addresses into GPU-readable page table
GPUVM	Map GPU virtual addresses to GART entries per VMID
TTM	Manage placement, eviction, and LRU of buffer objects

Each “zero copy” GPU read walks through all five. The CPU never touches the data — it just sets up the signposts. The GPU follows them.

And on the APU in my laptop, the whole thing runs on the same Infinity Fabric, with the same MOESI coherence protocol keeping the GPU’s view of memory consistent with the CPU’s. The GART is just the GPU’s page table. The IOMMU is just the fabric’s bouncer. Together, they make a malloc’d buffer visible to a shader core without moving a single byte.