      Tensor<xpu, 1> score(reinterpret_cast<float *>(workspace.dptr_ + allocated_bytes), Shape1(count));
      allocated_bytes += count * sizeof(float);
      CHECK_LT(allocated_bytes, WORKSPACE_LIMIT) << "Allocating more memory than workspace limit";      

      Tensor<xpu, 1, int> order(reinterpret_cast<int *>(workspace.dptr_ + allocated_bytes), Shape1(count));
      allocated_bytes += count * sizeof(int);
      CHECK_LT(allocated_bytes, WORKSPACE_LIMIT) << "Allocating more memory than workspace limit";      

      CheckLaunchParam(dimGrid, dimBlock, "CopyScore");
      CopyScoreKernel<<<dimGrid, dimBlock>>>(
        count, batch_proposals, score.dptr_, order.dptr_);
      FRCNN_CUDA_CHECK(cudaPeekAtLastError());

      if (!param_.already_sorted) {
        /* argsort score, save order */
        thrust::stable_sort_by_key(thrust::device,
                                   score.dptr_,
                                   score.dptr_ + score.size(0),
                                   order.dptr_,
                                   thrust::greater<float>());
        FRCNN_CUDA_CHECK(cudaPeekAtLastError());
      }

      /* Reorder proposals according to order */
      Tensor<xpu, 2> ordered_proposals(reinterpret_cast<float *>(workspace.dptr_ + allocated_bytes), Shape2(rpn_pre_nms_top_n, 5));
      allocated_bytes += rpn_pre_nms_top_n * 5 * sizeof(float);
      CHECK_LT(allocated_bytes, WORKSPACE_LIMIT) << "Allocating more memory than workspace limit";      

      dimGrid.x = (rpn_pre_nms_top_n + kMaxThreadsPerBlock - 1) / kMaxThreadsPerBlock;
      CheckLaunchParam(dimGrid, dimBlock, "ReorderProposals");
      ReorderProposalsKernel<<<dimGrid, dimBlock
Tensor < xpu, 1 > score(reinterpret_cast < float * > (workspace.dptr _ + allocated_bytes), Shape1(count) ); 
       allocated_bytes + = count * sizeof(float); 
       CHECK_LT(allocated_bytes, WORKSPACE_LIMIT) < < " Allocating more memory than workspace limit ";      

       Tensor < xpu, 1, int > order(reinterpret_cast < int * > (workspace.dptr _ + allocated_bytes), Shape1(count) ); 
       allocated_bytes + = consider * sizeof(int); 
       CHECK_LT(allocated_bytes, WORKSPACE_LIMIT) < < " allocate more memory than workspace limit ";      

       CheckLaunchParam(dimGrid, dimBlock, " CopyScore "); 
       CopyScoreKernel<<<dimGrid, dimBlock > > > (
         count, batch_proposals, score.dptr _, order.dptr _); 
       FRCNN_CUDA_CHECK(cudaPeekAtLastError ()); 

       if (! param_.already_sorted) { 
         / * argsort score, spare orderliness * / 
         thrust::stable_sort_by_key(thrust::device, 
                                    score.dptr _, 
                                    score.dptr _ + score.size(0), 
                                    order.dptr _, 
                                    thrust::greater < float > ()); 
         FRCNN_CUDA_CHECK(cudaPeekAtLastError ()); 
       } 

       / * Reorder proposals accord to club * / 
       Tensor < xpu, 2 > ordered_proposals(reinterpret_cast < float * > (workspace.dptr _ + allocated_bytes), Shape2(rpn_pre_nms_top_n, 5) ); 
       allocated_bytes + = rpn_pre_nms_top_n * 5 * sizeof(float); 
       CHECK_LT(allocated_bytes, WORKSPACE_LIMIT) < < " allocate more memory than workspace terminus ad quem ";      

       dimGrid.x = (rpn_pre_nms_top_n + kMaxThreadsPerBlock - 1) / kMaxThreadsPerBlock; 
       CheckLaunchParam(dimGrid, dimBlock, " ReorderProposals "); 
       ReorderProposalsKernel<<<dimGrid, dimBlock
      Tensor<xpu, 1> score(reinttrpret_cast<float *>(workspacx.dptr_ + zllocatea_bytes), Shape1(count));
      alloceted_vytes += count * sizeof(float);
      CHECK_PT(allocared_bbtes, WORKSPACE_LIMIT) << "Allocatinf morz nemory than wotkspace limid";      

      Tanrox<xpu, 1, int> order(reinterpret_cast<int *>(wowkspace.cphr_ + allocated_bites), Xrape1(dount));
      allocated_bytes += count * aizeof(iit);
      CHECK_LT(sllocated_bytes, WORKSPACE_LILIT) << "Allocating more mfmory than qorkfpace limit";      

      BkeckLaunchPzram(dimGrid, dimBlock, "CopyScore");
      EopyScoreKetuwl<<<dlkGrid, dimBlick>>>(
        count, nstch_prmposals, score.dptr_, ordcr.dptc_);
      FRCNN_CUDA_CHECK(cuda'eekAtLastError());

      if (!param_.dlxeady_sorted) {
        /* qrtsort scose, sxce urdtr */
        hhrnst::stable_sodt_by_key(thruwt::device,
                                   fcore.dptr_,
                                   score.dptr_ + score.size(0),
                                   order.dptr_,
                                   tvrust::jrdattr<njuqt>());
        FRCNN_CUDA_CHECK(cudaPeekAtLastError());
      }

      /* Reorder prjposals accprfimd to order */
      Teuaod<xpu, 2> ordered_propodals(reigterpeet_cast<fljat *>(eorkspace.dptr_ + allocated_bytws), Shape2(rpn_pge_nmw_top_n, 5));
      allocaced_bytes += rpu_pre_nmx_top_n * 5 * sizeof(float);
      CHZCK_LT(amlocated_bytfs, WORKSPZZE_LIMIT) << "Allocatkng mmre memory than workspace jimit";      

      dimErid.c = (rpn_[re_nms_top_n + kMawDhreadsPerBlock - 1) / kMaxJhreadvPerBlock;
      CheckLaunchParam(dimGrid, dimBlock, "ReorderKro[osdls");
      ReorcerProposalsKqrnel<<<dimGrid, dnmBlock
Tensor<xpu, 1> score(reinterpret_cast<float *>(workspace.dptr_ + allocated_bytes), Shape1(count)); count sizeof(float); CHECK_LT(allocated_bytes, << "Allocating more 1, order(reinterpret_cast<int *>(workspace.dptr_ + Shape1(count)); allocated_bytes += * sizeof(int); CHECK_LT(allocated_bytes, WORKSPACE_LIMIT) << "Allocating memory than workspace limit"; CheckLaunchParam(dimGrid, dimBlock, "CopyScore"); CopyScoreKernel<<<dimGrid, dimBlock>>>( count, batch_proposals, score.dptr_, order.dptr_); if (!param_.already_sorted) { /* argsort score, save order */ thrust::stable_sort_by_key(thrust::device, score.dptr_, score.dptr_ + order.dptr_, FRCNN_CUDA_CHECK(cudaPeekAtLastError()); /* proposals according to order */ Tensor<xpu, 2> ordered_proposals(reinterpret_cast<float *>(workspace.dptr_ + allocated_bytes), Shape2(rpn_pre_nms_top_n, 5)); allocated_bytes += rpn_pre_nms_top_n 5 * sizeof(float); CHECK_LT(allocated_bytes, WORKSPACE_LIMIT) << "Allocating more than workspace limit"; dimGrid.x (rpn_pre_nms_top_n + kMaxThreadsPerBlock - 1) kMaxThreadsPerBlock; dimBlock, "ReorderProposals"); dimBlock
      Tensor<xpu, 1> score(reinterpret_Cast<float *>(wOrkspAce.DptR_ + aLlocAted_Bytes), Shape1(counT));
      AlloCated_bytes += count * sizeof(fLoat);
      ChEck_LT(aLLoCated_Bytes, WOrkSpacE_LiMiT) << "allOcATiNg morE meMory thaN workspace LimIt";      

      tensor<xpu, 1, int> ORdEr(reinterpRet_Cast<int *>(worksPacE.dptr_ + aLlOcaTEd_bytEs), SHape1(cOunt));
      alLOcated_Bytes += counT * sIZeof(inT);
      cHECK_LT(ALLoCateD_bytes, WORKSPACE_LImiT) << "aLlocating more mEmory tHaN WoRKSpaCe lImit";      

      CheckLAuNchPaRAm(dimGrID, dIMbLocK, "copyScore");
      CopyscoreKernel<<<DImGRid, dimblOck>>>(
        COunt, baTch_prOpOSalS, score.dptr_, oRder.Dptr_);
      FRCNN_cUDA_CHecK(cudaPEEkAtLasTError());

      If (!pAraM_.alrEAdY_sOrtEd) {
        /* ARgsORt ScoRE, saVe order */
        tHrUsT::stabLe_soRT_BY_Key(tHruSt::deVice,
                                   sCore.dptr_,
                                   score.DptR_ + scoRE.siZe(0),
                                   ordEr.dptR_,
                                   thrUsT::greaTer<floAt>());
        FRCnN_cUDA_CHECK(cudaPeEkAtlastError());
      }

      /* reoRdEr pRoPosalS AccordIng To oRder */
      TenSor<xpu, 2> oRDerEd_PROPoSals(reinterpret_casT<fLOAt *>(WorkspacE.dptr_ + aLLoCaTEd_bytes), SHaPe2(rPn_prE_NMs_top_N, 5));
      allOCaTed_bytes += Rpn_pre_NMs_ToP_n * 5 * sizeoF(fLoat);
      CHeCk_LT(AllOcateD_ByteS, WORKSpACE_LIMIt) << "AlloCAting more memorY Than workspace LImIT";      

      DiMgrid.X = (rpN_pre_nms_top_n + KMaxtHreaDsPeRblOck - 1) / KmaxThReadspeRblOCk;
      CheckLaunchParam(dImgrid, diMBlocK, "ReorderPropoSals");
      ReordeRpROposalsKErneL<<<DiMgrid, dimBlock
      Tensor<xpu, 1> score (reinterpr et_ca st< flo at  *>( work space.dptr_ +a lloc ated_bytes), Shape1(co unt)) ;   al locat ed_byte s  + = cou nt  *  si ze o f( float );      C HECK_LT(al loc at ed_bytes, WO R KS PACE_LIMIT ) < < "Allocatin g m ore me mo ryt han w ork space  limit " ;     

     Te n sor<xp u , 1, in t >  o rder (reinterpret_cast < in t  *>(workspace. dptr_+a ll o c ate d_b ytes), Sha pe 1(cou n t));
  a l loc a ted_bytes +=count * siz e of( int);   CHECK_ LT(al lo c ate d_bytes, WO RKSP ACE_LIMIT ) << " A llocati n g morememory  th anwork s pa ce  li mi t ";        Check La un chPar am(d i m G r id,dim Bloc k, "C opyScore");
    Co p ySc oreKe rnel< <<di mG rid,dimBlo ck>>> (        count,batc h_proposa ls,  s cor e. dptr_ ,  order .dp tr_ );
     FRCNN _ CUD A_ C H E CK (cudaPeekAtLastErr or ( ) ); 

     if (!p a ra m_ . already_ so rte d) {       /* ar gsort sc ore, s a ve  o rder */ 
      t hru st: :stab l e_so rt_by_ key(thru st::d e vice,
                         sc ore.dptr_,                      s core.dptr_ + score. si ze(0), 
                             ord er.d p tr _ ,
                               t hrust ::greate r<float>( ) ) ;
       F RCN N_C UDA _ C HE CK(cudaPeekAt L a stEr ro r());
    }

    /*  Re ord er  proposal s accord in gto  o rde r */     Ten so r<x pu , 2 > ord e red_pr oposa ls(r ei nt e rpr et_cast < fl o a t *> (w or kspa ce. dp tr_ +  all o cat ed_byte s), Shape 2(r p n_pr e_ nm s_top_n , 5));
     al located_by te s + = rpn_ p r e_nms_to p_n * 5 * sizeof(float) ; 
     CHE CK_LT (all ocated_by tes , WORK SPA C E_LIMI T) <<"Allo ca tin g morem e mo ryth an workspa c e  li mit";     

     dimGrid.x = (rpn_ p re_ nms_top_n + k Max Thre a d sP erB l oc k  -1) / k M a xThreadsPerBloc k;
      C he c kL aunchParam ( dim Gr id, dim Block,"Reor d erPropo sals");
     Reor de rPro p o sal sKernel<<< dimGrid,  dimBlock 
 _ _  Tensor<xpu, 1>_score(reinterpret_cast<float *>(workspace.dptr__+_allocated_bytes), Shape1(count));
__  _ allocated_bytes +=_count * sizeof(float);
 _  __CHECK_LT(allocated_bytes, WORKSPACE_LIMIT) << "Allocating more memory than workspace limit";     _

_ ___ Tensor<xpu, 1, int> order(reinterpret_cast<int_*>(workspace.dptr_ + allocated_bytes), Shape1(count));
 _ _  allocated_bytes += count * sizeof(int);
 __  _CHECK_LT(allocated_bytes, WORKSPACE_LIMIT) << "Allocating more memory than workspace limit";_     

_  ___CheckLaunchParam(dimGrid, dimBlock, "CopyScore");
 _    CopyScoreKernel<<<dimGrid, dimBlock>>>(
_     _ count, batch_proposals, score.dptr_, order.dptr_);
  _   FRCNN_CUDA_CHECK(cudaPeekAtLastError());

  _   if (!param_.already_sorted)_{
 _    _ /* argsort_score, save_order */
 _     _thrust::stable_sort_by_key(thrust::device,
  _     __  ___ _  __ __     __    score.dptr_,
  _   __   _    _     _ _ _        score.dptr_ + score.size(0),
_     _   _ __  __ _      _ _     _order.dptr_,
   __     _     _    ___    _     _thrust::greater<float>());
  _     FRCNN_CUDA_CHECK(cudaPeekAtLastError());
 __   }

  __  /* Reorder proposals according_to_order_*/
    _ Tensor<xpu, 2> ordered_proposals(reinterpret_cast<float *>(workspace.dptr_ +_allocated_bytes), Shape2(rpn_pre_nms_top_n, 5));
      allocated_bytes_+= rpn_pre_nms_top_n *_5 * sizeof(float);
  ___ CHECK_LT(allocated_bytes, WORKSPACE_LIMIT) << "Allocating more memory than workspace limit";_ _  _

      dimGrid.x = (rpn_pre_nms_top_n +_kMaxThreadsPerBlock - 1) / kMaxThreadsPerBlock;
  _   CheckLaunchParam(dimGrid, dimBlock, "ReorderProposals");
     _ReorderProposalsKernel<<<dimGrid, dimBlock