Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

CUDA typing / lowering can't find contains_impl for char sequences #7059

Open
gmarkall opened this issue May 26, 2021 · 3 comments
Open

CUDA typing / lowering can't find contains_impl for char sequences #7059

gmarkall opened this issue May 26, 2021 · 3 comments
Labels
bug - failure to compile Bugs: failed to compile valid code CUDA CUDA related issue/PR
Projects
CUDA 0.57 overflow

Comments

@gmarkall
Copy link
Member

There are two reproducers here, which I include both of because I suspect the issues are related / the same.

With master + PR #7057 + PR #6917, the following:

from numba import cuda
import numpy as np

@cuda.jit
def k(result):
    result[0] = '123' in '12345'

result = np.zeros(1, dtype=np.bool_)
k[1, 1](result)

Gives the following traceback:

Traceback (most recent call last):
  File "/home/gmarkall/numbadev/numba/numba/core/errors.py", line 744, in new_error_context
    yield
  File "/home/gmarkall/numbadev/numba/numba/core/lowering.py", line 230, in lower_block
    self.lower_inst(inst)
  File "/home/gmarkall/numbadev/numba/numba/core/lowering.py", line 323, in lower_inst
    val = self.lower_assign(ty, inst)
  File "/home/gmarkall/numbadev/numba/numba/core/lowering.py", line 499, in lower_assign
    return self.lower_expr(ty, value)
  File "/home/gmarkall/numbadev/numba/numba/core/lowering.py", line 1003, in lower_expr
    return self.lower_binop(resty, expr, expr.fn)
  File "/home/gmarkall/numbadev/numba/numba/core/lowering.py", line 625, in lower_binop
    impl = self.context.get_function(op, sig)
  File "/home/gmarkall/numbadev/numba/numba/core/base.py", line 568, in get_function
    return self.get_function(fn, sig, _firstcall=False)
  File "/home/gmarkall/numbadev/numba/numba/core/base.py", line 570, in get_function
    raise NotImplementedError("No definition for lowering %s%s" % (key, sig))
NotImplementedError: No definition for lowering <built-in method contains_impl of _dynfunc._Closure object at 0x7f35e557bd00>(unicode_type, unicode_type) -> bool

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "/home/gmarkall/numbadev/issues/cuda-unicode-string-typing/repro_small.py", line 11, in <module>
    k[1, 1](result)
  File "/home/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 861, in __call__
    return self.dispatcher.call(args, self.griddim, self.blockdim,
  File "/home/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 994, in call
    kernel = _dispatcher.Dispatcher._cuda_call(self, *args)
  File "/home/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 1002, in _compile_for_args
    return self.compile(tuple(argtypes))
  File "/home/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 1099, in compile
    kernel = _Kernel(self.py_func, argtypes, link=self.link,
  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 538, 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 163, in compile_cuda
    cres = compiler.compile_extra(typingctx=typingctx,
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 675, in compile_extra
    return pipeline.compile_extra(func)
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 419, in compile_extra
    return self._compile_bytecode()
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 483, in _compile_bytecode
    return self._compile_core()
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 462, in _compile_core
    raise e
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 453, in _compile_core
    pm.run(self.state)
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 343, in run
    raise patched_exception
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 334, 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 289, in _runPass
    mutated |= check(pss.run_pass, internal_state)
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 262, in check
    mangled = func(compiler_state)
  File "/home/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 396, in run_pass
    lower.lower()
  File "/home/gmarkall/numbadev/numba/numba/core/lowering.py", line 136, in lower
    self.lower_normal_function(self.fndesc)
  File "/home/gmarkall/numbadev/numba/numba/core/lowering.py", line 190, in lower_normal_function
    entry_block_tail = self.lower_function_body()
  File "/home/gmarkall/numbadev/numba/numba/core/lowering.py", line 216, in lower_function_body
    self.lower_block(block)
  File "/home/gmarkall/numbadev/numba/numba/core/lowering.py", line 230, in lower_block
    self.lower_inst(inst)
  File "/home/gmarkall/miniconda3/envs/numbanp120/lib/python3.9/contextlib.py", line 135, in __exit__
    self.gen.throw(type, value, traceback)
  File "/home/gmarkall/numbadev/numba/numba/core/errors.py", line 754, in new_error_context
    raise newerr.with_traceback(tb)
numba.core.errors.LoweringError: Failed in cuda mode pipeline (step: native lowering)
No definition for lowering <built-in method contains_impl of _dynfunc._Closure object at 0x7f35e557bd00>(unicode_type, unicode_type) -> bool

File "repro_small.py", line 7:
def k(result):
    result[0] = '123' in '12345'
    ^

During: lowering "$6contains_op.2 = $const2.0 in $const4.1" at /home/gmarkall/numbadev/issues/cuda-unicode-string-typing/repro_small.py (7)

I believe the contains_impl method it can't lower is from:

numba/numba/cpython/charseq.py

Lines 512 to 529 in ae4a0f1

@overload(operator.contains)
def charseq_contains(a, b):
if not _same_kind(a, b):
return
left_code = _get_code_impl(a)
right_code = _get_code_impl(b)
if left_code is not None and right_code is not None:
if _is_bytes(a):
def contains_impl(a, b):
# Ideally, `return bytes(b) in bytes(a)` would be used
# here, but numba Bytes does not implement
# contains. So, using `unicode_type` implementation
# here:
return b._to_str() in a._to_str()
else:
def contains_impl(a, b):
return str(b) in str(a)
return contains_impl

If I pass in an array containing the string:

from numba import cuda
import numpy as np


@cuda.jit
def k(result, p2):
    result[0] = '123' in p2[()]


y = np.array('12345')
result = np.zeros(1, dtype=np.bool_)

k[1, 1](result, y)

then typing seems to go extremely slowly and eventually fail. If I apply the following:

diff --git a/numba/core/typing/context.py b/numba/core/typing/context.py
index c46360eb9..44ae4e721 100644
--- a/numba/core/typing/context.py
+++ b/numba/core/typing/context.py
@@ -244,6 +244,7 @@ class BaseContext(object):
 
         if isinstance(func, types.Callable):
             # XXX fold this into the __call__ attribute logic?
+            print(f"Resolving call type, {func} == {args} -- {kws}")
             return func.get_call_type(self, args, kws)

a large amount of output is produced, with a portion of it being:

Resolving call type, Function(<built-in function eq>) == (int64, Literal[int](0)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, AssertionError == (Literal[str](PY_UNICODE_WCHAR_KIND unsupported),) -- {}
Resolving call type, AssertionError == (Literal[str](Unexpected unicode encoding encountered),) -- {}
Resolving call type, Function(<intrinsic _malloc_string>) == (int64, int64, int64, bool) -- {}
Resolving call type, Function(uint32) == (Literal[int](0),) -- {}
Resolving call type, Function(<function _set_code_point at 0x7f9bcc0cc790>) == (unicode_type, int64, uint32) -- {}
Resolving call type, Function(<built-in function eq>) == (int32, Literal[int](1)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<intrinsic set_uint8>) == (void*, int64, uint32) -- {}
Resolving call type, Function(<built-in function eq>) == (int32, Literal[int](2)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<intrinsic set_uint16>) == (void*, int64, uint32) -- {}
Resolving call type, Function(<built-in function eq>) == (int32, Literal[int](4)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<intrinsic set_uint32>) == (void*, int64, uint32) -- {}
Resolving call type, AssertionError == (Literal[str](Unexpected unicode representation in _set_code_point),) -- {}
Resolving call type, Function(<built-in function eq>) == (int32, Literal[int](1)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<intrinsic set_uint8>) == (void*, int64, uint32) -- {}
Resolving call type, Function(<built-in function eq>) == (int32, Literal[int](2)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<intrinsic set_uint16>) == (void*, int64, uint32) -- {}
Resolving call type, Function(<built-in function eq>) == (int32, Literal[int](4)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<intrinsic set_uint32>) == (void*, int64, uint32) -- {}
Resolving call type, AssertionError == (Literal[str](Unexpected unicode representation in _set_code_point),) -- {}
Resolving call type, Function(<function _kind_to_byte_width at 0x7f9bcc0cc940>) == (int64,) -- {}
Resolving call type, Function(<intrinsic _malloc_string>) == (int64, int64, int64, bool) -- {}
Resolving call type, Function(uint32) == (Literal[int](0),) -- {}
Resolving call type, Function(<function _set_code_point at 0x7f9bcc0cc790>) == (unicode_type, int64, uint32) -- {}
Resolving call type, Function(<class 'range'>) == (int64,) -- {}
Resolving call type, Function(<function unicode_charseq_get_code at 0x7f9bcc01eb80>) == ([unichr x 5], int64) -- {}
Resolving call type, Function(<class 'str'>) == (unicode_type,) -- {}
Resolving call type, Function(<class 'str'>) == ([unichr x 5],) -- {}
Resolving call type, Function(<built-in function len>) == ([unichr x 5],) -- {}
Resolving call type, Function(<built-in function eq>) == (Literal[int](0), Literal[int](0)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<built-in function sub>) == (Literal[int](5), Literal[int](1)) -- {}
Resolving call type, Function(<built-in function lt>) == (int64, Literal[int](0)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<function unicode_charseq_get_code at 0x7f9bcc01eb80>) == ([unichr x 5], int64) -- {}
Resolving call type, Function(<class 'bool'>) == (Literal[int](1),) -- {}
Resolving call type, Function(<intrinsic deref_uint32>) == ([unichr x 5], int64) -- {}
Resolving call type, Function(<class 'bool'>) == (Literal[int](1),) -- {}
Resolving call type, Function(<intrinsic deref_uint32>) == ([unichr x 5], int64) -- {}
Resolving call type, Function(<built-in function add>) == (int64, Literal[int](1)) -- {}
Resolving call type, Function(<built-in function eq>) == (int64, Literal[int](0)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<built-in function sub>) == (int64, Literal[int](1)) -- {}
Resolving call type, Function(<built-in function lt>) == (int64, Literal[int](0)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<function unicode_charseq_get_code at 0x7f9bcc01eb80>) == ([unichr x 5], int64) -- {}
Resolving call type, Function(<built-in function add>) == (int64, Literal[int](1)) -- {}
Resolving call type, Function(<built-in function eq>) == (int64, Literal[int](0)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<built-in function sub>) == (int64, Literal[int](1)) -- {}
Resolving call type, Function(<built-in function lt>) == (int64, Literal[int](0)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<function unicode_charseq_get_code at 0x7f9bcc01eb80>) == ([unichr x 5], int64) -- {}
Resolving call type, Function(<built-in function add>) == (int64, Literal[int](1)) -- {}
Resolving call type, BoundFunction((<class 'numba.core.types.npytypes.UnicodeCharSeq'>, '_get_kind') for [unichr x 5]) == () -- {}
Resolving call type, Function(<built-in function len>) == ([unichr x 5],) -- {}
Resolving call type, Function(<built-in function eq>) == (Literal[int](0), Literal[int](0)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<built-in function sub>) == (Literal[int](5), Literal[int](1)) -- {}
Resolving call type, Function(<built-in function lt>) == (int64, Literal[int](0)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<function unicode_charseq_get_code at 0x7f9bcc01eb80>) == ([unichr x 5], int64) -- {}
Resolving call type, Function(<class 'bool'>) == (Literal[int](1),) -- {}
Resolving call type, Function(<intrinsic deref_uint32>) == ([unichr x 5], int64) -- {}
Resolving call type, Function(<class 'bool'>) == (Literal[int](1),) -- {}
Resolving call type, Function(<intrinsic deref_uint32>) == ([unichr x 5], int64) -- {}
Resolving call type, Function(<built-in function add>) == (int64, Literal[int](1)) -- {}
Resolving call type, Function(<built-in function eq>) == (int64, Literal[int](0)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<built-in function sub>) == (int64, Literal[int](1)) -- {}
Resolving call type, Function(<built-in function lt>) == (int64, Literal[int](0)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<function unicode_charseq_get_code at 0x7f9bcc01eb80>) == ([unichr x 5], int64) -- {}
Resolving call type, Function(<built-in function add>) == (int64, Literal[int](1)) -- {}
Resolving call type, Function(<built-in function eq>) == (int64, Literal[int](0)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<built-in function sub>) == (int64, Literal[int](1)) -- {}
Resolving call type, Function(<built-in function lt>) == (int64, Literal[int](0)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<function unicode_charseq_get_code at 0x7f9bcc01eb80>) == ([unichr x 5], int64) -- {}
Resolving call type, Function(<built-in function add>) == (int64, Literal[int](1)) -- {}
Resolving call type, Function(<class 'range'>) == (int64,) -- {}
Resolving call type, Function(<function unicode_charseq_get_code at 0x7f9bcc01eb80>) == ([unichr x 5], int64) -- {}
Resolving call type, Function(<class 'bool'>) == (Literal[int](1),) -- {}
Resolving call type, Function(<intrinsic deref_uint32>) == ([unichr x 5], int64) -- {}
Resolving call type, Function(<class 'bool'>) == (Literal[int](1),) -- {}
Resolving call type, Function(<intrinsic deref_uint32>) == ([unichr x 5], int64) -- {}
Resolving call type, Function(<built-in function gt>) == (uint32, Literal[int](0)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<built-in function gt>) == (Literal[int](0), Literal[int](65535)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<built-in function gt>) == (Literal[int](0), Literal[int](255)) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<built-in function len>) == ([unichr x 5],) -- {}
Resolving call type, Function(<class 'range'>) == (int64,) -- {}
Resolving call type, Function(<function unicode_charseq_get_code at 0x7f9bcc01eb80>) == ([unichr x 5], int64) -- {}
Resolving call type, Function(<built-in function gt>) == (uint32, int64) -- {}
Resolving call type, Function(<class 'bool'>) == (bool,) -- {}
Resolving call type, Function(<built-in function gt>) == (int64, Literal[int](65535)) -- {}

(this is just at a random point where I hit ctrl+C and copied the last output).

Eventually the typing fails with:

Traceback (most recent call last):
  File "/home/gmarkall/numbadev/issues/cuda-unicode-string-typing/repro.py", line 13, in <module>
    k[1, 1](result, y)
  File "/home/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 861, in __call__
    return self.dispatcher.call(args, self.griddim, self.blockdim,
  File "/home/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 994, in call
    kernel = _dispatcher.Dispatcher._cuda_call(self, *args)
  File "/home/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 1002, in _compile_for_args
    return self.compile(tuple(argtypes))
  File "/home/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 1099, in compile
    kernel = _Kernel(self.py_func, argtypes, link=self.link,
  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 538, 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 163, in compile_cuda
    cres = compiler.compile_extra(typingctx=typingctx,
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 675, in compile_extra
    return pipeline.compile_extra(func)
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 419, in compile_extra
    return self._compile_bytecode()
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 483, in _compile_bytecode
    return self._compile_core()
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 462, in _compile_core
    raise e
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 453, in _compile_core
    pm.run(self.state)
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 343, in run
    raise patched_exception
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 334, 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 289, in _runPass
    mutated |= check(pss.run_pass, internal_state)
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 262, in check
    mangled = func(compiler_state)
  File "/home/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 105, in run_pass
    typemap, return_type, calltypes, errs = type_inference_stage(
  File "/home/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 83, in type_inference_stage
    errs = infer.propagate(raise_errors=raise_errors)
  File "/home/gmarkall/numbadev/numba/numba/core/typeinfer.py", line 1071, in propagate
    raise errors[0]
numba.core.errors.TypingError: Failed in cuda mode pipeline (step: nopython frontend)
No implementation of function Function(<built-in function contains>) found for signature:
 
 >>> contains([unichr x 5], Literal[str](123))
 
There are 14 candidate implementations:
  - Of which 13 did not match due to:
  Overload of function 'contains': File: <numerous>: Line N/A.
    With argument(s): '([unichr x 5], unicode_type)':
   No match.
  - Of which 1 did not match due to:
  Overload in function 'charseq_contains': File: numba/cpython/charseq.py: Line 512.
    With argument(s): '([unichr x 5], unicode_type)':
   Rejected as the implementation raised a specific error:
     TypingError: Failed in nopython mode pipeline (step: nopython frontend)
   No implementation of function Function(<class 'str'>) found for signature:
    
    >>> str([unichr x 5])
    
   There are 6 candidate implementations:
       - Of which 4 did not match due to:
       Overload of function 'str': File: <numerous>: Line N/A.
         With argument(s): '([unichr x 5])':
        No match.
       - Of which 2 did not match due to:
       Overload in function 'charseq_str': File: numba/cpython/charseq.py: Line 600.
         With argument(s): '([unichr x 5])':
        Rejected as the implementation raised a specific error:
          TypingError: Failed in nopython mode pipeline (step: nopython frontend)
        No implementation of function Function(<function _empty_string at 0x7f700c75c280>) found for signature:
         
         >>> _empty_string(int64, int64, bool)
         
        There are 2 candidate implementations:
              - Of which 2 did not match due to:
              Overload in function 'register_jitable.<locals>.wrap.<locals>.ov_wrap': File: numba/core/extending.py: Line 159.
                With argument(s): '(int64, int64, bool)':
               Rejected as the implementation raised a specific error:
                 LoweringError: Failed in nopython mode pipeline (step: native lowering)
               NRT required but not enabled
               
               File "../../numba/numba/cpython/unicode.py", line 277:
               def _empty_string(kind, length, is_ascii=0):
                   <source elided>
                   char_width = _kind_to_byte_width(kind)
                   s = _malloc_string(kind, char_width, length, is_ascii)
                   ^
               
               During: lowering "s = call $10load_global.3(kind, char_width, length, is_ascii, func=$10load_global.3, args=[Var(kind, unicode.py:276), Var(char_width, unicode.py:276), Var(length, unicode.py:276), Var(is_ascii, unicode.py:276)], kws=(), vararg=None, target=None)" at /home/gmarkall/numbadev/numba/numba/cpython/unicode.py (277)
          raised from /home/gmarkall/numbadev/numba/numba/core/errors.py:754
        
        During: resolving callee type: Function(<function _empty_string at 0x7f700c75c280>)
        During: typing of call at /home/gmarkall/numbadev/numba/numba/cpython/charseq.py (609)
        
        
        File "../../numba/numba/cpython/charseq.py", line 609:
                def str_impl(s):
                    <source elided>
                    is_ascii = kind == 1 and s.isascii()
                    result = unicode._empty_string(kind, n, is_ascii)
                    ^
   
     raised from /home/gmarkall/numbadev/numba/numba/core/typeinfer.py:1071
   
   During: resolving callee type: Function(<class 'str'>)
   During: typing of call at /home/gmarkall/numbadev/numba/numba/cpython/charseq.py (528)
   
   
   File "../../numba/numba/cpython/charseq.py", line 528:
               def contains_impl(a, b):
                   return str(b) in str(a)
                   ^

  raised from /home/gmarkall/numbadev/numba/numba/core/typeinfer.py:1071

During: typing of intrinsic-call at /home/gmarkall/numbadev/issues/cuda-unicode-string-typing/repro.py (7)

File "repro.py", line 7:
def k(result, p2):
    result[0] = '123' in p2[()]
    ^

Related: #6897 - I think that might be the same issue, but the traceback / error has changed a little since that was reported.
@gmarkall gmarkall added bug - failure to compile Bugs: failed to compile valid code CUDA CUDA related issue/PR labels May 26, 2021
@gmarkall gmarkall added this to the Numba 0.55 RC milestone May 26, 2021
@gmarkall
Copy link
Member Author

gmarkall commented Jun 1, 2021

Character sequence implementations call unicode implementations, forcing allocations, which require NRT not present in CUDA. E.g. from above:

Resolving call type, Function(<intrinsic _malloc_string>) == (int64, int64, int64, bool) -- {}

@stuartarchibald
Copy link
Contributor

Is this fixed by #7057 ?

@gmarkall
Copy link
Member Author

gmarkall commented Nov 18, 2021 •
edited

Is this fixed by #7057 ?

I don't think so - #7057 is needed to get over issue #7041 in the reproducer so it can get to the point it currently does.

@esc esc modified the milestones: Numba 0.55 RC, Numba 0.55.1 Jan 15, 2022
@gmarkall gmarkall modified the milestones: Numba 0.55.1, Numba 0.56 RC Jan 17, 2022
@sklam sklam modified the milestones: Numba 0.56 RC, Numba 0.57 RC Jun 1, 2022
@gmarkall gmarkall added this to Imported in CUDA 0.57 overflow May 2, 2023
@gmarkall gmarkall removed this from the Numba 0.57 RC milestone May 2, 2023
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees
No one assigned
Labels
bug - failure to compile Bugs: failed to compile valid code CUDA CUDA related issue/PR
Projects
CUDA 0.57 overflow
Imported
Milestone
No milestone
Development

No branches or pull requests
4 participants
@esc @gmarkall @stuartarchibald @sklam