code: def reg_count(nthread, ngrp, nreg):
      reg_declr = "".join(
          [
              f"float reg_data{i} = (float)niter + {i};\n"
              for i in range(nreg)
          ]
      )
      reg_comp = "".join(
          [f"reg_data{i} *= {(i-1)%nreg};\n" for i in range(nreg)]
      )
      reg_reduce = "".join(
          [f"out_buf[{i}] = reg_data{i};\n" for i in range(nreg)]
      )
      prg = f"""__kernel void reg_count(
      __global float* out_buf,
      __private const int niter
    ) {{
      {reg_declr}
      int i = 0;
      for (; i < niter; ++i) {{
        {reg_comp}
      }}
      i = i >> 31;
      {reg_reduce}
    }}"""
      out_buf = CLBuffer(1, dtypes.float32)
      cl = CLProgram("reg_count", prg, argdtypes=[None, np.int32])
      return (
          min(
              [
                  cl(
                      [nthread, ngrp, 1],
                      [nthread, 1, 1],
                      out_buf,
                      20,
                      wait=True,
                  )
                  for _ in range(10)
              ]
          )
          * 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 4 for calling a global function with example usage: reg_count(nthread, ngrp, nreg) and returns: unknown