code: def __call__(
      self,
      *bufs: cl.cl_mem,
      global_size: Tuple[int, ...],
      local_size: Optional[Tuple[int, ...]] = None,
      vals: Tuple[int, ...] = (),
      wait=False,
  ) -> Optional[float]:  # noqa: E501
      for i, b in enumerate(bufs):
          cl.clSetKernelArg(
              self.kernel, i, ctypes.sizeof(b), ctypes.byref(b)
          )
      for i, b in enumerate(vals, start=len(bufs)):
          cl.clSetKernelArg(
              self.kernel, i, 4, ctypes.byref(ctypes.c_int32(b))
          )
      if local_size is not None:
          global_size = tuple(
              int(g * l) for g, l in zip(global_size, local_size)
          )
      event = cl.cl_event() if wait else None
      check(
          cl.clEnqueueNDRangeKernel(
              self.device.queue,
              self.kernel,
              len(global_size),
              None,
              (ctypes.c_size_t * len(global_size))(*global_size),
              (ctypes.c_size_t * len(local_size))(*local_size)
              if local_size
              else None,
              0,
              None,
              event,
          )
      )  # noqa: E501
      if wait:
          check(cl.clWaitForEvents(1, ctypes.byref(event)))
          start = init_c_var(
              ctypes.c_uint64(),
              lambda x: check(
                  cl.clGetEventProfilingInfo(
                      event,
                      cl.CL_PROFILING_COMMAND_START,
                      ctypes.sizeof(x),
                      ctypes.byref(x),
                      None,
                  )
              ),
          )  # noqa: E501
          end = init_c_var(
              ctypes.c_uint64(),
              lambda x: check(
                  cl.clGetEventProfilingInfo(
                      event,
                      cl.CL_PROFILING_COMMAND_END,
                      ctypes.sizeof(x),
                      ctypes.byref(x),
                      None,
                  )
              ),
          )  # noqa: E501
          return (
              float(end.value - start.value) * OSX_TIMING_RATIO * 1e-9
          )
      return None
  
imports: from __future__ import annotations
  from typing import Tuple, Optional, List
  import ctypes, functools, hashlib
  import gpuctypes.opencl
  from tinygrad.helpers import init_c_var, to_char_p_p, from_mv, OSX, DEBUG
  from tinygrad.dtype import ImageDType
  from tinygrad.codegen.kernel import LinearizerOptions
  from tinygrad.renderer.cstyle import OpenCLRenderer
  from tinygrad.device import Compiled, LRUAllocator
label_desc: Use image node_id 3 for calling the CLProgram obj's underlying member method code with example usage: obj.__call__() and returns: None, unknown