{"record":{"id":"3b57c512b45dd73b","repo":"jax-ml/jax","slug":"tmem-aliasing-only-supported-for-refs-with-the-sam","errorCode":null,"errorMessage":"TMEM aliasing only supported for Refs with the same first dimension, got {ref.shape[0]} != {transformed_shape[0]}.","messagePattern":"TMEM aliasing only supported for Refs with the same first dimension, got (.+?) != (.+?)\\.","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/lowering.py","lineNumber":1573,"sourceCode":") -> tuple[\n    ir.Value | tcgen05.TMEMRef,\n    state_types.AbstractRef,\n    Sequence[state_types.Transform],\n    Sequence[state_types.Transform],\n]:\n  # Looks for the first transform being an ExtractAliasedRef and pulls out the\n  # Ref there, updating the transforms.\n  match transforms:\n    case (\n        gpu_core.ExtractAliasedRef(dtype, transformed_shape, offset, alias_group_idx, layout) as t,\n        *other_transforms,\n    ):\n      ref_aval = t.transform_type(ref_aval)\n      mlir_dtype = mgpu_utils.dtype_to_ir_type(dtype)\n      if isinstance(ref, tcgen05.TMEMRef):\n        assert layout is not None\n        if ref.shape[0] != transformed_shape[0]:\n          raise ValueError(\n              \"TMEM aliasing only supported for Refs with the same first\"\n              f\" dimension, got {ref.shape[0]} != {transformed_shape[0]}.\"\n          )\n        address = arith_dialect.addi(ref.address, _i32_constant(offset))\n        ref = tcgen05.TMEMRef(\n            address=address,\n            shape=cast(tuple[int, int], transformed_shape),\n            dtype=mlir_dtype,\n            layout=layout,\n        )\n      else:\n        assert isinstance(ref, ir.Value)  # make pyrefly happy\n        input_ref_ty = ir.MemRefType(ref.type)\n        if input_ref_ty.memory_space == mgpu_utils.smem():\n          assert layout is None\n          ref_bits = math.prod(transformed_shape) * mgpu_utils.bitwidth(\n              mlir_dtype\n          )","sourceCodeStart":1555,"sourceCodeEnd":1591,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/lowering.py#L1555-L1591","documentation":"When aliasing tensor-memory (TMEM) refs, the lowering adjusts the base address by an offset while keeping the layout, which is only valid if the first dimension (the TMEM address/row dimension) is unchanged between the original ref and the transformed block. If ref.shape[0] != transformed_shape[0], aliasing would point at wrong rows.","triggerScenarios":"A Pallas kernel using TMEM (tcgen05) refs where a BlockMapping or transformation changes the first dimension of the block — e.g. transpose-like or reshaping transforms applied to TMEM aliased refs.","commonSituations":"Blackwell tcgen05 kernels with aliased accumulator buffers under block mappings that reshape the leading axis; migrating SMEM aliasing patterns to TMEM.","solutions":["Restructure the kernel so TMEM aliased refs keep the same leading dimension (apply the transform to other axes)","Move data to SMEM before applying shape-changing transforms","Split the kernel so the transform happens in a separate non-aliased step"],"exampleFix":"# before\n@plt_kernel  # BlockMapping transforms dim 0 of tmem ref\n# in_kernel: out_ref[...] aliases tmem_ref with transformed first dim\n# after: keep dim0 fixed, transform trailing dims\nout = tmem_ref[:, :k].view(...)  # dim0 unchanged","handlingStrategy":"validation","validationCode":"assert aliased.shape[0] == transformed.shape[0], 'TMEM alias must preserve the first dimension'","typeGuard":"def tmem_alias_ok(ref_shape, transformed_shape) -> bool:\n    return ref_shape[0] == transformed_shape[0]","tryCatchPattern":null,"preventionTips":["Keep block-mapping transforms off the leading TMEM dimension","Prefer SMEM for shape-changing reinterpretation","Add shape assertions in kernel setup for TMEM aliases"],"tags":["jax","pallas","tmem","tcgen05","aliasing","shape-mismatch"],"backgroundTag":"aliasing-shape-mismatch","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}