diff --git a/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/activation_gpu_kernel.cc b/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/activation_gpu_kernel.cc index 334550b213..157110d25c 100644 --- a/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/activation_gpu_kernel.cc +++ b/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/activation_gpu_kernel.cc @@ -23,6 +23,11 @@ MS_REG_GPU_KERNEL_ONE(ReLU, KernelAttr().AddInputAttr(kNumberTypeFloat32).AddOut MS_REG_GPU_KERNEL_ONE(ReLU, KernelAttr().AddInputAttr(kNumberTypeFloat16).AddOutputAttr(kNumberTypeFloat16), ActivationGpuFwdKernel, half) +MS_REG_GPU_KERNEL_ONE(ReLU6, KernelAttr().AddInputAttr(kNumberTypeFloat32).AddOutputAttr(kNumberTypeFloat32), + ActivationGpuFwdKernel, float) +MS_REG_GPU_KERNEL_ONE(ReLU6, KernelAttr().AddInputAttr(kNumberTypeFloat16).AddOutputAttr(kNumberTypeFloat16), + ActivationGpuFwdKernel, half) + MS_REG_GPU_KERNEL_ONE(Tanh, KernelAttr().AddInputAttr(kNumberTypeFloat32).AddOutputAttr(kNumberTypeFloat32), ActivationGpuFwdKernel, float) MS_REG_GPU_KERNEL_ONE(Tanh, KernelAttr().AddInputAttr(kNumberTypeFloat16).AddOutputAttr(kNumberTypeFloat16), diff --git a/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/activation_gpu_kernel.h b/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/activation_gpu_kernel.h index 14fc721889..b49e3b86a1 100644 --- a/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/activation_gpu_kernel.h +++ b/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/activation_gpu_kernel.h @@ -83,7 +83,8 @@ class ActivationGpuFwdKernel : public GpuKernel { return true; } std::vector shape; - CHECK_CUDNN_RET_WITH_EXCEPT(cudnnSetActivationDescriptor(activation_desc_, mode_, CUDNN_NOT_PROPAGATE_NAN, 0.0), + double coef = (mode_ == CUDNN_ACTIVATION_CLIPPED_RELU) ? 6.0 : 0.0; + CHECK_CUDNN_RET_WITH_EXCEPT(cudnnSetActivationDescriptor(activation_desc_, mode_, CUDNN_NOT_PROPAGATE_NAN, coef), "cudnnSetActivationDescriptor failed"); const int split_dim = 4; @@ -132,6 +133,7 @@ class ActivationGpuFwdKernel : public GpuKernel { } std::map kernel_map = {{"ReLU", CUDNN_ACTIVATION_RELU}, + {"ReLU6", CUDNN_ACTIVATION_CLIPPED_RELU}, {"Tanh", CUDNN_ACTIVATION_TANH}, {"ELU", CUDNN_ACTIVATION_ELU}, {"Sigmoid", CUDNN_ACTIVATION_SIGMOID}}; diff --git a/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/activation_grad_kernel.cc b/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/activation_grad_kernel.cc index 8fd486c08c..8e6f568031 100644 --- a/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/activation_grad_kernel.cc +++ b/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/activation_grad_kernel.cc @@ -27,6 +27,15 @@ MS_REG_GPU_KERNEL_ONE( KernelAttr().AddInputAttr(kNumberTypeFloat16).AddInputAttr(kNumberTypeFloat16).AddOutputAttr(kNumberTypeFloat16), ActivationGradGpuKernel, half) +MS_REG_GPU_KERNEL_ONE( + ReLU6Grad, + KernelAttr().AddInputAttr(kNumberTypeFloat32).AddInputAttr(kNumberTypeFloat32).AddOutputAttr(kNumberTypeFloat32), + ActivationGradGpuKernel, float) +MS_REG_GPU_KERNEL_ONE( + ReLU6Grad, + KernelAttr().AddInputAttr(kNumberTypeFloat16).AddInputAttr(kNumberTypeFloat16).AddOutputAttr(kNumberTypeFloat16), + ActivationGradGpuKernel, half) + MS_REG_GPU_KERNEL_ONE( TanhGrad, KernelAttr().AddInputAttr(kNumberTypeFloat32).AddInputAttr(kNumberTypeFloat32).AddOutputAttr(kNumberTypeFloat32), diff --git a/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/activation_grad_kernel.h b/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/activation_grad_kernel.h index e2f040f6be..c6fd1a0921 100644 --- a/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/activation_grad_kernel.h +++ b/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/activation_grad_kernel.h @@ -49,7 +49,7 @@ class ActivationGradGpuKernel : public GpuKernel { } T *dy = nullptr; T *y = nullptr; - if (mode_ == CUDNN_ACTIVATION_RELU || mode_ == CUDNN_ACTIVATION_ELU) { + if (mode_ == CUDNN_ACTIVATION_RELU || mode_ == CUDNN_ACTIVATION_ELU || mode_ == CUDNN_ACTIVATION_CLIPPED_RELU) { dy = GetDeviceAddress(inputs, 0); y = GetDeviceAddress(inputs, 1); } else { @@ -90,7 +90,8 @@ class ActivationGradGpuKernel : public GpuKernel { return true; } std::vector shape; - CHECK_CUDNN_RET_WITH_EXCEPT(cudnnSetActivationDescriptor(activation_desc_, mode_, CUDNN_PROPAGATE_NAN, 0.0), + double coef = (mode_ == CUDNN_ACTIVATION_CLIPPED_RELU) ? 5.999999 : 0.0; + CHECK_CUDNN_RET_WITH_EXCEPT(cudnnSetActivationDescriptor(activation_desc_, mode_, CUDNN_PROPAGATE_NAN, coef), "SetActivationDescriptor failed"); const int split_dim = 4; @@ -138,6 +139,7 @@ class ActivationGradGpuKernel : public GpuKernel { } std::map kernel_map = {{"ReluGrad", CUDNN_ACTIVATION_RELU}, + {"ReLU6Grad", CUDNN_ACTIVATION_CLIPPED_RELU}, {"TanhGrad", CUDNN_ACTIVATION_TANH}, {"ELUGrad", CUDNN_ACTIVATION_ELU}, {"SigmoidGrad", CUDNN_ACTIVATION_SIGMOID}}; diff --git a/mindspore/ccsrc/runtime/device/gpu/kernel_info_setter.h b/mindspore/ccsrc/runtime/device/gpu/kernel_info_setter.h index e6ef029fb2..2034220097 100644 --- a/mindspore/ccsrc/runtime/device/gpu/kernel_info_setter.h +++ b/mindspore/ccsrc/runtime/device/gpu/kernel_info_setter.h @@ -54,6 +54,8 @@ static std::map, std::vector> // Format insensitive. {prim::kPrimRelu->name(), {{0}, {0}}}, {prim::kPrimReluGrad->name(), {{0, 1}, {0}}}, + {prim::kPrimRelu6->name(), {{0}, {0}}}, + {prim::kPrimRelu6Grad->name(), {{0, 1}, {0}}}, {kSliceOpName, {{0}, {0}}}, {kTensorAddOpName, {{0, 1}, {0}}}, {prim::kPrimConcat->name(), {{}, {0}}}, diff --git a/mindspore/core/base/core_ops.h b/mindspore/core/base/core_ops.h index 48c07a8d03..58f15c79b4 100644 --- a/mindspore/core/base/core_ops.h +++ b/mindspore/core/base/core_ops.h @@ -127,6 +127,7 @@ inline const PrimitivePtr kPrimFusedBatchNormGradEx = std::make_shared("BatchNorm"); inline const PrimitivePtr kPrimBatchNormGrad = std::make_shared("BatchNormGrad"); inline const PrimitivePtr kPrimReluGrad = std::make_shared("ReluGrad"); +inline const PrimitivePtr kPrimRelu6Grad = std::make_shared("ReLU6Grad"); inline const PrimitivePtr kPrimConv2DBackpropInput = std::make_shared("Conv2DBackpropInput"); inline const PrimitivePtr kPrimConv2DBackpropFilter = std::make_shared("Conv2DBackpropFilter"); inline const PrimitivePtr kPrimDepthwiseConv2dNative = std::make_shared("DepthwiseConv2dNative"); @@ -151,6 +152,7 @@ inline const PrimitivePtr kPrimOneHot = std::make_shared("OneHot"); inline const PrimitivePtr kPrimGelu = std::make_shared("Gelu"); inline const PrimitivePtr kPrimGeluGrad = std::make_shared("GeluGrad"); inline const PrimitivePtr kPrimRelu = std::make_shared("ReLU"); +inline const PrimitivePtr kPrimRelu6 = std::make_shared("ReLU6"); inline const PrimitivePtr kPrimReluV2 = std::make_shared("ReLUV2"); inline const PrimitivePtr kPrimZerosLike = std::make_shared("ZerosLike"); inline const PrimitivePtr kPrimBpropCut = std::make_shared("bprop_cut");