{"record":{"id":"1524b36c931ba4bb","repo":"sgl-project/sglang","slug":"fn-name-dst-entry-dims-dims-entry-start-dim","errorCode":null,"errorMessage":"{fn_name}: dst entry dims (dims {entry_start_dim}..{dst.ndim - 1}) must be contiguous; got shape={tuple(dst.shape)} strides={tuple(dst.stride())}","messagePattern":"(.+?): dst entry dims \\(dims (.+?)\\.\\.(.+?)\\) must be contiguous; got shape=(.+?) strides=(.+?)","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/mamba/mamba_state_scatter_triton.py","lineNumber":25,"sourceCode":"\"\"\"\n\nimport torch\nimport triton\nimport triton.language as tl\n\n\ndef _require_entry_contiguous_dst(\n    dst: torch.Tensor, entry_start_dim: int, fn_name: str\n) -> None:\n    \"\"\"dst layout contract: the kernels index through the real layer/slot\n    strides (int64) plus a FLAT element offset within one (layer, slot)\n    entry — layer/slot strides may be arbitrary (envelope-strided unified\n    pool views), but the trailing entry dims must be contiguous.\n    \"\"\"\n    expected = 1\n    for i in range(dst.ndim - 1, entry_start_dim - 1, -1):\n        if dst.shape[i] != 1 and dst.stride(i) != expected:\n            raise ValueError(\n                f\"{fn_name}: dst entry dims (dims {entry_start_dim}..\"\n                f\"{dst.ndim - 1}) must be contiguous; got \"\n                f\"shape={tuple(dst.shape)} strides={tuple(dst.stride())}\"\n            )\n        expected *= dst.shape[i]\n\n\n@triton.jit\ndef track_mamba_state_if_needed_kernel(\n    conv_states_ptr,\n    ssm_states_ptr,\n    cache_indices_ptr,\n    mamba_track_mask_ptr,\n    mamba_track_indices_ptr,\n    conv_state_stride_0,  # stride for first dimension (batch/pool index)\n    ssm_state_stride_0,  # stride for first dimension (batch/pool index)\n    conv_state_numel_per_row: tl.constexpr,  # total elements per row\n    ssm_state_numel_per_row: tl.constexpr,  # total elements per row","sourceCodeStart":7,"sourceCodeEnd":43,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/mamba/mamba_state_scatter_triton.py#L7-L43","documentation":"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.","triggerScenarios":"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.","commonSituations":"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.","solutions":["Make dst contiguous in the entry dims: dst = dst.contiguous() (or build the view so only leading dims are strided)","Move any slicing/striding to the leading (layer/slot) dims, which the kernels explicitly support","If writing into a strided pool, materialize a contiguous buffer and scatter back, or restructure the pool layout"],"exampleFix":"// before\nfused_mamba_state_scatter_with_mask(dst=strided_pool[:, :, ::1].transpose(-1, -2), ...)\n\n// after\nfused_mamba_state_scatter_with_mask(dst=contiguous_dst, ...)","handlingStrategy":"validation","validationCode":"def entry_contiguous(t: torch.Tensor, entry_start_dim: int) -> bool:\n    exp = 1\n    for i in range(t.ndim - 1, entry_start_dim - 1, -1):\n        if t.shape[i] != 1 and t.stride(i) != exp:\n            return False\n        exp *= t.shape[i]\n    return True\nassert entry_contiguous(dst, entry_start_dim)","typeGuard":"def is_entry_contiguous_dst(t: torch.Tensor, entry_start_dim: int) -> bool:\n    return entry_contiguous(t, entry_start_dim)","tryCatchPattern":null,"preventionTips":["Construct unified pool views by striding/slicing only the leading layer/slot dims","When in doubt, call .contiguous() on dst before the scatter and copy back"],"tags":["mamba","contiguity","strides","triton"],"backgroundTag":"non-contiguous-tensor","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}