{"record":{"id":"e3a28b4bd844d441","repo":"jax-ml/jax","slug":"unsupported-packing-self-packing","errorCode":null,"errorMessage":"Unsupported packing: {self.packing}","messagePattern":"Unsupported packing: (.+?)","errorType":"exception","errorClass":"NotImplementedError","httpStatus":null,"severity":"error","filePath":"jax/experimental/mosaic/gpu/tcgen05.py","lineNumber":1540,"sourceCode":"  def _debug_print(self) -> None:\n    i32 = ir.IntegerType.get_signless(32)\n    num_cols = self.layout.cols_in_shape(self.shape, utils.bitwidth(self.dtype))\n    lane = arith.remui(utils.thread_idx(), arith.constant(i32, utils.WARPGROUP_SIZE))\n    for c in range(num_cols):\n      ptr = _tmem_addr_to_ptr(arith.addi(self.address, arith.constant(i32, c)))\n      i32_vec = ir.VectorType.get((1,), i32)\n      vec_val = nvvm.tcgen05_ld(i32_vec, nvvm.Tcgen05LdStShape.SHAPE_32X32B, ptr)\n      val = vector.extract(vec_val, [], [0])\n      dtype_bitwidth = utils.bitwidth(self.dtype)\n      full_packing = 32 // dtype_bitwidth\n      if self.packing == 1:\n        if dtype_bitwidth < 32:\n          val = arith.trunci(ir.IntegerType.get_signless(dtype_bitwidth), val)\n        val = utils.bitcast(val, self.dtype)\n      elif self.packing == full_packing:\n        val = utils.bitcast(val, ir.VectorType.get((full_packing,), self.dtype))\n      else:\n        raise NotImplementedError(f\"Unsupported packing: {self.packing}\")\n      # TODO(apaszke): Make this print logical, not physical location.\n      utils.debug_print(f\"[{{}}, {c}]: {{}}\", lane, val, uniform=False)\n\n\ndef _transfer_32xcols(\n    base_addr: ir.Value,\n    cols: int,\n    atom_shape: tuple[int, int],\n    tmem_packing: int,\n    reg_packing: int,\n) -> Iterator[tuple[ir.Value, int, int, slice]]:\n  \"\"\"Generates a sequence of parameters for a given TMEM read or write.\n\n  Arguments:\n    base_addr: The base address of the TMEM region.\n    cols: The number of logical columns to transfer.\n    atom_shape: The logical shape of the tile written by the warp in a single\n      TMEM transfer.","sourceCodeStart":1522,"sourceCodeEnd":1558,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/mosaic/gpu/tcgen05.py#L1522-L1558","documentation":"In _debug_print, TMEM words are unpacked according to self.packing; only the unpacking path (packing < full) and the full-packing bitcast are implemented. Any other packing value (e.g. packing inconsistent with the dtype) raises NotImplementedError.","triggerScenarios":"Calling tmem._debug_print() on a TensorMem whose packing is neither less than full_packing nor equal to full_packing — typically a mis-constructed allocation.","commonSituations":"Debugging kernels with unusual packing values; packing computed as 32//bitwidth producing values outside handled cases for exotic dtypes.","solutions":["Fix the packing used at allocation so it equals full_packing or a supported divisor","Print via a supported packing (store to a default-layout TMEM and debug_print that)","Avoid _debug_print for exotic configurations; inspect via SMEM round-trip instead"],"exampleFix":null,"handlingStrategy":"validation","validationCode":"full = 32 // utils.bitwidth(tmem.dtype)\nassert tmem.packing == full or tmem.packing < full, 'unsupported packing for debug_print'","typeGuard":null,"tryCatchPattern":null,"preventionTips":["_debug_print is a debugging aid; prefer it only on default allocations","Round-trip exotic data through a default-layout TMEM to inspect it"],"tags":["mosaic","gpu","tcgen05","debug","packing","not-implemented","jax"],"backgroundTag":"unsupported-operation-argument","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}