{"record":{"id":"e96e68542c1782e7","repo":"sgl-project/sglang","slug":"validate-failed-not-contiguous-on-dim-d","errorCode":null,"errorMessage":"Validate failed: not contiguous on dim D.","messagePattern":"Validate failed: not contiguous on dim D\\.","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/diffusion/norm/scale_residual_norm_cutedsl.py","lineNumber":198,"sourceCode":"        value = tNrN.load()\n        copy_if(tSCgSC, tSCrSC)  # gmem -> rmem\n        copy_if(tSHgSH, tSHrSH)  # gmem -> rmem\n        if cutlass.const_expr(isinstance(tSCrSC, cute.Tensor)):\n            value = value * (1 + tSCrSC.load())\n        if cutlass.const_expr(isinstance(tSHrSH, cute.Tensor)):\n            value = value + tSHrSH.load()\n        # Store: y\n        tYrY.store(value.to(tYrY.element_type))\n        copy_if(tYrY, tYgY)  # rmem -> gmem\n\n\ndef validate_x(t: torch.Tensor, B: int, S: int, D: int):\n    if t.dtype not in (torch.float16, torch.bfloat16, torch.float32):\n        raise ValueError(f\"Validate failed: unsupported dtype: {t.dtype}\")\n    if t.shape != (B, S, D):\n        raise ValueError(f\"Validate failed: unsupported tensor shape: {t.shape}.\")\n    if t.stride()[-1] != 1:\n        raise ValueError(\"Validate failed: not contiguous on dim D.\")\n\n\ndef validate_weight_bias(t: Optional[torch.Tensor], D: int):\n    if t is None:\n        return\n    if t.dtype not in (torch.float16, torch.bfloat16, torch.float32):\n        raise ValueError(f\"Validate failed: unsupported dtype: {t.dtype}\")\n    if t.shape != (D,):\n        raise ValueError(f\"Validate failed: unsupported tensor shape: {t.shape}.\")\n    if t.stride()[-1] != 1:\n        raise ValueError(\"Validate failed: not contiguous on dim D.\")\n\n\ndef validate_scale_shift(t: torch.Tensor, B: int, S: int, D: int):\n    if t.dtype not in (torch.float16, torch.bfloat16, torch.float32):\n        raise ValueError(f\"Validate failed: unsupported dtype: {t.dtype}\")\n    failed = False\n    if t.ndim == 1 and (t.shape[0] not in (1, D)):","sourceCodeStart":180,"sourceCodeEnd":216,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/diffusion/norm/scale_residual_norm_cutedsl.py#L180-L216","documentation":"The kernel requires the last (feature) dimension to be contiguous (stride[-1] == 1) so the D-dim loads/stores coalesce. validate_x rejects any x whose innermost stride is not 1 (e.g. transposed or sliced tensors).","triggerScenarios":"Passing x.t().transpose(...) style layouts, non-contiguous slices, or tensors from views whose last-dim stride != 1 into fused_norm_scale_shift / fused_scale_residual_norm_scale_shift.","commonSituations":"Reusing a transposed activation from an attention projection, or taking x[:, :, ::2] style strided slices without materializing.","solutions":["Call x = x.contiguous() (or ensure last-dim contiguity) before the fused op","Restructure upstream code to keep the hidden dim innermost","For unavoidable layouts, fall back to eager torch layer_norm/rms_norm"],"exampleFix":"# before\ny = fused_norm_scale_shift(x.transpose(-1, -2), ...)\n# after\ny = fused_norm_scale_shift(x.transpose(-1, -2).contiguous(), ...)","handlingStrategy":"validation","validationCode":"if x.stride(-1) != 1:\n    x = x.contiguous()","typeGuard":"def last_dim_contiguous(t: torch.Tensor) -> bool:\n    return t.stride(-1) == 1","tryCatchPattern":null,"preventionTips":["Keep hidden dim innermost in layout transforms","Call .contiguous() after transposes/slices feeding fused kernels"],"tags":["contiguity","stride","validation","cuda-kernel"],"backgroundTag":"non-contiguous-tensor","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}