def METHOD_NAME(x, axis=-1):
    """Perform log softmax activation on the data
    Parameters
    ----------
    data : tvm.te.Tensor
        N-D input data
    Returns
    -------
    output : tvm.te.Tensor
        N-D output with same shape
    """
    shape = x.shape
    if axis < 0:
        axis = len(shape) + axis
    if axis >= len(shape):
        ValueError("axis parameter should be less than input dim")
    k1 = te.reduce_axis((0, shape[axis]), name="k")
    k2 = te.reduce_axis((0, shape[axis]), name="k")
    def insert_reduce_index(indices, reduce_index):
        return indices[:axis] + (reduce_index,) + indices[axis:]
    def get_non_reduce_indices(indices):
        return tuple([var for (i, var) in enumerate(indices) if i != axis])
    def _compute_max(*indices):
        eval_range = insert_reduce_index(indices, k1)
        return tvm.te.max(x[eval_range], axis=k1)
    def _compute_expsum(max_elem, *indices):
        eval_range = insert_reduce_index(indices, k2)
        return te.sum(te.exp(x[eval_range] - max_elem[indices]), axis=k2)
    def _normalize(max_elem, expsum, *indices):
        non_reduce_indices = get_non_reduce_indices(indices)
        return x[indices] - max_elem[non_reduce_indices] - te.log(expsum[non_reduce_indices])
    reduced_shape = tuple([dim for (i, dim) in enumerate(shape) if i != axis])
    max_elem = te.compute(reduced_shape, _compute_max, name="T_softmax_maxelem")
    expsum = te.compute(reduced_shape, lambda *indices: _compute_expsum(max_elem, *indices))
    return te.compute(
        shape,
        lambda *indices: _normalize(max_elem, expsum, *indices),
        attrs={"axis": axis},
    )