def print_varargs(context, builder, sig, args):
"""This function is a generic 'print' wrapper for arbitrary types.
It dispatches to the appropriate 'print' implementations above
depending on the detected real types in the signature."""
vprint = nvvmutils.declare_vprint(builder.module)
formats = []
values = []
for i, (argtype, argval) in enumerate(zip(sig.args, args)):
argfmt, argvals = print_item(argtype, context, builder, argval)
formats.append(argfmt)
values.extend(argvals)
rawfmt = " ".join(formats) + "\n"
if len(args) > 32:
msg = ('CUDA print() cannot print more than 32 items. '
'The raw format string will be emitted by the kernel instead.')
warn(msg, NumbaWarning)
rawfmt = rawfmt.replace('%', '%%')
fmt = context.insert_string_const_addrspace(builder, rawfmt)
array = cgutils.make_anonymous_struct(builder, values)
arrayptr = cgutils.alloca_once_value(builder, array)
vprint = nvvmutils.declare_vprint(builder.module)
builder.call(vprint, (fmt, builder.bitcast(arrayptr, voidptr)))
return context.get_dummy_value()
def as_data(self, builder, values):
"""
Return the given values packed as a data structure.
"""
elems = [self._models[i].as_data(builder, values[i])
for i in self._pack_map]
return cgutils.make_anonymous_struct(builder, elems)
def stararg_handler(index, param, vars):
stararg_ty = signature.args[index]
assert isinstance(stararg_ty, types.BaseTuple), stararg_ty
values = [
self._cast_var(var, sigty)
for var, sigty in zip(vars, stararg_ty)
]
return cgutils.make_anonymous_struct(self.builder, values)
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, ir.IntType(value_type.bitwidth))
fname = 'llvm.nvvm.shfl.sync.i32'
lmod = builder.module
fnty = ir.FunctionType(
ir.LiteralStructType((ir.IntType(32), ir.IntType(1))),
(ir.IntType(32), ir.IntType(32), ir.IntType(32),
ir.IntType(32), ir.IntType(32))
)
func = cgutils.get_or_insert_function(lmod, fnty, 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, ir.FloatType())
ret = cgutils.make_anonymous_struct(builder, (fv, pred))
else:
value1 = builder.trunc(value, ir.IntType(32))
value_lshr = builder.lshr(value, context.get_constant(types.i8, 32))
value2 = builder.trunc(value_lshr, ir.IntType(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, ir.IntType(64))
rv2_64 = builder.zext(rv2, ir.IntType(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, ir.DoubleType())
ret = cgutils.make_anonymous_struct(builder, (rv, pred))
return ret
def frexp_impl(context, builder, sig, args):
val, = args
fltty = context.get_data_type(sig.args[0])
intty = context.get_data_type(sig.return_type[1])
expptr = cgutils.alloca_once(builder, intty, name='exp')
fnty = Type.function(fltty, (fltty, Type.pointer(intty)))
fname = {
"float": "numba_frexpf",
"double": "numba_frexp",
}[str(fltty)]
fn = cgutils.get_or_insert_function(builder.module, fnty, fname)
res = builder.call(fn, (val, expptr))
res = cgutils.make_anonymous_struct(builder, (res, builder.load(expptr)))
return impl_ret_untracked(context, builder, sig.return_type, res)
def lower_init_func(self, lower):
"""
Lower the generator's initialization function (which will fill up
the passed-by-reference generator structure).
"""
lower.setup_function(self.fndesc)
builder = lower.builder
# Insert the generator into the target context in order to allow
# calling from other Numba-compiled functions.
lower.context.insert_generator(self.gentype, self.gendesc,
[self.library])
# Init argument values
lower.extract_function_arguments()
lower.pre_lower()
# Initialize the return structure (i.e. the generator structure).
retty = self.context.get_return_type(self.gentype)
# Structure index #0: the initial resume index (0 == start of generator)
resume_index = self.context.get_constant(types.int32, 0)
# Structure index #1: the function arguments
argsty = retty.elements[1]
statesty = retty.elements[2]
lower.debug_print("# low_init_func incref")
# Incref all NRT arguments before storing into generator states
if self.context.enable_nrt:
for argty, argval in zip(self.fndesc.argtypes, lower.fnargs):
self.context.nrt.incref(builder, argty, argval)
# Filter out omitted arguments
argsval = self.arg_packer.as_data(builder, lower.fnargs)
# Zero initialize states
statesval = Constant(statesty, None)
gen_struct = cgutils.make_anonymous_struct(builder,
[resume_index, argsval,
statesval],
retty)
retval = self.box_generator_struct(lower, gen_struct)
lower.debug_print("# low_init_func before return")
self.call_conv.return_value(builder, retval)
lower.post_lower()
def print_varargs(context, builder, sig, args):
"""This function is a generic 'print' wrapper for arbitrary types.
It dispatches to the appropriate 'print' implementations above
depending on the detected real types in the signature."""
vprint = nvvmutils.declare_vprint(builder.module)
formats = []
values = []
for i, (argtype, argval) in enumerate(zip(sig.args, args)):
argfmt, argvals = print_item(argtype, context, builder, argval)
formats.append(argfmt)
values.extend(argvals)
rawfmt = " ".join(formats) + "\n"
fmt = context.insert_string_const_addrspace(builder, rawfmt)
array = cgutils.make_anonymous_struct(builder, values)
arrayptr = cgutils.alloca_once_value(builder, array)
vprint = nvvmutils.declare_vprint(builder.module)
builder.call(vprint, (fmt, builder.bitcast(arrayptr, voidptr)))
return context.get_dummy_value()
def range_to_range(context, builder, fromty, toty, val):
olditems = cgutils.unpack_tuple(builder, val, 3)
items = [
context.cast(builder, v, fromty.dtype, toty.dtype) for v in olditems
]
return cgutils.make_anonymous_struct(builder, items)
