RpSort16

函数功能

根据 Region Proposals 中的 score 域对其进行排序（score大的排前面），每次排 16 个Region Proposals。

函数原型

template <typename T> __aicore__ inline void RpSort16(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal, const int32_t repeatTimes)

参数说明

支持的型号

Atlas 训练系列产品

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

注意事项

• 用户需保证srcLocal和dstLocal中存储的Region Proposal数目大于实际所需数据，否则会存在tensor越界错误。
• 当存在proposal[i]与proposal[j]的score值相同时，如果i>j，则proposal[j]将首先被选出来，排在前面。
• 操作数地址偏移对齐要求请参见通用约束。

调用示例

• 接口使用样例

  // repeatTimes = 2, 对2个Region Proposal进行排序
  RpSort16(dstLocal, dstLocal, 2);

• 完整样例

  #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((__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()
      {
          LocalTensor<half> srcLocal = inQueueSrc.AllocTensor<half>();
          DataCopy(srcLocal, srcGlobal, srcDataSize);
          inQueueSrc.EnQue(srcLocal);
      }
      __aicore__ inline void PreProcess()
      {
          LocalTensor<half> srcLocal = inQueueSrc.DeQue<half>();
          LocalTensor<half> dstLocal = outQueueDst.AllocTensor<half>();
          Concat(dstLocal, srcLocal, repeat, mode); // sort排序是基于score的，此处先创建一个有score数据的proposal，需要注意的是，非score处的数据可能是随机值
          outQueueDst.EnQue<half>(dstLocal);
          inQueueSrc.FreeTensor(srcLocal);
      }
      __aicore__ inline void Compute()
      {
          LocalTensor<half> dstLocal = outQueueDst.DeQue<half>();
          RpSort16(dstLocal, dstLocal, repeat);
          outQueueDst.EnQue<half>(dstLocal);
      }
      __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::VECOUT, 1> outQueueDst;
      GlobalTensor<half> srcGlobal, dstGlobal;
      int srcDataSize = 32;
      int dstDataSize = 256;
      int repeat = srcDataSize / 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):
  [ -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.   ]