GPU hardware notes

The machines:

NVIDIA

A ring entry names one contiguous run of commands

A GPFIFO entry holds a pushbuffer segment's address and its length in words (NVC56F_GP_ENTRY0_GET, NVC56F_GP_ENTRY1_GET_HI and NVC56F_GP_ENTRY1_LENGTH in clc56f.h). The channel's DMA engine (PBDMA) reads that many words from that address, in address order. It knows nothing of the buffer the driver writes its segments into, so a segment that runs to that buffer's end continues past it.

On the RTX 5000 Ada, an entry that claimed 16 words starting 36 bytes before the end of a 1 MiB segment buffer ended both of the device's channels with the RM's channel error 32 (PBDMA_ERROR). A loop of 16-byte copies, each taking 64 bytes of segment, reached that position at exactly the 32,768th copy; the copies before it, and single copies of every memory kind, ran.

A driver that keeps its segments in a ring closes the open segment into its entry whenever its write position returns to the ring's start, also when the last command ended exactly at the ring's end, and opens a new segment there.

A channel's ring holds up to 2^31 entries

A channel's DMA engine holds the log2 of its GPFIFO's entries in a 5-bit field (LIMIT2, set from the allocation's gpFifoEntries), so a ring holds at most 2^31 entries. On the RTX 5000 Ada the RM made channels whose rings held 2^20, 2^24 and 2^31 entries, on the copy engine and, at 2^31, on the compute engine too, and 300,000 entries and a release then ran on each. A ring of 2^31 entries takes 16 GiB of host memory, which the RM made the channel over in about 3 s. Run in forked children on fudge, driver 615.

CUDA's compiled code reads the launch sizes from constant bank 0

Code that NVCC compiles reads blockDim and gridDim from the start of constant bank 0, an area the CUDA driver fills and no NVIDIA header describes. nvdisasm of NVCC 13.4's output shows where: for sm_89, blockDim x, y, z at c[0x0][0x0], [0x4], [0x8] and gridDim at [0xc] to [0x14]; for sm_120 (Blackwell), blockDim from c[0x0][0x360] and gridDim from [0x370]. On Ada the shared and local memory windows follow at 0x18 and 0x20, 64 bits each, and the stack limit at 0x28; kernels that use them run on the RTX 5000 Ada with these words.

A driver that launches through its own launch descriptors writes these words itself. Left at zero, a kernel that indexes by blockDim.x computes wrong addresses and does not fault: on the RTX 5000 Ada a 16-block launch wrote one block's elements, and 936 of 1,000 sums were missing. With the words written, a kernel that reads them saw the sizes set (test/conformance/test_conformance.ml, "a launch writes what its grid and parameters say, and a launch after it reads that"). The Blackwell offsets come from the disassembly alone; no Blackwell GPU ran them.

The kernel driver's structures change between releases

The parameters of the RM's escape ioctls and of the unified memory driver's ioctls change layout between release branches. Between 580, 610 and 615 these changed: the GPFIFO allocation (610), the channel group allocation (615), the schedule control (580), the framebuffer information control, the DMA map and the VA space allocation (580), and the unified memory driver's free and channel unregistration, which lost a field (610). Release 615's headers comment out two status codes. Established by generating the layouts from each release's headers, pinned by SHA-256, and comparing them (lib/nv/gen/gen.py). The headers also define some names twice under conditionals (the unified memory driver's ioctl base differs on Windows), so the generator evaluates them as a 64-bit Linux build does.

A driver reads the release from NV0000_CTRL_CMD_SYSTEM_GET_BUILD_VERSION_V2, uses that release's layouts, and refuses a release it has none for.

Unified memory needs a second file per process

UVM_MM_INITIALIZE ties the unified memory driver to the process's address space. It runs on a second open file of /dev/nvidia-uvm and names the first. Run on the first file, it is refused; registering a GPU's VA space then fails with NV_ERR_PAGE_TABLE_NOT_AVAIL, and no GPU opens (RTX 5000 Ada, test/nv/nvidia/test_nvidia.ml). The call is made once per process: a second one, such as CUDA's in the same process, is refused, and the first serves.

The unified memory driver also registers a GPU once per process. A driver keeps a GPU's RM device, subdevice, VA space and their registration for the life of the process, and a device's end releases only its channels' registrations (test/nv/nvidia/test_nvidia.ml, "a GPU has one device until it is stopped").

A map into an empty region pays for its page tables

On the RTX 5000 Ada, the RM's own map, NV_ESC_RM_MAP_MEMORY_DMA, of a 64 KiB allocation takes 21.4 µs when nothing else is mapped in its region of the GPU's address space, and 8.6 µs with the same parameters when one other allocation lives there. Measured per ioctl with an LD_PRELOAD counter. We read the difference as the RM building the region's page tables at its first map and freeing them at its last unmap. A 64 KiB allocation and free through that map took 73.5 µs alone and 46.0 µs beside one live 4 KiB allocation. No path makes that map now (the next note).

The unified memory driver's map pays less for an empty region. Through it alone, a 64 KiB allocation and free of Device memory takes 50.3 µs alone and 42.1 µs beside one live 4 KiB allocation, and of Pinned or Mapped memory 46.8 µs and 56.8 µs either way. Measured on fudge, medians of 7 runs of 200, in two runs.

A driver that allocates and frees small Device buffers in an otherwise empty region pays this every time. Long-lived memory of the driver in the same region keeps the tables; allocation benches hold one live allocation of the kind, as a running program does.

The RM's own mappings stop at about 39 GiB a GPU

NV_ESC_RM_MAP_MEMORY_DMA maps a memory object into a device's virtual memory object (NV01_MEMORY_VIRTUAL), in the RM's own page tables. An object made with no address space named (hVASpace 0) lies in the device's default address space, which the RM makes for a device whose client shares it with itself (vaspaceGetByHandleOrDeviceDefault, deviceInitClientShare; read in source, open-gpu-kernel-modules 570 and 615). A channel runs in the address space its context share names, and rig's name the externally owned FERMI_VASPACE_A that the unified memory driver registers and fills, so its work never reads the default one's mappings. On the RTX 5000 Ada (32 GiB) in a machine of 62 GiB, with every map in 4 KiB pages, the RM refuses it with NV_ERR_NO_MEMORY once about 39.04 GiB is mapped this way for the GPU, counting GPU memory and host memory alike, and every process's: 39 maps of 1 GiB of one shared memory file, 156 of 256 MiB or 2,498 of 16 MiB; 15 GiB of host memory beside 24 GiB of GPU memory; 9 GiB in a second process while the first held 30 GiB. A map succeeds again once others are freed. The unified memory driver's UVM_MAP_EXTERNAL_ALLOCATION maps the same objects into the channels' address space, from the memory's own description, and with no RM map it mapped 160 maps of 1 GiB of one file, and 70 beside 24 GiB of GPU memory. CUDA page-locked 61 maps of 1 GiB of one file, beside 24 GiB of its own GPU memory or not, then answered CUDA_ERROR_OUT_OF_MEMORY. Measured from forked children on fudge, driver 615.

The NVIDIA kernel path maps its memory through the unified memory driver alone, so no RM map bounds it.

Rings in host memory cost less host time than rings in BAR1

The RM accepts a channel's GPFIFO ring and USERD in host memory (an OS-descriptor allocation) as well as in GPU memory, and a one-page coherent host buffer as the channel's error notifier. On the RTX 5000 Ada:

Rings in BAR1 save 0.2-0.5 µs of latency, since the GPU fetches from its own memory, and cost 0.9 µs of host time per submission: uncached stores, and the BAR1 read that makes them land before the doorbell. The NVIDIA driver keeps its rings in host memory, and pays the BAR1 read only while memory the host writes through BAR1 is live: 0.66 µs per submission.

Device-to-host copies run at 60% of the link

On the RTX 5000 Ada (PCIe 4.0 x16, 31.5 GB/s), copy engine copies of 256 MiB between pinned host memory and GPU memory run at 25.2 GB/s host to device and 18.9 GB/s device to host, through the RM and through CUDA alike (CUDA: 25.1 and 19.1 GB/s). CUDA stages pageable host memory and copies it to the device at 18.5 GB/s. Within GPU memory CUDA copies at 241.8 GB/s, 84% of the 288 GB/s that 576 GB/s of DRAM bandwidth leaves for a read and a write. Copy benches gate on these figures, measured from C over the vendor's interface.

GSP firmware is an ELF file read by section

gsp-570.144.bin, the GSP-RM firmware of driver 570.144, is a 63.5 MB ELF file of 14 sections whose addresses (sh_addr) are all 0. Its payload is the .fwimage section, the only one allocated. Laid out as a loadable image, the file makes a 63 MB copy that no boot step reads. A loader reads the sections it needs by name and copies .fwimage once, into the memory the GSP boots from. fmc-570.144.bin is a 32-bit ELF file of 6 sections, also all at address 0. Read from the firmware files.

CUDA

After a sticky error, queued work does not run (driver 615)

CUDA's reference says that after a sticky error such as CUDA_ERROR_ILLEGAL_ADDRESS every later call in the context returns it. It does not say whether work queued behind the fault still runs. On the RTX 5000 Ada with driver 615.71.09 it does not. After a kernel stored to address 0, a write of the timeline word queued behind it and a copy into a watched pinned buffer queued behind that never ran; for 200 ms of sampling the word and the buffer stayed still. cuStreamQuery reported the error within the driver's 1 ms poll, a new context on the GPU in the same process was refused with the same error, and the GPU was back at its idle power state, with no memory in use, 11 s after the process exited. Every run of the fault test showed this, one of them under ASan and UBSan (test/cuda/test_fault.ml, "a fault of the work is CUDA's report").

A driver that reports a device stopped after a sticky error, promising that its work writes no more memory, rests on this behaviour of one driver release.

Module loads and unloads wait for all of the GPU's work

cuModuleUnload returns only once every kernel running on the GPU has finished, whichever module it belongs to: on the RTX 5000 Ada an unload took 950 ms behind a 1 s kernel of another module. cuModuleLoadData waits the same way. A binding to a runtime with a global lock releases it around both, or every other thread waits with them (test/cuda/test_cuda.ml, "unload lets other domains run while CUDA waits for the GPU (sampled)"). Measured while a kernel ran, cuModuleGetFunction took 31 µs and cuMemHostGetDevicePointer 1 µs: neither waits for the GPU.

A page-locked range may span unlocked pages

A driver that accepts host memory another owner page-locked must check that the whole range is locked. Asking CUDA about the range's two ends is not enough: on the RTX 5000 Ada, with pages 0 and 2 registered by two calls and page 1 not, both ends answered and the middle page returned CUDA_ERROR_INVALID_VALUE. CU_POINTER_ATTRIBUTE_RANGE_START_ADDR gives the start of the allocation that holds a pointer, for registered memory and cuMemHostAlloc memory alike, so the driver takes a range as locked only when both ends report the same start (test/cuda/test_cuda.ml, "map_host shares a host range both ways"). The GPU reports CAN_USE_HOST_POINTER_FOR_REGISTERED_MEM as 1, so one pointer from cuMemHostGetDevicePointer serves every device.

A stream memory-operation wait costs 1.6 µs of host time

On the RTX 5000 Ada, cuStreamWaitValue64 costs 1.62 µs of host time per call, cuStreamWaitEvent 61 ns and cuEventRecord 69 ns, for events made with CU_EVENT_DISABLE_TIMING and without CU_EVENT_BLOCKING_SYNC. A round trip that switches streams takes 4.86-4.93 µs ordered by a memory-operation wait and 3.64-3.71 µs ordered by an event. Four satisfied waits cost 6.46-6.54 µs as four calls and 1.85-1.87 µs as one cuStreamBatchMemOp, whose count the reference requires below 256. Measured from C over libcuda alone, medians of three runs.

The CUDA driver orders work within a device by events, and waits on other devices' words in batches of at most 255.

A copy beside a kernel runs at its own speed

On the RTX 5000 Ada, an upload of 256 MiB from page-locked host memory on a second stream took 10.69 ms beside a 10 ms launch of 600 groups of 256 threads, enough to fill every multiprocessor: as long as it took alone. The two one after the other took 20.69 ms. Measured through rig.cuda's two streams on fudge, medians of 9. A driver that runs its copies on a stream of their own beside its kernels halves such a pair's time.

AMD under KFD

KFD refuses a queue whose context save area is short of its rule

AMDKFD_IOC_CREATE_QUEUE refuses a compute queue whose context save area is smaller than KFD's own size (kfd_queue_ctx_save_restore_size in kfd_queue.c). That size is, for each die (XCC), the topology's cwsr_size plus a debugger area of 32 bytes per wave rounded up to 64 bytes, the total rounded up to a page. KFD counts 32 waves per compute unit from GFX 10.1, and before it 40 per compute unit, up to 512 per shader engine, whatever the compute unit runs: an area sized from the topology's max_waves_per_simd was refused on the R9700 (test/amd/amdgpu/test_amdgpu.ml, "an R9700's: 2048 waves of 64 compute units"). The queue's control stack size is the topology's ctl_stack_size.

A host data path flush from a GFX12 compute queue hangs it

On the R9700 under amdgpu, a compute queue that flushed the host data path (HDP) itself, writing GPU_HDP_FLUSH_REQ and waiting on GPU_HDP_FLUSH_DONE with WAIT_REG_MEM, hung within a few hundred batches run back to back, and the kernel driver then reset the GPU for every process on it. At the hang no wave ran and the MEC waited on a register read (CP_CPC_STALLED_STAT1.MEC1_WAIT_ON_RCIU_READ). The request sets every client's bit in one write, in a handshake the kernel driver's own engines use one bit each. Of four runs of 1,000 back-to-back calls with the flush in the queue, three hung; of five without it, none (tolk's test/runtime/ops_amd/test_ops_amd_exec.ml, "a thousand runs back to back while the host allocates"). A driver flushes the HDP from the host, as ROCr does.

Host writes through the BAR land after an HDP flush read back

GPU memory that the host writes through the BAR passes the HDP, which can hold the writes. amdgpu flushes it by writing the remapped HDP_MEM_FLUSH_CNTL register and reading it back (amdgpu_hdp_generic_flush); the read returns once the write has landed. The AMD driver does the same before a doorbell, and only while memory the host writes through the BAR is live: its rings, words and segments live in host memory, which the HDP does not carry. Read in source; the R9700's host has no host-visible GPU memory (below).

Without resizable BAR no GPU memory is host-visible

With resizable BAR off, the R9700's BAR0 is 256 MiB (512 MiB at most once resized, below) and KFD's topology lists no memory bank of the public heap type (heap_type 1), only private ones. Memory the host writes in GPU memory through the BAR does not exist on that machine, and a driver falls back to host memory: reading 1 MiB took 0.041 ms from both, where reading it across the bus would take milliseconds. A driver decides from the topology's heaps, and takes a BAR as small when it is smaller than the GPU's memory.

An AQL packet is live once its header is

An AQL queue's packet processor may take a packet as soon as its header type is valid, before any doorbell (the HSA runtime specification's packet header rule). A writer stores words 1 to 15 of each 64-byte packet, then word 0, which holds the header, with a release store. A writer that drops packets it already placed beyond the write position, after a failed fill, stores the INVALID type (1) into each dropped header, since the processor may already be reading them. The doorbell carries the index of the last packet written, one less than the write index. Unverified on hardware: the R9700 runs PM4 compute queues, and no GPU here takes AQL packets.

Write positions count dwords, bytes or packets

A queue's write position, and the value its doorbell takes, count dwords on a PM4 queue, bytes on an SDMA queue and packets on an AQL queue. The PM4 and SDMA units run on the R9700; the AQL unit is unverified there.

The copy engine's fence writes 32 bits

SDMA's FENCE packet writes 32 bits, and SDMA's POLL_REGMEM and PM4's WAIT_REG_MEM compare 32 bits. A 64-bit timeline word released by two fences, the low half and then the high half when the low half wraps, reads lower between the two writes: 0x0_ffffffff becomes 0x0_00000000 before 0x1_00000000, and a waiter comparing 64 bits sees the word move backwards. The AMD driver releases on the copy queue only values whose high half is unchanged, and sends the one value in 2^32 whose low half is 0 to the compute queue, whose RELEASE_MEM writes 64 bits. A wait through a 32-bit compare is exact only if the value it waits for cannot recur within 2^32 values. On the R9700, a device started near 2^32 released across it in order (test/amd/test_amd.ml, "values across 2^32 complete in order, ending on either queue").

A TRAP after the copy queue's fence makes it visible sooner, mostly

On the R9700, a copy-queue submission of a 16-byte copy and a FENCE of a word the host spins on completed in 10.8 µs without a following SDMA TRAP packet and in 3.2 µs with one, in the same run. The AMD driver ends every copy-queue release with TRAP.

The trap does not always work. Some copies with it take 11-12 µs, as late as a copy without it, against 3.0-3.6 µs for the rest. Slow copies come in clusters of 4 to 6 consecutive copies every 1.9-2.1 ms, in every process, and in episodes of 10 to 100 ms in which half to 97% of copies are slow. Another process holding KFD queues on the GPU, even idle, sets off such episodes: with one, 3 of 5 processes ran 25-50% slow for their first 1,300-1,500 copies; alone, none had more than one slow window of 100 copies. Episodes also came with no other process seen. Ruled out: the spinning CPU (pinned to seven different cores), load on the CPU that takes the GPU's interrupts, the teardown of a process that had just exited, the driver (raw KFD queues show the same modes), and SMU deep sleep, whose features KFD disables while a process has queues (pp_features). The cause lies below KFD, in what the firmware schedules, and its registers (MES, IH, SDMA) cannot be read without root. Measured with a raw KFD copy queue, 16-byte copies one at a time, slow meaning above 8 µs.

A latency measured on the copy queue has two modes between processes. Benches compare such rows in alternating pairs that report the range, and gate on the fastest worker's median.

A release does not stand in for the next acquire

A compute queue's RELEASE_MEM performs its cache actions at the end of the pipe, after the packets that follow it have started, and its actions cover the L2 and the caches between it and the shaders (GL2, GLM, GL1, GLV), never the scalar cache (GLK), which holds kernel arguments, nor the instruction cache (GLI). So the next submission's first dispatch can read an L2 line that the previous release has not yet invalidated, or stale arguments or code. Each submission that runs compute work starts with its own ACQUIRE_MEM; letting a release stand in for it would need a wait for the release's word, a PCIe round trip.

On the R9700 the acquire costs 0.23 µs per pipelined submission at either scope: at agent scope (the caches above the L2) it costs the same as at system scope (the L2 and instruction cache too). The cost is the packet. A release that only writes the L2 back (GL2_WB and SEQ, the fence amdgpu's gfx12 code writes) costs 1.10 µs per pipelined submission, against 1.28 µs for one that also invalidates. Measured on a raw KFD compute queue, medians of 100 batches, variants alternated in one process.

Dispatch and descriptor fields keep the low bits of what does not fit

PM4 dispatch registers and buffer descriptors take the low bits of a value too large for their field, without error, and the kernel runs with the truncated value. A driver bounds each before it writes:

• COMPUTE_PGM_RSRC2.LDS_SIZE is 9 bits, in units of 512 bytes on every generation the library encodes except GFX950, which counts units of 1280 bytes (AMDGPUUsage's LDS_SIZE row; LLVM's FeatureLDSEncodingGranularity1280). A larger request dispatches with almost no LDS (test/amd/abi/test_pm4.ml, "a GFX950 workgroup's LDS, in 1280-byte units" and "a workgroup's LDS is set up to 511 units and refused past them").
• COMPUTE_TMPRING_SIZE.WAVESIZE gives a wave's scratch in granules: 13 bits of 1 KiB on GFX9, 15 bits of 256 bytes on GFX11, 18 bits on GFX12, at most 131,056, 131,068 and 1,048,572 bytes per lane. Past it the waves' scratch wraps and they overwrite each other. LLVM compiles no more than this (getMaxWaveScratchSize in GCNSubtarget.h) (test/amd/abi/test_scratch.ml, "a lane's scratch is taken up to what WAVESIZE holds").
• A buffer descriptor's record count is 32 bits. A scratch share of 2^32 bytes becomes 0, which faults the kernels that use it (test/amd/abi/test_scratch.ml, "a descriptor holds the base, below 2^48, and n / xccs bytes, below 2^32 a die").

Two AMD sources disagree on counter events

AMD publishes performance counter event numbers in two places: the SOC enumerations that aqlprofile ships (vega10_enum.h, soc21_enum.h, soc24_enum.h) and rocprofiler-compute's counter_defs.yaml. They disagree. For gfx942 and gfx950, 60 of the 91 SQ event names the two share have different numbers; for GFX12, GL2C_HIT is 39 in the enumeration and 41 in the YAML. The AMD ABI library's counter table comes from counter_defs.yaml, the one AMD's profiler counts with (lib/amd/gen/gen.py). The files were compared; no gfx942 or gfx950 GPU ran the counters here.

CDNA drops completion interrupts whose context id is 0

On gfx9, KFD discards an end-of-pipe interrupt with context id 0 whenever it expects valid ids, which it does for every gfx9.4.x (MI200, MI300): a workaround for firmware that sent bogus signals with id 0 (kfd_int_process_v9.c:333-345). A process sleeping in AMDKFD_IOC_WAIT_EVENTS then wakes only at its timeout. On gfx11 and gfx12 an id of 0 misses the fast lookup and makes KFD scan the process's signal page (kfd_int_process_v11.c:355, kfd_events.c:707-735). A release that raises an interrupt carries its KFD signal event's id in INT_CTXID, as ROCr does. Read in source; no CDNA GPU here.

Unmapping registered host memory stalls every queue

When a process unmaps host memory that KFD still tracks before freeing it in KFD, the kernel invalidates the range, and KFD evicts every queue of the process and restores them later. On the R9700 the next work on any queue waited 5 to 10 ms. A driver frees the KFD allocation and its GPU mapping first, then unmaps. Memory that KFD allocates itself (GTT) is never evicted this way, so the AMD driver keeps its rings and words there.

KFD's wait timeouts round up to timer ticks

AMDKFD_IOC_WAIT_EVENTS converts its timeout in milliseconds to timer ticks and rounds up (user_timeout_to_jiffies in kfd_events.c). At 250 Hz a 1 ms timeout lasts at least a 4 ms tick when no event fires. A code load that polled a copy's completion with 1 ms sleeps took 8 to 12 ms on the R9700 for a copy of microseconds, since the copy raised no event. A driver spins on the word before it sleeps, and sleeps in KFD only on work that raises an event. A wake through the interrupt path costs about 36 µs more than spinning (wake/driver against release/driver in bench/amd).

A process has one GPU address space per GPU

KFD keeps one GPU address space per GPU per process, bound to one render node file by AMDKFD_IOC_ACQUIRE_VM; the process cannot hold a second one on the same GPU through another file. A memory fault in that address space reaches every user of the GPU in the process. The AMD bench opens each case's device in a worker process of its own for this reason.

The stable power state has one holder

amdgpu holds a GPU in its stable power state (AMDGPU_CTX_OP_SET_STABLE_PSTATE) for one context at a time and refuses another with EBUSY. The AMD driver takes it for thread traces, which need steady clocks and shader engines, once per process under a lock, so that two threads cannot both create a context and ask (lib/amd/amdgpu/rig_amd_amdgpu.ml, stable_power).

A GFX12 compute queue's 64-bit wait compares all 64 bits

WAIT_REG_MEM64, which the kernel's headers name but do not lay out (PAL's PM4_MEC_WAIT_REG_MEM64 does), compares the whole 64-bit word on the R9700 with MEC firmware 3010. A probe had a queue wait on a host word for 300 ms in five cases, releasing any that held with a write a 32-bit reading would also pass: 2^32 − 1 ≥ 2^32 held until the word became 2^32; 2^33 ≥ 2^32 + 5 and 2^32 + 5 ≥ 0xFFFFFFFF passed immediately; 5 ≥ 2^32 + 3 and 2^32 + 7 ≥ 3·2^32 held. A compare of the low 32 bits gets four of the five wrong. The AMD suite's waits group keeps the fact: a wait across 2^32 holds through a word whose low half is above the target's. The AMD driver waits in its queue on another device's word only from a firmware version seen to do this, per GC version (wait64_from in lib/amd/rig_amd.ml); elsewhere the wait is the host's.

A KFD process can outlive its process and pin amdgpu

On the R9700, KFD kept a process for 44 hours after it exited: a directory under /sys/class/kfd/kfd/proc/ with 320 KiB of GPU memory and no queues, and a client of the GPU's DRM device in debugfs's clients list, named <unknown>, whose process did not exist. No live process held a KFD, DRM or dma-buf file. While it stayed, Rig_amd_pci.detach refused the GPU (the kernel holds a file of its DRM device for that process) and modprobe -r amdgpu answered that the module is in use. Only a reboot cleared it. What left it is not known.

AMD over PCI

On the R9700's host BAR0 grows to 512 MiB at most

The R9700's BAR0 offers every size from 256 MiB to 32 GiB (resource0_resize reads 0xff00), and the host's firmware assigns it 256 MiB below 4 GiB. Asked for the largest size its bridges hold, as Rig_amd_pci.detach asks, Linux refused each size from 32 GiB down to 1 GiB, with "can't assign; no space" for every bridge up to the root port, and assigned 512 MiB at 0xc0000000, the bridges' prefetchable window becoming 0xb0000000–0xdfffffff. The host gives the GPU no 64-bit window above 4 GiB, likely because its firmware's Above 4G decoding is off; that setting was not read. amdgpu, bound after the resize, reports BAR=512M against 32,624 MiB of memory. A driver-less boot of this GPU therefore reads the discovery table a word at a time through MM_INDEX, MM_INDEX_HI and MM_DATA, and places its page tables in their own pool (Page_table.Pool).

The R9700's blocks

The R9700's discovery table, read through the index window after a detach, lists GC 12.0.1, SDMA0 and SDMA1 7.0.1, MP0 and MP1 14.0.3, MMHUB 4.1.0, OSSSYS 7.0.0, NBIF 6.3.1, HDP 7.0.0, ATHUB 4.1.0, DF 4.15.1, UMC 8.14.0, PCIE 6.1.0, LSDMA 7.0.1, DMU 4.1.0, SMUIO, THM, FUSE and CLKA 14.0.2, with no instance fused off. Its GC has 4 shader engines of 2 arrays of 8 compute units, 64 in all, 16 waves per SIMD and 64 KiB of LDS. Its register BAR (BAR 5) holds 131,072 words, so the registers a boot's plan reads (SCRATCH_REG6, SCRATCH_REG7, the GC hub's fault status) lie inside it, with no access through the RSMU window. The boot report finds the eight images rig_amd_pci.mli names for it, and each is laid out as its format says.

amdgpu leaves the security processor's OS running

Right after amdgpu unbinds from the R9700, the security processor's OS runs (its C2PMSG_81 is not zero), its bootloader is not ready, and SCRATCH_REG6, SCRATCH_REG7 and the GC hub's fault status read 0: no mark of this library, so a driver-less open plans a reset first (Booted in lib/amd/pci/boot.mli). Rig_amd_pci.reset resets it whole in 1.1 s with bus mastering off throughout, and its configuration reads back as saved. amdgpu's probe resets the GPU itself (mode 1) when it finds the OS running, as it does at every rebind and module load after its own session; after Rig_amd_pci.reset and Rig_amd_pci.attach its probe initialized the GPU with no reset of its own, which shows that no firmware ran.

Apple GPUs under Metal

Stores to unmapped GPU addresses do not fault

On the M1 Max, a compute kernel that stored 64 words at GPU address 0x10, at 0x7f0000000000, which nothing maps, or at the GPU address of a buffer freed and removed from the residency set before the submission, completed: its command buffer reported no error, later work ran, and the GPU stayed usable. Two guarded runs, one process each, on 2026-10-08. Metal tells a driver nothing of a stray store, and a test cannot make a page fault this way; the Metal driver's failure path is tested through its C completion ring without a GPU.

Completion handlers run in no documented order

Metal calls each command buffer's completed handler on a thread of its own and documents no order between the handlers of one queue. A handler that writes a timeline word directly can pass an earlier command buffer that then reports failure. The Metal driver gives each command buffer a slot before making it; a handler completes its slot, and the leading completed slots are released in commit order, each writing its value unless a slot before it failed (test/metal/test_metal.ml, "releases in commit order whatever order slots complete in"). There are as many slots as the queue holds command buffers, so taking a slot is the back-pressure, and Metal's commandBuffer never blocks.

An indirect command buffer's commands are autoreleased

indirectComputeCommandAtIndex: returns its command autoreleased. A thread without an autorelease pool of its own, such as an OCaml domain, never drains one, so each command lives until the thread ends. Code that records an indirect command buffer wraps the calls in its own @autoreleasepool (test/metal/test_metal.ml, "an icb's buffer and commands live until it is unreachable").

Buffers start at multiples of 256 bytes

Apple documents no alignment for a buffer's first byte. On the M1 Max, shared buffers smaller than a page start at multiples of 256 bytes inside a 16 KiB page (seen at offsets 7936, 8192, 8448 and 12544). The Metal driver promises 256 bytes and checks every buffer Metal returns (test/metal/test_metal.ml, "an allocation starts at a multiple of 256 bytes").

Mac GPU families publish no argument offset alignment

A dispatch's argument buffer offset must be a multiple of the GPU's minimum constant buffer offset alignment. Apple's feature set tables give 4 bytes for the Apple GPU families and nothing for the Mac families. The Metal driver uses 256 bytes on Mac families, a multiple of every smaller power of two, and the record compiled code reads carries the value (test/metal/test_metal.ml, "an Apple GPU aligns arguments to 4 bytes").

A dispatch costs about 200 µs round trip, mostly outside the GPU

On the M1 Max, one dispatch committed and waited with waitUntilCompleted took 206 µs at the median of 300 interleaved rounds: 98 µs from commit to GPU start, 9.6 µs on the GPU, 97 µs from GPU end until the host saw completion. A residency set, a fence, or a completion handler with the host spinning on a word each changed it by under 5 µs. A submission costs about 5 µs of CPU, of which creating the encoder takes 2.0 and the commit 1.8. Measured under load: an empty command buffer completes in 15-46 µs; pipelined, a command buffer costs about 25 µs; after 20 ms idle, commit to GPU start grows to about 390 µs.

A kernel can write a shared word that the host spins on, which the host sees 98-130 µs after the commit, earlier than the completion. The Metal driver does not take it as completion: Metal reports a failed command buffer only through its status.

A command buffer costs about 26 µs, a dispatch in one about 2 µs

On the M1 Max, command buffers of one one-thread dispatch each, committed back to back, start 31 µs apart on the GPU, each busy 5-7 µs: the host sustains one every 26 µs at most, at 6 µs of CPU each. With 3 in flight, as rig's replay rows run, each costs about 55-60 µs, a third of the round trip. The same dispatches in one serial encoder of one command buffer cost 2.3 µs each, 2.5 µs with an encoder each, and 30 µs with an encoder each that waits for a fence and updates it. A fence at each command buffer's encoder costs under 5 µs.

The GPU's clock follows its load. After 300 ms idle, 3,000 such command buffers with 3 in flight cost 80-145 µs each and never sped up; with 64 in flight they reached 29 µs within 1,000. Rows that run 100 command buffers after their setup idles measure the low clock: 77-124 µs.

An indirect command buffer adds about 16 µs of GPU time

On the M1 Max, running dispatches from an indirect command buffer takes about 16 µs more GPU time per command buffer than encoding them directly, 10-20 µs more latency on each side of the GPU, and 1.9 µs per dispatch that waits on the one before. None of these moved it: the default inherit flags, optimizeIndirectCommandBuffer, the buffer in the residency set or passed to useResource, a concurrent encoder. Direct serial encoding costs 0.1 µs of CPU per dispatch and wins in probes of 1 to 512 dependent dispatches.

On real steps of gpt-oss-20b the result depends on the size. Decode steps of 656 dispatches ran 3% faster encoded directly (31.3 against 32.2 ms); prefill steps of 3,452 dispatches ran 12-15% slower (1.21 s against 1.04-1.09 s). Encoding costs 0.5-0.7 µs per dispatch either way. The crossover lies between those sizes and was not located.

An indirect command buffer needs no pipeline set on the encoder

Some runtimes set each pipeline of an indirect command buffer on the encoder, with an empty dispatch, before running it on GPUs before the Apple9 family, after GPU faults seen on M1 and M2 Macs in 2024. On the M1 Max it is not needed. An indirect command buffer of five dispatches over three pipelines, each pipeline from an image loaded for that run and never set on an encoder, ran as the only command of its command buffer 520 times across 13 processes and wrote every word, in two of the processes under Metal's API or shader validation. A runtime that sets the pipelines ran, without that step, 988 indirect command buffers whose pipelines were made apart from those of its direct dispatches; each matched the direct dispatches' result. Neither the Metal driver nor the fills of its tests set an indirect command buffer's pipelines.

Memory is wired again after 1 to 3 seconds idle

After 1.2-3 s without work, the first submission on the M1 Max waits before the GPU starts, in proportion to the residency set's memory: 20-76 ms for 1 GiB, 2.5-8.6 ms for 64 MiB, 1.4-3.2 ms for 16 KiB, against 80-220 µs warm. The size-dependent part is macOS wiring the memory again; the 1-3 ms floor is the GPU waking. requestResidency called just before the commit changes nothing. Only work or a residency request well before the submission hides it.

Completion wakes are bimodal per process

On the M1 Max, a thread blocked on a condition that a completion handler signals wakes about 10 µs after the handler in most processes; in about one process in four, every wake comes 1.0-1.6 ms after it (p90 4 ms). GPU end to handler stays 92 µs in both modes, and waitUntilCompleted shows both. Raising the waiting thread to the user-interactive QoS class, by an override or by pthread_set_qos_class_self_np, moved neither mode; a command-line process's main thread already runs at that class. The cause was not found.

Host code

arm64 macOS writes JIT code through a listed callback

On arm64 macOS, code memory is MAP_JIT memory, mapped read-write-execute, whose write protection each thread lifts for itself. Under the hardened runtime's com.apple.security.cs.jit-write-allowlist entitlement, pthread_jit_write_protect_np ends the process at its first call. Code is written instead by a callback passed to pthread_jit_write_with_callback_np and listed with PTHREAD_JIT_WRITE_ALLOW_CALLBACKS_NP; without the entitlement the system lifts the protection around the callback as the pair of calls did. The system reads only the lists of images present at start-up, so a bytecode program that loads its C stubs at run time still ends at its first write, unless it holds jit-write-allowlist-freeze-late and calls pthread_jit_write_freeze_callbacks_np after loading them. Checked by hand on the M1 Max with an executable signed with the hardened runtime, allow-jit and jit-write-allowlist: killed with SIGKILL at the first link while the loader called pthread_jit_write_protect_np, passing with the callback.

arm64 macOS can start a child of fork with its JIT writes allowed

On arm64 macOS, a child of fork can start with its thread's MAP_JIT write protection lifted, although the forking thread had set it again. The child's first call of code in the mapping then dies of SIGBUS, an instruction abort on the writable mapping. Apple's pages for pthread_jit_write_protect_np(3) and fork(2) define the protection per thread and say nothing of fork. Checked by hand on the M1 Max at load ~68, with a C program that repeats a step: map 16 KiB of MAP_JIT memory, write mov x0, #42; ret between pthread_jit_write_protect_np(0) and pthread_jit_write_protect_np(1), invalidate the instruction cache, fork, and call the code in the child. It counts the children that end by a signal:

The second row shows that the child does not take the state its parent's thread ran with, and the third that the state stays wrong across a context switch. A second run with the last row's reset had no failure in 16,000 forks. The failures need a loaded machine; in rig's suite they killed the forked child of test/remote/proxy/test_link.ml twice in about 500 runs. The first code mapping of lib/rig_host_code.c registers a pthread_atfork child handler that makes that reset through the loader's listed callback, before the child runs anything: pthread_jit_write_protect_np would end a process that holds jit-write-allowlist.

arm64 cores need an isb to run code another core wrote

Cache maintenance after writing code (__builtin___clear_cache) reaches every core, but a core that fetched instructions earlier, such as those of a collected program once mapped at the same address, may run them until it takes a context synchronization, such as isb (Arm Architecture Reference Manual, instruction cache coherency). A pool worker running job after job takes none. The host loader counts installs with a release increment after the maintenance, and a thread executes isb before entering linked code, from a host launch or a driver's call of a fill (rig_code_sync), when the count moved since its last one, a single load when nothing changed. x86_64 keeps instruction fetch coherent with stores.

Windows commits a thread's stack one page at a time

Windows grows a thread's stack through a guard page, so a function whose stack frame exceeds 4 KiB must touch each page in order, which Windows compilers do with a stack probe (__chkstk). Code compiled for an ELF target has no probes: rune's incomplete-beta derivatives, with frames of 18 to 52 KiB, faulted on their first store. So before a thread first runs linked code, the host commits the rest of its stack but for its last 64 KiB, which stay reserved so that an overflow still faults (lib/rig_host_code.c, commit_stack), and a frame of any size within it runs. Clang's stack clash protection would probe in the code instead, but it crashed Ubuntu's Clang 18 on some kernels and would have to be asked of every compiler that targets rig. The host fixture deep.c, a frame of 256 KiB written first at its far end, runs on the caller's thread and the pool's.

Memory ordering

On x86, the barrier for device memory is mfence

A C11 sequentially consistent fence compiles to mfence with clang and to lock orq $0, (%rsp) with gcc 14. Linux's mb(), its barrier for device memory, is mfence, so rig's is mfence written out (lib/pci/rig_pci.h, rig_pci_barrier), whichever compiler builds it. It costs about 4 ns more: 8.2 against 4.0 ns per fence on the RTX 5000 Ada's host, 7.4 against 3.9 ns on the R9700's, beside bus round trips of hundreds of nanoseconds. On arm64 the barrier is dsb sy.

Before a doorbell: fence, publish, read back, ring

The GPU must see every store it will read before the doorbell. Uncached stores through a BAR are posted, and PCIe keeps them in order. Write-combined stores may reach the device in any order and after later stores, and AMD's HDP can hold writes after they arrive. A read through the BAR returns only after earlier posted writes to it have completed. The NVIDIA and AMD drivers submit with these steps, after UVM's uvm_hal_turing_host_write_gpu_put and amdgpu:

1. write the ring entries and command segments, then a fence: a store fence suffices (sfence on x86_64, dsb st on arm64: Linux's wmb()), and the AMD driver uses a full one;
2. publish the write position (NVIDIA: GP_PUT in USERD);
3. while memory the host writes through a BAR is live, a full fence and a read through that BAR (NVIDIA: BAR1, after step 2), or the HDP flush and its read back (AMD, before step 2);
4. ring the doorbell (NVIDIA: the channel's work submit token, written to NVC361_NOTIFY_CHANNEL_PENDING).

On the RTX 5000 Ada step 3 adds 0.66 µs per submission.

arm64 Device memory faults on the accesses memcpy makes

On arm64, an unaligned access to Device memory raises an alignment fault, as does DC ZVA (Arm Architecture Reference Manual). The C library's memcpy and memset use unaligned and overlapping accesses, and DC ZVA to zero, so they fault on a BAR mapped as Device memory. Accesses to register windows are loops of 32- or 64-bit loads and stores aligned by address, since a window may start off a word (test/pci/test_window.ml, "a mapped word is aligned by its address"). Not run: no arm64 host with a discrete GPU here.

Address spaces and sanitizers

NVIDIA command segments lie below 2^40

A GPFIFO entry holds 40 bits of segment address: bits 31:2 in GET and bits 7:0 of the high word in GET_HI (clc56f.h). The NVIDIA kernel path maps memory that the host also maps at the same address on both sides, and keeps those addresses below 2^40, reserving the range in the process so that nothing else maps there. Device memory lies at addresses the process reserves for each allocation wherever it has room. Unified memory maps a range for a GPU only in the lower half of the CPU's addresses (uvm_gpu_can_address), which bounds where the kernel places a mapping without a hint.

x86_64 AddressSanitizer owns 2.25 GiB to 2 TiB

On x86_64 Linux, ASan's shadow gap covers 0x00008fff7000 to 0x02008fff6fff, about 2.25 GiB to 2 TiB (compiler-rt's asan_mapping.h). By default ASan maps the gap inaccessible, so a process under ASan can reserve nothing there. With protect_shadow_gap=0 ASan maps the gap's own shadow instead, from 2.3 to 258 GiB, and leaves the rest free. NVIDIA segment addresses must lie below 2^40, inside the gap: with its GPU addresses from 64 GiB no GPU opened under ASan. The kernel path reserves them from 384 GiB, and the sanitize profile in the root dune sets protect_shadow_gap=0.

PCI functions on Linux

These matter to a driver that takes a GPU from its kernel driver. A note that a run as root on the R9700's host confirmed says so.

Bus mastering outlives the process

A function unbound from its kernel driver and driven through sysfs keeps bus mastering on after the process exits. Its GPU can then still write system memory the process gave back, which by then belongs to another process or to the kernel. The PCI library clears the command register's bus master bit at release and in an exit handler (test/pci/test_function.ml, "a function taken physically stops mastering the bus when released"). Behind an IOMMU, closing VFIO's device file clears it (vfio_pci_core_disable). Read in source.

flock locks a function's sysfs config file

flock on a descriptor of /sys/bus/pci/devices/<bus>/config locks: kernfs gives its files no flock operation, so the kernel's generic one locks the inode, and kernfs keeps one inode per sysfs node. Every process that opens the file sees the lock. Any user may read config, so any user can hold it (lib/pci/sysfs.ml, locked). VFIO's group file admits one process at a time and needs no lock. Read in source.

A BAR resize refused for room may fit smaller

Writing a size to resourceN_resize fails with ENOSPC when the bridge window above has no room for it, and a smaller size may then fit. Any other refusal (EBUSY, a bridge that cannot move, no privilege) holds for every size, and the BAR keeps its own (lib/pci/sysfs.ml, resize). Read in source, and seen on the R9700: Rig_pci.Driver.detach, run as root, was refused every size from 32 GiB down to 1 GiB, its bridges' prefetchable window lying below 4 GiB, and took 512 MiB, where amdgpu had left 256 MiB. The kernel moves the device's other BARs as it resizes: the doorbell BAR went from 0xa0000000 to 0xb0000000.

A GPU without an IOMMU outlives its process

A GPU that masters the bus without an IOMMU writes physical pages, and keeps writing after its process dies: SIGKILL, the OOM killer and a crash run no exit function, so nothing turns bus mastering off, and Linux gives the process's pages to others immediately. The PCI library keeps a physical take's memory in huge pages of a file under /dev/hugepages, which keep their pages allocated until each GPU that reached them is reset (Rig_pci.Driver.reset, which attach runs), and refuses the process's own memory to such a GPU (Function.pin). A kernel booted by kexec skips the platform's reset and inherits a GPU still writing.

Allocated is not in place: the kernel moves pages, freeing the old physical page, and a GPU without an IOMMU keeps writing it. Read in Linux 6.12:

• Only an extra reference stops a move: migration freezes the page's count at what its mappings explain and gives up otherwise (__folio_migrate_mapping, migrate_huge_page_move_mapping in mm/migrate.c). Any get_user_pages/pin_user_pages reference does; FOLL_LONGTERM moves the page out of ZONE_MOVABLE and CMA first (mm/gup.c). VFIO pins so behind an IOMMU; its no-IOMMU mode has no DMA mapping call and pins nothing (vfio_noiommu_ioctl, drivers/vfio/container.c). No pin a process holds outlives it: io_uring's registered buffers go with the ring's file.
• tmpfs pages are movable (GFP_HIGHUSER_MOVABLE, inode_init_always). While the process lives, mlock keeps them off the evictable list, which compaction skips only with vm.compact_unevictable_allowed = 0 (isolate_migratepages_block, mm/compaction.c); alloc_contig_range (CMA, virtio-mem), memory offlining and soft offlining move them whatever the setting. At the process's death the unmapping munlocks them (__folio_remove_rmap, __munlock_folio), and ordinary compaction, kcompactd and proactive compaction included, moves them.
• hugetlbfs pages are skipped by compaction (PageHuge without cc->alloc_contig) and never swapped, but alloc_contig_range migrates them (isolate_or_dissolve_huge_page, mm/hugetlb.c), as do memory offlining and soft offlining; the pool is allocated movable (htlb_alloc_mask). 2 MiB pages migrate on x86_64 and arm64 (ARCH_ENABLE_HUGEPAGE_MIGRATION).

So without an IOMMU nothing keeps a page in place once its process has died: a physical take's memory stays exposed to migration. Small pages would be moved by routine compaction, which is why the library gives such a GPU huge pages only, sharing each among the allocations of its 2 MiB block of addresses. That narrows the exposure to CMA allocations, virtio-mem, and memory going offline or a failing page taken out of service, which an administrator or a driver starts. An IOMMU is the only complete answer.

Unbinding amdgpu leaves KFD refusing every process

On Linux 6.12, after amdgpu is unbound from a GPU and bound again, amdgpu serves the render node, but every open("/dev/kfd") fails with EINVAL until the amdgpu module reloads. KFD refuses to create a process while its module-wide counter kfd_locked is above zero (kfd_process.c:860, kfd_locked at kfd_device.c:46). The unbind's amdgpu_device_ip_fini_early calls amdgpu_amdkfd_suspend(adev, false) (amdgpu_device.c:3390), whose kgd2kfd_suspend raises the counter (kfd_device.c:1021); the device is then freed without the matching resume, and nothing lowers it again. The kernel frees KFD's device first only for a surprise removal (amdgpu_device.c:4756). Seen on the R9700 after Rig_pci.Driver.detach then Rig_pci.Driver.attach, on Linux 6.12.111; source read at v6.12.111. Reloading amdgpu (modprobe -r amdgpu, then modprobe amdgpu) clears it, so a session that detaches an AMD GPU reloads the module before KFD's users come back.

Round trips on these machines

Medians of loops in C (Objective-C on the M1 Max) over each vendor's interface: what each operation costs with no runtime above it. The vendor scenario of bench/compare runs the releases and copies through CUDA, HIP and Metal. The RM and KFD columns came from loops over the kernel drivers' own interfaces, which no program in the tree runs; compare's CUDA and HIP cases stand in for them on those GPUs.

On the R9700 the copy engine's 206 GB/s within GPU memory is 32% of its 640 GB/s, and one copy queue runs the two directions one after the other: 2 × 256 MiB, one each way, take 38.48 ms, the sum of each alone. Pipelined, a release costs 1.3 µs; the rest of its 15 µs round trip is, as a guess, the command processor waking on an idle queue and fetching the ring. The M1 Max has one memory, so its copies are the host's.

Measuring on these machines

• The RTX 5000 Ada boosts between 2550 and 2805 MHz from run to run and within one, and an unprivileged process cannot pin the clock. A GPU-bound row moves with it: 64 chained launches took 47.6 µs at 2805 MHz and 51.8 µs at 2550. Rows record the P-state and clocks.
• On the R9700 rows that wait on a copy-queue fence are bimodal between processes (see the TRAP note above): a round trip that switches queues takes 14.5 or 18.8-19.8 µs, a 16-byte copy 3.9 to 6.1 µs. One session ran every latency row, the vendor loops included, 1.6 times slower; its cause was not found.
• The R9700's own link reads 32 GT/s to its PCIe switch, but the root port above runs PCIe 3.0 x16, and host copies stop at its 15.75 GB/s. A link budget reads every hop up to the root port.
• On the M1 Max, launch timings drift 10-30% between processes. Variants compared for small differences run interleaved in one process.

edit this page ↗