code: def _get_generic_filter1d(
      rk,
      length,
      n_lines,
      filter_size,
      origin,
      mode,
      cval,
      in_ctype,
      out_ctype,
      int_type,
  ):
      """
      The generic 1d filter is different than all other filters and thus is the
      only filter that doesn't use _generate_nd_kernel() and has a completely
      custom raw kernel.
      """
      in_length = length + filter_size - 1
      start = filter_size // 2 + origin
      end = start + length
  
      if mode == "constant":
          boundary, boundary_early = (
              "",
              """
          for (idx_t j = 0; j < {start}; ++j) {{ input_line[j] = {cval}; }}
          for (idx_t j = {end}; j<{in_length}; ++j) {{ input_line[j] = {cval}; }}
          """.format(
                  start=start, end=end, in_length=in_length, cval=cval
              ),
          )
      else:
          if length == 1:
              a = b = "j_ = 0;"
          elif mode == "reflect":
              j = (
                  "j_ = ({j}) % ({length} * 2);\n"
                  "j_ = min(j_, 2 * {length} - 1 - j_);"
              )
              a = j.format(j="-1 - j_", length=length)
              b = j.format(j="j_", length=length)
          elif mode == "mirror":
              j = (
                  "j_ = 1 + (({j}) - 1) % (({length} - 1) * 2);\n"
                  "j_ = min(j_, 2 * {length} - 2 - j_);"
              )
              a = j.format(j="-j_", length=length)
              b = j.format(j="j_", length=length)
          elif mode == "nearest":
              a, b = "j_ = 0;", "j_ = {length}-1;".format(length=length)
          elif mode == "wrap":
              a = "j_ = j_ % {length} + {length};".format(length=length)
              b = "j_ = j_ % {length};".format(length=length)
          loop = """for (idx_t j = {{}}; j < {{}}; ++j) {{{{
              idx_t j_ = j - {start};
              {{}}
              input_line[j] = input_line[j_ + {start}];
          }}}}""".format(
              start=start
          )
          boundary_early = ""
          boundary = (
              loop.format(0, start, a)
              + "\n"
              + loop.format(end, in_length, b)
          )
  
      name = "generic1d_{}_{}_{}".format(length, filter_size, rk.name)
      code = """#include "cupy/carray.cuh"
  #include "cupy/complex.cuh"
  {include_type_traits}  // let Jitify handle this
  
  namespace raw_kernel {{\n{rk_code}\n}}
  
  {CAST}
  
  typedef unsigned char byte;
  typedef {in_ctype} X;
  typedef {out_ctype} Y;
  typedef {int_type} idx_t;
  
  __device__ idx_t offset(idx_t i, idx_t axis, idx_t ndim,
                          const idx_t* shape, const idx_t* strides) {{
      idx_t index = 0;
      for (idx_t a = ndim; --a > 0; ) {{
          if (a == axis) {{ continue; }}
          index += (i % shape[a]) * strides[a];
          i /= shape[a];
      }}
      return index + strides[0] * i;
  }}
  
  extern "C" __global__
  void {name}(const byte* input, byte* output, const idx_t* x) {{
      const idx_t axis = x[0], ndim = x[1],
          *shape = x+2, *in_strides = x+2+ndim, *out_strides = x+2+2*ndim;
  
      const idx_t in_elem_stride = in_strides[axis];
      const idx_t out_elem_stride = out_strides[axis];
  
      double input_line[{in_length}];
      double output_line[{length}];
      {boundary_early}
  
      for (idx_t i = ((idx_t)blockIdx.x) * blockDim.x + threadIdx.x;
              i < {n_lines};
              i += ((idx_t)blockDim.x) * gridDim.x) {{
          // Copy line from input (with boundary filling)
          const byte* input_ = input + offset(i, axis, ndim, shape, in_strides);
          for (idx_t j = 0; j < {length}; ++j) {{
              input_line[j+{start}] = (double)*(X*)(input_+j*in_elem_stride);
          }}
          {boundary}
  
          raw_kernel::{rk_name}(input_line, {in_length}, output_line, {length});
  
          // Copy line to output
          byte* output_ = output + offset(i, axis, ndim, shape, out_strides);
          for (idx_t j = 0; j < {length}; ++j) {{
              *(Y*)(output_+j*out_elem_stride) = cast<Y>(output_line[j]);
          }}
      }}
  }}""".format(
          n_lines=n_lines,
          length=length,
          in_length=in_length,
          start=start,
          in_ctype=in_ctype,
          out_ctype=out_ctype,
          int_type=int_type,
          boundary_early=boundary_early,
          boundary=boundary,
          name=name,
          rk_name=rk.name,
          # Users can test RawKernel independently, but when passed to here
          # it must be used as a device function here. In fact, RawKernel
          # wouldn't compile if code only contains device functions, so this
          # is necessary.
          rk_code=rk.code.replace("__global__", "__device__"),
          include_type_traits=(
              "" if runtime.is_hip else "#include <type_traits>\n"
          ),
          CAST=_filters_core._CAST_FUNCTION,
      )
      return cupy.RawKernel(
          code, name, ("--std=c++11",) + rk.options, jitify=True
      )
  
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 6 for calling a global function with example usage: _get_generic_filter1d(rk, length, n_lines, filter_size, origin, mode, cval, in_ctype, out_ctype, int_type) and returns: cupy