feat(compile): composite emission seam, GELU legalization pass, @Lowering (#1342, part 2)

CHANGELOG.md

@@ -2,6 +2,11 @@
 
 ## [Unreleased]
 
+### Added
+- **Composite emission behind `OpGranularityPolicy`** (#1342 part 2): a `LoweringSpec` per fused-able op (`skainet.gelu`, `skainet.erf`) and `ConversionContext.emitComposite` emit `stablehlo.composite "<name>" {composite_attributes, decomposition = @fn, version}` plus one private decomposition function per (op, attributes, type) when a target's policy keeps the op fused; otherwise the same builder emits the decomposition inline. Helper functions go to a module trailer after the entry function so existing text tooling keeps finding `@main` first. First consumer of `OpGranularityPolicy.keepFused`.
+- **`GeluApproximationLegalizationPass`** (`skainet-compile-opt`): targets without any erf lowering rewrite `gelu(approximate = false)` to the tanh form explicitly, recording `legalizedFrom` and `maxAbsError = 4.8e-4` on the node and a `GELU_APPROX_LEGALIZED` diagnostic, instead of a silent default. Registered per target through `TargetOptimizers.registerDagPasses`.
+- **`@Lowering` annotation** and a *Lowering* column in the generated operator coverage matrix (`composite · decompose · native`); `operators.json` carries the field.
+
 ### Added
 - **`TensorOps.erf` and exact GELU** (#1342): `gelu(tensor, approximate: Boolean = true)` keeps the tanh form as the default and selects the exact `0.5·x·(1 + erf(x/√2))` with `approximate = false`, the form BERT-class encoders and the Laya decision models were trained with (the two differ by up to 4.7e-4 near |x| ≈ 2.7). CPU: Abramowitz–Stegun 7.1.26 erf evaluated in double (|err| < 1.5e-7) with a JVM Vector API lane kernel; `GELU(approximate)` module, `network {}` `gelu(approximate)`, Java `TensorJavaOps.gelu(a, approximate)` / `erf(a)` and `SequentialModelBuilder.gelu(approximate)`; exact backward rule and `erfBackward`.
 - **GELU carries its form as op metadata**: the flag is recorded as `parameters["approximate"]` on the graph node (`GeluOperation`, `ErfOperation`), read by `ComputeGraphExecutor`, the backward rule and the StableHLO converter, so the compiled path cannot silently diverge from eager execution. StableHLO lowers `erf` to `chlo.erf` and exact GELU to `0.5·x·(1 + chlo.erf(x·√½))`; one `GeluLowering` builder emits both forms.

build-logic/convention/src/main/kotlin/GenerateDocumentationTask.kt

@@ -130,7 +130,7 @@ abstract class GenerateDocumentationTask : DefaultTask() {
             appendLine("")
             appendLine("Generated from `operators.json` version `${module.version}` on ${formatTimestamp(module.timestamp)}.")
             appendLine("")
-            appendLine("Rows are `Operator.function` pairs. The `Validated` column shows whether the function's documentation has been DARC-validated by a reviewer (see xref:contributing/darc-workflow.adoc[DARC workflow]). Remaining columns are backends that appear in any function's `statusByBackend` map — a missing entry means the backend makes no claim about the function (treat it as \"unknown\", not \"not supported\").")
+            appendLine("Rows are `Operator.function` pairs. The `Validated` column shows whether the function's documentation has been DARC-validated by a reviewer (see xref:contributing/darc-workflow.adoc[DARC workflow]). The `Lowering` column lists the compile-path forms the function declares with `@Lowering` (`composite` = a `stablehlo.composite` when a target keeps it fused, `decompose` = StableHLO/CHLO primitives, `native` = a 1:1 target op); `—` means undeclared. Remaining columns are backends that appear in any function's `statusByBackend` map — a missing entry means the backend makes no claim about the function (treat it as \"unknown\", not \"not supported\").")
             appendLine("")
             if (rows.isEmpty() || allBackends.isEmpty()) {
                 appendLine("NOTE: No backend status information found in the source data.")
@@ -139,17 +139,18 @@ abstract class GenerateDocumentationTask : DefaultTask() {
             }
 
             // Header: row label, validation column, then one column per backend.
-            val colSpec = (listOf("2", "1") + List(allBackends.size) { "1" }).joinToString(",")
+            val colSpec = (listOf("2", "1", "1") + List(allBackends.size) { "1" }).joinToString(",")
             appendLine("[cols=\"$colSpec\", options=\"header\"]")
             appendLine("|===")
-            append("| Operator.function | Validated ")
+            append("| Operator.function | Validated | Lowering ")
             allBackends.forEach { append("| $it ") }
             appendLine("")
             appendLine("")
 
             rows.forEach { row ->
                 append("| `${row.operator}.${row.function.name}` ")
                 append("| ${darcCell(row.function, row.hasPartial)} ")
+                append("| ${row.function.lowering.ifEmpty { "—" }} ")
                 allBackends.forEach { backend ->
                     val raw = row.function.statusByBackend[backend]
                     val cell = if (raw == null) "—" else shortStatus(raw)
@@ -158,10 +159,11 @@ abstract class GenerateDocumentationTask : DefaultTask() {
                 appendLine("")
             }
 
-            // Totals footer: validated count followed by per-backend done count.
+            // Totals footer: validated count, declared-lowering count, then per-backend done count.
             appendLine("")
             val validatedCount = rows.count { it.function.validated }
-            append("| *Done* | *$validatedCount / ${rows.size}* ")
+            val loweringCount = rows.count { it.function.lowering.isNotEmpty() }
+            append("| *Done* | *$validatedCount / ${rows.size}* | *$loweringCount / ${rows.size}* ")
             allBackends.forEach { backend ->
                 val n = rows.count { isDone(it.function.statusByBackend[backend]) }
                 append("| *$n / ${rows.size}* ")

build-logic/convention/src/main/kotlin/models/DocumentationModels.kt

@@ -35,6 +35,8 @@ data class FunctionDoc(
     val validatedOn: String = "",
     val validatedCommit: String = "",
     val referencesChecked: Boolean = true,
+    /** Lowering forms from `@Lowering`: `composite`, `decompose`, `native` joined by ` · `; empty = undeclared. */
+    val lowering: String = "",
 )
 
 @Serializable

docs/modules/ROOT/pages/reference/ops-status-matrix.adoc

@@ -3,80 +3,80 @@
 
 Generated from `operators.json` version `0.58.0` on 2026-10-14.
 
-Rows are `Operator.function` pairs. The `Validated` column shows whether the function's documentation has been DARC-validated by a reviewer (see xref:contributing/darc-workflow.adoc[DARC workflow]). Remaining columns are backends that appear in any function's `statusByBackend` map — a missing entry means the backend makes no claim about the function (treat it as "unknown", not "not supported").
+Rows are `Operator.function` pairs. The `Validated` column shows whether the function's documentation has been DARC-validated by a reviewer (see xref:contributing/darc-workflow.adoc[DARC workflow]). The `Lowering` column lists the compile-path forms the function declares with `@Lowering` (`composite` = a `stablehlo.composite` when a target keeps it fused, `decompose` = StableHLO/CHLO primitives, `native` = a 1:1 target op); `—` means undeclared. Remaining columns are backends that appear in any function's `statusByBackend` map — a missing entry means the backend makes no claim about the function (treat it as "unknown", not "not supported").
 
-[cols="2,1,1,1,1", options="header"]
+[cols="2,1,1,1,1,1", options="header"]
 |===
-| Operator.function | Validated | apple | cpu | wasm 
+| Operator.function | Validated | Lowering | apple | cpu | wasm 
 
-| `TensorOps.add` | ✖ | — | — | — 
-| `TensorOps.subtract` | ✖ | — | — | — 
-| `TensorOps.multiply` | ✖ | — | — | — 
-| `TensorOps.divide` | ✖ | — | — | — 
-| `TensorOps.addScalar` | ✖ | — | — | — 
-| `TensorOps.subScalar` | ✖ | — | — | — 
-| `TensorOps.mulScalar` | ✖ | — | — | — 
-| `TensorOps.divScalar` | ✖ | — | — | — 
-| `TensorOps.rsubScalar` | ✖ | — | — | — 
-| `TensorOps.rdivScalar` | ✖ | — | — | — 
-| `TensorOps.matmul` | ✅ | — | — | — 
-| `TensorOps.matmulWeightTransposed` | ✖ | — | — | — 
-| `TensorOps.relayoutPackedWeightForKernels` | ✖ | — | — | — 
-| `TensorOps.transpose` | ✖ | — | — | — 
-| `TensorOps.permute` | ✖ | — | — | — 
-| `TensorOps.conv1d` | ✖ | — | — | — 
-| `TensorOps.conv2d` | ✖ | — | — | — 
-| `TensorOps.conv3d` | ✖ | — | — | — 
-| `TensorOps.convTranspose1d` | ✖ | — | — | — 
-| `TensorOps.maxPool2d` | ✖ | — | — | — 
-| `TensorOps.avgPool2d` | ✖ | — | — | — 
-| `TensorOps.upsample2d` | ✖ | — | — | — 
-| `TensorOps.reshape` | ✖ | — | — | — 
-| `TensorOps.flatten` | ✖ | — | — | — 
-| `TensorOps.concat` | ✖ | — | — | — 
-| `TensorOps.split` | ✖ | — | — | — 
-| `TensorOps.squeeze` | ✖ | — | — | — 
-| `TensorOps.unsqueeze` | ✖ | — | — | — 
-| `TensorOps.relu` | ✖ | — | — | — 
-| `TensorOps.leakyRelu` | ✖ | — | — | — 
-| `TensorOps.elu` | ✖ | — | — | — 
-| `TensorOps.softmax` | ✖ | — | — | — 
-| `TensorOps.logSoftmax` | ✖ | — | — | — 
-| `TensorOps.sigmoid` | ✖ | — | — | — 
-| `TensorOps.tanh` | ✖ | — | — | — 
-| `TensorOps.silu` | ✖ | — | — | — 
-| `TensorOps.gelu` | ✖ | — | — | — 
-| `TensorOps.sum` | ✖ | — | — | — 
-| `TensorOps.mean` | ✖ | — | — | — 
-| `TensorOps.variance` | ✖ | — | — | — 
-| `TensorOps.argMax` | ⚠ | — | — | — 
-| `TensorOps.sqrt` | ✖ | — | — | — 
-| `TensorOps.pow` | ✖ | — | — | — 
-| `TensorOps.powScalar` | ✖ | — | — | — 
-| `TensorOps.log` | ✖ | — | — | — 
-| `TensorOps.log2` | ✖ | — | — | — 
-| `TensorOps.log10` | ✖ | — | — | — 
-| `TensorOps.abs` | ✖ | — | — | — 
-| `TensorOps.sign` | ✖ | — | — | — 
-| `TensorOps.clamp` | ✖ | — | — | — 
-| `TensorOps.narrow` | ✖ | — | — | — 
-| `TensorOps.pad2d` | ✖ | — | — | — 
-| `TensorOps.unfold` | ✖ | — | — | — 
-| `TensorOps.lt` | ✖ | — | — | — 
-| `TensorOps.ge` | ✖ | — | — | — 
-| `TensorOps.tril` | ✖ | — | — | — 
-| `TensorOps.convert` | ✖ | — | — | — 
-| `TensorOps.gather` | ✖ | — | — | — 
-| `TensorOps.indexSelect` | ✖ | — | — | — 
-| `TensorOps.exp` | ✖ | — | — | — 
-| `TensorOps.expm1` | ✖ | — | — | — 
-| `TensorOps.erf` | ✖ | — | — | — 
-| `TensorOps.sin` | ✖ | — | — | — 
-| `TensorOps.cos` | ✖ | — | — | — 
-| `TensorOps.scaledDotProductAttention` | ✖ | — | — | — 
-| `Similarity.cosineDistance` | ✖ | ✅ | ✅ | ✅ 
+| `TensorOps.add` | ✖ | — | — | — | — 
+| `TensorOps.subtract` | ✖ | — | — | — | — 
+| `TensorOps.multiply` | ✖ | — | — | — | — 
+| `TensorOps.divide` | ✖ | — | — | — | — 
+| `TensorOps.addScalar` | ✖ | — | — | — | — 
+| `TensorOps.subScalar` | ✖ | — | — | — | — 
+| `TensorOps.mulScalar` | ✖ | — | — | — | — 
+| `TensorOps.divScalar` | ✖ | — | — | — | — 
+| `TensorOps.rsubScalar` | ✖ | — | — | — | — 
+| `TensorOps.rdivScalar` | ✖ | — | — | — | — 
+| `TensorOps.matmul` | ✅ | — | — | — | — 
+| `TensorOps.matmulWeightTransposed` | ✖ | — | — | — | — 
+| `TensorOps.relayoutPackedWeightForKernels` | ✖ | — | — | — | — 
+| `TensorOps.transpose` | ✖ | — | — | — | — 
+| `TensorOps.permute` | ✖ | — | — | — | — 
+| `TensorOps.conv1d` | ✖ | — | — | — | — 
+| `TensorOps.conv2d` | ✖ | — | — | — | — 
+| `TensorOps.conv3d` | ✖ | — | — | — | — 
+| `TensorOps.convTranspose1d` | ✖ | — | — | — | — 
+| `TensorOps.maxPool2d` | ✖ | — | — | — | — 
+| `TensorOps.avgPool2d` | ✖ | — | — | — | — 
+| `TensorOps.upsample2d` | ✖ | — | — | — | — 
+| `TensorOps.reshape` | ✖ | — | — | — | — 
+| `TensorOps.flatten` | ✖ | — | — | — | — 
+| `TensorOps.concat` | ✖ | — | — | — | — 
+| `TensorOps.split` | ✖ | — | — | — | — 
+| `TensorOps.squeeze` | ✖ | — | — | — | — 
+| `TensorOps.unsqueeze` | ✖ | — | — | — | — 
+| `TensorOps.relu` | ✖ | — | — | — | — 
+| `TensorOps.leakyRelu` | ✖ | — | — | — | — 
+| `TensorOps.elu` | ✖ | — | — | — | — 
+| `TensorOps.softmax` | ✖ | — | — | — | — 
+| `TensorOps.logSoftmax` | ✖ | — | — | — | — 
+| `TensorOps.sigmoid` | ✖ | — | — | — | — 
+| `TensorOps.tanh` | ✖ | — | — | — | — 
+| `TensorOps.silu` | ✖ | — | — | — | — 
+| `TensorOps.gelu` | ✖ | composite · decompose | — | — | — 
+| `TensorOps.sum` | ✖ | — | — | — | — 
+| `TensorOps.mean` | ✖ | — | — | — | — 
+| `TensorOps.variance` | ✖ | — | — | — | — 
+| `TensorOps.argMax` | ⚠ | — | — | — | — 
+| `TensorOps.sqrt` | ✖ | — | — | — | — 
+| `TensorOps.pow` | ✖ | — | — | — | — 
+| `TensorOps.powScalar` | ✖ | — | — | — | — 
+| `TensorOps.log` | ✖ | — | — | — | — 
+| `TensorOps.log2` | ✖ | — | — | — | — 
+| `TensorOps.log10` | ✖ | — | — | — | — 
+| `TensorOps.abs` | ✖ | — | — | — | — 
+| `TensorOps.sign` | ✖ | — | — | — | — 
+| `TensorOps.clamp` | ✖ | — | — | — | — 
+| `TensorOps.narrow` | ✖ | — | — | — | — 
+| `TensorOps.pad2d` | ✖ | — | — | — | — 
+| `TensorOps.unfold` | ✖ | — | — | — | — 
+| `TensorOps.lt` | ✖ | — | — | — | — 
+| `TensorOps.ge` | ✖ | — | — | — | — 
+| `TensorOps.tril` | ✖ | — | — | — | — 
+| `TensorOps.convert` | ✖ | — | — | — | — 
+| `TensorOps.gather` | ✖ | — | — | — | — 
+| `TensorOps.indexSelect` | ✖ | — | — | — | — 
+| `TensorOps.exp` | ✖ | — | — | — | — 
+| `TensorOps.expm1` | ✖ | — | — | — | — 
+| `TensorOps.erf` | ✖ | composite · decompose | — | — | — 
+| `TensorOps.sin` | ✖ | — | — | — | — 
+| `TensorOps.cos` | ✖ | — | — | — | — 
+| `TensorOps.scaledDotProductAttention` | ✖ | — | — | — | — 
+| `Similarity.cosineDistance` | ✖ | — | ✅ | ✅ | ✅ 
 
-| *Done* | *1 / 66* | *1 / 66* | *1 / 66* | *1 / 66* 
+| *Done* | *1 / 66* | *2 / 66* | *1 / 66* | *1 / 66* | *1 / 66* 
 |===
 
 Per-function detail including notes lives in xref:reference/operators/generated/index.adoc[Operator reference].

docs/modules/ROOT/pages/tutorials/hlo-getting-started.adoc

@@ -325,6 +325,32 @@ class RgbEnhanceOp : HloConverter {
 }
 ----
 
+==== Fused ops and `stablehlo.composite`
+
+Some ops have a `LoweringSpec` (`sk.ainet.compile.hlo.lowering`): a namespaced composite name, the node parameters that become `composite_attributes`, a version, and one decomposition builder. Whether a target sees the composite or the primitives is decided per op by the target's `OpGranularityPolicy`:
+
+[source,kotlin]
+----
+// A target that has a fused GELU kernel keeps the op; everything else decomposes.
+val policy = FusedOpAllowList("my-npu", setOf("gelu"))
+val module = toStableHlo(graph, "main", target = "my-npu", granularity = policy)
+----
+
+[source,mlir]
+----
+%v8 = stablehlo.composite "skainet.gelu" %v7 {composite_attributes = {approximate = false},
+        decomposition = @skainet_gelu_approximate_false_128x768xf32, version = 1 : i32}
+        : (tensor<128x768xf32>) -> tensor<128x768xf32>
+...
+func.func private @skainet_gelu_approximate_false_128x768xf32(%arg0: tensor<128x768xf32>) -> tensor<128x768xf32> {
+  // the same lines the portable emission inlines
+}
+----
+
+The private decomposition functions follow the entry function in the module. A composite is replaceable by its decomposition by definition, so a vendor compiler may pattern-match `skainet.gelu` while `stablehlo-legalize-composite-to-call` (or plain inlining) restores the exact formula everywhere else. Specs today: `GeluLowering.spec` (`skainet.gelu`, attribute `approximate`) and `ErfLowering.spec` (`skainet.erf`).
+
+Targets that cannot lower `erf` at all register `GeluApproximationLegalizationPass` through `TargetOptimizers.registerDagPasses`; it rewrites `gelu(approximate = false)` to the tanh form and records `legalizedFrom` / `maxAbsError` on the node, so the substitution shows up in the pass diagnostics instead of happening silently.
+
 ==== Debugging HLO
 
 Use SKaiNET's built-in debugging tools: