ProposalExtract

函数功能

与ProposalConcat功能相反，从Region Proposals内将相应位置的单个元素抽取后重排，每次迭代处理16个Region Proposals，抽取16个元素后连续排列。

函数原型

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

参数说明

返回值

无

支持的型号

Atlas 训练系列产品

Atlas推理系列产品AI Core

约束说明

• 用户需保证 srcLocal中存储的 proposal 数目大于等于实际所需数目，否则会存在 tensor 越界错误。
• 用户需保证 dstLocal中存储的 element 大于等于实际所需数目，否则会存在 tensor 越界错误。
• 操作数地址偏移对齐要求请参见通用约束。

调用示例

• 接口使用样例

  // repeatTimes = 2, modeNumber = 4, 把32个Region Proposal中的score域元素抽取出来排列成32个连续元素
  ProposalExtract(dstLocal, srcLocal, 2, 4);

• 完整样例

  #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();
          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> dstLocal = outQueueDst.AllocTensor<half>();
  
          ProposalExtract(dstLocal, srcLocal, repeat, mode);
  
          outQueueDst.EnQue<half>(dstLocal);
          inQueueSrc.FreeTensor(srcLocal);
      }
      __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 = 256;
      int dstDataSize = 32;
      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):
  [  0.      0.      0.      0.     33.3     0.      0.      0.      0.
     0.      0.      0.     67.56    0.      0.      0.      0.      0.
     0.      0.     68.5     0.      0.      0.      0.      0.      0.
     0.    -11.914   0.      0.      0.      0.      0.      0.      0.
    25.19    0.      0.      0.      0.      0.      0.      0.    -72.8
     0.      0.      0.      0.      0.      0.      0.     11.79    0.
     0.      0.      0.      0.      0.      0.    -49.47    0.      0.
     0.      0.      0.      0.      0.     49.44    0.      0.      0.
     0.      0.      0.      0.     84.4     0.      0.      0.      0.
     0.      0.      0.    -14.36    0.      0.      0.      0.      0.
     0.      0.     45.97    0.      0.      0.      0.      0.      0.
     0.     52.47    0.      0.      0.      0.      0.      0.      0.
    -5.387   0.      0.      0.      0.      0.      0.      0.    -13.12
     0.      0.      0.      0.      0.      0.      0.    -88.9     0.
     0.      0.      0.      0.      0.      0.     54.      0.      0.
     0.      0.      0.      0.      0.    -51.62    0.      0.      0.
     0.      0.      0.      0.    -20.67    0.      0.      0.      0.
     0.      0.      0.     59.56    0.      0.      0.      0.      0.
     0.      0.     35.72    0.      0.      0.      0.      0.      0.
     0.     -6.12    0.      0.      0.      0.      0.      0.      0.
   -39.4     0.      0.      0.      0.      0.      0.      0.    -11.46
     0.      0.      0.      0.      0.      0.      0.     -7.066   0.
     0.      0.      0.      0.      0.      0.     30.23    0.      0.
     0.      0.      0.      0.      0.    -11.18    0.      0.      0.
     0.      0.      0.      0.    -35.84    0.      0.      0.      0.
     0.      0.      0.    -40.88    0.      0.      0.      0.      0.
     0.      0.     60.9     0.      0.      0.      0.      0.      0.
     0.    -73.3     0.      0.      0.      0.      0.      0.      0.
    38.47    0.      0.      0.   ]
  输出数据(dst_gm):
  [ 33.3    67.56   68.5   -11.914  25.19  -72.8    11.79  -49.47   49.44
    84.4   -14.36   45.97   52.47   -5.387 -13.12  -88.9    54.    -51.62
   -20.67   59.56   35.72   -6.12  -39.4   -11.46   -7.066  30.23  -11.18
   -35.84  -40.88   60.9   -73.3    38.47 ]