DataCopyPad

功能说明

该接口提供数据非对齐搬运的功能，支持的数据传输通路如下：

GM->VECIN/VECOUT

VECIN/VECOUT->GM

VECIN/VECOUT->TSCM

其中从GM->VECIN/VECOUT进行数据搬运时，可以根据开发者的需要自行填充数据。

函数原型

• dataCopyParams为DataCopyExtParams类型，相比于DataCopyParams类型，支持的操作数步长等参数取值范围更大
  • 通路：GM->VECIN/VECOUT

    1 2	template <typename T> __aicore__ inline void DataCopyPad(const LocalTensor<T> &dstLocal, const GlobalTensor<T> &srcGlobal, const DataCopyExtParams &dataCopyParams, const DataCopyPadExtParams<T> &padParams)

  • 通路：VECIN/VECOUT->GM

    1 2	template <typename T> __aicore__ inline void DataCopyPad(const GlobalTensor<T> &dstGlobal, const LocalTensor<T> &srcLocal, const DataCopyExtParams &dataCopyParams)

  • 通路：VECIN/VECOUT->TSCM，实际搬运过程是VECIN/VECOUT->GM->TSCM

    1 2	template <typename T> __aicore__ inline void DataCopyPad(const LocalTensor<T> &dstLocal, const LocalTensor<T> &srcLocal, const DataCopyExtParams &dataCopyParams, const Nd2NzParams &nd2nzParams)

• dataCopyParams为DataCopyParams类型
  • 通路：GM->VECIN/VECOUT

    1 2	template<typename T> __aicore__ inline void DataCopyPad(const LocalTensor<T>& dstLocal, const GlobalTensor<T>& srcGlobal, const DataCopyParams& dataCopyParams, const DataCopyPadParams& padParams)

  • 通路：VECIN/VECOUT->GM

    1 2	template<typename T> __aicore__ inline void DataCopyPad(const GlobalTensor<T>& dstGlobal, const LocalTensor<T>& srcLocal,const DataCopyParams& dataCopyParams)

  • 通路：VECIN/VECOUT->TSCM，实际搬运过程是VECIN/VECOUT->GM->TSCM

    1 2	template<typename T> __aicore__ inline void DataCopyPad(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal, const DataCopyParams& dataCopyParams, const Nd2NzParams& nd2nzParams）

参数说明

template <typename T> struct DataCopyPadExtParams {
    bool isPad = false;
    uint8_t leftPadding = 0;
    uint8_t rightPadding = 0;
    T paddingValue = 0;
};

struct Nd2NzParams {
    uint16_t ndNum = 0;
    uint16_t nValue = 0;
    uint16_t dValue = 0;
    uint16_t srcNdMatrixStride = 0;
    uint16_t srcDValue = 0;
    uint16_t dstNzC0Stride = 0;
    uint16_t dstNzNStride = 0;
    uint16_t dstNzMatrixStride = 0;
};

• GM->VECIN/VECOUT
  • 参数解释
    • 当blockLen+leftPadding+rightPadding满足32字节对齐时，isPad为false，左右两侧填充的数据值会默认为随机值；否则为paddingValue。
    • 当blockLen+leftPadding+rightPadding不满足32字节对齐时，框架会填充一些假数据dummy，保证左右填充的数据和blockLen、假数据为32字节对齐。若leftPadding、rightPadding都为0：dummy会默认填充待搬运数据块的第一个元素值；若leftPadding/rightPadding不为0：isPad为false，左右两侧填充的数据值和dummy值均为随机值；否则为paddingValue。
  • 配置示例1：
    • blockLen为64，每个连续传输数据块包含64Bytes；srcStride为1，因为源操作数的逻辑位置为GM，srcStride的单位为Byte，也就是说源操作数相邻数据块之间间隔1Byte；dstStride为1，因为目的操作数的逻辑位置为VECIN/VECOUT，dstStride的单位为dataBlock(32Bytes)，也就是说目的操作数相邻数据块之间间隔1个dataBlock。
    • blockLen+leftPadding+rightPadding满足32字节对齐，isPad为false，左右两侧填充的数据值会默认为随机值；否则为paddingValue。此处示例中，leftPadding、rightPadding均为0，则不填充。
    • blockLen+leftPadding+rightPadding不满足32字节对齐时，框架会填充一些假数据dummy，保证左右填充的数据和blockLen、假数据为32字节对齐。leftPadding/rightPadding不为0：若isPad为false，左右两侧填充的数据值和dummy值均为随机值；否则为paddingValue。

    

  • 配置示例2：
    • blockLen为47，每个连续传输数据块包含47Bytes；srcStride为1，表示源操作数相邻数据块之间间隔1Byte；dstStride为1，表示目的操作数相邻数据块之间间隔1个dataBlock。
    • blockLen+leftPadding+rightPadding不满足32字节对齐，leftPadding、rightPadding均为0：dummy会默认填充待搬运数据块的第一个元素值。
    • blockLen+leftPadding+rightPadding不满足32字节对齐，leftPadding/rightPadding不为0：若isPad为false，左右两侧填充的数据值和dummy值均为随机值；否则为paddingValue。

    

返回值

无

支持的型号

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

约束说明

leftPadding、rightPadding的字节数均不能超过32Bytes。

调用示例

#include "kernel_operator.h"
namespace AscendC {
class TestDataCopyPad {
public:
    __aicore__ inline TestDataCopyPad()() {}
    __aicore__ inline void Init(__gm__ uint8_t* srcGm, __gm__ uint8_t* dstGm)
    {
        srcGlobal.SetGlobalBuffer((__gm__ half *)srcGm);
        dstGlobal.SetGlobalBuffer((__gm__ half *)dstGm);
        pipe.InitBuffer(inQueueSrc, 1, 32 * sizeof(half));
        pipe.InitBuffer(outQueueDst, 1, 32 * sizeof(half));
    }
    __aicore__ inline void Process()
    {
        CopyIn();
        Compute();
        CopyOut();
    }
private:
    __aicore__ inline void CopyIn()
    {
        LocalTensor<half> srcLocal = inQueueSrc.AllocTensor<half>();
        DataCopyExtParams copyParams{1, 20 * sizeof(half), 0, 0, 0}; // 结构体DataCopyExtParams最后一个参数是rsv保留位
        DataCopyPadExtParams<half> padParams{true, 0, 2, 0}; 
        DataCopyPad(srcLocal, srcGlobal, copyParams, padParams); // 从GM->VECIN搬运40Bytes
        inQueueSrc.EnQue<half>(srcLocal);
    }
    __aicore__ inline void Compute()
    {
        LocalTensor<half> srcLocal = inQueueSrc.DeQue<half>();
        LocalTensor<half> dstLocal = outQueueDst.AllocTensor<half>();
        Adds(dstLocal, srcLocal, scalar, 20);
        outQueueDst.EnQue(dstLocal);
        inQueueSrc.FreeTensor(srcLocal);

    }
    __aicore__ inline void CopyOut()
    {
        LocalTensor<half> dstLocal = outQueueDst.DeQue<half>();
        DataCopyExtParams copyParams{1, 20 * sizeof(half), 0, 0, 0};
        DataCopyPad(dstGlobal, ubLocal, copyParams); // 从VECIN->GM搬运40Bytes
        outQueueDst.FreeTensor(dstLocal);
    }
private:
    TPipe pipe;
    TQue<QuePosition::VECIN, 1> inQueueSrc;
    TQue<QuePosition::VECOUT, 1> outQueueDst;
    GlobalTensor<half> srcGlobal;
    GlobalTensor<half> dstGlobal;
    DataCopyPadExtParams<half> padParams;
    DataCopyExtParams copyParams;
    half scalar = 0;
};
}
extern "C" __global__ __aicore__ void kernel_data_copy_pad_kernel(__gm__ uint8_t* src_gm, __gm__ uint8_t* dst_gm)
{
    AscendC::TestDataCopyPad op;
    op.Init(src_gm, dst_gm);
    op.Process();
}

输入数据(src0Global): [1 2 3 ... 32]
输出数据(dstGlobal):[1 2 3 ... 20]