sgl-project/sglang · error · ValueError

{fn_name}: dst entry dims (dims {entry_start_dim}..{dst.ndim

Error message

{fn_name}: dst entry dims (dims {entry_start_dim}..{dst.ndim - 1}) must be contiguous; got shape={tuple(dst.shape)} strides={tuple(dst.stride())}

What it means

The fused mamba/conv state scatter Triton kernels require the trailing 'entry' dims of dst (e.g. the per-slot state tensor dims) to be contiguous; the leading envelope/layer dims may have arbitrary strides. This mirrors Triton's need for dense innermost addressing. Size-1 dims are exempt.

Source

Thrown at python/sglang/kernels/ops/mamba/mamba_state_scatter_triton.py:25

"""

import torch
import triton
import triton.language as tl


def _require_entry_contiguous_dst(
    dst: torch.Tensor, entry_start_dim: int, fn_name: str
) -> None:
    """dst layout contract: the kernels index through the real layer/slot
    strides (int64) plus a FLAT element offset within one (layer, slot)
    entry — layer/slot strides may be arbitrary (envelope-strided unified
    pool views), but the trailing entry dims must be contiguous.
    """
    expected = 1
    for i in range(dst.ndim - 1, entry_start_dim - 1, -1):
        if dst.shape[i] != 1 and dst.stride(i) != expected:
            raise ValueError(
                f"{fn_name}: dst entry dims (dims {entry_start_dim}.."
                f"{dst.ndim - 1}) must be contiguous; got "
                f"shape={tuple(dst.shape)} strides={tuple(dst.stride())}"
            )
        expected *= dst.shape[i]


@triton.jit
def track_mamba_state_if_needed_kernel(
    conv_states_ptr,
    ssm_states_ptr,
    cache_indices_ptr,
    mamba_track_mask_ptr,
    mamba_track_indices_ptr,
    conv_state_stride_0,  # stride for first dimension (batch/pool index)
    ssm_state_stride_0,  # stride for first dimension (batch/pool index)
    conv_state_numel_per_row: tl.constexpr,  # total elements per row
    ssm_state_numel_per_row: tl.constexpr,  # total elements per row

View on GitHub (pinned to 0132848349)

Solutions

1. Make dst contiguous in the entry dims: dst = dst.contiguous() (or build the view so only leading dims are strided)
2. Move any slicing/striding to the leading (layer/slot) dims, which the kernels explicitly support
3. If writing into a strided pool, materialize a contiguous buffer and scatter back, or restructure the pool layout

Example fix

// before
fused_mamba_state_scatter_with_mask(dst=strided_pool[:, :, ::1].transpose(-1, -2), ...)

// after
fused_mamba_state_scatter_with_mask(dst=contiguous_dst, ...)

def entry_contiguous(t: torch.Tensor, entry_start_dim: int) -> bool:
    exp = 1
    for i in range(t.ndim - 1, entry_start_dim - 1, -1):
        if t.shape[i] != 1 and t.stride(i) != exp:
            return False
        exp *= t.shape[i]
    return True
assert entry_contiguous(dst, entry_start_dim)

def is_entry_contiguous_dst(t: torch.Tensor, entry_start_dim: int) -> bool:
    return entry_contiguous(t, entry_start_dim)

When it happens

Trigger: Calling fused_mamba_state_scatter_with_mask or fused_conv_window_scatter_with_mask with a dst whose inner dims are non-contiguous — e.g. a transposed view, a strided slice like pool[:, :, ::2], or a tensor from .expand on an inner dim.

Common situations: Unified mamba pool views created with permute/slice where the slicing hit the state dims instead of layer dims; tests exercising envelope-strided views (accepted) vs entry-strided (rejected); refactoring cache layouts.

Related errors

• kv-canary: {name} must be contiguous
• `mixed_qkv` must be contiguous in the last dim.
• `mixed_qkv` must be a 2D tensor (got ndim={mixed_qkv.ndim}).
• `mixed_qkv` must be contiguous in the last dim.
• `a`/`b` must be contiguous in the last dim.

AI-assisted analysis of sgl-project/sglang@0132848349 (2026-08-28). Data as JSON: /api/errors/1524b36c931ba4bb. Report an issue: GitHub.