sgl-project/sglang · error · ValueError

q, k, and v must be CUDA tensors

Error message

q, k, and v must be CUDA tensors

What it means

pack_qkv_destination_major launches a Triton kernel and therefore requires q, k, and v to all be CUDA tensors. Supplying any CPU tensor (or a tensor on a non-CUDA device) fails this explicit check before kernel launch.

Source

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

    )
    output_base = head_slot * (3 * head_size) + dim
    tl.store(output_ptr + output_base, q, mask=mask)
    tl.store(output_ptr + output_base + head_size, k, mask=mask)
    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(

View on GitHub (pinned to 0132848349)

Solutions

1. Move all three tensors to the GPU: q, k, v = q.cuda(), k.cuda(), v.cuda()
2. In tests, always construct inputs with device='cuda'
3. Skip or mark tests xfail on CPU-only machines
4. Ensure the model runner places attention tensors on the correct device

Example fix

# before
q = torch.randn(128, 32, 64)  # CPU
# after
q = torch.randn(128, 32, 64, device='cuda', dtype=torch.bfloat16)

assert q.is_cuda and k.is_cuda and v.is_cuda

def all_cuda(*ts) -> bool:
    return all(t.is_cuda for t in ts)

When it happens

Trigger: Calling pack_qkv_destination_major with any of q, k, v on CPU — common in unit tests that build tensors with torch.zeros(...) without device='cuda', or when a tensor was moved to CPU for logging and not moved back.

Common situations: Writing tests without device='cuda'; mixed-device pipelines where activations were detached().cpu()'d for debugging; running on CPU-only environments where the kernel cannot work at all.

Related errors

• Triton sparse_attn_v4_paged_prefill requires CUDA/HIP tensor
• {name}_block_cnt and {name}_block_idx must be on the same de
• {name}_block tensors must live on CUDA
• FlashAttention combine kernel cannot be implemented with giv
• hd256 forward varlen expects q rank 3 or 5, got rank {q_rank

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