> 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_bwd

## Module Contents

### Functions

| Name                                                                                                                        | Description                                    |
| --------------------------------------------------------------------------------------------------------------------------- | ---------------------------------------------- |
| [`batched_indexer_bwd`](#nemo_automodel-components-models-deepseek_v4-kernels-tilelang_indexer_bwd-batched_indexer_bwd)     | Batched backward: loops over batch dim.        |
| [`indexer_bwd_interface`](#nemo_automodel-components-models-deepseek_v4-kernels-tilelang_indexer_bwd-indexer_bwd_interface) | Backward interface for a single batch element. |
| [`tl_indexer_bwd_impl`](#nemo_automodel-components-models-deepseek_v4-kernels-tilelang_indexer_bwd-tl_indexer_bwd_impl)     | -                                              |

### Data

[`BF16`](#nemo_automodel-components-models-deepseek_v4-kernels-tilelang_indexer_bwd-BF16)

[`FP32`](#nemo_automodel-components-models-deepseek_v4-kernels-tilelang_indexer_bwd-FP32)

[`INT32`](#nemo_automodel-components-models-deepseek_v4-kernels-tilelang_indexer_bwd-INT32)

[`pass_configs`](#nemo_automodel-components-models-deepseek_v4-kernels-tilelang_indexer_bwd-pass_configs)

### API

```python
nemo_automodel.components.models.deepseek_v4.kernels.tilelang_indexer_bwd.batched_indexer_bwd(
    index_q,
    weights,
    index_k,
    topk_indices,
    grad_scores
)
```

Batched backward: loops over batch dim.

**Parameters:**

\[seqlen, batch, heads, dim] bf16

\[seqlen, batch, heads] fp32

\[seqlen\_kv, batch, dim] bf16

\[batch, seqlen, topk] int32

\[batch, seqlen, topk] fp32

**Returns:**

\[seqlen, batch, heads, dim] bf16

```python
nemo_automodel.components.models.deepseek_v4.kernels.tilelang_indexer_bwd.indexer_bwd_interface(
    index_q: torch.Tensor,
    weights: torch.Tensor,
    index_k: torch.Tensor,
    topk_indices: torch.Tensor,
    grad_scores: torch.Tensor
)
```

Backward interface for a single batch element.

**Parameters:**

\[seq\_len, heads, dim] bf16

\[seq\_len, heads] fp32

\[seq\_len\_kv, dim] bf16

\[seq\_len, topk] int32

\[seq\_len, topk] fp32

**Returns:**

\[seq\_len, heads, dim] bf16

```python
nemo_automodel.components.models.deepseek_v4.kernels.tilelang_indexer_bwd.tl_indexer_bwd_impl(
    heads: int,
    dim: int,
    topk: int,
    block_I: int = 32,
    num_stages: int = 0,
    num_threads: int = 128
)
```

```python
nemo_automodel.components.models.deepseek_v4.kernels.tilelang_indexer_bwd.BF16 = T.bfloat16
```

```python
nemo_automodel.components.models.deepseek_v4.kernels.tilelang_indexer_bwd.FP32 = T.float32
```

```python
nemo_automodel.components.models.deepseek_v4.kernels.tilelang_indexer_bwd.INT32 = T.int32
```

```python
nemo_automodel.components.models.deepseek_v4.kernels.tilelang_indexer_bwd.pass_configs = {tl.PassConfigKey.TL_DISABLE_TMA_LOWER: True, tl.PassConfigKey.TL_DISABLE_WARP_S...
```