def impl(context, builder, sig, args):
        [tyinp, tyout] = sig.args
        [inp, out] = args
        ndim = tyinp.ndim

        iary = context.make_array(tyinp)(context, builder, inp)
        oary = context.make_array(tyout)(context, builder, out)

        if asfloat:
            sig = typing.signature(types.float64, types.float64)
        else:
            sig = typing.signature(tyout.dtype, tyinp.dtype)

        fnwork = context.get_function(funckey, sig)
        intpty = context.get_value_type(types.intp)

        # TODO handle differing shape by mimicking broadcasting
        shape = cgutils.unpack_tuple(builder, iary.shape, ndim)
        with cgutils.loop_nest(builder, shape, intp=intpty) as indices:
            pi = cgutils.get_item_pointer(builder, tyinp, iary, indices)
            po = cgutils.get_item_pointer(builder, tyout, oary, indices)

            ival = builder.load(pi)
            if asfloat:
                dval = context.cast(builder, ival, tyinp.dtype, types.float64)
                dres = fnwork(builder, [dval])
                res = context.cast(builder, dres, types.float64, tyout.dtype)
            elif tyinp.dtype != tyout.dtype:
                tempres = fnwork(builder, [ival])
                res = context.cast(builder, tempres, tyinp.dtype, tyout.dtype)
            else:
                res = fnwork(builder, [ival])
            builder.store(res, po)

        return out

def range_iter_len(typingctx, val):
    """
    An implementation of len(range_iter) for internal use.
    """
    if isinstance(val, types.RangeIteratorType):
        val_type = val.yield_type
        def codegen(context, builder, sig, args):
            (value,) = args
            iter_type = range_impl_map[val_type][1]
            iterobj = cgutils.create_struct_proxy(iter_type)(context, builder, value)
            int_type = iterobj.count.type
            return impl_ret_untracked(context, builder, int_type, builder.load(iterobj.count))
        return signature(val_type, val), codegen
    elif isinstance(val, types.ListIter):
        def codegen(context, builder, sig, args):
            (value,) = args
            intp_t = context.get_value_type(types.intp)
            iterobj = ListIterInstance(context, builder, sig.args[0], value)
            return impl_ret_untracked(context, builder, intp_t, iterobj.size)
        return signature(types.intp, val), codegen
    elif isinstance(val, types.ArrayIterator):
        def  codegen(context, builder, sig, args):
            (iterty,) = sig.args
            (value,) = args
            intp_t = context.get_value_type(types.intp)
            iterobj = context.make_helper(builder, iterty, value=value)
            arrayty = iterty.array_type
            ary = make_array(arrayty)(context, builder, value=iterobj.array)
            shape = cgutils.unpack_tuple(builder, ary.shape)
            # array iterates along the outer dimension
            return impl_ret_untracked(context, builder, intp_t, shape[0])
        return signature(types.intp, val), codegen

    def get_attribute(self, val, typ, attr):
        if isinstance(typ, types.Module):
            # Implement getattr for module-level globals.
            # We are treating them as constants.
            # XXX We shouldn't have to retype this
            attrty = self.typing_context.resolve_module_constants(typ, attr)
            if attrty is not None and not isinstance(attrty, types.Dummy):
                pyval = getattr(typ.pymod, attr)
                llval = self.get_constant(attrty, pyval)
                @impl_attribute(typ, attr, attrty)
                def imp(context, builder, typ, val):
                    return impl_ret_borrowed(context, builder, attrty, llval)
                return imp
            # No implementation required for dummies (functions, modules...),
            # which are dealt with later
            return None

        # Lookup specific attribute implementation for this type
        overloads = self.attrs[attr]
        try:
            return overloads.find(typing.signature(types.Any, typ))
        except NotImplementedError:
            pass
        # Lookup generic getattr implementation for this type
        overloads = self.attrs[None]
        try:
            return overloads.find(typing.signature(types.Any, typ))
        except NotImplementedError:
            raise Exception("No definition for lowering %s.%s" % (typ, attr))

    def add(self, sig=None, argtypes=None, restype=None):
        # Handle argtypes
        if argtypes is not None:
            warnings.warn("Keyword argument argtypes is deprecated",
                          DeprecationWarning)
            assert sig is None
            if restype is None:
                sig = tuple(argtypes)
            else:
                sig = restype(*argtypes)
        del argtypes
        del restype

        # compile core as device function
        args, return_type = sigutils.normalize_signature(sig)
        devfnsig = signature(return_type, *args)

        funcname = self.pyfunc.__name__
        kernelsource = self._get_kernel_source(self._kernel_template,
                                               devfnsig, funcname)
        corefn, return_type = self._compile_core(devfnsig)
        glbl = self._get_globals(corefn)
        sig = signature(types.void, *([a[:] for a in args] + [return_type[:]]))
        _exec(kernelsource, glbl)

        stager = glbl['__vectorized_%s' % funcname]
        kernel = self._compile_kernel(stager, sig)

        argdtypes = tuple(to_dtype(t) for t in devfnsig.args)
        resdtype = to_dtype(return_type)
        self.kernelmap[tuple(argdtypes)] = resdtype, kernel

def array_median(context, builder, sig, args):
    def partition(A, low, high):
        mid = (low + high) // 2
        # median of three {low, middle, high}
        LM = A[low] <= A[mid]
        MH = A[mid] <= A[high]
        LH = A[low] <= A[high]

        if LM == MH:
            median3 = mid
        elif LH != LM:
            median3 = low
        else:
            median3 = high

        # choose median3 as the pivot
        A[high], A[median3] = A[median3], A[high]

        x = A[high]
        i = low
        for j in range(low, high):
            if A[j] <= x:
                A[i], A[j] = A[j], A[i]
                i += 1
        A[i], A[high] = A[high], A[i]
        return i

    sig_partition = typing.signature(types.intp, *(sig.args[0], types.intp, types.intp))
    _partition = context.compile_subroutine(builder, partition, sig_partition)

    def select(arry, k):
        n = arry.shape[0]
        # XXX: assuming flat array till array.flatten is implemented
        # temp_arry = arry.flatten()
        temp_arry = arry.copy()
        high = n - 1
        low = 0
        # NOTE: high is inclusive
        i = _partition(temp_arry, low, high)
        while i != k:
            if i < k:
                low = i + 1
                i = _partition(temp_arry, low, high)
            else:
                high = i - 1
                i = _partition(temp_arry, low, high)
        return temp_arry[k]

    sig_select = typing.signature(sig.args[0].dtype, *(sig.args[0], types.intp))
    _select = context.compile_subroutine(builder, select, sig_select)

    def median(arry):
        n = arry.shape[0]
        if n % 2 == 0:
            return (_select(arry, n // 2 - 1) + _select(arry, n // 2)) / 2
        else:
            return _select(arry, n // 2)

    res = context.compile_internal(builder, median, sig, args)
    return impl_ret_untracked(context, builder, sig.return_type, res)

 def test_equality(self):
     self.assertEqual(types.int32, types.int32)
     self.assertEqual(types.uint32, types.uint32)
     self.assertEqual(types.complex64, types.complex64)
     self.assertEqual(types.float32, types.float32)
     # Different signedness
     self.assertNotEqual(types.int32, types.uint32)
     # Different width
     self.assertNotEqual(types.int64, types.int32)
     self.assertNotEqual(types.float64, types.float32)
     self.assertNotEqual(types.complex64, types.complex128)
     # Different domain
     self.assertNotEqual(types.int64, types.float64)
     self.assertNotEqual(types.uint64, types.float64)
     self.assertNotEqual(types.complex64, types.float64)
     # Same arguments but different return types
     get_pointer = None
     sig_a = typing.signature(types.intp, types.intp)
     sig_b = typing.signature(types.voidptr, types.intp)
     a = types.ExternalFunctionPointer(sig=sig_a, get_pointer=get_pointer)
     b = types.ExternalFunctionPointer(sig=sig_b, get_pointer=get_pointer)
     self.assertNotEqual(a, b)
     # Different call convention
     a = types.ExternalFunctionPointer(sig=sig_a, get_pointer=get_pointer)
     b = types.ExternalFunctionPointer(sig=sig_a, get_pointer=get_pointer,
                                       cconv='stdcall')
     self.assertNotEqual(a, b)
     # Different get_pointer
     a = types.ExternalFunctionPointer(sig=sig_a, get_pointer=get_pointer)
     b = types.ExternalFunctionPointer(sig=sig_a, get_pointer=object())
     self.assertNotEqual(a, b)

    def inline_array(array_var, expr, stmts, list_vars, dels):
        """Check to see if the given "array_var" is created from a list
        of constants, and try to inline the list definition as array
        initialization.

        Extra statements produced with be appended to "stmts".
        """
        callname = guard(find_callname, func_ir, expr)
        require(callname and callname[1] == 'numpy' and callname[0] == 'array')
        require(expr.args[0].name in list_vars)
        ret_type = calltypes[expr].return_type
        require(isinstance(ret_type, types.ArrayCompatible) and
                           ret_type.ndim == 1)
        loc = expr.loc
        list_var = expr.args[0]
        array_typ = typemap[array_var.name]
        debug_print("inline array_var = ", array_var, " list_var = ", list_var)
        dtype = array_typ.dtype
        seq, op = find_build_sequence(func_ir, list_var)
        size = len(seq)
        size_var = ir.Var(scope, mk_unique_var("size"), loc)
        size_tuple_var = ir.Var(scope, mk_unique_var("size_tuple"), loc)
        size_typ = types.intp
        size_tuple_typ = types.UniTuple(size_typ, 1)

        typemap[size_var.name] = size_typ
        typemap[size_tuple_var.name] = size_tuple_typ

        stmts.append(_new_definition(func_ir, size_var,
                 ir.Const(size, loc=loc), loc))

        stmts.append(_new_definition(func_ir, size_tuple_var,
                 ir.Expr.build_tuple(items=[size_var], loc=loc), loc))

        empty_func = ir.Var(scope, mk_unique_var("empty_func"), loc)
        fnty = get_np_ufunc_typ(np.empty)
        sig = context.resolve_function_type(fnty, (size_typ,), {})
        typemap[empty_func.name] = fnty #

        stmts.append(_new_definition(func_ir, empty_func,
                         ir.Global('empty', np.empty, loc=loc), loc))

        empty_call = ir.Expr.call(empty_func, [size_var], {}, loc=loc)
        calltypes[empty_call] = typing.signature(array_typ, size_typ)
        stmts.append(_new_definition(func_ir, array_var, empty_call, loc))

        for i in range(size):
            index_var = ir.Var(scope, mk_unique_var("index"), loc)
            index_typ = types.intp
            typemap[index_var.name] = index_typ
            stmts.append(_new_definition(func_ir, index_var,
                    ir.Const(i, loc), loc))
            setitem = ir.SetItem(array_var, index_var, seq[i], loc)
            calltypes[setitem] = typing.signature(types.none, array_typ,
                                                  index_typ, dtype)
            stmts.append(setitem)

        stmts.extend(dels)
        return True

 def test_call_notation(self):
     # Function call signature
     i = types.int32
     d = types.double
     self.assertEqual(i(), typing.signature(i))
     self.assertEqual(i(d), typing.signature(i, d))
     self.assertEqual(i(d, d), typing.signature(i, d, d))
     # Value cast
     self.assertPreciseEqual(i(42.5), 42)
     self.assertPreciseEqual(d(-5), -5.0)

def local_array(shape, dtype):
    ndim = 1
    if isinstance(shape, tuple):
        ndim = len(shape)

    fname = "ptx.lmem.alloc"
    restype = types.Array(dtype, ndim, 'C')
    if ndim == 1:
        sig = typing.signature(restype, types.intp, types.Any)
    else:
        sig = typing.signature(restype, types.UniTuple(types.intp, ndim),
                               types.Any)

    return ir.Intrinsic(fname, sig, args=(shape, dtype))

    def get_attribute(self, val, typ, attr):
        if isinstance(typ, types.Record):
            # Implement get attribute for records
            self.sentry_record_alignment(typ, attr)
            offset = typ.offset(attr)
            elemty = typ.typeof(attr)

            @impl_attribute(typ, attr, elemty)
            def imp(context, builder, typ, val):
                dptr = cgutils.get_record_member(builder, val, offset, self.get_data_type(elemty))
                return self.unpack_value(builder, elemty, dptr)

            return imp

        if isinstance(typ, types.Module):
            # Implement getattr for module-level globals.
            # We are treating them as constants.
            # XXX We shouldn't have to retype this
            attrty = self.typing_context.resolve_module_constants(typ, attr)
            if attrty is not None:
                try:
                    pyval = getattr(typ.pymod, attr)
                    llval = self.get_constant(attrty, pyval)
                except NotImplementedError:
                    # Module attribute is not a simple constant
                    # (e.g. it's a function), it will be handled later on.
                    pass
                else:

                    @impl_attribute(typ, attr, attrty)
                    def imp(context, builder, typ, val):
                        return llval

                    return imp
            # No implementation
            return None

        # Lookup specific attribute implementation for this type
        overloads = self.attrs[attr]
        try:
            return overloads.find(typing.signature(types.Any, typ))
        except NotImplementedError:
            pass
        # Lookup generic getattr implementation for this type
        overloads = self.attrs[None]
        try:
            return overloads.find(typing.signature(types.Any, typ))
        except NotImplementedError:
            raise Exception("No definition for lowering %s.%s" % (typ, attr))

    def test_cache(self):
        def times2(i):
            return 2*i

        def times3(i):
            return i*3

        i32 = lc.Type.int(32)
        llvm_fnty = lc.Type.function(i32, [i32])
        module = lc.Module.new("test_module")
        function = module.get_or_insert_function(llvm_fnty, name='test_fn')
        assert function.is_declaration
        entry_block = function.append_basic_block('entry')
        builder = lc.Builder.new(entry_block)
        
        sig = typing.signature(types.int32, types.int32)
        typing_context = typing.Context()
        context = cpu.CPUContext(typing_context).localized()

        # Ensure the cache is empty to begin with
        self.assertEqual(0, len(context.cached_internal_func))
        
        # After one compile, it should contain one entry
        context.compile_internal(builder, times2, sig, function.args)
        self.assertEqual(1, len(context.cached_internal_func))

        # After a second compilation of the same thing, it should still contain
        # one entry
        context.compile_internal(builder, times2, sig, function.args)
        self.assertEqual(1, len(context.cached_internal_func))

        # After compilation of another function, the cache should have grown by
        # one more.
        context.compile_internal(builder, times3, sig, function.args)
        self.assertEqual(2, len(context.cached_internal_func))

        sig2 = typing.signature(types.float64, types.float64)
        f64 = lc.Type.double()
        llvm_fnty2 = lc.Type.function(f64, [f64])
        function2 = module.get_or_insert_function(llvm_fnty2, name='test_fn_2')
        assert function2.is_declaration
        entry_block2 = function2.append_basic_block('entry')
        builder2 = lc.Builder.new(entry_block2)
        
        # Ensure that the same function with a different signature does not
        # reuse an entry from the cache in error
        context.compile_internal(builder2, times3, sig2, function2.args)
        self.assertEqual(3, len(context.cached_internal_func))

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 dot_2_vv(context, builder, sig, args, conjugate=False):
    """
    np.dot(vector, vector)
    np.vdot(vector, vector)
    """
    aty, bty = sig.args
    dtype = sig.return_type
    a = make_array(aty)(context, builder, args[0])
    b = make_array(bty)(context, builder, args[1])
    n, = cgutils.unpack_tuple(builder, a.shape)

    def check_args(a, b):
        m, = a.shape
        n, = b.shape
        if m != n:
            raise ValueError("incompatible array sizes for np.dot(a, b) "
                             "(vector * vector)")

    context.compile_internal(builder, check_args,
                             signature(types.none, *sig.args), args)
    check_c_int(context, builder, n)

    out = cgutils.alloca_once(builder, context.get_value_type(dtype))
    call_xxdot(context, builder, conjugate, dtype, n, a.data, b.data, out)
    return builder.load(out)

    def _backend(self, lowerfn, objectmode):
        """
        Back-end: Generate LLVM IR from Numba IR, compile to machine code
        """
        if self.library is None:
            codegen = self.targetctx.codegen()
            self.library = codegen.create_library(self.bc.func_qualname)
            # Enable object caching upfront, so that the library can
            # be later serialized.
            self.library.enable_object_caching()

        lowered = lowerfn()
        signature = typing.signature(self.return_type, *self.args)
        self.cr = compile_result(typing_context=self.typingctx,
                                 target_context=self.targetctx,
                                 entry_point=lowered.cfunc,
                                 typing_error=self.status.fail_reason,
                                 type_annotation=self.type_annotation,
                                 library=self.library,
                                 call_helper=lowered.call_helper,
                                 signature=signature,
                                 objectmode=objectmode,
                                 interpmode=False,
                                 lifted=self.lifted,
                                 fndesc=lowered.fndesc,
                                 environment=lowered.env,
                                 has_dynamic_globals=lowered.has_dynamic_globals,
                                 )

def _gauss_impl(context, builder, sig, args, state):
    # The type for all computations (either float or double)
    ty = sig.return_type
    llty = context.get_data_type(ty)

    state_ptr = get_state_ptr(context, builder, state)
    _random = {"py": random.random,
               "np": np.random.random}[state]

    ret = cgutils.alloca_once(builder, llty, name="result")

    gauss_ptr = get_gauss_ptr(builder, state_ptr)
    has_gauss_ptr = get_has_gauss_ptr(builder, state_ptr)
    has_gauss = cgutils.is_true(builder, builder.load(has_gauss_ptr))
    with builder.if_else(has_gauss) as (then, otherwise):
        with then:
            # if has_gauss: return it
            builder.store(builder.load(gauss_ptr), ret)
            builder.store(const_int(0), has_gauss_ptr)
        with otherwise:
            # if not has_gauss: compute a pair of numbers using the Box-Muller
            # transform; keep one and return the other
            pair = context.compile_internal(builder,
                                            _gauss_pair_impl(_random),
                                            signature(types.UniTuple(ty, 2)),
                                            ())

            first, second = cgutils.unpack_tuple(builder, pair, 2)
            builder.store(first, gauss_ptr)
            builder.store(second, ret)
            builder.store(const_int(1), has_gauss_ptr)

    mu, sigma = args
    return builder.fadd(mu,
                        builder.fmul(sigma, builder.load(ret)))

 def implementer(context, builder, sig, args):
     val, = args
     input_type = sig.args[0]
     fpval = context.cast(builder, val, input_type, types.float64)
     inner_sig = signature(types.float64, types.float64)
     res = wrapped_impl(context, builder, inner_sig, (fpval,))
     return context.cast(builder, res, types.float64, sig.return_type)

def impl_setitem(d, key, value):
    if not isinstance(d, types.DictType):
        return

    keyty, valty = d.key_type, d.value_type

    def impl(d, key, value):
        castedkey = _cast(key, keyty)
        castedval = _cast(value, valty)
        status = _dict_insert(d, castedkey, hash(castedkey), castedval)
        if status == Status.OK:
            return
        elif status == Status.OK_REPLACED:
            # replaced
            # XXX handle refcount
            return
        elif status == Status.ERR_CMP_FAILED:
            raise ValueError('key comparison failed')
        else:
            raise RuntimeError('dict.__setitem__ failed unexpectedly')

    if d.is_precise():
        # Handle the precise case.
        return impl
    else:
        # Handle the imprecise case.
        d = d.refine(key, value)
        # Re-bind the key type and value type to match the arguments.
        keyty, valty = d.key_type, d.value_type
        # Create the signature that we wanted this impl to have.
        sig = typing.signature(types.void, d, keyty, valty)
        return sig, impl

def grid_expand(ndim):
    """grid(ndim)

    ndim: [int] 1, 2 or 3

        if ndim == 1:
            return cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x
        elif ndim == 2:
            x = cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x
            y = cuda.threadIdx.y + cuda.blockIdx.y * cuda.blockDim.y
            return x, y
        elif ndim == 3:
            x = cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x
            y = cuda.threadIdx.y + cuda.blockIdx.y * cuda.blockDim.y
            z = cuda.threadIdx.z + cuda.blockIdx.z * cuda.blockDim.z
            return x, y, z
    """
    if ndim == 1:
        fname = "ptx.grid.1d"
        restype = types.int32
    elif ndim == 2:
        fname = "ptx.grid.2d"
        restype = types.UniTuple(types.int32, 2)
    elif ndim == 3:
        fname = "ptx.grid.3d"
        restype = types.UniTuple(types.int32, 3)
    else:
        raise ValueError('argument can only be 1, 2, 3')

    return ir.Intrinsic(fname, typing.signature(restype, types.intp),
                        args=[ndim])

def local_array(shape, dtype):
    shape = _legalize_shape(shape)
    ndim = len(shape)
    fname = "ptx.lmem.alloc"
    restype = types.Array(dtype, ndim, 'C')
    sig = typing.signature(restype, types.UniTuple(types.intp, ndim), types.Any)
    return ir.Intrinsic(fname, sig, args=(shape, dtype))

    def test_normalize_signature(self):
        f = sigutils.normalize_signature

        def check(sig, args, return_type):
            self.assertEqual(f(sig), (args, return_type))

        def check_error(sig, msg):
            with self.assertRaises(TypeError) as raises:
                f(sig)
            self.assertIn(msg, str(raises.exception))

        f32 = types.float32
        c64 = types.complex64
        i16 = types.int16
        a = types.Array(f32, 1, "C")

        check((c64,), (c64,), None)
        check((f32, i16), (f32, i16), None)
        check(a(i16), (i16,), a)
        check("int16(complex64)", (c64,), i16)
        check("(complex64, int16)", (c64, i16), None)
        check(typing.signature(i16, c64), (c64,), i16)

        check_error((types.Integer,), "invalid signature")
        check_error((None,), "invalid signature")
        check_error([], "invalid signature")