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.

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

Removes context.compile_internal where easy #8493

Open
wants to merge 15 commits into
base: main
Choose a base branch
from
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Jump to
Jump to file
Failed to load files.
Diff view
Diff view
37 changes: 0 additions & 37 deletions numba/core/typing/arraydecl.py
Original file line number Diff line number Diff line change
Expand Up @@ -414,29 +414,6 @@ def sentry_shape_scalar(ty):
retty = ary.copy(ndim=len(args))
return signature(retty, *args)

@bound_function("array.sort")
def resolve_sort(self, ary, args, kws):
assert not args
assert not kws
return signature(types.none)

@bound_function("array.argsort")
def resolve_argsort(self, ary, args, kws):
assert not args
kwargs = dict(kws)
kind = kwargs.pop('kind', types.StringLiteral('quicksort'))
if not isinstance(kind, types.StringLiteral):
raise TypingError('"kind" must be a string literal')
if kwargs:
msg = "Unsupported keywords: {!r}"
raise TypingError(msg.format([k for k in kwargs.keys()]))
if ary.ndim == 1:
def argsort_stub(kind='quicksort'):
pass
pysig = utils.pysignature(argsort_stub)
sig = signature(types.Array(types.intp, 1, 'C'), kind).replace(pysig=pysig)
return sig

@bound_function("array.view")
def resolve_view(self, ary, args, kws):
from .npydecl import parse_dtype
Expand Down Expand Up @@ -469,20 +446,6 @@ def resolve_astype(self, ary, args, kws):
retty = ary.copy(dtype=dtype, layout=layout, readonly=False)
return signature(retty, *args)

@bound_function("array.ravel")
def resolve_ravel(self, ary, args, kws):
# Only support no argument version (default order='C')
assert not kws
assert not args
return signature(ary.copy(ndim=1, layout='C'))

@bound_function("array.flatten")
def resolve_flatten(self, ary, args, kws):
# Only support no argument version (default order='C')
assert not kws
assert not args
return signature(ary.copy(ndim=1, layout='C'))

def generic_resolve(self, ary, attr):
# Resolution of other attributes, for record arrays
if isinstance(ary.dtype, types.Record):
Expand Down
21 changes: 0 additions & 21 deletions numba/core/typing/builtins.py
Original file line number Diff line number Diff line change
Expand Up @@ -654,17 +654,6 @@ def generic(self, args, kws):
sig = signature(ret, *args)
return sig

# Generic implementation for "not in"

@infer
class GenericNotIn(AbstractTemplate):
key = "not in"

def generic(self, args, kws):
args = args[::-1]
sig = self.context.resolve_function_type(operator.contains, args, kws)
return signature(sig.return_type, *sig.args[::-1])


#-------------------------------------------------------------------------------

Expand Down Expand Up @@ -920,16 +909,6 @@ class Min(MinMaxBase):
pass


@infer_global(round)
class Round(ConcreteTemplate):
cases = [
signature(types.intp, types.float32),
signature(types.int64, types.float64),
signature(types.float32, types.float32, types.intp),
signature(types.float64, types.float64, types.intp),
]


#------------------------------------------------------------------------------


Expand Down
6 changes: 1 addition & 5 deletions numba/core/typing/npydecl.py
Original file line number Diff line number Diff line change
Expand Up @@ -389,10 +389,6 @@ def sum_stub(arr, dtype):
def sum_stub(arr, axis, dtype):
pass
pysig = utils.pysignature(sum_stub)
elif self.method_name == 'argsort':
def argsort_stub(arr, kind='quicksort'):
pass
pysig = utils.pysignature(argsort_stub)
else:
fmt = "numba doesn't support kwarg for {}"
raise TypingError(fmt.format(self.method_name))
Expand All @@ -414,7 +410,7 @@ def _numpy_redirect(fname):
infer_global(numpy_function, types.Function(cls))


for func in ['sum', 'argsort', 'nonzero', 'ravel']:
for func in ['sum', 'nonzero']:
_numpy_redirect(func)


Expand Down
59 changes: 31 additions & 28 deletions numba/cpython/builtins.py
Original file line number Diff line number Diff line change
Expand Up @@ -238,27 +238,41 @@ def _round_intrinsic(tp):
# round() rounds half to even
return "llvm.rint.f%d" % (tp.bitwidth,)

@lower_builtin(round, types.Float)
def round_impl_unary(context, builder, sig, args):
fltty = sig.args[0]
llty = context.get_value_type(fltty)
module = builder.module
fnty = ir.FunctionType(llty, [llty])
fn = cgutils.get_or_insert_function(module, fnty, _round_intrinsic(fltty))
res = builder.call(fn, args)
# unary round() returns an int
res = builder.fptosi(res, context.get_value_type(sig.return_type))
return impl_ret_untracked(context, builder, sig.return_type, res)

@lower_builtin(round, types.Float, types.Integer)
def round_impl_binary(context, builder, sig, args):
fltty = sig.args[0]
@intrinsic
def round_impl_unary(typingctx, x):
sig = types.intp(x)

def codegen(context, builder, sig, args):
fltty = sig.args[0]
llty = context.get_value_type(fltty)
module = builder.module
fnty = ir.FunctionType(llty, [llty])
fn = cgutils.get_or_insert_function(module, fnty,
_round_intrinsic(fltty))
res = builder.call(fn, args)
# unary round() returns an int
res = builder.fptosi(res, context.get_value_type(sig.return_type))
return impl_ret_untracked(context, builder, sig.return_type, res)
return sig, codegen


@overload(round)
def round_impl_binary(x, ndigits=None):
if isinstance(x, types.Float) and cgutils.is_nonelike(ndigits):
def impl(x, ndigits=None):
return round_impl_unary(x)
return impl

if not (isinstance(x, types.Float) and isinstance(ndigits, types.Integer)):
return

# Allow calling the intrinsic from the Python implementation below.
# This avoids the conversion to an int in Python 3's unary round().
_round = types.ExternalFunction(
_round_intrinsic(fltty), typing.signature(fltty, fltty))
_round_intrinsic(x), typing.signature(x, x))

def round_ndigits(x, ndigits):
def round_ndigits(x, ndigits=None):
if math.isinf(x) or math.isnan(x):
return x

Expand All @@ -281,8 +295,7 @@ def round_ndigits(x, ndigits):
y = x / pow1
return _round(y) * pow1

res = context.compile_internal(builder, round_ndigits, sig, args)
return impl_ret_untracked(context, builder, sig.return_type, res)
return round_ndigits


#-------------------------------------------------------------------------------
Expand Down Expand Up @@ -401,16 +414,6 @@ def next_impl(context, builder, sig, args):

# -----------------------------------------------------------------------------

@lower_builtin("not in", types.Any, types.Any)
def not_in(context, builder, sig, args):
def in_impl(a, b):
return operator.contains(b, a)

res = context.compile_internal(builder, in_impl, sig, args)
return builder.not_(res)


# -----------------------------------------------------------------------------

@lower_builtin(len, types.ConstSized)
def constsized_len(context, builder, sig, args):
Expand Down
49 changes: 33 additions & 16 deletions numba/cuda/cudaimpl.py
Original file line number Diff line number Diff line change
Expand Up @@ -13,6 +13,7 @@
from numba import cuda
from numba.cuda import nvvmutils, stubs, errors
from numba.cuda.types import dim3, grid_group, CUDADispatcher
from numba.cuda.extending import overload, intrinsic


registry = Registry()
Expand Down Expand Up @@ -627,27 +628,43 @@ def ptx_min_f8(context, builder, sig, args):
])


@lower(round, types.f4)
@lower(round, types.f8)
def ptx_round(context, builder, sig, args):
fn = cgutils.get_or_insert_function(
builder.module,
ir.FunctionType(
ir.IntType(64),
(ir.DoubleType(),)),
'__nv_llrint')
return builder.call(fn, [
context.cast(builder, args[0], sig.args[0], types.double),
])
# @lower(round, types.f4)
# @lower(round, types.f8)
@intrinsic
def ptx_round(typingctx, x):
sig = types.int64(x)

def codegen(context, builder, sig, args):
fn = cgutils.get_or_insert_function(
builder.module,
ir.FunctionType(
ir.IntType(64),
(ir.DoubleType(),)),
'__nv_llrint')
return builder.call(fn, [
context.cast(builder, args[0], sig.args[0], types.double),
])

return sig, codegen


# This rounding implementation follows the algorithm used in the "fallback
# version" of double_round in CPython.
# https://github.com/python/cpython/blob/a755410e054e1e2390de5830befc08fe80706c66/Objects/floatobject.c#L964-L1007
@overload(round)
def round_to_impl(x):
if isinstance(x, types.Float):
def impl(x):
return ptx_round(x)
return impl


@overload(round)
def round_to_impl2(x, ndigits):
# XXX: it doesn't look like the CUDA target supports ndigits=None
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can it support ndigits=None with something like:

    if isinstance(ndigits, NoneType):
        def impl(x):
            return ptx_round(x)
        return impl

?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

No, it crashes with the following error:

...
numba.core.errors.LoweringError: Failed in cuda mode pipeline (step: native lowering)
No definition for lowering <built-in method impl of _dynfunc._Closure object at 0x7fe1ee519ee0>(float32, omitted(default=None)) -> int64

File "numba/cuda/tests/cudapy/test_intrinsics.py", line 207:
def simple_round(ary, c):
    ary[0] = round(c)

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think there's at least two problems going on here, one of which is #8528, the other is what is causing this lowering error. It seems like non-matching default kwargs might be part of the cause, it's also strange that the overload validation step doesn't catch this.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is also involved: #8530

if not isinstance(x, types.Float) and isinstance(ndigits, types.Integer):
return

@lower(round, types.f4, types.Integer)
@lower(round, types.f8, types.Integer)
def round_to_impl(context, builder, sig, args):
def round_ndigits(x, ndigits):
if math.isinf(x) or math.isnan(x):
return x
Expand Down Expand Up @@ -681,7 +698,7 @@ def round_ndigits(x, ndigits):

return z

return context.compile_internal(builder, round_ndigits, sig, args, )
return round_ndigits


def gen_deg_rad(const):
Expand Down
7 changes: 7 additions & 0 deletions numba/cuda/extending.py
Original file line number Diff line number Diff line change
Expand Up @@ -2,6 +2,13 @@
Added for symmetry with the core API
"""

from functools import partial
from numba.core.extending import intrinsic as _intrinsic
from numba.core.extending import overload as _overload
from numba.core.extending import overload_method as _overload_method
from numba.core.extending import overload_attribute as _overload_attribute

intrinsic = _intrinsic(target='cuda')
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Should I have used functools.partial for intrinsic as well?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't thin so, as the only argument you can feed to intrinsic is target.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is it worth doing functools.update_wrapper on these so as to inherit the docstrings etc?

overload = partial(_overload, target='cuda')
guilhermeleobas marked this conversation as resolved.
Show resolved Hide resolved
overload_method = partial(_overload_method, target='cuda')
overload_attribute = partial(_overload_attribute, target='cuda')