|
| 1 | +from outlines_core import Guide |
| 2 | + |
| 3 | +try: |
| 4 | + import mlx.core as mx |
| 5 | + import numpy as np |
| 6 | +except ImportError as e: |
| 7 | + missing_dep = "numpy" if "numpy" in str(e) else "mlx" |
| 8 | + raise ImportError( |
| 9 | + f"To use the kernels in `outlines_core.kernels.mlx`, {missing_dep} must be installed. You can install it with `pip install {missing_dep}`" |
| 10 | + ) from e |
| 11 | + |
| 12 | + |
| 13 | +def allocate_token_bitmask(vocab_size: int) -> np.ndarray: |
| 14 | + return np.full( |
| 15 | + (1, (vocab_size + 31) // 32), |
| 16 | + -1, |
| 17 | + dtype=np.int32, |
| 18 | + ) |
| 19 | + |
| 20 | + |
| 21 | +_KERNEL_SOURCE = r""" |
| 22 | +// Batch index |
| 23 | +uint batch = thread_position_in_grid.y; |
| 24 | +// Element index |
| 25 | +uint elem = thread_position_in_grid.x; |
| 26 | +
|
| 27 | +uint bit = ((elem >> 5) < mask_shape[1]) && |
| 28 | + ((mask[batch * mask_shape[1] + (elem >> 5)] >> (elem & 31)) & 1); |
| 29 | +
|
| 30 | +out[batch * inp_shape[1] + elem] = bit ? inp[batch * inp_shape[1] + elem] : -INFINITY; |
| 31 | +""" |
| 32 | + |
| 33 | +_KERNEL = mx.fast.metal_kernel( |
| 34 | + name="bitmask_apply_batched", |
| 35 | + input_names=["inp", "mask"], |
| 36 | + output_names=["out"], |
| 37 | + source=_KERNEL_SOURCE, |
| 38 | +) |
| 39 | + |
| 40 | + |
| 41 | +@mx.compile |
| 42 | +def _apply_token_bitmask_kernel(data: mx.array, mask: mx.array) -> mx.array: |
| 43 | + return _KERNEL( |
| 44 | + inputs=[data, mask], |
| 45 | + template=[("T", data.dtype)], |
| 46 | + grid=(data.shape[1], data.shape[0], 1), |
| 47 | + threadgroup=(256, 1, 1), |
| 48 | + output_shapes=[data.shape], |
| 49 | + output_dtypes=[data.dtype], |
| 50 | + )[0] |
| 51 | + |
| 52 | + |
| 53 | +def apply_token_bitmask(logits: mx.array, mask_np: np.ndarray) -> mx.array: |
| 54 | + """ |
| 55 | + Apply a logits bitmask inplace, setting the probability of invalid tokens |
| 56 | + to -infinity. |
| 57 | +
|
| 58 | + Arguments: |
| 59 | + logits (mx.array): The logits tensor. |
| 60 | +
|
| 61 | + mask (mx.array): The token bitmask representing the validity of each |
| 62 | + token in the logits tensor. |
| 63 | +
|
| 64 | + Raises: |
| 65 | + ValueError: If any of the following conditions are not met: |
| 66 | + - `mask.dtype` is not `mx.int32` |
| 67 | + - `mask` is not a 2D array |
| 68 | + - `logits` is not a 2D array |
| 69 | + - `mask.shape`shape does not match `logits.shape` |
| 70 | +
|
| 71 | + Returns: |
| 72 | + None: Modifies the mask array in place. |
| 73 | + """ |
| 74 | + # makes a copy - non consuming |
| 75 | + mask = mx.array(mask_np) |
| 76 | + |
| 77 | + logits = logits if len(logits.shape) != 1 else mx.expand_dims(logits, axis=0) |
| 78 | + mask = mask if len(mask.shape) != 1 else mx.expand_dims(mask, axis=0) |
| 79 | + |
| 80 | + if mask.dtype != mx.int32: |
| 81 | + raise ValueError( |
| 82 | + f"Invalid mask dtype: Expected `np.int32`, but got `{mask.dtype}`." |
| 83 | + ) |
| 84 | + elif len(mask.shape) != 2: |
| 85 | + raise ValueError( |
| 86 | + f"Invalid mask dimensions: Expected a 2D array, but got {mask.ndim}D." |
| 87 | + ) |
| 88 | + elif len(logits.shape) != 2: |
| 89 | + raise ValueError( |
| 90 | + f"Invalid logits dimensions: Expected a 2D array, but got {logits.ndim}D." |
| 91 | + ) |
| 92 | + elif mask.shape[0] != logits.shape[0]: |
| 93 | + raise ValueError( |
| 94 | + f"Invalid batch size: Expected `mask.shape[0]` ({mask.shape[0]}) to match `logits.shape[0]` ({logits.shape[0]})." |
| 95 | + ) |
| 96 | + return _apply_token_bitmask_kernel(logits, mask) |
| 97 | + |
| 98 | + |
| 99 | +def fill_next_token_bitmask(guide: Guide, mask: np.ndarray) -> None: |
| 100 | + """ |
| 101 | + Writes a bitmask to represent the tokens permissible by the current state of the `guide`. |
| 102 | + Each bit in the bitmask corresponds to a token ID, with a bit value of 1 indicating that |
| 103 | + the token is allowed and 0 indicating that it is disallowed. This function directly modifies |
| 104 | + the `mask` array in-place. |
| 105 | +
|
| 106 | + Arguments: |
| 107 | + guide (Guide): An instance of the `Guide` class that provides the current guidance state. |
| 108 | + mask (torch.Tensor): A 2D tensor of type `torch.int32` where the bitmask will be written. |
| 109 | + The tensor must be contiguous, have a single batch dimension |
| 110 | + (shape[0] == 1), and reside on the CPU. |
| 111 | +
|
| 112 | + Raises: |
| 113 | + ValueError: If any of the following conditions are not met: |
| 114 | + - `mask.dtype` is not `np.int32` |
| 115 | + - `mask` is not a 2D tensor |
| 116 | + - `mask` does not have a single batch dimension (shape[0] != 1) |
| 117 | + - `mask` is not contiguous in memory |
| 118 | + - `mask` is not on the CPU device |
| 119 | +
|
| 120 | + Returns: |
| 121 | + None: Modifies the `mask` tensor in-place. |
| 122 | + """ |
| 123 | + if mask.dtype != np.int32: |
| 124 | + raise ValueError( |
| 125 | + f"Invalid mask dtype: Expected `np.int32`, but got `{mask.dtype}`." |
| 126 | + ) |
| 127 | + elif mask.ndim != 2: |
| 128 | + raise ValueError( |
| 129 | + f"Invalid mask dimensions: Expected a 2D array, but got {mask.ndim}D." |
| 130 | + ) |
| 131 | + elif mask.shape[0] != 1: |
| 132 | + raise ValueError( |
| 133 | + f"Invalid batch size: Batch mask writes are not supported. Expected shape[0] == 1, but got shape {mask.shape}." |
| 134 | + ) |
| 135 | + elif not mask.flags["C_CONTIGUOUS"]: |
| 136 | + raise ValueError( |
| 137 | + "Mask array must be contiguous in memory. Use `np.ascontiguousarray(mask)`." |
| 138 | + ) |
| 139 | + |
| 140 | + return guide.write_mask_into(mask.ctypes.data, mask.size, mask.itemsize) |
0 commit comments