DataCopyPad

功能说明

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

GM->VECIN/VECOUT

VECIN/VECOUT->GM

VECIN/VECOUT->TSCM

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

函数原型

参数说明

表1 模板参数说明

参数名

描述

T

操作数以及paddingValue（待填充数据值）的数据类型。

Atlas A2训练系列产品/Atlas 800I A2推理产品，支持的数据类型为：half/bfloat16_t/int16_t/uint16_t/float/int32_t/uint32_t/int8_t/uint8_t/int64_t/uint64_t/double

Atlas 200/500 A2推理产品，支持的数据类型为：int8_t/uint8_t/half/bfloat16_t/int16_t/uint16_t/float/int32_t/uint32_t
表2 接口参数说明

参数名

输入/输出

描述

dstLocal/dstGlobal

输出

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

srcLocal/srcGlobal

输入

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

dataCopyParams

输入

搬运参数。

  • DataCopyExtParams类型，定义如下，具体参数说明请参考表3。
    1
2
3
4
5
6
7
    struct DataCopyExtParams {
    uint16_t blockCount = 0;
    uint32_t blockLen = 0;
    uint32_t srcStride = 0;
    uint32_t dstStride = 0;
    uint32_t rsv = 0; // reserved information
};
  • DataCopyParams类型，定义如下，具体参数说明请参考表4。
    1
2
3
4
5
6
    struct DataCopyParams {
    uint16_t blockCount = 0;
    uint16_t blockLen = 0;
    uint16_t srcStride = 0;
    uint16_t dstStride = 0;
};

padParams

输入

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

1
2
3
4
5
6
template <typename T> struct DataCopyPadExtParams {
    bool isPad = false;
    uint8_t leftPadding = 0;
    uint8_t rightPadding = 0;
    T paddingValue = 0;
};

nd2nzParams

输入

从VECIN/VECOUT->TSCM进行数据搬运时，可以进行ND到NZ的数据格式转换。nd2nzParams是用于控制数据格式转换的参数，Nd2NzParams类型，定义如下，具体参数请参考表7。

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
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;
};

注意：Nd2NzParams的ndNum仅支持设置为1。
表3 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

保留字段。
表4 DataCopyParams结构体参数定义

参数名称

含义

blockCount

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

blockLen

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

srcStride

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

dstStride

目的操作数，相邻连续数据块间的间隔（前面一个数据块的尾与后面数据块的头的间隔），如果目的操作数的逻辑位置为VECIN/VECOUT，则单位为dataBlock(32Bytes)，如果目的操作数的逻辑位置为GM，则单位为Byte。数据类型为uint16_t，dstStride不要超出该数据类型的取值范围。
表5 DataCopyPadExtParams结构体参数定义

参数名称

含义

isPad

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

true：填充padding value。

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

leftPadding

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

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

rightPadding

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

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

paddingValue

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

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

当数据类型长度为64位时，该参数只能设置为0。

返回值

无

支持的型号

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

Atlas 200/500 A2推理产品

约束说明

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

调用示例

本示例实现了GM->VECIN->GM的非对齐搬运过程。
 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
#include "kernel_operator.h"

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()
    {
        AscendC::LocalTensor<half> srcLocal = inQueueSrc.AllocTensor<half>();
        AscendC::DataCopyExtParams copyParams{1, 20 * sizeof(half), 0, 0, 0}; // 结构体DataCopyExtParams最后一个参数是rsv保留位
        AscendC::DataCopyPadExtParams<half> padParams{true, 0, 2, 0};
        AscendC::DataCopyPad(srcLocal, srcGlobal, copyParams, padParams); // 从GM->VECIN搬运40Bytes
        inQueueSrc.EnQue<half>(srcLocal);
    }
    __aicore__ inline void Compute()
    {
        AscendC::LocalTensor<half> srcLocal = inQueueSrc.DeQue<half>();
        AscendC::LocalTensor<half> dstLocal = outQueueDst.AllocTensor<half>();
        AscendC::Adds(dstLocal, srcLocal, scalar, 20);
        outQueueDst.EnQue(dstLocal);
        inQueueSrc.FreeTensor(srcLocal);
    }
    __aicore__ inline void CopyOut()
    {
        AscendC::LocalTensor<half> dstLocal = outQueueDst.DeQue<half>();
        AscendC::DataCopyExtParams copyParams{1, 20 * sizeof(half), 0, 0, 0};
        AscendC::DataCopyPad(dstGlobal, dstLocal, copyParams); // 从VECIN->GM搬运40Bytes
        outQueueDst.FreeTensor(dstLocal);
    }
private:
    AscendC::TPipe pipe;
    AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueSrc;
    AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueueDst;
    AscendC::GlobalTensor<half> srcGlobal;
    AscendC::GlobalTensor<half> dstGlobal;
    AscendC::DataCopyPadExtParams<half> padParams;
    AscendC::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)
{
    TestDataCopyPad op;
    op.Init(src_gm, dst_gm);
    op.Process();
}
结果示例：
输入数据(src0Global): [1 2 3 ... 32]
输出数据(dstGlobal):[1 2 3 ... 20]
父主题： 数据搬运