class_docstr: JIT CUDA kernel object.
  
  The decorator :func:``cupyx.jit.rawkernel`` converts the target function
  to an object of this class. This class is not inteded to be instantiated
  by users.
class_docstr_tok: ["JIT","CUDA","kernel","object",".","The","decorator",":","func",":","``","cupyx.jit.rawkernel","``","converts","the","target","function","to","an","object","of","this","class",".","This","class","is","not","inteded","to","be","instantiated","by","users","."]
docstr: Calls the CUDA kernel.
  
  The compilation will be deferred until the first function call.
  CuPy's JIT compiler infers the types of arguments at the call
  time, and will cache the compiled kernels for speeding up any
  subsequent calls.
  
  Args:
      grid (tuple of int): Size of grid in blocks.
      block (tuple of int): Dimensions of each thread block.
      args (tuple):
          Arguments of the kernel. The type of all elements must be
          ``bool``, ``int``, ``float``, ``complex``, NumPy scalar or
          ``cupy.ndarray``.
      shared_mem (int):
          Dynamic shared-memory size per thread block in bytes.
      stream (cupy.cuda.Stream): CUDA stream.
  
  .. seealso:: :ref:`jit_kernel_definition`
docstr_tok: ["Calls","the","CUDA","kernel",".","The","compilation","will","be","deferred","until","the","first","function","call",".","CuPy","'s","JIT","compiler","infers","the","types","of","arguments","at","the","call","time",",","and","will","cache","the","compiled","kernels","for","speeding","up","any","subsequent","calls",".","Args",":","grid","(","tuple","of","int",")",":","Size","of","grid","in","blocks",".","block","(","tuple","of","int",")",":","Dimensions","of","each","thread","block",".","args","(","tuple",")",":","Arguments","of","the","kernel",".","The","type","of","all","elements","must","be","``","bool","``",",","``","int","``",",","``","float","``",",","``","complex","``",",","NumPy","scalar","or","``","cupy.ndarray","``",".","shared_mem","(","int",")",":","Dynamic","shared-memory","size","per","thread","block","in","bytes",".","stream","(","cupy.cuda.Stream",")",":","CUDA","stream",".","..","seealso",":",":",":","ref",":","`","jit_kernel_definition","`"]
code: def __call__(self, grid, block, args, shared_mem=0, stream=None):
      """Calls the CUDA kernel.
  
      The compilation will be deferred until the first function call.
      CuPy's JIT compiler infers the types of arguments at the call
      time, and will cache the compiled kernels for speeding up any
      subsequent calls.
  
      Args:
          grid (tuple of int): Size of grid in blocks.
          block (tuple of int): Dimensions of each thread block.
          args (tuple):
              Arguments of the kernel. The type of all elements must be
              ``bool``, ``int``, ``float``, ``complex``, NumPy scalar or
              ``cupy.ndarray``.
          shared_mem (int):
              Dynamic shared-memory size per thread block in bytes.
          stream (cupy.cuda.Stream): CUDA stream.
  
      .. seealso:: :ref:`jit_kernel_definition`
      """
      in_types = []
      for x in args:
          if isinstance(x, cupy.ndarray):
              t = _cuda_types.CArray.from_ndarray(x)
          elif numpy.isscalar(x):
              t = _cuda_typerules.get_ctype_from_scalar(self._mode, x)
          else:
              raise TypeError(
                  f"{type(x)} is not supported for RawKernel"
              )
          in_types.append(t)
      in_types = tuple(in_types)
      device_id = cupy.cuda.get_device_id()
  
      kern, enable_cg = self._cache.get(
          (in_types, device_id), (None, None)
      )
      if kern is None:
          result = self._cached_codes.get(in_types)
          if result is None:
              result = _compile.transpile(
                  self._func,
                  ['extern "C"', "__global__"],
                  self._mode,
                  in_types,
                  _cuda_types.void,
              )
              self._cached_codes[in_types] = result
  
          fname = result.func_name
          enable_cg = result.enable_cooperative_groups
          options = result.options
          backend = result.backend
          if backend == "nvcc":
              options += ("-DCUPY_JIT_NVCC",)
          jitify = result.jitify
          module = core.compile_with_cache(
              source=result.code,
              options=options,
              backend=backend,
              jitify=jitify,
          )
          kern = module.get_function(fname)
          self._cache[(in_types, device_id)] = (kern, enable_cg)
  
      new_args = []
      for a, t in zip(args, in_types):
          if isinstance(t, Scalar):
              if t.dtype.char == "e":
                  a = numpy.float32(a)
              else:
                  a = t.dtype.type(a)
          new_args.append(a)
  
      kern(grid, block, tuple(new_args), shared_mem, stream, enable_cg)
  
imports: import functools
  import warnings
  import numpy
  import cupy
  from cupy._core import core
  from cupyx.jit import _compile
  from cupyx.jit import _cuda_typerules
  from cupyx.jit import _cuda_types
  from cupyx.jit import _internal_types
  from cupyx.jit._cuda_types import Scalar
label_desc: Use image node_id 2 for calling the _JitRawKernel obj's underlying member method code with example usage: obj.__call__(grid, block, args, shared_mem, stream) without return types