code: def group_reduction_impl_idx_max_or_min(
      context, builder, sig, args, function
  ):
      """
      Instruction boilerplate used for calling a groupby reduction
      __device__ function in the case where the function is either
      `idxmax` or `idxmin`. See `group_reduction_impl_basic` for
      details. This lowering differs from other reductions due to
      the presence of the index. This results in the forward
      declaration expecting an extra arg.
      """
      retty = sig.return_type
  
      grp = cgutils.create_struct_proxy(sig.args[0])(
          context, builder, value=args[0]
      )
      grp_type = sig.args[0]
  
      if grp_type.index_type != index_default_type:
          raise TypeError(
              f"Only inputs with default index dtype {index_default_type} "
              "are supported."
          )
  
      type_key = (index_default_type, grp_type.group_scalar_type)
      func = call_cuda_functions[function][type_key]
  
      return context.compile_internal(
          builder,
          func,
          nb_signature(
              retty,
              grp_type.group_data_type,
              grp_type.group_index_type,
              grp_type.group_size_type,
          ),
          (grp.group_data, grp.index, grp.size),
      )
  
file_path: cudf/python/cudf/cudf/core/udf/groupby_lowering.py
imports: from functools import partial
  from numba import types
  from numba.core import cgutils
  from numba.core.extending import lower_builtin
  from numba.core.typing import signature
  from numba.cuda.cudaimpl import lower
  from cudf.core.udf.groupby_typing import SUPPORTED_GROUPBY_NUMBA_TYPES, Group, GroupType, call_cuda_functions, group_size_type, index_default_type
label_desc: Use image node_id 4 for calling a global function with example usage: group_reduction_impl_idx_max_or_min(context, builder, sig, args, function) and returns: context