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.

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).
• 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).

• 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):

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).

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:

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):

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).
• 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).