sgl-project/sglang · error · ValueError

q, k, and v must be contiguous in head_size

Error message

q, k, and v must be contiguous in head_size

What it means

The Triton pack kernel indexes the head_size dimension with unit stride, so q, k, and v must be contiguous along their last dimension (stride(-1) == 1). Non-unit last-dim strides, typical of transposed or sliced views, are rejected.

Source

Thrown at python/sglang/kernels/ops/diffusion/layout/ulysses_qkv_triton.py:70

    tl.store(output_ptr + output_base + 2 * head_size, v, mask=mask)


def pack_qkv_destination_major(
    q: torch.Tensor,
    k: torch.Tensor,
    v: torch.Tensor,
    world_size: int,
    out: torch.Tensor | None = None,
) -> torch.Tensor:
    """Pack matching ``[rows, global_heads, head_size]`` Q/K/V tensors."""
    if q.dim() != 3 or q.shape != k.shape or q.shape != v.shape:
        raise ValueError("q, k, and v must have the same 3D shape")
    if not (q.is_cuda and k.is_cuda and v.is_cuda):
        raise ValueError("q, k, and v must be CUDA tensors")
    if not (q.device == k.device == v.device and q.dtype == k.dtype == v.dtype):
        raise ValueError("q, k, and v must have the same device and dtype")
    if q.stride(-1) != 1 or k.stride(-1) != 1 or v.stride(-1) != 1:
        raise ValueError("q, k, and v must be contiguous in head_size")
    if world_size < 1 or q.shape[1] % world_size != 0:
        raise ValueError("world_size must be positive and divide global_heads")

    rows, global_heads, head_size = q.shape
    local_heads = global_heads // world_size
    expected_shape = (world_size, rows, local_heads, 3 * head_size)
    if out is not None:
        if not (
            out.shape == expected_shape
            and out.is_contiguous()
            and out.dtype == q.dtype
            and out.device == q.device
        ):
            raise ValueError(
                "out must be a contiguous tensor with the expected shape, "
                "device, and dtype"
            )
        output = out

View on GitHub (pinned to 0132848349)

Solutions

1. Call .contiguous() on q, k, v (or at minimum ensure last-dim stride 1)
2. If tensors come from a transpose, apply .transpose(-1, -2).contiguous() before packing
3. Pre-check with q.stride(-1) == 1 in a debug assert
4. Use torch.empty + copy_ instead of as_strided when constructing inputs

Example fix

# before
q = q.transpose(1, 2)  # last dim now heads, not stride-1-friendly
packed = pack_qkv_destination_major(q, k, v, ws)
# after
q = q.transpose(1, 2).contiguous(); k = k.transpose(1, 2).contiguous(); v = v.transpose(1, 2).contiguous()
packed = pack_qkv_destination_major(q, k, v, ws)

if q.stride(-1) != 1: q = q.contiguous()
if k.stride(-1) != 1: k = k.contiguous()
if v.stride(-1) != 1: v = v.contiguous()

def last_dim_contiguous(*ts) -> bool:
    return all(t.stride(-1) == 1 for t in ts)

When it happens

Trigger: Passing tensors produced by .transpose(-1, -2), narrow/slice views along the last dim, or stride tricks (as_strided) where the last dimension is not stride-1; also expanded tensors with 0-stride dims.

Common situations: Attention implementations that keep Q/K/V in [rows, head_size, heads] transposed layout and forget to permute back; slicing off padding tokens along the head dimension; exporting from checkpoint with non-standard memory layout.

Related errors

• `mixed_qkv` must be contiguous in the last dim.
• timestep must be contiguous
• Validate failed: not contiguous on dim D.
• kv-canary: {name} must be contiguous
• {fn_name}: dst entry dims (dims {entry_start_dim}..{dst.ndim

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