code: def _reduction_kernel_code(rk, filter_size, out_dtype, in_dtype):
      # NOTE: differences from the code generated for real reduction kernels:
      #  * input is always 1D and always less than 2^31 elements
      #  * output is always 1D with a single element
      #  * never across threads (no _block_stride, _sdata, _sdata_raw, _REDUCE,
      #       _tid, _J, _i, _i_base, _j_offset, _J_offset, _j_stride, _J_stride)
      # Also, the code is moved into a namespace so that clashes are minimized
      # between the typedefs for the "template" variables.
  
      # figure out the types
      types = {}
      in_param, out_param = rk.in_params[0], rk.out_params[0]
      in_ctype = _get_type_info(in_param, in_dtype, types)
      out_ctype = _get_type_info(out_param, out_dtype, types)
      types = "\n".join(
          "typedef {} {};".format(typ, name)
          for name, typ in types.items()
      )
  
      return """namespace reduction_kernel {{
  {type_preamble}
  {preamble}
  __device__
  void {name}({in_const} CArray<{in_ctype}, 1, true, true>& _raw_{in_name},
              CArray<{out_ctype}, 1, true, true>& _raw_{out_name}) {{
      // these are just provided so if they are available for the RK
      CIndexer<1> _in_ind({{{size}}});
      CIndexer<0> _out_ind;
  
      #define REDUCE(a, b) ({reduce_expr})
      #define POST_MAP(a) ({post_map_expr})
      typedef {reduce_type} _type_reduce;
      _type_reduce _s = _type_reduce({identity});
      for (int _j = 0; _j < {size}; ++_j) {{
          _in_ind.set(_j);
          {in_const} {in_ctype}& {in_name} = _raw_{in_name}[_j];
          _type_reduce _a = static_cast<_type_reduce>({pre_map_expr});
          _s = REDUCE(_s, _a);
      }}
      _out_ind.set(0);
      {out_ctype} &{out_name} = _raw_{out_name}[0];
      POST_MAP(_s);
      #undef REDUCE
      #undef POST_MAP
  }}
  }}""".format(
          name=rk.name,
          type_preamble=types,
          preamble=rk.preamble,
          in_const="const" if in_param.is_const else "",
          in_ctype=in_ctype,
          in_name=in_param.name,
          out_ctype=out_ctype,
          out_name=out_param.name,
          pre_map_expr=rk.map_expr,
          identity="" if rk.identity is None else rk.identity,
          size=filter_size,
          reduce_type=rk.reduce_type,
          reduce_expr=rk.reduce_expr,
          post_map_expr=rk.post_map_expr,
      )
  
file_path: cupy/cupyx/scipy/ndimage/_filters_generic.py
imports: import cupy
  from cupy_backends.cuda.api import runtime
  from cupy import _util
  from cupyx.scipy.ndimage import _filters_core
label_desc: Use image node_id 3 for calling a global function with example usage: _reduction_kernel_code(rk, filter_size, out_dtype, in_dtype) and returns: str