def hypot_u64_impl(context, builder, sig, args):
[x, y] = args
y = builder.sitofp(y, Type.double())
x = builder.sitofp(x, Type.double())
fsig = signature(types.float64, types.float64, types.float64)
res = hypot_float_impl(context, builder, fsig, (x, y))
return impl_ret_untracked(context, builder, sig.return_type, res)
def int_utruediv_impl(context, builder, sig, args):
x, y = args
fx = builder.uitofp(x, Type.double())
fy = builder.uitofp(y, Type.double())
cgutils.guard_zero(context, builder, y,
(ZeroDivisionError, "division by zero"))
return builder.fdiv(fx, fy)
def hypot_u64_impl(context, builder, sig, args):
[x, y] = args
y = builder.sitofp(y, Type.double())
x = builder.sitofp(x, Type.double())
fsig = signature(types.float64, types.float64, types.float64)
res = hypot_float_impl(context, builder, fsig, (x, y))
return impl_ret_untracked(context, builder, sig.return_type, res)
def int_utruediv_impl(context, builder, sig, args):
x, y = args
fx = builder.uitofp(x, Type.double())
fy = builder.uitofp(y, Type.double())
cgutils.guard_zero(context, builder, y,
(ZeroDivisionError, "division by zero"))
return builder.fdiv(fx, fy)
def atan2_f64_impl(context, builder, sig, args):
assert len(args) == 2
mod = cgutils.get_module(builder)
fnty = Type.function(Type.double(), [Type.double(), Type.double()])
# Workaround atan2() issues under Windows
fname = "atan2_fixed" if sys.platform == "win32" else "atan2"
fn = mod.get_or_insert_function(fnty, name=fname)
return builder.call(fn, args)
def atan2_f64_impl(context, builder, sig, args):
assert len(args) == 2
mod = cgutils.get_module(builder)
fnty = Type.function(Type.double(), [Type.double(), Type.double()])
# Workaround atan2() issues under Windows
fname = "atan2_fixed" if sys.platform == "win32" else "atan2"
fn = mod.get_or_insert_function(fnty, name=fname)
return builder.call(fn, args)
def implementer(context, builder, sig, args):
[val] = args
input_type = sig.args[0]
if input_type.signed:
fpval = builder.sitofp(val, Type.double())
else:
fpval = builder.uitofp(val, Type.double())
sig = signature(types.float64, types.float64)
return wrapped_impl(context, builder, sig, [fpval])
def atan2_f64_impl(context, builder, sig, args):
assert len(args) == 2
mod = builder.module
fnty = Type.function(Type.double(), [Type.double(), Type.double()])
# Workaround atan2() issues under Windows
fname = "atan2_fixed" if sys.platform == "win32" else "atan2"
fn = cgutils.insert_pure_function(builder.module, fnty, name=fname)
res = builder.call(fn, args)
return impl_ret_untracked(context, builder, sig.return_type, res)
def implementer(context, builder, sig, args):
[val] = args
input_type = sig.args[0]
if input_type.signed:
fpval = builder.sitofp(val, Type.double())
else:
fpval = builder.uitofp(val, Type.double())
sig = signature(types.float64, types.float64)
return wrapped_impl(context, builder, sig, [fpval])
def round_impl_f64(context, builder, sig, args):
module = cgutils.get_module(builder)
fnty = Type.function(Type.double(), [Type.double()])
if utils.IS_PY3:
fn = module.get_or_insert_function(fnty, name="numba.round")
else:
fn = module.get_or_insert_function(fnty, name="round")
assert fn.is_declaration
return builder.call(fn, args)
def atan2_f64_impl(context, builder, sig, args):
assert len(args) == 2
mod = builder.module
fnty = Type.function(Type.double(), [Type.double(), Type.double()])
# Workaround atan2() issues under Windows
fname = "atan2_fixed" if sys.platform == "win32" else "atan2"
fn = cgutils.insert_pure_function(builder.module, fnty, name=fname)
res = builder.call(fn, args)
return impl_ret_untracked(context, builder, sig.return_type, res)
def __init__(self, context, builder):
"""
Note: Maybe called multiple times when lowering a function
"""
from numba.targets import boxing
self.context = context
self.builder = builder
self.module = builder.basic_block.function.module
# A unique mapping of serialized objects in this module
try:
self.module.__serialized
except AttributeError:
self.module.__serialized = {}
# Initialize types
self.pyobj = self.context.get_argument_type(types.pyobject)
self.voidptr = Type.pointer(Type.int(8))
self.long = Type.int(ctypes.sizeof(ctypes.c_long) * 8)
self.ulong = self.long
self.longlong = Type.int(ctypes.sizeof(ctypes.c_ulonglong) * 8)
self.ulonglong = self.longlong
self.double = Type.double()
self.py_ssize_t = self.context.get_value_type(types.intp)
self.cstring = Type.pointer(Type.int(8))
self.gil_state = Type.int(_helperlib.py_gil_state_size * 8)
self.py_buffer_t = ir.ArrayType(ir.IntType(8), _helperlib.py_buffer_size)
def __init__(self, context, builder):
"""
Note: Maybe called multiple times when lowering a function
"""
from numba.targets import boxing
self.context = context
self.builder = builder
self.module = builder.basic_block.function.module
# A unique mapping of serialized objects in this module
try:
self.module.__serialized
except AttributeError:
self.module.__serialized = {}
# Initialize types
self.pyobj = self.context.get_argument_type(types.pyobject)
self.voidptr = Type.pointer(Type.int(8))
self.long = Type.int(ctypes.sizeof(ctypes.c_long) * 8)
self.ulonglong = Type.int(ctypes.sizeof(ctypes.c_ulonglong) * 8)
self.longlong = self.ulonglong
self.double = Type.double()
self.py_ssize_t = self.context.get_value_type(types.intp)
self.cstring = Type.pointer(Type.int(8))
self.gil_state = Type.int(_helperlib.py_gil_state_size * 8)
self.py_buffer_t = ir.ArrayType(ir.IntType(8), _helperlib.py_buffer_size)
def is_scalar_neg(builder, value):
"""is _value_ negative?. Assumes _value_ is signed"""
nullval = Constant.null(value.type)
if value.type in (Type.float(), Type.double()):
isneg = builder.fcmp(lc.FCMP_OLT, value, nullval)
else:
isneg = builder.icmp(lc.ICMP_SLT, value, nullval)
return isneg
def is_scalar_neg(builder, value):
"""is _value_ negative?. Assumes _value_ is signed"""
nullval = Constant.null(value.type)
if value.type in (Type.float(), Type.double()):
isneg = builder.fcmp(lc.FCMP_OLT, value, nullval)
else:
isneg = builder.icmp(lc.ICMP_SLT, value, nullval)
return isneg
def is_scalar_zero(builder, value):
"""
Return a predicate representing whether *value* is equal to zero.
"""
assert not is_pointer(value.type)
assert not is_struct(value.type)
nullval = Constant.null(value.type)
if value.type in (Type.float(), Type.double()):
isnull = builder.fcmp(lc.FCMP_OEQ, nullval, value)
else:
isnull = builder.icmp(lc.ICMP_EQ, nullval, value)
return isnull
def is_scalar_zero(builder, value):
"""
Return a predicate representing whether *value* is equal to zero.
"""
assert not is_pointer(value.type)
assert not is_struct(value.type)
nullval = Constant.null(value.type)
if value.type in (Type.float(), Type.double()):
isnull = builder.fcmp(lc.FCMP_OEQ, nullval, value)
else:
isnull = builder.icmp(lc.ICMP_EQ, nullval, value)
return isnull
def is_not_scalar_zero(builder, value):
"""
Return a predicate representin whether a *value* is not equal to zero.
not exactly "not is_scalar_zero" because of nans
"""
assert not is_pointer(value.type)
assert not is_struct(value.type)
nullval = Constant.null(value.type)
if value.type in (Type.float(), Type.double()):
isnull = builder.fcmp(lc.FCMP_UNE, nullval, value)
else:
isnull = builder.icmp(lc.ICMP_NE, nullval, value)
return isnull
def is_not_scalar_zero(builder, value):
"""
Return a predicate representin whether a *value* is not equal to zero.
not exactly "not is_scalar_zero" because of nans
"""
assert not is_pointer(value.type)
assert not is_struct(value.type)
nullval = Constant.null(value.type)
if value.type in (Type.float(), Type.double()):
isnull = builder.fcmp(lc.FCMP_UNE, nullval, value)
else:
isnull = builder.icmp(lc.ICMP_NE, nullval, value)
return isnull
def ptx_shfl_sync_i32(context, builder, sig, args):
"""
The NVVM intrinsic for shfl only supports i32, but the cuda intrinsic
function supports both 32 and 64 bit ints and floats, so for feature parity,
i64, f32, and f64 are implemented. Floats by way of bitcasting the float to
an int, then shuffling, then bitcasting back. And 64-bit values by packing
them into 2 32bit values, shuffling thoose, and then packing back together.
"""
mask, mode, value, index, clamp = args
value_type = sig.args[2]
if value_type in types.real_domain:
value = builder.bitcast(value, Type.int(value_type.bitwidth))
fname = 'llvm.nvvm.shfl.sync.i32'
lmod = builder.module
fnty = Type.function(
Type.struct((Type.int(32), Type.int(1))),
(Type.int(32), Type.int(32), Type.int(32), Type.int(32), Type.int(32))
)
func = lmod.get_or_insert_function(fnty, name=fname)
if value_type.bitwidth == 32:
ret = builder.call(func, (mask, mode, value, index, clamp))
if value_type == types.float32:
rv = builder.extract_value(ret, 0)
pred = builder.extract_value(ret, 1)
fv = builder.bitcast(rv, Type.float())
ret = cgutils.make_anonymous_struct(builder, (fv, pred))
else:
value1 = builder.trunc(value, Type.int(32))
value_lshr = builder.lshr(value, context.get_constant(types.i8, 32))
value2 = builder.trunc(value_lshr, Type.int(32))
ret1 = builder.call(func, (mask, mode, value1, index, clamp))
ret2 = builder.call(func, (mask, mode, value2, index, clamp))
rv1 = builder.extract_value(ret1, 0)
rv2 = builder.extract_value(ret2, 0)
pred = builder.extract_value(ret1, 1)
rv1_64 = builder.zext(rv1, Type.int(64))
rv2_64 = builder.zext(rv2, Type.int(64))
rv_shl = builder.shl(rv2_64, context.get_constant(types.i8, 32))
rv = builder.or_(rv_shl, rv1_64)
if value_type == types.float64:
rv = builder.bitcast(rv, Type.double())
ret = cgutils.make_anonymous_struct(builder, (rv, pred))
return ret
def __init__(self, context, builder):
"""
Note: Maybe called multiple times when lowering a function
"""
fix_python_api()
self.context = context
self.builder = builder
self.module = builder.basic_block.function.module
# Initialize types
self.pyobj = self.context.get_argument_type(types.pyobject)
self.voidptr = Type.pointer(Type.int(8))
self.long = Type.int(ctypes.sizeof(ctypes.c_long) * 8)
self.ulonglong = Type.int(ctypes.sizeof(ctypes.c_ulonglong) * 8)
self.longlong = self.ulonglong
self.double = Type.double()
self.py_ssize_t = self.context.get_value_type(types.intp)
self.cstring = Type.pointer(Type.int(8))
self.gil_state = Type.int(_helperlib.py_gil_state_size * 8)
def __init__(self, context, builder):
"""
Note: Maybe called multiple times when lowering a function
"""
fix_python_api()
self.context = context
self.builder = builder
self.module = builder.basic_block.function.module
# Initialize types
self.pyobj = self.context.get_argument_type(types.pyobject)
self.voidptr = Type.pointer(Type.int(8))
self.long = Type.int(ctypes.sizeof(ctypes.c_long) * 8)
self.ulonglong = Type.int(ctypes.sizeof(ctypes.c_ulonglong) * 8)
self.longlong = self.ulonglong
self.double = Type.double()
self.py_ssize_t = self.context.get_value_type(types.intp)
self.cstring = Type.pointer(Type.int(8))
self.gil_state = Type.int(_helperlib.py_gil_state_size * 8)
def ptx_shfl_sync_i32(context, builder, sig, args):
"""
The NVVM intrinsic for shfl only supports i32, but the cuda intrinsic function supports
both 32 and 64 bit ints and floats, so for feature parity, i64, f32, and f64 are implemented.
Floats by way of bitcasting the float to an int, then shuffling, then bitcasting back.
And 64-bit values by packing them into 2 32bit values, shuffling thoose, and then packing back together.
"""
mask, mode, value, index, clamp = args
value_type = sig.args[2]
if value_type in types.real_domain:
value = builder.bitcast(value, Type.int(value_type.bitwidth))
fname = 'llvm.nvvm.shfl.sync.i32'
lmod = builder.module
fnty = Type.function(
Type.struct((Type.int(32), Type.int(1))),
(Type.int(32), Type.int(32), Type.int(32), Type.int(32), Type.int(32))
)
func = lmod.get_or_insert_function(fnty, name=fname)
if value_type.bitwidth == 32:
ret = builder.call(func, (mask, mode, value, index, clamp))
if value_type == types.float32:
rv = builder.extract_value(ret, 0)
pred = builder.extract_value(ret, 1)
fv = builder.bitcast(rv, Type.float())
ret = cgutils.make_anonymous_struct(builder, (fv, pred))
else:
value1 = builder.trunc(value, Type.int(32))
value_lshr = builder.lshr(value, context.get_constant(types.i8, 32))
value2 = builder.trunc(value_lshr, Type.int(32))
ret1 = builder.call(func, (mask, mode, value1, index, clamp))
ret2 = builder.call(func, (mask, mode, value2, index, clamp))
rv1 = builder.extract_value(ret1, 0)
rv2 = builder.extract_value(ret2, 0)
pred = builder.extract_value(ret1, 1)
rv1_64 = builder.zext(rv1, Type.int(64))
rv2_64 = builder.zext(rv2, Type.int(64))
rv_shl = builder.shl(rv2_64, context.get_constant(types.i8, 32))
rv = builder.or_(rv_shl, rv1_64)
if value_type == types.float64:
rv = builder.bitcast(rv, Type.double())
ret = cgutils.make_anonymous_struct(builder, (rv, pred))
return ret
def printf(builder, format_string, *values):
str_const = Constant.stringz(format_string)
global_str_const = get_module(builder).add_global_variable(
str_const.type, '')
global_str_const.initializer = str_const
idx = [Constant.int(Type.int(32), 0), Constant.int(Type.int(32), 0)]
str_addr = global_str_const.gep(idx)
args = []
for v in values:
if isinstance(v, int):
args.append(Constant.int(Type.int(), v))
elif isinstance(v, float):
args.append(Constant.real(Type.double(), v))
else:
args.append(v)
functype = Type.function(Type.int(32), [Type.pointer(Type.int(8))], True)
fn = get_module(builder).get_or_insert_function(functype, 'printf')
builder.call(fn, [str_addr] + args)
def printf(builder, format_string, *values):
str_const = Constant.stringz(format_string)
global_str_const = get_module(builder).add_global_variable(str_const.type,
'')
global_str_const.initializer = str_const
idx = [Constant.int(Type.int(32), 0), Constant.int(Type.int(32), 0)]
str_addr = global_str_const.gep(idx)
args = []
for v in values:
if isinstance(v, int):
args.append(Constant.int(Type.int(), v))
elif isinstance(v, float):
args.append(Constant.real(Type.double(), v))
else:
args.append(v)
functype = Type.function(Type.int(32), [Type.pointer(Type.int(8))], True)
fn = get_module(builder).get_or_insert_function(functype, 'printf')
builder.call(fn, [str_addr] + args)
int32 = TCon("Int32")
int64 = TCon("Int64")
float32 = TCon("Float")
double64 = TCon("Double")
void = TCon("Void")
array = lambda t: TApp(TCon("Array"), t)
array_int32 = array(int32)
array_int64 = array(int64)
array_double64 = array(double64)
pointer = Type.pointer
int_type = Type.int()
float_type = Type.float()
double_type = Type.double()
bool_type = Type.int(1)
void_type = Type.void()
void_ptr = pointer(Type.int(8))
struct_type = Type.struct([])
def array_type(elt_type):
return Type.struct([
pointer(elt_type), # data
int_type, # dimensions
pointer(int_type), # shape
'ndarray_' + str(elt_type), # name
])
def int64_as_f64(builder, val):
"""
Bitcast a 64-bit integer into a double.
"""
assert val.type == Type.int(64)
return builder.bitcast(val, Type.double())
def hypot_u64_impl(context, builder, sig, args):
[x, y] = args
y = builder.sitofp(y, Type.double())
x = builder.sitofp(x, Type.double())
fsig = signature(types.float64, types.float64, types.float64)
return hypot_float_impl(context, builder, fsig, (x, y))
def complex_from_doubles(self, realval, imagval):
fnty = Type.function(self.pyobj, [Type.double(), Type.double()])
fn = self._get_function(fnty, name="PyComplex_FromDoubles")
return self.builder.call(fn, [realval, imagval])
def atan2_u64_impl(context, builder, sig, args):
[y, x] = args
y = builder.uitofp(y, Type.double())
x = builder.uitofp(x, Type.double())
fsig = signature(types.float64, types.float64, types.float64)
return atan2_f64_impl(context, builder, fsig, (y, x))
def complex_imag_as_double(self, cobj):
fnty = Type.function(Type.double(), [self.pyobj])
fn = self._get_function(fnty, name="PyComplex_ImagAsDouble")
return self.builder.call(fn, [cobj])
def __init__(self):
"""Initialize the instance."""
super(MiddleIrTypeDouble, self).__init__(Type.double())
def hypot_u64_impl(context, builder, sig, args):
[x, y] = args
y = builder.sitofp(y, Type.double())
x = builder.sitofp(x, Type.double())
fsig = signature(types.float64, types.float64, types.float64)
return hypot_float_impl(context, builder, fsig, (x, y))
def complex_from_doubles(self, realval, imagval):
fnty = Type.function(self.pyobj, [Type.double(), Type.double()])
fn = self._get_function(fnty, name="PyComplex_FromDoubles")
return self.builder.call(fn, [realval, imagval])
def complex_imag_as_double(self, cobj):
fnty = Type.function(Type.double(), [self.pyobj])
fn = self._get_function(fnty, name="PyComplex_ImagAsDouble")
return self.builder.call(fn, [cobj])
def int_utruediv_impl(context, builder, sig, args):
x, y = args
fx = builder.uitofp(x, Type.double())
fy = builder.uitofp(y, Type.double())
cgutils.guard_zero(context, builder, y)
return builder.fdiv(fx, fy)
types.pyobject: PYOBJECT,
types.boolean: Type.int(8),
types.uint8: Type.int(8),
types.uint16: Type.int(16),
types.uint32: Type.int(32),
types.uint64: Type.int(64),
types.int8: Type.int(8),
types.int16: Type.int(16),
types.int32: Type.int(32),
types.int64: Type.int(64),
types.float32: Type.float(),
types.float64: Type.double(),
}
STRUCT_TYPES = {
types.complex64: builtins.Complex64,
types.complex128: builtins.Complex128,
types.slice3_type: builtins.Slice,
}
class Overloads(object):
def __init__(self):
# A list of (signature, implementation)
self.versions = []
def find(self, sig):
def atan2_u64_impl(context, builder, sig, args):
[y, x] = args
y = builder.uitofp(y, Type.double())
x = builder.uitofp(x, Type.double())
fsig = signature(types.float64, types.float64, types.float64)
return atan2_float_impl(context, builder, fsig, (y, x))
GENERIC_POINTER = Type.pointer(Type.int(8))
PYOBJECT = GENERIC_POINTER
LTYPEMAP = {
types.pyobject: PYOBJECT,
types.boolean: Type.int(8),
types.uint8: Type.int(8),
types.uint16: Type.int(16),
types.uint32: Type.int(32),
types.uint64: Type.int(64),
types.int8: Type.int(8),
types.int16: Type.int(16),
types.int32: Type.int(32),
types.int64: Type.int(64),
types.float32: Type.float(),
types.float64: Type.double(),
}
STRUCT_TYPES = {
types.complex64: builtins.Complex64,
types.complex128: builtins.Complex128,
types.slice3_type: builtins.Slice,
}
class Overloads(object):
def __init__(self):
self.versions = []
def find(self, sig):
for i, ver in enumerate(self.versions):
def int64_as_f64(builder, val):
"""
Bitcast a 64-bit integer into a double.
"""
assert val.type == Type.int(64)
return builder.bitcast(val, Type.double())
