Remove double precision math from LogSigmoid too

Author: apaszke

LogSigmoid.cu

@@ -6,22 +6,49 @@
 template <typename T>
 struct logSigmoid_updateOutput_functor
 {
-  __device__ void operator()(T *output, const T *input) const
-  {
-    T z = exp(-*input);
-    *output = ScalarConvert<double, T>::to(-log(1. + z));
+  __device__ void operator()(T *output, const T *input) const {
+    *output = -THCNumerics<T>::log(1.f + THCNumerics<T>::exp(- *input));
   }
 };
 
 template <typename T>
 struct logSigmoid_updateGradInput_functor
 {
-  __device__ void operator()(T *gradInput, const T *input, const T *gradOutput) const
-  {
-    T z = exp(-*input);
-    *gradInput = ScalarConvert<double, T>::to(*gradOutput * z / (1. + z));
+  __device__ void operator()(T *gradInput, const T *input, const T *gradOutput) const {
+    const T z = THCNumerics<T>::exp(- *input);
+    *gradInput = *gradOutput * z / (1.f + z);
   }
 };
 
+#ifdef CUDA_HALF_TENSOR
+template <>
+struct logSigmoid_updateOutput_functor<half> {
+  __device__ __forceinline__ void operator()(half* output, const half *input) const {
+#ifdef CUDA_HALF_INSTRUCTIONS
+    const half one = __float2half(1.f);
+    *output = __hneg(THCNumerics<half>::log(one + THCNumerics<half>::exp(__hneg(*input))));
+#else
+    float in = __half2float(*input);
+    *output = __float2half(-THCNumerics<float>::log(1.f + THCNumerics<float>::exp(-in)));
+#endif
+  }
+};
+
+template <>
+struct logSigmoid_updateGradInput_functor<half> {
+  __device__ __forceinline__ void operator()(half* gradInput, const half *input, const half *gradOutput) const {
+#ifdef CUDA_HALF_INSTRUCTIONS
+    const half one = __float2half(1.f);
+    const half in_exp = THCNumerics<half>::exp(__hneg(*input));
+    *gradInput = hdiv(__hmul(*gradOutput, in_exp), __hadd(one, in_exp));
+#else
+    const float in_exp = THCNumerics<float>::exp(-(__half2float(*input)));
+    const float go = __half2float(*gradOutput);
+    *gradInput = __float2half(go * in_exp / (1.f + in_exp));
+#endif
+  }
+};
+#endif
+
 #include "generic/LogSigmoid.cu"
 #include "THCGenerateFloatTypes.h"

Sigmoid.cu

@@ -4,22 +4,22 @@
 #include <THC/THCApply.cuh>
 
 template <typename T>
-struct SigmoidGradInputOp {
+struct sigmoid_updateGradInput_functor {
   __device__ __forceinline__ void operator()(T* gradInput, const T *output, const T *gradOutput) const {
     *gradInput = *gradOutput * (1.f - *output) * (*output);
   }
 };
 
 #ifdef CUDA_HALF_TENSOR
 template <>
-struct SigmoidGradInputOp<half> {
+struct sigmoid_updateGradInput_functor<half> {
   __device__ __forceinline__ void operator()(half* gradInput, const half *output, const half *gradOutput) const {
 #ifdef CUDA_HALF_INSTRUCTIONS
-    half one = __float2half(1.f);
+    const half one = __float2half(1.f);
     *gradInput = __hmul(*gradOutput, __hmul(__hadd(one, __hneg(*output)), *output));
 #else
-    float out = __half2float(*output);
-    float go = __half2float(*gradOutput);
+    const float out = __half2float(*output);
+    const float go = __half2float(*gradOutput);
     *gradInput = __float2half(go * (1.f - out) * out);
 #endif
   }

Tanh.cu

@@ -4,7 +4,7 @@
 #include <THC/THCApply.cuh>
 
 template <typename T>
-struct TanhGradInputOp
+struct tanh_updateGradInput_functor
 {
   __device__ __forceinline__ void operator()(T *gradInput,
           const T *output, const T *gradOutput) const {
@@ -14,7 +14,7 @@ struct TanhGradInputOp
 
 #ifdef CUDA_HALF_TENSOR
 template <>
-struct TanhGradInputOp<half>
+struct tanh_updateGradInput_functor<half>
 {
   __device__ __forceinline__ void operator()(half *gradInput,
           const half *output, const half *gradOutput) const {
@@ -23,8 +23,8 @@ struct TanhGradInputOp<half>
     const half out_square = __hmul(*output, *output);
     *gradInput = __hmul(*gradOutput, __hadd(one, __hneg(out_square)));
 #else
-    float out = __half2float(*output);
-    float go = __half2float(*gradOutput);
+    const float out = __half2float(*output);
+    const float go = __half2float(*gradOutput);
     *gradInput = __float2half(go * (1.f - out * out));
 #endif
   }

generic/Sigmoid.cu

@@ -20,10 +20,10 @@ void THNN_(Sigmoid_updateGradInput)(
            THCTensor *gradInput,
            THCTensor *output)
 {
-  THCUNN_check_nElement(state, input, gradOutput);
+  THCUNN_check_nElement(state, output, gradOutput);
   THCUNN_assertSameGPU(state, 3, output, gradOutput, gradInput);
   THCTensor_(resizeAs)(state, gradInput, output);
-  THC_pointwiseApply3(state, gradInput, output, gradOutput, SigmoidGradInputOp<real>());
+  THC_pointwiseApply3(state, gradInput, output, gradOutput, sigmoid_updateGradInput_functor<real>());
 }
 
 #endif

generic/THCUNN.h

@@ -911,7 +911,7 @@ TH_API void THNN_(Sigmoid_updateOutput)(
 
 TH_API void THNN_(Sigmoid_updateGradInput)(
                   THCState *state,
-                  THCTensor *input,
+                  THCTensor *input,          // [OPTIONAL]
                   THCTensor *gradOutput,
                   THCTensor *gradInput,
                   THCTensor *output);
@@ -1002,7 +1002,7 @@ TH_API void THNN_(Tanh_updateOutput)(
 
 TH_API void THNN_(Tanh_updateGradInput)(
                   THCState *state,
-                  THCTensor *input,
+                  THCTensor *input,          // [OPTIONAL]
                   THCTensor *gradOutput,
                   THCTensor *gradInput,
                   THCTensor *output);

generic/Tanh.cu

@@ -24,7 +24,7 @@ void THNN_(Tanh_updateGradInput)(
   THCUNN_check_shape(state, output, gradOutput);
   THCUNN_assertSameGPU(state, 3, output, gradOutput, gradInput);
   THCTensor_(resizeAs)(state, gradInput, output);
-  THC_pointwiseApply3(state, gradInput, output, gradOutput, TanhGradInputOp<real>());
+  THC_pointwiseApply3(state, gradInput, output, gradOutput, tanh_updateGradInput_functor<real>());
 }
 
 #endif