{"record":{"id":"299e74244cdfa23a","repo":"jax-ml/jax","slug":"arg-ty","errorCode":null,"errorMessage":"{arg_ty}","messagePattern":"\\{arg_ty\\}","errorType":"exception","errorClass":"NotImplementedError","httpStatus":null,"severity":"error","filePath":"jax/experimental/mosaic/gpu/fragmented_array.py","lineNumber":1985,"sourceCode":"        for i in range(vec_len):\n          vs = [\n              vector.extract(\n                  a,\n                  dynamic_position=[],\n                  static_position=ir.DenseI64ArrayAttr.get([i]),\n              )\n              for a in args\n          ]\n          vr = fast_instr(*vs)\n          result = vector.insert(\n              vr,\n              result,\n              dynamic_position=[],\n              static_position=ir.DenseI64ArrayAttr.get([i]),\n          )\n        return result\n      else:\n        raise NotImplementedError(arg_ty)\n    return fast_instr\n\n  @staticmethod\n  def _lift_fast_packed_instr(\n      packed_instr: str, single_instr: str,\n  ) -> Callable[[ir.Value, ir.Value], ir.Value]:\n    def fast_instr(*args):\n      arg_ty = original_arg_ty = args[0].type\n      assert all(a.type == arg_ty for a in args)\n      if not isinstance(arg_ty, ir.VectorType):\n        args = [vector.broadcast(ir.VectorType.get((1,), arg_ty), a) for a in args]\n      arg_ty = ir.VectorType(args[0].type)\n      [vec_len] = arg_ty.shape\n      vec_bitwidth = vec_len * utils.bitwidth(arg_ty.element_type)\n      if vec_len == 1 or vec_bitwidth == 32:\n        assert vec_bitwidth.bit_count() == 1\n        if vec_bitwidth == 32:\n          cstr = \"r\"","sourceCodeStart":1967,"sourceCodeEnd":2003,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/mosaic/gpu/fragmented_array.py#L1967-L2003","documentation":"Raised inside _lift_fast_instr's generated fast instruction: the register element type arg_ty is not a type the fast (PTX inline-asm) path supports for this operation, so no fast instruction can be emitted. This is an internal dispatch failure surfaced to the caller of approx pointwise ops.","triggerScenarios":"Using _lift_fast_instr-wrapped ops (sin.approx/cos.approx/tanh.approx/rsqrt.approx, exp2 approx variants) on fragments whose register element type is not one of the supported float/int widths in the fast path.","commonSituations":"Calling approx math on unusual vector element types or after bitcasts that changed register element types; typically an internal Mosaic limitation rather than user error.","solutions":["Avoid approx=True for this dtype and use the full-precision op","Ensure the array is f32 before approx ops","If you believe the type should be supported, file an issue with jax.experimental.mosaic maintainers"],"exampleFix":"// before\ny = x.sin(approx=True)  # unsupported register element type\n// after\ny = x.sin()","handlingStrategy":"fallback","validationCode":"try_approx = isinstance(fa.mlir_dtype, ir.FloatType) and fa.mlir_dtype == ir.F32Type.get()","typeGuard":null,"tryCatchPattern":"try:\n    y = op(approx=True)\nexcept NotImplementedError:\n    y = op(approx=False)","preventionTips":["Prefer full-precision ops unless profile shows approx is needed","Pin f32 for fast-math paths"],"tags":["mosaic","gpu","approx","internal"],"backgroundTag":"unsupported-dtype-operation","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}