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.

import * as Gpu from "gizhw:gpu";

Types

Vendor

type Vendor = { compiler: Option<Library>; metal: Option<Metal>; sources: Option<Sources>; }

What only one vendor has: HIP's comgr, which compiles; Metal's driver, whose calls are Objective-C messages.

• compiler: _gizlang_sdk_core.Option<gizlang_native.Library>
• metal: _gizlang_sdk_core.Option<M13.Metal>
• sources: _gizlang_sdk_core.Option<M10.Sources>

Sources

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

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

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

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

type Buffer

bytes of device memory (count f32s), freed when the last copy is dropped.

• count: i64
• bytes: i64

Selection

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

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

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

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

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

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

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

function parseSelection(value: string): Selection throws string

The Selection a name gives: "auto" (or empty), "cpu", "gpu", "hip", "cuda" or "metal".

Effects: fail.

open

function open(sys: System): Option<Gpu>

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

function openApple(nat: Native, objc: Objc): Option<Gpu>

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

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

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

function Buffer.address(self: Buffer): u64

A kernel's device address for launch arguments.

Buffer.view

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

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

function Gpu.others(self: Gpu): Vec<Gpu>

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

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

function GpuError.message(self: GpuError): string

The error as one line for a person.

Effects: fail.

Gpu.room

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

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

function Gpu.used(self: Gpu): i64

The device bytes of this Gpu's live buffers now (not its pool).

Effects: fail, read.

Gpu.peak

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

function Gpu.resetPeak(self: Gpu): void

Starts measuring Gpu.peak from the bytes used now.

Effects: fail, read, write.

Gpu.fits

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

function Gpu.buffer(self: Gpu, count: i64): Buffer throws GpuError

count f32s of device memory, contents undefined.

Effects: fail, read, write, io.

Gpu.bytes

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

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

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

function Gpu.upload(self: Gpu, data: Vec<f32>): Buffer throws GpuError

A device copy of data.

Effects: fail, read, write, io.

Gpu.download

function Gpu.download(self: Gpu, buffer: Buffer, count: i64 = -1, at: i64 = 0): Vec<f32>

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

function Pinned.address(self: Pinned): u64

The device-visible host address of the memory.

Gpu.pinned

function Gpu.pinned(self: Gpu, size: i64): Pinned

size bytes of page-locked host memory, contents undefined.

Effects: fail, read, io.

Gpu.readInto

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

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

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

function Gpu.now(self: Gpu): i64

Nanoseconds of a monotonic clock, for timing work on the GPU.

Effects: fail, io.

Gpu.stream

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

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

function Gpu.record<F>(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

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

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

function Gpu.write(self: Gpu, buffer: Buffer, data: Vec<f32>, 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

function Gpu.kernel<F>(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

function Gpu.launch(self: Gpu, kernel: u64, blocks: i64, threads: i64, arguments: Vec<u64>): 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

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

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

function Gpu.profiling(self: Gpu): bool

Whether startProfile is on (a recording cannot be made then).

Effects: fail, read.

Gpu.profile

function Gpu.profile(self: Gpu): Vec<KernelTime>

The time of each kernel launched since startProfile, the slowest first; profiling stops.

Effects: fail, read, write.

Gpu.slab

function Gpu.slab(self: Gpu, sizes: Vec<i64>): Vec<Buffer>

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

function Gpu.uploadBytes(self: Gpu, data: bytes): Buffer

A device copy of data's bytes.

Effects: fail, read, write, io.

Gpu.uploadFile

function Gpu.uploadFile(self: Gpu, fd: i64, offset: i64, size: i64): Result<Buffer, string>

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

function Gpu.readFiles(self: Gpu, fd: i64, reads: Vec<FileRead>): Result<bool, string>

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

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

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

function Gpu.compile<F>(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

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

function Gpu.run(self: Gpu, kernel: u64, blocks: i64, threads: i64, buffers: Vec<Buffer>, scalars: Vec<u64>): 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

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

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

function Gpu.tune<F>(self: Gpu, key: string, choices: Vec<i64>, 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

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

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

function Gpu.tuningsFile(self: Gpu, folder: string): string

The file under folder that keeps this GPU model's tunings.

Effects: fail.

Gpu.force

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

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

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

function Gpu.same(self: Gpu, other: Gpu): bool

Whether two Gpus are one opened GPU (copies of it).

Gpu.vendorMatmul

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

const DEVICE_ENV = "GIZLANG_DEVICE"

The name of the environment variable that chooses the device.