Calling a CUDA vectorized function from another CUDA vectorized function

[Andrea Cogliati]

I'm trying to speedup some Python code doing complex math, but since the complex power function is not available yet in Numba, I'm developing my own. This is what I tried

import numpy as np

import math

from numba import vectorize, complex128, float64

@vectorize([float64(complex128)],target='cuda')

def carg(z):

    return math.atan2(z.imag, z.real)

@vectorize([complex128(complex128,float64)],target='cuda')

def cpow(z, x):

    return np.complex((abs(z) ** x) * math.cos(x * carg(z)), (abs(z) ** x) * math.sin(x * carg(z)))

z = 1+2j

x = np.arange(0.0,10.0,0.5)

print(cpow(z,x)) 

However, it fails with an Untyped global name error:

Traceback (most recent call last):

  File "test.py", line 9, in <module>

    @vectorize([complex128(complex128,float64)],target='cuda')

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/npyufunc/decorators.py", line 120, in wrap

    vec.add(sig)

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/npyufunc/deviceufunc.py", line 384, in add

    corefn, return_type = self._compile_core(devfnsig)

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/cuda/vectorizers.py", line 17, in _compile_core

    cudevfn = cuda.jit(sig, device=True, inline=True)(self.pyfunc)

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/cuda/decorators.py", line 102, in device_jit

    debug=debug)

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/cuda/compiler.py", line 160, in compile_device

    return DeviceFunction(pyfunc, return_type, args, inline=True, debug=False)

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/cuda/compiler.py", line 191, in __init__

    debug=self.debug, inline=self.inline)

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/cuda/compiler.py", line 36, in core

    return fn(*args, **kwargs)

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/cuda/compiler.py", line 64, in compile_cuda

    locals={})

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/compiler.py", line 779, in compile_extra

    return pipeline.compile_extra(func)

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/compiler.py", line 362, in compile_extra

    return self._compile_bytecode()

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/compiler.py", line 738, in _compile_bytecode

    return self._compile_core()

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/compiler.py", line 725, in _compile_core

    res = pm.run(self.status)

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/compiler.py", line 248, in run

    raise patched_exception

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/compiler.py", line 240, in run

    stage()

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/compiler.py", line 454, in stage_nopython_frontend

    self.locals)

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/compiler.py", line 880, in type_inference_stage

    infer.build_constraint()

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/typeinfer.py", line 802, in build_constraint

    self.constrain_statement(inst)

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/typeinfer.py", line 961, in constrain_statement

    self.typeof_assign(inst)

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/typeinfer.py", line 1023, in typeof_assign

    self.typeof_global(inst, inst.target, value)

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/typeinfer.py", line 1119, in typeof_global

    typ = self.resolve_value_type(inst, gvar.value)

  File "/media/raid0/andrea/anaconda3/lib/python3.6/site-packages/numba/typeinfer.py", line 1042, in resolve_value_type

    raise TypingError(msg, loc=inst.loc)

numba.errors.TypingError: Failed at nopython (nopython frontend)

Untyped global name 'carg': cannot determine Numba type of <class 'numba.cuda.dispatcher.CUDAUFuncDispatcher'>

File "test.py", line 11

What am I doing wrong?

[Stanley Seibert]

You can't call a CUDA ufunc from another CUDA ufunc.  Since the body of the cpow() function should be written assuming scalar arguments (Numba takes care of the array broadcasting), what you want is a scalar CUDA device function for carg():

from numba import cuda

@cuda.jit(device=True)

def carg(z):
    return math.atan2(z.imag, z.real)

You don't need to specify data types for the device function since Numba can figure them out for you.

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users+unsubscribe@continuum.io.
To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/6a90d58f-cce8-41a8-9b21-f52a3c507fe7%40continuum.io.
For more options, visit https://groups.google.com/a/continuum.io/d/optout.