/**
 * Copyright (c) 2025 Huawei Technologies Co., Ltd.
 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
 * CANN Open Software License Agreement Version 2.0 (the "License").
 * Please refer to the License for details. You may not use this file except in compliance with the License.
 * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
 * See LICENSE in the root of the software repository for the full text of the License.
 */

#include "kernel_operator.h"
#include "kernel_add.h"

using namespace AscendC;

namespace {
constexpr int32_t BUFFER_NUM = 2; // tensor num for each queue

class KernelAdd {
public:
    __aicore__ inline KernelAdd() {}

    __aicore__ inline void Init(__gm__ float* srcA, __gm__ float* srcB, __gm__ float* dst, uint32_t totalSize)
    {
        blockIdx = GetBlockIdx();
        blockDim = GetBlockNum();
        elementsPerBlock = totalSize / blockDim;
        tileLength = elementsPerBlock / TILE_NUM / BUFFER_NUM;
        startIdx = blockIdx * elementsPerBlock;

        // Set the global buffer
        srcAGlobal.SetGlobalBuffer(srcA + startIdx, elementsPerBlock);
        srcBGlobal.SetGlobalBuffer(srcB + startIdx, elementsPerBlock);
        dstGlobal.SetGlobalBuffer(dst + startIdx, elementsPerBlock);

        // Initialize the pipe buffer
        pipe.InitBuffer(queueInA, BUFFER_NUM, tileLength * sizeof(float));
        pipe.InitBuffer(queueInB, BUFFER_NUM, tileLength * sizeof(float));
        pipe.InitBuffer(queueOut, BUFFER_NUM, tileLength * sizeof(float));
    }

    __aicore__ inline void Process()
    {
        int32_t loopCount = TILE_NUM * BUFFER_NUM;
        for (int32_t i = 0; i < loopCount; i++) {
            CopyIn(i);
            Compute(i);
            CopyOut(i);
        }
    }

private:
    __aicore__ inline void CopyIn(int32_t progress)
    {
        LocalTensor<float> localA = queueInA.AllocTensor<float>();
        LocalTensor<float> localB = queueInB.AllocTensor<float>();

        DataCopy(localA, srcAGlobal[progress * tileLength], tileLength);
        DataCopy(localB, srcBGlobal[progress * tileLength], tileLength);

        queueInA.EnQue(localA);
        queueInB.EnQue(localB);
    }

    __aicore__ inline void Compute(int32_t progress)
    {
        LocalTensor<float> localA = queueInA.DeQue<float>();
        LocalTensor<float> localB = queueInB.DeQue<float>();
        LocalTensor<float> localOut = queueOut.AllocTensor<float>();

        Add(localOut, localA, localB, tileLength);

        queueOut.EnQue<float>(localOut);
        queueInA.FreeTensor(localA);
        queueInB.FreeTensor(localB);
    }

    __aicore__ inline void CopyOut(int32_t progress)
    {
        LocalTensor<float> localOut = queueOut.DeQue<float>();
        DataCopy(dstGlobal[progress * tileLength], localOut, tileLength);
        queueOut.FreeTensor(localOut);
    }

private:
    TPipe pipe;
    TQue<QuePosition::VECIN, BUFFER_NUM> queueInA;
    TQue<QuePosition::VECIN, BUFFER_NUM> queueInB;
    TQue<QuePosition::VECOUT, BUFFER_NUM> queueOut;
    GlobalTensor<float> srcAGlobal;
    GlobalTensor<float> srcBGlobal;
    GlobalTensor<float> dstGlobal;

    uint32_t blockIdx;
    uint32_t blockDim;
    uint32_t elementsPerBlock;
    uint32_t startIdx;
    uint32_t tileLength;
    static constexpr int32_t TILE_NUM = 8; // split data into 8 tiles for each core
};
} // namespace

extern "C" __global__ __aicore__ void AddKernel(
    __gm__ float* srcA, __gm__ float* srcB, __gm__ float* dst, uint32_t totalSize)
{
    KernelAdd op;
    op.Init(srcA, srcB, dst, totalSize);
    op.Process();
}

void AddDo(uint32_t blockDim, void* stream, float* srcA, float* srcB, float* dst, uint32_t totalSize)
{
    AddKernel<<<blockDim, nullptr, stream>>>(srcA, srcB, dst, totalSize);
}