下载
中文
注册
我要评分
文档获取效率
文档正确性
内容完整性
文档易理解
在线提单
论坛求助
昇腾小AI

MrgSort4

功能说明

将已经排好序的最多4条region proposals队列,排列并合并成1条队列,结果按照score域由大到小排序。

函数原型

1
2
template <typename T>
__aicore__ inline void MrgSort4(const LocalTensor<T>& dstLocal, const MrgSortSrcList<T>& srcLocal, const MrgSort4Info& params)

参数说明

表1 接口参数说明

参数名称

输入/输出

含义

dstLocal

输出

目的操作数,存储经过排序后的Region Proposals。

类型为LocalTensor,支持的TPosition为VECIN/VECCALC/VECOUT。

LocalTensor的起始地址需要保证16字节对齐(针对half数据类型),32字节对齐(针对float数据类型)。

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

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

srcLocal

输入

源操作数,4个Region Proposals队列,并且每个Region Proposal队列都已经排好序,类型为MrgSortSrcList结构体,具体请参考表2

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

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

params

输入

排序所需参数,类型为MrgSort4Info结构体,具体请参考表3

表2 MrgSortSrcList参数说明

参数名称

输入/输出

含义

src1

输入

源操作数,第一个已经排好序的Region Proposals队列。

类型为LocalTensor,支持的TPosition为VECIN/VECCALC/VECOUT。

LocalTensor的起始地址需要保证16字节对齐(针对half数据类型),32字节对齐(针对float数据类型)。

数据类型与目的操作数保持一致。

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

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

src2

输入

源操作数,第二个已经排好序的Region Proposals队列。

类型为LocalTensor,支持的TPosition为VECIN/VECCALC/VECOUT。

LocalTensor的起始地址需要保证16字节对齐(针对half数据类型),32字节对齐(针对float数据类型)。

数据类型与目的操作数保持一致。

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

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

src3

输入

源操作数,第三个已经排好序的Region Proposals队列。

类型为LocalTensor,支持的TPosition为VECIN/VECCALC/VECOUT。

LocalTensor的起始地址需要保证16字节对齐(针对half数据类型),32字节对齐(针对float数据类型)。

数据类型与目的操作数保持一致。

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

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

src4

输入

源操作数,第四个已经排好序的Region Proposals队列。

类型为LocalTensor,支持的TPosition为VECIN/VECCALC/VECOUT。

LocalTensor的起始地址需要保证16字节对齐(针对half数据类型),32字节对齐(针对float数据类型)。

数据类型与目的操作数保持一致。

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

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

表3 MrgSort4Info参数说明

参数名称

输入/输出

含义

elementLengths

输入

四个源Region Proposals队列的长度(Region Proposal数目),类型为长度为4的uint16_t数据类型的数组,理论上每个元素取值范围[0, 4095],但不能超出UB的存储空间。

ifExhaustedSuspension

输入

某条队列耗尽后,指令是否需要停止,类型为bool,默认false。

validBit

输入

有效队列个数,取值如下:
  • 3:前两条队列有效
  • 7:前三条队列有效
  • 15:四条队列全部有效

repeatTimes

输入

迭代次数,每一次源操作数和目的操作数跳过四个队列总长度。取值范围:repeatTimes∈[1,255]。

repeatTimes参数生效是有条件的,需要同时满足以下四个条件:
  • 四个源Region Proposals队列的长度一致
  • 四个源Region Proposals队列连续存储
  • ifExhaustedSuspension = False
  • validBit=15

支持的型号

Atlas 训练系列产品

Atlas推理系列产品AI Core

约束说明

  • 当存在proposal[i]与proposal[j]的score值相同时,如果i>j,则proposal[j]将首先被选出来,排在前面。
  • 操作数地址偏移对齐要求请参见通用约束
  • 不支持源操作数与目的操作数之间存在地址重叠。

调用示例

  • 接口使用样例
    1
    2
    3
    4
    5
    // vconcatWorkLocal为已经创建并且完成排序的4个region proposal,每个Region Proposal数目是16个
    struct MrgSortSrcList<half> srcList(vconcatWorkLocal[0], vconcatWorkLocal[1], vconcatWorkLocal[2], vconcatWorkLocal[3]);
    uint16_t elementLengths[4] = {16, 16, 16, 16};
    struct MrgSort4Info srcInfo(elementLengths, false, 15, 1);
    AscendC::MrgSort4(dstLocal, srcList, srcInfo);
    
  • 完整样例
     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
    #include "kernel_operator.h"
    
    class KernelVecProposal {
    public:
        __aicore__ inline KernelVecProposal() {}
        __aicore__ inline void Init(__gm__ uint8_t* src, __gm__ uint8_t* dstGm)
        {
            srcGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ half*>(src), srcDataSize);
            dstGlobal.SetGlobalBuffer((__gm__ half*)dstGm);
    
            pipe.InitBuffer(inQueueSrc, 1, srcDataSize * sizeof(half));
            pipe.InitBuffer(workQueue, 1, dstDataSize * sizeof(half));
            pipe.InitBuffer(outQueueDst, 1, dstDataSize * sizeof(half));
        }
        __aicore__ inline void Process()
        {
            CopyIn();
            Compute();
            CopyOut();
        }
    
    private:
        __aicore__ inline void CopyIn()
        {
            AscendC::LocalTensor<half> srcLocal = inQueueSrc.AllocTensor<half>();
            AscendC::DataCopy(srcLocal, srcGlobal, srcDataSize);
            inQueueSrc.EnQue(srcLocal);
        }
        __aicore__ inline void Compute()
        {
            AscendC::LocalTensor<half> srcLocal = inQueueSrc.DeQue<half>();
            AscendC::LocalTensor<half> vconcatWorkLocal = workQueue.AllocTensor<half>();
            AscendC::LocalTensor<half> dstLocal = outQueueDst.AllocTensor<half>();
    
            // 先构造4个region proposal然后进行合并排序
            AscendC::ProposalConcat(vconcatWorkLocal[0], srcLocal[0], repeat, mode);
            AscendC::RpSort16(vconcatWorkLocal[0], vconcatWorkLocal[0], repeat);
    
            AscendC::ProposalConcat(vconcatWorkLocal[workDataSize], srcLocal[singleDataSize], repeat, mode);
            AscendC::RpSort16(vconcatWorkLocal[workDataSize], vconcatWorkLocal[workDataSize], repeat);
    
            AscendC::ProposalConcat(vconcatWorkLocal[workDataSize * 2], srcLocal[singleDataSize * 2], repeat, mode);
            AscendC::RpSort16(vconcatWorkLocal[workDataSize * 2], vconcatWorkLocal[workDataSize * 2], repeat);
    
            AscendC::ProposalConcat(vconcatWorkLocal[workDataSize * 3], srcLocal[singleDataSize * 3], repeat, mode);
            AscendC::RpSort16(vconcatWorkLocal[workDataSize * 3], vconcatWorkLocal[workDataSize * 3], repeat);
    
            AscendC::MrgSortSrcList<half> srcList(vconcatWorkLocal[0], vconcatWorkLocal[workDataSize],
                vconcatWorkLocal[workDataSize * 2], vconcatWorkLocal[workDataSize * 3]);
            uint16_t elementLengths[4] = {singleDataSize, singleDataSize, singleDataSize, singleDataSize};
            AscendC::MrgSort4Info srcInfo(elementLengths, false, 15, 1);
            AscendC::MrgSort4(dstLocal, srcList, srcInfo);
    
            outQueueDst.EnQue<half>(dstLocal);
            inQueueSrc.FreeTensor(srcLocal);
            workQueue.FreeTensor(vconcatWorkLocal);
        }
        __aicore__ inline void CopyOut()
        {
            AscendC::LocalTensor<half> dstLocal = outQueueDst.DeQue<half>();
            AscendC::DataCopy(dstGlobal, dstLocal, dstDataSize);
            outQueueDst.FreeTensor(dstLocal);
        }
    
    private:
        AscendC::TPipe pipe;
        AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueSrc;
        AscendC::TQue<AscendC::QuePosition::VECIN, 1> workQueue;
        AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueueDst;
        AscendC::GlobalTensor<half> srcGlobal, dstGlobal;
    
        int srcDataSize = 64;
        uint16_t singleDataSize = srcDataSize / 4;
        int dstDataSize = 512;
        int workDataSize = dstDataSize / 4;
        int repeat = srcDataSize / 4 / 16;
        int mode = 4;
    };
    
    extern "C" __global__ __aicore__ void vec_proposal_kernel(__gm__ uint8_t* src, __gm__ uint8_t* dstGm)
    {
        KernelVecProposal op;
        op.Init(src, dstGm);
        op.Process();
    }
    
    示例结果
    输入数据(src_gm):
    [-38.1    82.7   -40.75  -54.62   21.67  -58.53   25.94  -79.5   -61.44
      26.7   -27.45   48.78   86.75  -18.1   -58.8    62.38   46.38  -78.94
     -87.7   -13.81  -13.25   46.94  -47.8   -50.44   34.16   20.3    80.1
     -94.1    52.4   -42.75   83.4    80.44  -66.8   -82.7   -91.44  -95.6
      66.2   -30.97  -36.53   61.66   24.92  -45.1    38.97  -34.62  -69.8
      59.1    34.22   11.695 -33.47   52.1    -4.832  46.88   56.78   71.4
      13.29  -35.78   52.44  -46.03   83.8    83.56   71.3    -9.086 -65.06
      46.25 ]
    输出数据(dst_gm):
    [  0.      0.      0.      0.     86.75    0.      0.      0.      0.
       0.      0.      0.     83.8     0.      0.      0.      0.      0.
       0.      0.     83.56    0.      0.      0.      0.      0.      0.
       0.     83.4     0.      0.      0.      0.      0.      0.      0.
      82.7     0.      0.      0.      0.      0.      0.      0.     80.44
       0.      0.      0.      0.      0.      0.      0.     80.1     0.
       0.      0.      0.      0.      0.      0.     71.4     0.      0.
       0.      0.      0.      0.      0.     71.3     0.      0.      0.
       0.      0.      0.      0.     66.2     0.      0.      0.      0.
       0.      0.      0.     62.38    0.      0.      0.      0.      0.
       0.      0.     61.66    0.      0.      0.      0.      0.      0.
       0.     59.1     0.      0.      0.      0.      0.      0.      0.
      56.78    0.      0.      0.      0.      0.      0.      0.     52.44
       0.      0.      0.      0.      0.      0.      0.     52.4     0.
       0.      0.      0.      0.      0.      0.     52.1     0.      0.
       0.      0.      0.      0.      0.     48.78    0.      0.      0.
       0.      0.      0.      0.     46.94    0.      0.      0.      0.
       0.      0.      0.     46.88    0.      0.      0.      0.      0.
       0.      0.     46.38    0.      0.      0.      0.      0.      0.
       0.     46.25    0.      0.      0.      0.      0.      0.      0.
      38.97    0.      0.      0.      0.      0.      0.      0.     34.22
       0.      0.      0.      0.      0.      0.      0.     34.16    0.
       0.      0.      0.      0.      0.      0.     26.7     0.      0.
       0.      0.      0.      0.      0.     25.94    0.      0.      0.
       0.      0.      0.      0.     24.92    0.      0.      0.      0.
       0.      0.      0.     21.67    0.      0.      0.      0.      0.
       0.      0.     20.3     0.      0.      0.      0.      0.      0.
       0.     13.29    0.      0.      0.      0.      0.      0.      0.
      11.695   0.      0.      0.      0.      0.      0.      0.     -4.832
       0.      0.      0.      0.      0.      0.      0.     -9.086   0.
       0.      0.      0.      0.      0.      0.    -13.25    0.      0.
       0.      0.      0.      0.      0.    -13.81    0.      0.      0.
       0.      0.      0.      0.    -18.1     0.      0.      0.      0.
       0.      0.      0.    -27.45    0.      0.      0.      0.      0.
       0.      0.    -30.97    0.      0.      0.      0.      0.      0.
       0.    -33.47    0.      0.      0.      0.      0.      0.      0.
     -34.62    0.      0.      0.      0.      0.      0.      0.    -35.78
       0.      0.      0.      0.      0.      0.      0.    -36.53    0.
       0.      0.      0.      0.      0.      0.    -38.1     0.      0.
       0.      0.      0.      0.      0.    -40.75    0.      0.      0.
       0.      0.      0.      0.    -42.75    0.      0.      0.      0.
       0.      0.      0.    -45.1     0.      0.      0.      0.      0.
       0.      0.    -46.03    0.      0.      0.      0.      0.      0.
       0.    -47.8     0.      0.      0.      0.      0.      0.      0.
     -50.44    0.      0.      0.      0.      0.      0.      0.    -54.62
       0.      0.      0.      0.      0.      0.      0.    -58.53    0.
       0.      0.      0.      0.      0.      0.    -58.8     0.      0.
       0.      0.      0.      0.      0.    -61.44    0.      0.      0.
       0.      0.      0.      0.    -65.06    0.      0.      0.      0.
       0.      0.      0.    -66.8     0.      0.      0.      0.      0.
       0.      0.    -69.8     0.      0.      0.      0.      0.      0.
       0.    -78.94    0.      0.      0.      0.      0.      0.      0.
     -79.5     0.      0.      0.      0.      0.      0.      0.    -82.7
       0.      0.      0.      0.      0.      0.      0.    -87.7     0.
       0.      0.      0.      0.      0.      0.    -91.44    0.      0.
       0.      0.      0.      0.      0.    -94.1     0.      0.      0.
       0.      0.      0.      0.    -95.6     0.      0.      0.   ]
搜索结果
找到“0”个结果

当前产品无相关内容

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