sgl-project/sglang · error · ValueError

Unsupported type {type(data)}

Error message

Unsupported type {type(data)}

What it means

bf16x2_to_fp32x2 in the CuteDSL conversion utilities only accepts packed 16-bit pair containers (e.g. Uint32 holding two bf16 values). When the loop tries to convert each element of `data` and an element is not one of the supported packed types, the helper raises ValueError('Unsupported type {type(data)}').

Source

Thrown at python/sglang/kernels/ops/attention/cute_utils/cvt.py:50

            is_align_stack=False,
            loc=loc,
            ip=ip,
        )
        return (
            Float32(llvm.extractvalue(T.f32(), out, [0], loc=loc, ip=ip)),
            Float32(llvm.extractvalue(T.f32(), out, [1], loc=loc, ip=ip)),
        )

    elif isinstance(data, (cute.Tensor, cute.TensorSSA)):
        # NOTE: the output is always 1D
        size = cute.size(data.shape)
        out = cute.make_rmem_tensor(size * 2, Float32)
        for i in range(size):
            out[i * 2], out[i * 2 + 1] = bf16x2_to_fp32x2(data[i])
        return out

    else:
        raise ValueError(f"Unsupported type {type(data)}")


@dsl_user_op
def fp8x4_to_bf16x4(x: Uint32, *, loc=None, ip=None) -> cute.TensorSSA:
    # there is only fp8->fp16 conversion, hence we need to go
    # round trip through fp16.
    out = llvm.inline_asm(
        llvm.StructType.get_literal([T.i32()] * 2),
        [x.ir_value(loc=loc, ip=ip)],
        "{\n\t"
        ".reg .b16 x0, x1;\n\t"
        ".reg .b16 t00, t01, t10, t11;\n\t"
        "mov.b32 {x0, x1}, $2;\n\t"
        "cvt.rn.f16x2.e4m3x2 $0, x0;\n\t"
        "cvt.rn.f16x2.e4m3x2 $1, x1;\n\t"
        "mov.b32 {t00, t01}, $0;\n\t"
        "mov.b32 {t10, t11}, $1;\n\t"
        "cvt.rn.bf16.f16 t00, t00;\n\t"

View on GitHub (pinned to 0132848349)

Solutions

1. Ensure `data` elements are packed bf16x2 (typically Uint32) before calling; pack two bf16 values with cute.nv_cast or the packing helper
2. If input is already fp32, skip the conversion entirely
3. Reproduce with a tiny unit test that constructs Uint32 from two bf16 halves and confirm the helper returns Float32 pairs
4. Check the upstream kernel (e.g. in the same ops/attention module) for the canonical usage pattern and mirror its dtype flow

Example fix

# before
out = bf16x2_to_fp32x2(some_float32_tensor)  # ValueError: Unsupported type
# after
packed = cute.nv_cast(some_bf16_tensor, Uint32)  # two bf16 per element
out = bf16x2_to_fp32x2(packed)

from cute import Uint32
assert data.element_type == Uint32 or is_packed_bf16x2(data), f'bf16x2_to_fp32x2 needs packed bf16x2, got {data.element_type}'

def is_packed_bf16x2(x) -> bool:
    # packed pairs are stored in a 32-bit container
    return str(x.element_type) in ('Uint32', 'uint32', 'bfloat162')

When it happens

Trigger: Calling the vectorized wrapper with a tensor/SSA whose per-element type is not the packed bf16x2 (Uint32/ equivalent) type — e.g. passing Float32, BFloat16 (unpacked), or a wrong-width integer TensorSSA; also triggered from downstream kernels that feed mismatched dtypes into this helper.

Common situations: Writing a new CuteDSL kernel and wiring the wrong dtype into the conversion helper; refactoring a kernel to change element type without updating the conversion call; passing a scalar or a rmem tensor created with the wrong dtype string.

Related errors

• All tensors must have the same data type
• Reserved serve backend names cannot be used: {names}
• num_heads must be divisible by num_epi_subtiles
• num_heads // num_epi_subtiles must be divisible by 4 (FMA un
• Unexpected initial_state_source shape: {initial_state_source

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