We are currently solving the following issue within our repository. Here is the issue text:
--- BEGIN ISSUE ---
Cupy TypeError when using asarray with torch tensors
Hi,

I'm chaining a custom feature extraction model with a PyTorchWrapper model. In the feature extraction part, I'm using asarray function on a torch tensor but as you can see below, cupy raises an error when MemoryPointer is called. Do you have any idea why that line raises an exception?

Thanks,

```
Traceback (most recent call last): 
...
  File "/storage/myuser/myproject/main.py", line 180, in forward
    features.append(model.ops.asarray2f(attrs, dtype="float32"))
  File "/storage/miniconda/envs/safirx-ml/lib/python3.9/site-packages/thinc/backends/ops.py", line 494, in asarray2f
    return cast(Floats2d, self.asarray(data, dtype=dtype))
  File "/storage/miniconda/envs/safirx-ml/lib/python3.9/site-packages/thinc/backends/cupy_ops.py", line 67, in asarray
    pointer = cupy.cuda.MemoryPointer(data.data_ptr())
  File "cupy/cuda/memory.pyx", line 343, in cupy.cuda.memory.MemoryPointer.__init__
TypeError: __init__() takes exactly 2 positional arguments (1 given)```

--- END ISSUE ---

Below are some code segments, each from a relevant file. One or more of these files may contain bugs.
--- BEGIN FILES ---
Path: `thinc/backends/cupy_ops.py`
Content:
```
1 import numpy
2 
3 try:
4     import cupy
5     import cupyx
6     import cupy.cuda
7     from cupy.cuda.compiler import compile_with_cache  # noqa: F401
8 
9     has_cupy = True
10 
11     # We no longer have to set up the memory pool, fortunately.
12 except ImportError:
13     cupy = None
14     cupyx = None
15     has_cupy = False
16 
17 from .. import registry
18 from .ops import Ops
19 from .numpy_ops import NumpyOps
20 from . import _custom_kernels
21 from ..types import DeviceTypes
22 
23 
24 @registry.ops("CupyOps")
25 class CupyOps(Ops):
26     name = "cupy"
27     xp = cupy
28     _xp2 = cupyx
29 
30     def __init__(
31         self, device_type: DeviceTypes = "gpu", device_id: int = 0, **kwargs
32     ) -> None:
33         self.device_type = device_type
34         self.device_id = device_id
35 
36     def to_numpy(self, data, *, byte_order=None):
37         if not isinstance(data, numpy.ndarray):
38             data = data.get()
39         if byte_order:
40             dtype = data.dtype.newbyteorder(byte_order)
41             data = numpy.asarray(data, dtype=dtype)
42         return data
43 
44     def gelu(self, X, inplace=False):
45         if X.dtype == "float32":
46             return _custom_kernels.gelu(X, inplace=inplace, threshold=6.0)
47         else:
48             return super().gelu(X, inplace=inplace)
49 
50     def backprop_gelu(self, dY, X, inplace=False):
51         if X.dtype == "float32" and dY.dtype == "float32":
52             return _custom_kernels.backprop_gelu(dY, X, inplace=inplace, threshold=6.0)
53         else:
54             return super().backprop_gelu(dY, X, inplace=inplace)
55 
56     def gemm(self, x, y, out=None, trans1=False, trans2=False):
57         if isinstance(x, numpy.ndarray) or isinstance(y, numpy.ndarray):
58             raise ValueError(
59                 "Encountered a numpy array when processing with cupy. "
60                 "Did you call model.ops.asarray on your data?"
61             )
62         if trans1:
63             x = x.T
64         if trans2:
65             y = y.T
66         if out is None:
67             return self.xp.dot(x, y)
68         else:
69             self.xp.dot(x, y, out=out)
70             return out
71 
72     def asarray(self, data, dtype=None):
73         # This is sort of frustrating, but we can't easily otherwise pass
74         # forward "unset".
75         dtype = {"dtype": dtype} if dtype is not None else {}
76         if isinstance(data, cupy.ndarray):
77             return self.xp.asarray(data, **dtype)
78         elif hasattr(data, "data_ptr"):
79             # Handles PyTorch Tensors
80             pointer = cupy.cuda.MemoryPointer(data.data_ptr())
81             shape = data.stride()
82             array = self.xp.ndarray(shape, memptr=pointer, **dtype)
83             return array
84         else:
85             result = self.xp.array(data, **dtype)
86             return result
87 
88     def maxout(self, X):
89         return _custom_kernels.maxout(X)
90 
91     def backprop_maxout(self, dY, which, P):
92         return _custom_kernels.backprop_maxout(dY, which, P)
93 
94     def relu(self, X, inplace=False):
95         if not inplace:
96             return X * (X > 0)
97         else:
98             X *= X > 0
99             return X
100 
101     def backprop_relu(self, dY, Y, inplace=False):
102         if not inplace:
103             return dY * (Y > 0)
104         dY *= Y > 0
105         return dY
106 
107     def clipped_linear(
108         self,
109         X,
110         slope: float = 1.0,
111         offset: float = 0.0,
112         min_val: float = 0.0,
113         max_val: float = 1.0,
114         inplace: bool = False,
115     ):
116         if X.dtype == "float32":
117             return _custom_kernels.clipped_linear(
118                 X,
119                 inplace=inplace,
120                 slope=slope,
121                 offset=offset,
122                 min_val=min_val,
123                 max_val=max_val,
124             )
125         else:
126             return super().clipped_linear(
127                 X,
128                 inplace=inplace,
129                 slope=slope,
130                 offset=offset,
131                 min_val=min_val,
132                 max_val=max_val,
133             )
134 
135     def backprop_clipped_linear(
136         self,
137         dY,
138         X,
139         slope: float = 1.0,
140         offset: float = 0.0,
141         min_val: float = 0.0,
142         max_val: float = 1.0,
143         inplace: bool = False,
144     ):
145         if X.dtype == "float32" and dY.dtype == "float32":
146             return _custom_kernels.backprop_clipped_linear(
147                 dY=dY,
148                 X=X,
149                 slope=slope,
150                 offset=offset,
151                 min_val=min_val,
152                 max_val=max_val,
153                 inplace=inplace,
154             )
155         else:
156             return super().backprop_clipped_linear(
157                 dY=dY,
158                 X=X,
159                 slope=slope,
160                 offset=offset,
161                 min_val=min_val,
162                 max_val=max_val,
163                 inplace=inplace,
164             )
165 
166     def backprop_hard_swish(self, dY, X, inplace: bool = False):
167         if X.dtype == "float32" and dY.dtype == "float32":
168             return _custom_kernels.backprop_hard_swish(dY, X, inplace=inplace)
169         else:
170             return super().backprop_hard_swish(dY, X, inplace=inplace)
171 
172     def backprop_hard_swish_mobilenet(self, dY, X, inplace: bool = False):
173         if X.dtype == "float32" and dY.dtype == "float32":
174             return _custom_kernels.backprop_hard_swish_mobilenet(dY, X, inplace=inplace)
175         else:
176             return super().backprop_hard_swish_mobilenet(dY, X, inplace=inplace)
177 
178     def mish(self, X, threshold=20.0, inplace=False):
179         if X.dtype == "float32" and not inplace:
180             return _custom_kernels.mish(X, threshold=threshold, out=None)
181         else:
182             return super().mish(X, threshold, inplace)
183 
184     def backprop_mish(self, dY, X, threshold=20.0, inplace=False):
185         if dY.dtype == "float32" and X.dtype == "float32" and not inplace:
186             return _custom_kernels.backprop_mish(dY, X, threshold=threshold)
187         else:
188             return super().backprop_mish(dY, X, threshold, inplace)
189 
190     def swish(self, X, inplace=False):
191         if X.dtype == "float32":
192             return _custom_kernels.swish(X, inplace=inplace, threshold=17.0)
193         else:
194             return super().swish(X, inplace=inplace)
195 
196     def backprop_swish(self, dY, X, Y, inplace=False):
197         if X.dtype == "float32" and dY.dtype == "float32":
198             return _custom_kernels.backprop_swish(
199                 dY, X, Y, inplace=inplace, threshold=17.0
200             )
201         else:
202             return super().backprop_swish(dY, X, Y, inplace=inplace)
203 
204     def clip_gradient(self, gradient, threshold):
205         # We do not use CuPy's linalg.norm, since it uses scalar reductions
206         # using one CUDA block. This is a lot slower than the cuBLAS
207         # implementation.
208         def frobenius_norm(X):
209             X_vec = X.reshape(-1)
210             return cupy.cublas.nrm2(X_vec)
211 
212         grad_norm = cupy.maximum(frobenius_norm(gradient), 1e-12)
213         gradient *= cupy.minimum(threshold, grad_norm) / grad_norm
214         return gradient
215 
216     def seq2col(self, seq, nW, *, lengths=None):
217         """Given an (M, N) sequence of vectors, return an (M, N*(nW*2+1)) sequence.
218         The new sequence is constructed by concatenating nW preceding and succeeding
219         vectors onto each column in the sequence, to extract a window of features.
220         """
221         return _custom_kernels.seq2col(seq, nW, lengths=lengths)
222 
223     def backprop_seq2col(self, dY, nW, *, lengths=None):
224         return _custom_kernels.backprop_seq2col(dY, nW, lengths=lengths)
225 
226     def reduce_mean(self, X, lengths):
227         return _custom_kernels.reduce_mean(X, lengths)
228 
229     def backprop_reduce_mean(self, d_means, lengths):
230         return _custom_kernels.backprop_reduce_mean(d_means, lengths)
231 
232     def reduce_max(self, X, lengths):
233         return _custom_kernels.reduce_max(X, lengths)
234 
235     def backprop_reduce_max(self, d_maxes, which, lengths):
236         return _custom_kernels.backprop_reduce_max(d_maxes, which, lengths)
237 
238     def reduce_sum(self, X, lengths):
239         return _custom_kernels.reduce_sum(X, lengths)
240 
241     def backprop_reduce_sum(self, d_sums, lengths):
242         return _custom_kernels.backprop_reduce_sum(d_sums, lengths)
243 
244     def hash(self, ids, seed):
245         return _custom_kernels.hash(ids, seed)
246 
247     def scatter_add(self, table, indices, values):
248         self._xp2.scatter_add(table, indices, values)
249 
250     def adam(
251         self, weights, gradient, mom1, mom2, beta1, beta2, eps, learn_rate, mod_rate=1.0
252     ):
253         adam_kernel(
254             gradient, learn_rate, 1 - beta1, 1 - beta2, eps, weights, mom1, mom2
255         )
256         gradient.fill(0)
257         return weights, gradient, mom1, mom2
258 
259     def position_encode(self, N, D, period=10000, out=None):
260         positions = NumpyOps().position_encode(N, D, period=period, out=out)
261         return self.asarray(positions)
262 
263 
264 if cupy is not None:
265     adam_kernel = cupy.ElementwiseKernel(
266         "T grad, T lr, T one_minus_beta1, T one_minus_beta2, T eps",
267         "T param, T m, T v",
268         """m += one_minus_beta1 * (grad - m);
269         v += one_minus_beta2 * (grad * grad - v);
270         param -= lr * m / (sqrt(v) + eps);""",
271         "adam",
272     )
273 else:
274     adam_kernel = None
275
```
--- END FILES ---

Please first localize the bug based on the issue statement, and then generate a patch according to the `git diff` format fenced by three backticks.

Here is an example:

```diff
diff --git a/examples/server_async.py b/examples/server_async.py
--- a/examples/server_async.py
+++ b/examples/server_async.py
@@ -313,4 +313,4 @@
 
 
 if __name__ == "__main__":
-    asyncio.run(run_async_server("."), debug=True)
+    asyncio.run(run_async_server(), debug=True)
diff --git a/examples/server_sync.py b/examples/server_sync.py
--- a/examples/server_sync.py
+++ b/examples/server_sync.py
@@ -313,5 +313,5 @@
 
 
 if __name__ == "__main__":
-    server = run_sync_server(".")
+    server = run_sync_server()
     server.shutdown()

```

diff --git a/thinc/backends/cupy_ops.py b/thinc/backends/cupy_ops.py
--- a/thinc/backends/cupy_ops.py
+++ b/thinc/backends/cupy_ops.py
@@ -19,6 +19,8 @@
 from .numpy_ops import NumpyOps
 from . import _custom_kernels
 from ..types import DeviceTypes
+from ..util import torch2xp, tensorflow2xp, mxnet2xp
+from ..util import is_torch_array, is_tensorflow_array, is_mxnet_array
 
 
 @registry.ops("CupyOps")
@@ -73,17 +75,28 @@
         # This is sort of frustrating, but we can't easily otherwise pass
         # forward "unset".
         dtype = {"dtype": dtype} if dtype is not None else {}
+
+        # We'll try to perform a zero-copy conversion if possible.
+        array = None
+        cast_array = False
         if isinstance(data, cupy.ndarray):
-            return self.xp.asarray(data, **dtype)
-        elif hasattr(data, "data_ptr"):
-            # Handles PyTorch Tensors
-            pointer = cupy.cuda.MemoryPointer(data.data_ptr())
-            shape = data.stride()
-            array = self.xp.ndarray(shape, memptr=pointer, **dtype)
-            return array
+            array = self.xp.asarray(data, **dtype)
+        elif is_torch_array(data) and data.device.type == "cuda":
+            array = torch2xp(data)
+            cast_array = True
+        elif is_tensorflow_array(data) and "GPU:" in data.device:
+            array = tensorflow2xp(data)
+            cast_array = True
+        elif is_mxnet_array(data) and data.context.device_type != "cpu":
+            array = mxnet2xp(data)
+            cast_array = True
         else:
-            result = self.xp.array(data, **dtype)
-            return result
+            array = self.xp.array(data, **dtype)
+
+        if cast_array and dtype != {}:
+            array = array.astype(dtype["dtype"])
+
+        return array
 
     def maxout(self, X):
         return _custom_kernels.maxout(X)

{"golden_diff": "diff --git a/thinc/backends/cupy_ops.py b/thinc/backends/cupy_ops.py\n--- a/thinc/backends/cupy_ops.py\n+++ b/thinc/backends/cupy_ops.py\n@@ -19,6 +19,8 @@\n from .numpy_ops import NumpyOps\n from . import _custom_kernels\n from ..types import DeviceTypes\n+from ..util import torch2xp, tensorflow2xp, mxnet2xp\n+from ..util import is_torch_array, is_tensorflow_array, is_mxnet_array\n \n \n @registry.ops(\"CupyOps\")\n@@ -73,17 +75,28 @@\n         # This is sort of frustrating, but we can't easily otherwise pass\n         # forward \"unset\".\n         dtype = {\"dtype\": dtype} if dtype is not None else {}\n+\n+        # We'll try to perform a zero-copy conversion if possible.\n+        array = None\n+        cast_array = False\n         if isinstance(data, cupy.ndarray):\n-            return self.xp.asarray(data, **dtype)\n-        elif hasattr(data, \"data_ptr\"):\n-            # Handles PyTorch Tensors\n-            pointer = cupy.cuda.MemoryPointer(data.data_ptr())\n-            shape = data.stride()\n-            array = self.xp.ndarray(shape, memptr=pointer, **dtype)\n-            return array\n+            array = self.xp.asarray(data, **dtype)\n+        elif is_torch_array(data) and data.device.type == \"cuda\":\n+            array = torch2xp(data)\n+            cast_array = True\n+        elif is_tensorflow_array(data) and \"GPU:\" in data.device:\n+            array = tensorflow2xp(data)\n+            cast_array = True\n+        elif is_mxnet_array(data) and data.context.device_type != \"cpu\":\n+            array = mxnet2xp(data)\n+            cast_array = True\n         else:\n-            result = self.xp.array(data, **dtype)\n-            return result\n+            array = self.xp.array(data, **dtype)\n+\n+        if cast_array and dtype != {}:\n+            array = array.astype(dtype[\"dtype\"])\n+\n+        return array\n \n     def maxout(self, X):\n         return _custom_kernels.maxout(X)\n", "issue": "Cupy TypeError when using asarray with torch tensors\nHi,\r\n\r\nI'm chaining a custom feature extraction model with a PyTorchWrapper model. In the feature extraction part, I'm using asarray function on a torch tensor but as you can see below, cupy raises an error when MemoryPointer is called. Do you have any idea why that line raises an exception?\r\n\r\nThanks,\r\n\r\n```\r\nTraceback (most recent call last): \r\n...\r\n  File \"/storage/myuser/myproject/main.py\", line 180, in forward\r\n    features.append(model.ops.asarray2f(attrs, dtype=\"float32\"))\r\n  File \"/storage/miniconda/envs/safirx-ml/lib/python3.9/site-packages/thinc/backends/ops.py\", line 494, in asarray2f\r\n    return cast(Floats2d, self.asarray(data, dtype=dtype))\r\n  File \"/storage/miniconda/envs/safirx-ml/lib/python3.9/site-packages/thinc/backends/cupy_ops.py\", line 67, in asarray\r\n    pointer = cupy.cuda.MemoryPointer(data.data_ptr())\r\n  File \"cupy/cuda/memory.pyx\", line 343, in cupy.cuda.memory.MemoryPointer.__init__\r\nTypeError: __init__() takes exactly 2 positional arguments (1 given)```\n", "before_files": [{"content": "import numpy\n\ntry:\n    import cupy\n    import cupyx\n    import cupy.cuda\n    from cupy.cuda.compiler import compile_with_cache  # noqa: F401\n\n    has_cupy = True\n\n    # We no longer have to set up the memory pool, fortunately.\nexcept ImportError:\n    cupy = None\n    cupyx = None\n    has_cupy = False\n\nfrom .. import registry\nfrom .ops import Ops\nfrom .numpy_ops import NumpyOps\nfrom . import _custom_kernels\nfrom ..types import DeviceTypes\n\n\n@registry.ops(\"CupyOps\")\nclass CupyOps(Ops):\n    name = \"cupy\"\n    xp = cupy\n    _xp2 = cupyx\n\n    def __init__(\n        self, device_type: DeviceTypes = \"gpu\", device_id: int = 0, **kwargs\n    ) -> None:\n        self.device_type = device_type\n        self.device_id = device_id\n\n    def to_numpy(self, data, *, byte_order=None):\n        if not isinstance(data, numpy.ndarray):\n            data = data.get()\n        if byte_order:\n            dtype = data.dtype.newbyteorder(byte_order)\n            data = numpy.asarray(data, dtype=dtype)\n        return data\n\n    def gelu(self, X, inplace=False):\n        if X.dtype == \"float32\":\n            return _custom_kernels.gelu(X, inplace=inplace, threshold=6.0)\n        else:\n            return super().gelu(X, inplace=inplace)\n\n    def backprop_gelu(self, dY, X, inplace=False):\n        if X.dtype == \"float32\" and dY.dtype == \"float32\":\n            return _custom_kernels.backprop_gelu(dY, X, inplace=inplace, threshold=6.0)\n        else:\n            return super().backprop_gelu(dY, X, inplace=inplace)\n\n    def gemm(self, x, y, out=None, trans1=False, trans2=False):\n        if isinstance(x, numpy.ndarray) or isinstance(y, numpy.ndarray):\n            raise ValueError(\n                \"Encountered a numpy array when processing with cupy. \"\n                \"Did you call model.ops.asarray on your data?\"\n            )\n        if trans1:\n            x = x.T\n        if trans2:\n            y = y.T\n        if out is None:\n            return self.xp.dot(x, y)\n        else:\n            self.xp.dot(x, y, out=out)\n            return out\n\n    def asarray(self, data, dtype=None):\n        # This is sort of frustrating, but we can't easily otherwise pass\n        # forward \"unset\".\n        dtype = {\"dtype\": dtype} if dtype is not None else {}\n        if isinstance(data, cupy.ndarray):\n            return self.xp.asarray(data, **dtype)\n        elif hasattr(data, \"data_ptr\"):\n            # Handles PyTorch Tensors\n            pointer = cupy.cuda.MemoryPointer(data.data_ptr())\n            shape = data.stride()\n            array = self.xp.ndarray(shape, memptr=pointer, **dtype)\n            return array\n        else:\n            result = self.xp.array(data, **dtype)\n            return result\n\n    def maxout(self, X):\n        return _custom_kernels.maxout(X)\n\n    def backprop_maxout(self, dY, which, P):\n        return _custom_kernels.backprop_maxout(dY, which, P)\n\n    def relu(self, X, inplace=False):\n        if not inplace:\n            return X * (X > 0)\n        else:\n            X *= X > 0\n            return X\n\n    def backprop_relu(self, dY, Y, inplace=False):\n        if not inplace:\n            return dY * (Y > 0)\n        dY *= Y > 0\n        return dY\n\n    def clipped_linear(\n        self,\n        X,\n        slope: float = 1.0,\n        offset: float = 0.0,\n        min_val: float = 0.0,\n        max_val: float = 1.0,\n        inplace: bool = False,\n    ):\n        if X.dtype == \"float32\":\n            return _custom_kernels.clipped_linear(\n                X,\n                inplace=inplace,\n                slope=slope,\n                offset=offset,\n                min_val=min_val,\n                max_val=max_val,\n            )\n        else:\n            return super().clipped_linear(\n                X,\n                inplace=inplace,\n                slope=slope,\n                offset=offset,\n                min_val=min_val,\n                max_val=max_val,\n            )\n\n    def backprop_clipped_linear(\n        self,\n        dY,\n        X,\n        slope: float = 1.0,\n        offset: float = 0.0,\n        min_val: float = 0.0,\n        max_val: float = 1.0,\n        inplace: bool = False,\n    ):\n        if X.dtype == \"float32\" and dY.dtype == \"float32\":\n            return _custom_kernels.backprop_clipped_linear(\n                dY=dY,\n                X=X,\n                slope=slope,\n                offset=offset,\n                min_val=min_val,\n                max_val=max_val,\n                inplace=inplace,\n            )\n        else:\n            return super().backprop_clipped_linear(\n                dY=dY,\n                X=X,\n                slope=slope,\n                offset=offset,\n                min_val=min_val,\n                max_val=max_val,\n                inplace=inplace,\n            )\n\n    def backprop_hard_swish(self, dY, X, inplace: bool = False):\n        if X.dtype == \"float32\" and dY.dtype == \"float32\":\n            return _custom_kernels.backprop_hard_swish(dY, X, inplace=inplace)\n        else:\n            return super().backprop_hard_swish(dY, X, inplace=inplace)\n\n    def backprop_hard_swish_mobilenet(self, dY, X, inplace: bool = False):\n        if X.dtype == \"float32\" and dY.dtype == \"float32\":\n            return _custom_kernels.backprop_hard_swish_mobilenet(dY, X, inplace=inplace)\n        else:\n            return super().backprop_hard_swish_mobilenet(dY, X, inplace=inplace)\n\n    def mish(self, X, threshold=20.0, inplace=False):\n        if X.dtype == \"float32\" and not inplace:\n            return _custom_kernels.mish(X, threshold=threshold, out=None)\n        else:\n            return super().mish(X, threshold, inplace)\n\n    def backprop_mish(self, dY, X, threshold=20.0, inplace=False):\n        if dY.dtype == \"float32\" and X.dtype == \"float32\" and not inplace:\n            return _custom_kernels.backprop_mish(dY, X, threshold=threshold)\n        else:\n            return super().backprop_mish(dY, X, threshold, inplace)\n\n    def swish(self, X, inplace=False):\n        if X.dtype == \"float32\":\n            return _custom_kernels.swish(X, inplace=inplace, threshold=17.0)\n        else:\n            return super().swish(X, inplace=inplace)\n\n    def backprop_swish(self, dY, X, Y, inplace=False):\n        if X.dtype == \"float32\" and dY.dtype == \"float32\":\n            return _custom_kernels.backprop_swish(\n                dY, X, Y, inplace=inplace, threshold=17.0\n            )\n        else:\n            return super().backprop_swish(dY, X, Y, inplace=inplace)\n\n    def clip_gradient(self, gradient, threshold):\n        # We do not use CuPy's linalg.norm, since it uses scalar reductions\n        # using one CUDA block. This is a lot slower than the cuBLAS\n        # implementation.\n        def frobenius_norm(X):\n            X_vec = X.reshape(-1)\n            return cupy.cublas.nrm2(X_vec)\n\n        grad_norm = cupy.maximum(frobenius_norm(gradient), 1e-12)\n        gradient *= cupy.minimum(threshold, grad_norm) / grad_norm\n        return gradient\n\n    def seq2col(self, seq, nW, *, lengths=None):\n        \"\"\"Given an (M, N) sequence of vectors, return an (M, N*(nW*2+1)) sequence.\n        The new sequence is constructed by concatenating nW preceding and succeeding\n        vectors onto each column in the sequence, to extract a window of features.\n        \"\"\"\n        return _custom_kernels.seq2col(seq, nW, lengths=lengths)\n\n    def backprop_seq2col(self, dY, nW, *, lengths=None):\n        return _custom_kernels.backprop_seq2col(dY, nW, lengths=lengths)\n\n    def reduce_mean(self, X, lengths):\n        return _custom_kernels.reduce_mean(X, lengths)\n\n    def backprop_reduce_mean(self, d_means, lengths):\n        return _custom_kernels.backprop_reduce_mean(d_means, lengths)\n\n    def reduce_max(self, X, lengths):\n        return _custom_kernels.reduce_max(X, lengths)\n\n    def backprop_reduce_max(self, d_maxes, which, lengths):\n        return _custom_kernels.backprop_reduce_max(d_maxes, which, lengths)\n\n    def reduce_sum(self, X, lengths):\n        return _custom_kernels.reduce_sum(X, lengths)\n\n    def backprop_reduce_sum(self, d_sums, lengths):\n        return _custom_kernels.backprop_reduce_sum(d_sums, lengths)\n\n    def hash(self, ids, seed):\n        return _custom_kernels.hash(ids, seed)\n\n    def scatter_add(self, table, indices, values):\n        self._xp2.scatter_add(table, indices, values)\n\n    def adam(\n        self, weights, gradient, mom1, mom2, beta1, beta2, eps, learn_rate, mod_rate=1.0\n    ):\n        adam_kernel(\n            gradient, learn_rate, 1 - beta1, 1 - beta2, eps, weights, mom1, mom2\n        )\n        gradient.fill(0)\n        return weights, gradient, mom1, mom2\n\n    def position_encode(self, N, D, period=10000, out=None):\n        positions = NumpyOps().position_encode(N, D, period=period, out=out)\n        return self.asarray(positions)\n\n\nif cupy is not None:\n    adam_kernel = cupy.ElementwiseKernel(\n        \"T grad, T lr, T one_minus_beta1, T one_minus_beta2, T eps\",\n        \"T param, T m, T v\",\n        \"\"\"m += one_minus_beta1 * (grad - m);\n        v += one_minus_beta2 * (grad * grad - v);\n        param -= lr * m / (sqrt(v) + eps);\"\"\",\n        \"adam\",\n    )\nelse:\n    adam_kernel = None\n", "path": "thinc/backends/cupy_ops.py"}], "after_files": [{"content": "import numpy\n\ntry:\n    import cupy\n    import cupyx\n    import cupy.cuda\n    from cupy.cuda.compiler import compile_with_cache  # noqa: F401\n\n    has_cupy = True\n\n    # We no longer have to set up the memory pool, fortunately.\nexcept ImportError:\n    cupy = None\n    cupyx = None\n    has_cupy = False\n\nfrom .. import registry\nfrom .ops import Ops\nfrom .numpy_ops import NumpyOps\nfrom . import _custom_kernels\nfrom ..types import DeviceTypes\nfrom ..util import torch2xp, tensorflow2xp, mxnet2xp\nfrom ..util import is_torch_array, is_tensorflow_array, is_mxnet_array\n\n\n@registry.ops(\"CupyOps\")\nclass CupyOps(Ops):\n    name = \"cupy\"\n    xp = cupy\n    _xp2 = cupyx\n\n    def __init__(\n        self, device_type: DeviceTypes = \"gpu\", device_id: int = 0, **kwargs\n    ) -> None:\n        self.device_type = device_type\n        self.device_id = device_id\n\n    def to_numpy(self, data, *, byte_order=None):\n        if not isinstance(data, numpy.ndarray):\n            data = data.get()\n        if byte_order:\n            dtype = data.dtype.newbyteorder(byte_order)\n            data = numpy.asarray(data, dtype=dtype)\n        return data\n\n    def gelu(self, X, inplace=False):\n        if X.dtype == \"float32\":\n            return _custom_kernels.gelu(X, inplace=inplace, threshold=6.0)\n        else:\n            return super().gelu(X, inplace=inplace)\n\n    def backprop_gelu(self, dY, X, inplace=False):\n        if X.dtype == \"float32\" and dY.dtype == \"float32\":\n            return _custom_kernels.backprop_gelu(dY, X, inplace=inplace, threshold=6.0)\n        else:\n            return super().backprop_gelu(dY, X, inplace=inplace)\n\n    def gemm(self, x, y, out=None, trans1=False, trans2=False):\n        if isinstance(x, numpy.ndarray) or isinstance(y, numpy.ndarray):\n            raise ValueError(\n                \"Encountered a numpy array when processing with cupy. \"\n                \"Did you call model.ops.asarray on your data?\"\n            )\n        if trans1:\n            x = x.T\n        if trans2:\n            y = y.T\n        if out is None:\n            return self.xp.dot(x, y)\n        else:\n            self.xp.dot(x, y, out=out)\n            return out\n\n    def asarray(self, data, dtype=None):\n        # This is sort of frustrating, but we can't easily otherwise pass\n        # forward \"unset\".\n        dtype = {\"dtype\": dtype} if dtype is not None else {}\n\n        # We'll try to perform a zero-copy conversion if possible.\n        array = None\n        cast_array = False\n        if isinstance(data, cupy.ndarray):\n            array = self.xp.asarray(data, **dtype)\n        elif is_torch_array(data) and data.device.type == \"cuda\":\n            array = torch2xp(data)\n            cast_array = True\n        elif is_tensorflow_array(data) and \"GPU:\" in data.device:\n            array = tensorflow2xp(data)\n            cast_array = True\n        elif is_mxnet_array(data) and data.context.device_type != \"cpu\":\n            array = mxnet2xp(data)\n            cast_array = True\n        else:\n            array = self.xp.array(data, **dtype)\n\n        if cast_array and dtype != {}:\n            array = array.astype(dtype[\"dtype\"])\n\n        return array\n\n    def maxout(self, X):\n        return _custom_kernels.maxout(X)\n\n    def backprop_maxout(self, dY, which, P):\n        return _custom_kernels.backprop_maxout(dY, which, P)\n\n    def relu(self, X, inplace=False):\n        if not inplace:\n            return X * (X > 0)\n        else:\n            X *= X > 0\n            return X\n\n    def backprop_relu(self, dY, Y, inplace=False):\n        if not inplace:\n            return dY * (Y > 0)\n        dY *= Y > 0\n        return dY\n\n    def clipped_linear(\n        self,\n        X,\n        slope: float = 1.0,\n        offset: float = 0.0,\n        min_val: float = 0.0,\n        max_val: float = 1.0,\n        inplace: bool = False,\n    ):\n        if X.dtype == \"float32\":\n            return _custom_kernels.clipped_linear(\n                X,\n                inplace=inplace,\n                slope=slope,\n                offset=offset,\n                min_val=min_val,\n                max_val=max_val,\n            )\n        else:\n            return super().clipped_linear(\n                X,\n                inplace=inplace,\n                slope=slope,\n                offset=offset,\n                min_val=min_val,\n                max_val=max_val,\n            )\n\n    def backprop_clipped_linear(\n        self,\n        dY,\n        X,\n        slope: float = 1.0,\n        offset: float = 0.0,\n        min_val: float = 0.0,\n        max_val: float = 1.0,\n        inplace: bool = False,\n    ):\n        if X.dtype == \"float32\" and dY.dtype == \"float32\":\n            return _custom_kernels.backprop_clipped_linear(\n                dY=dY,\n                X=X,\n                slope=slope,\n                offset=offset,\n                min_val=min_val,\n                max_val=max_val,\n                inplace=inplace,\n            )\n        else:\n            return super().backprop_clipped_linear(\n                dY=dY,\n                X=X,\n                slope=slope,\n                offset=offset,\n                min_val=min_val,\n                max_val=max_val,\n                inplace=inplace,\n            )\n\n    def backprop_hard_swish(self, dY, X, inplace: bool = False):\n        if X.dtype == \"float32\" and dY.dtype == \"float32\":\n            return _custom_kernels.backprop_hard_swish(dY, X, inplace=inplace)\n        else:\n            return super().backprop_hard_swish(dY, X, inplace=inplace)\n\n    def backprop_hard_swish_mobilenet(self, dY, X, inplace: bool = False):\n        if X.dtype == \"float32\" and dY.dtype == \"float32\":\n            return _custom_kernels.backprop_hard_swish_mobilenet(dY, X, inplace=inplace)\n        else:\n            return super().backprop_hard_swish_mobilenet(dY, X, inplace=inplace)\n\n    def mish(self, X, threshold=20.0, inplace=False):\n        if X.dtype == \"float32\" and not inplace:\n            return _custom_kernels.mish(X, threshold=threshold, out=None)\n        else:\n            return super().mish(X, threshold, inplace)\n\n    def backprop_mish(self, dY, X, threshold=20.0, inplace=False):\n        if dY.dtype == \"float32\" and X.dtype == \"float32\" and not inplace:\n            return _custom_kernels.backprop_mish(dY, X, threshold=threshold)\n        else:\n            return super().backprop_mish(dY, X, threshold, inplace)\n\n    def swish(self, X, inplace=False):\n        if X.dtype == \"float32\":\n            return _custom_kernels.swish(X, inplace=inplace, threshold=17.0)\n        else:\n            return super().swish(X, inplace=inplace)\n\n    def backprop_swish(self, dY, X, Y, inplace=False):\n        if X.dtype == \"float32\" and dY.dtype == \"float32\":\n            return _custom_kernels.backprop_swish(\n                dY, X, Y, inplace=inplace, threshold=17.0\n            )\n        else:\n            return super().backprop_swish(dY, X, Y, inplace=inplace)\n\n    def clip_gradient(self, gradient, threshold):\n        # We do not use CuPy's linalg.norm, since it uses scalar reductions\n        # using one CUDA block. This is a lot slower than the cuBLAS\n        # implementation.\n        def frobenius_norm(X):\n            X_vec = X.reshape(-1)\n            return cupy.cublas.nrm2(X_vec)\n\n        grad_norm = cupy.maximum(frobenius_norm(gradient), 1e-12)\n        gradient *= cupy.minimum(threshold, grad_norm) / grad_norm\n        return gradient\n\n    def seq2col(self, seq, nW, *, lengths=None):\n        \"\"\"Given an (M, N) sequence of vectors, return an (M, N*(nW*2+1)) sequence.\n        The new sequence is constructed by concatenating nW preceding and succeeding\n        vectors onto each column in the sequence, to extract a window of features.\n        \"\"\"\n        return _custom_kernels.seq2col(seq, nW, lengths=lengths)\n\n    def backprop_seq2col(self, dY, nW, *, lengths=None):\n        return _custom_kernels.backprop_seq2col(dY, nW, lengths=lengths)\n\n    def reduce_mean(self, X, lengths):\n        return _custom_kernels.reduce_mean(X, lengths)\n\n    def backprop_reduce_mean(self, d_means, lengths):\n        return _custom_kernels.backprop_reduce_mean(d_means, lengths)\n\n    def reduce_max(self, X, lengths):\n        return _custom_kernels.reduce_max(X, lengths)\n\n    def backprop_reduce_max(self, d_maxes, which, lengths):\n        return _custom_kernels.backprop_reduce_max(d_maxes, which, lengths)\n\n    def reduce_sum(self, X, lengths):\n        return _custom_kernels.reduce_sum(X, lengths)\n\n    def backprop_reduce_sum(self, d_sums, lengths):\n        return _custom_kernels.backprop_reduce_sum(d_sums, lengths)\n\n    def hash(self, ids, seed):\n        return _custom_kernels.hash(ids, seed)\n\n    def scatter_add(self, table, indices, values):\n        self._xp2.scatter_add(table, indices, values)\n\n    def adam(\n        self, weights, gradient, mom1, mom2, beta1, beta2, eps, learn_rate, mod_rate=1.0\n    ):\n        adam_kernel(\n            gradient, learn_rate, 1 - beta1, 1 - beta2, eps, weights, mom1, mom2\n        )\n        gradient.fill(0)\n        return weights, gradient, mom1, mom2\n\n    def position_encode(self, N, D, period=10000, out=None):\n        positions = NumpyOps().position_encode(N, D, period=period, out=out)\n        return self.asarray(positions)\n\n\nif cupy is not None:\n    adam_kernel = cupy.ElementwiseKernel(\n        \"T grad, T lr, T one_minus_beta1, T one_minus_beta2, T eps\",\n        \"T param, T m, T v\",\n        \"\"\"m += one_minus_beta1 * (grad - m);\n        v += one_minus_beta2 * (grad * grad - v);\n        param -= lr * m / (sqrt(v) + eps);\"\"\",\n        \"adam\",\n    )\nelse:\n    adam_kernel = None\n", "path": "thinc/backends/cupy_ops.py"}]}