code: def buf_cache_hierarchy_pchase(ndata, stride=1, NCOMP=1, steps=65536):
      ndata //= NCOMP * 4  # ptr size
      prg = f"""__kernel void buf_cache_hierarchy_pchase(
      __global int{str(NCOMP) if NCOMP > 1 else ''}* src,
      __global int* dst,
      const int niter
    ) {{
      int idx = 0;
      for (int i = 0; i < niter; ++i) {{
        idx = src[idx]{'.x' if NCOMP > 1 else ''};
      }}
      *dst = idx;
    }}"""
      idx_buf = np.zeros(ndata * NCOMP, dtype=np.int32)
      for i in range(ndata):
          idx_buf[i * NCOMP] = (i + stride) % ndata
      in_buf = CLBuffer.fromCPU(idx_buf)
      out_buf = CLBuffer(1, dtypes.int32)
      cl = CLProgram(
          "buf_cache_hierarchy_pchase",
          prg,
          argdtypes=[None, None, np.int32],
      )
      return (
          min(
              [
                  cl(
                      [1, 1, 1],
                      [1, 1, 1],
                      in_buf,
                      out_buf,
                      steps,
                      wait=True,
                  )
                  / steps
                  for _ in range(5)
              ]
          )
          * 1e9
      )
  
file_path: tinygrad/extra/archprobe.py
imports: import numpy
  import pickle
  from tinygrad.runtime.ops_gpu import CLProgram, CLBuffer
  from tinygrad import dtypes
  from tqdm import trange, tqdm
  from matplotlib import pyplot
label_desc: Use image node_id 6 for calling a global function with example usage: buf_cache_hierarchy_pchase(ndata, stride, NCOMP, steps) and returns: unknown