{"record":{"id":"f0b91c5e05b05270","repo":"jax-ml/jax","slug":"ref-has-a-memory-space-mem-ref-ty-memory-space","errorCode":null,"errorMessage":"{ref} has a memory space {mem_ref_ty.memory_space} that is not TMEM.","messagePattern":"(.+?) has a memory space (.+?) that is not TMEM\\.","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/experimental/mosaic/gpu/dialect_lowering.py","lineNumber":2463,"sourceCode":"  with utils.when(ctx.single_warp_per_block_predicate):\n    tcgen05.tmem_dealloc(tmem_addr, ncols, collective, exact=False)\n\n  return []\n\n\ndef _tmem_ref_from_ir(\n    ref: ir.Value, expected_layout: ir.Attribute\n) -> tcgen05.TMEMRef:\n  \"\"\"Returns a TMEMRef from an IR value.\n\n  Throws an error if the annotated layout does not match the expected layout.\n  \"\"\"\n  if not isinstance(ref.type, ir.MemRefType):\n    raise ValueError(f\"{ref} is not a memref.\")\n  mem_ref_ty = ir.MemRefType(ref.type)\n\n  if mem_ref_ty.memory_space != utils.tmem():\n    raise ValueError(\n        f\"{ref} has a memory space {mem_ref_ty.memory_space} that is not TMEM.\"\n    )\n\n  i32 = ir.IntegerType.get_signless(32)\n  conversion_cast, [tmem_addr] = _undo_conversion_cast(ref, [i32])\n\n  assert mem_ref_ty.rank == 2\n  shape = cast(tuple[int, int], tuple(mem_ref_ty.shape))\n  el_ty = mem_ref_ty.element_type\n  layout_attr = conversion_cast.attributes[\"layout\"]\n  if layout_attr != expected_layout:\n    raise ValueError(\n        f\"{ref} has a layout {layout_attr} that does not match the expected\"\n        f\" layout {expected_layout}.\"\n    )\n  layout = layouts_lib.from_layout_attr(layout_attr)\n  assert isinstance(layout, fa.TiledLayout)\n  tmem_layout = tcgen05.TMEMLayout(","sourceCodeStart":2445,"sourceCodeEnd":2481,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/mosaic/gpu/dialect_lowering.py#L2445-L2481","documentation":"Mosaic GPU's TMEM lowering requires the operand memref's memory_space to equal utils.tmem(). This error means the value is a valid memref but lives in a different memory space (e.g. shared or global), so it cannot be treated as a tcgen05 TMEM reference.","triggerScenarios":"Passing an smem local_tensor or a global-tensor-derived memref into TMEM operands of tcgen05_mma / async_copy_smem_to_tmem / slice_tmem etc., where memory_space != tmem.","commonSituations":"Forgetting to set memory_space='tmem' on the allocation; using a shared-memory scratch buffer directly as an accumulator in a Blackwell MMA pipeline instead of copying it to TMEM first.","solutions":["Allocate the buffer with the TMEM memory space (e.g. mgpu.local_tensor(..., memory_space=utils.tmem()))","For data in smem use the async_copy_smem_to_tmem lowering path instead of feeding smem memrefs directly","Check any custom layout-cast / slice producers preserve the tmem memory space attribute"],"exampleFix":"// before\nacc = mgpu.local_tensor(shape, dtype, memory_space='shared')\n// after\nfrom jax.experimental.mosaic.gpu import utils\nacc = mgpu.local_tensor(shape, dtype, memory_space=utils.tmem())","handlingStrategy":"validation","validationCode":"from jaxlib.mlir import ir\nfrom jax.experimental.mosaic.gpu import utils\ndef in_tmem(v) -> bool:\n    t = v.type\n    return isinstance(t, ir.MemRefType) and t.memory_space == utils.tmem()","typeGuard":"def is_tmem_ref(v) -> bool:\n    from jaxlib.mlir import ir\n    return isinstance(v.type, ir.MemRefType) and str(v.type.memory_space) == '#gpu.memory_space<tmem>'","tryCatchPattern":null,"preventionTips":["Centralize TMEM allocation in one helper that always sets memory_space=utils.tmem()","Copy smem->tmem via the provided async copy path instead of direct operand reuse"],"tags":["mosaic","gpu","tmem","memory-space","mlir"],"backgroundTag":"wrong-memory-space","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}