文档
注册
评分
提单
论坛
小AI

LoadData

功能说明

LoadData分为Load2D和Load3D,其功能分别如下:

  • Load2D:

    源操作数/目的操作数的数据类型为uint8_t/int8_t时,分形矩阵大小在A1/A2上为16*32, 在B1/B2上为32*16。

    源操作数/目的操作数的数据类型为uint16_t/int16_t/half/bfloat16_t时,分形矩阵在A1/B1/A2/B2上的大小为16*16。

    源操作数/目的操作数的数据类型为uint32_t/int32_t/float时,分形矩阵大小在A1/A2上为16*8, 在B1/B2上为8*16。

    支持如下数据通路:

    GM->A1; GM->B1; GM->A2; GM->B2;A1->A2; B1->B2。

  • Load3D:

    image to column操作,将多维feature map转为二维矩阵。

    支持如下数据通路:

    A1->A2; B1->B2。

函数原型

  • Load2D接口
    1
    2
    3
    4
    template <typename T>
    void LoadData(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal, const LoadData2DParams& loadDataParam)
    template <typename T> 
    void LoadData(const LocalTensor<T>& dstLocal, const GlobalTensor<T>& srcLocal, const LoadData2DParams& loadDataParam)
    
  • Load3Dv1接口
    1
    2
    template <typename T, const IsResetLoad3dConfig &defaultConfig = IS_RESER_LOAD3D_DEFAULT_CONFIG>
    LoadData(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal,const LoadData3DParamsV1<T>& loadDataParams)
    
  • Load3Dv2接口
    1
    2
    template <typename T, const IsResetLoad3dConfig &defaultConfig = IS_RESER_LOAD3D_DEFAULT_CONFIG>
    LoadData(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal,const LoadData3DParamsV2<T>& loadDataParams)
    

参数说明

表1 模板参数说明

参数名称

含义

T

源操作数和目的操作数的数据类型。

load2d接口:

Atlas 训练系列产品,支持的数据类型为:uint8_t, int8_t, uint16_t, int16_t, half。

Atlas推理系列产品AI Core,支持的数据类型为:uint8_t, int8_t, uint16_t, int16_t, half。

Atlas A2训练系列产品/Atlas 800I A2推理产品,支持数据类型:uint8_t, int8_t, uint16_t, int16_t, half, bfloat16_t, uint32_t, int32_t, float。

Atlas 200/500 A2推理产品, 支持数据类型:uint8_t, int8_t, uint16_t, int16_t, half, bfloat16_t, uint32_t, int32_t, float。

load3dv1接口:

Atlas 训练系列产品,支持的数据类型为:uint8_t, int8_t, half。

Atlas推理系列产品AI Core,支持的数据类型为:uint8_t, int8_t, half。

load3dv2接口:

Atlas推理系列产品AI Core,支持的数据类型为:uint8_t, int8_t, half。

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

  • TPosition为A1/A2时,支持数据类型:uint8_t, int8_t, half, bfloat16_t, uint32_t, int32_t, float, int4b_t;
  • TPosition为B1/B2时,支持数据类型:half, bfloat16_t, uint32_t, int32_t, float。

Atlas 200/500 A2推理产品

  • TPosition为A1/A2时,支持数据类型:uint8_t, int8, half, bfloat16, uint32_t, int32_t, float, int4b_t;
  • TPosition为B1/B2时,支持数据类型:half, bfloat16_t, uint32_t, int32_t, float。

defaultConfig

控制是否在Load3Dv1/Load3Dv2接口内部设置相关属性。 IsResetLoad3dConfig类型。IsResetLoad3dConfig结构定义如下:

1
2
3
4
struct IsResetLoad3dConfig {
   bool isSetFMatrix = true;
   bool isSetPadding = true;
}; 

isSetFMatrix配置为true,表示在接口内部设置FeatureMap的属性描述(包括l1H、l1W、padList,参数介绍参考表4 LoadData3DParamsV1结构体内参数说明表5 LoadData3DParamsV2结构体内参数说明);设置为false,表示该接口传入的FeatureMap的属性描述不生效,开发者需要通过SetFmatrix进行设置。

isSetPadding配置为true, 表示在接口内部设置Pad属性描述(即padValue参数,参数介绍参考表4 LoadData3DParamsV1结构体内参数说明表5 LoadData3DParamsV2结构体内参数说明);设置为false,表示该接口传入的Pad属性不生效,开发者需要通过SetLoadDataPaddingValue进行设置。可参考样例调用示例

该参数的默认值如下:

1
constexpr IsResetLoad3dConfig IS_RESER_LOAD3D_DEFAULT_CONFIG = {true, true};
表2 通用参数说明

参数名称

输入/输出

含义

dstLocal

输出

目的操作数,类型为LocalTensor。

数据连续排列顺序由目的操作数所在TPosition决定,具体约束如下:

  • A2:小Z大Z格式;
  • B2:小N大Z格式;
  • A1/B1:无格式要求,一般情况下为小Z大N格式。

srcLocal

输入

源操作数,类型为LocalTensor或GlobalTensor。数据类型需要与dstLocal保持一致。

loadDataParams

输入

LoadData参数结构体,类型为:

  • LoadData2DParams,具体参考表3
  • LoadData3DParamsV1,具体参考表4
  • LoadData3DParamsV2,具体参考表5
表3 LoadData2DParams结构体内参数说明

参数名称

输入/输出

含义

startIndex

输入

分形矩阵ID,说明搬运起始位置为源操作数中第几个分形(0为源操作数中第1个分形矩阵)。取值范围:startIndex∈[0, 65535] 。单位:512B。默认为0。

repeatTimes

输入

迭代次数,每个迭代可以处理512B数据。取值范围:repeatTimes∈[1, 255]。

srcStride

输入

相邻迭代间,源操作数前一个分形与后一个分形起始地址的间隔,单位:512B。取值范围:src_stride∈[0, 65535]。默认为0。

sid

输入

预留参数,配置为0即可。

dstGap

输入

相邻迭代间,目的操作数前一个分形结束地址与后一个分形起始地址的间隔,单位:512B。取值范围:dstGap∈[0, 65535]。默认为0。

注:Atlas 训练系列产品此参数不使能。

ifTranspose

输入

是否启用转置功能,对每个分形矩阵进行转置,默认为false:

  • true:启用
  • false:不启用

注意:只有A1->A2和B1->B2通路才能使能转置,使能转置功能时,源操作数、目的操作数仅支持uint16_t, int16_t, half数据类型。

addrMode

输入

预留参数,配置为0即可。

表4 LoadData3DParamsV1结构体内参数说明

参数名称

输入/输出

含义

padList

输入

padding列表 [padding_left, padding_right, padding_top, padding_bottom],每个元素取值范围:[0,255]。默认为{0, 0, 0, 0}。

l1H

输入

源操作数 height,取值范围:l1H∈[1, 32767]。

l1W

输入

源操作数 width,取值范围:l1W∈[1, 32767] 。

c1Index

输入

该指令在源tensor C1维度的起点,取值范围:c1Index∈[0, 4095] 。默认为0。

fetchFilterW

输入

该指令在卷积核上w维度的起始位置,取值范围:fetchFilterW∈[0, 254] 。默认为0。

fetchFilterH

输入

该指令在filter上h维度的起始位置,取值范围:fetchFilterH∈[0, 254] 。默认为0。

leftTopW

输入

该指令在源操作数上w维度的起点,取值范围:leftTopW∈[-255, 32767] 。默认为0。如果padding_left = a,leftTopW配置为-a。

leftTopH

输入

该指令在源操作数上h维度的起点,取值范围:leftTopH∈[-255, 32767] 。默认为0。如果padding_top = a,leftTopH配置为-a。

strideW

输入

卷积核在源操作数w维度滑动的步长,取值范围:strideW∈[1, 63] 。

strideH

输入

卷积核在源操作数h维度滑动的步长,取值范围:strideH∈[1, 63] 。

filterW

输入

卷积核width,取值范围:filterW∈[1, 255] 。

filterH

输入

卷积核height,取值范围:filterH∈[1, 255] 。

dilationFilterW

输入

卷积核width膨胀系数,取值范围:dilationFilterW∈[1, 255] 。

dilationFilterH

输入

卷积核height膨胀系数,取值范围:dilationFilterH∈[1, 255] 。

jumpStride

输入

迭代之间,目的操作数首地址步长,取值范围:jumpStride∈[1, 127] 。

repeatMode

输入

迭代模式。
  • 模式0:每次迭代,增加卷积核窗口中的点,对应在目的矩阵上往w维度方向增长。
  • 模式1:每次迭代,增加滑动窗口左上坐标,对应在目的矩阵上往h维度方向增长。

取值范围:repeatMode∈[0, 1] 。默认为0。

repeatTime

输入

迭代次数,每一次源操作数和目的操作数的地址都会改变。取值范围:repeatTime∈[1,255] 。

cSize

输入

配置是否开启cSize = 4(b16) / cSize = 8(b8)优化,取值范围:cSize∈[0, 1] 。默认为0。

padValue

输入

Pad填充值的数值,数据类型需要与srcLocal保持一致。默认为0。若不想使能padding,可将padList设为全0。

表5 LoadData3DParamsV2结构体内参数说明

参数名称

输入/输出

含义

padList

输入

padding 列表 [padding_left, padding_right, padding_top, padding_bottom],每个元素取值范围:[0,255]。默认为{0, 0, 0, 0}。

l1H

输入

源操作数height,取值范围:l1H∈[1, 32767]。

l1W

输入

源操作数weight,取值范围:l1W∈[1, 32767] 。

channelSize

输入

源操作数的通道数,取值范围:channelSize∈[1, 63] 。

针对以下型号,channelSize的取值要求为:对于float16/half/bfloat16,channelSize可取值为4,8,16,N * 16 + 4,N * 16 + 8;对于int8_t/uint8_t,channelSize可取值为4,8,16,32,N * 32 + 4,N * 32 + 8,N * 32 + 16。N为正整数。

Atlas推理系列产品AI Core

针对以下型号,channelSize的取值要求为:对于uint32_t/int32_t/float32,channelSize可取值为4,N * 8,N * 8 + 4;对于float16/half/bfloat16,channelSize可取值为4,8,N * 16,N * 16 + 4,N * 16 + 8;对于int8_t/uint8_t,channelSize可取值为4,8,16, 32 * N,N * 32 + 4,N * 32 + 8,N * 32 + 16;对于int4b_t,ChannelSize可取值为8,16,32,N * 64,N * 64 + 8,N * 64 + 16,N * 64 + 32。N为正整数。

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

Atlas 200/500 A2推理产品

kExtension

输入

该指令在目的操作数width维度的传输长度,如果不覆盖最右侧的分形,对于half类型,应为16的倍数,对于int8_t/uint8_t应为32的倍数。取值范围: kExtension∈[1, 65535] 。

mExtension

输入

该指令在目的操作数height维度的传输长度,如果不覆盖最下侧的分形,对于half/int8_t/uint8_t,应为16的倍数。取值范围:mExtension∈[1, 65535] 。

kStartPt

输入

该指令在目的操作数width维度的起点,对于half类型,应为16的倍数,对于int8_t/uint8_t应为32的倍数。取值范围[0, 65535] 。默认为0。

mStartPt

输入

该指令在目的操作数height维度的起点,如果不覆盖最下侧的分形,对于half/int8_t/uint8_t,应为16的倍数。取值范围[0, 65535] 。默认为0。

strideW

输入

卷积核在源操作数width维度滑动的步长,取值范围:strideW∈[1, 63] 。

strideH

输入

卷积核在源操作数height 维度滑动的步长,取值范围:strideH∈[1, 63] 。

filterW

输入

卷积核width,取值范围:filterW∈[1, 255] 。

filterH

输入

卷积核height,取值范围:filterH∈[1, 255] 。

dilationFilterW

输入

卷积核width膨胀系数,取值范围:dilationFilterW∈[1, 255] 。

dilationFilterH

输入

卷积核height膨胀系数,取值范围:dilationFilterH∈[1, 255] 。

enTranspose

输入

是否启用转置功能,对整个目标矩阵进行转置,支持数据类型为 bool,仅在目的QuePosition为A2,且源操作数为half类型时有效。默认为false。

  • true:启用
  • false:不启用

enSmallK

输入

是否使能small k特性,每个分形矩阵大小为16*4,支持数据类型为 bool,默认为false。当前产品形态,该特性已不再支持。

  • true:使能
  • false:不使能

padValue

输入

Pad填充值的数值,数据类型需要与srcLocal保持一致。默认为0。若不想使能padding,可将padList设为全0。

支持的型号

  • Load2D接口

    Atlas 训练系列产品

    Atlas推理系列产品AI Core

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

    Atlas 200/500 A2推理产品

  • Load3Dv1

    Atlas 训练系列产品

    Atlas推理系列产品AI Core

  • Load3Dv2

    Atlas推理系列产品AI Core

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

    Atlas 200/500 A2推理产品

注意事项

  • 操作数地址偏移对齐要求请参见通用约束
  • LoadData3DParamsV1 cSize特性的开启,需要保证A1/B1中的 feature map 为 4 channel对齐的。
  • 不使用或者不想改变的配置,建议保持默认值,有助于性能提升。

load3d数据格式说明

要求输入的feature map和filter的格式是 NC1HWC0,其中 C0 是最低维度而且 C0 是固定值为 16(对于u8/s8类型为32),C1=C/C0。

为了简化场景,以下场景假设输入的 feature map 的 channel 为4,即 Ci=4。输入 feature maps 在 A1 中的形状为 (Hi,Wi,Ci),经过 load3dv1 处理后在 A2 的数据形状为(Wo*Ho, Hk*Wk*Ci)。其中 Wo 和 Ho 是卷积后输出的shape,Hk 和 Wk 是 filter 的 shape。

直观的来看,img2col 的过程就是 filter 在 feature map 上扫过,将对应 feature map 的数据展开成输出数据的每一行的过程。filter 首先在W方向上滑动 Wo 步,然后在 H 方向上走一步然后重复以上过程,最终输出 Wo*Ho 行数据。下图中红色和黄色的数据分别代表第一行和第二行。数字表示原始输入数据,filter 和输出数据三者之间的关联关系。可以看到,load3dv1 首先在输入数据的 Ci 维度搬运对应于 00 的 4 个数,然后搬运对应于 01 的四个数,最终这一行的大小为 Hk*Wk*Ci 即 3*3*4=36 个数。

对应的feature map格式如下图:

对应的 filter 的格式如下图:

其中 n 为 filter 的个数,可以看出维度排布为 (Hk,Wk,Ci,n),但是需要注意的是下图的格式还需要根据 mmad 中 B 矩阵的格式转换。

实际操作中,由于存储空间或者计算能力限制,我们通常会将整个卷积计算分块,一次只搬运并计算一小块数据。

对于 A2 的 feature map 来说有两种方案,水平分块和垂直分块。分别对应参数中 repeatMode 的 0 和 1。

注:下图中的分型矩阵大小为 4x4,实际应该为 16x16 (对于 u8/s8 类型为 16x32)

repeatMode =0 时,每次 repeat 会改变在 filter 窗口中读取数据点的位置,然后跳到下一个 C0 的位置。

repeatMode =1 的时候 filter 窗口中读取数据的位置保持不变,每个 repeat 在 feature map 中前进 C0 个元素。

返回值

调用示例

该调用示例支持的运行平台为Atlas推理系列产品AI Core

  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
 65
 66
 67
 68
 69
 70
 71
 72
 73
 74
 75
 76
 77
 78
 79
 80
 81
 82
 83
 84
 85
 86
 87
 88
 89
 90
 91
 92
 93
 94
 95
 96
 97
 98
 99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
#include "kernel_operator.h"

class KernelLoadData {
public:
    __aicore__ inline KernelLoadData()
    {
        coutBlocks = (Cout + 16 - 1) / 16;
        ho = (H + padTop + padBottom - dilationH * (Kh - 1) - 1) / strideH + 1;
        wo = (W + padLeft + padRight - dilationW * (Kw - 1) - 1) / strideW + 1;
        howo = ho * wo;
        howoRound = ((howo + 16 - 1) / 16) * 16;
        featureMapA1Size = C1 * H * W * C0;      // shape: [C1, H, W, C0]
        weightA1Size = C1 * Kh * Kw * Cout * C0; // shape: [C1, Kh, Kw, Cout, C0]
        featureMapA2Size = howoRound * (C1 * Kh * Kw * C0);
        weightB2Size = (C1 * Kh * Kw * C0) * coutBlocks * 16;
        m = howo;
        k = C1 * Kh * Kw * C0;
        n = Cout;
        dstSize = coutBlocks * howo * 16; // shape: [coutBlocks, howo, 16]
        dstCO1Size = coutBlocks * howoRound * 16;
        fmRepeat = featureMapA2Size / (16 * C0);
        weRepeat = weightB2Size / (16 * C0);
    }
    __aicore__ inline void Init(__gm__ uint8_t* fmGm, __gm__ uint8_t* weGm, __gm__ uint8_t* dstGm)
    {
        fmGlobal.SetGlobalBuffer((__gm__ half*)fmGm);
        weGlobal.SetGlobalBuffer((__gm__ half*)weGm);
        dstGlobal.SetGlobalBuffer((__gm__ half*)dstGm);
        pipe.InitBuffer(inQueueFmA1, 1, featureMapA1Size * sizeof(half));
        pipe.InitBuffer(inQueueFmA2, 1, featureMapA2Size * sizeof(half));
        pipe.InitBuffer(inQueueWeB1, 1, weightA1Size * sizeof(half));
        pipe.InitBuffer(inQueueWeB2, 1, weightB2Size * sizeof(half));
        pipe.InitBuffer(outQueueCO1, 1, dstCO1Size * sizeof(float));
        pipe.InitBuffer(outQueueUB, 1, dstSize * sizeof(half));
    }
    __aicore__ inline void Process()
    {
        CopyIn();
        Split();
        Compute();
        CopyUB();
        CopyOut();
    }

private:
    __aicore__ inline void CopyIn()
    {
        AscendC::LocalTensor<half> featureMapA1 = inQueueFmA1.AllocTensor<half>();
        AscendC::LocalTensor<half> weightB1 = inQueueWeB1.AllocTensor<half>();
        AscendC::DataCopy(featureMapA1, fmGlobal, { 1, static_cast<uint16_t>(featureMapA1Size * sizeof(half) / 32), 0, 0 });
        AscendC::DataCopy(weightB1, weGlobal, { 1, static_cast<uint16_t>(weightA1Size * sizeof(half) / 32), 0, 0 });
        inQueueFmA1.EnQue(featureMapA1);
        inQueueWeB1.EnQue(weightB1);
    }
    __aicore__ inline void Split()
    {
        AscendC::LocalTensor<half> featureMapA1 = inQueueFmA1.DeQue<half>();
        AscendC::LocalTensor<half> weightB1 = inQueueWeB1.DeQue<half>();
        AscendC::LocalTensor<half> featureMapA2 = inQueueFmA2.AllocTensor<half>();
        AscendC::LocalTensor<half> weightB2 = inQueueWeB2.AllocTensor<half>();
        uint8_t padList[4] = {padTop, padBottom, padLeft, padRight};
        AscendC::LoadData(featureMapA2, featureMapA1,
            { padList, H, W, 0, 0, 0, -1, -1, strideW, strideH, Kw, Kh, dilationW, dilationH, 1, 0, fmRepeat, 0, (half)(0)});
        AscendC::LoadData(weightB2, weightB1, { 0, weRepeat, 1, 0, 0, false, 0 });
        inQueueFmA2.EnQue<half>(featureMapA2);
        inQueueWeB2.EnQue<half>(weightB2);
        inQueueFmA1.FreeTensor(featureMapA1);
        inQueueWeB1.FreeTensor(weightB1);
    }
    __aicore__ inline void Compute()
    {
        AscendC::LocalTensor<half> featureMapA2 = inQueueFmA2.DeQue<half>();
        AscendC::LocalTensor<half> weightB2 = inQueueWeB2.DeQue<half>();
        AscendC::LocalTensor<float> dstCO1 = outQueueCO1.AllocTensor<float>();
        AscendC::Mmad(dstCO1, featureMapA2, weightB2, { m, n, k, 0, false, true });
        outQueueCO1.EnQue<float>(dstCO1);
        inQueueFmA2.FreeTensor(featureMapA2);
        inQueueWeB2.FreeTensor(weightB2);
    }
    __aicore__ inline void CopyUB()
    {
        AscendC::LocalTensor<float> dstCO1 = outQueueCO1.DeQue<float>();
        AscendC::LocalTensor<half> dstUB = outQueueUB.AllocTensor<half>();
        AscendC::DataCopyParams dataCopyParams;
        dataCopyParams.blockCount = 1;
        dataCopyParams.blockLen = m * n * sizeof(float) / 1024;
        AscendC::DataCopyEnhancedParams enhancedParams;
        enhancedParams.blockMode = AscendC::BlockMode::BLOCK_MODE_MATRIX;
        AscendC::DataCopy(dstUB, dstCO1, dataCopyParams, enhancedParams);
        outQueueUB.EnQue<half>(dstUB);
        outQueueCO1.FreeTensor(dstCO1);
    }
    __aicore__ inline void CopyOut()
    {
        AscendC::LocalTensor<half> dstUB = outQueueUB.DeQue<half>();
        AscendC::DataCopy(dstGlobal, dstUB, m * n);
        outQueueUB.FreeTensor(dstUB);
    }

private:
    AscendC::TPipe pipe;
    // feature map queue
    AscendC::TQue<AscendC::QuePosition::A1, 1> inQueueFmA1;
    AscendC::TQue<AscendC::QuePosition::A2, 1> inQueueFmA2;
    // weight queue
    AscendC::TQue<AscendC::QuePosition::B1, 1> inQueueWeB1;
    AscendC::TQue<AscendC::QuePosition::B2, 1> inQueueWeB2;
    // dst queue
    AscendC::TQue<AscendC::QuePosition::CO1, 1> outQueueCO1;
    AscendC::TQue<AscendC::QuePosition::CO2, 1> outQueueUB;
    AscendC::GlobalTensor<half> fmGlobal, weGlobal, dstGlobal;
    uint16_t C1 = 2;
    uint16_t H = 4, W = 4;
    uint8_t Kh = 2, Kw = 2;
    uint16_t Cout = 16;
    uint16_t C0 = 16;
    uint8_t dilationH = 2, dilationW = 2;
    uint8_t padTop = 1, padBottom = 1, padLeft = 1, padRight = 1;
    uint8_t strideH = 1, strideW = 1;
    uint16_t coutBlocks, ho, wo, howo, howoRound;
    uint32_t featureMapA1Size, weightA1Size, featureMapA2Size, weightB2Size, dstSize, dstCO1Size;
    uint16_t m, k, n;
    uint8_t fmRepeat, weRepeat;
};

extern "C" __global__ __aicore__ void load_data_simple_kernel(__gm__ uint8_t* fmGm, __gm__ uint8_t* weGm,
    __gm__ uint8_t* dstGm)
{
    KernelLoadData op;
    op.Init(fmGm, weGm, dstGm);
    op.Process();
}
搜索结果
找到“0”个结果

当前产品无相关内容

未找到相关内容,请尝试其他搜索词