code: def __init__(
      self,
      dtype_out,
      neutral,
      reduce_expr,
      map_expr=None,
      arguments=None,
      name="reduce_kernel",
      keep=False,
      options=None,
      preamble="",
  ):
      self.dtype_out = np.dtype(dtype_out)
  
      self.block_size = 512
  
      s1_func, self.stage1_arg_types = get_reduction_kernel_and_types(
          1,
          dtype_to_ctype(dtype_out),
          self.block_size,
          neutral,
          reduce_expr,
          map_expr,
          arguments,
          name=name + "_stage1",
          keep=keep,
          options=options,
          preamble=preamble,
      )
      self.stage1_func = s1_func.prepared_async_call
  
      # stage 2 has only one input and no map expression
      s2_func, self.stage2_arg_types = get_reduction_kernel_and_types(
          2,
          dtype_to_ctype(dtype_out),
          self.block_size,
          neutral,
          reduce_expr,
          arguments=arguments,
          name=name + "_stage2",
          keep=keep,
          options=options,
          preamble=preamble,
      )
      self.stage2_func = s2_func.prepared_async_call
  
      assert [
          i
          for i, arg_tp in enumerate(self.stage1_arg_types)
          if arg_tp == "P"
      ], (
          "ReductionKernel can only be used with functions that have at least one "
          "vector argument"
      )
  
file_path: pycuda/pycuda/reduction.py
imports: from pycuda.tools import context_dependent_memoize
  from pycuda.tools import dtype_to_ctype
  import numpy
label_desc: Use image node_id 1 to create a new ReductionKernel object  with example: obj = ReductionKernel(dtype_out, neutral, reduce_expr, map_expr, arguments, name, keep, options, preamble)