MrgSort4
功能说明
将已经排好序的最多4 条region proposals 队列，排列并合并成 1 条队列，结果按照 score 域由大到小排序。
函数原型
template <typename T> __aicore__ inline void MrgSort4(const LocalTensor<T>& dstLocal, const MrgSortSrcList<T>& srcLocal, const MrgSort4Info& params)
参数说明
参数名称 |
输入/输出 |
含义 |
|---|---|---|
dstLocal |
输出 |
目的操作数，存储经过排序后的 Region Proposals。 类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 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。 |
参数名称 |
输入/输出 |
含义 |
|---|---|---|
src1 |
输入 |
源操作数，第一个已经排好序的Region Proposals队列。 类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 数据类型与目的操作数保持一致。 Atlas 训练系列产品，支持的数据类型为：half Atlas推理系列产品AI Core，支持的数据类型为：half/float |
src2 |
输入 |
源操作数，第二个已经排好序的Region Proposals队列。 类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 数据类型与目的操作数保持一致。 Atlas 训练系列产品，支持的数据类型为：half Atlas推理系列产品AI Core，支持的数据类型为：half/float |
src3 |
输入 |
源操作数，第三个已经排好序的Region Proposals队列。 类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 数据类型与目的操作数保持一致。 Atlas 训练系列产品，支持的数据类型为：half Atlas推理系列产品AI Core，支持的数据类型为：half/float |
src4 |
输入 |
源操作数，第四个已经排好序的Region Proposals队列。 类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 数据类型与目的操作数保持一致。 Atlas 训练系列产品，支持的数据类型为：half Atlas推理系列产品AI Core，支持的数据类型为：half/float |
参数名称 |
输入/输出 |
含义 |
|---|---|---|
elementLengths |
输入 |
四个源Region Proposals队列的长度（Region Proposal数目），类型为长度为4的uint16_t数据类型的数组，理论上每个元素取值范围[0, 4095]，但不能超出UB的存储空间。 |
ifExhaustedSuspension |
输入 |
某条队列耗尽后，指令是否需要停止，类型为bool，默认false。 |
validBit |
输入 |
有效队列个数，取值如下：
|
repeatTimes |
输入 |
迭代次数，每一次源操作数和目的操作数跳过四个队列总长度。取值范围：repeatTimes∈[1,255]。
repeatTimes参数生效是有条件的，需要同时满足以下四个条件：
|
支持的型号
Atlas 训练系列产品
Atlas推理系列产品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. ]
