code: def gpu_stump(
      T_A,
      m,
      T_B=None,
      ignore_trivial=True,
      device_id=0,
      normalize=True,
      p=2.0,
      k=1,
      T_A_subseq_isconstant=None,
      T_B_subseq_isconstant=None,
  ):
      """
      Compute the z-normalized matrix profile with one or more GPU devices
  
      This is a convenience wrapper around the Numba `cuda.jit` `_gpu_stump` function
      which computes the matrix profile according to GPU-STOMP. The default number of
      threads-per-block is set to `512` and may be changed by setting the global parameter
      `config.STUMPY_THREADS_PER_BLOCK` to an appropriate number based on your GPU
      hardware.
  
      Parameters
      ----------
      T_A : numpy.ndarray
          The time series or sequence for which to compute the matrix profile
  
      m : int
          Window size
  
      T_B : numpy.ndarray, default None
          The time series or sequence that will be used to annotate T_A. For every
          subsequence in T_A, its nearest neighbor in T_B will be recorded. Default is
          `None` which corresponds to a self-join.
  
      ignore_trivial : bool, default True
          Set to `True` if this is a self-join. Otherwise, for AB-join, set this
          to `False`. Default is `True`.
  
      device_id : int or list, default 0
          The (GPU) device number to use. The default value is `0`. A list of
          valid device ids (int) may also be provided for parallel GPU-STUMP
          computation. A list of all valid device ids can be obtained by
          executing `[device.id for device in numba.cuda.list_devices()]`.
  
      normalize : bool, default True
          When set to `True`, this z-normalizes subsequences prior to computing distances.
          Otherwise, this function gets re-routed to its complementary non-normalized
          equivalent set in the `@core.non_normalized` function decorator.
  
      p : float, default 2.0
          The p-norm to apply for computing the Minkowski distance. Minkowski distance is
          typically used with `p` being 1 or 2, which correspond to the Manhattan distance
          and the Euclidean distance, respectively. This parameter is ignored when
          `normalize == True`.
  
      k : int, default 1
          The number of top `k` smallest distances used to construct the matrix profile.
          Note that this will increase the total computational time and memory usage
          when k > 1.
  
      T_A_subseq_isconstant : numpy.ndarray or function, default None
          A boolean array that indicates whether a subsequence in `T_A` is constant
          (True). Alternatively, a custom, user-defined function that returns a
          boolean array that indicates whether a subsequence in `T_A` is constant
          (True). The function must only take two arguments, `a`, a 1-D array,
          and `w`, the window size, while additional arguments may be specified
          by currying the user-defined function using `functools.partial`. Any
          subsequence with at least one np.nan/np.inf will automatically have its
          corresponding value set to False in this boolean array.
  
      T_B_subseq_isconstant : numpy.ndarray or function, default None
          A boolean array that indicates whether a subsequence in `T_B` is constant
          (True). Alternatively, a custom, user-defined function that returns a
          boolean array that indicates whether a subsequence in `T_B` is constant
          (True). The function must only take two arguments, `a`, a 1-D array,
          and `w`, the window size, while additional arguments may be specified
          by currying the user-defined function using `functools.partial`. Any
          subsequence with at least one np.nan/np.inf will automatically have its
          corresponding value set to False in this boolean array.
  
      Returns
      -------
      out : numpy.ndarray
          When k = 1 (default), the first column consists of the matrix profile,
          the second column consists of the matrix profile indices, the third column
          consists of the left matrix profile indices, and the fourth column consists
          of the right matrix profile indices. However, when k > 1, the output array
          will contain exactly 2 * k + 2 columns. The first k columns (i.e., out[:, :k])
          consists of the top-k matrix profile, the next set of k columns
          (i.e., out[:, k:2k]) consists of the corresponding top-k matrix profile
          indices, and the last two columns (i.e., out[:, 2k] and out[:, 2k+1] or,
          equivalently, out[:, -2] and out[:, -1]) correspond to the top-1 left
          matrix profile indices and the top-1 right matrix profile indices, respectively.
  
      See Also
      --------
      stumpy.stump : Compute the z-normalized matrix profile
      stumpy.stumped : Compute the z-normalized matrix profile with a distributed dask
          cluster
      stumpy.scrump : Compute an approximate z-normalized matrix profile
  
      Notes
      -----
      `DOI: 10.1109/ICDM.2016.0085 \
      <https://www.cs.ucr.edu/~eamonn/STOMP_GPU_final_submission_camera_ready.pdf>`__
  
      See Table II, Figure 5, and Figure 6
  
      Timeseries, T_A, will be annotated with the distance location
      (or index) of all its subsequences in another times series, T_B.
  
      Return: For every subsequence, Q, in T_A, you will get a distance
      and index for the closest subsequence in T_B. Thus, the array
      returned will have length T_A.shape[0]-m+1. Additionally, the
      left and right matrix profiles are also returned.
  
      Note: Unlike in the Table II where T_A.shape is expected to be equal
      to T_B.shape, this implementation is generalized so that the shapes of
      T_A and T_B can be different. In the case where T_A.shape == T_B.shape,
      then our algorithm reduces down to the same algorithm found in Table II.
  
      Additionally, unlike STAMP where the exclusion zone is m/2, the default
      exclusion zone for STOMP is m/4 (See Definition 3 and Figure 3).
  
      For self-joins, set `ignore_trivial = True` in order to avoid the
      trivial match.
  
      Note that left and right matrix profiles are only available for self-joins.
  
      Examples
      --------
      >>> import stumpy
      >>> import numpy as np
      >>> from numba import cuda
      >>> if __name__ == "__main__":
      ...     all_gpu_devices = [device.id for device in cuda.list_devices()]
      ...     stumpy.gpu_stump(
      ...         np.array([584., -11., 23., 79., 1001., 0., -19.]),
      ...         m=3,
      ...         device_id=all_gpu_devices)
      array([[0.11633857113691416, 4, -1, 4],
             [2.694073918063438, 3, -1, 3],
             [3.0000926340485923, 0, 0, 4],
             [2.694073918063438, 1, 1, -1],
             [0.11633857113691416, 0, 0, -1]], dtype=object)
      """
      if T_B is None:  # Self join!
          T_B = T_A
          ignore_trivial = True
          T_B_subseq_isconstant = T_A_subseq_isconstant
  
      T_A, μ_Q, σ_Q, Q_subseq_isconstant = core.preprocess(
          T_A, m, T_subseq_isconstant=T_A_subseq_isconstant
      )
      T_B, M_T, Σ_T, T_subseq_isconstant = core.preprocess(
          T_B, m, T_subseq_isconstant=T_B_subseq_isconstant
      )
  
      if T_A.ndim != 1:  # pragma: no cover
          raise ValueError(
              f"T_A is {T_A.ndim}-dimensional and must be 1-dimensional. "
              "For multidimensional STUMP use `stumpy.mstump` or `stumpy.mstumped`"
          )
  
      if T_B.ndim != 1:  # pragma: no cover
          raise ValueError(
              f"T_B is {T_B.ndim}-dimensional and must be 1-dimensional. "
              "For multidimensional STUMP use `stumpy.mstump` or `stumpy.mstumped`"
          )
  
      core.check_window_size(
          m, max_size=min(T_A.shape[0], T_B.shape[0])
      )
      ignore_trivial = core.check_ignore_trivial(
          T_A, T_B, ignore_trivial
      )
  
      n = T_B.shape[0]
      w = T_A.shape[0] - m + 1
      l = n - m + 1
      excl_zone = int(
          np.ceil(m / config.STUMPY_EXCL_ZONE_DENOM)
      )  # See Definition 3 and Figure 3
  
      T_A_fname = core.array_to_temp_file(T_A)
      T_B_fname = core.array_to_temp_file(T_B)
      μ_Q_fname = core.array_to_temp_file(μ_Q)
      σ_Q_fname = core.array_to_temp_file(σ_Q)
      M_T_fname = core.array_to_temp_file(M_T)
      Σ_T_fname = core.array_to_temp_file(Σ_T)
      Q_subseq_isconstant_fname = core.array_to_temp_file(
          Q_subseq_isconstant
      )
      T_subseq_isconstant_fname = core.array_to_temp_file(
          T_subseq_isconstant
      )
  
      if isinstance(device_id, int):
          device_ids = [device_id]
      else:
          device_ids = device_id
  
      profile = [None] * len(device_ids)
      indices = [None] * len(device_ids)
  
      profile_L = [None] * len(device_ids)
      indices_L = [None] * len(device_ids)
  
      profile_R = [None] * len(device_ids)
      indices_R = [None] * len(device_ids)
  
      for _id in device_ids:
          with cuda.gpus[_id]:
              if (
                  cuda.current_context().__class__.__name__
                  != "FakeCUDAContext"
              ):  # pragma: no cover
                  cuda.current_context().deallocations.clear()
  
      step = 1 + l // len(device_ids)
  
      # Start process pool for multi-GPU request
      if len(device_ids) > 1:  # pragma: no cover
          mp.set_start_method("spawn", force=True)
          pool = mp.Pool(processes=len(device_ids))
          results = [None] * len(device_ids)
  
      QT_fnames = []
      QT_first_fnames = []
  
      for idx, start in enumerate(range(0, l, step)):
          stop = min(l, start + step)
  
          QT, QT_first = core._get_QT(start, T_A, T_B, m)
          QT_fname = core.array_to_temp_file(QT)
          QT_first_fname = core.array_to_temp_file(QT_first)
          QT_fnames.append(QT_fname)
          QT_first_fnames.append(QT_first_fname)
  
          if (
              len(device_ids) > 1 and idx < len(device_ids) - 1
          ):  # pragma: no cover
              # Spawn and execute in child process for multi-GPU request
              results[idx] = pool.apply_async(
                  _gpu_stump,
                  (
                      T_A_fname,
                      T_B_fname,
                      m,
                      stop,
                      excl_zone,
                      μ_Q_fname,
                      σ_Q_fname,
                      QT_fname,
                      QT_first_fname,
                      M_T_fname,
                      Σ_T_fname,
                      Q_subseq_isconstant_fname,
                      T_subseq_isconstant_fname,
                      w,
                      ignore_trivial,
                      start + 1,
                      device_ids[idx],
                      k,
                  ),
              )
          else:
              # Execute last chunk in parent process
              # Only parent process is executed when a single GPU is requested
              (
                  profile[idx],
                  profile_L[idx],
                  profile_R[idx],
                  indices[idx],
                  indices_L[idx],
                  indices_R[idx],
              ) = _gpu_stump(
                  T_A_fname,
                  T_B_fname,
                  m,
                  stop,
                  excl_zone,
                  μ_Q_fname,
                  σ_Q_fname,
                  QT_fname,
                  QT_first_fname,
                  M_T_fname,
                  Σ_T_fname,
                  Q_subseq_isconstant_fname,
                  T_subseq_isconstant_fname,
                  w,
                  ignore_trivial,
                  start + 1,
                  device_ids[idx],
                  k,
              )
  
      # Clean up process pool for multi-GPU request
      if len(device_ids) > 1:  # pragma: no cover
          pool.close()
          pool.join()
  
          # Collect results from spawned child processes if they exist
          for idx, result in enumerate(results):
              if result is not None:
                  (
                      profile[idx],
                      profile_L[idx],
                      profile_R[idx],
                      indices[idx],
                      indices_L[idx],
                      indices_R[idx],
                  ) = result.get()
  
      os.remove(T_A_fname)
      os.remove(T_B_fname)
      os.remove(μ_Q_fname)
      os.remove(σ_Q_fname)
      os.remove(M_T_fname)
      os.remove(Σ_T_fname)
      os.remove(Q_subseq_isconstant_fname)
      os.remove(T_subseq_isconstant_fname)
      for QT_fname in QT_fnames:
          os.remove(QT_fname)
      for QT_first_fname in QT_first_fnames:
          os.remove(QT_first_fname)
  
      for idx in range(len(device_ids)):
          profile_fname = profile[idx]
          profile_L_fname = profile_L[idx]
          profile_R_fname = profile_R[idx]
          indices_fname = indices[idx]
          indices_L_fname = indices_L[idx]
          indices_R_fname = indices_R[idx]
  
          profile[idx] = np.load(profile_fname, allow_pickle=False)
          profile_L[idx] = np.load(profile_L_fname, allow_pickle=False)
          profile_R[idx] = np.load(profile_R_fname, allow_pickle=False)
          indices[idx] = np.load(indices_fname, allow_pickle=False)
          indices_L[idx] = np.load(indices_L_fname, allow_pickle=False)
          indices_R[idx] = np.load(indices_R_fname, allow_pickle=False)
  
          os.remove(profile_fname)
          os.remove(profile_L_fname)
          os.remove(profile_R_fname)
          os.remove(indices_fname)
          os.remove(indices_L_fname)
          os.remove(indices_R_fname)
  
      for i in range(1, len(device_ids)):  # pragma: no cover
          # Update (top-k) matrix profile and matrix profile indices
          core._merge_topk_PI(
              profile[0], profile[i], indices[0], indices[i]
          )
  
          # Update (top-1) left matrix profile and matrix profile indices
          mask = profile_L[0] > profile_L[i]
          profile_L[0][mask] = profile_L[i][mask]
          indices_L[0][mask] = indices_L[i][mask]
  
          # Update (top-1) right matrix profile and matrix profile indices
          mask = profile_R[0] > profile_R[i]
          profile_R[0][mask] = profile_R[i][mask]
          indices_R[0][mask] = indices_R[i][mask]
  
      out = np.empty(
          (w, 2 * k + 2), dtype=object
      )  # last two columns are to store
      # (top-1) left/right matrix profile indices
      out[:, :k] = profile[0]
      out[:, k:] = np.column_stack(
          (indices[0], indices_L[0], indices_R[0])
      )
  
      core._check_P(out[:, 0])
  
      return out
  
file_path: stumpy/stumpy/gpu_stump.py
imports: import math
  import multiprocessing
  import os
  import numpy
  from numba import cuda
  from .None import config, core
  from .gpu_aamp import gpu_aamp
label_desc: Use image node_id 3 for calling a global function with example usage: gpu_stump(T_A, m, T_B, ignore_trivial, device_id, normalize, p, k, T_A_subseq_isconstant, T_B_subseq_isconstant) and returns: out