{"record":{"id":"8f73b15d0f2a5051","repo":"xai-org/x-algorithm","slug":"only-float16-or-bfloat16-is-supported","errorCode":null,"errorMessage":"Only Float16 or BFloat16 is supported","messagePattern":"Only Float16 or BFloat16 is supported","errorType":"validation","errorClass":"TypeError","httpStatus":null,"severity":"error","filePath":"phoenix/xrex/cutedsl/ranker_fa4/flash_bwd_postprocess.py","lineNumber":221,"sourceCode":"                major_mode_size=self.tile_hdim // wg_d_dQ,\n            )\n        else:\n            self.sdQ_layout = sm100_utils_basic.make_smem_layout_epi(\n                self.dtype, LayoutEnum.ROW_MAJOR, (self.tile_m, self.tile_hdim), 1\n            )\n\n    @cute.jit\n    def __call__(\n        self,\n        mdQaccum: cute.Tensor,\n        mdQ: cute.Tensor,\n        scale: cutlass.Float32,\n        mCuSeqlensQ: Optional[cute.Tensor],\n        mSeqUsedQ: Optional[cute.Tensor],\n        stream: cuda.CUstream = None,\n    ):\n        if const_expr(mdQ.element_type not in [cutlass.Float16, cutlass.BFloat16]):\n            raise TypeError(\"Only Float16 or BFloat16 is supported\")\n        if const_expr(mdQaccum is not None):\n            if const_expr(mdQaccum.element_type not in [cutlass.Float32]):\n                raise TypeError(\"dQaccum tensor must be Float32\")\n\n        mdQaccum, mdQ = [assume_tensor_aligned(t) for t in (mdQaccum, mdQ)]\n\n        self.tiled_mma = self._get_tiled_mma()\n        self._setup_attributes()\n\n        smem_size = max(\n            cute.size_in_bytes(cutlass.Float32, self.sdQaccum_layout),\n            cute.size_in_bytes(self.dtype, self.sdQ_layout),\n        )\n\n        if const_expr(mCuSeqlensQ is not None):\n            TileScheduler = SingleTileVarlenScheduler\n            num_head = mdQ.shape[1]\n            num_batch = mCuSeqlensQ.shape[0] - 1","sourceCodeStart":203,"sourceCodeEnd":239,"githubUrl":"https://github.com/xai-org/x-algorithm/blob/24c60942c5c5fdad3a6addffb4c6e6d2f228f04f/phoenix/xrex/cutedsl/ranker_fa4/flash_bwd_postprocess.py#L203-L239","documentation":"The flash backward postprocess kernel (__call__ in flash_bwd_postprocess.py) only accepts Float16 or BFloat16 for the Q output tensor mdQ. Other dtypes (FP32, FP8, TF32) are rejected at compile time via const_expr.","triggerScenarios":"Calling the postprocess kernel with mdQ in torch.float32 or float64 — e.g. keeping Q/dQ in fp32 for numerical debugging or because autocast was disabled.","commonSituations":"Debugging with fp32 tensors then forgetting to cast back; a pipeline stage upcasting dQ to fp32 before postprocess; mismatched autocast contexts.","solutions":["Cast mdQ to half or bfloat16 before the call: mdQ.to(torch.bfloat16)","Run under autocast/bf16 so downstream tensors are already the right dtype","Keep the fp32 copy separately if needed for loss scaling"],"exampleFix":"# before\nmdQ = dq_accum_result.to(torch.float32)\npost(mdQaccum, mdQ, scale, ...)\n\n# after\nmdQ = dq_accum_result.to(torch.bfloat16)   # or torch.float16\npost(mdQaccum, mdQ, scale, ...)","handlingStrategy":"type-guard","validationCode":"assert mdQ.dtype in (torch.float16, torch.bfloat16), mdQ.dtype","typeGuard":"def is_supported_q_dtype(t: torch.Tensor) -> bool:\n    return t.dtype in (torch.float16, torch.bfloat16)","tryCatchPattern":null,"preventionTips":["Run under autocast so outputs are bf16/fp16","Cast debug fp32 tensors back before kernel calls"],"tags":["dtype-validation","flash-attention","cuda","backward"],"backgroundTag":"unsupported-dtype","analyzedSha":"24c60942c5c5fdad3a6addffb4c6e6d2f228f04f","analyzedAt":"2026-08-28T11:40:14.686Z","schemaVersion":2},"datasetVersion":"2026-08-28T16:17:29.566Z"}