We are currently solving the following issue within our repository. Here is the issue text:
--- BEGIN ISSUE ---
cmath linking issues in CUDA
From the mailing list:
https://groups.google.com/a/continuum.io/forum/#!topic/numba-users/oBcaFcsD-x4
This:
```python
from numba import cuda
import cmath
import numpy as np

@cuda.jit(device = True)
def quarticRoots(a, b, c, d, e):

	del_0 = c**2 - 3.*b*d + 12.*a*e
	del_1 = 2.*c**3 - 9.*b*c*d + 27*e*b**2 + 27*a*d**2 - 72*a*c*e

	p = (8.*a*c - 3.*b**2)/(8.*a**2)
	q = (b**3 - 4.*a*b*c + 8.*d*a**2)/(8.*a**3)

	Q = ((del_1 + cmath.sqrt(del_1**2 - 4.*del_0**3))/2.)**(1./3.)#**(1./3.) is the problem.
	S = 0.5*cmath.sqrt(-2./3.*p + 1/(3.*a)*(Q + del_0/Q))

	x1 = -b/(4.*a) - S + 0.5*cmath.sqrt(-4.*S**2 - 2.*p + q/S)
	x2 = -b/(4.*a) - S - 0.5*cmath.sqrt(-4.*S**2 - 2.*p + q/S)
	x3 = -b/(4.*a) + S + 0.5*cmath.sqrt(-4.*S**2 - 2.*p - q/S)
	x4 = -b/(4.*a) + S - 0.5*cmath.sqrt(-4.*S**2 - 2.*p - q/S)

	return (x1, x2, x3, x4)

@cuda.jit
def launcher(d_array):
	output = quarticRoots(2.,4.,2.,1.,1.)

	d_array[0] = output[0]
	d_array[1] = output[1]
	d_array[2] = output[2]
	d_array[3] = output[3]

def main():
	d_array = cuda.device_array(4, dtype = np.complex64)
	launcher[1,1](d_array)
	output = d_array.copy_to_host()
	print output

main()
```
yields:
```python
Traceback (most recent call last):
  File "<path>numba/numba/cuda/cudadrv/driver.py", line 1432, in complete
    driver.cuLinkComplete(self.handle, byref(cubin), byref(size))
  File "<path>numba/numba/cuda/cudadrv/driver.py", line 288, in safe_cuda_api_call
    self._check_error(fname, retcode)
  File "<path>numba/numba/cuda/cudadrv/driver.py", line 323, in _check_error
    raise CudaAPIError(retcode, msg)
numba.cuda.cudadrv.driver.CudaAPIError: [999] Call to cuLinkComplete results in CUDA_ERROR_UNKNOWN

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "issue_13.py", line 39, in <module>
    main()
  File "issue_13.py", line 35, in main
    launcher[1,1](d_array)
  File "<path>numba/numba/cuda/compiler.py", line 701, in __call__
    kernel = self.specialize(*args)
  File "<path>numba/numba/cuda/compiler.py", line 712, in specialize
    kernel = self.compile(argtypes)
  File "<path>numba/numba/cuda/compiler.py", line 730, in compile
    kernel.bind()
  File "<path>numba/numba/cuda/compiler.py", line 489, in bind
    self._func.get()
  File "<path>numba/numba/cuda/compiler.py", line 377, in get
    cubin, _size = linker.complete()
  File "<path>numba/numba/cuda/cudadrv/driver.py", line 1434, in complete
    raise LinkerError("%s\n%s" % (e, self.error_log))
numba.cuda.cudadrv.driver.LinkerError: [999] Call to cuLinkComplete results in CUDA_ERROR_UNKNOWN
error   : Undefined reference to 'numba_cpow' in '<cudapy-ptx>'
```
the `cmath.sqrt()` creates a complex domain result with is then raised to the power `1/3`, this ends up as a call to `numba_cpow` which is fine on the CPU but doesn't have an impl in CUDA, as a result there's a linking error.

Two issues:
1. Can the additional cmath function impls for the CPU be done for CUDA.
2. If not can this problem be caught earlier in the complication chain and a better message produced such that is doesn't appear as a link error.


--- END ISSUE ---

Below are some code segments, each from a relevant file. One or more of these files may contain bugs.
--- BEGIN FILES ---
Path: `numba/cuda/libdevice.py`
Content:
```
1 import math
2 from llvmlite.llvmpy.core import Type
3 from numba.core import types, cgutils
4 from numba.core.imputils import Registry
5 
6 registry = Registry()
7 lower = registry.lower
8 
9 float_set = types.float32, types.float64
10 
11 
12 def bool_implement(nvname, ty):
13     def core(context, builder, sig, args):
14         assert sig.return_type == types.boolean, nvname
15         fty = context.get_value_type(ty)
16         lmod = builder.module
17         fnty = Type.function(Type.int(), [fty])
18         fn = lmod.get_or_insert_function(fnty, name=nvname)
19         result = builder.call(fn, args)
20         return context.cast(builder, result, types.int32, types.boolean)
21 
22     return core
23 
24 
25 def unary_implement(nvname, ty):
26     def core(context, builder, sig, args):
27         fty = context.get_value_type(ty)
28         lmod = builder.module
29         fnty = Type.function(fty, [fty])
30         fn = lmod.get_or_insert_function(fnty, name=nvname)
31         return builder.call(fn, args)
32 
33     return core
34 
35 
36 def binary_implement(nvname, ty):
37     def core(context, builder, sig, args):
38         fty = context.get_value_type(ty)
39         lmod = builder.module
40         fnty = Type.function(fty, [fty, fty])
41         fn = lmod.get_or_insert_function(fnty, name=nvname)
42         return builder.call(fn, args)
43 
44     return core
45 
46 
47 def powi_implement(nvname):
48     def core(context, builder, sig, args):
49         [base, pow] = args
50         [basety, powty] = sig.args
51         lmod = builder.module
52         fty = context.get_value_type(basety)
53         ity = context.get_value_type(types.int32)
54         fnty = Type.function(fty, [fty, ity])
55         fn = lmod.get_or_insert_function(fnty, name=nvname)
56         return builder.call(fn, [base, pow])
57 
58     return core
59 
60 
61 lower(math.pow, types.float32, types.int32)(powi_implement('__nv_powif'))
62 lower(math.pow, types.float64, types.int32)(powi_implement('__nv_powi'))
63 
64 
65 def frexp_implement(nvname):
66     def core(context, builder, sig, args):
67         fracty, expty = sig.return_type
68         float_type = context.get_value_type(fracty)
69         int_type = context.get_value_type(expty)
70         fnty = Type.function(float_type, [float_type, Type.pointer(int_type)])
71 
72         fn = builder.module.get_or_insert_function(fnty, name=nvname)
73         expptr = cgutils.alloca_once(builder, int_type, name='exp')
74 
75         ret = builder.call(fn, (args[0], expptr))
76         return cgutils.pack_struct(builder, (ret, builder.load(expptr)))
77 
78     return core
79 
80 
81 lower(math.frexp, types.float32)(frexp_implement('__nv_frexpf'))
82 lower(math.frexp, types.float64)(frexp_implement('__nv_frexp'))
83 
84 lower(math.ldexp, types.float32, types.int32)(powi_implement('__nv_ldexpf'))
85 lower(math.ldexp, types.float64, types.int32)(powi_implement('__nv_ldexp'))
86 
87 
88 booleans = []
89 booleans += [('__nv_isnand', '__nv_isnanf', math.isnan)]
90 booleans += [('__nv_isinfd', '__nv_isinff', math.isinf)]
91 booleans += [('__nv_isfinited', '__nv_finitef', math.isfinite)]
92 
93 unarys = []
94 unarys += [('__nv_ceil', '__nv_ceilf', math.ceil)]
95 unarys += [('__nv_floor', '__nv_floorf', math.floor)]
96 unarys += [('__nv_fabs', '__nv_fabsf', math.fabs)]
97 unarys += [('__nv_exp', '__nv_expf', math.exp)]
98 unarys += [('__nv_expm1', '__nv_expm1f', math.expm1)]
99 unarys += [('__nv_erf', '__nv_erff', math.erf)]
100 unarys += [('__nv_erfc', '__nv_erfcf', math.erfc)]
101 unarys += [('__nv_tgamma', '__nv_tgammaf', math.gamma)]
102 unarys += [('__nv_lgamma', '__nv_lgammaf', math.lgamma)]
103 unarys += [('__nv_sqrt', '__nv_sqrtf', math.sqrt)]
104 unarys += [('__nv_log', '__nv_logf', math.log)]
105 unarys += [('__nv_log10', '__nv_log10f', math.log10)]
106 unarys += [('__nv_log1p', '__nv_log1pf', math.log1p)]
107 unarys += [('__nv_acosh', '__nv_acoshf', math.acosh)]
108 unarys += [('__nv_acos', '__nv_acosf', math.acos)]
109 unarys += [('__nv_cos', '__nv_cosf', math.cos)]
110 unarys += [('__nv_cosh', '__nv_coshf', math.cosh)]
111 unarys += [('__nv_asinh', '__nv_asinhf', math.asinh)]
112 unarys += [('__nv_asin', '__nv_asinf', math.asin)]
113 unarys += [('__nv_sin', '__nv_sinf', math.sin)]
114 unarys += [('__nv_sinh', '__nv_sinhf', math.sinh)]
115 unarys += [('__nv_atan', '__nv_atanf', math.atan)]
116 unarys += [('__nv_atanh', '__nv_atanhf', math.atanh)]
117 unarys += [('__nv_tan', '__nv_tanf', math.tan)]
118 unarys += [('__nv_tanh', '__nv_tanhf', math.tanh)]
119 
120 binarys = []
121 binarys += [('__nv_copysign', '__nv_copysignf', math.copysign)]
122 binarys += [('__nv_atan2', '__nv_atan2f', math.atan2)]
123 binarys += [('__nv_pow', '__nv_powf', math.pow)]
124 binarys += [('__nv_fmod', '__nv_fmodf', math.fmod)]
125 binarys += [('__nv_hypot', '__nv_hypotf', math.hypot)]
126 
127 
128 for name64, name32, key in booleans:
129     impl64 = bool_implement(name64, types.float64)
130     lower(key, types.float64)(impl64)
131     impl32 = bool_implement(name32, types.float32)
132     lower(key, types.float32)(impl32)
133 
134 
135 for name64, name32, key in unarys:
136     impl64 = unary_implement(name64, types.float64)
137     lower(key, types.float64)(impl64)
138     impl32 = unary_implement(name32, types.float32)
139     lower(key, types.float32)(impl32)
140 
141 for name64, name32, key in binarys:
142     impl64 = binary_implement(name64, types.float64)
143     lower(key, types.float64, types.float64)(impl64)
144     impl32 = binary_implement(name32, types.float32)
145     lower(key, types.float32, types.float32)(impl32)
146 
147 
148 def modf_implement(nvname, ty):
149     def core(context, builder, sig, args):
150         arg, = args
151         argty, = sig.args
152         fty = context.get_value_type(argty)
153         lmod = builder.module
154         ptr = cgutils.alloca_once(builder, fty)
155         fnty = Type.function(fty, [fty, fty.as_pointer()])
156         fn = lmod.get_or_insert_function(fnty, name=nvname)
157         out = builder.call(fn, [arg, ptr])
158         ret = context.make_tuple(builder, types.UniTuple(argty, 2),
159                                  [out, builder.load(ptr)])
160         return ret
161     return core
162 
163 
164 for (ty, intrin) in ((types.float64, '__nv_modf',),
165                      (types.float32, '__nv_modff',)):
166     lower(math.modf, ty)(modf_implement(intrin, ty))
167
```
--- 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/numba/cuda/libdevice.py b/numba/cuda/libdevice.py
--- a/numba/cuda/libdevice.py
+++ b/numba/cuda/libdevice.py
@@ -1,4 +1,5 @@
 import math
+import operator
 from llvmlite.llvmpy.core import Type
 from numba.core import types, cgutils
 from numba.core.imputils import Registry
@@ -164,3 +165,41 @@
 for (ty, intrin) in ((types.float64, '__nv_modf',),
                      (types.float32, '__nv_modff',)):
     lower(math.modf, ty)(modf_implement(intrin, ty))
+
+
+# Complex power implementations - translations of _Py_c_pow from CPython
+# https://github.com/python/cpython/blob/a755410e054e1e2390de5830befc08fe80706c66/Objects/complexobject.c#L123-L151
+#
+# The complex64 variant casts all constants and some variables to ensure that
+# as much computation is done in single precision as possible. A small number
+# of operations are still done in 64-bit, but these come from libdevice code.
+
+def cpow_implement(fty, cty):
+    def core(context, builder, sig, args):
+        def cpow_internal(a, b):
+
+            if b.real == fty(0.0) and b.imag == fty(0.0):
+                return cty(1.0) + cty(0.0j)
+            elif a.real == fty(0.0) and b.real == fty(0.0):
+                return cty(0.0) + cty(0.0j)
+
+            vabs = math.hypot(a.real, a.imag)
+            len = math.pow(vabs, b.real)
+            at = math.atan2(a.imag, a.real)
+            phase = at * b.real
+            if b.imag != fty(0.0):
+                len /= math.exp(at * b.imag)
+                phase += b.imag * math.log(vabs)
+
+            return len * (cty(math.cos(phase)) +
+                          cty(math.sin(phase) * cty(1.0j)))
+
+        return context.compile_internal(builder, cpow_internal, sig, args)
+
+    lower(operator.pow, cty, cty)(core)
+    lower(operator.ipow, cty, cty)(core)
+    lower(pow, cty, cty)(core)
+
+
+cpow_implement(types.float32, types.complex64)
+cpow_implement(types.float64, types.complex128)

{"golden_diff": "diff --git a/numba/cuda/libdevice.py b/numba/cuda/libdevice.py\n--- a/numba/cuda/libdevice.py\n+++ b/numba/cuda/libdevice.py\n@@ -1,4 +1,5 @@\n import math\n+import operator\n from llvmlite.llvmpy.core import Type\n from numba.core import types, cgutils\n from numba.core.imputils import Registry\n@@ -164,3 +165,41 @@\n for (ty, intrin) in ((types.float64, '__nv_modf',),\n                      (types.float32, '__nv_modff',)):\n     lower(math.modf, ty)(modf_implement(intrin, ty))\n+\n+\n+# Complex power implementations - translations of _Py_c_pow from CPython\n+# https://github.com/python/cpython/blob/a755410e054e1e2390de5830befc08fe80706c66/Objects/complexobject.c#L123-L151\n+#\n+# The complex64 variant casts all constants and some variables to ensure that\n+# as much computation is done in single precision as possible. A small number\n+# of operations are still done in 64-bit, but these come from libdevice code.\n+\n+def cpow_implement(fty, cty):\n+    def core(context, builder, sig, args):\n+        def cpow_internal(a, b):\n+\n+            if b.real == fty(0.0) and b.imag == fty(0.0):\n+                return cty(1.0) + cty(0.0j)\n+            elif a.real == fty(0.0) and b.real == fty(0.0):\n+                return cty(0.0) + cty(0.0j)\n+\n+            vabs = math.hypot(a.real, a.imag)\n+            len = math.pow(vabs, b.real)\n+            at = math.atan2(a.imag, a.real)\n+            phase = at * b.real\n+            if b.imag != fty(0.0):\n+                len /= math.exp(at * b.imag)\n+                phase += b.imag * math.log(vabs)\n+\n+            return len * (cty(math.cos(phase)) +\n+                          cty(math.sin(phase) * cty(1.0j)))\n+\n+        return context.compile_internal(builder, cpow_internal, sig, args)\n+\n+    lower(operator.pow, cty, cty)(core)\n+    lower(operator.ipow, cty, cty)(core)\n+    lower(pow, cty, cty)(core)\n+\n+\n+cpow_implement(types.float32, types.complex64)\n+cpow_implement(types.float64, types.complex128)\n", "issue": "cmath linking issues in CUDA\nFrom the mailing list:\r\nhttps://groups.google.com/a/continuum.io/forum/#!topic/numba-users/oBcaFcsD-x4\r\nThis:\r\n```python\r\nfrom numba import cuda\r\nimport cmath\r\nimport numpy as np\r\n\r\n@cuda.jit(device = True)\r\ndef quarticRoots(a, b, c, d, e):\r\n\r\n\tdel_0 = c**2 - 3.*b*d + 12.*a*e\r\n\tdel_1 = 2.*c**3 - 9.*b*c*d + 27*e*b**2 + 27*a*d**2 - 72*a*c*e\r\n\r\n\tp = (8.*a*c - 3.*b**2)/(8.*a**2)\r\n\tq = (b**3 - 4.*a*b*c + 8.*d*a**2)/(8.*a**3)\r\n\r\n\tQ = ((del_1 + cmath.sqrt(del_1**2 - 4.*del_0**3))/2.)**(1./3.)#**(1./3.) is the problem.\r\n\tS = 0.5*cmath.sqrt(-2./3.*p + 1/(3.*a)*(Q + del_0/Q))\r\n\r\n\tx1 = -b/(4.*a) - S + 0.5*cmath.sqrt(-4.*S**2 - 2.*p + q/S)\r\n\tx2 = -b/(4.*a) - S - 0.5*cmath.sqrt(-4.*S**2 - 2.*p + q/S)\r\n\tx3 = -b/(4.*a) + S + 0.5*cmath.sqrt(-4.*S**2 - 2.*p - q/S)\r\n\tx4 = -b/(4.*a) + S - 0.5*cmath.sqrt(-4.*S**2 - 2.*p - q/S)\r\n\r\n\treturn (x1, x2, x3, x4)\r\n\r\n@cuda.jit\r\ndef launcher(d_array):\r\n\toutput = quarticRoots(2.,4.,2.,1.,1.)\r\n\r\n\td_array[0] = output[0]\r\n\td_array[1] = output[1]\r\n\td_array[2] = output[2]\r\n\td_array[3] = output[3]\r\n\r\ndef main():\r\n\td_array = cuda.device_array(4, dtype = np.complex64)\r\n\tlauncher[1,1](d_array)\r\n\toutput = d_array.copy_to_host()\r\n\tprint output\r\n\r\nmain()\r\n```\r\nyields:\r\n```python\r\nTraceback (most recent call last):\r\n  File \"<path>numba/numba/cuda/cudadrv/driver.py\", line 1432, in complete\r\n    driver.cuLinkComplete(self.handle, byref(cubin), byref(size))\r\n  File \"<path>numba/numba/cuda/cudadrv/driver.py\", line 288, in safe_cuda_api_call\r\n    self._check_error(fname, retcode)\r\n  File \"<path>numba/numba/cuda/cudadrv/driver.py\", line 323, in _check_error\r\n    raise CudaAPIError(retcode, msg)\r\nnumba.cuda.cudadrv.driver.CudaAPIError: [999] Call to cuLinkComplete results in CUDA_ERROR_UNKNOWN\r\n\r\nDuring handling of the above exception, another exception occurred:\r\n\r\nTraceback (most recent call last):\r\n  File \"issue_13.py\", line 39, in <module>\r\n    main()\r\n  File \"issue_13.py\", line 35, in main\r\n    launcher[1,1](d_array)\r\n  File \"<path>numba/numba/cuda/compiler.py\", line 701, in __call__\r\n    kernel = self.specialize(*args)\r\n  File \"<path>numba/numba/cuda/compiler.py\", line 712, in specialize\r\n    kernel = self.compile(argtypes)\r\n  File \"<path>numba/numba/cuda/compiler.py\", line 730, in compile\r\n    kernel.bind()\r\n  File \"<path>numba/numba/cuda/compiler.py\", line 489, in bind\r\n    self._func.get()\r\n  File \"<path>numba/numba/cuda/compiler.py\", line 377, in get\r\n    cubin, _size = linker.complete()\r\n  File \"<path>numba/numba/cuda/cudadrv/driver.py\", line 1434, in complete\r\n    raise LinkerError(\"%s\\n%s\" % (e, self.error_log))\r\nnumba.cuda.cudadrv.driver.LinkerError: [999] Call to cuLinkComplete results in CUDA_ERROR_UNKNOWN\r\nerror   : Undefined reference to 'numba_cpow' in '<cudapy-ptx>'\r\n```\r\nthe `cmath.sqrt()` creates a complex domain result with is then raised to the power `1/3`, this ends up as a call to `numba_cpow` which is fine on the CPU but doesn't have an impl in CUDA, as a result there's a linking error.\r\n\r\nTwo issues:\r\n1. Can the additional cmath function impls for the CPU be done for CUDA.\r\n2. If not can this problem be caught earlier in the complication chain and a better message produced such that is doesn't appear as a link error.\r\n\n", "before_files": [{"content": "import math\nfrom llvmlite.llvmpy.core import Type\nfrom numba.core import types, cgutils\nfrom numba.core.imputils import Registry\n\nregistry = Registry()\nlower = registry.lower\n\nfloat_set = types.float32, types.float64\n\n\ndef bool_implement(nvname, ty):\n    def core(context, builder, sig, args):\n        assert sig.return_type == types.boolean, nvname\n        fty = context.get_value_type(ty)\n        lmod = builder.module\n        fnty = Type.function(Type.int(), [fty])\n        fn = lmod.get_or_insert_function(fnty, name=nvname)\n        result = builder.call(fn, args)\n        return context.cast(builder, result, types.int32, types.boolean)\n\n    return core\n\n\ndef unary_implement(nvname, ty):\n    def core(context, builder, sig, args):\n        fty = context.get_value_type(ty)\n        lmod = builder.module\n        fnty = Type.function(fty, [fty])\n        fn = lmod.get_or_insert_function(fnty, name=nvname)\n        return builder.call(fn, args)\n\n    return core\n\n\ndef binary_implement(nvname, ty):\n    def core(context, builder, sig, args):\n        fty = context.get_value_type(ty)\n        lmod = builder.module\n        fnty = Type.function(fty, [fty, fty])\n        fn = lmod.get_or_insert_function(fnty, name=nvname)\n        return builder.call(fn, args)\n\n    return core\n\n\ndef powi_implement(nvname):\n    def core(context, builder, sig, args):\n        [base, pow] = args\n        [basety, powty] = sig.args\n        lmod = builder.module\n        fty = context.get_value_type(basety)\n        ity = context.get_value_type(types.int32)\n        fnty = Type.function(fty, [fty, ity])\n        fn = lmod.get_or_insert_function(fnty, name=nvname)\n        return builder.call(fn, [base, pow])\n\n    return core\n\n\nlower(math.pow, types.float32, types.int32)(powi_implement('__nv_powif'))\nlower(math.pow, types.float64, types.int32)(powi_implement('__nv_powi'))\n\n\ndef frexp_implement(nvname):\n    def core(context, builder, sig, args):\n        fracty, expty = sig.return_type\n        float_type = context.get_value_type(fracty)\n        int_type = context.get_value_type(expty)\n        fnty = Type.function(float_type, [float_type, Type.pointer(int_type)])\n\n        fn = builder.module.get_or_insert_function(fnty, name=nvname)\n        expptr = cgutils.alloca_once(builder, int_type, name='exp')\n\n        ret = builder.call(fn, (args[0], expptr))\n        return cgutils.pack_struct(builder, (ret, builder.load(expptr)))\n\n    return core\n\n\nlower(math.frexp, types.float32)(frexp_implement('__nv_frexpf'))\nlower(math.frexp, types.float64)(frexp_implement('__nv_frexp'))\n\nlower(math.ldexp, types.float32, types.int32)(powi_implement('__nv_ldexpf'))\nlower(math.ldexp, types.float64, types.int32)(powi_implement('__nv_ldexp'))\n\n\nbooleans = []\nbooleans += [('__nv_isnand', '__nv_isnanf', math.isnan)]\nbooleans += [('__nv_isinfd', '__nv_isinff', math.isinf)]\nbooleans += [('__nv_isfinited', '__nv_finitef', math.isfinite)]\n\nunarys = []\nunarys += [('__nv_ceil', '__nv_ceilf', math.ceil)]\nunarys += [('__nv_floor', '__nv_floorf', math.floor)]\nunarys += [('__nv_fabs', '__nv_fabsf', math.fabs)]\nunarys += [('__nv_exp', '__nv_expf', math.exp)]\nunarys += [('__nv_expm1', '__nv_expm1f', math.expm1)]\nunarys += [('__nv_erf', '__nv_erff', math.erf)]\nunarys += [('__nv_erfc', '__nv_erfcf', math.erfc)]\nunarys += [('__nv_tgamma', '__nv_tgammaf', math.gamma)]\nunarys += [('__nv_lgamma', '__nv_lgammaf', math.lgamma)]\nunarys += [('__nv_sqrt', '__nv_sqrtf', math.sqrt)]\nunarys += [('__nv_log', '__nv_logf', math.log)]\nunarys += [('__nv_log10', '__nv_log10f', math.log10)]\nunarys += [('__nv_log1p', '__nv_log1pf', math.log1p)]\nunarys += [('__nv_acosh', '__nv_acoshf', math.acosh)]\nunarys += [('__nv_acos', '__nv_acosf', math.acos)]\nunarys += [('__nv_cos', '__nv_cosf', math.cos)]\nunarys += [('__nv_cosh', '__nv_coshf', math.cosh)]\nunarys += [('__nv_asinh', '__nv_asinhf', math.asinh)]\nunarys += [('__nv_asin', '__nv_asinf', math.asin)]\nunarys += [('__nv_sin', '__nv_sinf', math.sin)]\nunarys += [('__nv_sinh', '__nv_sinhf', math.sinh)]\nunarys += [('__nv_atan', '__nv_atanf', math.atan)]\nunarys += [('__nv_atanh', '__nv_atanhf', math.atanh)]\nunarys += [('__nv_tan', '__nv_tanf', math.tan)]\nunarys += [('__nv_tanh', '__nv_tanhf', math.tanh)]\n\nbinarys = []\nbinarys += [('__nv_copysign', '__nv_copysignf', math.copysign)]\nbinarys += [('__nv_atan2', '__nv_atan2f', math.atan2)]\nbinarys += [('__nv_pow', '__nv_powf', math.pow)]\nbinarys += [('__nv_fmod', '__nv_fmodf', math.fmod)]\nbinarys += [('__nv_hypot', '__nv_hypotf', math.hypot)]\n\n\nfor name64, name32, key in booleans:\n    impl64 = bool_implement(name64, types.float64)\n    lower(key, types.float64)(impl64)\n    impl32 = bool_implement(name32, types.float32)\n    lower(key, types.float32)(impl32)\n\n\nfor name64, name32, key in unarys:\n    impl64 = unary_implement(name64, types.float64)\n    lower(key, types.float64)(impl64)\n    impl32 = unary_implement(name32, types.float32)\n    lower(key, types.float32)(impl32)\n\nfor name64, name32, key in binarys:\n    impl64 = binary_implement(name64, types.float64)\n    lower(key, types.float64, types.float64)(impl64)\n    impl32 = binary_implement(name32, types.float32)\n    lower(key, types.float32, types.float32)(impl32)\n\n\ndef modf_implement(nvname, ty):\n    def core(context, builder, sig, args):\n        arg, = args\n        argty, = sig.args\n        fty = context.get_value_type(argty)\n        lmod = builder.module\n        ptr = cgutils.alloca_once(builder, fty)\n        fnty = Type.function(fty, [fty, fty.as_pointer()])\n        fn = lmod.get_or_insert_function(fnty, name=nvname)\n        out = builder.call(fn, [arg, ptr])\n        ret = context.make_tuple(builder, types.UniTuple(argty, 2),\n                                 [out, builder.load(ptr)])\n        return ret\n    return core\n\n\nfor (ty, intrin) in ((types.float64, '__nv_modf',),\n                     (types.float32, '__nv_modff',)):\n    lower(math.modf, ty)(modf_implement(intrin, ty))\n", "path": "numba/cuda/libdevice.py"}], "after_files": [{"content": "import math\nimport operator\nfrom llvmlite.llvmpy.core import Type\nfrom numba.core import types, cgutils\nfrom numba.core.imputils import Registry\n\nregistry = Registry()\nlower = registry.lower\n\nfloat_set = types.float32, types.float64\n\n\ndef bool_implement(nvname, ty):\n    def core(context, builder, sig, args):\n        assert sig.return_type == types.boolean, nvname\n        fty = context.get_value_type(ty)\n        lmod = builder.module\n        fnty = Type.function(Type.int(), [fty])\n        fn = lmod.get_or_insert_function(fnty, name=nvname)\n        result = builder.call(fn, args)\n        return context.cast(builder, result, types.int32, types.boolean)\n\n    return core\n\n\ndef unary_implement(nvname, ty):\n    def core(context, builder, sig, args):\n        fty = context.get_value_type(ty)\n        lmod = builder.module\n        fnty = Type.function(fty, [fty])\n        fn = lmod.get_or_insert_function(fnty, name=nvname)\n        return builder.call(fn, args)\n\n    return core\n\n\ndef binary_implement(nvname, ty):\n    def core(context, builder, sig, args):\n        fty = context.get_value_type(ty)\n        lmod = builder.module\n        fnty = Type.function(fty, [fty, fty])\n        fn = lmod.get_or_insert_function(fnty, name=nvname)\n        return builder.call(fn, args)\n\n    return core\n\n\ndef powi_implement(nvname):\n    def core(context, builder, sig, args):\n        [base, pow] = args\n        [basety, powty] = sig.args\n        lmod = builder.module\n        fty = context.get_value_type(basety)\n        ity = context.get_value_type(types.int32)\n        fnty = Type.function(fty, [fty, ity])\n        fn = lmod.get_or_insert_function(fnty, name=nvname)\n        return builder.call(fn, [base, pow])\n\n    return core\n\n\nlower(math.pow, types.float32, types.int32)(powi_implement('__nv_powif'))\nlower(math.pow, types.float64, types.int32)(powi_implement('__nv_powi'))\n\n\ndef frexp_implement(nvname):\n    def core(context, builder, sig, args):\n        fracty, expty = sig.return_type\n        float_type = context.get_value_type(fracty)\n        int_type = context.get_value_type(expty)\n        fnty = Type.function(float_type, [float_type, Type.pointer(int_type)])\n\n        fn = builder.module.get_or_insert_function(fnty, name=nvname)\n        expptr = cgutils.alloca_once(builder, int_type, name='exp')\n\n        ret = builder.call(fn, (args[0], expptr))\n        return cgutils.pack_struct(builder, (ret, builder.load(expptr)))\n\n    return core\n\n\nlower(math.frexp, types.float32)(frexp_implement('__nv_frexpf'))\nlower(math.frexp, types.float64)(frexp_implement('__nv_frexp'))\n\nlower(math.ldexp, types.float32, types.int32)(powi_implement('__nv_ldexpf'))\nlower(math.ldexp, types.float64, types.int32)(powi_implement('__nv_ldexp'))\n\n\nbooleans = []\nbooleans += [('__nv_isnand', '__nv_isnanf', math.isnan)]\nbooleans += [('__nv_isinfd', '__nv_isinff', math.isinf)]\nbooleans += [('__nv_isfinited', '__nv_finitef', math.isfinite)]\n\nunarys = []\nunarys += [('__nv_ceil', '__nv_ceilf', math.ceil)]\nunarys += [('__nv_floor', '__nv_floorf', math.floor)]\nunarys += [('__nv_fabs', '__nv_fabsf', math.fabs)]\nunarys += [('__nv_exp', '__nv_expf', math.exp)]\nunarys += [('__nv_expm1', '__nv_expm1f', math.expm1)]\nunarys += [('__nv_erf', '__nv_erff', math.erf)]\nunarys += [('__nv_erfc', '__nv_erfcf', math.erfc)]\nunarys += [('__nv_tgamma', '__nv_tgammaf', math.gamma)]\nunarys += [('__nv_lgamma', '__nv_lgammaf', math.lgamma)]\nunarys += [('__nv_sqrt', '__nv_sqrtf', math.sqrt)]\nunarys += [('__nv_log', '__nv_logf', math.log)]\nunarys += [('__nv_log10', '__nv_log10f', math.log10)]\nunarys += [('__nv_log1p', '__nv_log1pf', math.log1p)]\nunarys += [('__nv_acosh', '__nv_acoshf', math.acosh)]\nunarys += [('__nv_acos', '__nv_acosf', math.acos)]\nunarys += [('__nv_cos', '__nv_cosf', math.cos)]\nunarys += [('__nv_cosh', '__nv_coshf', math.cosh)]\nunarys += [('__nv_asinh', '__nv_asinhf', math.asinh)]\nunarys += [('__nv_asin', '__nv_asinf', math.asin)]\nunarys += [('__nv_sin', '__nv_sinf', math.sin)]\nunarys += [('__nv_sinh', '__nv_sinhf', math.sinh)]\nunarys += [('__nv_atan', '__nv_atanf', math.atan)]\nunarys += [('__nv_atanh', '__nv_atanhf', math.atanh)]\nunarys += [('__nv_tan', '__nv_tanf', math.tan)]\nunarys += [('__nv_tanh', '__nv_tanhf', math.tanh)]\n\nbinarys = []\nbinarys += [('__nv_copysign', '__nv_copysignf', math.copysign)]\nbinarys += [('__nv_atan2', '__nv_atan2f', math.atan2)]\nbinarys += [('__nv_pow', '__nv_powf', math.pow)]\nbinarys += [('__nv_fmod', '__nv_fmodf', math.fmod)]\nbinarys += [('__nv_hypot', '__nv_hypotf', math.hypot)]\n\n\nfor name64, name32, key in booleans:\n    impl64 = bool_implement(name64, types.float64)\n    lower(key, types.float64)(impl64)\n    impl32 = bool_implement(name32, types.float32)\n    lower(key, types.float32)(impl32)\n\n\nfor name64, name32, key in unarys:\n    impl64 = unary_implement(name64, types.float64)\n    lower(key, types.float64)(impl64)\n    impl32 = unary_implement(name32, types.float32)\n    lower(key, types.float32)(impl32)\n\nfor name64, name32, key in binarys:\n    impl64 = binary_implement(name64, types.float64)\n    lower(key, types.float64, types.float64)(impl64)\n    impl32 = binary_implement(name32, types.float32)\n    lower(key, types.float32, types.float32)(impl32)\n\n\ndef modf_implement(nvname, ty):\n    def core(context, builder, sig, args):\n        arg, = args\n        argty, = sig.args\n        fty = context.get_value_type(argty)\n        lmod = builder.module\n        ptr = cgutils.alloca_once(builder, fty)\n        fnty = Type.function(fty, [fty, fty.as_pointer()])\n        fn = lmod.get_or_insert_function(fnty, name=nvname)\n        out = builder.call(fn, [arg, ptr])\n        ret = context.make_tuple(builder, types.UniTuple(argty, 2),\n                                 [out, builder.load(ptr)])\n        return ret\n    return core\n\n\nfor (ty, intrin) in ((types.float64, '__nv_modf',),\n                     (types.float32, '__nv_modff',)):\n    lower(math.modf, ty)(modf_implement(intrin, ty))\n\n\n# Complex power implementations - translations of _Py_c_pow from CPython\n# https://github.com/python/cpython/blob/a755410e054e1e2390de5830befc08fe80706c66/Objects/complexobject.c#L123-L151\n#\n# The complex64 variant casts all constants and some variables to ensure that\n# as much computation is done in single precision as possible. A small number\n# of operations are still done in 64-bit, but these come from libdevice code.\n\ndef cpow_implement(fty, cty):\n    def core(context, builder, sig, args):\n        def cpow_internal(a, b):\n\n            if b.real == fty(0.0) and b.imag == fty(0.0):\n                return cty(1.0) + cty(0.0j)\n            elif a.real == fty(0.0) and b.real == fty(0.0):\n                return cty(0.0) + cty(0.0j)\n\n            vabs = math.hypot(a.real, a.imag)\n            len = math.pow(vabs, b.real)\n            at = math.atan2(a.imag, a.real)\n            phase = at * b.real\n            if b.imag != fty(0.0):\n                len /= math.exp(at * b.imag)\n                phase += b.imag * math.log(vabs)\n\n            return len * (cty(math.cos(phase)) +\n                          cty(math.sin(phase) * cty(1.0j)))\n\n        return context.compile_internal(builder, cpow_internal, sig, args)\n\n    lower(operator.pow, cty, cty)(core)\n    lower(operator.ipow, cty, cty)(core)\n    lower(pow, cty, cty)(core)\n\n\ncpow_implement(types.float32, types.complex64)\ncpow_implement(types.float64, types.complex128)\n", "path": "numba/cuda/libdevice.py"}]}