{"record":{"id":"659c330f470fa88b","repo":"jax-ml/jax","slug":"unsupported-transforms-a-sparse-metadata-transfo","errorCode":null,"errorMessage":"Unsupported transforms: {a_sparse_metadata_transforms}","messagePattern":"Unsupported transforms: (.+?)","errorType":"exception","errorClass":"NotImplementedError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/primitives.py","lineNumber":2837,"sourceCode":"    if b_scale_transforms:\n      raise NotImplementedError(f\"Unsupported transforms: {b_scale_transforms}\")\n  if a_sparse_metadata_transforms_tree is not None:\n    a_sparse_metadata_transforms = a_sparse_metadata_transforms_tree.unflatten(\n        a_sparse_metadata_transforms_leaves\n    )\n    a_sparse_metadata_transform_avals = (\n        a_sparse_metadata_transforms_tree.unflatten(\n            a_sparse_metadata_transforms_leaves_avals\n        )\n    )\n    assert isinstance(a_sparse_metadata_ref_aval, state_types.AbstractRef)\n    a_sparse_metadata_ref, _, a_sparse_metadata_transforms = (\n        lowering._handle_transforms(  # pyrefly: ignore[bad-specialization]\n            ctx, a_sparse_metadata_ref_aval, a_sparse_metadata_ref,\n            a_sparse_metadata_transform_avals, a_sparse_metadata_transforms)\n    )\n    if a_sparse_metadata_transforms:\n      raise NotImplementedError(\n          f\"Unsupported transforms: {a_sparse_metadata_transforms}\"\n      )\n\n  predicate = ctx.module_ctx.single_lane_predicate\n  if collective_axis is not None:\n    assert predicate is not None\n    is_leader_block = _collective_mma_predicate(ctx, collective_axis)\n    predicate = arith_dialect.andi(predicate, is_leader_block)\n    collective = True\n  else:\n    collective = False\n\n  with mgpu.when(predicate):\n    tcgen05.mma(\n        acc,\n        a_ref,\n        b_ref,\n        a_swizzle=int(lhs_swizzle),","sourceCodeStart":2819,"sourceCodeEnd":2855,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/primitives.py#L2819-L2855","documentation":"For sparse MMA on Blackwell, the sparse-metadata reference (index/end pointers) must be passed without residual layout transforms; _handle_transforms must consume all of them or lowering fails with NotImplementedError.","triggerScenarios":"Supplying a_sparse_metadata_ref with transform trees (transposes etc.) to tcgen05_mma such that transforms remain after handling.","commonSituations":"Writing 2:4 sparse GEMM kernels and routing metadata through transformed references.","solutions":["Pass the sparse metadata reference without transforms","Rearrange metadata layout manually before the MMA call"],"exampleFix":"// before\ntcgen05_mma(a_sparse, b, acc, a_sparse_metadata=meta_ref, a_sparse_metadata_transforms=transforms)\n// after\ntcgen05_mma(a_sparse, b, acc, a_sparse_metadata=meta_ref)","handlingStrategy":"validation","validationCode":"assert not a_sparse_metadata_transforms","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Pass sparse metadata without ref transforms","Check transforms survive _handle_transforms in tests"],"tags":["jax","pallas","tcgen05","sparse","not-implemented"],"backgroundTag":"unsupported-transform-on-operand","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}