# gizhw:gpu a GPU found at run time through its vendor's driver library, opened with gizlang:native (nothing is linked or required at build time). AMD GPUs run LLVM IR compiled in-process by ROCm's code object manager (comgr, no clang); NVIDIA GPUs run PTX the CUDA driver compiles; Apple GPUs run the MSL that the system compiler turns into a pipeline (gizhw:metal, unified memory: a copy to the device is a plain copy). A program uses one `Gpu` for device memory (`Buffer`, freed when its last copy is dropped), kernels compiled once per source and launches; gizml:tensor generates the kernels. ```gizlang import * as Gpu from "gizhw:gpu"; ``` ## Types ### `Vendor` ```gizlang type Vendor = { compiler: Option; metal: Option; sources: Option; } ``` What only one vendor has: HIP's comgr, which compiles; Metal's driver, whose calls are Objective-C messages. - `compiler`: `_gizlang_sdk_core.Option` - `metal`: `_gizlang_sdk_core.Option` - `sources`: `_gizlang_sdk_core.Option` ### `Sources` ```gizlang type Sources = { sys: System; folder: string; } ``` GIZLANG_GPU_SOURCES: the folder each compiled kernel's source is written to (to read its PTX, SASS or IR). - `sys`: `gizlang_system.System` - `folder`: `string` ### `KernelTime` ```gizlang type KernelTime = { key: string; entry: u64; blocks: i64; launches: i64; ns: i64; } ``` The time of one kernel's launches on `blocks` workgroups while profiling: `key` as compiled, `entry` its address. - `key`: `string` - `entry`: `u64` - `blocks`: `i64` - `launches`: `i64` - `ns`: `i64` ### `Launching` ```gizlang type Launching = { entry: u64; shared: i64; after: bool; } ``` A kernel's launch facts: `shared` bytes of dynamic workgroup memory, `after` its programmatic dependency (awaitGrid). - `entry`: `u64` - `shared`: `i64` - `after`: `bool` ### `Gpu` ```gizlang type Gpu ``` An opened GPU: `kind` is "hip" (AMD, ROCm), "cuda" (NVIDIA) or "metal" (Apple), `arch` its instruction set ("gfx1151", "sm_120", "apple9"), the features its kernels are made for. - `kind`: `string` - `arch`: `string` - `name`: `string` ### `Buffer` ```gizlang type Buffer ``` `bytes` of device memory (`count` f32s), freed when the last copy is dropped. - `count`: `i64` - `bytes`: `i64` ### `Selection` ```gizlang type Selection = | { kind: "auto"; } | { kind: "cpu"; } | { kind: "gpu"; } | { kind: "hip"; } | { kind: "cuda"; } | { kind: "metal"; } ``` Placement policy, distinct from whether a driver can be opened. - `auto` - `cpu` - `gpu` - `hip` - `cuda` - `metal` ### `GpuUsage` ```gizlang type GpuUsage = { taken: i64; used: i64; peak: i64; } ``` Device bytes of a Gpu: `taken` from the driver now (in use or pooled), `used` by live buffers now, and `peak` the most used since Gpu.resetPeak. A pooled block is not used: a peak of taken bytes counted the pool's idle blocks of other sizes and grew with the free memory (flux3-klein refused every job, 10/21). - `taken`: `i64` - `used`: `i64` - `peak`: `i64` ### `GpuError` ```gizlang type GpuError = | { kind: "memory"; requested: i64; room: i64; } | { kind: "driver"; what: string; code: i64; } ``` Why a device call failed: the card has no room for `requested` bytes (`room` it has: its free bytes less the reserve), or the driver gave the status `code` for the call `what`. A program that cannot go on without the card lets it fault (R004 at the call); one that can, catches it: `try { ... } catch (e) { ... }`. - `memory` - `driver` ### `DeviceMemory` ```gizlang type DeviceMemory = { free: i64; total: i64; } ``` Bytes of device memory: `free` for new buffers now (the pool of dropped buffers not counted), `total` of the GPU. - `free`: `i64` - `total`: `i64` ### `Pinned` ```gizlang type Pinned ``` Page-locked host memory of `bytes`: the GPU copies it by DMA at the bus's rate (Gpu.copyAhead). Freed when the last copy is dropped. - `bytes`: `i64` ### `Recording` ```gizlang type Recording ``` Kernels that Gpu.record launched once, launched again all together by Gpu.replay (CUDA and HIP graphs; Metal keeps the launches and encodes them again). Opaque: create and use it through this module's functions. ### `FileRead` ```gizlang type FileRead = { buffer: Buffer; at: i64; offset: i64; size: i64; } ``` One read of Gpu.readFiles: `size` bytes of the file from byte `offset` into `buffer` from byte `at`. - `buffer`: `M10.Buffer` - `at`: `i64` - `offset`: `i64` - `size`: `i64` ## Functions ### `parseSelection` ```gizlang function parseSelection(value: string): Selection throws string ``` The Selection a name gives: "auto" (or empty), "cpu", "gpu", "hip", "cuda" or "metal". Effects: fail. ### `open` ```gizlang function open(sys: System): Option ``` The GPU that `GIZLANG_DEVICE` selects (`auto` or unset: HIP, else CUDA, else Metal, else none; `cpu`: none). Tensors go there with Tensor.on, and each operation runs where its operands are. An invalid value, or a required GPU that does not open, faults rather than falling back to the CPU. Effects: fail, read, write, io. ### `openApple` ```gizlang function openApple(nat: Native, objc: Objc): Option ``` The Apple GPU that `objc` (gizhw:metal's Objc: the system's, or a test's) finds, opened as a Gpu. Effects: fail, read, write, io. ### `loadTunings` ```gizlang function loadTunings(sys: System, gpu: Gpu, folder: string): void ``` Takes this GPU model's tunings kept under `folder`, when there are any. Effects: fail, read, write, io. ### `saveTunings` ```gizlang function saveTunings(sys: System, gpu: Gpu, folder: string): bool throws string ``` Keeps what this GPU has measured under `folder` for the next process (loadTunings). Effects: fail, read, write, console, io. ## Methods ### `Buffer.address` ```gizlang function Buffer.address(self: Buffer): u64 ``` A kernel's device address for launch arguments. ### `Buffer.view` ```gizlang function Buffer.view(self: Buffer, at: i64, size: i64): Buffer ``` Bytes [at, at + size) of the buffer as a buffer of their own on the same memory (no copy); it keeps this one. Effects: fail. ### `Selection.required` ```gizlang function Selection.required(self: Selection): bool ``` Whether the selection fails without a GPU: `gpu`, `hip`, `cuda` and `metal` do; `auto` and `cpu` do not. ### `Gpu.others` ```gizlang function Gpu.others(self: Gpu): Vec ``` The driver's other GPUs, in its order (CUDA: each device CUDA_VISIBLE_DEVICES shows but this one; none elsewhere yet). Each opened is its own Gpu, current on this thread after it opens: Gpu.enter makes one current again. One of this GPU's model shares its tunings (Gpu.tune). Effects: fail, read, write, io. ### `Gpu.enter` ```gizlang function Gpu.enter(self: Gpu): void ``` Makes this GPU the one the thread's calls go to (CUDA: its context current; nothing elsewhere, one GPU only). Effects: fail, read, io. ### `GpuError.message` ```gizlang function GpuError.message(self: GpuError): string ``` The error as one line for a person. Effects: fail. ### `Gpu.room` ```gizlang function Gpu.room(self: Gpu): i64 throws GpuError ``` The device bytes a new allocation may take now: the free bytes less the card's reserve (what a caller that sizes its work by free memory must compare with: Gpu.memory's free includes the reserve). Effects: fail, read, write, io. ### `Gpu.taken` ```gizlang function Gpu.taken(self: Gpu): i64 ``` The device bytes this Gpu holds from the driver now: its buffers and its pool of dropped ones. Effects: fail, read. ### `Gpu.used` ```gizlang function Gpu.used(self: Gpu): i64 ``` The device bytes of this Gpu's live buffers now (not its pool). Effects: fail, read. ### `Gpu.peak` ```gizlang function Gpu.peak(self: Gpu): i64 ``` The most device bytes of live buffers since the last Gpu.resetPeak: what a piece of work needs, measured. Effects: fail, read. ### `Gpu.resetPeak` ```gizlang function Gpu.resetPeak(self: Gpu): void ``` Starts measuring Gpu.peak from the bytes used now. Effects: fail, read, write. ### `Gpu.fits` ```gizlang function Gpu.fits(self: Gpu, size: i64): bool throws GpuError ``` Whether `size` more bytes fit on the card now: its free memory, with the pool of dropped buffers given back, less the reserve kept for the driver and the card's other processes. Where they do not, the values no one reaches are collected first (collected). Effects: fail, read, write, io. ### `Gpu.buffer` ```gizlang function Gpu.buffer(self: Gpu, count: i64): Buffer throws GpuError ``` `count` f32s of device memory, contents undefined. Effects: fail, read, write, io. ### `Gpu.bytes` ```gizlang function Gpu.bytes(self: Gpu, size0: i64): Buffer throws GpuError ``` `size` bytes of device memory, contents undefined (`count` is size / 4). Effects: fail, read, write, io. ### `Gpu.trim` ```gizlang function Gpu.trim(self: Gpu): void ``` Gives the blocks of dropped buffers back to the GPU (Gpu.memory then counts them free), after the work before. Effects: fail, read, write, io. ### `Gpu.memory` ```gizlang function Gpu.memory(self: Gpu): DeviceMemory throws GpuError ``` The free and total bytes of the GPU's memory now (DeviceMemory). Effects: fail, read, write, io. ### `Gpu.upload` ```gizlang function Gpu.upload(self: Gpu, data: Vec): Buffer throws GpuError ``` A device copy of `data`. Effects: fail, read, write, io. ### `Gpu.download` ```gizlang function Gpu.download(self: Gpu, buffer: Buffer, count: i64 = -1, at: i64 = 0): Vec ``` `count` f32s of the buffer from f32 `at` (to its end when -1), after every kernel launched before. Effects: fail, read, write, io. ### `Pinned.address` ```gizlang function Pinned.address(self: Pinned): u64 ``` The device-visible host address of the memory. ### `Gpu.pinned` ```gizlang function Gpu.pinned(self: Gpu, size: i64): Pinned ``` `size` bytes of page-locked host memory, contents undefined. Effects: fail, read, io. ### `Gpu.readInto` ```gizlang function Gpu.readInto(self: Gpu, buffer: Buffer, p: Pinned, at: i64): void ``` Copies `buffer` to `p` from byte `at`, after every kernel launched before. Effects: fail, read, write, io. ### `Gpu.copyAhead` ```gizlang function Gpu.copyAhead(self: Gpu, target: Buffer, p: Pinned, at: i64): void ``` Starts a copy of `target.bytes` bytes of `p` from byte `at` to `target` once the work launched before is done (so the kernels before may still read `target`). The copy runs beside the kernels launched after it: the first that reads `target` comes after Gpu.awaitCopies. Effects: fail, read, write, io. ### `Gpu.awaitCopies` ```gizlang function Gpu.awaitCopies(self: Gpu): void ``` The kernels launched after this run once every copy Gpu.copyAhead started before is done. Effects: fail, read, io. ### `Gpu.now` ```gizlang function Gpu.now(self: Gpu): i64 ``` Nanoseconds of a monotonic clock, for timing work on the GPU. Effects: fail, io. ### `Gpu.stream` ```gizlang function Gpu.stream(self: Gpu): u64 ``` The stream every kernel launches on (CUDA, HIP), for a vendor library whose work must take its place in order with them; 0 on Metal. ### `Gpu.sync` ```gizlang function Gpu.sync(self: Gpu): void throws GpuError ``` Waits for every launched kernel (a fault in one fails here). Effects: fail, read, write, io. ### `Gpu.record` ```gizlang function Gpu.record(self: Gpu, work: F): Recording ``` Records the kernels that `work` launches, after the work before it is done. Each replay runs them again on the same buffers: the buffers made in `work` stay the recording's until release, and the ones made before that it reads must live as long. `work` must not wait for the GPU: no download, sync, profile, or tune that is not known yet (run `work` once before). An upload in `work` is copied once, while recording: its values are the recording's. ### `Gpu.replay` ```gizlang function Gpu.replay(self: Gpu, r: Recording): void ``` Runs the recorded kernels again, after the work before (returns at once, as a launch does). Effects: fail, read, write, io. ### `Gpu.release` ```gizlang function Gpu.release(self: Gpu, r: Recording): void ``` Frees a recording after the work before. Its buffers that were dropped go back to the pool, and the others do when they are dropped. Effects: fail, read, write, io. ### `Gpu.write` ```gizlang function Gpu.write(self: Gpu, buffer: Buffer, data: Vec, at: i64 = 0): void ``` Copies `data` into the f32s of `buffer` from f32 `at`, after the work before. Effects: fail, read, write, io. ### `Gpu.kernel` ```gizlang function Gpu.kernel(self: Gpu, key: string, name: string, source: F): u64 ``` The kernel `name` of the program `source()` makes (LLVM IR for "hip", PTX for "cuda"), made and compiled the first time this Gpu sees `key`: a key names one program, and costs less to compare than the program to make. ### `Gpu.launch` ```gizlang function Gpu.launch(self: Gpu, kernel: u64, blocks: i64, threads: i64, arguments: Vec): void throws GpuError ``` Launches KERNEL on BLOCKS blocks of THREADS threads with ARGUMENTS (the addresses of the argument values). Effects: fail, read, write, io. ### `Gpu.tensorMap` ```gizlang function Gpu.tensorMap(self: Gpu, b: Buffer, rowBytes: i64, rows: i64, pitch: i64, boxBytes: i64, boxRows: i64, swizzle: i64): Buffer throws GpuError ``` The copy engine's map (NVIDIA, Kernel.hasTma) of `rows` rows of `rowBytes` bytes, `pitch` bytes apart, in `b`: Kernel.tensorLoad copies boxes of `boxRows` rows of `boxBytes` bytes (a multiple of 16, at most 256) from it, with zeros past its edges. `swizzle` 0, 32, 64 or 128 permutes the 16-byte chunks of each box row in that span (boxBytes at most `swizzle`), as the chunk XOR of ldmatrix reads wants. The map is in a buffer of its own; its address is the kernel's argument. Effects: fail, read, write, io. ### `Gpu.startProfile` ```gizlang function Gpu.startProfile(self: Gpu): void ``` Starts profiling (and clears the times): each launch from now on waits for the work before it and for itself, so the times are the kernels' own, and the program runs slower. Effects: fail, write. ### `Gpu.profiling` ```gizlang function Gpu.profiling(self: Gpu): bool ``` Whether startProfile is on (a recording cannot be made then). Effects: fail, read. ### `Gpu.profile` ```gizlang function Gpu.profile(self: Gpu): Vec ``` The time of each kernel launched since startProfile, the slowest first; profiling stops. Effects: fail, read, write. ### `Gpu.slab` ```gizlang function Gpu.slab(self: Gpu, sizes: Vec): Vec ``` Buffers of `sizes` bytes in one block of device memory, each from a multiple of 256 bytes. The driver gives each block in pages of 2 MiB: many large buffers of their own lose up to a page each (the 1146 weights of q27-4 lost 475 MiB). Effects: fail, read, write, io. ### `Gpu.uploadBytes` ```gizlang function Gpu.uploadBytes(self: Gpu, data: bytes): Buffer ``` A device copy of `data`'s bytes. Effects: fail, read, write, io. ### `Gpu.uploadFile` ```gizlang function Gpu.uploadFile(self: Gpu, fd: i64, offset: i64, size: i64): Result ``` A device buffer of `size` bytes of the open file `fd` from byte `offset`: pread straight into page-locked memory a piece at a time, each piece copied by DMA while the next one reads (no other copy on the host). Effects: fail, read, write, io. ### `Gpu.readFiles` ```gizlang function Gpu.readFiles(self: Gpu, fd: i64, reads: Vec): Result ``` Reads each of `reads` from the open file `fd`, after every kernel launched before. On unified memory (Metal) the reads go straight into the buffers, else through the staging memory. Effects: fail, read, write, io. ### `Gpu.downloadBytes` ```gizlang function Gpu.downloadBytes(self: Gpu, buffer: Buffer, size: i64 = -1, at: i64 = 0): bytes ``` `size` bytes of the buffer from byte `at` (to its end when -1), after every kernel launched before. Effects: fail, read, write, io. ### `Gpu.writeBytes` ```gizlang function Gpu.writeBytes(self: Gpu, buffer: Buffer, data: bytes, at: i64 = 0): void ``` Copies `data` into the buffer from byte `at`, after the work before. Effects: fail, read, write, io. ### `Gpu.compile` ```gizlang function Gpu.compile(self: Gpu, key: string, make: F): u64 ``` The Kernel `make()` builds, compiled for this GPU the first time it sees `key` (a key names one kernel, and costs less to compare than the kernel to build). ### `Gpu.slots` ```gizlang function Gpu.slots(self: Gpu, kernel: u64, threads: i64): i64 ``` The workgroups of `threads` threads of `kernel` that run at once on the whole GPU (one wave), or 0 where the driver does not say (AMD). Effects: fail, read, io. ### `Gpu.run` ```gizlang function Gpu.run(self: Gpu, kernel: u64, blocks: i64, threads: i64, buffers: Vec, scalars: Vec): void ``` Runs `kernel` on `blocks` workgroups of `threads` threads: its buffer parameters are `buffers` in order, then its scalars. Effects: fail, read, write, io. ### `Gpu.copyRows` ```gizlang function Gpu.copyRows(self: Gpu, dtype: i64, dst: Buffer, src: Buffer, outer: i64, len: i64, srcRow: i64, srcOff: i64, dstRow: i64, dstOff: i64): void ``` dst[o * dstRow + dstOff + j] = src[o * srcRow + srcOff + j] for o < outer, j < len (`dtype` elements). Effects: fail, read, write, console, io. ### `Gpu.zero` ```gizlang function Gpu.zero(self: Gpu, b: Buffer, first: i64, count: i64): void ``` Zeroes `count` words of 4 bytes of b from word `first`. Effects: fail, read, write, console, io. ### `Gpu.tune` ```gizlang function Gpu.tune(self: Gpu, key: string, choices: Vec, work: F): i64 ``` The fastest candidate, measured after compilation and cached on this opened GPU. `key` names one operation, shape, layout and dtype; candidates must give the same result. `work` may overwrite its output and runs once per choice to warm up, then three timed times per choice, the choices in turn: work of another process on the GPU then slows each choice alike. ### `Gpu.tunings` ```gizlang function Gpu.tunings(self: Gpu): string ``` The choices tune measured on this GPU, one "CHOICE KEY" line each: a later process takes them (Gpu.learn) instead of measuring again. Effects: fail, read. ### `Gpu.learn` ```gizlang function Gpu.learn(self: Gpu, text: string): void ``` Takes choices of Gpu.tunings' text (lines it cannot read are skipped): tune measures only keys it has not seen. A key names an operation and its shape, not the GPU, so keep a text for each GPU model (tuningsFile). Effects: fail, read, write. ### `Gpu.tuningsFile` ```gizlang function Gpu.tuningsFile(self: Gpu, folder: string): string ``` The file under `folder` that keeps this GPU model's tunings. Effects: fail. ### `Gpu.force` ```gizlang function Gpu.force(self: Gpu, position: i64): void ``` From now on each tune picks its candidate at `position` (the last where it has fewer), measures nothing and keeps nothing; -1 tunes again. A test runs each candidate with it. Effects: fail, write. ### `Gpu.useVendorBlas` ```gizlang function Gpu.useVendorBlas(self: Gpu, on: bool): void ``` On: from now on a product without an epilogue may also run as cuBLAS's (cublasGemmEx, as PyTorch's matmul), where the dynamic loader finds libcublas.so.12 and the tuner measures it faster. Its order of additions is NVIDIA's, so its bits can differ from the CPU's. Off (the default): every product has the CPU's bits. Effects: fail, read, write, io. ### `Gpu.vendorBlas` ```gizlang function Gpu.vendorBlas(self: Gpu): bool ``` Whether products may run as cuBLAS's: useVendorBlas(true) was called and the library loaded. Effects: fail, read. ### `Gpu.same` ```gizlang function Gpu.same(self: Gpu, other: Gpu): bool ``` Whether two Gpus are one opened GPU (copies of it). ### `Gpu.vendorMatmul` ```gizlang function Gpu.vendorMatmul(self: Gpu, c: Buffer, a: Buffer, b: Buffer, input: string, output: string, m: i64, k: i64, n: i64, transposeA: bool, transposeB: bool): void ``` C = op(A) op(B) by cuBLAS (Gpu.vendorBlas must be true): row-major m x n C of `output`, A m x k (stored k x m if transposeA) and B k x n (stored n x k if transposeB) of `input` ("f32", "f16" or "bf16"), added in f32. Effects: fail, read, io. ## Constants ### `DEVICE_ENV` ```gizlang const DEVICE_ENV = "GIZLANG_DEVICE" ``` The name of the environment variable that chooses the device.