{"record":{"id":"4e091e18f9bce71b","repo":"jax-ml/jax","slug":"unsupported-dtype-for-reduction-self-dtype","errorCode":null,"errorMessage":"Unsupported dtype for reduction: {self.dtype}","messagePattern":"Unsupported dtype for reduction: (.+?)","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/experimental/mosaic/gpu/tcgen05.py","lineNumber":1352,"sourceCode":"      else:\n        raise ValueError(f\"TMEM layout {self.layout} is not supported\")\n    if reduce is not None:\n      if isinstance(self.dtype, ir.IntegerType) and bitwidth == 32:\n        if reduce not in (\"min\", \"max\"):\n          raise ValueError(\n              \"Unsupported reduction for i32. Only min and max are supported,\"\n              f\" got: {reduce}\"\n          )\n        if not is_signed:\n          reduce = \"abs\" + reduce  # type: ignore\n      elif isinstance(self.dtype, ir.F32Type):\n        if reduce not in (\"min\", \"max\", \"absmin\", \"absmax\"):\n          raise ValueError(\n              \"Unsupported reduction for f32. Only min, max, absmin, and\"\n              f\" absmax are supported, got: {reduce}\"\n          )\n      else:\n        raise ValueError(f\"Unsupported dtype for reduction: {self.dtype}\")\n\n    has_default_layout = self.layout == tmem_default_layout(packing)\n    regs_shape = layout.registers_shape(self.shape)\n    # TODO(olechwierowicz): `sparse_meta_layout()` does not really describe the\n    # actual TMEM layout of the result of `async_copy_sparse_smem_to_tmem`.\n    # As a result storing through SMEM -> Reg -> TMEM is not equivalent to\n    # SMEM -> TMEM. We raise in this case to prevent inconsistent behaviour.\n    # This restriction can be lifted if `TiledLayout` supports multiple\n    # vector dims.\n    if self.layout == sparse_meta_layout():\n      raise NotImplementedError(\"Sparse meta layout loads unsupported.\")\n    if regs_shape[0] != 1:  # We'll need to issue multiple loads below.\n      raise NotImplementedError(\"Loading multiple row tiles\")\n    if (\n        layout == LAYOUT\n        and self.layout == tmem_default_layout(packing)\n        and is_at_least_16b\n    ):","sourceCodeStart":1334,"sourceCodeEnd":1370,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/mosaic/gpu/tcgen05.py#L1334-L1370","documentation":"Raised by TensorMem.load when a fused load-reduce is requested but the tensor memory dtype is neither an integer type nor f32. The reduction path only knows how to lower reductions for integer and f32 element types.","triggerScenarios":"Calling tmem.load(..., reduce=...) where the TensorMem was allocated with a dtype like f16, bf16, or f64 — anything other than an ir.IntegerType or ir.F32Type.","commonSituations":"Allocating TMEM with half-precision dtypes (common in attention kernels) and then attempting a fused load-reduce; recent dtype additions to Mosaic that lack reduction support.","solutions":["Drop the reduce argument and perform the reduction manually in registers","Redo arithmetic so the reduced tensor lives in f32 or an integer dtype","Check isinstance(tmem.dtype, (ir.IntegerType, ir.F32Type)) before requesting a reduction"],"exampleFix":"// before\narr = tmem.load(layout, reduce='max')  # tmem dtype is f16\n// after\narr = tmem.load(layout)\nred = mx.maximum(arr, axis=...)  # explicit reduction","handlingStrategy":"validation","validationCode":"from mlir import ir\nif reduce is not None and not isinstance(tmem.dtype, (ir.IntegerType, ir.F32Type)):\n    raise ValueError('fused load-reduce needs integer or f32 dtype')","typeGuard":"def can_fused_reduce(tmem) -> bool:\n    return isinstance(tmem.dtype, (ir.IntegerType, ir.F32Type))","tryCatchPattern":"try:\n    arr, red = tmem.load(layout, reduce=reduce)\nexcept ValueError:\n    arr, _ = tmem.load(layout)\n    red = reduce_in_registers(arr, reduce)","preventionTips":["Allocate reduced results in f32 or integer TMEM","Centralize dtype policy for reductions in kernel config"],"tags":["mosaic","gpu","tcgen05","dtype","reduction","jax"],"backgroundTag":"unsupported-dtype","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}