{"record":{"id":"b2b9713714b877ad","repo":"jax-ml/jax","slug":"data-type-bitcast-is-only-supported-for-contiguous","errorCode":null,"errorMessage":"Data type bitcast is only supported for contiguous 1D arrays, but got stride={stride}.","messagePattern":"Data type bitcast is only supported for contiguous 1D arrays, but got stride=(.+?)\\.","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/lowering.py","lineNumber":1492,"sourceCode":"  Raises:\n    ValueError: if the source ref is not in SMEM.\n  \"\"\"\n  if src_dtype == dst_dtype:\n    return ref\n  if src_dtype != ir.IntegerType.get_signless(8):\n    raise NotImplementedError(\n        \"Data type bitcast is only supported from i8 to other types.\"\n    )\n  ref_ty = ir.MemRefType(ref.type)\n  if not mgpu_utils.is_smem_ref(ref_ty):\n    raise ValueError(f\"Only workgroup memory is supported but got {ref}.\")\n  if len(ref_ty.shape) != 1:\n    raise NotImplementedError(\n        \"Data type bitcast is only supported for 1D arrays.\"\n    )\n  [stride], _ = ref_ty.get_strides_and_offset()\n  if stride != 1:\n    raise ValueError(\n        \"Data type bitcast is only supported for contiguous 1D arrays, but got \"\n        f\"stride={stride}.\"\n    )\n  [shape_bytes] = ref_ty.shape\n  shape_bitwidth = shape_bytes * 8\n  target_bitwidth = mgpu_utils.bitwidth(dst_dtype)\n\n  if shape_bitwidth % target_bitwidth:\n    raise ValueError(\n        f\"Can not bitcast memory region of size {shape_bitwidth} bits to dtype \"\n        f\"with {target_bitwidth} bits.\"\n    )\n\n  result_type = ir.MemRefType.get(\n      shape=(shape_bitwidth // target_bitwidth,),\n      element_type=dst_dtype,\n      memory_space=ref_ty.memory_space,\n  )","sourceCodeStart":1474,"sourceCodeEnd":1510,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/lowering.py#L1474-L1510","documentation":"Bitcasting bytes to a wider dtype requires physically contiguous memory, because the lowering produces a new memref with stride 1 in the target dtype. If the source i8 memref has a stride other than 1 (e.g. a strided slice of SMEM), the reinterpretation would be incorrect and is rejected.","triggerScenarios":"Aliasing a non-contiguous (strided) i8 SMEM ref with a different dtype — e.g. taking a column slice or every-other-element view then calling .view(dtype).","commonSituations":"Slicing scratch buffers before reinterpretation; strided views produced by block transformations during lowering of aliased Refs.","solutions":["Materialize a contiguous copy of the data before bitcasting","Adjust slicing so the resulting memref is contiguous (slice the innermost dim fully)","Use gather/scatter loads instead of memory reinterpretation"],"exampleFix":"// before\ns = buf[::2]           # stride 2\nv = s.view(jnp.float32)\n// after\ns = buf[:n]            # contiguous\nv = s.view(jnp.float32)","handlingStrategy":"validation","validationCode":"import numpy as np\nassert all(s == 0 or np.prod(buf.shape[max(i+1,0):]) == buf.shape[-1] for i, s in enumerate(getattr(buf, 'strides', (1,)))), 'ref must be contiguous before view'","typeGuard":"def is_contiguous_1d(buf) -> bool:\n    return len(buf.shape) == 1 and (getattr(buf, 'strides', (1,))[0] in (1, None))","tryCatchPattern":null,"preventionTips":["Avoid strided slices of buffers you intend to bitcast","Copy into a fresh contiguous buffer when in doubt","Keep reinterpretation on the innermost contiguous axis"],"tags":["jax","pallas","mosaic-gpu","bitcast","stride","contiguity"],"backgroundTag":"non-contiguous-tensor-operation","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}