 
   // for relu backward when T is double
   __device__ __forceinline__ typename CudaVecType<T>::type Compute(
-      const typename CudaVecType<T>::type* out,
-      const typename CudaVecType<T>::type* dout);
+      const typename CudaVecType<T>::type out,
+      const typename CudaVecType<T>::type dout) {
+    return out > zero_ ? dout : zero_;
+  }
 
   // when num % vecsize != 0 this func will be used
   __device__ __forceinline__ T ComputeRemainder(const T out, const T dout) {

class BaseGPUFunctor {
 public:
  using ELEMENT_TYPE = T;
};

/* ========================================================================== */

};

template <>
__device__ __forceinline__ CudaVecType<float>::type
ReluGradGPUFunctor<float>::Compute(const CudaVecType<float>::type out,
                                   const CudaVecType<float>::type dout) {
  // relu backward : dx = out > 0 ? dout : 0;
  return make_float4((out.x > zero_) * (dout.x), (out.y > zero_) * (dout.y),
                     (out.z > zero_) * (dout.z), (out.w > zero_) * (dout.w));
}

template <>
__device__ __forceinline__ CudaVecType<float16>::type
ReluGradGPUFunctor<float16>::Compute(const CudaVecType<float16>::type out,
                                     const CudaVecType<float16>::type dout) {
// relu backward : dx = out > 0 ? dout : 0;
#ifdef __HIPCC__ || CUDA_ARCH_FP16_SUPPORTED(__CUDA_ARCH__)
  const half2 kzero = __float2half2_rn(0.0f);
  return __hmul2(__hgt2(out, kzero), dout);
#else
  const float2 xx = __half22float2(out);
  const float2 yy = __half22float2(dout);
  return __floats2half2_rn((xx.x > 0.0f) * static_cast<float>(yy.x),
                           (xx.y > 0.0f) * static_cast<float>(yy.y));
#endif
}

/* ========================================================================== */
/* ========================    leaky relu forward    ========================
 */
template <typename T>
class LeakyReluGPUFunctor : public BaseGPUFunctor<T> {
 private:
  T zero_;
  float alpha_;

 public:
  LeakyReluGPUFunctor() { zero_ = static_cast<T>(0.0f); }

  typename BaseActivationFunctor<T>::AttrPair GetAttrs() {
    return {{"alpha", &alpha_}};
  }
  // leakyrelu forward : out = x > 0 ? x : x * alpha
  __device__ __forceinline__ typename CudaVecType<T>::type Compute(
      const typename CudaVecType<T>::type in) {
    return in > zero_ ? in : static_cast<T>(alpha_) * in;
  }

  __device__ __forceinline__ T ComputeRemainder(const T in) {
    // leakyrelu forward : out = x > 0 ? x : x * alpha
    return in > zero_ ? in : static_cast<T>(alpha_) * in;
  }
};

template <>
__device__ __forceinline__ CudaVecType<float>::type
LeakyReluGPUFunctor<float>::Compute(const CudaVecType<float>::type in) {
  // leakyrelu forward : out = x > 0 ? x : x * alpha
  return make_float4((in.x > zero_) ? (in.x) : (in.x) * alpha_,
                     (in.y > zero_) ? (in.y) : (in.y) * alpha_,
                     (in.z > zero_) ? (in.z) : (in.z) * alpha_,
                     (in.w > zero_) ? (in.w) : (in.w) * alpha_);
}

template <>
__device__ __forceinline__ CudaVecType<float16>::type
LeakyReluGPUFunctor<float16>::Compute(const CudaVecType<float16>::type in) {
  // leakyrelu forward : out = x > 0 ? x : x * alpha
  const float2 xx = __half22float2(in);
  return __floats2half2_rn((xx.x > 0.0f) ? xx.x : xx.x * alpha_,
                           (xx.y > 0.0f) ? xx.y : xx.y * alpha_);
}
/* ========================================================================== */

/* ===========================  leaky relu backward   =======================
 */
template <typename T>
class LeakyReluGradGPUFunctor : public BaseGPUFunctor<T> {
 private:
  T zero_;
  float alpha_;

 public:
  LeakyReluGradGPUFunctor() { zero_ = static_cast<T>(0.0f); }

  typename BaseActivationFunctor<T>::AttrPair GetAttrs() {
    return {{"alpha", &alpha_}};
  }

  // for leaky relu backward when T is double
  __device__ __forceinline__ typename CudaVecType<T>::type Compute(
      const typename CudaVecType<T>::type in,
      const typename CudaVecType<T>::type dout) {
    // leakyrelu backward : dx = x > 0 ? dout : alpha * dout
    return in > zero_ ? dout : static_cast<T>(alpha_) * dout;
  }

  // when num % vecsize != 0 this func will be used
  __device__ __forceinline__ T ComputeRemainder(const T in, const T dout) {
    // leakyrelu backward : dx = x > 0 ? dout : alpha * dout
    return in > zero_ ? dout : static_cast<T>(alpha_) * dout;
  }

  static constexpr ActBwdOpFwdDeps FwdDeps() { return kDepX; }
};

template <>
__device__ __forceinline__ CudaVecType<float>::type
LeakyReluGradGPUFunctor<float>::Compute(const CudaVecType<float>::type in,
                                        const CudaVecType<float>::type dout) {
  // leakyrelu backward : dx = x > 0 ? dout : alpha * dout
  return make_float4((in.x > zero_) ? (dout.x) : alpha_ * (dout.x),
                     (in.y > zero_) ? (dout.y) : alpha_ * (dout.y),
                     (in.z > zero_) ? (dout.z) : alpha_ * (dout.z),
                     (in.w > zero_) ? (dout.w) : alpha_ * (dout.w));
}

template <>
__device__ __forceinline__ CudaVecType<float16>::type LeakyReluGradGPUFunctor<
    float16>::Compute(const CudaVecType<float16>::type in,
                      const CudaVecType<float16>::type dout) {
  // leakyrelu backward : dx = x > 0 ? dout : alpha * dout
  const float2 xx = __half22float2(in);
  const float2 yy = __half22float2(dout);
  return __floats2half2_rn((xx.x > 0.0f) ? yy.x : alpha_ * yy.x,
                           (xx.y > 0.0f) ? yy.y : alpha_ * yy.y);
}

/* ========================================================================== */
