{"record":{"id":"7650d32f412da573","repo":"jax-ml/jax","slug":"stored-array-has-shape-value-shape-but-tmem-has","errorCode":null,"errorMessage":"Stored array has shape {value.shape}, but TMEM has shape {self.shape}","messagePattern":"Stored array has shape (.+?), but TMEM has shape (.+?)","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/experimental/mosaic/gpu/tcgen05.py","lineNumber":1461,"sourceCode":"      # that we didn't forget to handle reduce in any of the cases.\n      assert reduced_reg is None\n      return result\n    reduced_layout = layout.reduce((len(layout.base_tile_shape) - 1,))\n    assert reduced_layout.vector_length == 1\n    reduced_regs_shape = reduced_layout.registers_shape(self.shape[:-1])\n    assert math.prod(reduced_regs_shape) == 1\n    reduced_result = fa.FragmentedArray(\n        _registers=np.asarray(reduced_reg, dtype=object).reshape(reduced_regs_shape),\n        _layout=reduced_layout,\n        _is_signed=is_signed,\n    )\n    return result, reduced_result\n\n  def store(self, value: fa.FragmentedArray):\n    if not isinstance(value, fa.FragmentedArray):\n      raise TypeError(f\"TMEM stores expect a FragmentedArray, got: {value}\")\n    if value.shape != self.shape:\n      raise ValueError(\n          f\"Stored array has shape {value.shape}, but TMEM has shape\"\n          f\" {self.shape}\"\n      )\n    if value.mlir_dtype != self.dtype:\n      raise ValueError(\n          f\"Stored array has dtype {value.mlir_dtype}, but TMEM has dtype\"\n          f\" {self.dtype}\"\n      )\n    if not isinstance(value.layout, fa.TiledLayout):\n      raise TypeError(f\"Stored array has layout {value.layout}, but TMEM stores expect a TiledLayout\")\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 stores unsupported.\")","sourceCodeStart":1443,"sourceCodeEnd":1479,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/mosaic/gpu/tcgen05.py#L1443-L1479","documentation":"TensorMem.store requires the stored FragmentedArray to have exactly the same logical shape as the TMEM allocation. A shape mismatch means the allocation does not cover the data (or covers more), so the store is rejected.","triggerScenarios":"Allocating TMEM of shape (128, N) but storing a FragmentedArray of shape (64, N) or (128, N*2), e.g. after slicing or reshaping intermediate results.","commonSituations":"Splitting a computation over TMEM halves and forgetting to slice the array; changing tile sizes without resizing TMEM allocations.","solutions":["Make value.shape match self.shape exactly (slice or pad the FragmentedArray)","Re-allocate TMEM with the shape of the value you produce","Assert shapes match before the store in kernel scaffolding"],"exampleFix":"// before\ntmem_128x64.store(arr_64x64)\n// after\ntmem = TensorMem.alloc(shape=arr.shape, ...)\ntmem.store(arr)","handlingStrategy":"validation","validationCode":"if value.shape != tmem.shape:\n    raise ValueError(f'{value.shape} vs {tmem.shape}')","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Derive TMEM allocation shapes from the producing computation","Add shape assertions in kernel scaffolding before stores"],"tags":["mosaic","gpu","tcgen05","shape-mismatch","tensor-memory","jax"],"backgroundTag":"shape-mismatch","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}