RpSort16
函数功能
根据Region Proposals中的score域对其进行排序（score大的排前面），每次排16个Region Proposals。
函数原型
1 2 |
template <typename T> __aicore__ inline void RpSort16(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal, const int32_t repeatTimes) |
参数说明
参数名称 |
输入/输出 |
含义 |
|---|---|---|
dstLocal |
输出 |
目的操作数，存储经过排序后的Region Proposals。 类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 LocalTensor的起始地址需要32字节对齐。 Atlas 训练系列产品，支持的数据类型为：half Atlas推理系列产品AI Core，支持的数据类型为：half/float |
srcLocal |
输入 |
源操作数，存储未经过排序的Region Proposals。 类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 LocalTensor的起始地址需要32字节对齐。 Atlas 训练系列产品，支持的数据类型为：half Atlas推理系列产品AI Core，支持的数据类型为：half/float |
repeatTimes |
输入 |
重复迭代次数，int32_t类型，每次排16个proposals。取值范围：repeatTimes∈[0,255]。 |
支持的型号
Atlas 训练系列产品
Atlas推理系列产品AI Core
注意事项
- 用户需保证srcLocal和dstLocal中存储的Region Proposal数目大于实际所需数据，否则会存在tensor越界错误。
- 当存在proposal[i]与proposal[j]的score值相同时，如果i>j，则proposal[j]将首先被选出来，排在前面。
- 操作数地址偏移对齐要求请参见通用约束。
调用示例
- 接口使用样例
1 2
// repeatTimes = 2, 对2个Region Proposal进行排序 AscendC::RpSort16(dstLocal, dstLocal, 2);
- 完整样例
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
#include "kernel_operator.h" class KernelVecProposal { public: __aicore__ inline KernelVecProposal() {} __aicore__ inline void Init(__gm__ uint8_t* src, __gm__ uint8_t* dstGm) { srcGlobal.SetGlobalBuffer((__gm__ half*)src); dstGlobal.SetGlobalBuffer((__gm__ half*)dstGm); pipe.InitBuffer(inQueueSrc, 1, srcDataSize * sizeof(half)); pipe.InitBuffer(outQueueDst, 1, dstDataSize * sizeof(half)); } __aicore__ inline void Process() { CopyIn(); PreProcess(); 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 PreProcess() { AscendC::LocalTensor<half> srcLocal = inQueueSrc.DeQue<half>(); AscendC::LocalTensor<half> dstLocal = outQueueDst.AllocTensor<half>(); AscendC::ProposalConcat(dstLocal, srcLocal, repeat, mode); // sort排序是基于score的，此处先创建一个有score数据的proposal，需要注意的是，非score处的数据可能是随机值 outQueueDst.EnQue<half>(dstLocal); inQueueSrc.FreeTensor(srcLocal); } __aicore__ inline void Compute() { AscendC::LocalTensor<half> dstLocal = outQueueDst.DeQue<half>(); AscendC::RpSort16(dstLocal, dstLocal, repeat); outQueueDst.EnQue<half>(dstLocal); } __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::VECOUT, 1> outQueueDst; AscendC::GlobalTensor<half> srcGlobal, dstGlobal; int srcDataSize = 32; int dstDataSize = 256; int repeat = srcDataSize / 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): [ -1.624 -42.3 -54.12 91.25 -99.4 36.72 67.44 -66.3 -52.53 3.377 -62.47 -15.85 -31.47 3.143 58.47 -83.75 21.58 63.47 7.234 35.16 -39.72 37.8 73.06 -98.7 44.1 -77.2 67.2 19.62 -87.9 -14.875 15.86 -77.75] 输出数据(dst_gm): [ 0. 0. 0. 0. 91.25 0. 0. 0. 0. 0. 0. 0. 67.44 0. 0. 0. 0. 0. 0. 0. 58.47 0. 0. 0. 0. 0. 0. 0. 36.72 0. 0. 0. 0. 0. 0. 0. 3.377 0. 0. 0. 0. 0. 0. 0. 3.143 0. 0. 0. 0. 0. 0. 0. -1.624 0. 0. 0. 0. 0. 0. 0. -15.85 0. 0. 0. 0. 0. 0. 0. -31.47 0. 0. 0. 0. 0. 0. 0. -42.3 0. 0. 0. 0. 0. 0. 0. -52.53 0. 0. 0. 0. 0. 0. 0. -54.12 0. 0. 0. 0. 0. 0. 0. -62.47 0. 0. 0. 0. 0. 0. 0. -66.3 0. 0. 0. 0. 0. 0. 0. -83.75 0. 0. 0. 0. 0. 0. 0. -99.4 0. 0. 0. 0. 0. 0. 0. 73.06 0. 0. 0. 0. 0. 0. 0. 67.2 0. 0. 0. 0. 0. 0. 0. 63.47 0. 0. 0. 0. 0. 0. 0. 44.1 0. 0. 0. 0. 0. 0. 0. 37.8 0. 0. 0. 0. 0. 0. 0. 35.16 0. 0. 0. 0. 0. 0. 0. 21.58 0. 0. 0. 0. 0. 0. 0. 19.62 0. 0. 0. 0. 0. 0. 0. 15.86 0. 0. 0. 0. 0. 0. 0. 7.234 0. 0. 0. 0. 0. 0. 0. -14.875 0. 0. 0. 0. 0. 0. 0. -39.72 0. 0. 0. 0. 0. 0. 0. -77.2 0. 0. 0. 0. 0. 0. 0. -77.75 0. 0. 0. 0. 0. 0. 0. -87.9 0. 0. 0. 0. 0. 0. 0. -98.7 0. 0. 0. ]
父主题： 矢量计算
