message: Dropout layer updated to use binary_op
diff: */
  #include "HugeCTR/include/layers/dropout_layer.hpp"
  +#include <linalg/binary_op.cuh>
  #include <algorithm>
  #include <cstdio>
  namespace HugeCTR {
  -namespace {
  -
  -template <typename T>
  -__global__ void dropout_kernel(const T* in, const float* mask, T* out, const int len,
  - const float rate, const float scale) {
  - for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < len; i += blockDim.x * gridDim.x) {
  - out[i] = TypeConvertFunc<T, float>::convert(((1.f - mask[i]) >= rate) *
  - TypeConvertFunc<float, T>::convert(in[i]) * scale);
  - }
  -}
  -
  -} // end namespace
  -
  template <typename T>
  DropoutLayer<T>::DropoutLayer(const std::shared_ptr<Tensor<T>>& in_tensor,
  const std::shared_ptr<Tensor<T>>& out_tensor, float rate,
  @@ -112,9 +100,10 @@ template <typename T>
  void DropoutLayer<T>::prop_common(const T* in, T* out, cudaStream_t stream) {
  int len = in_tensors_[0]->get_num_elements();
  - int grid_size = n_sms_ * 16;
  - int block_size = 256;
  - dropout_kernel<<<grid_size, block_size, 0, stream>>>(in, mask_, out, len, rate_, scale_);
  + float r = rate_;
  + float s = scale_;
  + MLCommon::LinAlg::binaryOp(out, in, mask_, len,
  + [r, s] __device__(T a, T b) { return ((T(1.0) - b) >= r) * a * s; }, stream);
  #ifndef NDEBUG
  cudaDeviceSynchronize();
  