Improve GPU path for logsigmoid

Author: bartekkuncer

3rdparty/onednn

@@ -36,9 +36,10 @@
 8. Activation
     8.1 relu
     8.2 sigmoid
-    8.3 softrelu
-    8.4 softsign
-    8.5 tanh
+    8.3 log_sigmoid
+    8.4 softrelu
+    8.5 softsign
+    8.6 tanh
 
 """
 

benchmark/opperf/nd_operations/nn_activation_operators.py

@@ -375,7 +375,7 @@
 
 # For NN operators
 DEFAULT_ACT_TYPE_LR = ['leaky', 'elu', 'selu', 'gelu']
-DEFAULT_ACT_TYPE_ACTIVATION = ['relu', 'sigmoid', 'softrelu', 'softsign', 'tanh']
+DEFAULT_ACT_TYPE_ACTIVATION = ['relu', 'sigmoid', 'log_sigmoid', 'softrelu', 'softsign', 'tanh']
 DEFAULT_LABEL_SOFTMAX = [(1024, 1024), (10000, 1), (10000, 100)]
 
 DEFAULT_LABEL_SOFTMAX_LARGE_TENSOR = [(2**32, 1)]

benchmark/opperf/rules/default_params.py

@@ -225,7 +225,6 @@ struct ActivationParam : public dmlc::Parameter<ActivationParam> {
     DMLC_DECLARE_FIELD(act_type)
     .add_enum("relu", activation::kReLU)
     .add_enum("sigmoid", activation::kSigmoid)
-    .add_enum("log_sigmoid", activation::kLogSigmoid)
     .add_enum("tanh", activation::kTanh)
     .add_enum("softrelu", activation::kSoftReLU)
     .describe("Activation function to be applied.");

docs/static_site/src/pages/api/faq/new_op.md

@@ -2171,6 +2171,14 @@ def log1p(self, *args, **kwargs):
         """
         return op.log1p(self, *args, **kwargs)
 
+    def log_sigmoid(self, *args, **kwargs):
+        """Convenience fluent method for :py:func:`log_sigmoid`.
+
+        The arguments are the same as for :py:func:`log_sigmoid`, with
+        this array as data.
+        """
+        return op.log_sigmoid(self, *args, **kwargs)
+
     def sqrt(self, *args, **kwargs):
         """Convenience fluent method for :py:func:`sqrt`.
 

python/mxnet/ndarray/ndarray.py

@@ -2260,6 +2260,14 @@ def log1p(self, *args, **kwargs):
         """
         raise AttributeError('mxnet.numpy.ndarray object has no attribute log1p')
 
+    def log_sigmoid(self, *args, **kwargs):
+        """Convenience fluent method for :py:func:`log_sigmoid`.
+
+        The arguments are the same as for :py:func:`log_sigmoid`, with
+        this array as data.
+        """
+        raise AttributeError('mxnet.numpy.ndarray object has no attribute log_sigmoid')
+
     def sqrt(self, *args, **kwargs):
         """Convenience fluent method for :py:func:`sqrt`.
 

python/mxnet/numpy/multiarray.py

@@ -2519,6 +2519,14 @@ def log1p(self, *args, **kwargs):
         """
         return op.log1p(self, *args, **kwargs)
 
+    def log_sigmoid(self, *args, **kwargs):
+        """Convenience fluent method for :py:func:`log_sigmoid`.
+
+        The arguments are the same as for :py:func:`log_sigmoid`, with
+        this array as data.
+        """
+        return op.log_sigmoid(self, *args, **kwargs)
+
     def sqrt(self, *args, **kwargs):
         """Convenience fluent method for :py:func:`sqrt`.
 
@@ -2583,14 +2591,6 @@ def sigmoid(self, *args, **kwargs):
         """
         return op.sigmoid(self, *args, **kwargs)
 
-    def log_sigmoid(self, *args, **kwargs):
-        """Convenience fluent method for :py:func:`log_sigmoid`.
-
-        The arguments are the same as for :py:func:`log_sigmoid`, with
-        this array as data.
-        """
-        return op.log_sigmoid(self, *args, **kwargs)
-
     def softmax(self, *args, **kwargs):
         """Convenience fluent method for :py:func:`softmax`.
 

python/mxnet/symbol/symbol.py

@@ -34,6 +34,8 @@ inline int String2MXNetActType(const std::string& s) {
     return activation::kReLU;
   } else if (s == "sigmoid") {
     return activation::kSigmoid;
+  } else if (s == "log_sigmoid") {
+    return activation::kLogSigmoid;
   } else if (s == "tanh") {
     return activation::kTanh;
   } else if (s == "softrelu") {

src/api/operator/numpy_extension/npx_activation_op.cc

@@ -40,8 +40,14 @@ backward_relu(const DTypeGrad grad, const DType val) {
 
 template <typename DType, typename DTypeGrad>
 __device__ inline mixed_type<DTypeGrad, DType>
-backward_sigmoid(const DTypeGrad grad, const DType out) {
-  return grad * out * (1 - out);
+backward_sigmoid(const DTypeGrad grad, const DType val) {
+  return grad * val * (1 - val);
+}
+
+template <typename DType, typename DTypeGrad>
+__device__ inline mixed_type<DTypeGrad, DType>
+backward_log_sigmoid(const DTypeGrad grad, const DType val) {
+  return grad * 1 / (1 + op::exp(val));
 }
 
 template <typename DType, typename DTypeGrad>

src/common/cuda/rtc/backward_functions-inl.h

@@ -685,6 +685,15 @@ __device__ inline DType sigmoid(const DType val) {
   }
 }
 
+template <typename DType>
+__device__ inline DType log_sigmoid(const DType val) {
+  if (type_util::has_double_or_integral<DType>::value) {
+    return ::log(1./(1 + ::exp(-val)));
+  } else {
+    return ::logf(1.f/(1 + expf(-val)));
+  }
+}
+
 template <typename DType>
 __device__ inline DType softrelu(const DType val) {
   if (type_util::has_double_or_integral<DType>::value) {