|
14 | 14 | """xla template on python side.""" |
15 | 15 |
|
16 | 16 | from collections import namedtuple |
17 | | -from functools import partial |
18 | 17 | from typing import Any, Callable, List, Tuple, Union, cast |
19 | 18 |
|
20 | 19 | import numpy as np |
21 | | -from jax import core, dtypes |
| 20 | +from jax import ShapeDtypeStruct, dtypes, ffi |
22 | 21 | from jax import numpy as jnp |
23 | | -from jax._src.lib import xla_client |
24 | | -from jax.core import ShapedArray |
25 | | -from jax.interpreters import xla |
26 | | - |
27 | | -_core = cast(Any, core) |
28 | | -_xla = cast(Any, xla) |
29 | | -_xla_client = cast(Any, xla_client) |
30 | | - |
31 | | - |
32 | | -def _shape_with_layout( |
33 | | - specs: Tuple[Tuple[List[int], Any], ...] |
34 | | -) -> Tuple[xla_client.Shape, ...]: |
35 | | - return tuple( |
36 | | - xla_client.Shape |
37 | | - .array_shape(dtype, shape, tuple(range(len(shape) - |
38 | | - 1, -1, -1))) if len(shape) > |
39 | | - 0 else xla_client.Shape.scalar_shape(dtype) for shape, dtype in specs |
40 | | - ) |
41 | 22 |
|
42 | 23 |
|
43 | 24 | def _normalize_specs( |
44 | 25 | specs: Tuple[Tuple[Any, List[int]], ...] |
45 | | -) -> Tuple[Tuple[List[int], Any], ...]: |
| 26 | +) -> Tuple[Tuple[Tuple[int, ...], Any], ...]: |
46 | 27 | return tuple( |
47 | | - (shape, dtypes.canonicalize_dtype(dtype)) for dtype, shape in specs |
| 28 | + (tuple(shape), dtypes.canonicalize_dtype(dtype)) for dtype, shape in specs |
48 | 29 | ) |
49 | 30 |
|
50 | 31 |
|
| 32 | +def _shape_dtype_struct(shape: Tuple[int, ...], dtype: Any) -> ShapeDtypeStruct: |
| 33 | + return ShapeDtypeStruct(shape, dtype) |
| 34 | + |
| 35 | + |
| 36 | +def _layout(shape: Tuple[int, ...]) -> Tuple[int, ...]: |
| 37 | + return tuple(range(len(shape))) |
| 38 | + |
| 39 | + |
51 | 40 | def _make_xla_function( |
52 | 41 | obj: Any, |
53 | 42 | handle: bytes, |
54 | 43 | name: str, |
55 | 44 | specs: Tuple[Tuple[Any, ...], Tuple[Any, ...]], |
56 | 45 | capsules: Tuple[Any, Any], |
57 | 46 | ) -> Callable: |
58 | | - if not ( |
59 | | - hasattr(_core, "Primitive") and |
60 | | - hasattr(_xla, "backend_specific_translations") |
61 | | - ): |
62 | | - raise RuntimeError( |
63 | | - "XLA is unavailable because this JAX version removed the legacy " |
64 | | - "primitive/translation APIs used by envpool." |
65 | | - ) |
66 | 47 | in_specs, out_specs = specs |
67 | 48 | in_specs = _normalize_specs(in_specs) |
68 | 49 | out_specs = _normalize_specs(out_specs) |
69 | 50 | cpu_capsule, gpu_capsule = capsules |
70 | | - _xla_client.register_custom_call_target( |
71 | | - f"{type(obj).__name__}_{id(obj)}_{name}_cpu".encode(), |
| 51 | + call_target_name = f"{type(obj).__name__}_{id(obj)}_{name}" |
| 52 | + ffi.register_ffi_target( |
| 53 | + call_target_name, |
72 | 54 | cpu_capsule, |
73 | | - platform="cpu" |
| 55 | + platform="cpu", |
| 56 | + api_version=0, |
74 | 57 | ) |
75 | | - _xla_client.register_custom_call_target( |
76 | | - f"{type(obj).__name__}_{id(obj)}_{name}_gpu".encode(), |
| 58 | + ffi.register_ffi_target( |
| 59 | + call_target_name, |
77 | 60 | gpu_capsule, |
78 | 61 | platform="gpu", |
| 62 | + api_version=0, |
79 | 63 | ) |
80 | | - |
81 | | - def abstract( |
82 | | - *args: List[jnp.ndarray] |
83 | | - ) -> Union[ShapedArray, Tuple[ShapedArray, ...]]: |
84 | | - if len(out_specs) > 1: |
85 | | - return tuple(ShapedArray(*spec) for spec in out_specs) |
86 | | - else: |
87 | | - return ShapedArray(*out_specs[0]) |
88 | | - |
89 | | - def translation(c: Any, *args: Any, platform: str = "cpu") -> Any: |
90 | | - output_shape_with_layout = _shape_with_layout(out_specs) |
91 | | - if len(out_specs) == 1: |
92 | | - output_shape = output_shape_with_layout[0] |
93 | | - else: |
94 | | - output_shape = xla_client.Shape.tuple_shape(output_shape_with_layout) |
95 | | - return _xla_client.ops.CustomCallWithLayout( |
96 | | - c, |
97 | | - f"{type(obj).__name__}_{id(obj)}_{name}_{platform}".encode(), |
98 | | - operands=args, |
99 | | - operand_shapes_with_layout=_shape_with_layout(in_specs), |
100 | | - shape_with_layout=output_shape, |
101 | | - opaque=handle, |
102 | | - has_side_effect=True, |
103 | | - ) |
104 | | - |
105 | | - prim = _core.Primitive(f"{type(obj).__name__}_{id(obj)}_{name}") |
106 | | - prim.multiple_results = (len(out_specs) > 1) |
107 | | - prim.def_impl(partial(xla.apply_primitive, prim)) |
108 | | - prim.def_abstract_eval(abstract) |
109 | | - _xla.backend_specific_translations["cpu"][prim] = partial( |
110 | | - translation, platform="cpu" |
111 | | - ) |
112 | | - _xla.backend_specific_translations["gpu"][prim] = partial( |
113 | | - translation, platform="gpu" |
| 64 | + result_specs = tuple(_shape_dtype_struct(*spec) for spec in out_specs) |
| 65 | + xla_func = ffi.ffi_call( |
| 66 | + call_target_name, |
| 67 | + result_specs if len(result_specs) > 1 else result_specs[0], |
| 68 | + has_side_effect=True, |
| 69 | + input_layouts=tuple(_layout(shape) for shape, _ in in_specs), |
| 70 | + output_layouts=( |
| 71 | + tuple(_layout(shape) for shape, _ in out_specs) |
| 72 | + if len(out_specs) > 1 else _layout(out_specs[0][0]) |
| 73 | + ), |
| 74 | + custom_call_api_version=0, |
| 75 | + legacy_backend_config=cast(Any, handle), |
114 | 76 | ) |
115 | 77 |
|
116 | 78 | def call(*args: Any) -> Any: |
117 | | - return prim.bind(*args) |
| 79 | + return xla_func(*args) |
118 | 80 |
|
119 | 81 | return call |
120 | 82 |
|
|
0 commit comments