code: def _kernel_connect():
      return _core.ElementwiseKernel(
          "raw int32 shape, raw int32 dirs, int32 ndirs, int32 ndim",
          "raw Y y",
          """
          if (y[i] < 0) continue;
          for (int dr = 0; dr < ndirs; dr++) {
              int j = i;
              int rest = j;
              int stride = 1;
              int k = 0;
              for (int dm = ndim-1; dm >= 0; dm--) {
                  int pos = rest % shape[dm] + dirs[dm + dr * ndim];
                  if (pos < 0 || pos >= shape[dm]) {
                      k = -1;
                      break;
                  }
                  k += pos * stride;
                  rest /= shape[dm];
                  stride *= shape[dm];
              }
              if (k < 0) continue;
              if (y[k] < 0) continue;
              while (1) {
                  while (j != y[j]) { j = y[j]; }
                  while (k != y[k]) { k = y[k]; }
                  if (j == k) break;
                  if (j < k) {
                      int old = atomicCAS( &y[k], k, j );
                      if (old == k) break;
                      k = old;
                  }
                  else {
                      int old = atomicCAS( &y[j], j, k );
                      if (old == j) break;
                      j = old;
                  }
              }
          }
          """,
          "cupyx_scipy_ndimage_label_connect",
      )
  
imports: import warnings
  import numpy
  import cupy
  from cupy import _core
  from cupy import _util
label_desc: Use image node_id 5 for calling a global function with example usage: _kernel_connect() and returns: _core