code: def gpu_aamp(
      T_A, m, T_B=None, ignore_trivial=True, device_id=0, p=2.0, k=1
  ):
      """
      Compute the non-normalized (i.e., without z-normalization) matrix profile with
      one or more GPU devices
  
      This is a convenience wrapper around the Numba `cuda.jit` `_gpu_aamp` function
      which computes the non-normalized (top-k) matrix profile according to modified
      version 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 contain your query subsequences
          of interest. 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()]`.
  
      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.
  
      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.
  
      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.
  
      Notes
      -----
      `arXiv:1901.05708 \
      <https://arxiv.org/pdf/1901.05708.pdf>`__
  
      See Algorithm 1
  
      Note that we have extended this algorithm for AB-joins as well.
  
      `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
      """
      if T_B is None:  # Self join!
          T_B = T_A
          ignore_trivial = True
  
      T_A, T_A_subseq_isfinite = core.preprocess_non_normalized(T_A, m)
      T_B, T_B_subseq_isfinite = core.preprocess_non_normalized(T_B, m)
  
      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)
      T_A_subseq_isfinite_fname = core.array_to_temp_file(
          T_A_subseq_isfinite
      )
      T_B_subseq_isfinite_fname = core.array_to_temp_file(
          T_B_subseq_isfinite
      )
  
      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)
  
      p_norm_fnames = []
      p_norm_first_fnames = []
  
      for idx, start in enumerate(range(0, l, step)):
          stop = min(l, start + step)
  
          p_norm = np.power(
              core.mass_absolute(T_B[start : start + m], T_A, p=p), p
          )
          p_norm_first = np.power(
              core.mass_absolute(T_A[:m], T_B, p=p), p
          )
  
          p_norm_fname = core.array_to_temp_file(p_norm)
          p_norm_first_fname = core.array_to_temp_file(p_norm_first)
          p_norm_fnames.append(p_norm_fname)
          p_norm_first_fnames.append(p_norm_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_aamp,
                  (
                      T_A_fname,
                      T_B_fname,
                      m,
                      stop,
                      excl_zone,
                      T_A_subseq_isfinite_fname,
                      T_B_subseq_isfinite_fname,
                      p,
                      p_norm_fname,
                      p_norm_first_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_aamp(
                  T_A_fname,
                  T_B_fname,
                  m,
                  stop,
                  excl_zone,
                  T_A_subseq_isfinite_fname,
                  T_B_subseq_isfinite_fname,
                  p,
                  p_norm_fname,
                  p_norm_first_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(T_A_subseq_isfinite_fname)
      os.remove(T_B_subseq_isfinite_fname)
      for p_norm_fname in p_norm_fnames:
          os.remove(p_norm_fname)
      for p_norm_first_fname in p_norm_first_fnames:
          os.remove(p_norm_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_aamp.py
imports: import math
  import multiprocessing
  import os
  import numpy
  from numba import cuda
  from .None import config, core
label_desc: Use image node_id 3 for calling a global function with example usage: gpu_aamp(T_A, m, T_B, ignore_trivial, device_id, p, k) and returns: out