code: def warp_size2(nthread):
      prg = """__kernel void warp_size2(
      __global float* src,
      __global int* dst,
      const int niter,
      const int prime_number
    ) {
      int drain = 0;
      for (int j = 0; j < niter; ++j) {
        drain += j / prime_number;
        barrier(0);
      }
      dst[get_local_id(0)] = drain;
    }"""
      src_buf = CLBuffer(1, dtypes.float32)
      dst_buf = CLBuffer(1, dtypes.int32)
      cl = CLProgram(
          "warp_size2", prg, argdtypes=[None, None, np.int32, np.int32]
      )
      return (
          min(
              [
                  cl(
                      [nthread, 1024, 1],
                      [nthread, 1, 1],
                      src_buf,
                      dst_buf,
                      10,
                      3,
                      wait=True,
                  )
                  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 2 for calling a global function with example usage: warp_size2(nthread) and returns: unknown