def __init__(self, elements, name=''):
"""Initialize the instance."""
element_types = [element._ptr for element in elements]
super(MiddleIrTypeStruct, self).__init__(
Type.struct(element_types))#, name)) #FIXME name is unused in LLVMLite?
self.type = list(elements)
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 ptx_vote_sync(context, builder, sig, args):
fname = 'llvm.nvvm.vote.sync'
lmod = builder.module
fnty = Type.function(Type.struct((Type.int(32), Type.int(1))),
(Type.int(32), Type.int(32), Type.int(1)))
func = lmod.get_or_insert_function(fnty, name=fname)
return builder.call(func, args)
def ptx_vote_sync(context, builder, sig, args):
fname = 'llvm.nvvm.vote.sync'
lmod = builder.module
fnty = Type.function(Type.struct((Type.int(32), Type.int(1))),
(Type.int(32), Type.int(32), Type.int(1)))
func = lmod.get_or_insert_function(fnty, name=fname)
return builder.call(func, args)
def make_anonymous_struct(builder, values):
"""
Create an anonymous struct containing the given LLVM *values*.
"""
struct_type = Type.struct([v.type for v in values])
struct_val = Constant.undef(struct_type)
for i, v in enumerate(values):
struct_val = builder.insert_value(struct_val, v, i)
return struct_val
def make_anonymous_struct(builder, values):
"""
Create an anonymous struct constant containing the given LLVM *values*.
"""
struct_type = Type.struct([v.type for v in values])
struct_val = Constant.undef(struct_type)
for i, v in enumerate(values):
struct_val = builder.insert_value(struct_val, v, i)
return struct_val
def ptx_match_all_sync(context, builder, sig, args):
mask, value = args
width = sig.args[1].bitwidth
if sig.args[1] in types.real_domain:
value = builder.bitcast(value, Type.int(width))
fname = 'llvm.nvvm.match.all.sync.i{}'.format(width)
lmod = builder.module
fnty = Type.function(Type.struct((Type.int(32), Type.int(1))),
(Type.int(32), Type.int(width)))
func = lmod.get_or_insert_function(fnty, name=fname)
return builder.call(func, (mask, value))
def get_value_type(self, ty):
if ty == types.boolean:
return Type.int(1)
dataty = self.get_data_type(ty)
if isinstance(ty, types.Record):
# Record data are passed by refrence
memory = dataty.elements[0]
return Type.struct([Type.pointer(memory)])
return dataty
def get_value_type(self, ty):
if ty == types.boolean:
return Type.int(1)
dataty = self.get_data_type(ty)
if isinstance(ty, types.Record):
# Record data are passed by refrence
memory = dataty.elements[0]
return Type.struct([Type.pointer(memory)])
return dataty
def ptx_match_all_sync(context, builder, sig, args):
mask, value = args
width = sig.args[1].bitwidth
if sig.args[1] in types.real_domain:
value = builder.bitcast(value, Type.int(width))
fname = 'llvm.nvvm.match.all.sync.i{}'.format(width)
lmod = builder.module
fnty = Type.function(Type.struct((Type.int(32), Type.int(1))),
(Type.int(32), Type.int(width)))
func = lmod.get_or_insert_function(fnty, name=fname)
return builder.call(func, (mask, value))
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 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
_trace_refs_ = hasattr(sys, "getobjects")
_plat_bits = struct_.calcsize("@P") * 8
_int8 = Type.int(8)
_int32 = Type.int(32)
_void_star = Type.pointer(_int8)
_int8_star = _void_star
_sizeof_py_ssize_t = ctypes.sizeof(getattr(ctypes, "c_size_t"))
_llvm_py_ssize_t = Type.int(_sizeof_py_ssize_t * 8)
if _trace_refs_:
_pyobject_head = Type.struct(
[_void_star, _void_star, _llvm_py_ssize_t, _void_star])
_pyobject_head_init = Constant.struct([
Constant.null(_void_star), # _ob_next
Constant.null(_void_star), # _ob_prev
Constant.int(_llvm_py_ssize_t, 1), # ob_refcnt
Constant.null(_void_star), # ob_type
])
else:
_pyobject_head = Type.struct([_llvm_py_ssize_t, _void_star])
_pyobject_head_init = Constant.struct([
Constant.int(_llvm_py_ssize_t, 1), # ob_refcnt
Constant.null(_void_star), # ob_type
])
_pyobject_head_p = Type.pointer(_pyobject_head)
def get_data_type(self, ty):
"""
Get a LLVM data representation of the Numba type *ty* that is safe
for storage. Record data are stored as byte array.
The return value is a llvmlite.ir.Type object, or None if the type
is an opaque pointer (???).
"""
try:
fac = type_registry.match(ty)
except KeyError:
pass
else:
return fac(self, ty)
if (isinstance(ty, types.Dummy) or isinstance(ty, types.Module)
or isinstance(ty, types.Function)
or isinstance(ty, types.Dispatcher)
or isinstance(ty, types.Object)
or isinstance(ty, types.Macro)):
return PYOBJECT
elif isinstance(ty, types.CPointer):
dty = self.get_data_type(ty.dtype)
return Type.pointer(dty)
elif isinstance(ty, types.Optional):
return self.get_struct_type(self.make_optional(ty))
elif isinstance(ty, types.Array):
return self.get_struct_type(self.make_array(ty))
elif isinstance(ty, types.UniTuple):
dty = self.get_value_type(ty.dtype)
return Type.array(dty, ty.count)
elif isinstance(ty, types.Tuple):
dtys = [self.get_value_type(t) for t in ty]
return Type.struct(dtys)
elif isinstance(ty, types.Record):
# Record are represented as byte array
return Type.struct([Type.array(Type.int(8), ty.size)])
elif isinstance(ty, types.UnicodeCharSeq):
charty = Type.int(numpy_support.sizeof_unicode_char * 8)
return Type.struct([Type.array(charty, ty.count)])
elif isinstance(ty, types.CharSeq):
charty = Type.int(8)
return Type.struct([Type.array(charty, ty.count)])
elif ty in STRUCT_TYPES:
return self.get_struct_type(STRUCT_TYPES[ty])
else:
try:
impl = struct_registry.match(ty)
except KeyError:
pass
else:
return self.get_struct_type(impl(ty))
if isinstance(ty, types.Pair):
pairty = self.make_pair(ty.first_type, ty.second_type)
return self.get_struct_type(pairty)
else:
return LTYPEMAP[ty]
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
])
int32_array = pointer(array_type(int_type))
int64_array = pointer(array_type(Type.int(64)))
double_array = pointer(array_type(double_type))
def get_data_type(self, ty):
"""
Get a LLVM data representation of the Numba type *ty* that is safe
for storage. Record data are stored as byte array.
The return value is a llvmlite.ir.Type object, or None if the type
is an opaque pointer (???).
"""
try:
fac = type_registry.match(ty)
except KeyError:
pass
else:
return fac(self, ty)
if (
isinstance(ty, types.Dummy)
or isinstance(ty, types.Module)
or isinstance(ty, types.Function)
or isinstance(ty, types.Dispatcher)
or isinstance(ty, types.Object)
or isinstance(ty, types.Macro)
):
return PYOBJECT
elif isinstance(ty, types.CPointer):
dty = self.get_data_type(ty.dtype)
return Type.pointer(dty)
elif isinstance(ty, types.Optional):
return self.get_struct_type(self.make_optional(ty))
elif isinstance(ty, types.Array):
return self.get_struct_type(self.make_array(ty))
elif isinstance(ty, types.UniTuple):
dty = self.get_value_type(ty.dtype)
return Type.array(dty, ty.count)
elif isinstance(ty, types.Tuple):
dtys = [self.get_value_type(t) for t in ty]
return Type.struct(dtys)
elif isinstance(ty, types.Record):
# Record are represented as byte array
return Type.struct([Type.array(Type.int(8), ty.size)])
elif isinstance(ty, types.UnicodeCharSeq):
charty = Type.int(numpy_support.sizeof_unicode_char * 8)
return Type.struct([Type.array(charty, ty.count)])
elif isinstance(ty, types.CharSeq):
charty = Type.int(8)
return Type.struct([Type.array(charty, ty.count)])
elif ty in STRUCT_TYPES:
return self.get_struct_type(STRUCT_TYPES[ty])
else:
try:
impl = struct_registry.match(ty)
except KeyError:
pass
else:
return self.get_struct_type(impl(ty))
if isinstance(ty, types.Pair):
pairty = self.make_pair(ty.first_type, ty.second_type)
return self.get_struct_type(pairty)
else:
return LTYPEMAP[ty]
def get_struct_type(self, struct):
"""
Get the LLVM struct type for the given Structure class *struct*.
"""
fields = [self.get_value_type(v) for _, v in struct._fields]
return Type.struct(fields)
def get_struct_type(self, struct):
"""
Get the LLVM struct type for the given Structure class *struct*.
"""
fields = [self.get_value_type(v) for _, v in struct._fields]
return Type.struct(fields)
_trace_refs_ = hasattr(sys, "getobjects")
_plat_bits = struct_.calcsize("@P") * 8
_int8 = Type.int(8)
_int32 = Type.int(32)
_void_star = Type.pointer(_int8)
_int8_star = _void_star
_sizeof_py_ssize_t = ctypes.sizeof(getattr(ctypes, "c_size_t"))
_llvm_py_ssize_t = Type.int(_sizeof_py_ssize_t * 8)
if _trace_refs_:
_pyobject_head = Type.struct([_void_star, _void_star, _llvm_py_ssize_t, _void_star])
_pyobject_head_init = Constant.struct(
[
Constant.null(_void_star), # _ob_next
Constant.null(_void_star), # _ob_prev
Constant.int(_llvm_py_ssize_t, 1), # ob_refcnt
Constant.null(_void_star), # ob_type
]
)
else:
_pyobject_head = Type.struct([_llvm_py_ssize_t, _void_star])
_pyobject_head_init = Constant.struct(
[Constant.int(_llvm_py_ssize_t, 1), Constant.null(_void_star)] # ob_refcnt # ob_type
)
