sgl-project/sglang · error · TypeError

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

Error message

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

What it means

The kernel validates that Q and K share the same dtype; tcgen05 MMA requires both operands of the same element type. The dtypes were derived from the passed tensors and differ.

Source

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

                self.cta_tiler,
                self.is_persistent,
            )

        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,

View on GitHub (pinned to 0132848349)

Solutions

1. Cast q to k's dtype (or vice versa) so both match
2. If using FP8 KV cache, use the FP8-specific kernel entry point rather than the generic one
3. Check server args for accidental mixed precision settings

Example fix

# before
out = fmha(q, k_fp8, v_fp8)
# after
out = fmha(q, k.to(q.dtype), v.to(q.dtype))  # or use the fp8 kernel variant

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

def dtypes_match(*tensors: torch.Tensor) -> bool:
    return len({t.dtype for t in tensors}) == 1

When it happens

Trigger: Passing q of one dtype (e.g. bfloat16) and k of another (e.g. float16 or float8) to hd256 2cta fmha forward.

Common situations: Mixing a bf16 query with an FP8 KV cache without a quantization-aware kernel variant; partial dtype conversion in a custom backend.

Related errors

• Type mismatch: {self.q_dtype} != {self.k_dtype}
• Type mismatch: {self.q_dtype} != {self.v_dtype}
• hd256 forward varlen expects k rank 3 or 5, got rank {k_rank
• hd256 forward non-varlen expects k rank 4 or 5, got rank {k_
• The layout of q is not supported

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