{"record":{"id":"9f854c4bb77e17ee","repo":"sgl-project/sglang","slug":"out-must-be-a-contiguous-tensor-with-the-expected","errorCode":null,"errorMessage":"out must be a contiguous tensor with the expected shape, device, and dtype","messagePattern":"out must be a contiguous tensor with the expected shape, device, and dtype","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/diffusion/layout/ulysses_qkv_triton.py","lineNumber":84,"sourceCode":"        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\n    else:\n        output = torch.empty(\n            expected_shape,\n            dtype=q.dtype,\n            device=q.device,\n        )\n    total_elements = rows * global_heads * head_size\n    if total_elements == 0:\n        return output\n\n    block_size = 1024\n    with torch.get_device_module().device(q.device):\n        _pack_qkv_destination_major_kernel[(triton.cdiv(total_elements, block_size),)](\n            output,","sourceCodeStart":66,"sourceCodeEnd":102,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/diffusion/layout/ulysses_qkv_triton.py#L66-L102","documentation":"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.","triggerScenarios":"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.","commonSituations":"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.","solutions":["Allocate out as torch.empty((world_size, rows, global_heads // world_size, 3 * head_size), device=q.device, dtype=q.dtype)","Or pass out=None and let the function allocate","If reusing a workspace, take an exact narrow/view so the resulting tensor is contiguous with the right shape","Double-check dtype/device match after any model dtype migration"],"exampleFix":"# before\nout = torch.empty((ws, rows, local_heads, head_size), device='cuda', dtype=torch.float32)  # wrong: missing *3 and wrong dtype\n# after\nout = torch.empty((ws, rows, local_heads, 3 * head_size), device=q.device, dtype=q.dtype)\npacked = pack_qkv_destination_major(q, k, v, ws, out=out)","handlingStrategy":"validation","validationCode":"expected = (world_size, q.shape[0], q.shape[1] // world_size, 3 * q.shape[2])\nif out is not None:\n    assert out.shape == expected and out.is_contiguous() and out.dtype == q.dtype and out.device == q.device\nelse:\n    out = None  # let the function allocate","typeGuard":"def out_ok(out, q, ws) -> bool:\n    r, h, d = q.shape\n    return out.shape == (ws, r, h // ws, 3 * d) and out.is_contiguous() and out.dtype == q.dtype and out.device == q.device","tryCatchPattern":null,"preventionTips":["Reallocate buffers after config changes","Prefer out=None unless profiling shows allocation cost matters"],"tags":["out-buffer","shape-mismatch","allocation","ulysses"],"backgroundTag":"output-buffer-validation-failed","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}