{"record":{"id":"023b50e8e2f7a3ea","repo":"sgl-project/sglang","slug":"unsupported-type-type-data","errorCode":null,"errorMessage":"Unsupported type {type(data)}","messagePattern":"Unsupported type (.+?)","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/attention/cute_utils/cvt.py","lineNumber":50,"sourceCode":"            is_align_stack=False,\n            loc=loc,\n            ip=ip,\n        )\n        return (\n            Float32(llvm.extractvalue(T.f32(), out, [0], loc=loc, ip=ip)),\n            Float32(llvm.extractvalue(T.f32(), out, [1], loc=loc, ip=ip)),\n        )\n\n    elif isinstance(data, (cute.Tensor, cute.TensorSSA)):\n        # NOTE: the output is always 1D\n        size = cute.size(data.shape)\n        out = cute.make_rmem_tensor(size * 2, Float32)\n        for i in range(size):\n            out[i * 2], out[i * 2 + 1] = bf16x2_to_fp32x2(data[i])\n        return out\n\n    else:\n        raise ValueError(f\"Unsupported type {type(data)}\")\n\n\n@dsl_user_op\ndef fp8x4_to_bf16x4(x: Uint32, *, loc=None, ip=None) -> cute.TensorSSA:\n    # there is only fp8->fp16 conversion, hence we need to go\n    # round trip through fp16.\n    out = llvm.inline_asm(\n        llvm.StructType.get_literal([T.i32()] * 2),\n        [x.ir_value(loc=loc, ip=ip)],\n        \"{\\n\\t\"\n        \".reg .b16 x0, x1;\\n\\t\"\n        \".reg .b16 t00, t01, t10, t11;\\n\\t\"\n        \"mov.b32 {x0, x1}, $2;\\n\\t\"\n        \"cvt.rn.f16x2.e4m3x2 $0, x0;\\n\\t\"\n        \"cvt.rn.f16x2.e4m3x2 $1, x1;\\n\\t\"\n        \"mov.b32 {t00, t01}, $0;\\n\\t\"\n        \"mov.b32 {t10, t11}, $1;\\n\\t\"\n        \"cvt.rn.bf16.f16 t00, t00;\\n\\t\"","sourceCodeStart":32,"sourceCodeEnd":68,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/cute_utils/cvt.py#L32-L68","documentation":"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)}').","triggerScenarios":"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.","commonSituations":"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.","solutions":["Ensure `data` elements are packed bf16x2 (typically Uint32) before calling; pack two bf16 values with cute.nv_cast or the packing helper","If input is already fp32, skip the conversion entirely","Reproduce with a tiny unit test that constructs Uint32 from two bf16 halves and confirm the helper returns Float32 pairs","Check the upstream kernel (e.g. in the same ops/attention module) for the canonical usage pattern and mirror its dtype flow"],"exampleFix":"# before\nout = bf16x2_to_fp32x2(some_float32_tensor)  # ValueError: Unsupported type\n# after\npacked = cute.nv_cast(some_bf16_tensor, Uint32)  # two bf16 per element\nout = bf16x2_to_fp32x2(packed)","handlingStrategy":"type-guard","validationCode":"from cute import Uint32\nassert data.element_type == Uint32 or is_packed_bf16x2(data), f'bf16x2_to_fp32x2 needs packed bf16x2, got {data.element_type}'","typeGuard":"def is_packed_bf16x2(x) -> bool:\n    # packed pairs are stored in a 32-bit container\n    return str(x.element_type) in ('Uint32', 'uint32', 'bfloat162')","tryCatchPattern":null,"preventionTips":["Construct inputs with the same dtype flow as existing kernels in ops/attention","Add dtype asserts at kernel entry in debug builds"],"tags":["cutedsl","dtype","type-mismatch","kernel"],"backgroundTag":"dtype-mismatch","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}