Slice Data Transfer

Applicability

Function

Supports slice data transfer, which extracts subsets of multi-dimensional tensor data for transfer.

Prototype

• Global Memory -> Local Memory

  1 2	template <typename T> __aicore__ inline void DataCopy(const LocalTensor<T>& dst, const GlobalTensor<T>& src, const SliceInfo dstSliceInfo[], const SliceInfo srcSliceInfo[], const uint32_t dimValue = 1)

• Local Memory -> Global Memory

  1 2	template <typename T> __aicore__ inline void DataCopy(const GlobalTensor<T> &dst, const LocalTensor<T> &src, const SliceInfo dstSliceInfo[], const SliceInfo srcSliceInfo[], const uint32_t dimValue = 1)

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

Parameters

The preceding parameters are parsed by using specific examples. The following figure shows the details.

Figure 1 Parameter parsing

• dimValue of 2 indicates that the operand has 2 dimensions.
• srcSliceInfo is {{16, 70, 7, 3, 87}, {0, 2, 1, 1, 3}}.
  • {16, 70, 7, 3, 87} configures parameters for a single row in the one-dimensional view, where each value corresponds to a specific parameter:

    startIndex = 16 indicates that the valid data segment starts from the 16th element.

    endIndex = 70 indicates that the valid data segment ends at the 70th element.

    stride = 7, in units of elements, indicates that there are 7 zero-valued elements between two adjacent sliced data segments.

    burstLen = 3, in units of 32 bytes, indicates that in this valid data segment, the size of a sliced data segment is 3 data blocks.

    shapeValue = 87 indicates the total length of the single row in elements, which equals 8 × 10 + 7 = 87 elements.

  • {0, 2, 1, 1, 3} configures parameters for multiple rows in the two-dimensional view, where each value corresponds to a row-related parameter:

    startIndex = 0 indicates that the valid data segment starts from row 0.

    endIndex = 2 indicates that the valid data segment ends at row 2.

    stride = 1 indicates that there is one row between two adjacent sliced data segments.

    burstLen = 1, which is mandatory when dimValue is greater than 1.

    shapeValue = 3 indicates that there are three rows in total.

• dstSliceInfo is {{0, 47, 0, 3, 48}, {0, 1, 0, 1, 2}}.
  • {0, 47, 0, 3, 48} configures parameters for a single row in the one-dimensional view, where each value corresponds to a specific parameter:

    startIndex = 0 indicates that the valid data segment starts from the 0th element.

    endIndex = 47 indicates that the valid data segment ends at the 47th element.

    stride = 0, in units of elements, indicates that there is no gap between two adjacent sliced data segments.

    burstLen = 3, in units of 32 bytes, indicates that in this valid data segment, the size of a sliced data segment is 3 data blocks.

    shapeValue = 48 indicates the total length of the single row in elements, which equals 8 × 6 = 48 elements.

  • {0, 1, 0, 1, 2} configures parameters for multiple rows in the two-dimensional view, where each value corresponds to a row-related parameter:

    startIndex = 0 indicates that the valid data segment starts from row 0.

    endIndex = 1 indicates that the valid data segment ends at row 1.

    stride = 0 indicates that there is no gap between two adjacent sliced data segments.

    burstLen = 1, which is mandatory when dimValue is greater than 1.

    shapeValue = 2 indicates that there are two rows in total.

Returns

None

Restrictions

• When configuring the horizontal burstLen for slice data transfer, calculate it using the formula: Number of horizontal slice elements × sizeof(T)/32 bytes. The product of the number of horizontal slice elements and sizeof(T) must be a multiple of 32 bytes.
• The array size of the SliceInfo structure must match dimValue and must not exceed 8.
• The array sizes of srcSliceInfo and dstSliceInfo must be identical, and their respective burstLen values must be equal, that is, srcSliceInfo[i].burstLen = dstSliceInfo[i].burstLen.
• Slice data transfer has certain requirements on parameters. You are advised to refer to the calling examples and verify correctness through CPU simulation before executing on the NPU.

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
56
57
58
59
60
61
62
63
64

#include "kernel_operator.h"
// In this example, the tensor data type is float.
template <typename T>
class KernelDataCopySliceGM2UB {
public:
    __aicore__ inline KernelDataCopySliceGM2UB()
    {}
    __aicore__ inline void Init(__gm__ uint8_t* dstGm, __gm__ uint8_t* srcGm)
    {
        AscendC::SliceInfo srcSliceInfoIn[] = {{16, 70, 7, 3, 87}, {0, 2, 1, 1, 3}};// startIndex is 16, endIndex is 70, burstLen is 3, stride is 7, and shapeValue is 87.
        AscendC::SliceInfo dstSliceInfoIn[] = {{0, 47, 0, 3, 48}, {0, 1, 0, 1, 2}};// UB memory is relatively limited. You are advised to set stride to 0.
        uint32_t dimValueIn = 2;
        uint32_t dstDataSize = 96;
        uint32_t srcDataSize = 261;
        dimValue = dimValueIn;
        for (uint32_t i = 0; i < dimValueIn; i++) {
            srcSliceInfo[i].startIndex = srcSliceInfoIn[i].startIndex;
            srcSliceInfo[i].endIndex = srcSliceInfoIn[i].endIndex;
            srcSliceInfo[i].stride = srcSliceInfoIn[i].stride;
            srcSliceInfo[i].burstLen = srcSliceInfoIn[i].burstLen;
            srcSliceInfo[i].shapeValue = srcSliceInfoIn[i].shapeValue;
            dstSliceInfo[i].startIndex = dstSliceInfoIn[i].startIndex;
            dstSliceInfo[i].endIndex = dstSliceInfoIn[i].endIndex;
            dstSliceInfo[i].stride = dstSliceInfoIn[i].stride;
            dstSliceInfo[i].burstLen = dstSliceInfoIn[i].burstLen;
            dstSliceInfo[i].shapeValue = dstSliceInfoIn[i].shapeValue;
        }
        srcGlobal.SetGlobalBuffer((__gm__ T *)srcGm);
        dstGlobal.SetGlobalBuffer((__gm__ T *)dstGm);
        pipe.InitBuffer(inQueueSrcVecIn, 1, dstDataSize * sizeof(T));
    }
    __aicore__ inline void Process()
    {
        CopyIn();
        CopyOut();
    }
private:
    __aicore__ inline void CopyIn()
    {
        AscendC::LocalTensor<T> srcLocal = inQueueSrcVecIn.AllocTensor<T>();
        AscendC::DataCopy(srcLocal, srcGlobal,  dstSliceInfo, srcSliceInfo, dimValue);
        inQueueSrcVecIn.EnQue(srcLocal);
    }
    __aicore__ inline void CopyOut()
    {
        AscendC::LocalTensor<T> srcOutLocal = inQueueSrcVecIn.DeQue<T>();
        AscendC::DataCopy(dstGlobal, srcOutLocal, dstSliceInfo, dstSliceInfo, dimValue);
        inQueueSrcVecIn.FreeTensor(srcOutLocal);
    }
private:
    AscendC::TPipe pipe;
    AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueueSrcVecIn;
    AscendC::GlobalTensor<T> srcGlobal;
    AscendC::GlobalTensor<T> dstGlobal;
    AscendC::SliceInfo dstSliceInfo[K_MAX_DIM];
    AscendC::SliceInfo srcSliceInfo[K_MAX_DIM]; // K_MAX_DIM = 8
    uint32_t dimValue;
};
extern "C" __global__ __aicore__ void kernel_data_copy_slice_out2ub(__gm__ uint8_t* src_gm, __gm__ uint8_t* dst_gm)
{
    KernelDataCopySliceGM2UB<TYPE> op;
    op.Init(dst_gm, src_gm);
    op.Process();
}

For details about the result example, see Figure 1.