> For clean Markdown of any page, append .md to the page URL.
> For a complete documentation index, see https://docs.nvidia.com/nemo/automodel/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.deepseek_v4.kernels.tilelang_indexer_fwd

## Module Contents

### Functions

| Name                                                                                                                            | Description                                                            |
| ------------------------------------------------------------------------------------------------------------------------------- | ---------------------------------------------------------------------- |
| [`_make_causal_cu_seqlens`](#nemo_automodel-components-models-deepseek_v4-kernels-tilelang_indexer_fwd-_make_causal_cu_seqlens) | Generate cu\_seqlens for causal masking on compressed KV positions.    |
| [`batched_indexer_fwd`](#nemo_automodel-components-models-deepseek_v4-kernels-tilelang_indexer_fwd-batched_indexer_fwd)         | Batched forward: loops over batch dim.                                 |
| [`clean_logits_`](#nemo_automodel-components-models-deepseek_v4-kernels-tilelang_indexer_fwd-clean_logits_)                     | -                                                                      |
| [`indexer_fwd_interface`](#nemo_automodel-components-models-deepseek_v4-kernels-tilelang_indexer_fwd-indexer_fwd_interface)     | Forward interface matching GLM-5's API but for a single batch element. |
| [`tl_indexer_fwd_impl`](#nemo_automodel-components-models-deepseek_v4-kernels-tilelang_indexer_fwd-tl_indexer_fwd_impl)         | -                                                                      |

### API

```python
nemo_automodel.components.models.deepseek_v4.kernels.tilelang_indexer_fwd._make_causal_cu_seqlens(
    seq_len_q,
    seq_len_kv,
    compress_ratio,
    device
)
```

Generate cu\_seqlens for causal masking on compressed KV positions.

For query at position p, valid compressed groups are \[0, (p+1) // compress\_ratio).

```python
nemo_automodel.components.models.deepseek_v4.kernels.tilelang_indexer_fwd.batched_indexer_fwd(
    q,
    k,
    weights,
    cu_seqlen_ks,
    cu_seqlen_ke
)
```

Batched forward: loops over batch dim.

**Parameters:**

\[seqlen, batch, heads, dim] bf16

\[seqlen\_kv, batch, dim] bf16

\[seqlen, batch, heads] fp32

\[seqlen] int32

\[seqlen] int32

**Returns:**

\[batch, seqlen, seqlen\_kv] fp32

```python
nemo_automodel.components.models.deepseek_v4.kernels.tilelang_indexer_fwd.clean_logits_(
    threads: int = 512,
    block_K: int = 4096
)
```

```python
nemo_automodel.components.models.deepseek_v4.kernels.tilelang_indexer_fwd.indexer_fwd_interface(
    q,
    kv,
    weights,
    cu_seqlen_ks,
    cu_seqlen_ke,
    clean_logits = True
)
```

Forward interface matching GLM-5's API but for a single batch element.

**Parameters:**

\[seq\_len, heads, index\_dim] bf16

\[seq\_len\_kv, index\_dim] bf16

\[seq\_len, heads] fp32

\[seq\_len] int32 — start of valid KV range per query

\[seq\_len] int32 — end of valid KV range per query

**Returns:**

\[seq\_len, seq\_len\_kv] fp32

```python
nemo_automodel.components.models.deepseek_v4.kernels.tilelang_indexer_fwd.tl_indexer_fwd_impl(
    heads,
    index_dim,
    block_N = 256,
    num_stages = 3,
    threads = 512,
    block_Q = None
)
```