下载
中文
注册

DataCopyPad

功能说明

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

GM->VECIN/VECOUT

VECIN/VECOUT->GM

VECIN/VECOUT->TSCM

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

函数原型

  • 通路:GM->VECIN/VECOUT

    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

    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

    template <typename T>

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

以下接口不推荐使用,新开发内容不要使用以下接口。

template<typename T>

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

template<typename T>

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

template<typename T>

__aicore__ inline void DataCopyPad(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal, const DataCopyParams& dataCopyParams, const Nd2NzParams& nd2nzParams

参数说明

表1 接口参数说明

参数名

输入/输出

描述

dstLocal, dstGlobal

输出

目的操作数,类型为LocalTensor或GlobalTensor。

Atlas A2训练系列产品/Atlas 800I A2推理产品,支持的数据类型为:half/int16_t/uint16_t/float/int32_t/uint32_t/int8_t/uint8_t

srcLocal, srcGlobal

输入

源操作数,类型为LocalTensor或GlobalTensor。

Atlas A2训练系列产品/Atlas 800I A2推理产品,支持的数据类型为:half/int16_t/uint16_t/float/int32_t/uint32_t/int8_t/uint8_t

dataCopyParams

输入

搬运参数。DataCopyExtParams类型,定义如下:

struct DataCopyExtParams {
    uint16_t blockCount = 0;
    uint32_t blockLen = 0;
    uint32_t srcStride = 0;
    uint32_t dstStride = 0;
    uint32_t rsv = 0; // reserved information
};

DataCopyExtParams具体参数说明请参考表2

padParams

输入

从GM->VECIN/VECOUT进行数据搬运时,可以根据开发者需要,在搬运数据左边或右边填充数据。padParams是用于控制数据填充过程的参数,DataCopyPadExtParams类型,定义如下:

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

DataCopyPadExtParams结构定义请参考表3

nd2nzParams

输入

搬运参数。Nd2NzParams类型,Nd2NzParams结构定义请参考表7

注意:Nd2NzParams的ndNum仅支持设置为1

表2 DataCopyExtParams结构体参数定义

参数名称

含义

blockCount

指定该指令包含的连续传输数据块个数,数据类型为uint16_t,取值范围:blockCount∈[1, 4095]。

blockLen

指定该指令每个连续传输数据块长度,该指令支持非对齐搬运每个连续传输数据块长度单位为Byte。数据类型为uint32_t,blockLen不要超出该数据类型的取值范围。

srcStride

源操作数,相邻连续数据块的间隔(前面一个数据块的尾与后面数据块的头的间隔),如果源操作数的逻辑位置为VECIN/VECOUT,则单位为dataBlock(32Bytes), 如果源操作数的逻辑位置为GM,则单位为Byte。数据类型为uint32_t,srcStride不要超出该数据类型的取值范围。

dstStride

目的操作数,相邻连续数据块间的间隔(前面一个数据块的尾与后面数据块的头的间隔),如果目的操作数的逻辑位置为VECIN/VECOUT,则单位为dataBlock(32Bytes),如果目的操作数的逻辑位置为GM,则单位为Byte。数据类型为uint32_t,dstStride不要超出该数据类型的取值范围。

rsv

保留字段。

表3 DataCopyPadExtParams结构体参数定义

参数名称

含义

isPad

是否需要填充用户自定义的数据,取值范围:true,false。

true:填充padding value。

false:表示用户不需要指定填充值,会默认填充随机值。

leftPadding

连续搬运数据块左侧需要补充的数据范围,单位为元素个数。

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

rightPadding

连续搬运数据块右侧需要补充的数据范围,单位为元素个数。

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

paddingValue

左右两侧需要填充的数据值,需要保证在数据占用字节范围内。

数据类型和源操作数保持一致。

Atlas A2训练系列产品/Atlas 800I A2推理产品,支持的数据类型为:half/int16_t/uint16_t/float/int32_t/uint32_t/int8_t/uint8_t

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。

GM->VECIN/VECOUT情况下,配置示例如下:

  • 配置示例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。

VECIN/VECOUT->GM情况下,参数相关解释和配置示例如下:
  • 当每个连续传输数据块长度blockLen为32字节对齐时,下图呈现了需要传入的DataCopyParams示例,blockLen为64,每个连续传输数据块包含64Bytes;srcStride为1,因为源操作数的逻辑位置为VECIN/VECOUT,srcStride的单位为dataBlock(32Bytes),也就是说源操作数相邻数据块之间间隔1个dataBlock;dstStride为1,因为目的操作数的逻辑位置为GM,dstStride的单位为Byte,也就是说目的操作数相邻数据块之间间隔1Byte。

  • 当每个连续传输数据块长度blockLen不满足32字节对齐,由于Unified Buffer要求32字节对齐,框架在搬出时会自动补充一些假数据来保证对齐,但在当搬到GM时会自动将填充的假数据丢弃掉。下图呈现了该场景下需要传入的DataCopyParams示例和假数据补齐的原理。blockLen为47,每个连续传输数据块包含47Bytes,不满足32字节对齐;srcStride为1,表示源操作数相邻数据块之间间隔1个dataBlock;dstStride为1,表示目的操作数相邻数据块之间间隔1Byte。框架在搬出时会自动补充17Bytes的假数据来保证对齐,搬到GM时再自动将填充的假数据丢弃掉。

VECIN/VECOUT->TSCM情况下,参数相关解释和配置示例如下:

注意:因为内部实现涉及AIC和AIV之间的通信,实际搬运路径为VECIN / VECOUT -> GM -> TSCM,因此发送通信消息会有开销,性能会受到影响

如下图所示,展示了从VECIN / VECOUT 搬运到GM,再搬运到TSCM的数据流格式。以数据类型half为例,单个datablock(32B)含有16个half元素。从存储于VECIN / VECOUT 中的src来看,datablock1: A1~A6, datablock2: B1~B6, datablock3: C1~C6为需要进行DataCopy的datablock。

针对于从 VECIN / VECOUT -> GM的搬运,blockCount为3表明有3个数据块需要搬运; blockLen为一个数据块的大小(单位为Byte),为 6 * 32 = 192;srcStride表明src的第x个block的尾和第x+1个block的头的间隔,因为源操作数逻辑位置为VECIN / VECOUT所以单位为block, 例如针对block1和block2,中间相隔1个A7 block, 因此srtStride为1个block;dstStride表明dst的第x个block的尾与第x+1个block的头的间隔,因为目的操作数逻辑位置为GM所以单位为Byte,例如针对block1和block2,中间相隔2个空白的block,因此dstStride为64 Byte。数据存储格式没有发生转变,依然是ND。

针对于从 GM -> TSCM的搬运,ndNum固定为1,即3个block视作一整个ndMatrix;nValue表明ndMatrix的行数,即为3行,对应3个block;dValue表明ndMatrix中的其中一行所含的元素个数,因此为 6 * 16 = 96个元素;srcNdMatrixStride表明相邻ndMatrixStride之间的距离,因为仅涉及1个ndMatrix,所以可以填为0;srcDValue表明ndMatrix的第x行和第x+1行所相隔的元素个数,如A1~B1的距离,即为8个block 8 * 16 = 128个元素;dstNzC0Stride为src同一行的相邻block在NZ矩阵中相隔block数,如A1~A2的距离,即为7个block (A1 + 空 + B1 + 空 + C1 + 空*2);dstNzNStride为src中ndMatrix的相邻行在Nz矩阵中相隔多少个block,如A1~B1的距离,即为2个block (A1 + 空) ;dstNzMatrixStride为为相邻NZ矩阵之间的元素个数,因为仅涉及1个NZ矩阵,所以可以填为1。TSCM中的数据存储格式为NZ。

返回值

支持的型号

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

约束说明

leftPadding+rightPadding的字节数之和不能超过32Bytes。

调用示例

本演示示例实现了GM->VECIN->GM的非对齐搬运过程。
#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]