def get_exec_cond_and_kernel(
    func_attrs,
    reduce_op,
    reduction_axis,
    prologue_code,
    epilogue_scalar_code,
    input_type,
    output_type,
    acc_type,
    output_accessors,
    output_alignment,
    reduction_identity,
) -> str:
    """return a pair that contains the execution condition for this special
       reduction kernel and the source code of this reduction kernel

    Parameters
    ----------
    func_attrs : Dict[str, Any]
        holds attributes of this reduce op
    reduce_op : str
        the reduce op's string representation such as cutlass::plus
    reduction_axis : int
        the axis along which the reduction is performed
    prologue_code : str
        prologue code to process input before reduction
    epilogue_scalar_code : str
        epilogue code to process each final reduced element
    input_type : str
        specifies the input data type
    output_type : str
        specifies the output data type
    acc_type : str,
        specifies the data type for accumulation
    output_accessors : List[TensorAccessor]
        output TensorAccessor(s)
    output_alignment : int
        max alignment value that meets the requirement for accessing strided output

    Returns
    -------
    str
        returns the rendered code for the complete implementation of the reduce op
    """
    x = func_attrs["inputs"][0]
    x_shape = x._attrs["shape"]
    reduction_dim = x_shape[reduction_axis]
    # TODO: support dynamic reduction axis
    if not isinstance(reduction_dim, IntImm):
        return ("", "")

    reduction_dim_upperbound = 128
    if reduction_dim._attrs["values"][0] > reduction_dim_upperbound:
        return ("", "")

    rank = len(x_shape)
    # TODO: support reduction_axis = rank - 2
    valid_static_small_reduction_dim = reduction_axis == rank - 1
    reduction_dim_val = reduction_dim._attrs["values"][0]

    exec_cond = EXEC_COND_TEMPLATE.render(
        indent="  ",
        func_name=func_attrs["name"],
        elem_output_type=output_type,
        elem_input_type=input_type,
        elem_compute_type=acc_type,
        reduction_dim_upperbound=reduction_dim_upperbound,
        reduction_dim_val=reduction_dim_val,
        static_small_reduction_dim=valid_static_small_reduction_dim,
    )

    if output_accessors[0].is_contiguous:
        read_vec_type = _get_read_vector_type(x_shape, input_type)
    else:
        # For strided accesses, we force to take minimal vector type to
        # reduce uncoalesced device memory accesses. The perf penalty
        # of uncoalesced accesses (both read and write) seems to outweigh the
        # benefit from vector load.
        read_vec_type = _get_read_vector_type(
            x_shape, input_type, force_min_vec_type=True
        )
    kernel_src = KERNEL_SRC_TEMPLATE.render(
        reduce_op=reduce_op,
        reduction_identity=reduction_identity,
        prologue_code=prologue_code,
        epilogue_scalar_code=epilogue_scalar_code,
        read_vec_type=read_vec_type,
        output_accessor=output_accessors[0],
        output_access_alignment=output_alignment,
    )
    return (exec_cond, kernel_src)
return a pair that contains the execution condition for this special
reduction kernel and the source code of this reduction kernel

Parameters
----------
func_attrs : Dict[str, Any]
holds attributes of this reduce op
reduce_op : str
the reduce op's string representation such as cutlass::plus
reduction_axis : int
the axis along which the reduction is performed
prologue_code : str
prologue code to process input before reduction
epilogue_scalar_code : str
epilogue code to process each final reduced element
input_type : str
specifies the input data type
output_type : str
specifies the output data type
acc_type : str,
specifies the data type for accumulation
output_accessors : List[TensorAccessor]
output TensorAccessor(s)
output_alignment : int
max alignment value that meets the requirement for accessing strided output

Returns
-------
str
returns the rendered code for the complete implementation of the reduce op
python/aitemplate/backend/cuda/reduce/reduce_small_axis.py