Skip to content

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

Description

@RemiLehe

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)

Activity

  1. gmarkall commented on Aug 3, 2022

    @gmarkall
    Member

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

  2. gmarkall commented on Aug 3, 2022

    @gmarkall
    Member

    The culprit is 0401550 from PR #7999.

  3. gmarkall commented on Aug 3, 2022

    @gmarkall
    Member

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

  4. gmarkall commented on Aug 3, 2022

    @gmarkall
    Member

    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)
  5. added this to the Numba 0.56.1 milestone on Aug 3, 2022
  6. gmarkall commented on Aug 3, 2022

    @gmarkall
    Member

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

  7. RemiLehe commented on Aug 3, 2022

    @RemiLehe
    Author

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

  8. gmarkall commented on Aug 4, 2022

    @gmarkall
    Member

    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)

  9. gmarkall commented on Aug 4, 2022

    @gmarkall
    Member

    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
    
  10. gmarkall commented on Aug 4, 2022

    @gmarkall
    Member

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

  11. added a commit that references this issue on Aug 4, 2022
  12. gmarkall commented on Aug 4, 2022

    @gmarkall
    Member

    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

  13. added a commit that references this issue on Aug 17, 2022
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Labels

    CUDACUDA related issue/PRbug - failure to compileBugs: failed to compile valid codebug - regressionA regression against a previous version of Numba

    Type

    No type

    Projects

    No projects

      Milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions