> For clean Markdown of any page, append .md to the page URL.
> For a complete documentation index, see https://docs.nvidia.com/nemo/automodel/latest/nemo-automodel/nemo_automodel/components/models/glm_moe_dsa/kernels/llms.txt.
> For AI client integration (Claude Code, Cursor, etc.), connect to the MCP server at https://docs.nvidia.com/nemo/automodel/_mcp/server.

# nemo_automodel.components.models.glm_moe_dsa.kernels

GLM-5.2 DSA kernel adapters and vendored TileLang kernels.

`cudnn_dsa.py` is AutoModel's adapter for the optional cuDNN Frontend DSA
and FlashMLA runtimes. The other kernel sources documented below are vendored.

**slime GLM-5.2 kernels**

The vendored lighting-indexer and sparse-MLA kernels were adapted from THUDM's
slime GLM-5.2 plugin:

* Upstream project: [https://github.com/THUDM/slime](https://github.com/THUDM/slime)
* Upstream revision: `8f5e2151943e9ed0bbffaed93741d3473abb58d9`
* Upstream source tree:
  [https://github.com/THUDM/slime/tree/8f5e2151943e9ed0bbffaed93741d3473abb58d9/slime\_plugins/models/glm5/ops](https://github.com/THUDM/slime/tree/8f5e2151943e9ed0bbffaed93741d3473abb58d9/slime_plugins/models/glm5/ops)
* Upstream license: Apache License 2.0

The slime kernels are themselves adapted from the tile-ai/tilelang DeepSeek-V3.2
examples (per-file upstream links are preserved in each file header).

Per-file source mapping:

\===============================  ==============================================================
Local file                       Upstream file (slime\_plugins/models/glm5/ops/)
\===============================  ==============================================================
`indexer.py`                   `indexer.py`
`sparse_mla.py`                `sparse_mla.py`
`tilelang_indexer_fwd.py`      `tilelang_indexer_fwd.py`
`tilelang_indexer_bwd.py`      `tilelang_indexer_bwd.py`
`tilelang_sparse_mla_fwd.py`   `tilelang_sparse_mla_fwd.py`
`tilelang_sparse_mla_bwd.py`   `tilelang_sparse_mla_bwd.py`
\===============================  ==============================================================

Local modifications: each raw kernel file imports `T`/`tilelang` through the
local `_tilelang.py` lazy shim, matching the DeepSeek-V4 kernels. This lets
AutoModel import without importing the optional TileLang runtime; real TileLang
is loaded only when a TileLang kernel is called. The kernels are wired into
AutoModel's GLM-5.2 DSA layers via `optimized_kernels.py` and gated behind
`backend.attn == "tilelang"`.

These kernels require `tilelang` (an optional dependency). Note `tilelang`
0.1.11 must be paired with `apache-tvm-ffi==0.1.11`; `apache-tvm-ffi` 0.1.12
breaks `import tilelang` with a tvm-ffi type double-registration error.

## Submodules

* **[`nemo_automodel.components.models.glm_moe_dsa.kernels._tilelang`](/nemo-automodel/nemo_automodel/components/models/glm_moe_dsa/kernels/_tilelang)**
* **[`nemo_automodel.components.models.glm_moe_dsa.kernels.cudnn_dsa`](/nemo-automodel/nemo_automodel/components/models/glm_moe_dsa/kernels/cudnn_dsa)**
* **[`nemo_automodel.components.models.glm_moe_dsa.kernels.indexer`](/nemo-automodel/nemo_automodel/components/models/glm_moe_dsa/kernels/indexer)**
* **[`nemo_automodel.components.models.glm_moe_dsa.kernels.sparse_mla`](/nemo-automodel/nemo_automodel/components/models/glm_moe_dsa/kernels/sparse_mla)**
* **[`nemo_automodel.components.models.glm_moe_dsa.kernels.tilelang_indexer_bwd`](/nemo-automodel/nemo_automodel/components/models/glm_moe_dsa/kernels/tilelang_indexer_bwd)**
* **[`nemo_automodel.components.models.glm_moe_dsa.kernels.tilelang_indexer_fwd`](/nemo-automodel/nemo_automodel/components/models/glm_moe_dsa/kernels/tilelang_indexer_fwd)**
* **[`nemo_automodel.components.models.glm_moe_dsa.kernels.tilelang_sparse_mla_bwd`](/nemo-automodel/nemo_automodel/components/models/glm_moe_dsa/kernels/tilelang_sparse_mla_bwd)**
* **[`nemo_automodel.components.models.glm_moe_dsa.kernels.tilelang_sparse_mla_fwd`](/nemo-automodel/nemo_automodel/components/models/glm_moe_dsa/kernels/tilelang_sparse_mla_fwd)**

## Docs

- [nemo_automodel.components.models.glm_moe_dsa.kernels.cudnn_dsa](https://docs.nvidia.com/nemo/automodel/latest/nemo-automodel/nemo_automodel/components/models/glm_moe_dsa/kernels/cudnn_dsa.md)
- [nemo_automodel.components.models.glm_moe_dsa.kernels.indexer](https://docs.nvidia.com/nemo/automodel/latest/nemo-automodel/nemo_automodel/components/models/glm_moe_dsa/kernels/indexer.md)
- [nemo_automodel.components.models.glm_moe_dsa.kernels.sparse_mla](https://docs.nvidia.com/nemo/automodel/latest/nemo-automodel/nemo_automodel/components/models/glm_moe_dsa/kernels/sparse_mla.md)
- [nemo_automodel.components.models.glm_moe_dsa.kernels.tilelang_indexer_bwd](https://docs.nvidia.com/nemo/automodel/latest/nemo-automodel/nemo_automodel/components/models/glm_moe_dsa/kernels/tilelang_indexer_bwd.md)
- [nemo_automodel.components.models.glm_moe_dsa.kernels.tilelang_indexer_fwd](https://docs.nvidia.com/nemo/automodel/latest/nemo-automodel/nemo_automodel/components/models/glm_moe_dsa/kernels/tilelang_indexer_fwd.md)
- [nemo_automodel.components.models.glm_moe_dsa.kernels.tilelang_sparse_mla_bwd](https://docs.nvidia.com/nemo/automodel/latest/nemo-automodel/nemo_automodel/components/models/glm_moe_dsa/kernels/tilelang_sparse_mla_bwd.md)
- [nemo_automodel.components.models.glm_moe_dsa.kernels.tilelang_sparse_mla_fwd](https://docs.nvidia.com/nemo/automodel/latest/nemo-automodel/nemo_automodel/components/models/glm_moe_dsa/kernels/tilelang_sparse_mla_fwd.md)