LoadData

功能说明

LoadData分为Load2D和Load3D，其功能分别如下：

• Load2D：

  源操作数/目的操作数的数据类型为uint8_t/int8_t时，分形矩阵大小在A1/A2上为16*32， 在B1/B2上为32*16。

  源操作数/目的操作数的数据类型为uint16_t/int16_t/half/bfloat16_t时，分形矩阵在A1/B1/A2/B2上的大小为16*16。

  源操作数/目的操作数的数据类型为uint32_t/int32_t/float时，分形矩阵大小在A1/A2上为16*8， 在B1/B2上为8*16。

  支持如下数据通路：

  GM->A1; GM->B1; GM->A2; GM->B2;A1->A2; B1->B2。

• Load3D：

  image to column操作，将多维feature map转为二维矩阵。

  支持如下数据通路：

  A1->A2; B1->B2。

函数原型

• Load2D接口

  template <typename T> void LoadData(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal,const LoadData2DParams& loadDataParam);
  template <typename T> void LoadData(const LocalTensor<T>& dstLocal, const GlobalTensor<T>& srcLocal,const LoadData2DParams& loadDataParam);

• Load3Dv1接口

  template <typename T, const IsResetLoad3dConfig &defaultConfig = IS_RESER_LOAD3D_DEFAULT_CONFIG> LoadData(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal,const LoadData3DParamsV1<T>& loadDataParams);

• Load3Dv2接口

  template <typename T, const IsResetLoad3dConfig &defaultConfig = IS_RESER_LOAD3D_DEFAULT_CONFIG> LoadData(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal,const LoadData3DParamsV2<T>& loadDataParams);

参数说明

struct IsResetLoad3dConfig {
   bool isSetFMatrix = true;
   bool isSetPadding = true;
}; 

支持的型号

Atlas 训练系列产品

Atlas推理系列产品AI Core

Atlas A2训练系列产品/Atlas 800I A2推理产品

注意事项

• 操作数地址偏移对齐要求请参见通用约束。
• LoadData3DParamsV1 cSize 特性的开启，需要保证 A1/B1 中的 feature map 为 4 channel 对齐的。
• 不使用或者不想改变的配置，建议保持默认值，有助于性能提升。
• Atlas 训练系列产品不支持Load3Dv2接口。

load3d数据格式说明

要求输入的feature map和filter的格式是 NC1HWC0，其中 C0 是最低维度而且 C0 是固定值为 16（对于u8/s8类型为32），C1=C/C0。

为了简化场景，以下场景假设输入的 feature map 的 channel 为4，即 Ci=4。输入 feature maps 在 A1 中的形状为 (Hi,Wi,Ci)，经过 load3dv1 处理后在 A2 的数据形状为(Wo*Ho, Hk*Wk*Ci)。其中 Wo 和 Ho 是卷积后输出的shape，Hk 和 Wk 是 filter 的 shape。

直观的来看，img2col 的过程就是 filter 在 feature map 上扫过，将对应 feature map 的数据展开成输出数据的每一行的过程。filter 首先在W方向上滑动 Wo 步,然后在 H 方向上走一步然后重复以上过程，最终输出 Wo*Ho 行数据。下图中红色和黄色的数据分别代表第一行和第二行。数字表示原始输入数据，filter 和输出数据三者之间的关联关系。可以看到，load3dv1 首先在输入数据的 Ci 维度搬运对应于 00 的 4 个数，然后搬运对应于 01 的四个数，最终这一行的大小为 Hk*Wk*Ci 即 3*3*4=36 个数。

对应的feature map格式如下图：

对应的 filter 的格式如下图：

其中 n 为 filter 的个数，可以看出维度排布为 (Hk,Wk,Ci,n)，但是需要注意的是下图的格式还需要根据 mmad 中 B 矩阵的格式转换。

实际操作中，由于存储空间或者计算能力限制，我们通常会将整个卷积计算分块，一次只搬运并计算一小块数据。

对于 A2 的 feature map 来说有两种方案，水平分块和垂直分块。分别对应参数中 repeatMode 的 0 和 1。

注：下图中的分型矩阵大小为 4x4，实际应该为 16x16 (对于 u8/s8 类型为 16x32)

repeatMode =0 时，每次 repeat 会改变在 filter 窗口中读取数据点的位置，然后跳到下一个 C0 的位置。

repeatMode =1 的时候 filter 窗口中读取数据的位置保持不变，每个 repeat 在 feature map 中前进 C0 个元素。

返回值

无

调用示例

该调用示例支持的运行平台为Atlas推理系列产品AI Core。

#include "kernel_operator.h"
namespace AscendC {
class KernelLoadData {
public:
    __aicore__ inline KernelLoadData()
    {
        coutBlocks = (Cout + 16 - 1) / 16;
        ho = (H + padTop + padBottom - dilationH * (Kh - 1) - 1) / strideH + 1;
        wo = (W + padLeft + padRight - dilationW * (Kw - 1) - 1) / strideW + 1;
        howo = ho * wo;
        howoRound = ((howo + 16 - 1) / 16) * 16;
        featureMapA1Size = C1 * H * W * C0;      // shape: [C1, H, W, C0]
        weightA1Size = C1 * Kh * Kw * Cout * C0; // shape: [C1, Kh, Kw, Cout, C0]
        featureMapA2Size = howoRound * (C1 * Kh * Kw * C0);
        weightB2Size = (C1 * Kh * Kw * C0) * coutBlocks * 16;
        m = howo;
        k = C1 * Kh * Kw * C0;
        n = Cout;
        dstSize = coutBlocks * howo * 16; // shape: [coutBlocks, howo, 16]
        dstCO1Size = coutBlocks * howoRound * 16;
        fmRepeat = featureMapA2Size / (16 * C0);
        weRepeat = weightB2Size / (16 * C0);
    }
    __aicore__ inline void Init(__gm__ uint8_t* fmGm, __gm__ uint8_t* weGm, __gm__ uint8_t* dstGm)
    {
        fmGlobal.SetGlobalBuffer((__gm__ half*)fmGm);
        weGlobal.SetGlobalBuffer((__gm__ half*)weGm);
        dstGlobal.SetGlobalBuffer((__gm__ half*)dstGm);
        pipe.InitBuffer(inQueueFmA1, 1, featureMapA1Size * sizeof(half));
        pipe.InitBuffer(inQueueFmA2, 1, featureMapA2Size * sizeof(half));
        pipe.InitBuffer(inQueueWeB1, 1, weightA1Size * sizeof(half));
        pipe.InitBuffer(inQueueWeB2, 1, weightB2Size * sizeof(half));
        pipe.InitBuffer(outQueueCO1, 1, dstCO1Size * sizeof(float));
        pipe.InitBuffer(outQueueUB, 1, dstSize * sizeof(half));
    }
    __aicore__ inline void Process()
    {
        CopyIn();
        Split();
        Compute();
        CopyUB();
        CopyOut();
    }
private:
    __aicore__ inline void CopyIn()
    {
        LocalTensor<half> featureMapA1 = inQueueFmA1.AllocTensor<half>();
        LocalTensor<half> weightB1 = inQueueWeB1.AllocTensor<half>();
        DataCopy(featureMapA1, fmGlobal, { 1, static_cast<uint16_t>(featureMapA1Size * sizeof(half) / 32), 0, 0 });
        DataCopy(weightB1, weGlobal, { 1, static_cast<uint16_t>(weightA1Size * sizeof(half) / 32), 0, 0 });
        inQueueFmA1.EnQue(featureMapA1);
        inQueueWeB1.EnQue(weightB1);
    }
    __aicore__ inline void Split()
    {
        LocalTensor<half> featureMapA1 = inQueueFmA1.DeQue<half>();
        LocalTensor<half> weightB1 = inQueueWeB1.DeQue<half>();
        LocalTensor<half> featureMapA2 = inQueueFmA2.AllocTensor<half>();
        LocalTensor<half> weightB2 = inQueueWeB2.AllocTensor<half>();
        uint8_t padList[4] = {padTop, padBottom, padLeft, padRight};
        LoadData(featureMapA2, featureMapA1,
            { padList, H, W, 0, 0, 0, -1, -1, strideW, strideH, Kw, Kh, dilationW, dilationH, 1, 0, fmRepeat, 0, (half)(0)});
        LoadData(weightB2, weightB1, { 0, weRepeat, 1, 0, 0, false, 0 });
        inQueueFmA2.EnQue<half>(featureMapA2);
        inQueueWeB2.EnQue<half>(weightB2);
        inQueueFmA1.FreeTensor(featureMapA1);
        inQueueWeB1.FreeTensor(weightB1);
    }
    __aicore__ inline void Compute()
    {
        LocalTensor<half> featureMapA2 = inQueueFmA2.DeQue<half>();
        LocalTensor<half> weightB2 = inQueueWeB2.DeQue<half>();
        LocalTensor<float> dstCO1 = outQueueCO1.AllocTensor<float>();
        Mmad(dstCO1, featureMapA2, weightB2, { m, n, k, 0, false, true });
        outQueueCO1.EnQue<float>(dstCO1);
        inQueueFmA2.FreeTensor(featureMapA2);
        inQueueWeB2.FreeTensor(weightB2);
    }
    __aicore__ inline void CopyUB()
    {
        LocalTensor<float> dstCO1 = outQueueCO1.DeQue<float>();
        LocalTensor<half> dstUB = outQueueUB.AllocTensor<half>();
        DataCopyParams dataCopyParams;
        dataCopyParams.blockCount = 1;
        dataCopyParams.blockLen = m * n * sizeof(float) / 1024;
        DataCopyEnhancedParams enhancedParams;
        enhancedParams.blockMode = BlockMode::BLOCK_MODE_MATRIX;
        DataCopy(dstUB, dstCO1, dataCopyParams, enhancedParams);
        outQueueUB.EnQue<half>(dstUB);
        outQueueCO1.FreeTensor(dstCO1);
    }
    __aicore__ inline void CopyOut()
    {
        LocalTensor<half> dstUB = outQueueUB.DeQue<half>();
        DataCopy(dstGlobal, dstUB, m * n);
        outQueueUB.FreeTensor(dstUB);
    }
private:
    TPipe pipe;
    // feature map queue
    TQue<QuePosition::A1, 1> inQueueFmA1;
    TQue<QuePosition::A2, 1> inQueueFmA2;
    // weight queue
    TQue<QuePosition::B1, 1> inQueueWeB1;
    TQue<QuePosition::B2, 1> inQueueWeB2;
    // dst queue
    TQue<QuePosition::CO1, 1> outQueueCO1;
    TQue<QuePosition::CO2, 1> outQueueUB;
    GlobalTensor<half> fmGlobal, weGlobal, dstGlobal;
    uint16_t C1 = 2;
    uint16_t H = 4, W = 4;
    uint8_t Kh = 2, Kw = 2;
    uint16_t Cout = 16;
    uint16_t C0 = 16;
    uint8_t dilationH = 2, dilationW = 2;
    uint8_t padTop = 1, padBottom = 1, padLeft = 1, padRight = 1;
    uint8_t strideH = 1, strideW = 1;
    uint16_t coutBlocks, ho, wo, howo, howoRound;
    uint32_t featureMapA1Size, weightA1Size, featureMapA2Size, weightB2Size, dstSize, dstCO1Size;
    uint16_t m, k, n;
    uint8_t fmRepeat, weRepeat;
};
} // namespace AscendC
extern "C" __global__ __aicore__ void load_data_simple_kernel(__gm__ uint8_t* fmGm, __gm__ uint8_t* weGm,
    __gm__ uint8_t* dstGm)
{
    AscendC::KernelLoadData op;
    op.Init(fmGm, weGm, dstGm);
    op.Process();
}