# 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.
# pylint: disable=invalid-name, missing-docstring

import tvm
import tvm.testing
from tvm.script import ir as I
from tvm.script import tirx as T


def test_rewrite_to_shuffle_0():
    transform = tvm.tirx.transform.PointerValueTypeRewrite()

    @I.ir_module
    class Before:
        @T.prim_func(s_tir=True)
        def main(A: T.Buffer((16,), "float32"), B: T.Buffer((4,), "float32")):
            A_local = T.alloc_buffer((16,), scope="local")
            for i in range(4):
                A_local[i * 4 : i * 4 + 4] = A[i * 4 : i * 4 + 4]
            for i in range(4):
                B[i] = A_local[i * 4] + A_local[i * 4 + 1] + A_local[i * 4 + 2] + A_local[i * 4 + 3]

    @I.ir_module
    class Expected:
        @T.prim_func(s_tir=True)
        def main(A: T.Buffer((4,), "float32x4"), B: T.Buffer((4,), "float32")):
            A_local = T.alloc_buffer((4,), "float32x4", scope="local")
            for i in range(4):
                A_local[T.Div(i * 4, 4)] = A[T.Div(i * 4, 4)]
            for i in range(4):
                B[i] = (
                    T.Shuffle([A_local[T.Div(i * 4, 4)]], [0])
                    + T.Shuffle([A_local[T.Div(i * 4 + 1, 4)]], [1])
                    + T.Shuffle([A_local[T.Div(i * 4 + 2, 4)]], [2])
                    + T.Shuffle([A_local[T.Div(i * 4 + 3, 4)]], [3])
                )

    After = transform(Before)
    tvm.ir.assert_structural_equal(After, Expected)


def test_rewrite_to_shuffle_1():
    transform = tvm.tirx.transform.PointerValueTypeRewrite()

    @I.ir_module
    class Before:
        @T.prim_func(s_tir=True)
        def main(A: T.Buffer((8,), "float32"), B: T.Buffer((1,), "float32")):
            A_local = T.alloc_buffer((8,), scope="local")
            A_local[0:4] = A[0:4]
            A_local[4:8] = A[4:8]
            B[0] = (
                A_local[0]
                + A_local[1]
                + A_local[2]
                + A_local[3]
                + A_local[4]
                + A_local[5]
                + A_local[6]
                + A_local[7]
            )

    @I.ir_module
    class Expected:
        @T.prim_func(s_tir=True)
        def main(A: T.Buffer((2,), "float32x4"), B: T.Buffer((1,), "float32")):
            A_local = T.alloc_buffer((2,), "float32x4", scope="local")
            A_local[0] = A[0]
            A_local[1] = A[1]
            B[0] = (
                T.Shuffle([A_local[0]], [0])
                + T.Shuffle([A_local[0]], [1])
                + T.Shuffle([A_local[0]], [2])
                + T.Shuffle([A_local[0]], [3])
                + T.Shuffle([A_local[1]], [0])
                + T.Shuffle([A_local[1]], [1])
                + T.Shuffle([A_local[1]], [2])
                + T.Shuffle([A_local[1]], [3])
            )

    After = transform(Before)
    tvm.ir.assert_structural_equal(After, Expected)


def test_address_of():
    transform = tvm.tirx.transform.PointerValueTypeRewrite()

    @I.ir_module
    class Before:
        @T.prim_func(s_tir=True)
        def main(A: T.Buffer((16,), "float32"), B: T.Buffer((16,), "float32")):
            for i in range(4):
                T.evaluate(T.address_of(A[i * 4]))
                B[i * 4 : i * 4 + 4] = A[i * 4 : i * 4 + 4]

    @I.ir_module
    class Expected:
        @T.prim_func(s_tir=True)
        def main(A: T.Buffer((16,), "float32"), B: T.Buffer((4,), "float32x4")):
            for i in range(4):
                T.evaluate(T.address_of(A[i * 4]))
                B[T.Div(i * 4, 4)] = A[i * 4 : i * 4 + 4]

    After = transform(Before)
    tvm.ir.assert_structural_equal(After, Expected)


def test_scalar_read_without_write():
    transform = tvm.tirx.transform.PointerValueTypeRewrite()

    @I.ir_module
    class Before:
        @T.prim_func(s_tir=True)
        def main(A: T.Buffer((16,), "float32")):
            for i in range(4):
                T.evaluate(A[i * 4])

    # Expected is the same as Before - no transformation
    @I.ir_module
    class Expected:
        @T.prim_func(s_tir=True)
        def main(A: T.Buffer((16,), "float32")):
            for i in range(4):
                T.evaluate(A[i * 4])

    After = transform(Before)
    tvm.ir.assert_structural_equal(After, Expected)


def test_decl_buffer_alias_chain_uses_flat_root_map():
    transform = tvm.tirx.transform.PointerValueTypeRewrite()

    @I.ir_module
    class Before:
        @T.prim_func(s_tir=True)
        def main(A: T.Buffer((16,), "float32")):
            A_view = T.decl_buffer((16,), "float32", data=A.data)
            A_view_2 = T.decl_buffer((16,), "float32", data=A_view.data)
            for i in range(4):
                A_view_2[i * 4 : i * 4 + 4] = T.broadcast(T.float32(1), 4)

    After = transform(Before)
    assert tvm.tirx.analysis.verify_well_formed(After)
    func = After["main"]
    assert func.params[0].ty.dtype == tvm.ir.PrimType("float32x4")

    decl_buffers = []
    buffer_stores = []
    tvm.tirx.stmt_functor.post_order_visit(
        func.body,
        lambda node: (
            decl_buffers.append(node)
            if isinstance(node, tvm.tirx.DeclBuffer)
            else buffer_stores.append(node)
            if isinstance(node, tvm.tirx.BufferStore)
            else None
        ),
    )
    assert len(decl_buffers) == 2
    assert all(decl.buffer.ty.dtype == tvm.ir.PrimType("float32x4") for decl in decl_buffers)
    assert len(buffer_stores) == 1
    assert buffer_stores[0].buffer.ty.dtype == tvm.ir.PrimType("float32x4")


if __name__ == "__main__":
    tvm.testing.main()
