🧠 D: DEFINE — Problem & Opportunity

TensorOps.gelu is the tanh approximation everywhere: DefaultCpuOps computes 0.5·x·(1 + tanh(√(2/π)·(x + 0.044715·x³))) and the HLO converter lowers gelu "using approximation". There is no erf op and no exact GELU.

Encoder checkpoints trained with hidden_act = "gelu" (ModernBERT, mmBERT, BERT-class models from Hugging Face, the Laya decision models) use the exact 0.5·x·(1 + erf(x/√2)). Measured on the Laya multilingual checkpoint while building skainet-transformers-inference-laya (transformers #487): with the engine's tanh GELU the decision-head outputs move by about 7.5e-4 against the reference; the parity gate for that model is 5e-5. The adapter therefore ships a workaround, LayaGelu, which copies the activation to a FloatArray, applies erf on the host and copies back — correct, but a host round trip per FFN per layer (22 layers × 2 per decision) on the eager path, outside every kernel tier and invisible to the compiled path.

🔍 A: ASSESS — Options

1. gelu(approximate = false) + erf in TensorOps (proposed): a boolean on the existing op plus a standalone erf, default unchanged (tanh), so no existing model changes numerically.
2. A separate geluExact op without erf: smaller surface, but erf is useful on its own (normal CDF, GELU variants) and is what StableHLO/CHLO expose (chlo.erf).
3. Keep host-side workarounds in model modules: what Laya does today; every BERT-class consumer re-implements it and the compiled path cannot see it.

✅ R: RESOLVE — Decision

Option 1. Scope:

• TensorOps: fun erf(x), fun gelu(x, approximate: Boolean = true); VoidTensorOps, RecordingExecution, ComputeGraphExecutor follow.
• CPU: scalar kernel (kotlin.math has no erf; use the Abramowitz-Stegun 7.1.26 form with max error 1.5e-7, or the erf of the FFM/libm tier where available) and a Vector API variant; JNI/FFM tiers optional in the first PR.
• Lowering: the node carries the approximate flag as op metadata; the StableHLO emitter chooses between a decomposed form (chlo.erf) and a stablehlo.composite through OpGranularityPolicy — see Lowering metadata below.
• NN DSL: GELU(approximate = …) module parameter; docs operator/kernel matrix row.
• Tests: erf against a table (|err| < 2e-7 in [-6, 6]), exact-vs-tanh GELU difference (max ≈ 2e-3 near |x| ≈ 2), parity of GELU(approximate = false) with the host formula.

🧩 Lowering metadata — one op, exact on every path

The point of a native op is not the CPU kernel alone: gelu(approximate = false) and erf must survive direct execution, the recorded/optimized graph and every export with their semantics intact, and carry enough metadata that a lowering transformation can pick the best legal form for its target without re-deriving what the op means.

What travels with the op

• Graph node: Operation.parameters = { "approximate": Boolean } on gelu; erf is its own op name. Operation.parameters is already serialized by BaseOperation.serialize() and the JSON export, so the flag reaches every emitter for free.

• Recording path gap (must be fixed in the same PR): RecordingExecution.gelu delegates to base.gelu without recording a node (skainet-compile-core/.../tape/RecordingExecution.kt:443, the "unrecorded/less critical ops" block). A compiled graph therefore never sees GELU today — the activation is computed eagerly and vanishes. The PR records it like PowOperation(parameters = mapOf("scalar_exponent" to n)), and ComputeGraphExecutor dispatches "gelu" -> ops.gelu(x, approximate = node.parameters["approximate"]), "erf" -> ops.erf(x). DefaultExecutionTape.geluBackward already receives attributes and has both derivative formulas in comments; it picks one from the flag.

• Lowering descriptor (new, declarative): one small record per fused-able op, registered next to its converter:

  LoweringSpec(
      op = "gelu",
      compositeName = "skainet.gelu",            // namespaced, as the StableHLO spec requires
      attributes = listOf("approximate"),        // copied into composite_attributes
      version = 1,
      decomposition = GeluDecomposition,         // the one "composite constructor"
  )

  GeluDecomposition is a single Kotlin builder that emits the formula as StableHLO primitives. The same builder writes the body of the composite's @decomposition function and the inline fallback, so eager semantics, composite and decomposed form cannot drift. For approximate = false the builder uses chlo.erf when the target toolchain accepts CHLO, otherwise the engine's own rational approximation with the same coefficients as the CPU kernel (Abramowitz–Stegun 7.1.26), which keeps DIRECT and lowered output within the 2e-7 kernel bound.

• Source annotation (docs): the operator matrix is generated by the docs KSP processor from @Backend/@InProgress. A @Lowering(composite = "skainet.gelu", decomposes = true) annotation on the TensorOps function (same package as @TensorOp, @DslOp) lets the matrix grow a Lowering column (composite · decompose · native) without hand edits, and gives tooling one place to read what an op can lower to.

Emission rules (StableHLO)

Checked against the StableHLO spec and ChloOps.td on 2026-10-14: the core op set has no erf (the alphabetical list jumps from dynamic_update_slice to exp2); chlo.erf, chlo.erfc, chlo.erf_inv live in CHLO and chlo-legalize-to-stablehlo provides their polynomial expansion; stablehlo.composite takes name, composite_attributes, decomposition, version and by definition "can be replaced with its decomposition without changing program semantics"; stablehlo-legalize-composite-to-call does exactly that.

A vendor compiler pattern-matches the composite name + attributes and maps it to its kernel; everyone else inlines or runs legalize-composite-to-call and gets the exact formula back. erf follows the same two rows with "skainet.erf".

This is also the first real consumer of OpGranularityPolicy: the policy is threaded into ConversionContext.granularity but no converter calls keepFused yet. The GELU PR adds the generic ConversionContext.emitComposite(spec, operands, attrs) helper (composite op + deduplicated decomposition function) that layernorm and scaled-dot-product attention can then reuse, which is what the policy KDoc already promises.

Degrade visibly, never silently

Today the exact→tanh substitution happens implicitly on every path, which is how the 7.5e-4 Laya drift stayed invisible. After this change a target that has no erf anywhere (no CHLO, no composite match) gets an explicit TargetOptimization pass in skainet-compile-opt that rewrites gelu(approximate = false) → gelu(approximate = true) and records it on the node (parameters["legalizedFrom"] = "gelu(approximate=false)", parameters["maxAbsError"] ≈ 2e-3). The export report lists these rewrites, so a numerics regression is a visible legalization decision instead of a surprise.

Other emitters

• C codegen: erff() from <math.h>; approximate = true keeps the current tanh form.
• Minerva: MinervaCompatibilityValidator.activationOperations is relu, sigmoid, tanh; add gelu/erf or reject with a remediation that names the legalization pass above.
• ONNX (if an exporter ever lands): Gelu (opset 20) has the same approximate attribute ("none" / "tanh"), Erf is a core op — the metadata maps 1:1, which is the test that the descriptor is target-neutral.

🚀 C: COMMIT — Acceptance

• TensorOps.erf and gelu(approximate = false) on the CPU path, scalar and Vector API
• RecordingExecution records gelu/erf with the approximate parameter; ComputeGraphExecutor and geluBackward dispatch on it; DIRECT vs OPTIMIZED on a Laya FFN block within 1e-6
• StableHLO: decomposed (chlo.erf) and composite ("skainet.gelu" / "skainet.erf", version = 1) emission selected by OpGranularityPolicy; ConversionContext.emitComposite helper; test runs the composite through stablehlo-legalize-composite-to-call and compares with the decomposed module
• One GeluDecomposition builder feeds both the composite body and the inline fallback (asserted by a golden-text test)
• Explicit exact→tanh legalization pass with recorded legalizedFrom/maxAbsError, surfaced in the export report
• @Lowering annotation and the Lowering column in the generated operator matrix
• GELU module exposes the flag; operator matrix documents it
• transformers: LayaGelu replaced by ctx.ops.gelu(x, approximate = false) with the 5e-5 parity gate still green (follow-up PR there, linked here)

Related: #1343 (SDPA key lengths — reuses the LoweringSpec / emitComposite seam defined here; together the first consumers of OpGranularityPolicy), SKaiNET-transformers#490.

Consumer evidence: transformers llm-inference/laya (LayaModules.kt, LayaGelu), measured 2026-10-12/13.