Numba 0.56 does not support atomic.add on arrays of complexs anymore

Reporting a bug

I have tried using the latest released version of Numba (most recent is
visible in the change log (https://github.com/numba/numba/blob/main/CHANGE_LOG).
I have included a self contained code sample to reproduce the problem.
i.e. it's possible to run as 'python bug.py'.

The following code snippet uses an atomic.add on an array of complex numbers.
It runs without any issues with numba 0.55.2, but fails with numba 0.56

from numba import cuda
import numpy as np

@cuda.jit
def atomic_add_one(values):
    """
    Adds 1 to each element values, with an atomic add

    Parameters:
    -----------
    values: 1darray of `complex128`
    """
    tid = cuda.threadIdx.x
    bid = cuda.blockIdx.x
    bdim = cuda.blockDim.x
    i = (bid * bdim) + tid
    cuda.atomic.add(values.real, i, 1)

arr = np.random.rand(16384) + 1.j*np.random.rand(16384)
print(arr.dtype)

atomic_add_one[256,64](arr)

Thanks for the reproducer - I can reproduce this issue, and I'm looking into it.

The culprit is 0401550 from PR #7999.

I have just tested the reproducer with #8271 - I had hoped it would fix this but it seems not to.

The following also does not help, neither with nor without #8271:

diff --git a/numba/np/arrayobj.py b/numba/np/arrayobj.py
index 7be350ba0..83bbc0579 100644
--- a/numba/np/arrayobj.py
+++ b/numba/np/arrayobj.py
@@ -2617,7 +2617,7 @@ def array_flags_f_contiguous(context, builder, typ, value):
 # ------------------------------------------------------------------------------
 # .real / .imag
 
-@overload_attribute(types.Array, 'real')
+@overload_attribute(types.Array, 'real',)
 def ol_array_real(arr):
     typ = arr
     if typ.dtype in types.complex_domain:
@@ -2647,7 +2647,7 @@ def _force_readonly(tyctx, arr):
     return sig, impl
 
 
-@overload_attribute(types.Array, 'imag')
+@overload_attribute(types.Array, 'imag',)
 def ol_array_imag(arr):
     typ = arr
     if typ.dtype in types.complex_domain:
@@ -2740,8 +2740,8 @@ _array_real_attr = _generate_real_imag_attr("real")
 _array_imag_attr = _generate_real_imag_attr("imag")
 
 
-@overload_method(types.Array, 'conj')
-@overload_method(types.Array, 'conjugate')
+@overload_method(types.Array, 'conj',)
+@overload_method(types.Array, 'conjugate',)
 def array_conj(arr):
     def impl(arr):
         return np.conj(arr)

One hypothesis: this is happening because the lowering for atomic addition is using the low-level API, which can't see overloads written with the high-level API (as .real and .imag now are as of the commit above).

Thanks for the information, and thanks a lot for looking into this!

OK, it's not an issue with typing / resolving the function. The overload implementation is a function that returns an array, which is not supported by the CUDA calling convention:

numba.core.errors.NumbaTypeError: Failed in nopython mode pipeline (step: nopython frontend)
Only accept returning of array passed into the function as argument

(from the traceback this produces)

Ah, even if I permit the array return type (by commenting out the check in BaseTypeInference.run_pass.legalize_return_type()) I hit an error resolving the real atrtibute:

Traceback (most recent call last):
  File "/home/gmarkall/numbadev/issues/8309/repro.py", line 22, in <module>
    atomic_add_one[256,64](arr)
  File "/home/gmarkall/numbadev/numba/numba/cuda/dispatcher.py", line 498, in __call__
    return self.dispatcher.call(args, self.griddim, self.blockdim,
  File "/home/gmarkall/numbadev/numba/numba/cuda/dispatcher.py", line 632, in call
    kernel = _dispatcher.Dispatcher._cuda_call(self, *args)
  File "/home/gmarkall/numbadev/numba/numba/cuda/dispatcher.py", line 640, in _compile_for_args
    return self.compile(tuple(argtypes))
  File "/home/gmarkall/numbadev/numba/numba/cuda/dispatcher.py", line 820, in compile
    kernel = _Kernel(self.py_func, argtypes, **self.targetoptions)
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 35, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/gmarkall/numbadev/numba/numba/cuda/dispatcher.py", line 75, in __init__
    cres = compile_cuda(self.py_func, types.void, self.argtypes,
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 35, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 210, in compile_cuda
    cres = compiler.compile_extra(typingctx=typingctx,
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 716, in compile_extra
    return pipeline.compile_extra(func)
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 452, in compile_extra
    return self._compile_bytecode()
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 520, in _compile_bytecode
    return self._compile_core()
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 495, in _compile_core
    raise e
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 486, in _compile_core
    pm.run(self.state)
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 364, in run
    raise e
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 356, in run
    self._runPass(idx, pass_inst, state)
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 35, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 311, in _runPass
    mutated |= check(pss.run_pass, internal_state)
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 273, in check
    mangled = func(compiler_state)
  File "/home/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 394, in run_pass
    lower.lower()
  File "/home/gmarkall/numbadev/numba/numba/core/lowering.py", line 168, in lower
    self.lower_normal_function(self.fndesc)
  File "/home/gmarkall/numbadev/numba/numba/core/lowering.py", line 222, in lower_normal_function
    entry_block_tail = self.lower_function_body()
  File "/home/gmarkall/numbadev/numba/numba/core/lowering.py", line 251, in lower_function_body
    self.lower_block(block)
  File "/home/gmarkall/numbadev/numba/numba/core/lowering.py", line 265, in lower_block
    self.lower_inst(inst)
  File "/home/gmarkall/numbadev/numba/numba/core/lowering.py", line 439, in lower_inst
    val = self.lower_assign(ty, inst)
  File "/home/gmarkall/numbadev/numba/numba/core/lowering.py", line 626, in lower_assign
    return self.lower_expr(ty, value)
  File "/home/gmarkall/numbadev/numba/numba/core/lowering.py", line 1260, in lower_expr
    res = impl(self.context, self.builder, ty, val, expr.attr)
  File "/home/gmarkall/numbadev/numba/numba/core/imputils.py", line 152, in res
    return real_impl(context, builder, typ, value, attr)
  File "/home/gmarkall/numbadev/numba/numba/np/arrayobj.py", line 2777, in array_record_getattr
    raise NotImplementedError("attribute %r of %s not defined"
NotImplementedError: attribute 'real' of array(complex128, 1d, C) not defined

Progress towards resolving this issue can be observed in PR #8310.

Add test based on Issue numba#8309 reproducer

One potential fix is now ready for consideration in #8310. Whether this is the right fix and the general path forward is a complex decision, so I've added it to the agenda for next week's dev meeting along with the considerations I can see as of now: https://hackmd.io/P1y0q3BcStKfIZky7CLHHg?view

Merge pull request #8310 from gmarkall/issue-8309

Navigation Menu

Repository navigation

Numba 0.56 does not support `atomic.add` on arrays of complexs anymore #8309

Reporting a bug

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Numba 0.56 does not support atomic.add on arrays of complexs anymore #8309

Description

Reporting a bug

Activity

gmarkall commented on Aug 3, 2022

gmarkall commented on Aug 3, 2022

gmarkall commented on Aug 3, 2022

gmarkall commented on Aug 3, 2022

gmarkall commented on Aug 3, 2022

RemiLehe commented on Aug 3, 2022

gmarkall commented on Aug 4, 2022

gmarkall commented on Aug 4, 2022

gmarkall commented on Aug 4, 2022

gmarkall commented on Aug 4, 2022

Metadata

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Issue actions

Numba 0.56 does not support `atomic.add` on arrays of complexs anymore #8309