def _assign_kernel(ndim):
    """
    A separate method so we don't need to compile code every assignment (!).

    :param ndim: We need to have static array sizes for cuda.local.array, so
        bake in the number of dimensions into the kernel
    """
    from numba import cuda  # circular!

    if ndim == 0:
        # the (2, ndim) allocation below is not yet supported, so avoid it
        @cuda.jit
        def kernel(lhs, rhs):
            lhs[()] = rhs[()]
        return kernel

    @cuda.jit
    def kernel(lhs, rhs):
        location = cuda.grid(1)

        n_elements = 1
        for i in range(lhs.ndim):
            n_elements *= lhs.shape[i]
        if location >= n_elements:
            # bake n_elements into the kernel, better than passing it in
            # as another argument.
            return

        # [0, :] is the to-index (into `lhs`)
        # [1, :] is the from-index (into `rhs`)
        idx = cuda.local.array(
            shape=(2, ndim),
            dtype=types.int64)

        for i in range(ndim - 1, -1, -1):
            idx[0, i] = location % lhs.shape[i]
            idx[1, i] = (location % lhs.shape[i]) * (rhs.shape[i] > 1)
            location //= lhs.shape[i]

        lhs[to_fixed_tuple(idx[0], ndim)] = rhs[to_fixed_tuple(idx[1], ndim)]
    return kernel

    A separate method so we don't need to compile code every assignment (!).

    :param ndim: We need to have static array sizes for cuda.local.array, so
        bake in the number of dimensions into the kernel
    
numba/cuda/cudadrv/devicearray.py