DataCopyPad(ISASI)

Function Usage

Enables data non-aligned movement. The supported channels are as follows:

GM->VECIN/VECOUT

VECIN/VECOUT->GM

VECIN/VECOUT->TSCM

During data movement from GM to VECIN/VECOUT, you can pad the data as required.

Prototype

• dataCopyParams is of the DataCopyExtParams type. It supports a larger value range for parameters such as the operand stride compared to the DataCopyParams type.
  • Path: 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)

  • Path: VECIN/VECOUT -> GM

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

  • Channel: VECIN/VECOUT-> TSCM. The actual channel is VECIN/VECOUT->GM->TSCM.

    The Atlas 200I/500 A2 inference products does not support this API prototype.

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

• dataCopyParams is of the DataCopyParams type.
  • Path: 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)

  • Path: VECIN/VECOUT -> GM

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

  • Channel: VECIN/VECOUT-> TSCM. The actual channel is VECIN/VECOUT->GM->TSCM.

    The Atlas 200I/500 A2 inference products does not support this API prototype.

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

Parameters

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;
};

 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;
};

• GM -> VECIN/VECOUT
  • Description
    • When blockLen + leftPadding + rightPadding meets the 32-byte alignment requirement, if isPad is false, the values padded on the left and right sides are random numbers by default; if not, the values padded on the left and right sides are configured padding values.
    • When blockLen + leftPadding + rightPadding does not meet the 32-byte alignment requirement, the framework pads some dummy data to ensure the 32-byte alignment. When both leftPadding and rightPadding are 0, the first element value of the data chunk to be transferred is padded by default. When leftPadding or rightPadding is not 0, if isPad is set to false, the data values padded on the left and right sides and dummy values are random numbers; if not, the data values padded on the left and right sides and dummy values are configured padding values.
  • Configuration example 1
    • The value of blockLen is 64, meaning each data chunk contains 64 bytes. The value of srcStride is 1. Because the logical position of the source operand is GM and the unit of srcStride is byte, there is 1-byte interval between adjacent data chunks of the source operand. The value of dstStride is 1. Because the logical position of the destination operand is VECIN/VECOUT and the unit of dstStride is a 32-byte data block, there is an interval of one data block between adjacent data chunks of the destination operand.
    • When blockLen + leftPadding + rightPadding meets the 32-byte alignment requirement, if isPad is set to false, the data values padded on the left and right sides are random numbers by default; if not, the data values padded on the left and right sides are configured padding values. In this example, both leftPadding and rightPadding are 0, indicating that padding is not performed.
    • When blockLen + leftPadding + rightPadding does not meet the 32-byte alignment requirement, the framework pads some dummy data to ensure the 32-byte alignment. When leftPadding or rightPadding is not 0, if isPad is set to false, the data values padded on the left and right sides and dummy values are random numbers; if not, the data values padded on the left and right sides are configured padding values.

    

  • Configuration example 2
    • The value of blockLen is 47, meaning each data chunk contains 47 bytes. The value of srcStride is 1, indicating that there is 1-byte interval between adjacent data chunks of the source operand. The value of dstStride is 1, indicating that there is an interval of one data block between adjacent data chunks of the destination operand.
    • When blockLen + leftPadding + rightPadding does not meet the 32-byte alignment requirement, and both leftPadding and rightPadding are 0, the first element value of the data chunk to be transferred is padded with dummy values by default.
    • When blockLen + leftPadding + rightPadding does not meet the 32-byte alignment requirement, and leftPadding or rightPadding is not 0, if isPad is false, the data values padded on the left and right sides and dummy values are random numbers; if not, the data values padded on the left and right sides and dummy values are configured padding values.

    

Returns

None

Availability

Atlas A2 training products / Atlas A2 inference products

Atlas A3 training products / Atlas A3 inference products

Atlas 200I/500 A2 inference products

Constraints

The values of leftPadding and rightPadding cannot exceed 32 bytes.

Example

 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}; // The last parameter of DataCopyExtParams is reserved.
        AscendC::DataCopyPadExtParams<half> padParams{true, 0, 2, 0};
        AscendC::DataCopyPad(srcLocal, srcGlobal, copyParams, padParams); // Move 40 bytes from GM to VECIN.
        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); // Move 40 bytes from VECIN to GM.
        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();
}

Input data (src0Global): [1 2 3 ... 32]
Output data (dstGlobal):[1 2 3 ... 20]