code: def __call__(
      self,
      *bufs,
      global_size: Tuple[int, int, int],
      local_size: Tuple[int, int, int],
      vals: Tuple[int, ...] = (),
      wait=False,
  ):
      if not CUDACPU:
          check(cuda.cuCtxSetCurrent(self.device.context))
      c_kernel_input_config = (
          encode_args_cuda_style(
              bufs, vals, cuda.CUdeviceptr_v2, (1, 2, 0)
          )[0]
          if not CUDACPU
          else (bufs + tuple(vals))
      )
      return cu_time_execution(
          lambda: check(
              cuda.cuLaunchKernel(
                  self.prg,
                  *global_size,
                  *local_size,
                  0,
                  None,
                  None,
                  c_kernel_input_config,
              )
          ),
          enable=wait,
      )
  
imports: from __future__ import annotations
  import subprocess, hashlib, tempfile, ctypes, ctypes.util, functools, re
  from pathlib import Path
  from typing import Tuple, Optional
  import gpuctypes.cuda
  from tinygrad.helpers import DEBUG, getenv, from_mv, init_c_var, colored, cpu_time_execution, compile_cuda_style, encode_args_cuda_style, time_execution_cuda_style
  from tinygrad.device import Compiled, LRUAllocator, MallocAllocator
  from tinygrad.codegen.kernel import LinearizerOptions
  from tinygrad.renderer.cstyle import CUDARenderer
label_desc: Use image node_id 3 for calling the CUDAProgram obj's underlying member method code with example usage: obj.__call__() and returns: cu_time_execution