def iternext_specific(self, context, builder, arrty, arr, result): zero = context.get_constant(types.intp, 0) one = context.get_constant(types.intp, 1) ndim = arrty.ndim nitems = arr.nitems index = builder.load(self.index) is_valid = builder.icmp(lc.ICMP_SLT, index, nitems) result.set_valid(is_valid) with cgutils.if_likely(builder, is_valid): ptr = builder.load(self.pointer) value = context.unpack_value(builder, arrty.dtype, ptr) if kind == 'flat': result.yield_(value) else: # ndenumerate(): fetch and increment indices indices = self.indices idxvals = [ builder.load(cgutils.gep(builder, indices, dim)) for dim in range(ndim) ] idxtuple = cgutils.pack_array(builder, idxvals) result.yield_( cgutils.make_anonymous_struct( builder, [idxtuple, value])) _increment_indices_array(context, builder, arrty, arr, indices) index = builder.add(index, one) builder.store(index, self.index) ptr = cgutils.pointer_add(builder, ptr, self.stride) builder.store(ptr, self.pointer)
def to_native_tuple(self, obj, typ): """ Convert tuple *obj* to a native array (if homogenous) or structure. """ n = len(typ) values = [] cleanups = [] is_error = cgutils.false_bit for i, eltype in enumerate(typ): elem = self.tuple_getitem(obj, i) native = self.to_native_value(elem, eltype) values.append(native.value) is_error = self.builder.or_(is_error, native.is_error) if native.cleanup is not None: cleanups.append(native.cleanup) if cleanups: def cleanup(): for func in reversed(cleanups): func() else: cleanup = None if isinstance(typ, types.UniTuple): value = cgutils.pack_array(self.builder, values) else: value = cgutils.make_anonymous_struct(self.builder, values) return NativeValue(value, is_error=is_error, cleanup=cleanup)
def iternext_specific(self, context, builder, arrty, arr, result): zero = context.get_constant(types.intp, 0) one = context.get_constant(types.intp, 1) ndim = arrty.ndim nitems = arr.nitems index = builder.load(self.index) is_valid = builder.icmp(lc.ICMP_SLT, index, nitems) result.set_valid(is_valid) with cgutils.if_likely(builder, is_valid): ptr = builder.load(self.pointer) value = context.unpack_value(builder, arrty.dtype, ptr) if kind == 'flat': result.yield_(value) else: # ndenumerate(): fetch and increment indices indices = self.indices idxvals = [builder.load(cgutils.gep(builder, indices, dim)) for dim in range(ndim)] idxtuple = cgutils.pack_array(builder, idxvals) result.yield_( cgutils.make_anonymous_struct(builder, [idxtuple, value])) _increment_indices_array(context, builder, arrty, arr, indices) index = builder.add(index, one) builder.store(index, self.index) ptr = cgutils.pointer_add(builder, ptr, self.stride) builder.store(ptr, self.pointer)
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 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 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 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 = builder.module.get_or_insert_function(fnty, name=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 to_native_tuple(self, obj, typ): """ Convert tuple *obj* to a native array (if homogenous) or structure. """ n = len(typ) values = [] for i, eltype in enumerate(typ): elem = self.tuple_getitem(obj, i) values.append(self.to_native_value(elem, eltype)) if isinstance(typ, types.UniTuple): return cgutils.pack_array(self.builder, values) else: return cgutils.make_anonymous_struct(self.builder, values)
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_module(builder).get_or_insert_function(fnty, name=fname) res = builder.call(fn, (val, expptr)) return cgutils.make_anonymous_struct(builder, (res, builder.load(expptr)))
def iternext_enumerate(context, builder, sig, args, result): [enumty] = sig.args [enum] = args enumcls = make_enumerate_cls(enumty) enum = enumcls(context, builder, value=enum) count = builder.load(enum.count) ncount = builder.add(count, context.get_constant(types.intp, 1)) builder.store(ncount, enum.count) srcres = call_iternext(context, builder, enumty.source_type, enum.iter) is_valid = srcres.is_valid() result.set_valid(is_valid) with builder.if_then(is_valid): srcval = srcres.yielded_value() result.yield_(cgutils.make_anonymous_struct(builder, [count, srcval]))
def iternext_enumerate(context, builder, sig, args, result): [enumty] = sig.args [enum] = args enumcls = make_enumerate_cls(enumty) enum = enumcls(context, builder, value=enum) count = builder.load(enum.count) ncount = builder.add(count, context.get_constant(types.intp, 1)) builder.store(ncount, enum.count) srcres = call_iternext(context, builder, enumty.source_type, enum.iter) is_valid = srcres.is_valid() result.set_valid(is_valid) with cgutils.ifthen(builder, is_valid): srcval = srcres.yielded_value() result.yield_(cgutils.make_anonymous_struct(builder, [count, srcval]))
def iternext_zip(context, builder, sig, args, result): [zip_type] = sig.args [zipobj] = args zipcls = make_zip_cls(zip_type) zipobj = zipcls(context, builder, value=zipobj) if len(zipobj) == 0: # zip() is an empty iterator result.set_exhausted() return is_valid = context.get_constant(types.boolean, True) values = [] for iterobj, srcty in zip(zipobj, zip_type.source_types): srcres = call_iternext(context, builder, srcty, iterobj) is_valid = builder.and_(is_valid, srcres.is_valid()) values.append(srcres.yielded_value()) result.set_valid(is_valid) with builder.if_then(is_valid): result.yield_(cgutils.make_anonymous_struct(builder, values))
def iternext_zip(context, builder, sig, args, result): [zip_type] = sig.args [zipobj] = args zipcls = make_zip_cls(zip_type) zipobj = zipcls(context, builder, value=zipobj) if len(zipobj) == 0: # zip() is an empty iterator result.set_exhausted() return is_valid = context.get_constant(types.boolean, True) values = [] for iterobj, srcty in zip(zipobj, zip_type.source_types): srcres = call_iternext(context, builder, srcty, iterobj) is_valid = builder.and_(is_valid, srcres.is_valid()) values.append(srcres.yielded_value()) result.set_valid(is_valid) with cgutils.ifthen(builder, is_valid): result.yield_(cgutils.make_anonymous_struct(builder, values))
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 cast(self, builder, val, fromty, toty): if fromty == toty or toty == types.Any or isinstance(toty, types.Kind): return val elif isinstance(fromty, types.Integer) and isinstance(toty, types.Integer): if toty.bitwidth == fromty.bitwidth: # Just a change of signedness return val elif toty.bitwidth < fromty.bitwidth: # Downcast return builder.trunc(val, self.get_value_type(toty)) elif fromty.signed: # Signed upcast return builder.sext(val, self.get_value_type(toty)) else: # Unsigned upcast return builder.zext(val, self.get_value_type(toty)) elif fromty in types.real_domain and toty in types.real_domain: lty = self.get_value_type(toty) if fromty == types.float32 and toty == types.float64: return builder.fpext(val, lty) elif fromty == types.float64 and toty == types.float32: return builder.fptrunc(val, lty) elif fromty in types.real_domain and toty in types.complex_domain: if fromty == types.float32: if toty == types.complex128: real = self.cast(builder, val, fromty, types.float64) else: real = val elif fromty == types.float64: if toty == types.complex64: real = self.cast(builder, val, fromty, types.float32) else: real = val if toty == types.complex128: imag = self.get_constant(types.float64, 0) elif toty == types.complex64: imag = self.get_constant(types.float32, 0) else: raise Exception("unreachable") cmplx = self.make_complex(toty)(self, builder) cmplx.real = real cmplx.imag = imag return cmplx._getvalue() elif fromty in types.integer_domain and toty in types.real_domain: lty = self.get_value_type(toty) if fromty in types.signed_domain: return builder.sitofp(val, lty) else: return builder.uitofp(val, lty) elif toty in types.integer_domain and fromty in types.real_domain: lty = self.get_value_type(toty) if toty in types.signed_domain: return builder.fptosi(val, lty) else: return builder.fptoui(val, lty) elif fromty in types.integer_domain and toty in types.complex_domain: cmplxcls, flty = builtins.get_complex_info(toty) cmpl = cmplxcls(self, builder) cmpl.real = self.cast(builder, val, fromty, flty) cmpl.imag = self.get_constant(flty, 0) return cmpl._getvalue() elif fromty in types.complex_domain and toty in types.complex_domain: srccls, srcty = builtins.get_complex_info(fromty) dstcls, dstty = builtins.get_complex_info(toty) src = srccls(self, builder, value=val) dst = dstcls(self, builder) dst.real = self.cast(builder, src.real, srcty, dstty) dst.imag = self.cast(builder, src.imag, srcty, dstty) return dst._getvalue() elif (isinstance(fromty, (types.UniTuple, types.Tuple)) and isinstance(toty, (types.UniTuple, types.Tuple)) and len(toty) == len(fromty)): olditems = cgutils.unpack_tuple(builder, val, len(fromty)) items = [self.cast(builder, i, f, t) for i, f, t in zip(olditems, fromty, toty)] return cgutils.make_anonymous_struct(builder, items) elif toty == types.boolean: return self.is_true(builder, fromty, val) elif fromty == types.boolean: # first promote to int32 asint = builder.zext(val, Type.int()) # then promote to number return self.cast(builder, asint, types.int32, toty) elif fromty == types.none and isinstance(toty, types.Optional): return self.make_optional_none(builder, toty.type) elif isinstance(toty, types.Optional): casted = self.cast(builder, val, fromty, toty.type) return self.make_optional_value(builder, toty.type, casted) elif isinstance(fromty, types.Optional): optty = self.make_optional(fromty) optval = optty(self, builder, value=val) validbit = cgutils.as_bool_bit(builder, optval.valid) with cgutils.if_unlikely(builder, builder.not_(validbit)): msg = "expected %s, got None" % (fromty.type,) self.call_conv.return_user_exc(builder, TypeError, (msg,)) return optval.data elif (isinstance(fromty, types.Array) and isinstance(toty, types.Array)): # Type inference should have prevented illegal array casting. assert toty.layout == 'A' return val elif (isinstance(fromty, types.List) and isinstance(toty, types.List)): # Casting from non-reflected to reflected assert fromty.dtype == toty.dtype return val elif (isinstance(fromty, types.RangeType) and isinstance(toty, types.RangeType)): olditems = cgutils.unpack_tuple(builder, val, 3) items = [self.cast(builder, v, fromty.dtype, toty.dtype) for v in olditems] return cgutils.make_anonymous_struct(builder, items) elif fromty in types.integer_domain and toty == types.voidptr: return builder.inttoptr(val, self.get_value_type(toty)) raise NotImplementedError("cast", val, fromty, toty)
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)
def iternext_specific(self, context, builder, arrty, arr, result): ndim = arrty.ndim data = arr.data shapes = cgutils.unpack_tuple(builder, arr.shape, ndim) strides = cgutils.unpack_tuple(builder, arr.strides, ndim) indices = self.indices pointers = self.pointers zero = context.get_constant(types.intp, 0) one = context.get_constant(types.intp, 1) bbend = cgutils.append_basic_block(builder, 'end') # Catch already computed iterator exhaustion is_exhausted = cgutils.as_bool_bit( builder, builder.load(self.exhausted)) with cgutils.if_unlikely(builder, is_exhausted): result.set_valid(False) builder.branch(bbend) result.set_valid(True) # Current pointer inside last dimension last_ptr = cgutils.gep(builder, pointers, ndim - 1) ptr = builder.load(last_ptr) value = context.unpack_value(builder, arrty.dtype, ptr) if kind == 'flat': result.yield_(value) else: # ndenumerate() => yield (indices, value) idxvals = [builder.load(cgutils.gep(builder, indices, dim)) for dim in range(ndim)] idxtuple = cgutils.pack_array(builder, idxvals) result.yield_( cgutils.make_anonymous_struct(builder, [idxtuple, value])) # Update indices and pointers by walking from inner # dimension to outer. for dim in reversed(range(ndim)): idxptr = cgutils.gep(builder, indices, dim) idx = builder.add(builder.load(idxptr), one) count = shapes[dim] stride = strides[dim] in_bounds = builder.icmp(lc.ICMP_SLT, idx, count) with cgutils.if_likely(builder, in_bounds): # Index is valid => pointer can simply be incremented. builder.store(idx, idxptr) ptrptr = cgutils.gep(builder, pointers, dim) ptr = builder.load(ptrptr) ptr = cgutils.pointer_add(builder, ptr, stride) builder.store(ptr, ptrptr) # Reset pointers in inner dimensions for inner_dim in range(dim + 1, ndim): ptrptr = cgutils.gep(builder, pointers, inner_dim) builder.store(ptr, ptrptr) builder.branch(bbend) # Reset index and continue with next dimension builder.store(zero, idxptr) # End of array builder.store(cgutils.true_byte, self.exhausted) builder.branch(bbend) builder.position_at_end(bbend)
def iternext_specific(self, context, builder, arrty, arr, result): ndim = arrty.ndim data = arr.data shapes = cgutils.unpack_tuple(builder, arr.shape, ndim) strides = cgutils.unpack_tuple(builder, arr.strides, ndim) indices = self.indices pointers = self.pointers zero = context.get_constant(types.intp, 0) one = context.get_constant(types.intp, 1) bbend = cgutils.append_basic_block(builder, 'end') # Catch already computed iterator exhaustion is_exhausted = cgutils.as_bool_bit( builder, builder.load(self.exhausted)) with cgutils.if_unlikely(builder, is_exhausted): result.set_valid(False) builder.branch(bbend) result.set_valid(True) # Current pointer inside last dimension last_ptr = cgutils.gep(builder, pointers, ndim - 1) ptr = builder.load(last_ptr) value = context.unpack_value(builder, arrty.dtype, ptr) if kind == 'flat': result.yield_(value) else: # ndenumerate() => yield (indices, value) idxvals = [ builder.load(cgutils.gep(builder, indices, dim)) for dim in range(ndim) ] idxtuple = cgutils.pack_array(builder, idxvals) result.yield_( cgutils.make_anonymous_struct(builder, [idxtuple, value])) # Update indices and pointers by walking from inner # dimension to outer. for dim in reversed(range(ndim)): idxptr = cgutils.gep(builder, indices, dim) idx = builder.add(builder.load(idxptr), one) count = shapes[dim] stride = strides[dim] in_bounds = builder.icmp(lc.ICMP_SLT, idx, count) with cgutils.if_likely(builder, in_bounds): # Index is valid => pointer can simply be incremented. builder.store(idx, idxptr) ptrptr = cgutils.gep(builder, pointers, dim) ptr = builder.load(ptrptr) ptr = cgutils.pointer_add(builder, ptr, stride) builder.store(ptr, ptrptr) # Reset pointers in inner dimensions for inner_dim in range(dim + 1, ndim): ptrptr = cgutils.gep(builder, pointers, inner_dim) builder.store(ptr, ptrptr) builder.branch(bbend) # Reset index and continue with next dimension builder.store(zero, idxptr) # End of array builder.store(cgutils.true_byte, self.exhausted) builder.branch(bbend) builder.position_at_end(bbend)
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)