From be1a6e22157e698898ecca51c8fd5b766a33e29e Mon Sep 17 00:00:00 2001 From: Andrey Stotskiy Date: Mon, 16 Mar 2020 16:47:14 +0300 Subject: [PATCH 1/2] fix numerically unstable fused softrelu op --- src/operator/fusion/fused_op-inl.h | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/src/operator/fusion/fused_op-inl.h b/src/operator/fusion/fused_op-inl.h index 7373cd07400a..0504149dfc68 100644 --- a/src/operator/fusion/fused_op-inl.h +++ b/src/operator/fusion/fused_op-inl.h @@ -550,7 +550,10 @@ __device__ inline DType sigmoid(const DType val) { template __device__ inline DType softrelu(const DType val) { - return logf(1 + expf(val)); + // Avoid overflow of exp for large inputs. + // The threshold 20 is chosen such that softrelu(a) = a + // for a > 20 using floating precision. + return val > 20 ? val : logf(1 + expf(val)); } template From 31fae17ed0e3023f4439f85cdf6a76db48f57e12 Mon Sep 17 00:00:00 2001 From: Andrey Stotskiy Date: Mon, 16 Mar 2020 16:59:12 +0300 Subject: [PATCH 2/2] implement test for softrelu numerical stability --- tests/python/gpu/test_fusion.py | 3 +++ 1 file changed, 3 insertions(+) diff --git a/tests/python/gpu/test_fusion.py b/tests/python/gpu/test_fusion.py index 9a37c0e844a0..c0a8bdbb0807 100644 --- a/tests/python/gpu/test_fusion.py +++ b/tests/python/gpu/test_fusion.py @@ -136,6 +136,9 @@ def announce_check(op_name): for act_type in ['relu', 'sigmoid', 'tanh', 'softrelu', 'softsign']: announce_check("Activation(act_type='{}')".format(act_type)) check_fused_symbol(mx.sym.Activation(a, act_type=act_type), a=arr) + if act_type == 'softrelu': + # Check that softrelu implementation doesn't overflow on large inputs + check_fused_symbol(mx.sym.Activation(a, act_type=act_type), a=1000 * arr) # Cast requires dtype for dtype in ['float16', 'float32', 'float64', 'int32']: