Im running a jupyter notebook frome here : https://github.com/noahgift/nuclear_powered_command_line_tools/blob/master/notebooks/numba-cuda.ipynb

The docs of current numba/cuda is here : https://numba.readthedocs.io/en/stable/cuda-reference/kernel.html#kernel-declaration

Im running this line in numba.cuda in python :

from numba import cuda
from numba import *

I have a mandel function as follow:

@jit
def mandel(x, y, max_iters):"""Given the real and imaginary parts of a complex number,determine if it is a candidate for membership in the Mandelbrotset given a fixed number of iterations."""c = complex(x, y)z = 0.0jfor i in range(max_iters):z = z*z + cif (z.real*z.real + z.imag*z.imag) >= 4:return ireturn max_iters

I run cuda with

mandel_gpu = cuda.jit(func_or_sig=mandel,device=True)
# previously  mandel_gpu = cuda.jit(restype=uint32, argtypes=[f8, f8, uint32], device=True)(mandel)

with mandel defined as

  @cuda.jit#(argtypes=[f8, f8, f8, f8, uint8[:,:], uint32])
def mandel_kernel(min_x, max_x, min_y, max_y, image, iters):height = image.shape[0]width = image.shape[1]pixel_size_x = (max_x - min_x) / widthpixel_size_y = (max_y - min_y) / heightstartX, startY = cuda.grid(2)gridX = cuda.gridDim.x * cuda.blockDim.x;gridY = cuda.gridDim.y * cuda.blockDim.y;for x in range(startX, width, gridX):real = min_x + x * pixel_size_xfor y in range(startY, height, gridY):imag = min_y + y * pixel_size_y image[y, x] = mandel_gpu(real, imag, iters)

I get then the following call :

gimage = np.zeros((1024, 1536), dtype = np.uint8)
blockdim = (32, 8)
griddim = (32,16)start = timer()
d_image = cuda.to_device(gimage)
mandel_kernel[griddim, blockdim](-2.0, 1.0, -1.0, 1.0, d_image, 20) 
d_image.to_host()
dt = timer() - startprint("Mandelbrot created on GPU in %f s" % dt)imshow(gimage)
show()

I get then the error :

    ---------------------------------------------------------------------------NotImplementedError                       Traceback (most recent call last)/usr/local/lib/python3.7/dist-packages/numba/core/errors.py in new_error_context(fmt_, *args, **kwargs)822     try:
--> 823         yield824     except NumbaError as e:33 framesNotImplementedError: No definition for lowering <built-in method mandel of _dynfunc._Closure object at 0x7efde36f8590>(float64, float64, int64) -> int64During handling of the above exception, another exception occurred:LoweringError                             Traceback (most recent call last)/usr/local/lib/python3.7/dist-packages/numba/core/errors.py in new_error_context(fmt_, *args, **kwargs)835             else:836                 tb = None
--> 837             raise newerr.with_traceback(tb)838         elif use_new_style_errors():839             raise eLoweringError: Failed in cuda mode pipeline (step: native lowering)
No definition for lowering <built-in method mandel of _dynfunc._Closure object at 0x7efde36f8590>(float64, float64, int64) -> int64File "<ipython-input-14-82f426a4058e>", line 17:
def mandel_kernel(min_x, max_x, min_y, max_y, image, iters):<source elided>imag = min_y + y * pixel_size_y image[y, x] = mandel_gpu(real, imag, iters)^During: lowering "$160call_function.12 = call $152load_global.8(real, imag, iters, func=$152load_global.8, args=[Var(real, <ipython-input-14-82f426a4058e>:14), Var(imag, <ipython-input-14-82f426a4058e>:16), Var(iters, <ipython-input-14-82f426a4058e>:3)], kws=(), vararg=None, varkwarg=None, target=None)" at <ipython-input-14-82f426a4058e> (17)

I have tried to update the code with the comments stating the updated api doesnt take argtypes and restypes but i still get an error and i have no clues how to correct it... Thank you Thanks for your help

I would recommend taking a few minutes to review how python decorators work, because this is at the core of your problem.

When you do this:

@jit
def mandel(x, y, max_iters):...

what is happening is this:

def mandel(x, y, max_iters):...mandel = jit(mandel)

i.e. your original function is replaced by a Numba processed, jit compiled version of your original Python code. Plainly the Numba GPU compiler can't compile a previously Numba compiled function, thus the error when you try and compile a device function from that already compiled code.

If you did something like this:

def mandel(x, y, max_iters):...mandel_gpu = cuda.jit(mandel, device=True)
mandel = jit(mandel)

I suspect the compilation would work correctly (although I haven't tried it, that is up to you).