{"record":{"id":"61b1eb93923fef31","repo":"sgl-project/sglang","slug":"timestep-must-be-contiguous","errorCode":null,"errorMessage":"timestep must be contiguous","messagePattern":"timestep must be contiguous","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/diffusion/modulate/ltx2_ada_values_triton.py","lineNumber":147,"sourceCode":"    tl.store(out2_ptr + base, (table2 + temb2).to(tl.bfloat16), mask=mask)\n    tl.store(out3_ptr + base, (table3 + temb3).to(tl.bfloat16), mask=mask)\n    tl.store(out4_ptr + base, (table4 + temb4).to(tl.bfloat16), mask=mask)\n    tl.store(out5_ptr + base, (table5 + temb5).to(tl.bfloat16), mask=mask)\n    tl.store(out6_ptr + base, (table6 + temb6).to(tl.bfloat16), mask=mask)\n    tl.store(out7_ptr + base, (table7 + temb7).to(tl.bfloat16), mask=mask)\n    tl.store(out8_ptr + base, (table8 + temb8).to(tl.bfloat16), mask=mask)\n\n\ndef ltx2_ada_values9(\n    scale_shift_table: torch.Tensor,\n    timestep: torch.Tensor,\n) -> tuple[torch.Tensor, ...]:\n    if timestep.ndim != 3:\n        raise ValueError(\"timestep must have shape [B, S, 9 * D]\")\n    if not timestep.is_cuda or timestep.dtype != torch.bfloat16:\n        raise ValueError(\"timestep must be a CUDA bfloat16 tensor\")\n    if not timestep.is_contiguous():\n        raise ValueError(\"timestep must be contiguous\")\n    if scale_shift_table.ndim != 2 or scale_shift_table.shape[0] != 9:\n        raise ValueError(\"scale_shift_table must have shape [9, D]\")\n    if (\n        not scale_shift_table.is_cuda\n        or scale_shift_table.dtype not in (torch.bfloat16, torch.float32)\n        or scale_shift_table.stride(-1) != 1\n    ):\n        raise ValueError(\n            \"scale_shift_table must be CUDA, bf16/fp32, last-dim contiguous\"\n        )\n\n    total_params = int(scale_shift_table.shape[0])\n    hidden = int(scale_shift_table.shape[1])\n    if hidden <= 0 or timestep.shape[-1] != total_params * hidden:\n        raise ValueError(\"timestep last dim must equal 9 * hidden\")\n    if hidden % 256 != 0 or hidden > 8192:\n        raise ValueError(\"hidden size is outside the supported LTX2 fast-path range\")\n","sourceCodeStart":129,"sourceCodeEnd":165,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/diffusion/modulate/ltx2_ada_values_triton.py#L129-L165","documentation":"The Triton kernel for ltx2_ada_values9 assumes a contiguous timestep tensor so it can compute flat row indices. Non-contiguous inputs (transposed, sliced, or expanded views) are rejected before launch.","triggerScenarios":"Passing an expanded view (timestep[:, None, :].expand(...)) without materializing, a transposed tensor, or a slice along batch/seq that breaks contiguity.","commonSituations":"Broadcasting a [B, 1, 9*D] embedding across sequence positions with expand (stride-0 seq dim); memory-layout tricks from a checkpoint; sliced batch views in pipeline parallelism.","solutions":["Call .contiguous() on timestep before the call","Materialize expanded views: emb.expand(B, S, D).contiguous()","Use .reshape instead of views/slices that yield non-contiguous results"],"exampleFix":"# before\nemb = emb[:, None, :].expand(B, S, 9*D)\nvals = ltx2_ada_values9(table, emb)\n# after\nemb = emb[:, None, :].expand(B, S, 9*D).contiguous()\nvals = ltx2_ada_values9(table, emb)","handlingStrategy":"validation","validationCode":"if not timestep.is_contiguous():\n    timestep = timestep.contiguous()","typeGuard":"def contiguous_or_fix(t):\n    return t if t.is_contiguous() else t.contiguous()","tryCatchPattern":null,"preventionTips":["Materialize expanded views before kernel calls","Prefer .reshape over stride-0 expand for kernel inputs"],"tags":["contiguity","triton","ltx2"],"backgroundTag":"non-contiguous-tensor","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}