NZ2ND Transfer with On-the-Fly Conversion

Applicability

Function

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

Prototype

1
2

template <typename T>
__aicore__ inline void DataCopy(const GlobalTensor<T>& dst, const LocalTensor<T>& src, const Nz2NdParamsFull& 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 NZ2ND conversion of the half type. The parameter settings in the example are described as follows:

• ndNum = 2 indicates that the number of source NZ matrices is 2 (NZ matrix 1 is A1–A4 + B1–B4, and NZ matrix 2 is C1–C4 + D1–D4)
• nValue = 4 indicates the number of rows in the NZ matrix, that is, the height of the matrix is 4.
• dValue = 32 indicates the number of columns in the NZ matrix, that is, the width of the matrix is 32 elements.
• srcNdMatrixStride = 1 indicates the offset between the start addresses of adjacent NZ matrices, that is, the distance between A1 and C1 (16 data blocks × 16 elements = 256 elements).
• srcNStride = 4 indicates the offset between adjacent Z-tiles in the same source NZ matrix, that is, the distance between A1 and B1 (4 data blocks × 16 elements per block = 64 elements).
• dstDStride = 160 indicates the offset between adjacent rows in a destination ND matrix, that is, the distance between A1 and A2, which is 10 data blocks (10 × 16 = 160 elements).
• dstNdMatrixStride = 48 indicates the offset between the start points of the x th and (x+1) th destination ND matrices in dst, that is, the distance between A1 and C1, which is three data blocks (3 × 16 = 48 elements).

Figure 1 NZ2ND conversion (half type)

The following figure shows the NZ2ND conversion of the float type. The parameter settings in the example are described as follows:

• ndNum = 2 indicates that the number of source NZ matrices is 2 (NZ matrix 1 is A1–A8 + B1–B8, and NZ matrix 2 is C1–C8 + D1–D8).
• nValue = 4 indicates the number of rows in the NZ matrix, that is, the height of the matrix is 4.
• dValue = 32 indicates the number of columns in the NZ matrix, that is, the width of the matrix is 32 elements.
• srcNdMatrixStride = 1 indicates the offset between the start addresses of adjacent NZ matrices, that is, the distance from A1 to C1, which is 256 elements (32 data blocks × 8 elements per block).
• srcNStride = 4 indicates the offset between adjacent Z-tiles in the same source NZ matrix, that is, the distance between A1 and B1, which is 64 elements (8 data blocks × 8 elements per block).
• dstDStride = 144 indicates the offset between adjacent rows in a destination ND matrix, that is, the distance between A1 and A3, which is 18 data blocks (18 × 8 = 144 elements).
• dstNdMatrixStride = 40 indicates the offset between the start points of the x th and (x+1) th destination ND matrices in dst, that is, the distance between A1 and C1, which is five data blocks (5 × 8 = 40 elements).

Figure 2 NZ2ND conversion (float type)

Returns

None

Restrictions

None

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 KernelDataCopyUb2GmNz2Nd {
public:
    __aicore__ inline KernelDataCopyUb2GmNz2Nd()
    {}
    __aicore__ inline void Init(__gm__ uint8_t* dstGm, __gm__ uint8_t* srcGm)
    {
        AscendC::Nz2NdParamsFull intriParamsIn{1, 32, 32, 1, 32, 32, 1};
        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.nValue * intriParams.dValue);
        inQueueSrcVecIn.EnQue(srcLocal);
    }
    __aicore__ inline void Compute()
    {
        AscendC::LocalTensor<half> dstLocal = inQueueSrcVecIn.DeQue<half>();
        AscendC::LocalTensor<half> srcOutLocal = inQueueSrcVecOut.AllocTensor<half>();
        AscendC::DataCopy(srcOutLocal, dstLocal, intriParams.nValue * intriParams.dValue);
        inQueueSrcVecOut.EnQue(srcOutLocal);
        inQueueSrcVecIn.FreeTensor(dstLocal);
    }
    __aicore__ inline void CopyOut()
    {
        AscendC::LocalTensor<half> srcOutLocalDe = inQueueSrcVecOut.DeQue<half>();
        AscendC::DataCopy(dstGlobal, srcOutLocalDe, intriParams);
        inQueueSrcVecOut.FreeTensor(srcOutLocalDe);
    }
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::Nz2NdParamsFull intriParams;
};
extern "C" __global__ __aicore__ void kernel_data_copy_nz2nd_ub2out(__gm__ uint8_t* src_gm, __gm__ uint8_t* dst_gm)
{
    KernelDataCopyUb2GmNz2Nd op;
    op.Init(dst_gm, src_gm);
    op.Process();
}

Result example:

Input (srcGlobal): [1 2 3 ... 1024]
Output (dstGlobal): [1 2 ... 15 16 513 514 ... 527 528 17 18 ... 31 32 529 530 ... 543 544 ...497 498 ...  511 512  1009 1010... 1023 1024]