sgl-project/sglang · error · ValueError

cos and sin must be CUDA and last-dim contiguous

Error message

cos and sin must be CUDA and last-dim contiguous

What it means

cos and sin must be CUDA tensors with unit stride in the last dimension; the Triton kernel loads them directly from GPU memory assuming a dense innermost dimension. CPU tables or non-contiguous slices are rejected.

Source

Thrown at python/sglang/kernels/ops/diffusion/rope/hunyuan_qkv_pack_triton.py:185

        raise ValueError("QKV tensors must be on the same CUDA device")
    batch, img_tokens, num_heads, head_dim = img_q.shape
    txt_tokens = txt_q.shape[1]
    expected_img = (batch, img_tokens, num_heads, head_dim)
    expected_txt = (batch, txt_tokens, num_heads, head_dim)
    if any(tuple(x.shape) != expected_img for x in (img_q, img_k, img_v)):
        raise ValueError("image QKV shapes must match")
    if any(tuple(x.shape) != expected_txt for x in (txt_q, txt_k, txt_v)):
        raise ValueError("text QKV shapes must match")
    if any(x.stride(-1) != 1 for x in tensors):
        raise ValueError("QKV last dimensions must be contiguous")
    if head_dim <= 0 or head_dim > 128 or head_dim % 2:
        raise ValueError("head_dim must be positive, even, and <= 128")
    if cos.ndim != 2 or sin.ndim != 2 or cos.shape != sin.shape:
        raise ValueError("cos and sin must have matching [S, D/2] shapes")
    if cos.shape[0] < img_tokens or cos.shape[1] != head_dim // 2:
        raise ValueError("cos/sin shape does not cover image tokens and head_dim")
    if not cos.is_cuda or not sin.is_cuda or cos.stride(-1) != 1 or sin.stride(-1) != 1:
        raise ValueError("cos and sin must be CUDA and last-dim contiguous")
    if cos.device != img_q.device or sin.device != img_q.device:
        raise ValueError("QKV and cos/sin tensors must be on the same CUDA device")

    total_tokens = img_tokens + txt_tokens
    storage = torch.empty(
        (3, batch, total_tokens, num_heads, head_dim),
        device=img_q.device,
        dtype=img_q.dtype,
    )
    args = []
    for x in tensors:
        args.extend((x.stride(0), x.stride(1), x.stride(2)))
    with torch.cuda.device(img_q.device):
        _hunyuan_qkv_rope_pack_kernel[
            lambda meta: (
                batch * total_tokens,
                triton.cdiv(num_heads, meta["BLOCK_HEADS"]),
            )

View on GitHub (pinned to 0132848349)

Solutions

1. Move tables to the same CUDA device: cos = cos.cuda() (or .to(img_q.device)).
2. Call .contiguous() if the table was sliced.
3. Precompute tables once on GPU at model setup.

Example fix

// before
cos, sin = build_rope_table(seq, head_dim)  # CPU tensors
// after
cos, sin = build_rope_table(seq, head_dim)
cos, sin = cos.to('cuda').contiguous(), sin.to('cuda').contiguous()

assert cos.is_cuda and sin.is_cuda and cos.stride(-1) == 1 and sin.stride(-1) == 1

When it happens

Trigger: Passing cos/sin computed on CPU (not moved with .cuda()), or a non-contiguous slice like freqs[:, ::2] as the table.

Common situations: Precomputing RoPE tables on CPU at model init and forgetting .to(device), or column-subsampled frequency tables.

Related errors

• q must be contiguous
• kv must be contiguous
• unsupported input for causal Conv3D cat/pad CUDA
• QKV and cos/sin tensors must be on the same CUDA device
• unsupported input for LTX2 QKNorm split-RoPE CUDA

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