sgl-project/sglang · error · ValueError

Validate failed: unsupported dtype: {t.dtype}

Error message

Validate failed: unsupported dtype: {t.dtype}

What it means

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.

Source

Thrown at python/sglang/kernels/ops/diffusion/norm/scale_residual_norm_cutedsl.py:194

        tNrN = cute.make_rmem_tensor_like(tXrX, tXrX.element_type)
        tNrN.store(value.to(tNrN.element_type))
        tNrN = norm(tNrN, tWrW, tBrB)
        # Compute: value = value * (1 + <scale>) + <shift>
        value = tNrN.load()
        copy_if(tSCgSC, tSCrSC)  # gmem -> rmem
        copy_if(tSHgSH, tSHrSH)  # gmem -> rmem
        if cutlass.const_expr(isinstance(tSCrSC, cute.Tensor)):
            value = value * (1 + tSCrSC.load())
        if cutlass.const_expr(isinstance(tSHrSH, cute.Tensor)):
            value = value + tSHrSH.load()
        # Store: y
        tYrY.store(value.to(tYrY.element_type))
        copy_if(tYrY, tYgY)  # rmem -> gmem


def validate_x(t: torch.Tensor, B: int, S: int, D: int):
    if t.dtype not in (torch.float16, torch.bfloat16, torch.float32):
        raise ValueError(f"Validate failed: unsupported dtype: {t.dtype}")
    if t.shape != (B, S, D):
        raise ValueError(f"Validate failed: unsupported tensor shape: {t.shape}.")
    if t.stride()[-1] != 1:
        raise ValueError("Validate failed: not contiguous on dim D.")


def validate_weight_bias(t: Optional[torch.Tensor], D: int):
    if t is None:
        return
    if t.dtype not in (torch.float16, torch.bfloat16, torch.float32):
        raise ValueError(f"Validate failed: unsupported dtype: {t.dtype}")
    if t.shape != (D,):
        raise ValueError(f"Validate failed: unsupported tensor shape: {t.shape}.")
    if t.stride()[-1] != 1:
        raise ValueError("Validate failed: not contiguous on dim D.")


def validate_scale_shift(t: torch.Tensor, B: int, S: int, D: int):

View on GitHub (pinned to 0132848349)

Solutions

1. Cast x (and residual) to float16/bfloat16/float32 before calling the fused op
2. Verify the model config dtype (e.g. torch_dtype) matches what the caller produces
3. Check upstream ops are not silently upcasting to float64 (e.g. Python float scalars in eager ops)

Example fix

# before
y = fused_norm_scale_shift(x.double(), w, b, scale, shift, "rms")
# after
y = fused_norm_scale_shift(x.to(torch.bfloat16), w.to(torch.bfloat16), b, scale, shift, "rms")

if x.dtype not in (torch.float16, torch.bfloat16, torch.float32):
    x = x.to(torch.bfloat16)

def is_supported_dtype(t: torch.Tensor) -> bool:
    return t.dtype in (torch.float16, torch.bfloat16, torch.float32)

When it happens

Trigger: 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).

Common situations: 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.

Related errors

• Validate failed: unsupported tensor shape: {t.shape}.
• O tensor must match dtype
• unsupported dtype for causal Conv3D cat/pad: {x.dtype}
• Validate failed: not contiguous on dim D.
• D={D} not supported, must be multiple of 256 and <= 8192

AI-assisted analysis of sgl-project/sglang@0132848349 (2026-08-28). Data as JSON: /api/errors/1f1ef6ca855877eb. Report an issue: GitHub.