# Compute: all of the hardware from one program

This page is the design of the hardware-independent compute engine (SIMD, tensors, GPU, AI/ML) and the measured state
of each part.

The aim: Gizlang above Mojo and Julia on performance and modernity, from first principles,
with few orthogonal concepts. There are three pillars. The third is the differentiator.

1. **A compute engine independent of hardware**, at the level of CubeCL and IREE: one source,
   automatic SIMD, specialized GPU code, autotune.
2. **Numerics and AI/ML** at the level of JAX and Julia: automatic differentiation, tensors,
   scientific computing, training and inference. This replaces Python for that work.
3. **A live semantic graph**, the own pillar of Gizlang: functions, state, data and GPU kernels
   updated by the impact of a change, and swapped safely while they run. An advantage counts only
   when it is measured.

## References

Studied in this order, for what they settled and what they left open. That gap is the work of Gizlang.

| Reference | Settled | Left open |
| --- | --- | --- |
| **CubeCL** | One Rust source to CUDA, ROCm, Vulkan/WGPU and CPU; comptime specialization; autotune | A host language and kernels are two worlds; no live replacement |
| **IREE** | One MLIR program for every device; a HAL; ahead of time | Compile then deploy; no incremental change |
| **Mojo** | SIMD types in the language; parametric widths; MLIR; compiler open under Apache 2.0 (2026-08) | MAX binaries under the Modular Community License; no graph of meaning |
| **JuliaGPU** | One generic function to CPU and GPU by specialization | Compile latency; invalidation by method tables, not by meaning |
| **JAX/XLA** | Functional transforms `grad`/`vmap`/`jit`; fused kernels | Trace in Python; compile again by shape; no swap of a running program |

## Design

- **One semantic IR.** The checked graph is the program for every device.
  - The optimizer works on one IR at tensor level: loops, maps, reductions, layouts.
  - Each backend only generates code: LLVM for CPUs (x86 AVX2/AVX-512/AMX, ARM NEON/SVE, Wasm
    SIMD) and NVPTX/AMDGPU; Metal Shading Language for Apple GPUs; SPIR-V for Vulkan and WGSL from their own generators.
- **Automatic placement that you can override.**
  - The runtime picks CPU lanes, cores, GPU or tensor cores by a cost model,
    `T_compile + T_transfer + T_launch + T_compute + T_sync`, fitted from measurements on the machine.
  - `on: Device.cpu | Device.gpu(0)` and option flags override it.
  - A kernel runs on a device only when it is proven executable there (types, memory, features).
    Otherwise it runs on the CPU and says why: a diagnostic, never a silent change of result.
- **Features, not generations.** Detected instruction features choose the code paths (`gizlang:machine`,
  capped by `GIZLANG_CPU`; CUDA compute capability as features). Kernel units compile for the host level
  (`HASH-TOOLCHAIN-v4.so`), and a process never loads a unit above its level.
  - sm_120 (RTX 50) tensor cores are `mma.sync`. `tcgen05` is sm_100 only.
  - AVX-512 VBMI decides 64-lane byte lookups.
- **One owner for each hardware fact** (package `sdk-hw`, `hw/`; `sdk/core` never imports it).
  - `gizlang:machine`: what does not change while the process runs: CPU features and level, the CPUs it may use, caches,
    memory nodes. Runtime code reads the same facts in C (`runtime/machine.c`).
  - `gizhw:device`: the devices a program opens and every reading that changes: sysfs classes, hwmon sensors,
    RAPL energy, and the GPU boards' readings and controls (`GpuBoards`, through NVML).
  - `gizhw:gpu`: compute only (memory, kernels, streams). Management does not need a compute context.
  - Memory placement (NUMA nodes, huge pages, limits) is `gizlang:memory`'s, and io_uring rings are `runtime/uring.c`'s.
    A database or library measures its own costs from these facts and does not detect them again.
- **One `Tensor` over one storage layer** (host memory, device buffers, mapped files).
  - Storage, compute and accumulate precision are separate: f32, f16, bf16, FP8, and FP4 storage with f32 accumulate.
  - Reduction orders (`Reduction`):

    | `Reduction` | Meaning |
    | --- | --- |
    | `strict` | Order reproducible bit for bit |
    | `deterministic` | Same result on each machine |
    | `fast` | Any order |
- **Fixed and scalable SIMD.** Types such as `u8x32` are fixed widths. `Simd<T>` is the natural
  width of the machine, and scalable lengths such as SVE and RVV map to it.
- **Specialized patterns** (matmul, convolution, attention, scan) get their own tiled kernels and
  autotuned tiles. Vendor libraries (cuBLAS, cuDNN) are optional faster paths, loaded at run time, never required.
- **Kernel profile.** `gpu.startProfile()` then `gpu.profile()` gives the GPU time of each kernel and grid size
  (`KernelTime`: name, blocks, launches, ns), slowest first. Each launch then waits for the GPU, so use it to find
  the cost, not to time a whole step.
- **Recordings (CUDA and HIP graphs).** Every launch goes to one created stream. `gpu.record(work)` records the
  kernels `work` launches; `gpu.replay(r)` launches them all again in one driver call, on the same buffers.
  `Tensor.assign(x)` writes new values into a tensor's memory, so a recorded training step updates its weights in
  place, and new input goes in the same way. `Tensor.total` and `Tensor.average` keep a sum on the GPU (Graph.sum and
  Graph.mean use them), and so does Graph.crossEntropy (its loss is `T.crossEntropyRows(z, labels).average`), so a
  step does not wait for the GPU. An upload in `work` (such as the labels) is copied once, while recording. `work` must not wait for the GPU (no download or
  untuned product): run it once first. MLP 64 x 256-512-512-256, RTX 5090 idle: a BF16 training step takes 0.29 eager
  and 0.0451 ms replayed (PyTorch 0.634 eager, 0.0533 with a CUDA graph); the forward pass 0.051 and 0.0144 ms
  (PyTorch 0.080 and 0.0164). F32: 0.076 ms replayed (PyTorch graph 0.069). The products of this MLP have 4 to 8
  tiles of C: one workgroup each would wait on its loop over k, so they split k (split-K, matmulGpu): on the matrix
  cores a product of at most 32 tiles of 64 x 64 and k a multiple of 64 runs as 2, 4 or 8 slices of k (the most in 128
  workgroups) into F32, and one pass adds the slices in order and runs the chain. The shape alone sets the slices, so
  the bits are the same in each run. F32 products keep the CPU's order of sums, so they do not split. A strict sum adds blocks of 512 in f64 and joins
  the block sums pairwise (a tree of fixed shape), so one workgroup adds a block in 9 levels; the CPU adds in the
  same order.
- **Copies to and from the device** go through one page-locked host block (cuMemAllocHost, hipHostMalloc) that the
  GPU reads and writes by DMA. RTX 5090, 64 MiB: upload 18.5 GB/s (was 12.9 through pageable memory), download
  5.5 GB/s into a new Vec (was 2.0): the rest is the copy from that block, and the first writes of the new Vec, which
  runtime/memory.c maps fresh past 8 MiB with transparent huge pages (2.6 GB/s without them). PyTorch uploads 29.5
  GB/s from pageable memory, and downloads 3.3 GB/s into a new tensor (`.cpu()`) and 53.4 into a pinned one.
- **cuBLAS as a choice (CUDA).** `gpu.useVendorBlas(true)` adds cuBLAS's `cublasGemmEx` (the call PyTorch's matmul makes)
  to the tuner's choices for each product, where the dynamic loader finds libcublas.so.12. A product with an epilogue
  is then measured both fused and as cuBLAS's product in f32 followed by one pass of the chain. Its order of additions
  is NVIDIA's, so the bits can differ from the CPU's; it is off by default. RTX 5090: F32 8192 16.7 ms (own kernel
  21.1, PyTorch 16.6), the F32 MLP step 9.89 ms (own 11.6, PyTorch 9.60).
- **Matmul epilogue.** `T.lazyMatmul(a, ta, b, tb)` is a `Lazy` product. Elementwise operations on it (bias,
  activation, residual, mask, scale, cast) run in the matmul's store, with their own inputs; `eval` writes C once.
  `gizml:grad` builds forward values and gradients this way: an activation and its gradient mask go into the
  product's store, and SGD without momentum writes `w - rate * g` in one pass.
- **Differentiation and randomness as transforms.** Forward and reverse mode over the IR, custom
  gradients, `scan` for loops, checkpoints (compute again, not store), and functional RNG (a key
  split for each use, so results are reproducible across devices).
- **The live graph.** A compiled kernel is cached by its semantic hash, shapes, dtypes, hardware
  features and options. An edit compiles again only the kernels whose meaning changed, and swaps
  them between launches, after its data and state migrate (`giz live`).
- **No Python** anywhere in the compiler, SDK or runtime. We buy no GPU. The test devices are the
  Radeon 8060S of the development machine, the RTX 5090 host, and the CPU Vulkan lavapipe.

## Tile programs

`gizml:tile` (kernels more modern than Triton and Mojo, on every device, at top speed)
is how a person writes a kernel.

- **What the user says:** what one output block is, in ordinary Gizlang over tiles, and the block sizes. Nothing
  else: no thread ids, index arithmetic, masks, workgroup memory, barriers or fragment layouts.
- **One IR, every device.** The body is traced once into the tile IR. Host tensors run it on the CPU's cores; a GPU's
  run one `gizhw:kernel` program made from it (PTX, AMDGPU LLVM, MSL), cached by the IR and the storage types. A new
  device is a new lowering of the same IR.
- **Inferred layouts.** Each value gets a layout (Uniform, Blocked, matrix-core fragments, or a Slice of one) from
  its producer; free values (loads, iota, math on them) are made directly in the layout their consumer needs; other
  meetings go through workgroup memory. The epilogue of a product (bias, activation, cast, store) runs in the
  fragments' layout with no data movement.
- **Operands and fragments.** `dot` operands read ahead in blocks of 16 bytes (`cp.async` on NVIDIA sm_80+, `Program.stages(n)`
  to choose 1..4, 2 by default); reductions directly on matrix-core fragments (a row maximum of a product never leaves
  its registers); operators on tiles and indices (`q.dot(kt) * scale` after `const scale = q.k(s)`); `Tile.ir(...)`
  prints the IR with each value's layout; `Tile.tune` choices kept per device (`Gpu.tuningsFile`, `loadTunings`,
  `saveTunings`); a tile program with its gradient programs is one `gizml:grad` operation (`Graph.custom`).
  CPU, AMD gfx1151 and NVIDIA sm_89 (cp.async) give the goldens.
- **Not supported in tile programs:** FP8/FP4 matrix cores, `wgmma`/`tcgen05`, Vulkan, swizzled workgroup memory (the `+8` row
  padding already avoids bank conflicts). Each changes one lowering function, never a user's program.

## Parts

Each part is measured in `bench/compiler/perf`: Gizlang against C, or against the reference library on the
same work, with compile time beside it. The rule: within 1.3x of C, or the C stays.

### CPU

Done 2026-10-10:

- Programs are built for any x86-64 or ARM64 CPU; optimized x86-64 programs also carry v3 and v4 code and run the highest level the CPU has.
- `push` vectorizes.
- base64 by index within 1.0-1.2x of C.
- SIMD types and builtins in the language ([the language reference](../reference/language.md#simd)).
- `gizlang:native`: dynamic libraries, raw memory, a kernel on every core, owned native resources.

Done 2026-10-11:

- UTF-8 validation and search in Gizlang SIMD: `utf8Valid` of `gizlang:runtime` (which
  `string(bytes)` calls), and `bytes.indexOf` and `indexOfAny`. At 1 MB, validation is 2.7x (ASCII)
  and 9x (mixed) faster than the C that it replaced. Search is 0.9-2.3x.
- Base64 stays C: in Gizlang, encode takes 2.0x and decode 1.8x the time of the AVX2 C. Integer
  text also stays C: `string()` validates what it gets, and a new buffer starts zeroed.
- On x86-64, each function that uses vector types (the tensor kernels) also has copies for x86-64-v3 and v4, and
  its first call takes the best one (tools/cli/levels.giz, runtime/levels.c). Other code is baseline only. 1024^3
  matmul with copies of the whole program (2026-10-12): 140 ms on the baseline, 2.4 ms chosen, 1.8 ms native.
- A `parallel` whose body took less time than the wake of the threads (20 us) runs on its caller
  (runtime/parallel.c keeps the last cost of each body). 20000 small sections: 300-330 ms to 9-11 ms.
  Large sections do not change. `GIZLANG_THREADS=N` limits the threads.

### Tensors

On the CPU: `gizml:tensor` ([the language reference](../reference/language.md)).

- f32 tensors and NumPy broadcasts.
- `lazy` chains fused into one blocked pass on every core.
- `exp`, `log`, `tanh` and `sigmoid` within 1-3 ulp.
- Sums under three policies, softmax and transpose.
- matmul by packed 6x32 FMA tiles on every core: 1024^3 in approximately 3 ms, 0.7 TFLOPS on 16 Zen 5 cores.
- Each core writes its part of a result in place (`Task.fill`, 2026-10-11).

Against clang -O3 -march=native C on one core (4M floats, 4096x1024; under load):

| Operation | Gizlang | C |
| --- | --- | --- |
| exp | 0.5-2.0 ms | 5.6 ms |
| sum | 0.1-0.2 ms | 1.9 ms |
| softmax | 4.0-4.6 ms | 13.8 ms |
| `bench/compiler/perf/tensor` | 0.11-0.14 s | 0.34 s |
| fused relu(a*b+c) | 0.6-1.9 ms (2.0 ms before, joined by one serial copy) | 1.2-1.6 ms |
| transpose (32-column bands of 8x8 register blocks on every core) | 0.4-0.7 ms | 4.4-7.0 ms |

### GPU

Done 2026-10-11 for AMD and NVIDIA.

- `gizhw:gpu` opens the driver at run time: ROCm HIP with LLVM IR compiled in process by comgr
  (no clang), and CUDA with PTX.
- `gizml:tensor` writes each fused chain one time, in a small register IR printed as LLVM IR or
  PTX, with the polynomials and order of the CPU. Thus GPU results are the bits of the CPU
  (tests/hw/gpu, edge values included; PTX checked on the RTX 5090).
- Placement is explicit (`Tensor.on`, `Tensor.host`): each operation runs where its operands are. `GIZLANG_DEVICE`
  selects the GPU that `open` gives.
- F32 matmul in FMA tiles through workgroup memory: 32x32, 64x64 and 128x128, and where copies run in the background
  (CUDA) also 128x256 and 256x128 (256 threads of 128 sums each, as CUTLASS's SIMT kernel 256x128). The tuner picks
  one for each shape, and the 32x32 tiles (64 threads) give a small C enough workgroups.
- **The last wave in parts.** `gpu.slots(kernel, threads)` is the number of workgroups of a kernel that run at once
  (the driver's occupancy times the multiprocessors; 0 on HIP). When a product's last wave fills at most half of the
  slots, matmulGpu runs the tiles of that wave as tiles of the smallest plan (64 x 64 on the matrix cores, 32 x 32 for
  FMA), on more workgroups. Each sum is the same, in the same order, so the bits do not change (tests/hw/gpu). BF16
  8192 on the RTX 5090: 5.29 -> 5.07 ms.
- Sums, row sums (`sumTo`) and softmax on the device (2026-10-11) keep the chains of the CPU: one thread for
  each chain, one thread for each row that joins them in f64. Rows of 2048 or more go in tiles of
  8192 elements through workgroup memory, with the loads of the next tile in flight. The bits are
  those of the CPU on both GPUs (strict, fast, inf/NaN/denormal rows).

| Measurement | Radeon 8060S (gfx1151, the GPU of the machine) | RTX 5090 | CPU |
| --- | --- | --- | --- |
| matmul | 4.0-4.3 TFLOPS at 1024-4096 (rocBLAS 3.0-3.2) | F32 65 TFLOPS at 4096 and 8192 (cuBLAS 70-71, 2026-10-17) | 1.0-1.2 TFLOPS |
| exp-tanh-mul chain, 4M elements | 0.18-0.22 ms on the device | — | 0.9-2.3 ms |
| softmax 4096x1024 | 0.27 ms | 0.061 ms | 4.0-6.4 ms |
| row sums (`sumTo`) | 0.02-0.08 ms | — | 0.3-0.4 ms |
| One row of 4M elements (stays in one workgroup) | 5.2 ms | 2.8 ms | 3.9-6.4 ms |

Such long rows belong on the CPU.

### AI/ML

Done on the CPU (2026-10-11; the encoder and the tokenizer 2026-10-16):

- `gizml:grad`: reverse mode in one walk of a tape, forward mode by tangents, custom gradients,
  checkpoints, constants without gradients.
- Philox keys in `gizml:tensor` (`key`, `split`, `uniform`, `normal`; the known answer of Random123).
- `gizml:nn`: layers, MSE and cross-entropy, SGD with momentum, Adam and AdamW in fused passes,
  safetensors F32/F16/BF16/F64 and integer tensors; the BERT encoder and sentence-transformers pooling
  (`bert`, `encoderOf`, `Encoder.embed`).
- `gizml:tokenizer`: the ids of a HuggingFace `tokenizer.json` with a WordPiece model (BERT normalizer and
  pre-tokenizer, template post-processor, added tokens, truncation), and the langchain text splitters.
- `gizlang:hub`: the files of Hugging Face Hub repositories in a local folder, downloaded once.
- `gizml:onnx`: ONNX files of linear, tree, PCA and k-means models, and models of the ai.onnx.ml linear and
  tree operators.
- Gradients agree with central differences to 1e-4, and forward mode agrees with reverse to 1e-7
  (tests/ml/grad). AdamW fits two spirals to 99.75% (tests/ml/nn).

Against libtorch 2.8 C++ (no Python) on the same 16 CPUs of the RTX 5090 host, a 784-512-512-10 ReLU MLP at batch 256 in f32:

| Work | Gizlang | libtorch |
| --- | --- | --- |
| Training step (forward, backward, Adam) | 2.9-3.2 ms | 3.9-4.2 ms |
| Inference | 0.68-0.78 ms | 0.75-0.83 ms |
| 256x784x512 matmul | 0.30 ms | 0.32 ms |

Both take 2.4 ms for the matmul on one core. The changes that got there:

- threads as many as the CPU quota of the container (it saw 512 CPUs, and ran 32 threads for each CPU);
- a pool that spins 100 us before it sleeps;
- matmul that reads transposes in place, with B packed one time.

### NPU access

`gizhw:npu` has one shape for every NPU: `devices(sys)` lists the accel nodes a backend serves, and `info(sys, npu)`
reads the driver's name and version (`gizhw:drm`, shared by all NPUs) and its firmware version (the backend's own query).
`supports(driver)` gates every open, so an unknown driver is a typed error. Runs go through the same module: `open(sys)`
gives a `Session` (buffers, `load` of a kernel, command buffers), and `submit`, `finish` and `run` start runs and wait
for them; a kernel's firmware image and control code are in its backend's formats. A new vendor is one backend file in
`sdk/hw` (`amdxdna.giz` is the AMD XDNA one) and one case in each function of `npu.giz`.

Verified on the Ryzen AI NPU (driver amdxdna, firmware 1.1.2.65): discovery and both reads (tests/hw/device/npu/run.giz),
and the marshalling against fixed bytes (tests/hw/device/npu.test.giz). The reads create no context or buffer.
XRT 2.27's own validate runs (gemm, latency, throughput) pass on a shared server.

Execution (`amdxdna.giz`: open, load, run) gives XRT's results bit-exact on a shared server for both kernel forms:
- add_one, an xclbin's instruction list (`insts.bin`), started with an ERT START_CU packet: 0 of 64 words wrong.
- XRT's validate gemm, a full ELF (`gemm.elf`, its `.ctrltext` is the control code), started with START_NPU: the debug
  buffer after 200 runs equals XRT's (0 of 131072 bytes differ), 5.7 ms a run with its wait.

The NPU does not snoop the CPU caches and reaches every buffer through this process's page tables (the heap by its
process address, a SHARE buffer by its own map). So each host write to a buffer is followed by a cache flush
(`Native.flush`: clflush, as XRT's shim does), a read drops the lines first, and a new buffer is written once (zeros) so
each page is mapped. Without the flush the NPU ran stale control code; without the first write a run right after the
buffer was made lost its output (20 ms later it did not). NPU kernels generated from Gizlang's own IR (an AIE
code-generation backend) are not started.

`amdxdna.giz` also reads an xclbin (`xclbin`: its PDIs and kernel names) and writes AIE transaction control code (`Txn`:
run-time parameter writes, shim DMA copies with the firmware's address patch and task tokens, waits). `Txn` gives the
MLIR-AIE add_one instruction list byte for byte (tests/hw/device/npu.test.giz). A second kernel loaded in a session gets
its own hardware context on the same heap (`more`); `Npu.variant` gives a program its own control code, `Npu.recoded`
replaces it between runs, and `submit` queues a run in its own command buffer so the next can go before it ends.

LLM inference on the NPU: `gizml:npullm` runs Gemma 4 E2B with the layer and lm_head kernels of FastFlowLM (their
xclbins, used unmodified; the runtime is MIT). A token is 35 runs of the layer kernel and one of the lm_head kernel,
each with control code that `npullm` writes for the context length (a port of FastFlowLM's sequence generator). The CPU
does the q8_0 embeddings and the rotary tables; the token is chosen by `gizml:llm`'s sampling (`pick`). The 35 layer
runs are queued at once and waited for once. Measured on a Ryzen AI MAX+ 395 (
XDNA 2, LPDDR5X-8000 256-bit):

| Run | Result |
| --- | --- |
| Load (model.q4nx, 1.5 GB to the NPU, page cache warm) | 2.7-3.3 s |
| "What is the capital of France? Answer in one sentence." | "The capital of France is Paris." |
| 900 tokens of an essay (935 positions, past the 512-position window) | coherent text, 22.8 tokens/s (waits each run) |
| 200 tokens of an essay, the layer runs queued | 24.8 tokens/s (23.7 with a wait each run), the same tokens |
| Prompt of 621 tokens, one pass | 1.25-1.27 s (489-497 tokens/s); FastFlowLM 1.15-1.29 s on the same machine |
| Prompt of 1491 tokens, two passes | 2.95 s (505 tokens/s); a token at a time 63.4 s, the same top tokens |
| Prompt of 21 tokens, one pass | 0.71 s |
| One token: 35 layer runs / lm_head run / CPU | 31.5 ms / 7.6 ms / 0.5 ms |

A token reads all its weights (1.2 GB of layers, 0.25 GB of lm_head) once: 38 GB/s in the layers and 33 GB/s in the
lm_head, about a seventh of the memory's 256 GB/s. The limit is the shim DMA of these kernels, not the memory or the
CPU: a faster token needs kernels that read on more channels, or fewer weight reads a token (several tokens a run: the
prompt, or drafted tokens to check).

The prompt goes in passes of up to 1024 tokens (`Model.pass`; 16 or fewer tokens left go a token at a time) on
FastFlowLM's prefill xclbins: dequant (q4 weights to bf16), the matrix multiply in 256-row blocks
over the 32 cores, and the global and sliding-window attention. `npullm` writes all their control code (`mmCode`,
`dequantCode`, `attentionCode`). The CPU does the norms, rotary embedding, the GELU product and the residual adds a row
at a time in reused buffers (`Session.downloadInto`). Of the 621-token pass, the NPU runs take about 1.2 s. A change of
kernel costs about 3 ms (the same run again: 0.3 ms), so dequant makes the weights of 4 layers in one round (4 sets
of bf16 weights, 512 MB).

### Placement, GPU networks and GPT training

The tensor library has:

- one Tensor placement API (`on`/`host`, GPU F16/BF16 storage);
- GPU network operations;
- reverse rules for embedding, layer norm, GELU, label cross-entropy, layout, dropout, conv and pool;
- decoder-only GPT training and KV-cache inference.

Tests:

- The existing finite-difference step and tolerances, the Philox known answer and the CPU goldens
  stay fixed. Nonfinite oracle values now fail through the shared `Math.isFinite`.
- CPU tensor, grad and nn tests pass, and so do AMD gfx1151 GPU operations, also typed BF16 kernel execution.
- CPU GPT loss decline, counts and KV-cache equivalence pass.

### Apple GPUs: Metal

Written without a Mac host: it is checked by tests that run anywhere, and **no kernel has run on an Apple GPU yet**.


- **Driver (`gizhw:metal`, opened by `gizhw:gpu` as `kind = "metal"`; `GIZLANG_DEVICE=metal`).** Metal.framework
  through Objective-C's runtime with `gizlang:native`. Memory is unified: a buffer is one shared `MTLBuffer`, its
  `gpuAddress` is the address a kernel gets (a view is the address plus an offset, as on CUDA and HIP), and a copy
  to or from the device is a plain copy after the work before. A kernel gets its arguments as one struct of 8-byte
  words (`constant GizArgs&`), so no buffer is bound by index. Every live buffer is made resident for the commands
  after it (`useResources`). Launches go into compute encoders of 128 launches; `sync` commits and waits. `record`
  keeps the launches and encodes them again at `replay`. `slab` fills
  blocks of at most `maxBufferLength`.
- **Printer (`Kernel.msl`).** The kernel IR as MSL: `simd_shuffle*`, `threadgroup` arrays, `threadgroup_barrier`, and
  the 16x16 matrix ops as 2 x 2 `simdgroup_matrix` of 8 x 8 (F16 or BF16 inputs, F32 sums). `hasMma("metal")` is true;
  the integer, FP8 and FP4 matrix ops are not, so those products dequantize to F16 first. Workgroup memory is 32 KB:
  the GEMM takes 2 tiles of rows and the attention tile 32 KB (`gemmTiles`, `tileHeads`); in `gizml:tensor` the
  matrix plans drop stages until they fit `K.sharedLimit` (`mmaPlans`), and the attention takes the widest key block
  that fits, with Q, K and V in one array where separate ones do not (`attentionShape`). Apple GPUs have no f64: an
  F64 value is a double-float pair of f32 (48-bit mantissa, error-free sums and products), and an f64 buffer holds
  IEEE doubles, so the f64 sums (groupNorm, softmax, strict sums) print unchanged.
- **Images and video (`gizml:diffusion`, `gizml:video`).** Without INT8 matrix cores a block runs as BF16 on
  `simdgroup_matrix`: `Block.quantized` keeps it, and `Block.dequantized` turns ComfyUI's INT8 FLUX.3 file and H3's W4
  file into BF16 as they load (ConvRot's rotation turned back). Calculated sizes: FLUX.3 klein about 11 GB (5.6B x 2
  bytes), H3's transformer about 38.5 GB (19.3B x 2), beside their text encoders. H.264 goes through ffmpeg's libx264
  where NVENC is absent (`V.h264`), from the same NV12 kernel.
- **Tested without a Mac:** golden MSL lines; the double-float functions, built for this host, against its f64 (within
  2^-44); and clang's syntax and type check of every kernel the printer makes, of every kernel a tiny language
  model's engine launches (4 and 8-bit weights), and of every kernel FLUX.3's joint block, H3's block from W4,
  FLUX.3's image decoder and H.264's NV12 kernel launch, against a stand-in `metal_stdlib`
  (tests/hw/kernel/metal; the engines on the fake device of `kernelsOnMetal`: tests/ml/llm/metal.test.giz,
  tests/ml/diffusion/metal.test.giz, tests/ml/video/metal.test.giz). On CUDA, H3's W4 block as BF16 gives the
  host's W4 rows within INT8 error (tests/ml/diffusion, testW4BlockMatchesHost). The driver's call sequence is
  checked against a fake Objective-C runtime that answers with host memory (tests/hw/gpu/metal.test.giz). None of
  the Metal checks runs a kernel on a GPU.
- **To measure on a Mac (unverified until then):** that the system compiler accepts the MSL (`#pragma METAL fp`,
  `bfloat`, `simdgroup_matrix` of mixed types); the lane layout of `simdgroup_matrix` elements that `MMA_ROW`,
  `MMA_COL` and the attention's `rowLanes` assume; F32 division and square root rounding against the CPU's bits;
  the cost of a dispatch and of a commit; that the GPU address of a buffer works as a kernel argument; the
  tokens a second of the language models and the seconds a step of FLUX.3 against the memory bandwidth limit.

### Inference: language models, images and video

Three engines run on the same tensors and kernels. Each model file keeps only what is its own, with a reason.

- **Language models (`gizml:llm`, the ML maintainers).** Hugging Face models are converted once to rotated codebooks of 2-8 bits.
  The same code runs on NVIDIA, AMD and the CPU: paged keys and values in 16, 8, 4 or 4/2 bits, many sequences a
  step, MTP drafts where they give tokens faster, continuous batching, and the OpenAI chat API with tools
  (`serve`, `Engine.warm`). Qwen3.8-27B in 4 bits on an RTX 4070 Ti SUPER: 95 tokens/s for one sequence, 395 in
  total for eight, prefill 1062 tokens/s ([bench/ml/llm](../../bench/ml/llm/README.md)).
- **Images and video (`gizml:diffusion`, the ML maintainers).** FLUX.3 klein (text to image, edits, prompt upsampler; one pipeline, `D.Klein`), Z-Image Turbo (`D.ZImagePipeline`) and MiniMax H3 (W4A8 as stored). A FLUX.3
  step at 1024² takes 0.175 s on an RTX 4080 SUPER (ComfyUI 0.172). An H3 request of 832 x 480, 124 frames and
  6 steps takes about 31.5 s warm on a 16 GB RTX 4070 Ti SUPER (WanGP 85 s).
- **Shared parts they use.** In `gizhw:kernel`: s8 mma (m16n8k32, I32 sums) on sm_80 and later; E4M3 mma with F32
  sums, or F16 sums (`mmaHalf`, twice the rate on GeForce, PTX 8.7) on sm_89 and later; `Kernel.hadamard`, also with
  `regular` (ConvRot's matrix). In `gizml:tensor`: INT8 products whose epilogues take scales, bias, SwiGLU pairs,
  a gated residual and a LoRA's up product; `attentionRows` (BF16 or 8-bit, with the head norm and rotary in its
  quantizing pass); `modulatedNorm`, `int8Rows`, `modulatedInt8RowsInto`; `lazyConv` and `conv1d`.
- **Measured limits.** The INT8 products and the 8-bit attention run at 65-80% of the matrix cores' peak. On the
  RTX 4070 Ti SUPER, the card holds its 285 W cap at about 2.63 GHz. The attention kernel waits on the mma latency
  of its P·V chains (255 registers, 2 workgroups an SM).

### Autotune and patterns

Tile search:

- matmul K depths 16, 32 and 64 are measured after warm-up, and cached on the opened GPU by shape,
  layout and dtype. The original order of accumulation stays.
- The shared kernel IR prints LLVM and PTX.
- Cache reuse and CPU/GPU numerical equality pass on AMD gfx1151. The performance improvement is not measured.

Attention 2026-10-16:

- `attention(q, k, v, scale, causal, rate, key)` is one operation for training (with its reverse rule) and for
  inference with a KV cache. It keeps no (tq, tk) score matrix: keys go in blocks of 16 with an online softmax, in
  f32x16 lanes on every core or in one GPU kernel (16 query rows of 16 threads per workgroup). The CPU and the GPU
  take the same f32 steps, so the bits are the same (tests/hw/gpu on AMD gfx1151 and the RTX 5090).
- `gpt`, `gptStep` and `bert` use it in place of three composed copies of masked softmax and batch matmul.
- `geluErf` (the erf form, torch's default) runs fused on the CPU and the GPU with the same bits, in place of a
  host-only f64 loop in `bert`.

Tensor cores 2026-10-16:

- gizhw:kernel has one matrix multiply-accumulate (`mmaZero`, `mmaLoadA`, `mmaLoadB`, `mma`, and `mmaGet`,
  `mmaRow`, `mmaCol` for the elements of each lane): a 16x16x16 F16 or BF16 product into f32. It prints `llvm.amdgcn.wmma` on RDNA 3 (gfx11, 16 elements a lane) and RDNA 4
  (gfx12, 8 elements a lane) and two `mma.sync.m16n8k16` on sm_80 and later. `hasMma(kind, arch)` tells where it
  exists. gfx12 is checked by `llc` only: no RDNA 4 GPU is here.
- An F16 or BF16 matmul on such a GPU uses it. Tune picks one of three tiles of C: 128x128 (8 waves), 128x256 (8
  waves) and 64x64 (4 waves, for a C of too few large tiles). An operand is kept in workgroup memory along k or
  across it, as its 16-bit patterns, and each lane loads its fragments with `ldmatrix` (`.trans` across k). When
  the operand's length along its consecutive dimension is a multiple of 8, the copy is 16 bytes (`copy`: cp.async in
  the background on sm_80 and later, through 4 stages for 128x128 and 3 for 128x256); else it is one 2-byte load an
  element. Tiles go to the workgroups in groups of 8 rows of tiles (`tileOrigin`, also for F32). The fragments of the
  first half of the next step are loaded before its barrier, so the matrix cores do not wait for them.
- C is of the operands' type (`T.productType`, as in torch): the f32 sums are rounded once at the end. On NVIDIA each
  lane stores two elements of a row at once (`Kernel.store2`: `cvt.rn.bf16x2.f32` and one 32-bit store).
- The F32 matmul reads 16 bytes at a time where the dimension read along is a multiple of 4. On CUDA it copies
  through 3 or 4 stages in the background (`Kernel.copy`, `copyElement`: `cp.async`); a whole step of k copies with
  no checks. Elsewhere each thread loads the next step into registers.
- `Gpu.tune` warms each choice up, then times the choices in turn, three times each, and keeps each one's fastest:
  work of another process slows every choice alike. `Gpu.force(position)` makes each tune pick one candidate, so a
  test runs every tile plan.
- The matrix-core plans are 128x128 (8 warps, 4 stages), 128x256 (3 stages), and 64x64, 128x64, 64x128 and 128x128
  of 4 warps for a C of too few large tiles to fill the SMs (1024: 64x128, 22.7 -> 16.7 us).
- The products are exact, but the sums are the hardware's: gfx1151 gives 1 - 2^-24 for an identity product when a
  term is negative. Tests check the f32 bound 2 k 2^-24 sum |ab| for each element. The F32 path keeps its bits.
- gfx1151, 4096x4096x4096, after warm-up, 20 runs: F32 36.6 ms (3.8 TFLOP/s), F16 7.8-8.6 ms (16-18 TFLOP/s), BF16
  7.6-8.3 ms (17-18 TFLOP/s); with 2-byte copies F16 and BF16 were 9.0-9.3 ms (15 TFLOP/s). 1024 and 2048: 14-16
  TFLOP/s. In one step of k the BF16 kernel does 4 `global_load_b128` in place of 32 `global_load_u16`.
- **FP8 (E4M3).** `T.FP8` tensors hold 1-byte E4M3 values (`K.E_E4M3`: at most ±448, no infinities). A cast rounds to
  nearest even and saturates; a load decodes with integer operations, so every kernel takes FP8 inputs. A product of
  FP8 operands is F32. On NVIDIA sm_89 and later (`K.hasFp8Mma`, PTX 8.4) it runs on the matrix cores: the mma kernel reads
  pairs of E4M3 values as 16-bit elements (`K.E_E4M3X2`), so with k even the tiles, copies and `ldmatrix` are the BF16
  ones, and `mma` prints `mma.sync.m16n8k32...e4m3.e4m3`. Both operands must lie along k (A by rows, B by columns);
  a 2-D operand that does not is transposed first. Elsewhere, or with k odd, the FMA kernel decodes the bytes. A
  cast has no scale. `T.fp8Scaled(t)` scales as tensorwise scaling does (Transformer Engine): it gives t / s as FP8 and
  s = max |t| / 448, a tensor of shape [] on the GPU (no copy to the host). `T.lazyMatmulIn(a, ta, b, tb, FP8)` of other
  types scales both operands and multiplies the product by sa sb, so values of 1e-4 or 1e3 keep 3 bits of mantissa.
  `Graph.autocast(FP8)` trains so: each product of a step, forward and backward, scales its operands; the gradients
  are F32. RTX 5090, B by columns, F32 out: 2048 / 4096 / 8192 take 0.0554 /
  0.381 / 2.70 ms (310 / 360 / 407 TFLOP/s), 1.96 times BF16; `torch._scaled_mm` with scale 1 takes 0.0514 / 0.366 /
  2.51 ms.
- **FP4 (NVFP4).** `T.lazyMatmulIn(a, ta, b, tb, T.FP4)` multiplies in NVFP4. `T.fp4Blocks(t, across)` gives
  [q, s]: q holds E2M1 values (at most ±6) in pairs, then one E4M3 scale for each block of 16 along k; s =
  max |t| / 2688 is a GPU tensor of shape [], so a block scale fits E4M3. The product is multiplied by sa sb. On
  NVIDIA sm_120 and sm_121 (`K.hasFp4Mma`, target sm_120a) with k a multiple of 64, the mma kernel reads E2M1 values
  in groups of 4 (`K.E_E2M1X4`) and `K.mmaScaled` prints `mma.sync...m16n8k64.kind::mxf4nvf4.block_scale`. Each
  stage copies the scale words of its tiles to workgroup memory with the operands. Elsewhere `fp4Values` decodes q
  to F32 first. The error of NVFP4 on normal values is about 13% RMS (FP8 4%). RTX 5090, F32 out: 4096 / 8192 take
  0.180 / 1.13 ms (762 / 970 TFLOP/s); `torch._scaled_mm` NVFP4 takes 0.121 / 0.813 ms (1135 / 1353 TFLOP/s).
- `Graph.autocast(BF16)` runs the products of a training step on the matrix cores. On a two-layer step (tests/hw/gpu)
  the loss and the gradients are within 0.5% of the largest f32 value.
- RTX 5090 (sm_120), 2026-10-17: tests/hw/kernel and tests/hw/gpu pass (mma.sync within the f32 bound, autocast).
  Against PyTorch 2.8 eager (cuBLAS, TF32 off) on the same host, the same types in and out, 20-200 runs after
  warm-up, with the GPU idle:

  | Work | Gizlang | PyTorch |
  | --- | --- | --- |
  | matmul F32 4096 / 8192 | 2.11 / 16.9 ms (65 / 65 TFLOP/s) | 1.96 / 15.4 ms (70 / 71) |
  | matmul F16 4096 / 8192 | 0.80 / 5.10 ms (171 / 215 TFLOP/s) | 0.71 / 5.00 ms (193 / 220) |
  | matmul F16 / BF16 1024 | 0.0167 / 0.0167 ms (129 TFLOP/s) | 0.0144 / 0.0226 ms (149 / 95) |
  | matmul F16 / BF16 2048 | 0.103 / 0.103 ms (167 TFLOP/s) | 0.100 / 0.100 ms (171) |
  | matmul BF16 4096 / 8192 | 0.726 / 5.07 ms (189 / 217 TFLOP/s) | 0.784 / 4.84 ms (175 / 227) |
  | MLP 1024-4096-4096-1024 step, batch 4096, SGD: F32 / BF16 autocast | 11.6 / 3.50 ms | 9.6 / 3.9 ms |

  BF16 4096 is faster than cuBLAS, the other matrix-core sizes are within 2-13%. F32 is 91-93% of cuBLAS, with the
  tiles of 256 x 128 (it was 78% with 128 x 128). The RTX 5090 runs at 2932 MHz for a short time only, then at 2400 MHz
  or less: the F32 rows are the best single run of 60 in a new process on both sides, with no tuning before it. An F32
  operand read along k (A, or B transposed) lies along k in workgroup memory: its copies are 16 bytes, not 4, and a
  thread's elements of it are 16 apart, so a warp's reads meet each bank once. That made A B, A Bt and At Bt 5-15%
  faster (2026-10-17; At B, with neither operand along k, is the same kernel as before). The
  kernel has 193 registers (one workgroup on each SM) and no spills. ptxas already makes the workgroup-memory
  addresses 32-bit with immediate offsets (I32 indices in the kernel IR measured the same).

### Live graph

`giz live` swaps functions and GPU kernels by semantic hash while the program runs: an edit compiles again only the
kernels whose meaning changed ([live code swap](runtime.md#live-code-swap)).