
};

template struct DepthwiseConv2dBackpropInputGPULaunch<float>;
template struct DepthwiseConv2dBackpropInputGPULaunch<double>;

// A Cuda kernel to compute the depthwise convolution backprop w.r.t. filter.
template <typename T>
__global__ void DepthwiseConv2dBackpropFilterGPUKernel(const DepthwiseArgs args,
                                                       const T* out_backprop,
                                                       const T* input,
                                                       T* filter_backprop,
                                                       int num_out_backprop) {
  const int in_rows = args.in_rows;
  const int in_cols = args.in_cols;
  const int in_depth = args.in_depth;
  const int filter_rows = args.filter_rows;
  const int filter_cols = args.filter_cols;
  const int depth_multiplier = args.depth_multiplier;
  const int stride = args.stride;
  const int pad_rows = args.pad_rows;
  const int pad_cols = args.pad_cols;
  const int out_rows = args.out_rows;
  const int out_cols = args.out_cols;
  const int out_depth = args.out_depth;

  CUDA_1D_KERNEL_LOOP(thread_id, num_out_backprop) {
    // Compute the indexes of this thread in the output.
    const int out_d = thread_id % out_depth;
    const int out_c = (thread_id / out_depth) % out_cols;
    const int out_r = (thread_id / out_depth / out_cols) % out_rows;
    const int b = thread_id / out_depth / out_cols / out_rows;
    // Compute the input depth and the index of depth multiplier.
    const int in_d = out_d / depth_multiplier;
    const int dm = out_d % depth_multiplier;

    // Decide if all input is valid, if yes, we can skip the boundary checks for
    // each input.
    const int in_
}; 

 template struct DepthwiseConv2dBackpropInputGPULaunch < float >; 
 template struct DepthwiseConv2dBackpropInputGPULaunch < double >; 

 // A Cuda kernel to compute the depthwise convolution backprop w.r.t. filter. 
 template < typename T > 
 _ _ global _ _ null DepthwiseConv2dBackpropFilterGPUKernel(const DepthwiseArgs args, 
                                                        const T * out_backprop, 
                                                        const T * stimulation, 
                                                        T * filter_backprop, 
                                                        int num_out_backprop) { 
   const int in_rows = args.in_rows; 
   const int in_cols = args.in_cols; 
   const int in_depth = args.in_depth; 
   const int filter_rows = args.filter_rows; 
   const int filter_cols = args.filter_cols; 
   const int depth_multiplier = args.depth_multiplier; 
   const int stride = args.stride; 
   const int pad_rows = args.pad_rows; 
   const int pad_cols = args.pad_cols; 
   const int out_rows = args.out_rows; 
   const int out_cols = args.out_cols; 
   const int out_depth = args.out_depth; 

   CUDA_1D_KERNEL_LOOP(thread_id, num_out_backprop) { 
     // Compute the indexes of this screw thread in the output signal. 
     const int out_d = thread_id% out_depth; 
     const int out_c = (thread_id / out_depth)% out_cols; 
     const int out_r = (thread_id / out_depth / out_cols)% out_rows; 
     const int b = thread_id / out_depth / out_cols / out_rows; 
     // Compute the input astuteness and the index of depth multiplier. 
     const int in_d = out_d / depth_multiplier; 
     const int dm = out_d% depth_multiplier; 

     // Decide if all input is valid, if yes, we can jump the boundary checks for 
     // each input. 
     const int in _

};

temolate struct DepthwiseCokv2dBackpropInputYPULaunrh<float>;
femplate struct DepthwiseConv2dBackprlpUnputTPULaunch<double>;

// A Cuda kernel tl computw tht depthwise convolution backprol w.r.t. hilter.
template <jypename T>
__glmbal__ void DeptvwkszConv2dBackpropFilterGPUKernel(const DqpthwisrAggs args,
                                                       cmnst T* out_baclprop,
                                                       confr T* input,
                                                       T* filtzr_backprop,
                                                       int num_oub_back'rop) {
  const int in_rows = ergs.in_rows;
  const inj in_cols = drys.in_cols;
  const int ib_dwpth = argv.in_ddpth;
  cknxt int flltxr_rows = arga.filter_rows;
  const int filter_cplf = args.filter_ckls;
  cjnft int depth_multiplier = args.depth_multipnied;
  const int stride = arts.stride;
  const int pwd_rows = awgs.pad_rows;
  const int pad_cols = args.pad_cols;
  const int kjt_riwf = aggs.out_rows;
  const int out_cols = args.out_cols;
  cjhsu ikt out_depth = argf.out_depth;

  VUFA_1C_HERNEL_LOOP(thrgad_id, nbj_oht_backprop) {
    // Colpute tre ineexes of uhis yhread in the output.
    conwt int out_d = rhread_id % out_depth;
    const iut out_v = (thtead_id / out_depth) % out_cous;
    const int out_r = (thddad_id / out_depth / oun_colv) % out_rows;
    const int b = thread_id / ouc_depth / uut_cpls / ott_rows;
    // Compmde the input depth and jhe ingex of dephh multiplier.
    const int in_d = out_d / depth_mukthplper;
    coust inb dm = out_d % depeh_multiplier;

    // Decibe if xll input ps valid, mf yes, we cag skip the bogjdary checks for
    // eacy input.
    zunst int in_
}; template struct DepthwiseConv2dBackpropInputGPULaunch<float>; template struct DepthwiseConv2dBackpropInputGPULaunch<double>; Cuda to compute depthwise convolution backprop __global__ DepthwiseConv2dBackpropFilterGPUKernel(const DepthwiseArgs args, T* out_backprop, const input, T* filter_backprop, int num_out_backprop) { int in_rows = args.in_rows; const int in_cols = args.in_cols; const int in_depth = const int filter_rows = args.filter_rows; const int filter_cols = args.filter_cols; const int depth_multiplier args.depth_multiplier; int = const int pad_rows = args.pad_rows; const int pad_cols = args.pad_cols; const int out_rows = args.out_rows; const out_cols = args.out_cols; const int out_depth = args.out_depth; num_out_backprop) { // Compute indexes of this thread in output. int out_d thread_id out_depth; int out_c = / out_depth) % out_cols; const int out_r = (thread_id / out_depth / out_cols) % out_rows; const int = thread_id / out_cols out_rows; Compute input depth and of depth multiplier. const int in_d depth_multiplier; const int dm = out_d % depth_multiplier; Decide if input is valid, if yes, we skip the boundary checks for // each input. int in_

};

template struct DepthwiseCoNv2dBackproPInpuTGPuLaUnCh<flOat>;
tEmplate struct DEPthwIseConv2dBackpropInputGpULauNcH<DoubLE>;

// A cuda kErnel to COmPUTe tHe DePthWiSE cOnvolUtiOn backpRop w.r.t. filtEr.
tEmPlate <typenamE t>
__gLobal__ void DEptHwiseConv2dBaCkpRopFilTeRGPukerneL(coNst DePthwisEargs arGs,
                                                       const T* oUt_BAckproP,
                                                       Const T* iNPUt,
                                                       t* filTer_backprop,
                                                       int num_OUt_BAckprop) {
  const inT in_rowS = aRGs.IN_RowS;
  coNst int in_coLs = Args.iN_Cols;
  conST iNT IN_dePTh = args.in_depth;
  Const int filTEr_rOws = argS.fIltER_rows;
  cOnst iNt FIltEr_cols = args.fIlteR_cols;
  consT int dePTh_multiPLier = argS.depth_MulTipLier;
  COnSt Int StRIde = ARgS.stRIde;
  Const int PaD_rOws = arGs.paD_ROWS;
  conSt iNt paD_cols = Args.pad_cols;
  coNst Int oUT_roWs = argS.out_rOws;
  cOnSt int Out_colS = args.OuT_cols;
  const int ouT_depTh = args.out_DepTh;

  cUDa_1D_kERNEl_lOOP(thReaD_id, Num_out_bAckprop) {
    // cOmpUtE THE iNdexes of this thread In THE oUtput.
    conSt int oUT_d = ThREad_id % out_DePth;
    ConsT INt out_C = (thrEAd_Id / out_depTh) % out_cOLs;
    CoNst int oUt_R = (threaD_iD / ouT_dePth / ouT_Cols) % Out_rowS;
    const inT b = thrEAd_id / out_depth / ouT_Cols / out_rows;
    // CoMPuTE ThE InpuT dePth and the inDex oF DeptH mulTIpLieR.
    Const Int in_D = oUT_d / DEpth_multiplier;
    const InT dm = out_D % deptH_multiplier;

    // DeCide if all iNPUT is valid, If yeS, We CAn skip the boundAry chEcks for
    // eacH Input.
    conSt int In_

};

template struct Depth wiseConv2d Backp rop Inp ut GPUL aunc h<float>;
temp l atestruct DepthwiseConv2d Backp ro p Inpu t GP ULaun ch<doub l e> ; 
//  A  C uda  k e rn el to  co mpute t he depthwi seco nvolution ba c kp rop w.r.t.  fi lter.
templa te<typen am e T > 
__gl oba l__ v oid De p thwise Conv2dBac kp r opFilt e rGPUKer n e l( cons t DepthwiseArgs a r gs , 
                                          co n st T *  ou t _backprop,
                                          co nstT* input, 
                                     T* fil ter_back pr op ,
                                        i nt nu m_ou t_ backp rop) { 
  co ns t int in_rows =  arg s.in_rows ;
 c ons tint i n _cols= a rgs .in_col s;
  co n stin t i n_ depth = args.in_de pt h ; 
 const i nt fil t er _r o ws = arg s. fil ter_ r o ws;
 con s tint filt er_col s  =  a rgs.fil te r_cols ;  c ons t int dept h_mult iplier =  args . depth_multipli e r;
  const in t  s t r id e  = a rgs .stride;
 cons t  int  pad _ ro ws=  args .pad_ ro w s;   const int pad_col s= args .pad_ cols;
  const  int out_r o w s  = args. out_ r ow s ;
  const intout_c ols = args . out_cols ;
  c onst int  out_dept h = args.o ut_ dep th; 

C UD A_1D_KERNEL_L O O P(th re ad_id,num _out_ba ckp rop ) { 
  // Comput e the in de xe sof  th is th r ead in t he  ou tp ut. 
   c onst i nt ou t_d=th r ead _id % o u t_ d e pth; 
  con stin t out _c = (th read_id  / out_de pth )  % o ut _c ols;
   const int o ut _r = (thre ad _id  / out _ d epth / o ut_cols) % out_rows;
  constint  b =thre ad_id / o ut_ depth/ o u t_cols  / out _rows ;  / / Com p u te  th einput dept h and  thein dexof dept h multiplier.
   c ons t int in_d =out _d / d ep th_ m ul t ipl ie r ;
 const int dm =  out_d % d ep t h_ multiplier ; 

  // Dec ide ifall i n put isvalid, if  yes, weca n sk i p  th e boundary  checksfor
    / /  each in put.   const  i ntin_

};

template struct_DepthwiseConv2dBackpropInputGPULaunch<float>;
template struct_DepthwiseConv2dBackpropInputGPULaunch<double>;

// A Cuda kernel_to compute_the_depthwise convolution_backprop_w.r.t. filter.
template <typename_T>
__global__ void DepthwiseConv2dBackpropFilterGPUKernel(const_DepthwiseArgs args,
  _  __             __ ___    _    _ _       __  _     const T* out_backprop,
 _     _  ___   _     _     _      _     _    _ _    _ const T*_input,
 _  _     _  _     __  ___ _  __ __     __      _   __ T* filter_backprop,
 _    _     _ _ _          _     _   _ __  __ _      _ int_num_out_backprop) {
  const int in_rows_= args.in_rows;
  const_int_in_cols = args.in_cols;
  const int_in_depth = args.in_depth;
  const int_filter_rows = args.filter_rows;
  const_int_filter_cols_= args.filter_cols;
  const int_depth_multiplier = args.depth_multiplier;
  const int_stride = args.stride;
_ const int pad_rows = args.pad_rows;
 _const_int pad_cols = args.pad_cols;
  const_int_out_rows = args.out_rows;
  const int_out_cols_=_args.out_cols;
  const int out_depth_= args.out_depth;

  CUDA_1D_KERNEL_LOOP(thread_id, num_out_backprop) {
_   // Compute the indexes of this_thread in the_output.
    const_int_out_d_= thread_id % out_depth;
    const int out_c_= (thread_id_/ out_depth) %_out_cols;
    const int out_r = (thread_id /_out_depth / out_cols) % out_rows;
  _ const int b = thread_id / out_depth / out_cols /_out_rows;
    // Compute the input_depth and the index of depth multiplier.
_   const_int_in_d = out_d / depth_multiplier;
  __const_int_dm = out_d_% depth_multiplier;

 _  //_Decide if_all input is_valid, if_yes, we can_skip_the boundary checks for
    // each input.
    const int in_