{"record":{"id":"19b6a360721eb407","repo":"jax-ml/jax","slug":"can-not-bitcast-memory-region-of-size-shape-bitwi","errorCode":null,"errorMessage":"Can not bitcast memory region of size {shape_bitwidth} bits to dtype with {target_bitwidth} bits.","messagePattern":"Can not bitcast memory region of size (.+?) bits to dtype with (.+?) bits\\.","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/lowering.py","lineNumber":1501,"sourceCode":"  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  )\n\n  # Do a memref_ptr/ptr_as_memref roundtrip instead of using `memref.view`,\n  # which refuses to take in our source ref. This is because `memref.view` only\n  # works on a super restricted set of `memref`s. E.g., it does not work if an\n  # offset is specified, which can be the case for our SMEM refs.\n  return mgpu_utils.ptr_as_memref(mgpu_utils.memref_ptr(ref), result_type)\n\n\n@overload","sourceCodeStart":1483,"sourceCodeEnd":1519,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/lowering.py#L1483-L1519","documentation":"After bitcasting, the total size in bits of the memory region (num bytes * 8) must be evenly divisible by the target dtype's bitwidth, so the lowering can compute shape = shape_bits / target_bits. If the byte count doesn't divide evenly (e.g. 6 bytes to f32), the reshape is impossible and ValueError is raised.","triggerScenarios":"Viewing an i8 SMEM buffer whose length in bytes is not a multiple of the target dtype size — e.g. buf of length 6 viewed as f32 (4 bytes each), or odd-sized buffers viewed as i16/f32/f64.","commonSituations":"Dynamic block sizes that produce non-multiple buffer lengths; packing sub-byte types then reinterpreting at a width that doesn't divide the total.","solutions":["Round the buffer allocation up to a multiple of the target dtype size (pad and slice after view)","Choose a target dtype whose bitwidth divides the region size (e.g. view as i8/i16 instead of f32)","Assert/validate byte length divisibility in kernel setup"],"exampleFix":"# before\nbuf = alloc_smem((6,), jnp.uint8))\nv = buf.view(jnp.float32)  # 48 bits % 32 != 0\n# after\nbuf = alloc_smem((8,), jnp.uint8))\nv = buf.view(jnp.float32)  # 64 bits / 32 = 2 elements","handlingStrategy":"validation","validationCode":"n_bytes = int(buf.shape[0]) if buf.dtype == jnp.uint8 else None\ntarget_bits = jnp.dtype(target_dtype).itemsize * 8\nassert n_bytes is not None and (n_bytes * 8) % target_bits == 0, f'{n_bytes}B not divisible by {target_dtype}'","typeGuard":"def fits_bitcast(n_bytes: int, target: jnp.dtype) -> bool:\n    return (n_bytes * 8) % (jnp.dtype(target).itemsize * 8) == 0","tryCatchPattern":null,"preventionTips":["Allocate buffers in multiples of the widest target dtype size","Validate sizes once at kernel-launch config time","Pad dynamically-sized buffers to dtype multiples"],"tags":["jax","pallas","mosaic-gpu","bitcast","size-mismatch","alignment"],"backgroundTag":"size-mismatch-in-view","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}