# Licensed to the Apache Software Foundation (ASF) under one
# or more contributor license agreements.  See the NOTICE file
# distributed with this work for additional information
# regarding copyright ownership.  The ASF licenses this file
# to you under the Apache License, Version 2.0 (the
# "License"); you may not use this file except in compliance
# with the License.  You may obtain a copy of the License at
#
#   http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing,
# software distributed under the License is distributed on an
# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
# KIND, either express or implied.  See the License for the
# specific language governing permissions and limitations
# under the License.

"""FIFO buffer op"""
from __future__ import absolute_import as _abs
import tvm
from tvm import te
from .. import tag
from ..transform import concatenate, strided_slice



[文档]
@tvm.te.tag_scope(tag=tag.INJECTIVE + ",fifo_buffer")
def fifo_buffer(data, buffer, axis):
    """
    FIFO buffer to enable computation reuse in CNNs with sliding indow input

    Compute equivalent of

    .. code-block:: python

        concat(buffer, data, axis=axis)
        .slice_axis(axis=axis,
                    begin=data.shape[axis],
                    end=data.shape[axis]+buffer.shape[axis])

    Useful for

    * Encoding explicit re-use of computation in convolution ops operated on a sliding window input
    * Implementing a FIFO queue to cache intermediate results, e.g. as in Fast WaveNet.

    Parameters
    ----------
    data : tvm.te.Tensor
        The input data
    buffer : tvm.te.Tensor
        Previous value of the FIFO buffer
    axis : int
        Specify which axis should be used for buffering

    Returns
    -------
    result : tvm.te.Tensor
        Updated value for the buffer
    """
    assert len(data.shape) == len(buffer.shape), (
        f"buffer and data must have same number of dimensions, "
        f"buffer.shape = {buffer.shape}, data.shape = {data.shape}"
    )
    assert len(buffer.shape) >= 1, "Zero-dimension tensor not supported"
    assert 0 <= axis < len(buffer.shape), "buffer axis out of range"
    for i in range(len(data.shape)):
        if i == axis:
            assert int(str(data.shape[i])) <= int(str(buffer.shape[i]))
        else:
            assert int(str(data.shape[i])) == int(str(buffer.shape[i]))

    buflen = buffer.shape[axis]
    data_size = data.shape[axis]

    # Explicitly write out formula up to 4D, and then use concat+slice combo for 5D and higher
    if len(buffer.shape) == 1:
        return te.compute(
            buffer.shape,
            lambda i: tvm.tir.if_then_else(
                i < buflen - data_size, buffer[i + data_size], data[i - buflen + data_size]
            ),
            name="new_buffer",
        )
    if len(buffer.shape) == 2:
        if axis == 0:
            return te.compute(
                buffer.shape,
                lambda i, j: tvm.tir.if_then_else(
                    i < buflen - data_size,
                    buffer[i + data_size, j],
                    data[i - buflen + data_size, j],
                ),
                name="new_buffer",
            )
        if axis == 1:
            return te.compute(
                buffer.shape,
                lambda i, j: tvm.tir.if_then_else(
                    j < buflen - data_size,
                    buffer[i, j + data_size],
                    data[i, j - buflen + data_size],
                ),
                name="new_buffer",
            )
        assert False, f"Invalid value for axis; it should be at most {len(buffer.shape)}"
    elif len(buffer.shape) == 3:
        if axis == 0:
            return te.compute(
                buffer.shape,
                lambda i, j, k: tvm.tir.if_then_else(
                    i < buflen - data_size,
                    buffer[i + data_size, j, k],
                    data[i - buflen + data_size, j, k],
                ),
                name="new_buffer",
            )
        if axis == 1:
            return te.compute(
                buffer.shape,
                lambda i, j, k: tvm.tir.if_then_else(
                    j < buflen - data_size,
                    buffer[i, j + data_size, k],
                    data[i, j - buflen + data_size, k],
                ),
                name="new_buffer",
            )
        if axis == 2:
            return te.compute(
                buffer.shape,
                lambda i, j, k: tvm.tir.if_then_else(
                    k < buflen - data_size,
                    buffer[i, j, k + data_size],
                    data[i, j, k - buflen + data_size],
                ),
                name="new_buffer",
            )
        assert False, f"Invalid value for axis; it should be at most {len(buffer.shape)}"
    elif len(buffer.shape) == 4:
        if axis == 0:
            return te.compute(
                buffer.shape,
                lambda i, j, k, l: tvm.tir.if_then_else(
                    i < buflen - data_size,
                    buffer[i + data_size, j, k, l],
                    data[i - buflen + data_size, j, k, l],
                ),
                name="new_buffer",
            )
        if axis == 1:
            return te.compute(
                buffer.shape,
                lambda i, j, k, l: tvm.tir.if_then_else(
                    j < buflen - data_size,
                    buffer[i, j + data_size, k, l],
                    data[i, j - buflen + data_size, k, l],
                ),
                name="new_buffer",
            )
        if axis == 2:
            return te.compute(
                buffer.shape,
                lambda i, j, k, l: tvm.tir.if_then_else(
                    k < buflen - data_size,
                    buffer[i, j, k + data_size, l],
                    data[i, j, k - buflen + data_size, l],
                ),
                name="new_buffer",
            )
        if axis == 3:
            return te.compute(
                buffer.shape,
                lambda i, j, k, l: tvm.tir.if_then_else(
                    l < buflen - data_size,
                    buffer[i, j, k, l + data_size],
                    data[i, j, k, l - buflen + data_size],
                ),
                name="new_buffer",
            )
        assert False, f"Invalid value for axis; it should be at most {len(buffer.shape)}"
    else:
        # Implement FIFO buffer as combination of concat and slice
        begin = [0] * len(buffer.shape)
        begin[axis] = data.shape[axis]
        end = list(buffer.shape[:])
        end[axis] += data.shape[axis]
        return strided_slice(concatenate((buffer, data), axis=axis), begin=begin, end=end)
    return None