// Copyright (c) 2026 Huawei Technologies Co., Ltd
// All rights reserved.
//
// Licensed under the BSD 3-Clause License  (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// 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.

#include <torch/extension.h>
#include "torch_npu/csrc/core/npu/NPUStream.h"
#include "torch_npu/csrc/framework/OpCommand.h"
#include "kernel_operator.h"

constexpr uint32_t BUFFER_NUM = 2;

// ------------------------------ KernelAdd算子类 ------------------------------
class KernelAdd {
public:
    __aicore__ inline KernelAdd() {}
    
    // 初始化方法：配置计算参数和内存资源
    // 参数：
    //   x, y, z: 输入输出张量的全局内存地址
    //   totalLength: 张量的总元素数量
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength)
    {
        // 计算每个block的长度
        this->blockLength = totalLength / AscendC::GetBlockNum();
        // 设置tile数量
        this->tileNum = 8;
        // 计算每个tile的长度
        this->tileLength = this->blockLength / this->tileNum / BUFFER_NUM;
        
        // 设置全局内存缓冲区
        xGm.SetGlobalBuffer((__gm__ int32_t *)x + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
        yGm.SetGlobalBuffer((__gm__ int32_t *)y + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
        zGm.SetGlobalBuffer((__gm__ int32_t *)z + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
        
        // 初始化输入输出队列
        pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(int32_t));
        pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileLength * sizeof(int32_t));
        pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(int32_t));
    }
    
    // Process方法
    __aicore__ inline void Process()
    {
        // 计算总循环次数 = tile数量 × 缓冲区数量
        int32_t loopCount = this->tileNum * BUFFER_NUM;
        
        // 执行流水线操作：CopyIn → Compute → CopyOut
        for (int32_t i = 0; i < loopCount; i++) {
            CopyIn(i);   // 数据从Global Memory复制到Local Memory
            Compute(i);  // 执行加法计算
            CopyOut(i);  // 结果从Local Memory复制回Global Memory
        }
    }

private:
    __aicore__ inline void CopyIn(int32_t progress)
    {
        AscendC::LocalTensor<int32_t> xLocal = inQueueX.AllocTensor<int32_t>();
        AscendC::LocalTensor<int32_t> yLocal = inQueueY.AllocTensor<int32_t>();
        
        AscendC::DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength);
        AscendC::DataCopy(yLocal, yGm[progress * this->tileLength], this->tileLength);

        inQueueX.EnQue(xLocal);
        inQueueY.EnQue(yLocal);
    }
    
    __aicore__ inline void Compute(int32_t progress)
    {
        // 从输入队列获取待处理的本地张量
        AscendC::LocalTensor<int32_t> xLocal = inQueueX.DeQue<int32_t>();
        AscendC::LocalTensor<int32_t> yLocal = inQueueY.DeQue<int32_t>();
        
        AscendC::LocalTensor<int32_t> zLocal = outQueueZ.AllocTensor<int32_t>();
        
        AscendC::Add(zLocal, xLocal, yLocal, this->tileLength);
        
        outQueueZ.EnQue<int32_t>(zLocal);
        
        inQueueX.FreeTensor(xLocal);
        inQueueY.FreeTensor(yLocal);
    }
    
    __aicore__ inline void CopyOut(int32_t progress)
    {
        AscendC::LocalTensor<int32_t> zLocal = outQueueZ.DeQue<int32_t>();
        
        AscendC::DataCopy(zGm[progress * this->tileLength], zLocal, this->tileLength);
        
        outQueueZ.FreeTensor(zLocal);
    }

private:
    AscendC::TPipe pipe;  // 任务管道，用于协调数据传输和计算
    
    AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueX;  // 输入张量x的队列
    AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueY;  // 输入张量y的队列
    
    AscendC::TQue<AscendC::TPosition::VECOUT, BUFFER_NUM> outQueueZ;  // 输出张量z的队列
    
    AscendC::GlobalTensor<int32_t> xGm;  // 输入张量x的全局内存映射
    AscendC::GlobalTensor<int32_t> yGm;  // 输入张量y的全局内存映射
    AscendC::GlobalTensor<int32_t> zGm;  // 输出张量z的全局内存映射
    
    // 计算参数
    uint32_t blockLength;  // 每个计算块的长度
    uint32_t tileNum;      // 每个块内的tile数量
    uint32_t tileLength;   // 每个tile的长度
};

// ------------------------------ 核函数 ------------------------------
// add_custom: 加法核函数，Ascend C算子设备侧实现入口
// 参数：
//   x, y, z: 输入输出张量的全局内存地址
//   totalLength: 张量的总元素数量
__global__ __vector__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength)
{
    // 初始化算子类
    KernelAdd op;
    
    // 初始化函数
    op.Init(x, y, z, totalLength);
    
    // 执行计算流程
    op.Process();
}

// ------------------------------ C++接口函数 ------------------------------
// ascendc_add: PyTorch和自定义内核之间的接口函数
// 参数：
//   x, y: 输入张量
// 返回值：
//   加法结果张量
namespace ascendc_ops {
at::Tensor run_ascendc_add(const at::Tensor &x, const at::Tensor &y)
{
    // 获取当前NPU流
    auto acl_stream = c10_npu::getCurrentNPUStream().stream(true);
    
    // 创建与输入形状相同的空张量作为输出
    at::Tensor z = at::empty_like(x);
    
    // 设置块维度
    uint32_t blockDim = 8;
    
    // 计算输入张量的总元素数量
    uint32_t totalLength = 1;
    for (uint32_t size : x.sizes()) {
        totalLength *= size;
    }
    
    // 定义内核启动lambda函数
    auto acl_call = [=]() -> int{
        // 启动自定义内核
        add_custom<<<blockDim, nullptr, acl_stream>>>((uint8_t *)(x.mutable_data_ptr()), 
                                                      (uint8_t *)(y.mutable_data_ptr()),
                                                      (uint8_t *)(z.mutable_data_ptr()), 
                                                      totalLength);
        return 0;
    };
    
    // 通过OpCommand运行内核
    at_npu::native::OpCommand::RunOpApiV2("ascendc_add", acl_call);
    
    // 返回计算结果
    return z;
}

}  // namespace ascendc_ops