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:

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

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

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:

result = enigma.where(false_val, true_val, condition)

Example — clamp-free maximum:

bigger = enigma.where(b, a, enigma.cmp_gt(a, b))

Vector operations

Constructors

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

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:

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.