sgl-project/sglang · error · ValueError

timestep must be a CUDA bfloat16 tensor

Error message

timestep must be a CUDA bfloat16 tensor

What it means

The LTX2 fused AdaLN Triton kernel is hard-coded for bfloat16 timestep inputs on CUDA. Tensors on CPU, other devices, or in other dtypes (fp16/fp32) fail this check.

Source

Thrown at python/sglang/kernels/ops/diffusion/modulate/ltx2_ada_values_triton.py:145

    tl.store(out0_ptr + base, (table0 + temb0).to(tl.bfloat16), mask=mask)
    tl.store(out1_ptr + base, (table1 + temb1).to(tl.bfloat16), mask=mask)
    tl.store(out2_ptr + base, (table2 + temb2).to(tl.bfloat16), mask=mask)
    tl.store(out3_ptr + base, (table3 + temb3).to(tl.bfloat16), mask=mask)
    tl.store(out4_ptr + base, (table4 + temb4).to(tl.bfloat16), mask=mask)
    tl.store(out5_ptr + base, (table5 + temb5).to(tl.bfloat16), mask=mask)
    tl.store(out6_ptr + base, (table6 + temb6).to(tl.bfloat16), mask=mask)
    tl.store(out7_ptr + base, (table7 + temb7).to(tl.bfloat16), mask=mask)
    tl.store(out8_ptr + base, (table8 + temb8).to(tl.bfloat16), mask=mask)


def ltx2_ada_values9(
    scale_shift_table: torch.Tensor,
    timestep: torch.Tensor,
) -> tuple[torch.Tensor, ...]:
    if timestep.ndim != 3:
        raise ValueError("timestep must have shape [B, S, 9 * D]")
    if not timestep.is_cuda or timestep.dtype != torch.bfloat16:
        raise ValueError("timestep must be a CUDA bfloat16 tensor")
    if not timestep.is_contiguous():
        raise ValueError("timestep must be contiguous")
    if scale_shift_table.ndim != 2 or scale_shift_table.shape[0] != 9:
        raise ValueError("scale_shift_table must have shape [9, D]")
    if (
        not scale_shift_table.is_cuda
        or scale_shift_table.dtype not in (torch.bfloat16, torch.float32)
        or scale_shift_table.stride(-1) != 1
    ):
        raise ValueError(
            "scale_shift_table must be CUDA, bf16/fp32, last-dim contiguous"
        )

    total_params = int(scale_shift_table.shape[0])
    hidden = int(scale_shift_table.shape[1])
    if hidden <= 0 or timestep.shape[-1] != total_params * hidden:
        raise ValueError("timestep last dim must equal 9 * hidden")
    if hidden % 256 != 0 or hidden > 8192:

View on GitHub (pinned to 0132848349)

Solutions

1. Cast: timestep = timestep.to(torch.bfloat16).cuda()
2. Align the timestep projection layer to bf16 output
3. Construct test inputs with device='cuda', dtype=torch.bfloat16
4. If fp32 is required, use the non-fused reference implementation

Example fix

# before
vals = ltx2_ada_values9(table, t_emb_fp32)
# after
t_emb = t_emb_fp32.to(device='cuda', dtype=torch.bfloat16)
vals = ltx2_ada_values9(table, t_emb)

timestep = timestep.to(device='cuda', dtype=torch.bfloat16)
assert timestep.is_cuda and timestep.dtype == torch.bfloat16

def t_ok(t: torch.Tensor) -> bool:
    return t.is_cuda and t.dtype == torch.bfloat16

When it happens

Trigger: Passing a float32 timestep embedding (common when the projection layer outputs fp32), a fp16 tensor, or a CPU tensor in tests.

Common situations: Models kept in fp32 or fp16 precision; tests that create embeddings without device='cuda' or dtype=torch.bfloat16; upstream components returning fp32 projections that weren't cast.

Related errors

• SplitKV partial output (mO) must be Float32
• All tensors must have the same data type
• Only Float16 or BFloat16 is supported
• LSE tensor must be Float32
• cu_seqlens_q tensor must be Int32

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