[Docs] Follow TileOPs' release renames, and drop register_kernel_type from the backend guides

docs/api/attention.md

@@ -7,15 +7,15 @@ what runs when you call `op(...)`.
 
 ## Multi-Head Attention
 
-::: tileops.attention.MHADecodePagedWithKVCacheFwdOp
+::: tileops.attention.MHAPagedFwdOp
     options:
       show_root_heading: true
       heading_level: 3
       members: ["__init__", "forward"]
 
 ## Grouped-Query Attention
 
-::: tileops.attention.GQABwdOp
+::: tileops.attention.GQADenseBwdOp
     options:
       show_root_heading: true
       heading_level: 3
@@ -41,7 +41,7 @@ what runs when you call `op(...)`.
 
 ## Multi-Head Latent Attention
 
-::: tileops.attention.MLADecodeWithKVCacheFwdOp
+::: tileops.attention.MLADenseDecodeFwdOp
     options:
       show_root_heading: true
       heading_level: 3
@@ -93,7 +93,7 @@ what runs when you call `op(...)`.
 
 ## DeepSeek Sparse Attention
 
-::: tileops.attention.DSADecodeWithKVCacheFwdOp
+::: tileops.attention.DSADenseFwdOp
     options:
       show_root_heading: true
       heading_level: 3

docs/api/elementwise.md

@@ -179,37 +179,37 @@ what runs when you call `op(...)`.
 
 ## Activations
 
-::: tileops.elementwise.ReluFwdOp
+::: tileops.elementwise.ReLUFwdOp
     options:
       show_root_heading: true
       heading_level: 3
       members: ["__init__", "forward"]
 
-::: tileops.elementwise.LeakyReluFwdOp
+::: tileops.elementwise.LeakyReLUFwdOp
     options:
       show_root_heading: true
       heading_level: 3
       members: ["__init__", "forward"]
 
-::: tileops.elementwise.EluFwdOp
+::: tileops.elementwise.ELUFwdOp
     options:
       show_root_heading: true
       heading_level: 3
       members: ["__init__", "forward"]
 
-::: tileops.elementwise.SeluFwdOp
+::: tileops.elementwise.SELUFwdOp
     options:
       show_root_heading: true
       heading_level: 3
       members: ["__init__", "forward"]
 
-::: tileops.elementwise.GeluFwdOp
+::: tileops.elementwise.GELUFwdOp
     options:
       show_root_heading: true
       heading_level: 3
       members: ["__init__", "forward"]
 
-::: tileops.elementwise.SiluFwdOp
+::: tileops.elementwise.SiLUFwdOp
     options:
       show_root_heading: true
       heading_level: 3
@@ -257,7 +257,7 @@ what runs when you call `op(...)`.
       heading_level: 3
       members: ["__init__", "forward"]
 
-::: tileops.elementwise.PreluFwdOp
+::: tileops.elementwise.PReLUFwdOp
     options:
       show_root_heading: true
       heading_level: 3
@@ -271,19 +271,19 @@ what runs when you call `op(...)`.
 
 ## Gated activations
 
-::: tileops.elementwise.SiluAndMulFwdOp
+::: tileops.elementwise.SiLUAndMulFwdOp
     options:
       show_root_heading: true
       heading_level: 3
       members: ["__init__", "forward"]
 
-::: tileops.elementwise.GeluAndMulFwdOp
+::: tileops.elementwise.GELUAndMulFwdOp
     options:
       show_root_heading: true
       heading_level: 3
       members: ["__init__", "forward"]
 
-::: tileops.elementwise.GeluTanhAndMulFwdOp
+::: tileops.elementwise.GELUTanhAndMulFwdOp
     options:
       show_root_heading: true
       heading_level: 3

docs/api/index.md

@@ -7,9 +7,9 @@ dtypes — and the call takes the tensors. Both appear under each op as `__init_
 
 ```python
 import torch
-from tileops.gemm import GemmFwdOp
+from tileops.gemm import GEMMFwdOp
 
-op = GemmFwdOp()                        # construct once, reuse
+op = GEMMFwdOp()                        # construct once, reuse
 d = op(a, b)                         # the specialized kernel is built on first call
 ```
 
@@ -51,4 +51,4 @@ Two things this reference does not carry:
 
 These pages are generated from the docstrings in TileOPs, so an op whose docstring is
 thin reads thin here. The fix belongs upstream, in
-[`src/tileops/ops/`](https://github.com/tile-ai/TileOPs/tree/main/src/tileops/ops).
+[`src/tileops/_ops/`](https://github.com/tile-ai/TileOPs/tree/main/src/tileops/_ops).

docs/api/linear-algebra.md

@@ -7,39 +7,39 @@ what runs when you call `op(...)`.
 
 ## Dense GEMM
 
-::: tileops.gemm.GemmFwdOp
+::: tileops.gemm.GEMMFwdOp
     options:
       show_root_heading: true
       heading_level: 3
       members: ["__init__", "forward"]
       merge_init_into_class: false
       show_signature_annotations: false
 
-::: tileops.gemm.GemmFP8FwdOp
+::: tileops.gemm.GEMMFP8FwdOp
     options:
       show_root_heading: true
       heading_level: 3
       members: ["__init__", "forward"]
       merge_init_into_class: false
       show_signature_annotations: false
 
-::: tileops.gemm.GemmW4A16FwdOp
+::: tileops.gemm.GEMMW4A16FwdOp
     options:
       show_root_heading: true
       heading_level: 3
       members: ["__init__", "forward"]
       merge_init_into_class: false
       show_signature_annotations: false
 
-::: tileops.gemm.GemmW4A8FwdOp
+::: tileops.gemm.GEMMW4A8FwdOp
     options:
       show_root_heading: true
       heading_level: 3
       members: ["__init__", "forward"]
       merge_init_into_class: false
       show_signature_annotations: false
 
-::: tileops.gemm.GemmINT8W8A8FwdOp
+::: tileops.gemm.GEMMINT8W8A8FwdOp
     options:
       show_root_heading: true
       heading_level: 3
@@ -49,15 +49,15 @@ what runs when you call `op(...)`.
 
 ## Batched GEMM
 
-::: tileops.gemm.BmmFwdOp
+::: tileops.gemm.BMMFwdOp
     options:
       show_root_heading: true
       heading_level: 3
       members: ["__init__", "forward"]
       merge_init_into_class: false
       show_signature_annotations: false
 
-::: tileops.gemm.BmmFP8FwdOp
+::: tileops.gemm.BMMFP8FwdOp
     options:
       show_root_heading: true
       heading_level: 3
@@ -67,7 +67,7 @@ what runs when you call `op(...)`.
 
 ## Grouped GEMM
 
-::: tileops.gemm.GroupedGemmFwdOp
+::: tileops.gemm.GroupedGEMMFwdOp
     options:
       show_root_heading: true
       heading_level: 3

docs/api/moe.md

@@ -8,7 +8,7 @@ what runs when you call `op(...)`.
 A routed mixture-of-experts layer is available two ways here. `FusedMoEFwdOp` runs
 the whole FFN, and `FusedMoESharedExpertFwdOp` adds a shared expert beside the routed ones. The rest are its stages, callable on their own: the op that picks
 each token's experts, the ops that move tokens into an expert-contiguous layout and
-back, and the expert GEMMs that run on it. `MoEGroupedGemmFwdOp` is one grouped
+back, and the expert GEMMs that run on it. `MoEGroupedGEMMFwdOp` is one grouped
 GEMM; `MoEExpertMLPFwdOp` is the pair of them with the gated activation fused into
 the first; `FusedMoEExpertsFwdOp` is that MLP with the permutes around it, on the
 tight (no-pad) layout, and `IndexedExpertMLPFwdOp` is the backend it picks instead when
@@ -56,13 +56,13 @@ the routes are few enough to read the weights once per route rather than once pe
 
 ## Expert GEMMs
 
-::: tileops.moe.MoEGroupedGemmFwdOp
+::: tileops.moe.MoEGroupedGEMMFwdOp
     options:
       show_root_heading: true
       heading_level: 3
       members: ["__init__", "forward"]
 
-::: tileops.moe.MoEGroupedGemmFP8FwdOp
+::: tileops.moe.MoEGroupedGEMMFP8FwdOp
     options:
       show_root_heading: true
       heading_level: 3

docs/api/quantization.md

@@ -19,7 +19,7 @@ what runs when you call `op(...)`.
       heading_level: 3
       members: ["__init__", "forward"]
 
-::: tileops.quantization.INT8QuantPerBlockFwdOp
+::: tileops.quantization.INT8QuantPerTokenGroupFwdOp
     options:
       show_root_heading: true
       heading_level: 3
@@ -33,7 +33,7 @@ what runs when you call `op(...)`.
 
 ## FP8 quantization
 
-::: tileops.quantization.FP8QuantFwdOp
+::: tileops.quantization.FP8QuantPerTokenGroupFwdOp
     options:
       show_root_heading: true
       heading_level: 3

docs/api/trace.md

@@ -6,7 +6,7 @@ wire the markers into a `@tilelang.jit` builder, run, and export a timeline.
 For a walkthrough, see the
 [In-Kernel Timeline Trace guide](../performance-guides/trace-timeline.md).
 
-::: tileops.trace.api._Trace
+::: tileops.trace._api._Trace
     options:
       show_root_heading: false
       show_root_toc_entry: false

docs/backends.md

@@ -31,20 +31,13 @@ The second half explains why the protocol is designed this way:
 - the interfaces available to a caller;
 - what the protocol deliberately does not support.
 
-## Two ways to extend dispatch {#two-ways}
+## How a backend extends dispatch {#extend}
 
-A package outside TileOPs picks one of two mechanisms, by how much of an op it takes
-over:
-
-1. `register_kernel_type`: adds an implementation to a kernel interface, which takes
-   part in selection alongside the in-tree implementations;
-1. target: takes over every call of the op.
-
-In the first, the kernel class follows the kernel interface, the same contract the
-in-tree implementations follow; how to write one is in
-[How a backend joins TileOPs](user-guide/dispatch/backends.md). A target follows the op's
-manifest signature instead, and the backend writes a `build_kernel` for it. The rest of
-this page covers the target only.
+A package outside TileOPs takes over every call of an op through a target. A target
+follows the op's manifest signature, and the backend writes a `build_kernel` for it. The
+kernel interfaces and call specs of TileOPs' ops are not public API, so a backend does not
+add an implementation to a single kernel interface; see
+[How a backend joins TileOPs](user-guide/dispatch/backends.md).
 
 ## Four things to write
 
@@ -126,10 +119,10 @@ this call arrives as `None`.
 
 ```python
 # GroupNormFwdOp's spec: weight and bias are optional, and arrive as None when absent
-def build_group_norm(x, weight, bias, *, num_groups, eps):
+def build_group_norm(input, weight, bias, *, num_groups, eps):
     if weight is None:                                   # presence read off the slot
-        return AcmeGroupNorm(num_groups, eps, x.dtype)
-    return AcmeGroupNormAffine(num_groups, eps, x.dtype)
+        return AcmeGroupNorm(num_groups, eps, input.dtype)
+    return AcmeGroupNormAffine(num_groups, eps, input.dtype)
 ```
 
 The op layer calls `build_kernel` with the call's device made current through
@@ -187,8 +180,8 @@ call's device: `arch`, `calibration`, `sm_count` and `smem_budget`.
 | Another registered type | The registered `read`, which takes the call's `torch.device` and returns `DeviceFacts` |
 | An unregistered type, such as `cpu` | Nowhere: it reports `-1 / None / 0 / 0` |
 
-- A backend registers facts when it adds implementations with `register_kernel_type` that
-  are selected by architecture, SM count or shared memory. An implementation that declares
+- Facts are read by implementations that are selected by architecture, SM count or shared
+  memory. An implementation that declares
   `supported_archs` is never available on an unregistered device type.
 - A target that serves whole ops does not need them: its kernels make their own choices.
 - Facts belong to a device type, not to a target. If several targets serve one device
@@ -266,7 +259,7 @@ The steps of one call, from user code to a backend's `build_kernel`:
 
 ```python
 # ── the caller ───────────────────────────────────────────────────────
-op = GemmFwdOp()                 # no target= in the constructor, so the inputs' device decides
+op = GEMMFwdOp()                 # no target= in the constructor, so the inputs' device decides
                                  #   target="acme" skips detection and uses it directly;
                                  #   target=BUILTIN forces the in-tree kernels
 a = torch.randn(4096, 4096, dtype=torch.float16, device="acme:0")
@@ -282,7 +275,7 @@ d = op(a, b)                     # every input on one device: a.device == b.devi
 #   two or more True   → AmbiguousTargetError, asking for an explicit target=
 
 # ── op layer: run the checks generated from the manifest signature, then hand the whole op to the target ──
-#   GemmFwdOp's own forward and kernel_for serve the in-tree path only, and do not run
+#   GEMMFwdOp's own forward and kernel_for serve the in-tree path only, and do not run
 #   the tensors go to the target in signature.inputs order, inputs it does not write made contiguous
 
 # ── op layer: look up the external memo table — device, then input signature ──
@@ -311,7 +304,7 @@ def build_gemm(a: TensorSpec, b: TensorSpec, *, trans_a, trans_b):
     return AcmeGemm(a, b, trans_a, trans_b)
 
 
-register_kernel_builder(op="GemmFwdOp", target="acme", build_kernel=build_gemm)
+register_kernel_builder(op="GEMMFwdOp", target="acme", build_kernel=build_gemm)
 ```
 
 The op layer calls `build_gemm`; the backend never calls it itself. Importing the backend
@@ -409,7 +402,7 @@ every op the target takes over, with each key spelled as in the manifest. The pa
 ```python
 BUILDERS = {
     "RMSNormFwdOp": build_rms_norm,
-    "GemmFwdOp": build_gemm,
+    "GEMMFwdOp": build_gemm,
 }
 ```
 
@@ -658,8 +651,7 @@ print(load_failures())
 from its own set.** The second layer happens inside `build_kernel`, without the protocol:
 this path has no kernel-level concept, no capability negotiation and no candidate
 filtering. The candidate filtering TileOPs does run (availability, applicability,
-precedence) belongs to the in-tree path and to the smaller mechanism, `register_kernel_type` (see
-[Two ways to extend dispatch](#two-ways)); a target bypasses it.
+precedence) belongs to the in-tree path only; a target bypasses it.
 
 `detect` answers only which devices belong to the backend, and nothing finer.
 **Whether this call is supported (dtype, shape, parameter combination) is answered by
@@ -762,7 +754,7 @@ A claim comes with three sets of numbers:
   `pytest tests/ops --collect-only -q` runs: one with the default options, one with
   `--tileops-target` and `--tileops-device`.
 
-The in-tree kernels ([`src/tileops/kernels/`](https://github.com/tile-ai/TileOPs/tree/main/src/tileops/kernels))
+The in-tree kernels ([`src/tileops/_kernels/`](https://github.com/tile-ai/TileOPs/tree/main/src/tileops/_kernels))
 are the **default implementation**: they have no target name and are not in the registry.
 
 **The in-tree implementations run by default.** With no backend installed, no `target=`