Sum

#include "kernel_operator.h"
template <typename T>
class KernelSumCustom
{
public:
    __aicore__ inline KernelSumCustom() {}
    __aicore__ inline void Init(__gm__ uint8_t *src1Gm, __gm__ uint8_t *dstGm, uint32_t tmpSize, uint32_t outter, uint32_t inner, uint32_t n)
    {
        elementNumPerBlk = AscendC::ONE_BLK_SIZE / sizeof(T); // half=16 float=8
        elementNumPerRep = AscendC::ONE_REPEAT_BYTE_SIZE / sizeof(T);
        width = inner;
        height = outter;
        nNum = n;
        src1Global.SetGlobalBuffer((__gm__ T *)src1Gm);
        dstGlobal.SetGlobalBuffer((__gm__ T *)dstGm);
        pipe.InitBuffer(inQueueSrc1, 1, height * width * sizeof(T));
        pipe.InitBuffer(outQueueDst, 1, (height * sizeof(T) + AscendC::ONE_BLK_SIZE - 1) / AscendC::ONE_BLK_SIZE * AscendC::ONE_BLK_SIZE);
        pipe.InitBuffer(workQueue, 1, tmpSize);
    }

    __aicore__ inline void Process()
    {
        CopyIn();
        Compute();
        CopyOut();
    }

private:
    __aicore__ inline void CopyIn()
    {
        AscendC::LocalTensor<T> srcLocal1 = inQueueSrc1.AllocTensor<T>();
        AscendC::DataCopy(srcLocal1, src1Global, height * width);
        inQueueSrc1.EnQue(srcLocal1);
    }

    __aicore__ inline void Compute()
    {
        AscendC::LocalTensor<T> srcLocal1 = inQueueSrc1.DeQue<T>();
        AscendC::LocalTensor<uint8_t> workLocal = workQueue.AllocTensor<uint8_t>();
        AscendC::LocalTensor<T> dstLocal = outQueueDst.AllocTensor<T>();

        AscendC::SumParams params{height, width, nNum};
        T scalar(0);
        AscendC::Duplicate<T>(dstLocal, scalar, (height * sizeof(T) + AscendC::ONE_BLK_SIZE - 1) / AscendC::ONE_BLK_SIZE * AscendC::ONE_BLK_SIZE / sizeof(T));
        AscendC::Sum(dstLocal, srcLocal1, params);

        outQueueDst.EnQue<T>(dstLocal);
        workQueue.FreeTensor(workLocal);
        inQueueSrc1.FreeTensor(srcLocal1);
    }
    __aicore__ inline void CopyOut()
    {
        AscendC::LocalTensor<T> dstLocal = outQueueDst.DeQue<T>();
        AscendC::DataCopy(dstGlobal, dstLocal, (height * sizeof(T) + AscendC::ONE_BLK_SIZE - 1) / AscendC::ONE_BLK_SIZE * AscendC::ONE_BLK_SIZE / sizeof(T));
        outQueueDst.FreeTensor(dstLocal);
    }

private:
    AscendC::TPipe pipe;
    AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueSrc1;
    AscendC::TQue<AscendC::QuePosition::VECIN, 1> workQueue;
    AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueueDst;
    AscendC::GlobalTensor<T> src1Global, dstGlobal;
    uint32_t elementNumPerBlk = 0;
    uint32_t elementNumPerRep = 0;
    uint32_t width = 0;
    uint32_t height = 0;
    uint32_t nNum = 0;
};

extern "C" __global__ __aicore__ void sum_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
{
    GET_TILING_DATA(tiling_data, tiling);
    KernelSumCustom<DTYPE_X> op;
    op.Init(x, y, tiling_data.size, tiling_data.outter, tiling_data.inner, tiling_data.n);
    op.Process();
}

#include "register/tilingdata_base.h"

namespace optiling
{
  BEGIN_TILING_DATA_DEF(SumCustomTilingData)
  TILING_DATA_FIELD_DEF(uint32_t, size);
  TILING_DATA_FIELD_DEF(uint32_t, outter);
  TILING_DATA_FIELD_DEF(uint32_t, inner);
  TILING_DATA_FIELD_DEF(uint32_t, n);
  END_TILING_DATA_DEF;

  REGISTER_TILING_DATA_CLASS(SumCustom, SumCustomTilingData)
}

#include "sum_custom_tiling.h"
#include "register/op_def_registry.h"
#include "tiling/tiling_api.h"

namespace optiling
{
    static ge::graphStatus TilingFunc(gert::TilingContext *context)
    {
        SumCustomTilingData tiling;
        const gert::RuntimeAttrs *sumattrs = context->GetAttrs();
        const uint32_t tmp = *(sumattrs->GetAttrPointer<uint32_t>(0));
        const uint32_t outter = *(sumattrs->GetAttrPointer<uint32_t>(1));
        const uint32_t inner = *(sumattrs->GetAttrPointer<uint32_t>(2));
        uint32_t n = *(sumattrs->GetAttrPointer<uint32_t>(3));
        auto dt = context->GetInputTensor(0)->GetDataType();

        tiling.set_outter(outter); // 往结构体设置
        tiling.set_inner(inner);
        tiling.set_n(n);
        uint32_t maxValue = 0;
        uint32_t minValue = 0;
        uint32_t dtypesize;
        if (dt == ge::DT_FLOAT16) {
            dtypesize = 2;
        } else {
            dtypesize = 4;
        }
        AscendC::GetSumMaxMinTmpSize(n, dtypesize, false, maxValue, minValue);
        tiling.set_size(minValue);

        context->SetBlockDim(8);
        tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());
        context->GetRawTilingData()->SetDataSize(tiling.GetDataSize());
        size_t *currentWorkspace = context->GetWorkspaceSizes(1);
        currentWorkspace[0] = 0;
        return ge::GRAPH_SUCCESS;
    }
}

namespace ge
{
    static ge::graphStatus InferShape(gert::InferShapeContext *context)
    {
        const gert::Shape *x1_shape = context->GetInputShape(0);
        gert::Shape *y_shape = context->GetOutputShape(0);
        *y_shape = *x1_shape;
        return GRAPH_SUCCESS;
    }
}

namespace ops
{
    class SumCustom : public OpDef
    {
    public:
        explicit SumCustom(const char *name) : OpDef(name)
        {
            this->Input("x")
                .ParamType(REQUIRED)
                .DataType({ge::DT_FLOAT16})
                .Format({ge::FORMAT_ND});
            this->Output("y")
                .ParamType(REQUIRED)
                .DataType({ge::DT_FLOAT16})
                .Format({ge::FORMAT_ND});
            this->Attr("size")
                .AttrType(REQUIRED)
                .Int(0);
            this->Attr("outter")
                .AttrType(REQUIRED)
                .Int(0);
            this->Attr("inner")
                .AttrType(REQUIRED)
                .Int(0);
            this->Attr("n")
                .AttrType(REQUIRED)
                .Int(0);
            this->SetInferShape(ge::InferShape);
            this->AICore()
                .SetTiling(optiling::TilingFunc);
            this->AICore().AddConfig("ascendxxx"); // ascendxxx请修改为对应的昇腾AI处理器型号。
        }
    };
    OP_ADD(SumCustom);
} // namespace ops