Use int64 for indexing

For indexed kernels, int32 is too small for large domain sizes. Thus the coordinates are cast to int64 in this MR to allow huge domain sizes.

As an example of the adaption the generated code for a Neumann boundary is shown. Before:

FUNC_PREFIX void kernel(double * RESTRICT _data_C, uint8_t * RESTRICT const _data_indexField, int64_t const _size_indexField_0, int64_t const _stride_indexField_0)
{
   #pragma omp parallel
   {
      #pragma omp for schedule(static)
      for (int64_t ctr_0 = 0; ctr_0 < _size_indexField_0; ctr_0 += 1)
      {
         const int32_t x = *((int32_t *)(& _data_indexField[12*_stride_indexField_0*ctr_0]));
         const int32_t y = *((int32_t *)(& _data_indexField[12*_stride_indexField_0*ctr_0 + 4]));
         
         
         const int64_t cx [] = { 0, 0, 0, -1, 1, -1, 1, -1, 1 };
         const int64_t cy [] = { 0, 1, -1, 0, 0, 1, 1, -1, -1 };
         const int invdir [] = { 0, 2, 1, 4, 3, 8, 7, 6, 5 };
         
         const int32_t dir = *((int32_t *)(& _data_indexField[12*_stride_indexField_0*ctr_0 + 8]));
         _data_C[x + 11*y] = _data_C[x + 11*y + cx[dir] + 11*cy[dir]];
      }
   }
}

After:

FUNC_PREFIX void kernel(double * RESTRICT _data_C, uint8_t * RESTRICT const _data_indexField, int64_t const _size_indexField_0, int64_t const _stride_indexField_0)
{
   #pragma omp parallel
   {
      #pragma omp for schedule(static)
      for (int64_t ctr_0 = 0; ctr_0 < _size_indexField_0; ctr_0 += 1)
      {
         const int64_t x = *((int32_t *)(& _data_indexField[12*_stride_indexField_0*ctr_0]));
         const int64_t y = *((int32_t *)(& _data_indexField[12*_stride_indexField_0*ctr_0 + 4]));
         
         
         const int64_t cx [] = { 0, 0, 0, -1, 1, -1, 1, -1, 1 };
         const int64_t cy [] = { 0, 1, -1, 0, 0, 1, 1, -1, -1 };
         const int64_t invdir [] = { 0, 2, 1, 4, 3, 8, 7, 6, 5 };
         
         const int64_t dir = *((int32_t *)(& _data_indexField[12*_stride_indexField_0*ctr_0 + 8]));
         _data_C[x + 11*y] = _data_C[x + 11*y + cx[dir] + 11*cy[dir]];
      }
   }
}

pystencils/boundaries/boundaryhandling.py

@@ -444,7 +444,6 @@ class BoundaryOffsetInfo(CustomCodeNode):
 
 def create_boundary_kernel(field, index_field, stencil, boundary_functor, target='cpu', **kernel_creation_args):
     elements = [BoundaryOffsetInfo(stencil)]
-    index_arr_dtype = index_field.dtype.numpy_dtype
     dir_symbol = TypedSymbol("dir", np.int64)
     elements += [Assignment(dir_symbol, index_field[0]('dir'))]
     elements += boundary_functor(field, direction_symbol=dir_symbol, index_field=index_field)