Taking the 1/3 power causing LinkerError error

[Chris Uchytil]

Attempting to find the root of a quartic polynomial. For some reason the **(1./3.) is causing issues. If I comment it out everything runs, but obviously you get the wrong answer.

CODE:

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()

ERROR:

Traceback (most recent call last):

  File "root_test.py", line 39, in <module>

    main()

  File "root_test.py", line 35, in main

    launcher[1,1](d_array)

  File "/home/uchytilc/anaconda2/lib/python2.7/site-packages/numba/cuda/compiler.py", line 701, in __call__

    kernel = self.specialize(*args)

  File "/home/uchytilc/anaconda2/lib/python2.7/site-packages/numba/cuda/compiler.py", line 712, in specialize

    kernel = self.compile(argtypes)

  File "/home/uchytilc/anaconda2/lib/python2.7/site-packages/numba/cuda/compiler.py", line 730, in compile

    kernel.bind()

  File "/home/uchytilc/anaconda2/lib/python2.7/site-packages/numba/cuda/compiler.py", line 489, in bind

    self._func.get()

  File "/home/uchytilc/anaconda2/lib/python2.7/site-packages/numba/cuda/compiler.py", line 377, in get

    cubin, _size = linker.complete()

  File "/home/uchytilc/anaconda2/lib/python2.7/site-packages/numba/cuda/cudadrv/driver.py", line 1426, 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>'

[Chris Uchytil]

After playing around a little more this seems to make it work. Not sure why though.

Q = ((del_1 + cmath.sqrt(del_1**2 - 4.*del_0**3))/2.).real**(1./3.)

[stuart]

Thanks for the report.

I think this is because the `cmath.sqrt()` will create a result in the complex domain and then doing `**(1./3.)` will be interpreted as "use complex domain power function to raise expression to 1/3", there is no `cmath.pow`. To get around this, on CPUs a Numba implemented function is used to perform `cmath.pow()` like functionality hence the function ref in to `numba_cpow`, this is clearly leaking into the CUDA code without a definition on the CUDA side. The use of `.real` hits a supported CUDA `math.pow` equivalent, hence works.

Essentially I think there's two issues here:
1. cmath support in CUDA not matching CPU.
2. That the lack of CUDA implementation for `numba_cpow()` should have been picked up before it appeared as a link error.

I've moved this to the Numba issue tracker https://github.com/numba/numba/issues/2495.

Thanks,

--
stuart