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)