Is your feature request related to a problem? Please describe.

I want to have a function to tile array, something like that:

from numba import cuda
from numba.extending import overload
import numpy as np


def slice_array(a, x_id, x_size, y_id, y_size):
    pass


@overload(slice_array, inline="always", target="cuda")
def ol_slice_array(a, x_id, x_size, y_id, y_size):
    def impl(a, x_id, x_size, y_id, y_size):
        return a[
            x_id * x_size : (x_id + 1) * x_size : 1,
            y_id * y_size : (y_id + 1) * y_size : 1,
        ]

    return impl


@cuda.jit
def kernel(a):
    a_slice = slice_array(a, 1, 2, 2, 2)
    a_slice[0, 0] = 1


a = np.zeros((4, 6), dtype=np.int32)

kernel[1, 1](a)

assert a[2, 4] == 1

Describe the solution you'd like

Above code is working)

Describe alternatives you've considered

If I don't wrap tiling in a function it works. However I have to use tiling in my code more than ones and it increases chance of error.

Another option is to create intrinsic with llvm codegen, but it feels a bit overkill for the trivial task...

Additional context

I guess it comes from numa's type check and the idea there is to prevent returning array from nopython to python, however this restriction is also applied to functions that work only inside nopython...