MrgSort4

功能说明

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

函数原型

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。 Atlas 训练系列产品，支持的数据类型为：half Atlas推理系列产品（Ascend 310P处理器）AI Core，支持的数据类型为：half/float
srcLocal	输入	源操作数，4个Region Proposals队列，并且每个Region Proposal队列都已经排好序，类型为MrgSortSrcList结构体，具体请参考表2。 Atlas 训练系列产品，支持的数据类型为：half Atlas推理系列产品（Ascend 310P处理器）AI Core，支持的数据类型为：half/float
params	输入	排序所需参数，类型为MrgSort4Info结构体，具体请参考表3。

表2 MrgSortSrcList参数说明
参数名称	输入/输出	含义
src1	输入	源操作数，第一个已经排好序的Region Proposals队列。类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。数据类型与目的操作数保持一致。 Atlas 训练系列产品，支持的数据类型为：half Atlas推理系列产品（Ascend 310P处理器）AI Core，支持的数据类型为：half/float
src2	输入	源操作数，第二个已经排好序的Region Proposals队列。类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。数据类型与目的操作数保持一致。 Atlas 训练系列产品，支持的数据类型为：half Atlas推理系列产品（Ascend 310P处理器）AI Core，支持的数据类型为：half/float
src3	输入	源操作数，第三个已经排好序的Region Proposals队列。类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。数据类型与目的操作数保持一致。 Atlas 训练系列产品，支持的数据类型为：half Atlas推理系列产品（Ascend 310P处理器）AI Core，支持的数据类型为：half/float
src4	输入	源操作数，第四个已经排好序的Region Proposals队列。类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。数据类型与目的操作数保持一致。 Atlas 训练系列产品，支持的数据类型为：half Atlas推理系列产品（Ascend 310P处理器）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 valid_bit=15

支持的型号

Atlas 训练系列产品

Atlas推理系列产品（Ascend 310P处理器）AI Core

约束说明

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

不支持源操作数与目的操作数之间存在地址重叠。

调用示例

接口使用样例

// 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);
MrgSort4(dstLocal, srcList, srcInfo);

完整样例

#include "kernel_operator.h"

namespace AscendC {
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()
    {
        LocalTensor<half> srcLocal = inQueueSrc.AllocTensor<half>();
        DataCopy(srcLocal, srcGlobal, srcDataSize);
        inQueueSrc.EnQue(srcLocal);
    }
    __aicore__ inline void Compute()
    {
        LocalTensor<half> srcLocal = inQueueSrc.DeQue<half>();
        LocalTensor<half> vconcatWorkLocal = workQueue.AllocTensor<half>();
        LocalTensor<half> dstLocal = outQueueDst.AllocTensor<half>();

        // 先构造4个region proposal然后进行合并排序
        Concat(vconcatWorkLocal[0], srcLocal[0], repeat, mode);
        RpSort16(vconcatWorkLocal[0], vconcatWorkLocal[0], repeat);

        Concat(vconcatWorkLocal[workDataSize], srcLocal[singleDataSize], repeat, mode);
        RpSort16(vconcatWorkLocal[workDataSize], vconcatWorkLocal[workDataSize], repeat);
 
        Concat(vconcatWorkLocal[workDataSize * 2], srcLocal[singleDataSize * 2], repeat, mode);
        RpSort16(vconcatWorkLocal[workDataSize * 2], vconcatWorkLocal[workDataSize * 2], repeat);

        Concat(vconcatWorkLocal[workDataSize * 3], srcLocal[singleDataSize * 3], repeat, mode);
        RpSort16(vconcatWorkLocal[workDataSize * 3], vconcatWorkLocal[workDataSize * 3], repeat);

        struct MrgSortSrcList<half> srcList(vconcatWorkLocal[0], vconcatWorkLocal[workDataSize],
            vconcatWorkLocal[workDataSize * 2], vconcatWorkLocal[workDataSize * 3]);
        uint16_t elementLengths[4] = {singleDataSize, singleDataSize, singleDataSize, singleDataSize};
        struct MrgSort4Info srcInfo(elementLengths, false, 15, 1);
        MrgSort4(dstLocal, srcList, srcInfo);

        outQueueDst.EnQue<half>(dstLocal);
        inQueueSrc.FreeTensor(srcLocal);
        workQueue.FreeTensor(vconcatWorkLocal);
    }
    __aicore__ inline void CopyOut()
    {
        LocalTensor<half> dstLocal = outQueueDst.DeQue<half>();
        DataCopy(dstGlobal, dstLocal, dstDataSize);
        outQueueDst.FreeTensor(dstLocal);
    }

private:
    TPipe pipe;
    TQue<QuePosition::VECIN, 1> inQueueSrc;
    TQue<QuePosition::VECIN, 1> workQueue;
    TQue<QuePosition::VECOUT, 1> outQueueDst;
    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;
};
} // namespace AscendC

extern "C" __global__ __aicore__ void vec_proposal_kernel(__gm__ uint8_t* src, __gm__ uint8_t* dstGm)
{
    AscendC::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.   ]

父主题： 矢量计算