{"record":{"id":"a8d3fb9d7cb9cb3f","repo":"sgl-project/sglang","slug":"q-k-and-v-must-be-contiguous-in-head-size","errorCode":null,"errorMessage":"q, k, and v must be contiguous in head_size","messagePattern":"q, k, and v must be contiguous in head_size","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/diffusion/layout/ulysses_qkv_triton.py","lineNumber":70,"sourceCode":"    tl.store(output_ptr + output_base + 2 * head_size, v, mask=mask)\n\n\ndef pack_qkv_destination_major(\n    q: torch.Tensor,\n    k: torch.Tensor,\n    v: torch.Tensor,\n    world_size: int,\n    out: torch.Tensor | None = None,\n) -> torch.Tensor:\n    \"\"\"Pack matching ``[rows, global_heads, head_size]`` Q/K/V tensors.\"\"\"\n    if q.dim() != 3 or q.shape != k.shape or q.shape != v.shape:\n        raise ValueError(\"q, k, and v must have the same 3D shape\")\n    if not (q.is_cuda and k.is_cuda and v.is_cuda):\n        raise ValueError(\"q, k, and v must be CUDA tensors\")\n    if not (q.device == k.device == v.device and q.dtype == k.dtype == v.dtype):\n        raise ValueError(\"q, k, and v must have the same device and dtype\")\n    if q.stride(-1) != 1 or k.stride(-1) != 1 or v.stride(-1) != 1:\n        raise ValueError(\"q, k, and v must be contiguous in head_size\")\n    if world_size < 1 or q.shape[1] % world_size != 0:\n        raise ValueError(\"world_size must be positive and divide global_heads\")\n\n    rows, global_heads, head_size = q.shape\n    local_heads = global_heads // world_size\n    expected_shape = (world_size, rows, local_heads, 3 * head_size)\n    if out is not None:\n        if not (\n            out.shape == expected_shape\n            and out.is_contiguous()\n            and out.dtype == q.dtype\n            and out.device == q.device\n        ):\n            raise ValueError(\n                \"out must be a contiguous tensor with the expected shape, \"\n                \"device, and dtype\"\n            )\n        output = out","sourceCodeStart":52,"sourceCodeEnd":88,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/diffusion/layout/ulysses_qkv_triton.py#L52-L88","documentation":"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.","triggerScenarios":"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.","commonSituations":"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.","solutions":["Call .contiguous() on q, k, v (or at minimum ensure last-dim stride 1)","If tensors come from a transpose, apply .transpose(-1, -2).contiguous() before packing","Pre-check with q.stride(-1) == 1 in a debug assert","Use torch.empty + copy_ instead of as_strided when constructing inputs"],"exampleFix":"# before\nq = q.transpose(1, 2)  # last dim now heads, not stride-1-friendly\npacked = pack_qkv_destination_major(q, k, v, ws)\n# after\nq = q.transpose(1, 2).contiguous(); k = k.transpose(1, 2).contiguous(); v = v.transpose(1, 2).contiguous()\npacked = pack_qkv_destination_major(q, k, v, ws)","handlingStrategy":"validation","validationCode":"if q.stride(-1) != 1: q = q.contiguous()\nif k.stride(-1) != 1: k = k.contiguous()\nif v.stride(-1) != 1: v = v.contiguous()","typeGuard":"def last_dim_contiguous(*ts) -> bool:\n    return all(t.stride(-1) == 1 for t in ts)","tryCatchPattern":null,"preventionTips":["Call .contiguous() after transpose-based layouts","Avoid as_strided for attention inputs"],"tags":["stride","contiguity","ulysses","triton"],"backgroundTag":"non-contiguous-tensor","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}