{"record":{"id":"5905e03960cebbf9","repo":"sgl-project/sglang","slug":"type-mismatch-self-q-dtype-self-v-dtype-5905e0","errorCode":null,"errorMessage":"Type mismatch: {self.q_dtype} != {self.v_dtype}","messagePattern":"Type mismatch: (.+?) != (.+?)","errorType":"exception","errorClass":"TypeError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/attention/flash_attn/cute/sm100_hd256_2cta_fmha_forward.py","lineNumber":437,"sourceCode":"            )\n\n        self.q_major_mode = utils.LayoutEnum.from_tensor(q).mma_major_mode()\n        self.k_major_mode = utils.LayoutEnum.from_tensor(k).mma_major_mode()\n        self.v_major_mode = utils.LayoutEnum.from_tensor(v).mma_major_mode()\n        self.o_layout = utils.LayoutEnum.from_tensor(o)\n\n        if cutlass.const_expr(self.q_major_mode != tcgen05.OperandMajorMode.K):\n            raise RuntimeError(\"The layout of q is not supported\")\n        if cutlass.const_expr(self.k_major_mode != tcgen05.OperandMajorMode.K):\n            raise RuntimeError(\"The layout of k is not supported\")\n        if cutlass.const_expr(self.v_major_mode != tcgen05.OperandMajorMode.MN):\n            raise RuntimeError(\"The layout of v is not supported\")\n\n        # check type consistency\n        if cutlass.const_expr(self.q_dtype != self.k_dtype):\n            raise TypeError(f\"Type mismatch: {self.q_dtype} != {self.k_dtype}\")\n        if cutlass.const_expr(self.q_dtype != self.v_dtype):\n            raise TypeError(f\"Type mismatch: {self.q_dtype} != {self.v_dtype}\")\n        self._setup_attributes()\n\n        cta_group = tcgen05.CtaGroup.TWO\n        # the intermediate tensor p is from tmem & k-major\n        p_source = tcgen05.OperandSource.TMEM\n        p_major_mode = tcgen05.OperandMajorMode.K\n        qk_tiled_mma = sm100_utils.make_trivial_tiled_mma(\n            self.q_dtype,\n            self.q_major_mode,\n            self.k_major_mode,\n            self.qk_acc_dtype,\n            cta_group,\n            self.qk_mma_tiler[:2],\n        )\n        pv_tiled_mma = sm100_utils.make_trivial_tiled_mma(\n            self.v_dtype,\n            p_major_mode,\n            self.v_major_mode,","sourceCodeStart":419,"sourceCodeEnd":455,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/flash_attn/cute/sm100_hd256_2cta_fmha_forward.py#L419-L455","documentation":"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.","triggerScenarios":"Passing v with a different dtype than q (e.g. FP8 V cache with bf16 q) to the hd256 forward.","commonSituations":"FP8 KV cache with non-FP8 kernel path; upgrading to FP8 without switching to the FP8 kernel variant.","solutions":["Match v's dtype to q's (cast or use FP8-aware kernel)","Verify kv_cache_dtype server arg is compatible with the selected attention backend"],"exampleFix":"# before\nout = fmha(q_bf16, k_bf16, v_fp8)\n# after\nout = fmha(q_bf16, k_bf16, v_bf16)","handlingStrategy":"validation","validationCode":"if q.dtype != v.dtype:\n    v = v.to(q.dtype)","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Pair fp8 KV caches only with fp8-aware kernel variants","Add a dtype-consistency assert in the decode wrapper"],"tags":["cuda","flash-attention","dtype-mismatch","kv-cache"],"backgroundTag":"dtype-mismatch","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}