sgl-project/sglang · error · ValueError

Validate failed: not contiguous on dim D.

Error message

Validate failed: not contiguous on dim D.

What it means

The kernel requires the last (feature) dimension to be contiguous (stride[-1] == 1) so the D-dim loads/stores coalesce. validate_x rejects any x whose innermost stride is not 1 (e.g. transposed or sliced tensors).

Source

Thrown at python/sglang/kernels/ops/diffusion/norm/scale_residual_norm_cutedsl.py:198

        value = tNrN.load()
        copy_if(tSCgSC, tSCrSC)  # gmem -> rmem
        copy_if(tSHgSH, tSHrSH)  # gmem -> rmem
        if cutlass.const_expr(isinstance(tSCrSC, cute.Tensor)):
            value = value * (1 + tSCrSC.load())
        if cutlass.const_expr(isinstance(tSHrSH, cute.Tensor)):
            value = value + tSHrSH.load()
        # Store: y
        tYrY.store(value.to(tYrY.element_type))
        copy_if(tYrY, tYgY)  # rmem -> gmem


def validate_x(t: torch.Tensor, B: int, S: int, D: int):
    if t.dtype not in (torch.float16, torch.bfloat16, torch.float32):
        raise ValueError(f"Validate failed: unsupported dtype: {t.dtype}")
    if t.shape != (B, S, D):
        raise ValueError(f"Validate failed: unsupported tensor shape: {t.shape}.")
    if t.stride()[-1] != 1:
        raise ValueError("Validate failed: not contiguous on dim D.")


def validate_weight_bias(t: Optional[torch.Tensor], D: int):
    if t is None:
        return
    if t.dtype not in (torch.float16, torch.bfloat16, torch.float32):
        raise ValueError(f"Validate failed: unsupported dtype: {t.dtype}")
    if t.shape != (D,):
        raise ValueError(f"Validate failed: unsupported tensor shape: {t.shape}.")
    if t.stride()[-1] != 1:
        raise ValueError("Validate failed: not contiguous on dim D.")


def validate_scale_shift(t: torch.Tensor, B: int, S: int, D: int):
    if t.dtype not in (torch.float16, torch.bfloat16, torch.float32):
        raise ValueError(f"Validate failed: unsupported dtype: {t.dtype}")
    failed = False
    if t.ndim == 1 and (t.shape[0] not in (1, D)):

View on GitHub (pinned to 0132848349)

Solutions

1. Call x = x.contiguous() (or ensure last-dim contiguity) before the fused op
2. Restructure upstream code to keep the hidden dim innermost
3. For unavoidable layouts, fall back to eager torch layer_norm/rms_norm

Example fix

# before
y = fused_norm_scale_shift(x.transpose(-1, -2), ...)
# after
y = fused_norm_scale_shift(x.transpose(-1, -2).contiguous(), ...)

if x.stride(-1) != 1:
    x = x.contiguous()

def last_dim_contiguous(t: torch.Tensor) -> bool:
    return t.stride(-1) == 1

When it happens

Trigger: Passing x.t().transpose(...) style layouts, non-contiguous slices, or tensors from views whose last-dim stride != 1 into fused_norm_scale_shift / fused_scale_residual_norm_scale_shift.

Common situations: Reusing a transposed activation from an attention projection, or taking x[:, :, ::2] style strided slices without materializing.

Related errors

• out must have stride 1 in the last dimension
• O tensor must match dtype
• q, k, and v must be contiguous in head_size
• Validate failed: unsupported dtype: {t.dtype}
• Validate failed: unsupported tensor shape: {t.shape}.

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