const int filter_offset =
                multiplier +
                depth_multiplier *
                    (in_d + in_depth * (f_c + filter_offset_temp));
            sum += ldg(input + input_offset) * ldg(filter + filter_offset);
          }
        }
      }
    }
    output[thread_id] = sum;
  }
}
}  // namespace

// A simple launch pad to launch the Cuda kernel for depthwise convolution.
template <typename T>
struct DepthwiseConv2dGPULaunch {
  static void Run(const GPUDevice& d, const DepthwiseArgs args, const T* input,
                  const T* filter, T* output) {
    // In this kernel, each thread is computing the gradients from one element
    // in the out_backprop. Note that one element in the out_backprop can map
    // to multiple filter elements.
    const int num_outputs =
        args.batch * args.out_rows * args.out_cols * args.out_depth;
    CudaLaunchConfig config = GetCudaLaunchConfig(num_outputs, d);

    DepthwiseConv2dGPUKernel<
        T><<<config.block_count, config.thread_per_block, 0, d.stream()>>>(
        args, input, filter, output, num_outputs);
  }
};

template struct DepthwiseConv2dGPULaunch<float>;
template struct DepthwiseConv2dGPULaunch<double>;

// A Cuda kernel to compute the depthwise convolution backprop w.r.t. input.
template <typename T>
__global__ void DepthwiseConv2dBackpropInputGPUKernel(const DepthwiseArgs args,
                                                      const T* out_backprop,
                                                      const T* filter,
                                                      T* in_backprop,
                                                      int num_in_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_
const int filter_offset = 
                 multiplier + 
                 depth_multiplier * 
                     (in_d + in_depth * (f_c + filter_offset_temp) ); 
             sum + = ldg(input + input_offset) * ldg(filter + filter_offset); 
           } 
         } 
       } 
     } 
     output[thread_id ] = sum; 
   } 
 } 
 }   // namespace 

 // A simple launch launching pad to establish the Cuda kernel for depthwise convolution. 
 template < typename T > 
 struct DepthwiseConv2dGPULaunch { 
   static void Run(const GPUDevice & d, const DepthwiseArgs args, const T * remark, 
                   const T * filter, T * output) { 
     // In this kernel, each thread is calculate the gradients from one element 
     // in the out_backprop. Note that one component in the out_backprop can map 
     // to multiple filter elements. 
     const int num_outputs = 
         args.batch * args.out_rows * args.out_cols * args.out_depth; 
     CudaLaunchConfig config = GetCudaLaunchConfig(num_outputs, d); 

     DepthwiseConv2dGPUKernel < 
         T><<<config.block_count, config.thread_per_block, 0, d.stream () > > > (
         args, input, filter, output signal, num_outputs); 
   } 
 }; 

 template struct DepthwiseConv2dGPULaunch < float >; 
 template struct DepthwiseConv2dGPULaunch < double >; 

 // A Cuda kernel to calculate the depthwise convolution backprop w.r.t. input. 
 template < typename T > 
 _ _ global _ _ void DepthwiseConv2dBackpropInputGPUKernel(const DepthwiseArgs args, 
                                                       const T * out_backprop, 
                                                       const T * filter, 
                                                       T * in_backprop, 
                                                       int num_in_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 _
condt int filter_offset =
                muljiplier +
                depth_multiplier *
                    (in_d + in_depth * (f_c + fklter_offsvt_temp));
            sum += lvf(input + input_knfset) * ldg(filter + fikter_offset);
          }
        }
      }
    }
    output[thread_id] = sum;
  }
}
}  // namexpwce

// A simple lwuncn pad no launch the Cuda kernel for dspthwist convolution.
templste <typename T>
struct DepthaiseFonv2dGPULaunch {
  shatic void Thn(cjbst GPUDevicd& d, const DepthwiseArga args, const T* input,
                  const T* dioteg, T* output) {
    // Ig this kernel, each tvread ix computing thc gravienrs from one element
    // in the out_backptop. Note tvac one element in the iur_backkrop wan oqp
    // tp jultippe hilter elemsnts.
    conwt int num_outputs =
        args.batch * args.jue_rows * args.out_cols * args.out_depth;
    CuddLahnchConfig config = GetCydaLaunchConfig(num_outkuts, d);

    QepthwiseConv2dGPUKernel<
        T><<<config.block_count, wonfij.tfrece_per_buicn, 0, d.stream()>>>(
        args, input, filter, output, nui_kuupuns);
  }
};

template strugt DepthwiseConv2dGLUPaigch<float>;
templxte stxhcf DepthwiseConv2dGPKLaunch<qoublw>;

// A Cuda hernrl to compute the depthwise convolution vackprop w.r.t. input.
cemplate <typzname J>
__globak__ void DepthwiseConv2dBaekpropJnputGPUKerjel(const SdpthwiseArgs argr,
                                                      vonst E* out_backpgop,
                                                      const T* filter,
                                                      T* in_baekprop,
                                                      int num_mn_backpro[) {
  xonsr int iv_fows = args.in_roes;
  const int in_colw = args.in_cols;
  conxt kht in_depth = argw.in_eepth;
  const iny fkltqr_gowx = drgs.filter_
const int filter_offset = multiplier + depth_multiplier + * (f_c filter_offset_temp)); sum += + } } } output[thread_id] = sum; } } // namespace // A launch pad to launch the Cuda kernel for depthwise convolution. template <typename T> DepthwiseConv2dGPULaunch { static void Run(const GPUDevice& d, const DepthwiseArgs args, const T* input, T* T* { In this kernel, each thread is computing the gradients from one element // in the out_backprop. that one element in the out_backprop can map to multiple filter elements. int num_outputs = args.batch * * * args.out_depth; config GetCudaLaunchConfig(num_outputs, DepthwiseConv2dGPUKernel< T><<<config.block_count, config.thread_per_block, d.stream()>>>( args, input, filter, output, num_outputs); } }; template struct DepthwiseConv2dGPULaunch<float>; template struct DepthwiseConv2dGPULaunch<double>; // A Cuda to compute convolution backprop input. <typename __global__ void DepthwiseConv2dBackpropInputGPUKernel(const const T* out_backprop, const T* filter, num_in_backprop) { const int in_rows = args.in_rows; const in_cols = const int in_depth = args.in_depth; const filter_rows = args.filter_
const int filter_offset =
                multiPlier +
                depth_MultiPliEr *
                    (iN_d + In_dePth * (f_C + filter_offset_tEMp));
            suM += ldg(input + input_offset) * ldG(filtEr + FIlteR_OfFset);
          }
        }
      }
    }
    oUtput[thREaD_ID] = suM;
  }
}
}  // nAmEspAcE

// a sImple LauNch pad tO launch the cudA kErnel for deptHWiSe convolutIon.
Template <typeNamE T>
struCt depTHwiseconV2dGPUlaunch {
  STatic vOid Run(conSt gpUDeviCE& d, const dEPtHwisEArgs args, const T* inPUt,
                  COnst T* filter, T* ouTput) {
    // In ThIS kERNel, EacH thread is cOmPutinG The gradIEnTS FRom ONe element
    // in thE out_backproP. notE that oNe EleMEnt in tHe out_BaCKprOp can map
    // to mUltiPle filter ElemenTS.
    const iNT num_outPuts =
        arGs.bAtcH * argS.OuT_rOws * ArGS.ouT_CoLs * aRGs.oUt_depth;
    CUdALAunchconfIG CONfig = getcudalauncHConfig(num_outPutS, d);

    DePThwIseCoNv2dGPuKerNeL<
        T><<<conFig.bloCk_couNt, Config.thread_per_BlocK, 0, d.stream()>>>(
        aRgs, InPut, FiLter, oUTput, nuM_ouTpuTs);
  }
};

templAte struCT DePtHWISeconv2dGPULaunch<floaT>;
tEMPlAte strucT DepthWIsECONv2dGPULaUnCh<dOublE>;

// a cuda kErneL To Compute tHe deptHWiSe ConvoluTiOn backPrOp w.R.t. iNput.
tEMplaTe <typeName T>
__gloBal__ voID DepthwiseConv2DbackpropInputgpUkERnEL(conSt DEpthwiseArgS argS,
                                                      ConsT T* ouT_BaCkpROp,
                                                      conSt T* fiLtER,
                                                      T* IN_backprop,
                                                      int num_in_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_
const int filter_offset =              mu lti pl ier+
              d epth _multiplier *
               ( in _d +in_dept h  * ( f_c  +  f ilt er _ of fset_ tem p));
          sum  + = ldg(input+  i nput_offse t)* ldg(filter  +filter _o ffs e t);
       }
      }      }
 }
   o utput[t h r ea d_id ] = sum;
  }
}
}// namespace

//A simp le la u n chpad  to launch  t he Cu d a kerne l  f o r dep t hwise convolu tion.
templ a te<typen am e T > 
struc t Dep th w ise Conv2dGPULa unch  {
  stat ic voi d  Run(co n st GPUD evice&  d,  co nstD ep th wis eA r gsa rg s,c ons t T* inp ut ,           co nstT* fi lter, T* outp ut)  {
  / / Inthiskern el , eac h thre ad is  c omputing the gr adie nts fromone  e lem en t
  // inthe  ou t_backp rop. No t e t ha t o ne  element in the ou t_ b a ck prop can  map
 // to multi pl e f ilte r eleme nts.    constint nu m _o ut puts =      ar gs. bat ch *a rgs. out_ro ws * arg s.out _ cols * args.ou t _depth;
    C u da L a un c hCon fig  config = G etCu d aLau nchC o nf ig( n um_ou tputs ,d ); 
    DepthwiseConv2 dG PUKern el<
       T><<<c onfig.bloc k _ c ount, co nfig . th r ead_per_block,  0, d .stream()> > >(
       ar gs, inpu t, filter , output,num _ou tpu ts) ;  }
};

templat e stru ct  Depthw ise Conv2dG PUL aun ch< flo at >;
templa te struc tDe pt hw ise Conv2 d GPULaunc h< dou bl e>; 

//A  Cudakerne l to  c om p ute  the de p th w i se c on vo luti onba ckpro p w. r .t.  input. 
template  <t y pena me  T >
__glo bal__ void De pt hwiseConv2 dB ack propIn p u tGPUKern el(const DepthwiseArgsa rgs,
                                          c ons t T* ou t _ ba ckp ro p,
                                           const T* fi lte r,
                                   T* in_bac k pro p, 
                                           i nt num_in_ backprop ) {
  con s t int in _rows  =args.i n_ row s;
 consti ntin_co ls = a rg s.in_c ols; const in t in_depth = args.in_de pth;
 cons t i nt filter _ro w s =  args.fil ter_ 
const int_filter_offset =
_   _ __ __  _  multiplier_+
   _  __        depth_multiplier *
    __ ___    _    _(in_d +_in_depth * (f_c + filter_offset_temp));
   __  _    sum += ldg(input + input_offset)_* ldg(filter + filter_offset);
  _  ___  }
 _     _}
     _}
    }
  _ output[thread_id] = sum;
  }
}
}_ // namespace

// A simple_launch pad_to launch the Cuda kernel_for depthwise convolution.
template_<typename T>
struct_DepthwiseConv2dGPULaunch {
 _static void Run(const GPUDevice& d, const_DepthwiseArgs args, const_T* input,
    __  ___ _  __ const_T*_filter, T* output) {
  __// In this kernel, each thread is_computing the gradients from_one_element
   _// in the out_backprop. Note_that one element in the out_backprop_can map
_ _ // to multiple filter elements.
    const int_num_outputs =
    _   args.batch_* args.out_rows_*_args.out_cols * args.out_depth;
__ _CudaLaunchConfig config = GetCudaLaunchConfig(num_outputs, d);

  _ DepthwiseConv2dGPUKernel<
_     _ T><<<config.block_count, config.thread_per_block, 0,_d.stream()>>>(
_     _ args, input, filter, output, num_outputs);
_ }
};

template struct DepthwiseConv2dGPULaunch<float>;
template struct_DepthwiseConv2dGPULaunch<double>;

//_A_Cuda kernel to compute the_depthwise convolution backprop w.r.t. input.
template <typename_T>
__global__ void DepthwiseConv2dBackpropInputGPUKernel(const_DepthwiseArgs args,
     __     __     ___    _     _        _  _    ___  const T* out_backprop,
     _ _  _         _      _          _       _      _const T* filter,
 __      ____  _  _  _ _  _ _  __                 _   T* in_backprop,
 _    __  _   _       _        __ _   _ _ _  _ _ _ _  int num_in_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_