{"record":{"id":"7564a784c08865c1","repo":"xai-org/x-algorithm","slug":"dqaccum-tensor-must-be-float32","errorCode":null,"errorMessage":"dQaccum tensor must be Float32","messagePattern":"dQaccum tensor must be Float32","errorType":"validation","errorClass":"TypeError","httpStatus":null,"severity":"error","filePath":"phoenix/xrex/cutedsl/ranker_fa4/flash_bwd_postprocess.py","lineNumber":224,"sourceCode":"            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\n            num_block = cute.ceil_div(mdQ.shape[0], self.tile_m)\n        else:\n            TileScheduler = SingleTileScheduler","sourceCodeStart":206,"sourceCodeEnd":242,"githubUrl":"https://github.com/xai-org/x-algorithm/blob/24c60942c5c5fdad3a6addffb4c6e6d2f228f04f/phoenix/xrex/cutedsl/ranker_fa4/flash_bwd_postprocess.py#L206-L242","documentation":"When mdQaccum is provided to the flash backward postprocess kernel, it must be Float32 — the accumulator is kept in higher precision and downcast only at the end. Any other dtype fails the const_expr type check.","triggerScenarios":"Passing a bf16/fp16 dQ accumulation tensor to the postprocess kernel, e.g. allocating dQaccum with the same dtype as Q to save memory.","commonSituations":"Memory optimizations that shrink the accumulator; dtype propagation from autocast allocating bf16 buffers; adapting a pipeline that assumed fp16 accumulators.","solutions":["Allocate mdQaccum as torch.float32 (torch.zeros(..., dtype=torch.float32, device=...))","Don't let autocast dictate the accumulator dtype — create it explicitly","If memory-bound, keep fp32 accumulator but free it promptly after postprocess"],"exampleFix":"# before\ndqaccum = torch.zeros(q_shape, dtype=q.dtype, device='cuda')  # bf16\n\n# after\ndqaccum = torch.zeros(q_shape, dtype=torch.float32, device='cuda')","handlingStrategy":"type-guard","validationCode":"assert mdQaccum.dtype == torch.float32, mdQaccum.dtype","typeGuard":"def is_fp32(t: torch.Tensor) -> bool:\n    return t.dtype == torch.float32","tryCatchPattern":null,"preventionTips":["Allocate accumulators with explicit dtype=torch.float32","Never reuse a bf16 buffer as the dQ accumulator"],"tags":["dtype-validation","fp32-accumulator","flash-attention","backward"],"backgroundTag":"unsupported-dtype","analyzedSha":"24c60942c5c5fdad3a6addffb4c6e6d2f228f04f","analyzedAt":"2026-08-28T11:40:14.686Z","schemaVersion":2},"datasetVersion":"2026-08-28T16:17:29.566Z"}