ND2NZ Transfer with On-the-Fly Conversion

Applicability

Function

Supports ND-to-NZ format conversion during data transfer.

Prototype

• Global Memory -> Local Memory

  1 2	template <typename T> __aicore__ inline void DataCopy(const LocalTensor<T>& dst, const GlobalTensor<T>& src, const Nd2NzParams& intriParams)

• Local Memory -> Local Memory

  1 2	template <typename T> __aicore__ inline void DataCopy(const LocalTensor<T>& dst, const LocalTensor<T>& src, const Nd2NzParams& intriParams)

For details about supported transfer paths and data types of each prototype, see Supported Paths and Data Types.

Parameters

The following figure shows the ND2NZ conversion. The parameter settings in the example are described as follows:

• ndNum = 2 indicates that the number of ND matrices to transfer is 2 (ND matrix 1 is A1 to A2 + B1 to B2, and ND matrix 2 is C1 to C2 + D1 to D2).
• nValue = 2 indicates that the number of rows in the ND matrix is 2, that is, the height of the matrix is 2.
• dValue = 24 indicates that the number of columns in the ND matrix is 24, that is, the width of the matrix is 24 elements. If dValue is not 32-byte aligned, the insufficient part in the destination operand is padded with 0s. For example, the blank part in the data block where A2 is located is padded with 0s.
• srcNdMatrixStride = 144 indicates the offset between the start addresses of adjacent ND matrices, that is, the distance between A1 and C1, which is nine data blocks (9 × 16 = 144 elements).
• srcDValue = 48 indicates the number of elements in a row, that is, the distance between A1 and B1, which is three data blocks (3 × 16 = 48 elements).
• dstNzC0Stride = 11 indicates that after conversion from ND to NZ format, a single row in the source operand is split into multiple rows in the destination operand. For example, A1 and A2 that belong to one row in src are divided into two separate rows in dst. The offset between the start addresses of these rows corresponds to the offset between A1 and A2 in dst, which is 11 data blocks.
• dstNzNStride = 2 indicates the offset between the xth and (x+1)th rows of an ND matrix in src after conversion to the NZ format in dst. That is, the offset between A1 and B1 in dst is two data blocks.
• dstNzMatrixStride = 96 indicates the offset between the start points of the x th and (x+1) th ND matrices in dst, that is, the distance between A1 and C1, which is six data blocks (6 × 16 = 96 elements).

Figure 1 ND2NZ conversion (half type)

Returns

None

Restrictions

For the Atlas inference product's AI Core, when the ND2NZ transfer API of the Global Memory -> Local Memory path is used, 8 KB UB space must be reserved as the temporary data storage area of the API.

Supported Paths and Data Types

The following data paths are expressed using logical positions TPosition, with the corresponding physical paths noted. For details about the mapping between TPosition and the physical memory, see Table 1.

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

#include "kernel_operator.h"
class KernelDataCopyGm2UbNd2Nz{
public:
    __aicore__ inline KernelDataCopyGm2UbNd2Nz()
    {}
    __aicore__ inline void Init(__gm__ uint8_t* dstGm, __gm__ uint8_t* srcGm)
    {
        AscendC::Nd2NzParams intriParamsIn{1, 32, 32, 0, 32, 32, 1, 0};
        intriParams = intriParamsIn;
        srcGlobal.SetGlobalBuffer((__gm__ half *)srcGm);
        dstGlobal.SetGlobalBuffer((__gm__ half *)dstGm);
        pipe.InitBuffer(inQueueSrcVecIn, 1, intriParams.nValue * intriParams.dValue * sizeof(half));
        pipe.InitBuffer(inQueueSrcVecOut, 1, intriParams.nValue * intriParams.dValue * sizeof(half));
    }
    __aicore__ inline void Process()
    {
        CopyIn();
        Compute();
        CopyOut();
    }
private:
    __aicore__ inline void CopyIn()
    {
        AscendC::LocalTensor<half> srcLocal = inQueueSrcVecIn.AllocTensor<half>();
        AscendC::DataCopy(srcLocal, srcGlobal, intriParams);
        inQueueSrcVecIn.EnQue(srcLocal);
    }
    __aicore__ inline void Compute()
    {
        AscendC::LocalTensor<half> srcLocal = inQueueSrcVecIn.DeQue<half>();
        AscendC::LocalTensor<half> dstLocal = inQueueSrcVecOut.AllocTensor<half>();
        AscendC::DataCopy(dstLocal, srcLocal, intriParams.nValue * intriParams.dValue);
        inQueueSrcVecOut.EnQue(dstLocal);
        inQueueSrcVecIn.FreeTensor(srcLocal);
    }
    __aicore__ inline void CopyOut()
    {
        AscendC::LocalTensor<half> dstLocal = inQueueSrcVecOut.DeQue<half>();
        AscendC::DataCopy(dstGlobal, dstLocal, intriParams.nValue * intriParams.dValue);
        inQueueSrcVecOut.FreeTensor(dstLocal);
    }
private:
    AscendC::TPipe pipe;
    AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueueSrcVecIn;
    AscendC::TQue<AscendC::TPosition::VECOUT, 1> inQueueSrcVecOut;
    AscendC::GlobalTensor<half> srcGlobal;
    AscendC::GlobalTensor<half> dstGlobal;
    AscendC::Nd2NzParams intriParams;
};
extern "C" __global__ __aicore__ void kernel_data_copy_nd2nz_ub2out(__gm__ uint8_t* src_gm, __gm__ uint8_t* dst_gm)
{
    KernelDataCopyGm2UbNd2Nz op;
    op.Init(dst_gm, src_gm);
    op.Process();
}

Result example:

Input (srcGlobal): [1 2 3 ... 1024]
Output (dstGlobal): [1 2 ... 15 16 33 34 ... 47 48 65 66 ... 79 80 97 98 ... 111 112 ... 1009 1010... 1023 1024]