Conv2D

本样例中，输入矩阵形状为[C1, H, W, C0]，特征矩阵的形状为[C1, Kh, Kw, Cout, C0], 计算结果搬出至GM，目的矩阵无需初始化。

#include "kernel_operator.h"

namespace AscendC {
class KernelCubeConv2D {
public:
    __aicore__ inline KernelCubeConv2D() {}
    __aicore__ inline void Init(__gm__ half* fmGm, __gm__ half* weGm, __gm__ float* biasGm,
        __gm__ float* dstGm, Conv2dParams params, LoopMode mode)
    {
        conv2dParams.imgShape[0]        = params.imgShape[0];
        conv2dParams.imgShape[1]        = params.imgShape[1];
        conv2dParams.kernelShape[0]     = params.kernelShape[0];
        conv2dParams.kernelShape[1]     = params.kernelShape[1];
        conv2dParams.stride[0]          = params.stride[0];
        conv2dParams.stride[1]          = params.stride[1];
        conv2dParams.cin                = params.cin;
        conv2dParams.cout               = params.cout;
        conv2dParams.padList[0]         = params.padList[0];
        conv2dParams.padList[1]         = params.padList[1];
        conv2dParams.padList[2]         = params.padList[2];
        conv2dParams.padList[3]         = params.padList[3];
        conv2dParams.dilation[0]        = params.dilation[0];
        conv2dParams.dilation[1]        = params.dilation[1];
        conv2dParams.initY              = params.initY;
        conv2dParams.partialSum         = params.partialSum;

        loopMode = mode;
        Conv2dTilling tilling = GetConv2dTiling<half>(conv2dParams);

        roundm = tilling.roundM;
        roundn = tilling.roundN;
        roundk = tilling.roundK;

        m = tilling.mNum;
        n = tilling.nNum;

        featureMapA1Size = conv2dParams.cin * conv2dParams.imgShape[0] * conv2dParams.imgShape[1];
        weightB1Size = conv2dParams.cin * conv2dParams.cout * conv2dParams.kernelShape[0] * conv2dParams.kernelShape[1];
        biasSize = conv2dParams.cout;
        dstCO1Size = roundm * roundn;
        dstUBSize = m * n;

        fmGlobal.SetGlobalBuffer((__gm__ half*)fmGm);
        weGlobal.SetGlobalBuffer((__gm__ half*)weGm);
        biasGlobal.SetGlobalBuffer((__gm__ float*)biasGm);
        dstGlobal.SetGlobalBuffer((__gm__ float*)dstGm);

        pipe.InitBuffer(inQueueFmA1, 1, featureMapA1Size * sizeof(half));
        pipe.InitBuffer(inQueueWeB1, 1, weightB1Size * sizeof(half));
        pipe.InitBuffer(inQueueBiasA1, 1, biasSize * sizeof(float));
        pipe.InitBuffer(outQueueCO1, 1, dstCO1Size * sizeof(float));
        pipe.InitBuffer(outQueueUB, 1, dstUBSize * sizeof(float));
    }
    __aicore__ inline void Process()
    {
        CopyIn();
        Compute();
        CopyUB();
        CopyOut();
    }

private:
    __aicore__ inline void CopyIn()
    {
        LocalTensor<half> featureMapA1 = inQueueFmA1.AllocTensor<half>();
        LocalTensor<half> weightB1 = inQueueWeB1.AllocTensor<half>();
        LocalTensor<float> biasA1 = inQueueBiasA1.AllocTensor<float>();

        DataCopy(featureMapA1, fmGlobal, featureMapA1Size);
        DataCopy(weightB1, weGlobal, weightB1Size);
        DataCopy(biasA1, biasGlobal, biasSize);

        inQueueFmA1.EnQue(featureMapA1);
        inQueueWeB1.EnQue(weightB1);
        inQueueBiasA1.EnQue(biasA1);
    }

    __aicore__ inline void Compute()
    {
        LocalTensor<half> featureMapA1 = inQueueFmA1.DeQue<half>();
        LocalTensor<half> weightB1 = inQueueWeB1.DeQue<half>();
        LocalTensor<float> biasA1 = inQueueBiasA1.DeQue<float>();
        LocalTensor<float> dstCO1 = outQueueCO1.AllocTensor<float>();

        Conv2dTilling tilling = GetConv2dTiling<half>(conv2dParams);
        tilling.loopMode = loopMode;
        // 输入矩阵形状为[C1, H, W, C0],特征矩阵的形状为[C1, Kh, Kw, Cout, C0], 计算结果搬出至GM，目的矩阵无需初始化
        Conv2D(dstCO1, biasA1, featureMapA1, weightB1, conv2dParams, tilling);

        outQueueCO1.EnQue<float>(dstCO1);
        inQueueFmA1.FreeTensor(featureMapA1);
        inQueueWeB1.FreeTensor(weightB1);
        inQueueBiasA1.FreeTensor(biasA1);
    }

    __aicore__ inline void CopyUB() 
    {
        LocalTensor<float> dstCO1 = outQueueCO1.DeQue<float>();
        LocalTensor<float> dstUB = outQueueUB.AllocTensor<float>();

        DataCopyParams dataCopyParams;
        dataCopyParams.blockCount = 1;
        dataCopyParams.blockLen = dstCO1Size * sizeof(float) / 1024;
        DataCopyEnhancedParams enhancedParams;
        enhancedParams.blockMode = BlockMode::BLOCK_MODE_MATRIX;

        DataCopy(dstUB, dstCO1, dataCopyParams, enhancedParams);

        outQueueUB.EnQue<float>(dstUB);
        outQueueCO1.FreeTensor(dstCO1);
    }

    __aicore__ inline void CopyOut()
    {
        LocalTensor<float> dstUB = outQueueUB.DeQue<float>();
        DataCopy(dstGlobal, dstUB, dstUBSize);
        outQueueUB.FreeTensor(dstUB);
    }

private:
    TPipe pipe;
    // feature map queue
    TQue<QuePosition::A1, 1> inQueueFmA1;
    // weight queue
    TQue<QuePosition::B1, 1> inQueueWeB1;
    // bias queue
    TQue<QuePosition::A1, 1> inQueueBiasA1;
    // dst queue
    TQue<QuePosition::CO1, 1> outQueueCO1;

    TQue<QuePosition::VECOUT, 1> outQueueUB;

    GlobalTensor<half> fmGlobal, weGlobal;
    GlobalTensor<float> biasGlobal, dstGlobal;

    LoopMode loopMode = LoopMode::ModeNm;
    Conv2dParams conv2dParams;

    uint32_t m, n;
    uint32_t roundm, roundn, roundk;
    uint32_t featureMapA1Size, weightB1Size, biasSize, dstCO1Size, dstUBSize;
};
} // namespace AscendC

extern "C" __global__ __aicore__ void kernel_conv2d_operator(__gm__ half* fmGm, __gm__ half* weGm, __gm__ float* biasGm,
    __gm__ float* dstGm, Conv2dParams params, LoopMode mode)
{   
   // 上方示例参数为：
   // imgShape[0] = 4; imgShape[1] = 4; kernelShape[0] = 2; kernelShape[1] = 2; stride[0] = 1; stride[1] = 1;
   // cin = 32; cout = 16; padList[0] = 0; padList[1] = 0; padList[2] = 0; padList[3] = 0;
   // dilation[0] = 2; dilation[1] = 2; initY = 0; partialSum = false;
    AscendC::KernelCubeConv2D op;
    op.Init(fmGm, weGm, biasGm, dstGm, params, mode);
    op.Process();
}