{"record":{"id":"e6a716d91a69ef13","repo":"jax-ml/jax","slug":"tcgen05-cp-only-supports-fully-packed-tmem-referen","errorCode":null,"errorMessage":"tcgen05.cp only supports fully packed TMEM references (packing={32 // bitwidth}), but got packing={tmem_ref.packing}","messagePattern":"tcgen05\\.cp only supports fully packed TMEM references \\(packing=(.+?)\\), but got packing=(.+?)","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/experimental/mosaic/gpu/tcgen05.py","lineNumber":2009,"sourceCode":"    )\n\n\ndef async_copy_smem_to_tmem(\n    smem_ref: ir.Value,\n    tmem_ref: TMEMRef,\n    swizzle: int,\n    collective: bool = False,\n) -> None:\n  i8 = ir.IntegerType.get_signless(8)\n  i32 = ir.IntegerType.get_signless(32)\n  smem_ty = ir.MemRefType(smem_ref.type)\n  if (dtype := smem_ty.element_type) != tmem_ref.dtype:\n    raise ValueError(f\"Incompatible dtypes: SMEM has {dtype}, TMEM has {tmem_ref.dtype}\")\n  if swizzle not in {16, 32, 64, 128}:\n    raise ValueError(f\"Unsupported swizzle, expected 16, 32, 64 or 128, but got: {swizzle}\")\n  bitwidth = utils.bitwidth(dtype)\n  if tmem_ref.packing != 32 // bitwidth:\n    raise ValueError(\n        \"tcgen05.cp only supports fully packed TMEM references\"\n        f\" (packing={32 // bitwidth}), but got packing={tmem_ref.packing}\"\n    )\n  if tmem_ref.shape[0] != TMEM_ROWS:\n    raise ValueError(\n        f\"TMEM reference must have {TMEM_ROWS} rows, but got {tmem_ref.shape[0]}\"\n    )\n  if tmem_ref.layout != tmem_default_layout(packing=tmem_ref.packing):\n    raise ValueError(\n        f\"Only standard TMEM layout is supported, got: {tmem_ref.layout}\"\n    )\n  swizzle_elems = 8 * swizzle // bitwidth\n  expected_smem_shape = utils.tile_shape(tmem_ref.shape, (8, swizzle_elems))\n  smem_shape = tuple(smem_ty.shape)\n  if smem_shape != expected_smem_shape:\n    raise ValueError(\n        f\"SMEM has shape {smem_shape}, but expected {expected_smem_shape} for\"\n        f\" TMEM shape {tmem_ref.shape} with swizzle={swizzle}\"","sourceCodeStart":1991,"sourceCodeEnd":2027,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/mosaic/gpu/tcgen05.py#L1991-L2027","documentation":"Raised by async_copy_smem_to_tmem when the TMEM reference's packing factor does not equal 32/bitwidth. tcgen05.cp requires narrow types (f16, bf16, f8, i8...) to be fully packed into 32-bit TMEM lanes; a partially packed or unpacked TMEM reference cannot be a destination of this instruction.","triggerScenarios":"Calling tcgen05.async_copy_smem_to_tmem with a tmem_ref whose packing attribute (from tmem_alloc or a tmem slice) is not 32/bitwidth of the element dtype — e.g. packing=1 with f16 elements (requires packing=2), or a tmem ref refined via tmem_ref with a non-default packing.","commonSituations":"Allocating TMEM without specifying packing for sub-32-bit types, or slicing/re-laying-out a TMEM reference with explicit_packing and passing the result to the SMEM-to-TMEM copy.","solutions":["Allocate the TMEM reference with fully packed layout: packing = 32 // bitwidth(dtype) (e.g. 2 for f16/bf16, 4 for i8/f8)","Use tmem_default_layout/packing helpers rather than hand-computing packing","Verify no intermediate op (slice, cast) reset packing to 1"],"exampleFix":"// before\ntmem = tmem_alloc(f16, shape, packing=1)\n// after\ntmem = tmem_alloc(f16, shape, packing=32 // 16)  # packing=2","handlingStrategy":"validation","validationCode":"bw = utils.bitwidth(dtype)\nrequired_packing = 32 // bw\nassert tmem_ref.packing == required_packing, f'need packing={required_packing}'","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Compute packing as 32 // bitwidth(dtype) whenever allocating TMEM for narrow types","Centralize TMEM allocation in one helper that enforces full packing"],"tags":["gpu","mosaic","tcgen05","packing","tensor-memory"],"backgroundTag":"invalid-layout-configuration","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}