class_docstr: A kernel that performs a generic reduction on arrays.
  
  Generate a kernel that takes a number of scalar or vector *arguments*
  (at least one vector argument), performs the *map_expr* on each entry of
  the vector argument and then the *reduce_expr* on the outcome of that.
  *neutral* serves as an initial value. *preamble* offers the possibility
  to add preprocessor directives and other code (such as helper functions)
  to be added before the actual reduction kernel code.
  
  Vectors in *map_expr* should be indexed by the variable *i*. *reduce_expr*
  uses the formal values "a" and "b" to indicate two operands of a binary
  reduction operation. If you do not specify a *map_expr*, ``in[i]`` is
  automatically assumed and treated as the only one input argument.
  
  *dtype_out* specifies the :class:`numpy.dtype` in which the reduction is
  performed and in which the result is returned. *neutral* is specified as
  float or integer formatted as string. *reduce_expr* and *map_expr* are
  specified as string formatted operations and *arguments* is specified as a
  string formatted as a C argument list. *name* specifies the name as which
  the kernel is compiled. *options* are passed unmodified to
  :meth:`pyopencl.Program.build`. *preamble* specifies a string of code that
  is inserted before the actual kernels.
  
  .. automethod:: __init__
  .. automethod:: __call__
class_docstr_tok: ["A","kernel","that","performs","a","generic","reduction","on","arrays",".","Generate","a","kernel","that","takes","a","number","of","scalar","or","vector","*","arguments","*","(","at","least","one","vector","argument",")",",","performs","the","*","map_expr","*","on","each","entry","of","the","vector","argument","and","then","the","*","reduce_expr","*","on","the","outcome","of","that",".","*","neutral","*","serves","as","an","initial","value",".","*","preamble","*","offers","the","possibility","to","add","preprocessor","directives","and","other","code","(","such","as","helper","functions",")","to","be","added","before","the","actual","reduction","kernel","code",".","Vectors","in","*","map_expr","*","should","be","indexed","by","the","variable","*","i","*",".","*","reduce_expr","*","uses","the","formal","values","``","a","''","and","``","b","''","to","indicate","two","operands","of","a","binary","reduction","operation",".","If","you","do","not","specify","a","*","map_expr","*",",","``","in","[","i","]","``","is","automatically","assumed","and","treated","as","the","only","one","input","argument",".","*","dtype_out","*","specifies","the",":","class",":","`","numpy.dtype","`","in","which","the","reduction","is","performed","and","in","which","the","result","is","returned",".","*","neutral","*","is","specified","as","float","or","integer","formatted","as","string",".","*","reduce_expr","*","and","*","map_expr","*","are","specified","as","string","formatted","operations","and","*","arguments","*","is","specified","as","a","string","formatted","as","a","C","argument","list",".","*","name","*","specifies","the","name","as","which","the","kernel","is","compiled",".","*","options","*","are","passed","unmodified","to",":","meth",":","`","pyopencl.Program.build","`",".","*","preamble","*","specifies","a","string","of","code","that","is","inserted","before","the","actual","kernels",".","..","automethod",":",":","__init__","..","automethod",":",":","__call__"]
code: def __init__(
      self,
      ctx: cl.Context,
      dtype_out: Any,
      neutral: str,
      reduce_expr: str,
      map_expr: Optional[str] = None,
      arguments: Optional[Union[str, List[DtypedArgument]]] = None,
      name: str = "reduce_kernel",
      options: Any = None,
      preamble: str = "",
  ) -> None:
      if arguments is None:
          raise ValueError("arguments must not be None")
  
      from pyopencl.tools import parse_arg_list
  
      arguments = parse_arg_list(arguments, with_offset=True)
  
      dtype_out = self.dtype_out = np.dtype(dtype_out)
  
      max_group_size = None
      trip_count = 0
  
      while True:
          self.stage_1_inf = get_reduction_kernel(
              1,
              ctx,
              dtype_out,
              neutral,
              reduce_expr,
              map_expr,
              arguments,
              name=f"{name}_stage1",
              options=options,
              preamble=preamble,
              max_group_size=max_group_size,
          )
  
          kernel_max_wg_size = (
              self.stage_1_inf.kernel.get_work_group_info(
                  cl.kernel_work_group_info.WORK_GROUP_SIZE,
                  ctx.devices[0],
              )
          )
  
          if self.stage_1_inf.group_size <= kernel_max_wg_size:
              break
          else:
              max_group_size = kernel_max_wg_size
  
          trip_count += 1
          assert trip_count <= 2
  
      self.stage_2_inf = get_reduction_kernel(
          2,
          ctx,
          dtype_out,
          neutral,
          reduce_expr,
          arguments=arguments,
          name=f"{name}_stage2",
          options=options,
          preamble=preamble,
          max_group_size=max_group_size,
      )
  
imports: from dataclasses import dataclass
  from typing import Any, List, Optional, Tuple, Union
  import numpy
  import pyopencl
  from pyopencl.tools import DtypedArgument, KernelTemplateBase, context_dependent_memoize, dtype_to_ctype, _process_code_for_macro
label_desc: Use image node_id 1 to create a new ReductionKernel object  with example: obj = ReductionKernel(ctx, dtype_out, neutral, reduce_expr, map_expr, arguments, name, options, preamble)