from tilelang.tileop.gemm_sp.gemm_sp_base import GemmSPBase
from tilelang.layout import make_swizzled_layout
from tilelang.cuda.intrinsics.macro.mma_sp_macro_generator import SparseTensorCoreIntrinEmitter
from tilelang.utils.language import is_shared, is_fragment
from tilelang import tvm as tvm
from tvm.target import Target
from tvm.ir import Range
from tvm import tirx
from tilelang import language as T
from tilelang.transform.simplify import _Simplify


GEMM_SP_INST_MMA_SP = "cuda.mma.sp"


class GemmSPMMA(GemmSPBase):
    def infer_layout(self, target: Target, thread_nums: int):
        m_warp, n_warp = self.policy.compute_warp_partition(self.M, self.N, thread_nums, target, GEMM_SP_INST_MMA_SP)
        warp_row_tiles = int(self.M // m_warp)
        warp_col_tiles = int(self.N // n_warp)
        mma_emitter = SparseTensorCoreIntrinEmitter(
            a_dtype=self.a_dtype,
            e_dtype=self.e_dtype,
            b_dtype=self.b_dtype,
            accum_dtype=self.accum_dtype,
            a_transposed=self.trans_A,
            b_transposed=self.trans_B,
            e_transposed=self.trans_E,
            block_row_warps=m_warp,
            block_col_warps=n_warp,
            warp_row_tiles=warp_row_tiles,
            warp_col_tiles=warp_col_tiles,
            warp_k=self.K,
        )
        if self.is_gemm_ss():
            return {
                self.A: make_swizzled_layout(self.A),
                self.B: make_swizzled_layout(self.B),
                self.C: mma_emitter.make_mma_store_layout(self.C),
            }
        elif self.is_gemm_sr():
            return {
                self.A: make_swizzled_layout(self.A),
                self.B: mma_emitter.make_mma_load_layout(self.B, matrix="B"),
                self.C: mma_emitter.make_mma_store_layout(self.C),
            }
        elif self.is_gemm_rs():
            return {
                self.A: mma_emitter.make_mma_load_layout(self.A, matrix="A"),
                self.B: make_swizzled_layout(self.B),
                self.C: mma_emitter.make_mma_store_layout(self.C),
            }
        elif self.is_gemm_rr():
            return {
                self.A: mma_emitter.make_mma_load_layout(self.A, matrix="A"),
                self.B: mma_emitter.make_mma_load_layout(self.B, matrix="B"),
                self.C: mma_emitter.make_mma_store_layout(self.C),
            }
        else:
            raise ValueError(f"Unsupported gemm combination, A: {self.A.scope()}, B: {self.B.scope()}")

    def lower(self, layout_map: dict, target: Target, thread_bounds: Range, thread_var: tirx.Var):
        thread_nums = thread_bounds.extent
        # Emitter lane/warp math uses zero-based ids within the current thread bounds.
        local_thread_var = thread_var - thread_bounds.min
        m_warp, n_warp = self.policy.compute_warp_partition(self.M, self.N, thread_nums, target, GEMM_SP_INST_MMA_SP)
        warp_row_tiles = int(self.M // m_warp)
        warp_col_tiles = int(self.N // n_warp)
        mma_emitter = SparseTensorCoreIntrinEmitter(
            a_dtype=self.a_dtype,
            b_dtype=self.b_dtype,
            e_dtype=self.e_dtype,
            accum_dtype=self.accum_dtype,
            a_transposed=self.trans_A,
            b_transposed=self.trans_B,
            e_transposed=self.trans_E,
            block_row_warps=m_warp,
            block_col_warps=n_warp,
            warp_row_tiles=warp_row_tiles,
            warp_col_tiles=warp_col_tiles,
            warp_k=self.K,
            thread_var=local_thread_var,
        )

        a_dtype = self.a_dtype
        b_dtype = self.b_dtype
        warp_rows = mma_emitter.warp_rows
        warp_cols = mma_emitter.warp_cols
        local_size_a = mma_emitter.local_size_a
        local_size_e = mma_emitter.local_size_e
        local_size_b = mma_emitter.local_size_b
        micro_size_k = mma_emitter.micro_size_k
        A_shared = self.ARegion
        E_shared = self.ERegion
        B_shared = self.BRegion
        C_local = self.C
        clear_accum = self.clear_accum
        assert micro_size_k <= self.K, f"K dimension {self.K} should be >= micro size k {micro_size_k}"
        if self.is_gemm_ss():

            @T.prim_func
            def _gemm_ssr() -> None:
                """
                The inner macro that loads data from shared buffers A_shared and
                B_shared into local fragments, then issues Tensor Core mma ops,
                accumulating into C_local.
                """
                A_local = T.alloc_local((warp_rows * local_size_a), a_dtype)
                E_local = T.alloc_local((warp_rows * local_size_e), self.e_dtype)
                B_local = T.alloc_local((warp_cols * local_size_b), b_dtype)

                if clear_accum:
                    T.clear(C_local)

                for ki in T.serial(0, (self.K // micro_size_k)):
                    # Load A into fragment
                    mma_emitter.ldmatrix_a(
                        A_local,
                        A_shared,
                        ki,
                    )

                    # Load E into fragment
                    mma_emitter.ldmatrix_e(
                        E_local,
                        E_shared,
                        ki,
                    )

                    # Load B into fragment
                    mma_emitter.ldmatrix_b(
                        B_local,
                        B_shared,
                        ki,
                    )

                    # Perform Matrix Multiplication
                    mma_emitter.mma_sp(A_local, E_local, B_local, C_local, ki)

            # Simplify to optimize the index computing
            # Must inline let statements to simplify the analysis
            return _Simplify(_gemm_ssr, inline_let=True)
        elif self.is_gemm_sr():
            B_local = self.B

            @T.prim_func
            def _gemm_srr() -> None:
                """
                The inner macro that loads data from shared buffers A_shared and
                B_shared into local fragments, then issues Tensor Core mma ops,
                accumulating into C_local.
                """
                A_local = T.alloc_local((warp_rows * local_size_a), a_dtype)
                E_local = T.alloc_local((warp_rows * local_size_e), self.e_dtype)

                if clear_accum:
                    T.clear(C_local)

                for ki in T.serial(0, (self.K // micro_size_k)):
                    # Load A into fragment
                    mma_emitter.ldmatrix_a(
                        A_local,
                        A_shared,
                        ki,
                    )

                    # Load E into fragment
                    mma_emitter.ldmatrix_e(
                        E_local,
                        E_shared,
                        ki,
                    )

                    # Perform Matrix Multiplication
                    mma_emitter.mma_sp(A_local, E_local, B_local, C_local, ki)

            # Simplify to optimize the index computing
            # Must inline let statements to simplify the analysis
            # alloc_buffers body
            # insert into parent block
            return _Simplify(_gemm_srr, inline_let=True)
        elif self.is_gemm_rs():
            A_local = self.A

            @T.prim_func
            def _gemm_rsr() -> None:
                """
                The inner macro that loads data from shared buffers A_shared and
                B_shared into local fragments, then issues Tensor Core mma ops,
                accumulating into C_local.
                """
                E_local = T.alloc_local((warp_rows * local_size_e), self.e_dtype)
                B_local = T.alloc_local((warp_cols * local_size_b), b_dtype)

                if clear_accum:
                    T.clear(C_local)

                for ki in T.serial(0, (self.K // micro_size_k)):
                    # Load E into fragment
                    mma_emitter.ldmatrix_e(
                        E_local,
                        E_shared,
                        ki,
                    )

                    # Load B into fragment
                    mma_emitter.ldmatrix_b(
                        B_local,
                        B_shared,
                        ki,
                    )

                    # Perform Matrix Multiplication
                    mma_emitter.mma_sp(A_local, E_local, B_local, C_local, ki)

            # Simplify to optimize the index computing
            # Must inline let statements to simplify the analysis
            return _Simplify(_gemm_rsr, inline_let=True)
        elif self.is_gemm_rr():
            A_local = self.A
            B_local = self.B

            @T.prim_func
            def _gemm_rrr() -> None:
                """
                The inner macro that loads data from shared buffers A_shared and
                B_shared into local fragments, then issues Tensor Core mma ops,
                accumulating into C_local.
                """
                E_local = T.alloc_local((warp_rows * local_size_e), self.e_dtype)

                if clear_accum:
                    T.clear(C_local)

                for ki in T.serial(0, (self.K // micro_size_k)):
                    # Load E into fragment
                    mma_emitter.ldmatrix_e(
                        E_local,
                        E_shared,
                        ki,
                    )

                    # Perform Matrix Multiplication
                    mma_emitter.mma_sp(A_local, E_local, B_local, C_local, ki)

            # Simplify to optimize the index computing
            # Must inline let statements to simplify the analysis
            return _Simplify(_gemm_rrr, inline_let=True)
        else:
            raise ValueError(f"Unsupported gemm combination, A: {self.A.scope()}, B: {self.B.scope()}")

    def is_gemm_ss(self) -> bool:
        return is_shared(self.A) and is_shared(self.B)

    def is_gemm_sr(self) -> bool:
        return is_shared(self.A) and is_fragment(self.B)

    def is_gemm_rs(self) -> bool:
        return is_fragment(self.A) and is_shared(self.B)

    def is_gemm_rr(self) -> bool:
        return is_fragment(self.A) and is_fragment(self.B)
