{"record":{"id":"1f1ef6ca855877eb","repo":"sgl-project/sglang","slug":"validate-failed-unsupported-dtype-t-dtype","errorCode":null,"errorMessage":"Validate failed: unsupported dtype: {t.dtype}","messagePattern":"Validate failed: unsupported dtype: (.+?)","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/diffusion/norm/scale_residual_norm_cutedsl.py","lineNumber":194,"sourceCode":"        tNrN = cute.make_rmem_tensor_like(tXrX, tXrX.element_type)\n        tNrN.store(value.to(tNrN.element_type))\n        tNrN = norm(tNrN, tWrW, tBrB)\n        # Compute: value = value * (1 + <scale>) + <shift>\n        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):","sourceCodeStart":176,"sourceCodeEnd":212,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/diffusion/norm/scale_residual_norm_cutedsl.py#L176-L212","documentation":"The CuTe-DSL fused norm/scale/shift kernel only accepts fp16/bf16/fp32 activations. validate_x checks the input tensor dtype before dispatching to the compiled CUDA kernel because the kernel is only instantiated for those element types.","triggerScenarios":"Calling fused_norm_scale_shift or fused_scale_residual_norm_scale_shift with an x (or residual) tensor whose dtype is not torch.float16, torch.bfloat16, or torch.float32 (e.g. float64 or uint8).","commonSituations":"Passing autocast-off fp64 debug tensors, quantized/int tensors, or a residual stored in a different dtype than x after a dtype change elsewhere in the pipeline.","solutions":["Cast x (and residual) to float16/bfloat16/float32 before calling the fused op","Verify the model config dtype (e.g. torch_dtype) matches what the caller produces","Check upstream ops are not silently upcasting to float64 (e.g. Python float scalars in eager ops)"],"exampleFix":"# before\ny = fused_norm_scale_shift(x.double(), w, b, scale, shift, \"rms\")\n# after\ny = fused_norm_scale_shift(x.to(torch.bfloat16), w.to(torch.bfloat16), b, scale, shift, \"rms\")","handlingStrategy":"validation","validationCode":"if x.dtype not in (torch.float16, torch.bfloat16, torch.float32):\n    x = x.to(torch.bfloat16)","typeGuard":"def is_supported_dtype(t: torch.Tensor) -> bool:\n    return t.dtype in (torch.float16, torch.bfloat16, torch.float32)","tryCatchPattern":null,"preventionTips":["Standardize activations to bf16/fp16 early in the pipeline","Assert dtype before fused kernel dispatch"],"tags":["dtype","validation","cuda-kernel","diffusion"],"backgroundTag":"unsupported-tensor-dtype","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}