code: def get_sum_kernel(ctx, dtype_out, dtype_in):
      if dtype_out is None:
          dtype_out = dtype_in
  
      reduce_expr = "a+b"
      neutral_expr = "0"
      if dtype_out.kind == "c":
          from pyopencl.elementwise import complex_dtype_to_name
  
          dtname = complex_dtype_to_name(dtype_out)
          reduce_expr = f"{dtname}_add(a, b)"
          neutral_expr = f"{dtname}_new(0, 0)"
  
      return ReductionKernel(
          ctx,
          dtype_out,
          neutral_expr,
          reduce_expr,
          arguments="const {} *in".format(dtype_to_ctype(dtype_in)),
      )
  
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 5 for calling a global function with example usage: get_sum_kernel(ctx, dtype_out, dtype_in) and returns: ReductionKernel