sgl-project/sglang · error · ValueError

out must be a contiguous tensor with the expected shape, dev

Error message

out must be a contiguous tensor with the expected shape, device, and dtype

What it means

When an out tensor is supplied, pack_qkv_destination_major requires it to exactly match the expected output shape [world_size, rows, local_heads, 3*head_size], be contiguous, and share the input's dtype and device. Otherwise it refuses to write into it.

Source

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

        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
    else:
        output = torch.empty(
            expected_shape,
            dtype=q.dtype,
            device=q.device,
        )
    total_elements = rows * global_heads * head_size
    if total_elements == 0:
        return output

    block_size = 1024
    with torch.get_device_module().device(q.device):
        _pack_qkv_destination_major_kernel[(triton.cdiv(total_elements, block_size),)](
            output,

View on GitHub (pinned to 0132848349)

Solutions

1. Allocate out as torch.empty((world_size, rows, global_heads // world_size, 3 * head_size), device=q.device, dtype=q.dtype)
2. Or pass out=None and let the function allocate
3. If reusing a workspace, take an exact narrow/view so the resulting tensor is contiguous with the right shape
4. Double-check dtype/device match after any model dtype migration

Example fix

# before
out = torch.empty((ws, rows, local_heads, head_size), device='cuda', dtype=torch.float32)  # wrong: missing *3 and wrong dtype
# after
out = torch.empty((ws, rows, local_heads, 3 * head_size), device=q.device, dtype=q.dtype)
packed = pack_qkv_destination_major(q, k, v, ws, out=out)

expected = (world_size, q.shape[0], q.shape[1] // world_size, 3 * q.shape[2])
if out is not None:
    assert out.shape == expected and out.is_contiguous() and out.dtype == q.dtype and out.device == q.device
else:
    out = None  # let the function allocate

def out_ok(out, q, ws) -> bool:
    r, h, d = q.shape
    return out.shape == (ws, r, h // ws, 3 * d) and out.is_contiguous() and out.dtype == q.dtype and out.device == q.device

When it happens

Trigger: Passing a preallocated out buffer with the wrong shape (e.g. forgetting the factor 3 on head_size or swapping dims), a non-contiguous buffer (sliced view), or one allocated with a different dtype/device than q.

Common situations: Memory-reuse optimizations where buffers are allocated once with a stale shape after config changes; buffers created with torch.empty on the default device while inputs live on another GPU; passing a broader workspace slice instead of an exact-shaped view.

Related errors

• q, k, and v must have the same 3D shape
• {tensor_name}{context_clause} with shape {tensor.shape} cann
• Block sparse tensors{context} must have shapes (B, H, M) and
• Block sparse tensors{context} {dim_name} dim must be {tgt} o
• Block sparse tensors{context} must share the same m-block di

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