# gizhw:kernel

GPU kernels written once and printed for every GPU. A Kernel is a small structured program (values, variables, if, loop, buffers, workgroup memory, barriers) built by calls; `llvm` prints it as LLVM IR for AMD (compiled by comgr), `ptx` as PTX for NVIDIA. A value is used only inside the block that made it and the blocks within; what crosses blocks goes through a variable.

```gizlang
import * as Kernel from "gizhw:kernel";
```

## Types

### `KIns`

```gizlang
type KIns = { op: i64; t: i64; a: i64; b: i64; c: i64; e: i64; n: i64; k: f64; }
```

One instruction: `t` the type it makes, `a`, `b`, `c` operands (values, or a
buffer/param/shared index), `e` an element type, `n` an integer constant, `k` a float one.

- `op`: `i64`
- `t`: `i64`
- `a`: `i64`
- `b`: `i64`
- `c`: `i64`
- `e`: `i64`
- `n`: `i64`
- `k`: `f64`

### `Kernel`

```gizlang
type Kernel = { name: string; threads: i64; params: Vec<Param>; shared: Vec<Shared>; code: Vec<KIns>; counters: Vec<i64>; }
```

A kernel: one dimension of workgroups of `threads` threads.

- `name`: `string`
- `threads`: `i64`
- `params`: `Vec<M12.Param>`
- `shared`: `Vec<M12.Shared>`
- `code`: `Vec<M12.KIns>`
- `counters`: `Vec<i64>`

## Functions

### `valueOf`

```gizlang
function valueOf(elem: i64): i64
```

The value type an element loads as.

### `sizeOf`

```gizlang
function sizeOf(elem: i64): i64
```

Bytes of one element.

### `kernel`

```gizlang
function kernel(name: string, threads: i64): Kernel
```

An empty kernel NAME that runs THREADS threads in each workgroup.

### `regularSign`

```gizlang
function regularSign(j: i64): i64
```

-1 where an odd number of the base-4 digits of index j are 3, else 1 (hadamard's `regular`).

Effects: fail.

### `regularSource`

```gizlang
function regularSource(j: i64): i64
```

Index j with the two bits of each base-4 digit swapped (hadamard's `regular`; its own inverse).

Effects: fail.

### `hasTma`

```gizlang
function hasTma(kind: string, arch: string): bool
```

Whether the copy engine ops run on this GPU: the tensor and bulk loads with their barriers (mbarrierInit to
tensorLoad) and the grid dependency (awaitGrid, releaseGrid). NVIDIA sm_90 and later.

Effects: fail.

### `hasMmaPack`

```gizlang
function hasMmaPack(kind: string, arch: string): bool
```

Whether mmaPackA runs on this GPU: NVIDIA sm_80 and later, where a lane's accumulator elements are an A tile's
(the m16n8k16 fragments).

Effects: fail.

### `hasMmaWords`

```gizlang
function hasMmaWords(kind: string, arch: string): bool
```

Whether mmaWords runs on this GPU: AMD gfx11 (RDNA 3), where a lane holds all 16 elements of its row (A) or column
(B) l % 16.

Effects: fail.

### `sharedLimit`

```gizlang
function sharedLimit(kind: string, arch: string): i64
```

The workgroup memory a workgroup may use on a GPU of `kind` ("cuda", "hip", "metal") and `arch`, in bytes.

Effects: fail.

### `hasMma`

```gizlang
function hasMma(kind: string, arch: string): bool
```

Whether the printers emit mma for this GPU: AMD gfx11 and gfx120x (RDNA 3 and 4, WMMA in waves of 32), NVIDIA sm_80 and later, or any Apple GPU (simdgroup_matrix, 8x8 in a simdgroup of 32).

Effects: fail.

### `hasFp8Mma`

```gizlang
function hasFp8Mma(kind: string, arch: string): bool
```

Whether mma takes E4M3 pairs (E_E4M3X2) on this GPU: NVIDIA sm_89 and later (such a kernel is PTX 8.4 or later).

Effects: fail.

### `hasI8Mma`

```gizlang
function hasI8Mma(kind: string, arch: string): bool
```

Whether mma takes pairs of 8-bit integers (E_I8X2) on this GPU: NVIDIA sm_80 and later.

Effects: fail.

### `hasI4Mma`

```gizlang
function hasI4Mma(kind: string, arch: string): bool
```

Whether mma takes quads of 4-bit integers (E_I4X4) on this GPU: NVIDIA sm_80 to sm_89 (Hopper and later drop or
emulate it).

Effects: fail.

### `hasFp4Mma`

```gizlang
function hasFp4Mma(kind: string, arch: string): bool
```

Whether mmaScaled runs on this GPU: NVIDIA sm_120 and sm_121 (its PTX targets sm_120a, sm_121a).

### `ptxVersion`

```gizlang
function ptxVersion(arch: string): string
```

The PTX version a target needs.

Effects: fail.

## Methods

### `Kernel.buffer`

```gizlang
function Kernel.buffer(self: mut Kernel, name: string, elem: i64): i64
```

A buffer parameter of `elem` elements: its index for load and store. Buffers come before scalars.

Effects: fail.

### `Kernel.scalar`

```gizlang
function Kernel.scalar(self: mut Kernel, name: string): i64
```

A 64-bit integer parameter, as a value.

Effects: fail.

### `Kernel.workgroup`

```gizlang
function Kernel.workgroup(self: mut Kernel, elem: i64, count: i64): i64
```

Workgroup memory of `count` elements: its index for sload and sstore.

Effects: fail.

### `Kernel.float`

```gizlang
function Kernel.float(self: mut Kernel, x: f64): i64
```

The F32 constant X.

Effects: fail.

### `Kernel.double`

```gizlang
function Kernel.double(self: mut Kernel, x: f64): i64
```

The F64 constant X.

Effects: fail.

### `Kernel.int`

```gizlang
function Kernel.int(self: mut Kernel, x: i64): i64
```

The I32 constant X.

Effects: fail.

### `Kernel.long`

```gizlang
function Kernel.long(self: mut Kernel, x: i64): i64
```

The I64 constant X.

Effects: fail.

### `Kernel.truth`

```gizlang
function Kernel.truth(self: mut Kernel, x: bool): i64
```

The BOOL constant X.

Effects: fail.

### `Kernel.tid`

```gizlang
function Kernel.tid(self: mut Kernel): i64
```

The thread's index in its workgroup, and the workgroup's index (I64).

Effects: fail.

### `Kernel.gid`

```gizlang
function Kernel.gid(self: mut Kernel): i64
```

The workgroup's index in the grid (I64).

Effects: fail.

### `Kernel.global`

```gizlang
function Kernel.global(self: mut Kernel): i64
```

The global index: gid * threads + tid.

Effects: fail.

### `Kernel.add`

```gizlang
function Kernel.add(self: mut Kernel, a: i64, b: i64): i64
```

A + B (both of one type).

Effects: fail.

### `Kernel.sub`

```gizlang
function Kernel.sub(self: mut Kernel, a: i64, b: i64): i64
```

A - B (both of one type).

Effects: fail.

### `Kernel.mul`

```gizlang
function Kernel.mul(self: mut Kernel, a: i64, b: i64): i64
```

A * B (both of one type).

Effects: fail.

### `Kernel.div`

```gizlang
function Kernel.div(self: mut Kernel, a: i64, b: i64): i64
```

Division: IEEE rounded for floats on every GPU, truncating for signed integers.

Effects: fail.

### `Kernel.rem`

```gizlang
function Kernel.rem(self: mut Kernel, a: i64, b: i64): i64
```

The remainder of A / B, with the sign of A (signed integers and floats).

Effects: fail.

### `Kernel.udiv`

```gizlang
function Kernel.udiv(self: mut Kernel, a: i64, b: i64): i64
```

Unsigned integer division, remainder and compare. By a constant power of two they are a shift and a mask (a
64-bit division is a called routine on NVIDIA).

Effects: fail.

### `Kernel.urem`

```gizlang
function Kernel.urem(self: mut Kernel, a: i64, b: i64): i64
```

The remainder of A / B for unsigned integers.

Effects: fail.

### `Kernel.ult`

```gizlang
function Kernel.ult(self: mut Kernel, a: i64, b: i64): i64
```

A < B for unsigned integers.

Effects: fail.

### `Kernel.shuffleXor`

```gizlang
function Kernel.shuffleXor(self: mut Kernel, v: i64, mask: i64): i64
```

The value V of the lane whose index in the warp (32 lanes) is this lane's xor MASK (a butterfly exchange; F32 or
I32). Every lane of the warp must run it.

Effects: fail.

### `Kernel.warpSum`

```gizlang
function Kernel.warpSum(self: mut Kernel, v: i64): i64
```

The sum, the greatest and the least of V (F32 or I32) over the warp's 32 lanes, in every lane (butterfly exchanges).
Every lane of the warp must run them.

Effects: fail.

### `Kernel.warpMax`

```gizlang
function Kernel.warpMax(self: mut Kernel, v: i64): i64
```

The greatest of V over the warp's 32 lanes, in every lane. Every lane of the warp must run it.

Effects: fail.

### `Kernel.warpMin`

```gizlang
function Kernel.warpMin(self: mut Kernel, v: i64): i64
```

The least of V over the warp's 32 lanes, in every lane. Every lane of the warp must run it.

Effects: fail.

### `Kernel.groupArgmax`

```gizlang
function Kernel.groupArgmax(self: mut Kernel, v: i64, i: i64): Vec<i64>
```

The largest of the workgroup's values v (F32) with its index i (I32), the lowest index of equal ones, in every
thread: [value, index]. A NaN never wins. Every thread of the workgroup must run it.

Effects: fail.

### `Kernel.slicesSum`

```gizlang
function Kernel.slicesSum(self: mut Kernel, b: i64, i: i64, total: i64, slices: i64): i64
```

The sum of element i of the `slices` slices of `total` elements of b, added in order (split-K's partial sums).

Effects: fail.

### `Kernel.shuffle`

```gizlang
function Kernel.shuffle(self: mut Kernel, v: i64, from: i64): i64
```

The value V of lane `from` (I32 or I64, 0 to 31) of the warp (F32 or I32). Every lane of the warp must run it.

Effects: fail.

### `Kernel.hadamard`

```gizlang
function Kernel.hadamard(self: mut Kernel, v: Vec<i64>, lane: i64, block: i64, regular: bool = false): Vec<i64>
```

The Hadamard transform of each block of `block` values of a warp, divided by sqrt(block): lane l holds the n values
nl to nl + n - 1 (n = v.length a power of 2, block a power of 2 from n to 32 n). The transform keeps dot products
and spreads a block's outliers over it (rotation before quantization, as QuaRot). With `regular` (block a power of
4), each value first takes its regularSign: output j is then regularSign(j) times output regularSource(j) of the
regular Hadamard matrix (the Kronecker power of H4, whose row a has -1 at column 3 - a; ConvRot's), so a product
with weights rotated by that matrix takes them in that order and sign.

Effects: fail.

### `Kernel.regularIndex`

```gizlang
function Kernel.regularIndex(self: mut Kernel, j: i64, block: i64): Vec<i64>
```

regularSource of index j (I64) within its block of `block` (a power of 4), and whether its regularSign is -1
(BOOL): their GPU form.

Effects: fail.

### `Kernel.fma`

```gizlang
function Kernel.fma(self: mut Kernel, a: i64, b: i64, c: i64): i64
```

a * b + c rounded once.

Effects: fail.

### `Kernel.min`

```gizlang
function Kernel.min(self: mut Kernel, a: i64, b: i64): i64
```

Floats: the number when the other is NaN (maxnum); integers signed.

Effects: fail.

### `Kernel.max`

```gizlang
function Kernel.max(self: mut Kernel, a: i64, b: i64): i64
```

The larger of A and B. Floats: the number when the other is NaN (maxnum); integers signed.

Effects: fail.

### `Kernel.neg`

```gizlang
function Kernel.neg(self: mut Kernel, a: i64): i64
```

-A.

Effects: fail.

### `Kernel.abs`

```gizlang
function Kernel.abs(self: mut Kernel, a: i64): i64
```

The absolute value of A.

Effects: fail.

### `Kernel.sqrt`

```gizlang
function Kernel.sqrt(self: mut Kernel, a: i64): i64
```

The square root of A.

Effects: fail.

### `Kernel.rint`

```gizlang
function Kernel.rint(self: mut Kernel, a: i64): i64
```

Round to the nearest integer, ties to even.

Effects: fail.

### `Kernel.floor`

```gizlang
function Kernel.floor(self: mut Kernel, a: i64): i64
```

A rounded down to an integer.

Effects: fail.

### `Kernel.lt`

```gizlang
function Kernel.lt(self: mut Kernel, a: i64, b: i64): i64
```

A < B: ordered for floats (false with a NaN).

Effects: fail.

### `Kernel.le`

```gizlang
function Kernel.le(self: mut Kernel, a: i64, b: i64): i64
```

A <= B: ordered for floats (false with a NaN).

Effects: fail.

### `Kernel.gt`

```gizlang
function Kernel.gt(self: mut Kernel, a: i64, b: i64): i64
```

A > B: ordered for floats (false with a NaN).

Effects: fail.

### `Kernel.ge`

```gizlang
function Kernel.ge(self: mut Kernel, a: i64, b: i64): i64
```

A >= B: ordered for floats (false with a NaN).

Effects: fail.

### `Kernel.eq`

```gizlang
function Kernel.eq(self: mut Kernel, a: i64, b: i64): i64
```

Ordered float compares (false with a NaN); eq/ne also on integers and bools.

Effects: fail.

### `Kernel.ne`

```gizlang
function Kernel.ne(self: mut Kernel, a: i64, b: i64): i64
```

A != B; for floats, true with a NaN.

Effects: fail.

### `Kernel.isNan`

```gizlang
function Kernel.isNan(self: mut Kernel, a: i64): i64
```

Whether A is NaN.

Effects: fail.

### `Kernel.select`

```gizlang
function Kernel.select(self: mut Kernel, c: i64, a: i64, b: i64): i64
```

A where C is true, else B (A and B of one type).

Effects: fail.

### `Kernel.and`

```gizlang
function Kernel.and(self: mut Kernel, a: i64, b: i64): i64
```

A AND B: bitwise for integers, logical for bools.

Effects: fail.

### `Kernel.both`

```gizlang
function Kernel.both(self: mut Kernel, a: i64, b: i64): i64
```

Whether truths a and b both hold.

Effects: fail.

### `Kernel.or`

```gizlang
function Kernel.or(self: mut Kernel, a: i64, b: i64): i64
```

A OR B: bitwise for integers, logical for bools.

Effects: fail.

### `Kernel.xor`

```gizlang
function Kernel.xor(self: mut Kernel, a: i64, b: i64): i64
```

A XOR B: bitwise for integers, logical for bools.

Effects: fail.

### `Kernel.not`

```gizlang
function Kernel.not(self: mut Kernel, a: i64): i64
```

NOT A: bitwise for integers, logical for bools.

Effects: fail.

### `Kernel.shl`

```gizlang
function Kernel.shl(self: mut Kernel, a: i64, b: i64): i64
```

A shifted left by B bits.

Effects: fail.

### `Kernel.shr`

```gizlang
function Kernel.shr(self: mut Kernel, a: i64, b: i64): i64
```

Arithmetic and logical shifts right.

Effects: fail.

### `Kernel.ushr`

```gizlang
function Kernel.ushr(self: mut Kernel, a: i64, b: i64): i64
```

A shifted right by B bits, with zeros (logical).

Effects: fail.

### `Kernel.to`

```gizlang
function Kernel.to(self: mut Kernel, t: i64, a: i64): i64
```

A conversion: floats round to nearest even, floats to integers toward zero,
integers widen signed; BOOL to a number is 1 or 0.

Effects: fail.

### `Kernel.bits`

```gizlang
function Kernel.bits(self: mut Kernel, t: i64, a: i64): i64
```

The same bits as another type of the same size (F32/I32, F64/I64).

Effects: fail.

### `Kernel.load`

```gizlang
function Kernel.load(self: mut Kernel, b: i64, index: i64): i64
```

Element `index` (I64) of buffer `b`, as its value type.

Effects: fail.

### `Kernel.e4m3`

```gizlang
function Kernel.e4m3(self: mut Kernel, x: i64): i64
```

The E4M3 byte (I32, 0 to 255) of F32 x, as a store into E4M3 memory gives it: two of them make an E4M3X2 element.

Effects: fail.

### `Kernel.toByte`

```gizlang
function Kernel.toByte(self: mut Kernel, v: i64): i64
```

v (F32) rounded and clamped to a byte, as I32.

Effects: fail.

### `Kernel.roundTo`

```gizlang
function Kernel.roundTo(self: mut Kernel, e: i64, x: i64): i64
```

F32 X rounded to the values of element type E and back to F32, as a store and a load of E give it: F16 and BF16 to
nearest even (F16 past 65504 to infinity, its subnormals exact), E4M3 as store encodes it; F32 unchanged. NaN stays
NaN. gizml:tile's casts use it, and its CPU side has the same arithmetic.

Effects: fail.

### `Kernel.store`

```gizlang
function Kernel.store(self: mut Kernel, b: i64, index: i64, value: i64): void
```

Stores `value` (the element's value type) as element `index` of buffer `b`; F16 and BF16 round to nearest even. An
I32 value into F16 or BF16 elements stores its low 16 bits unchanged (the bits of a U16 load), so a copy converts
nothing; sstore does the same.

Effects: fail.

### `Kernel.store2`

```gizlang
function Kernel.store2(self: mut Kernel, b: i64, index: i64, x: i64, y: i64): void
```

Stores F32 values x and y as elements `index` (even) and index + 1 of F32, F16 or BF16 buffer `b`, rounded as
store does: one store on NVIDIA (sm_80 and later for F16 and BF16).

Effects: fail.

### `Kernel.sstore2`

```gizlang
function Kernel.sstore2(self: mut Kernel, s: i64, index: i64, x: i64, y: i64): void
```

As store2, into F16 or BF16 workgroup memory `s`.

Effects: fail.

### `Kernel.sload`

```gizlang
function Kernel.sload(self: mut Kernel, s: i64, index: i64): i64
```

The element INDEX of the workgroup memory S.

Effects: fail.

### `Kernel.sstore`

```gizlang
function Kernel.sstore(self: mut Kernel, s: i64, index: i64, value: i64): void
```

Stores VALUE at element INDEX of the workgroup memory S.

Effects: fail.

### `Kernel.sload4`

```gizlang
function Kernel.sload4(self: mut Kernel, s: i64, index: i64): i64
```

16 bytes of workgroup memory `s` from element `index` (its byte offset a multiple of 16) in one load. Its 4
lanes (lane(v, j)) are F32 elements for F32 memory, else 32-bit words of raw bits (I32: two F16, BF16 or U16
elements, the first in the low half).

Effects: fail.

### `Kernel.load4`

```gizlang
function Kernel.load4(self: mut Kernel, b: i64, index: i64, ok: i64): i64
```

16 bytes of buffer `b` from element `index` (its byte offset a multiple of 16) in one load, with lanes as
sload4. Where `ok` is false nothing is read and every lane is 0 (the edge of a tile).

Effects: fail.

### `Kernel.lane`

```gizlang
function Kernel.lane(self: Kernel, v: i64, j: i64): i64
```

Lane J (0 to 3) of the wide load V.

Effects: fail.

### `Kernel.sstore4`

```gizlang
function Kernel.sstore4(self: mut Kernel, s: i64, index: i64, v: i64): void
```

The 16 bytes of `v` (a load4 or sload4) unchanged at element `index` of workgroup memory `s` (its byte offset
a multiple of 16).

Effects: fail.

### `Kernel.copy`

```gizlang
function Kernel.copy(self: mut Kernel, s: i64, slot: i64, b: i64, index: i64, ok: i64): void
```

Copies 16 bytes of buffer `b` from element `index` to workgroup memory `s` at element `slot` (both byte offsets
multiples of 16), or 16 zero bytes where `ok` is false. On NVIDIA sm_80 and later the copy runs in the background
(cp.async): it is done after `commit` closes its group and `wait` lets at most that many later groups run, and the
others see it after a barrier. Elsewhere it is done at once.

Effects: fail.

### `Kernel.copyElement`

```gizlang
function Kernel.copyElement(self: mut Kernel, s: i64, slot: i64, b: i64, index: i64, ok: i64): void
```

As copy, for one 4-byte element (its byte offsets multiples of 4).

Effects: fail.

### `Kernel.commit`

```gizlang
function Kernel.commit(self: mut Kernel): void
```

Closes the group of the copies started since the last commit.

Effects: fail.

### `Kernel.wait`

```gizlang
function Kernel.wait(self: mut Kernel, pending: i64): void
```

Waits until at most `pending` groups of this thread's copies are still running.

Effects: fail.

### `Kernel.mbarrierInit`

```gizlang
function Kernel.mbarrierInit(self: mut Kernel, s: i64, slot: i64, count: i64): void
```

Sets barrier `slot` of `s` (a workgroup array of E_I64) to wait for `count` arrivals; then fenceBarriers once,
and a workgroup barrier before another thread uses it.

Effects: fail.

### `Kernel.fenceBarriers`

```gizlang
function Kernel.fenceBarriers(self: mut Kernel): void
```

Makes the barriers set by mbarrierInit visible to the copy engine.

Effects: fail.

### `Kernel.mbarrierExpect`

```gizlang
function Kernel.mbarrierExpect(self: mut Kernel, s: i64, slot: i64, bytes: i64): void
```

Arrives at barrier `slot` of `s` and adds `bytes` to the bytes its phase waits for (the copies to it).

Effects: fail.

### `Kernel.mbarrierArrive`

```gizlang
function Kernel.mbarrierArrive(self: mut Kernel, s: i64, slot: i64): void
```

Arrives at barrier `slot` of `s`.

Effects: fail.

### `Kernel.mbarrierWait`

```gizlang
function Kernel.mbarrierWait(self: mut Kernel, s: i64, slot: i64, parity: i64): void
```

Waits until the phase of barrier `slot` of `s` with parity `parity` (0 or 1, a value) has completed.

Effects: fail.

### `Kernel.tensorLoad`

```gizlang
function Kernel.tensorLoad(self: mut Kernel, s: i64, slot: i64, map: i64, x: i64, y: i64, bars: i64, bar: i64): void
```

Starts the copy engine's load of the box at byte x of row y of the 2-D tensor whose map is at global address
`map` (Gpu.tensorMap) into `s` from element `slot`; barrier `bar` of `bars` counts its bytes.
Elements outside the tensor are zeros. One thread issues it.

Effects: fail.

### `Kernel.bulkLoad`

```gizlang
function Kernel.bulkLoad(self: mut Kernel, s: i64, slot: i64, b: i64, index: i64, bytes: i64, bars: i64, bar: i64): void
```

Starts the copy engine's load of `bytes` (a multiple of 16) from element `index` of buffer `b` into `s` from
element `slot` (both 16-byte aligned); barrier `bar` of `bars` counts them. One thread issues it.

Effects: fail.

### `Kernel.awaitGrid`

```gizlang
function Kernel.awaitGrid(self: mut Kernel): void
```

Waits for the grid before this one on the stream to finish or call releaseGrid. A kernel with awaitGrid starts
while that grid still runs (Gpu.launch gives it the programmatic dependency): the work before awaitGrid, such as
its barriers and prefetches, overlaps the end of the grid before.

Effects: fail.

### `Kernel.awaitsGrid`

```gizlang
function Kernel.awaitsGrid(self: Kernel): bool
```

Whether the kernel waits at awaitGrid.

Effects: fail.

### `Kernel.releaseGrid`

```gizlang
function Kernel.releaseGrid(self: mut Kernel): void
```

Lets the next grid on the stream start (it still waits at its awaitGrid).

Effects: fail.

### `Kernel.sharedBytes`

```gizlang
function Kernel.sharedBytes(self: Kernel): i64
```

Bytes of workgroup memory the kernel uses.

Effects: fail.

### `Kernel.barrier`

```gizlang
function Kernel.barrier(self: mut Kernel): void
```

Every thread of the workgroup waits here, and sees the others' workgroup memory writes.

Effects: fail.

### `Kernel.variable`

```gizlang
function Kernel.variable(self: mut Kernel, init: i64): i64
```

A variable of `init`'s type (a tile's element type too), set to it.

Effects: fail.

### `Kernel.get`

```gizlang
function Kernel.get(self: mut Kernel, v: i64): i64
```

The value of the variable V.

Effects: fail.

### `Kernel.set`

```gizlang
function Kernel.set(self: mut Kernel, v: i64, value: i64): void
```

Sets the variable V to VALUE.

Effects: fail.

### `Kernel.when`

```gizlang
function Kernel.when(self: mut Kernel, c: i64): void
```

if (c) { ... } [else { ... }] end

Effects: fail.

### `Kernel.otherwise`

```gizlang
function Kernel.otherwise(self: mut Kernel): void
```

Starts the else block of the open when.

Effects: fail.

### `Kernel.end`

```gizlang
function Kernel.end(self: mut Kernel): void
```

Ends the open when.

Effects: fail.

### `Kernel.loop`

```gizlang
function Kernel.loop(self: mut Kernel): void
```

loop { ... breakIf(c) ... } endLoop: breakIf only directly in the loop's body.

Effects: fail.

### `Kernel.breakIf`

```gizlang
function Kernel.breakIf(self: mut Kernel, c: i64): void
```

Leaves the loop when C is true; only directly in the loop's body.

Effects: fail.

### `Kernel.endLoop`

```gizlang
function Kernel.endLoop(self: mut Kernel): void
```

Ends the loop.

Effects: fail.

### `Kernel.repeat`

```gizlang
function Kernel.repeat(self: mut Kernel, from: i64, to: i64, step: i64): i64
```

for (i = from; i < to; i += step) { ... } next(): returns i (I64); `to` is read each turn.

Effects: fail.

### `Kernel.unroll`

```gizlang
function Kernel.unroll(self: mut Kernel, count: i64): void
```

Asks the compiler to unroll the loop just ended `count` times (a hint).

Effects: fail.

### `Kernel.next`

```gizlang
function Kernel.next(self: mut Kernel): void
```

The end of the innermost repeat.

Effects: fail.

### `Kernel.mmaZero`

```gizlang
function Kernel.mmaZero(self: mut Kernel): i64
```

The 16x16 F32 accumulator of zeros.

Effects: fail.

### `Kernel.mmaLoadA`

```gizlang
function Kernel.mmaLoadA(self: mut Kernel, s: i64, at: i64, stride: i64, swizzled: bool = false, elem: i64 = -1): i64
```

The 16x16 tile of A at element `at` of F16, BF16 or packed (E4M3X2, E2M1X4, I8X2, I4X4) workgroup memory `s`, row-major: A[i][p] at at + i stride + p.
`swizzled` (stride 32, `at` a multiple of 16 rows): the 8-element chunk of p is XOR (i / 2) % 4 in its row, so an
8x8 matrix meets every bank once with no padding. `elem` (-1: s's) reads s's 16-bit elements as another type.

Effects: fail.

### `Kernel.mmaLoadB`

```gizlang
function Kernel.mmaLoadB(self: mut Kernel, s: i64, at: i64, stride: i64, swizzled: bool = false, elem: i64 = -1): i64
```

The 16x16 tile of B at element `at` of F16, BF16 or packed (E4M3X2, E2M1X4, I8X2, I4X4) workgroup memory `s`, stored by columns: B[p][j] at at + j stride + p;
`swizzled` and `elem` as mmaLoadA's, by column j.

Effects: fail.

### `Kernel.mmaLoadAT`

```gizlang
function Kernel.mmaLoadAT(self: mut Kernel, s: i64, at: i64, stride: i64): i64
```

The tile of A stored by columns (along k): A[i][p] at at + p stride + i.

Effects: fail.

### `Kernel.mmaLoadBT`

```gizlang
function Kernel.mmaLoadBT(self: mut Kernel, s: i64, at: i64, stride: i64): i64
```

The tile of B stored by rows: B[p][j] at at + p stride + j.

Effects: fail.

### `Kernel.mma`

```gizlang
function Kernel.mma(self: mut Kernel, a: i64, b: i64, c: i64): i64
```

c + a b, with the f32 sums of the hardware (not the order of fma): A and B tiles of one element type.

Effects: fail.

### `Kernel.mmaHalf`

```gizlang
function Kernel.mmaHalf(self: mut Kernel, a: i64, b: i64, c: i64): i64
```

c + a b of E4M3X2 tiles with the F16 sums of the hardware (twice the F32 rate on GeForce Ada): c is mmaZero's
or mmaHalf's, and mmaGet gives the sums as F32. Each sum must stay within F16's range (65504).

Effects: fail.

### `Kernel.mmaScaled`

```gizlang
function Kernel.mmaScaled(self: mut Kernel, a: i64, b: i64, c: i64, scales: i64): i64
```

c + a b of E2M1X4 tiles (16 rows of 64 FP4 values), each value times the scale of its block of 16 along k:
FP8 E4M3 bytes (positive), 4 to a 32-bit word in the order of k. `scales` (I64) holds two words: in its low half
the word of the tile's row g + 8 (t % 2) of A, in its high half that of column g + 8 (t % 2) of B (lane l is in
group g = l / 4 at t = l % 4).

Effects: fail.

### `Kernel.mmaGet`

```gizlang
function Kernel.mmaGet(self: mut Kernel, c: i64, r: i64, integer: bool = false): i64
```

This lane's element r (0 to 7) of accumulator `c` as F32: C[mmaRow(r)][mmaCol(r)]. With `integer` the
accumulator holds the I32 sums of I8X2 or I4X4 products, converted here.

Effects: fail.

### `Kernel.mmaRow`

```gizlang
function Kernel.mmaRow(self: mut Kernel, r: i64): i64
```

The row (I64, 0 to 15) of this lane's accumulator element r.

Effects: fail.

### `Kernel.mmaCol`

```gizlang
function Kernel.mmaCol(self: mut Kernel, r: i64): i64
```

The column (I64, 0 to 15) of this lane's accumulator element r.

Effects: fail.

### `Kernel.mmaPackA`

```gizlang
function Kernel.mmaPackA(self: mut Kernel, values: Vec<i64>, e: i64): i64
```

The F16 or BF16 A tile of the 16x16 values the lanes hold in accumulator order (`values[r]` at mmaRow(r),
mmaCol(r)), rounded to nearest even, in registers: FlashAttention's weights P go from the scores' accumulators to
the next product without workgroup memory. Only where hasMmaPack. E_E4M3X2 takes 16 values, the 32 x 16 tile of
m16n8k32 (hasFp8Mma): register j holds values 4 j to 4 j + 3, one row's 4 adjacent inputs, so a caller passes
elements 0, 1, 4, 5 of a 16-column accumulator for row g and 2, 3, 6, 7 for row g + 8, and orders the B rows to match.

Effects: fail.

### `Kernel.mmaWords`

```gizlang
function Kernel.mmaWords(self: mut Kernel, words: Vec<i64>, e: i64): i64
```

The F16 or BF16 tile whose row (A) or column (B) l % 16 holds, in lane l, the 16 elements of the 8 I32 `words` in the
order of k, element 2 j in the low half of word j: a tile made in registers. Only where hasMmaWords. On NVIDIA
(hasFp8Mma) an E4M3X2 A tile of 4 `words`, register j as mmaPackA's: its E4M3 bytes as they are, low byte first.

Effects: fail.

### `Kernel.exp`

```gizlang
function Kernel.exp(self: mut Kernel, x: i64): i64
```

e^x within 2 ulp (Cephes expf), as tensor's vexp.

Effects: fail.

### `Kernel.sin`

```gizlang
function Kernel.sin(self: mut Kernel, x: i64): i64
```

sin x (Cephes sinf: x reduced by multiples of pi/4 in three parts, then the sine or cosine polynomial), within
2 ulp for |x| below 8192.

Effects: fail.

### `Kernel.log`

```gizlang
function Kernel.log(self: mut Kernel, x: i64): i64
```

ln x within 2 ulp (Cephes logf), as tensor's vlog.

Effects: fail.

### `Kernel.gelu`

```gizlang
function Kernel.gelu(self: mut Kernel, x: i64): i64
```

GELU's tanh form, as tensor's vgelu.

Effects: fail.

### `Kernel.geluGrad`

```gizlang
function Kernel.geluGrad(self: mut Kernel, x: i64, dy: i64): i64
```

dy times gelu's derivative at x, as tensor's vgeluGrad.

Effects: fail.

### `Kernel.erf`

```gizlang
function Kernel.erf(self: mut Kernel, x0: i64): i64
```

erf x, as tensor's verf.

Effects: fail.

### `Kernel.geluErf`

```gizlang
function Kernel.geluErf(self: mut Kernel, x: i64): i64
```

GELU's erf form, as tensor's vgeluErf.

Effects: fail.

### `Kernel.geluErfGrad`

```gizlang
function Kernel.geluErfGrad(self: mut Kernel, x: i64, dy: i64): i64
```

dy times its derivative, as tensor's vgeluErfGrad.

Effects: fail.

### `Kernel.exp2`

```gizlang
function Kernel.exp2(self: mut Kernel, x: i64): i64
```

2^x by the GPU's approximation (NVIDIA ex2.approx within 2 ulp, subnormal results 0; AMD v_exp_f32): a few
times faster than exp, and not the CPU's bits. For softmax weights, with log2(e) folded into the scale.

Effects: fail.

### `Kernel.sigmoid`

```gizlang
function Kernel.sigmoid(self: mut Kernel, x: i64): i64
```

The sigmoid 1 / (1 + e^-x) of the F32 X, without overflow.

Effects: fail.

### `Kernel.silu`

```gizlang
function Kernel.silu(self: mut Kernel, x: i64): i64
```

x sigmoid(x) (SiLU, swish).

Effects: fail.

### `Kernel.tanh`

```gizlang
function Kernel.tanh(self: mut Kernel, x: i64): i64
```

The hyperbolic tangent of the F32 X.

Effects: fail.

### `Kernel.llvm`

```gizlang
function Kernel.llvm(self: Kernel, arch: string): string
```

The kernel as LLVM IR for AMD `arch` (gfx1151): the matrix operations take the layout of its generation.

Effects: fail.

### `Kernel.ptx`

```gizlang
function Kernel.ptx(self: Kernel, arch: string): string
```

The kernel as PTX for `arch` (sm_XY).

Effects: fail.

### `Kernel.msl`

```gizlang
function Kernel.msl(self: Kernel, arch: string): string
```

The kernel as Metal Shading Language for an Apple GPU `arch` ("apple7" to "apple10", the GPU family). Compiled at
run time with fast math off: the arithmetic is IEEE and a result equals the CPU's, except exp2, sums of mma and
the F32 division and square root, whose rounding Apple's compiler states and the Mac tests check. An F64 is a
double-float (msDouble).

Effects: fail.

## Constants

### `F32`

```gizlang
const F32 = 0
```

The value type of 32-bit floats.

### `F64`

```gizlang
const F64 = 1
```

The value type of 64-bit floats.

### `I32`

```gizlang
const I32 = 2
```

The value type of 32-bit integers.

### `I64`

```gizlang
const I64 = 3
```

The value type of 64-bit integers.

### `BOOL`

```gizlang
const BOOL = 4
```

The value type of booleans.

### `TILE`

```gizlang
const TILE = 5
```

The value type of a 16x16 F16 or BF16 operand tile of matrix multiply-accumulate (see mma).

### `ACC`

```gizlang
const ACC = 6
```

The value type of a 16x16 F32 accumulator tile of matrix multiply-accumulate (see mma).

### `E_F32`

```gizlang
const E_F32 = 0
```

The element type of 32-bit floats.

### `E_F64`

```gizlang
const E_F64 = 1
```

The element type of 64-bit floats.

### `E_I32`

```gizlang
const E_I32 = 2
```

The element type of 32-bit integers.

### `E_I64`

```gizlang
const E_I64 = 3
```

The element type of 64-bit integers.

### `E_F16`

```gizlang
const E_F16 = 4
```

The element type of IEEE half-precision floats; it loads as F32.

### `E_BF16`

```gizlang
const E_BF16 = 5
```

The element type of bfloat16 floats; it loads as F32.

### `E_U8`

```gizlang
const E_U8 = 6
```

The element type of unsigned bytes; it loads as I32.

### `E_I8`

```gizlang
const E_I8 = 7
```

The element type of signed bytes; it loads as I32.

### `E_U16`

```gizlang
const E_U16 = 8
```

The element type of unsigned 16-bit integers; it loads as I32.

### `E_E4M3`

```gizlang
const E_E4M3 = 9
```

FP8 E4M3 (OCP: bias 7, no infinities, 0x7F the NaN, largest 448): loads as F32; an F32 stores rounded to nearest
even and saturated to ±448 (NaN stays NaN), an I32 stores its low 8 bits unchanged. E_E4M3X2 is a pair of them as
one 16-bit element of raw bits (I32, as U16): the operand of mma on NVIDIA (hasFp8Mma).

### `E_E4M3X2`

```gizlang
const E_E4M3X2 = 10
```

A pair of FP8 E4M3 values as one 16-bit element of raw bits: the operand of mma on NVIDIA (hasFp8Mma).

### `E_E2M1X4`

```gizlang
const E_E2M1X4 = 11
```

Four FP4 E2M1 values (1 sign, 2 exponent and 1 mantissa bits: 0, 0.5, 1, 1.5, 2, 3, 4, 6), two a byte with the
first in the low half, as one 16-bit element of raw bits (I32, as U16): the operand of mmaScaled (hasFp4Mma).

### `E_I8X2`

```gizlang
const E_I8X2 = 12
```

A pair of signed 8-bit integers as one 16-bit element of raw bits (I32, as U16): the operand of mma on NVIDIA
(hasI8Mma), whose accumulator holds I32 sums (mmaGet with `integer`).

### `E_I4X4`

```gizlang
const E_I4X4 = 13
```

Four signed 4-bit integers, two a byte with the first in the low half, as one 16-bit element of raw bits (I32, as
U16): the operand of mma on NVIDIA (hasI4Mma, m16n8k64), whose accumulator holds I32 sums as I8X2's.

### `STATIC_SHARED_BYTES`

```gizlang
const STATIC_SHARED_BYTES = 49152
```

The workgroup memory that a kernel declares statically on NVIDIA; more is dynamic, opted in at launch.

### `APPLE_SHARED_BYTES`

```gizlang
const APPLE_SHARED_BYTES = 32768
```

Workgroup memory a threadgroup may use on an Apple GPU (32 KiB on every generation shipped).