sgl-project/sglang · error · TypeError

Type mismatch: {self.q_dtype} != {self.v_dtype}

Error message

Type mismatch: {self.q_dtype} != {self.v_dtype}

What it means

The kernel validates that Q and V share the same dtype, since the attention output accumulates in Q's type from V operands. The tensors passed had differing dtypes.

Source

Thrown at python/sglang/kernels/ops/attention/flash_attn/cute/sm100_hd256_2cta_fmha_forward.py:437

            )

        self.q_major_mode = utils.LayoutEnum.from_tensor(q).mma_major_mode()
        self.k_major_mode = utils.LayoutEnum.from_tensor(k).mma_major_mode()
        self.v_major_mode = utils.LayoutEnum.from_tensor(v).mma_major_mode()
        self.o_layout = utils.LayoutEnum.from_tensor(o)

        if cutlass.const_expr(self.q_major_mode != tcgen05.OperandMajorMode.K):
            raise RuntimeError("The layout of q is not supported")
        if cutlass.const_expr(self.k_major_mode != tcgen05.OperandMajorMode.K):
            raise RuntimeError("The layout of k is not supported")
        if cutlass.const_expr(self.v_major_mode != tcgen05.OperandMajorMode.MN):
            raise RuntimeError("The layout of v is not supported")

        # check type consistency
        if cutlass.const_expr(self.q_dtype != self.k_dtype):
            raise TypeError(f"Type mismatch: {self.q_dtype} != {self.k_dtype}")
        if cutlass.const_expr(self.q_dtype != self.v_dtype):
            raise TypeError(f"Type mismatch: {self.q_dtype} != {self.v_dtype}")
        self._setup_attributes()

        cta_group = tcgen05.CtaGroup.TWO
        # the intermediate tensor p is from tmem & k-major
        p_source = tcgen05.OperandSource.TMEM
        p_major_mode = tcgen05.OperandMajorMode.K
        qk_tiled_mma = sm100_utils.make_trivial_tiled_mma(
            self.q_dtype,
            self.q_major_mode,
            self.k_major_mode,
            self.qk_acc_dtype,
            cta_group,
            self.qk_mma_tiler[:2],
        )
        pv_tiled_mma = sm100_utils.make_trivial_tiled_mma(
            self.v_dtype,
            p_major_mode,
            self.v_major_mode,

View on GitHub (pinned to 0132848349)

Solutions

1. Match v's dtype to q's (cast or use FP8-aware kernel)
2. Verify kv_cache_dtype server arg is compatible with the selected attention backend

Example fix

# before
out = fmha(q_bf16, k_bf16, v_fp8)
# after
out = fmha(q_bf16, k_bf16, v_bf16)

if q.dtype != v.dtype:
    v = v.to(q.dtype)

When it happens

Trigger: Passing v with a different dtype than q (e.g. FP8 V cache with bf16 q) to the hd256 forward.

Common situations: FP8 KV cache with non-FP8 kernel path; upgrading to FP8 without switching to the FP8 kernel variant.

Related errors

• The layout of k is not supported
• The layout of v is not supported
• Type mismatch: {self.q_dtype} != {self.k_dtype}
• v_cache must be provided
• k_cache can only be None when only_qv=True

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