> ## Documentation Index
> Fetch the complete documentation index at: https://klyne-research.mintlify.site/llms.txt
> Use this file to discover all available pages before exploring further.

# Operations

> Arithmetic, math, vector, comparison, and selection operations available in kernel bodies.

This page covers the scalar and vector operations available inside `@enigma.kernel` bodies. All operations work on `IRValue` objects and are lowered to their Metal Shading Language equivalents.

## Arithmetic

Standard Python operators work on `IRValue` objects and produce new `IRValue` nodes:

```python theme={null}
c = a + b     # add
c = a - b     # subtract
c = a * b     # multiply
c = a / b     # divide
c = a % b     # modulo
```

Integer literals are auto-promoted to `u32`. Float literals become `f32`.

## Float math intrinsics

### Unary

| Function          | Metal   | Description            |
| ----------------- | ------- | ---------------------- |
| `enigma.sqrt(x)`  | `sqrt`  | Square root            |
| `enigma.rsqrt(x)` | `rsqrt` | Reciprocal square root |
| `enigma.abs(x)`   | `abs`   | Absolute value         |
| `enigma.ceil(x)`  | `ceil`  | Round up               |
| `enigma.floor(x)` | `floor` | Round down             |
| `enigma.round(x)` | `round` | Round to nearest       |
| `enigma.trunc(x)` | `trunc` | Truncate toward zero   |
| `enigma.sin(x)`   | `sin`   | Sine                   |
| `enigma.cos(x)`   | `cos`   | Cosine                 |
| `enigma.tan(x)`   | `tan`   | Tangent                |
| `enigma.asin(x)`  | `asin`  | Arc sine               |
| `enigma.acos(x)`  | `acos`  | Arc cosine             |
| `enigma.atan(x)`  | `atan`  | Arc tangent            |
| `enigma.exp(x)`   | `exp`   | Natural exponent       |
| `enigma.exp2(x)`  | `exp2`  | Base-2 exponent        |
| `enigma.log(x)`   | `log`   | Natural logarithm      |
| `enigma.log2(x)`  | `log2`  | Base-2 logarithm       |
| `enigma.sign(x)`  | `sign`  | Sign (−1, 0, 1)        |
| `enigma.fract(x)` | `fract` | Fractional part        |

### Binary

| Function                | Metal      | Description            |
| ----------------------- | ---------- | ---------------------- |
| `enigma.fmin(x, y)`     | `fmin`     | Float minimum          |
| `enigma.fmax(x, y)`     | `fmax`     | Float maximum          |
| `enigma.pow(x, y)`      | `pow`      | Power                  |
| `enigma.atan2(y, x)`    | `atan2`    | Arc tangent of y/x     |
| `enigma.step(edge, x)`  | `step`     | 0 if x \< edge, else 1 |
| `enigma.copysign(x, y)` | `copysign` | x with sign of y       |

### Ternary

| Function                       | Metal        | Description                        |
| ------------------------------ | ------------ | ---------------------------------- |
| `enigma.fma(a, b, c)`          | `fma`        | Fused multiply-add: a×b + c        |
| `enigma.clamp(x, lo, hi)`      | `clamp`      | Clamp to \[lo, hi]                 |
| `enigma.mix(x, y, a)`          | `mix`        | Linear interpolation: x + a\*(y-x) |
| `enigma.smoothstep(e0, e1, x)` | `smoothstep` | Smooth step function               |

## Integer math

| Function                                        | Description                  |
| ----------------------------------------------- | ---------------------------- |
| `enigma.imin(x, y)` / `enigma.imax(x, y)`       | Signed min/max               |
| `enigma.iclamp(x, lo, hi)`                      | Integer clamp                |
| `enigma.abs_diff(x, y)`                         | Absolute difference          |
| `enigma.add_sat(x, y)` / `enigma.sub_sat(x, y)` | Saturating add/subtract      |
| `enigma.mul_hi(x, y)`                           | High half of 32-bit multiply |
| `enigma.mad_sat(a, b, c)`                       | Saturating multiply-add      |
| `enigma.rotate(x, n)`                           | Bit rotate                   |

## Integer bit operations

| Function                                  | Metal          | Description          |
| ----------------------------------------- | -------------- | -------------------- |
| `enigma.popcount(x)`                      | `popcount`     | Count set bits       |
| `enigma.clz(x)`                           | `clz`          | Count leading zeros  |
| `enigma.ctz(x)`                           | `ctz`          | Count trailing zeros |
| `enigma.reverse_bits(x)`                  | `reverse_bits` | Reverse bit pattern  |
| `enigma.extract_bits(x, offset, count)`   | `extract_bits` | Extract bit field    |
| `enigma.insert_bits(x, y, offset, count)` | `insert_bits`  | Insert bit field     |

## Predicates

```python theme={null}
enigma.isnan(x)      # True if x is NaN
enigma.isinf(x)      # True if x is infinite
enigma.isfinite(x)   # True if x is finite
enigma.signbit(x)    # True if x is negative
enigma.isnormal(x)   # True if x is normal (not zero, subnormal, inf, or NaN)
```

## Comparisons

All comparisons return an `IRValue` of type `b1` (boolean):

```python theme={null}
enigma.cmp_eq(a, b)   # a == b
enigma.cmp_ne(a, b)   # a != b
enigma.cmp_lt(a, b)   # a < b  (signed)
enigma.cmp_le(a, b)   # a <= b
enigma.cmp_gt(a, b)   # a > b
enigma.cmp_ge(a, b)   # a >= b
```

## Selection

### `enigma.where`

Ternary select — returns `true_val` when `condition` is true, `false_val` otherwise:

```python theme={null}
result = enigma.where(false_val, true_val, condition)
```

Example — clamp-free maximum:

```python theme={null}
bigger = enigma.where(b, a, enigma.cmp_gt(a, b))
```

## Vector operations

### Constructors

```python theme={null}
v2 = enigma.make_float2(x, y)
v3 = enigma.make_float3(x, y, z)
v4 = enigma.make_float4(x, y, z, w)
v  = enigma.make_vec(x, y, z, w)   # generic, infers type from components
```

### Component access

```python theme={null}
px = v.x
py = v.y
pz = v.z
pw = v.w
lane = enigma.vec_extract(v, 2)   # component by index
```

### Geometry

| Function                         | Description                 |
| -------------------------------- | --------------------------- |
| `enigma.dot(a, b)`               | Dot product                 |
| `enigma.length(v)`               | Euclidean length            |
| `enigma.distance(a, b)`          | Distance between two points |
| `enigma.normalize(v)`            | Unit vector                 |
| `enigma.cross(a, b)`             | Cross product (float3 only) |
| `enigma.reflect(v, n)`           | Reflection around normal n  |
| `enigma.refract(v, n, eta)`      | Refraction                  |
| `enigma.faceforward(n, i, nref)` | Face-forward normal         |

## Simdgroup matrix operations

For 8×8 matrix multiply-accumulate on Apple Silicon:

```python theme={null}
matA = enigma.simdgroup_matrix_load(buf_A, elements_per_row=8)
matB = enigma.simdgroup_matrix_load(buf_B, elements_per_row=8)
matC = enigma.make_filled_simdgroup_matrix(0.0)
matC = enigma.simdgroup_multiply_accumulate(matA, matB, matC)
enigma.simdgroup_matrix_store(matC, buf_C, elements_per_row=8)
```

Requires `caps.supports_simdgroup_matrix` to be `True`. See [Device Capabilities](/concepts/memory-model#device-capabilities).