{"record":{"id":"edb849f5f9eab24f","repo":"xai-org/x-algorithm","slug":"type-mismatch-self-q-dtype-self-k-dtype","errorCode":null,"errorMessage":"Type mismatch: {self.q_dtype} != {self.k_dtype}","messagePattern":"Type mismatch: (.+?) != (.+?)","errorType":"validation","errorClass":"TypeError","httpStatus":null,"severity":"error","filePath":"phoenix/xrex/cutedsl/ranker_fa4/flash_fwd_sm100.py","lineNumber":490,"sourceCode":"                [2, 4, 3, 1, 0] if const_expr(mCuSeqlensQ is None) else [1, 3, 2, 0]\n            )\n            LSE_layout_transpose = [3, 2, 1, 0] if const_expr(mCuSeqlensQ is None) else [2, 1, 0]\n            num_splits = mO.shape[0]\n        else:\n            O_layout_transpose = [1, 3, 2, 0] if const_expr(mCuSeqlensQ is None) else [0, 2, 1]\n            LSE_layout_transpose = [2, 1, 0] if const_expr(mCuSeqlensQ is None) else [1, 0]\n            num_splits = Int32(1)\n        mO = cute.make_tensor(mO.iterator, cute.select(mO.layout, mode=O_layout_transpose))\n        mLSE = (\n            cute.make_tensor(mLSE.iterator, cute.select(mLSE.layout, mode=LSE_layout_transpose))\n            if const_expr(mLSE is not None)\n            else None\n        )\n        V_layout_transpose = [1, 0, 2, 3] if const_expr(mCuSeqlensK is None) else [1, 0, 2]\n        mV = cute.make_tensor(mV.iterator, cute.select(mV.layout, mode=V_layout_transpose))\n\n        if const_expr(self.q_dtype != self.k_dtype):\n            raise TypeError(f\"Type mismatch: {self.q_dtype} != {self.k_dtype}\")\n        if const_expr(self.q_dtype != self.v_dtype):\n            raise TypeError(f\"Type mismatch: {self.q_dtype} != {self.v_dtype}\")\n        if const_expr(self.q_dtype.width == 8):\n            paged_kv_non_tma = not self.use_tma_KV\n            if const_expr(self.head_dim_padded < 96):\n                fp8_regs = _FP8_SMALL_HDIM_REGS[paged_kv_non_tma]\n                self.num_regs_softmax = fp8_regs[\"num_regs_softmax\"]\n                self.num_regs_correction = fp8_regs[\"num_regs_correction\"]\n                self.num_regs_other = fp8_regs[\"num_regs_other\"]\n            else:\n                fp8_tune = _FP8_TUNING_CONFIG.get(\n                    (self.use_2cta_instrs, self.is_causal, self.head_dim_padded, self.is_sm103), {}\n                )\n                if const_expr(\"ex2_emu_freq\" in fp8_tune):\n                    self._tune = {**self._tune, **fp8_tune}\n                    self.enable_ex2_emu = self._tune[\"ex2_emu_freq\"] > 0\n                if const_expr(not paged_kv_non_tma and \"num_regs_softmax\" in fp8_tune):\n                    self.num_regs_softmax = fp8_tune[\"num_regs_softmax\"]","sourceCodeStart":472,"sourceCodeEnd":508,"githubUrl":"https://github.com/xai-org/x-algorithm/blob/24c60942c5c5fdad3a6addffb4c6e6d2f228f04f/phoenix/xrex/cutedsl/ranker_fa4/flash_fwd_sm100.py#L472-L508","documentation":"The SM100 flash-attention forward kernel requires Q, K, and V to share the same element type because the tcgen05 MMA and TMA paths are compiled for one dtype. Before launching the inner kernel, __call__ compares self.q_dtype with self.k_dtype and raises TypeError on mismatch.","triggerScenarios":"Calling the forward kernel with Q in one dtype (e.g. bfloat16) and K in another (e.g. float16 or fp8), so const_expr(q_dtype != k_dtype) is true at compile time.","commonSituations":"Mixing a bf16 Q embedding with fp16 KV projections, enabling fp8 KV cache while Q stays bf16, or a dataloader/config change that cast one tensor but not the others.","solutions":["Cast Q to the same dtype as K (usually the KV dtype is the reference): q = q.to(k.dtype)","Make all of Q/K/V consistent, e.g. all bfloat16 or all float16, in the model config before calling the kernel","If fp8 Q was intended, ensure K and V are also fp8 with the same variant (E4M3FN vs E5M2)","Add an assertion in the wrapper that q.dtype == k.dtype == v.dtype to fail early with a clearer message"],"exampleFix":"// before\nout = flash_attn(q, k, v)  # q: bf16, k: fp16\n// after\nq = q.to(k.dtype)\nout = flash_attn(q, k, v)","handlingStrategy":"type-guard","validationCode":"assert q.dtype == k.dtype, f'q {q.dtype} != k {k.dtype}'","typeGuard":"def same_dtype(*ts):\n    d = ts[0].dtype\n    return all(t.dtype == d for t in ts)","tryCatchPattern":null,"preventionTips":["Normalize dtypes at model boundaries so Q/K/V enter attention with one dtype","Assert dtype equality in wrappers before the expensive kernel launch"],"tags":["dtype","flash-attention","type-mismatch","cuda"],"backgroundTag":"dtype-mismatch","analyzedSha":"24c60942c5c5fdad3a6addffb4c6e6d2f228f04f","analyzedAt":"2026-08-28T11:40:14.686Z","schemaVersion":2},"datasetVersion":"2026-08-28T16:17:29.566Z"}