[Metal] fix_relu and relu6 Add_exp (#8082)

* fix_relu_exp_relu6

* fix_relu_exp_relu6

* swish

Author: xiaoxiaohehe001

lite/backends/metal/metal_kernel/texture/ActivationKernel.metal

@@ -107,15 +107,17 @@ kernel void hard_sigmoid(texture2d_array<ftype, access::sample> inTexture[[textu
     outTexture.write(output, gid.xy, gid.z);
 }
 
+// activation function: swish
 kernel void swish(texture2d_array<ftype, access::sample> inTexture[[texture(0)]],
     texture2d_array<ftype, access::write> outTexture[[texture(1)]],
+    constant SwishParam& param[[buffer(0)]],
     uint3 gid[[thread_position_in_grid]]) {
     if (gid.x >= outTexture.get_width() || gid.y >= outTexture.get_height() ||
         gid.z >= outTexture.get_array_size())
         return;
     constexpr sampler s(coord::pixel, filter::nearest, address::clamp_to_zero);
     const ftype4 input = inTexture.read(gid.xy, gid.z);
-    const ftype4 output = input / (1.0 + exp(-input));
+    const ftype4 output = input / (1.0 + exp(-(input * param.beta)));
     outTexture.write(output, gid.xy, gid.z);
 }
 

lite/backends/metal/metal_kernel/texture/Common.metal

@@ -196,6 +196,10 @@ struct HardSigmoidParam {
     float offset;
 };
 
+struct SwishParam {
+    float beta;
+};
+
 struct HardSwishParam {
     float offset;
     float threshold;

lite/kernels/metal/image_op/activation_image_compute.mm

@@ -64,9 +64,16 @@
         case 5:
             function_name_ = "sigmoid";
             break;
-        case 7:
+        case 7: {
+            SwishMetalParam metal_param{param.Swish_beta};
+            param_buffer_ =
+                std::make_shared<MetalBuffer>(metal_context_, sizeof(metal_param), &metal_param);
+        }
             function_name_ = "swish";
             break;
+        case 8:
+            function_name_ = "exp";
+            break;
         case 10: {
             HardSwishMetalParam metal_param{
                 param.hard_swish_offset, param.hard_swish_threshold, param.hard_swish_scale};
@@ -106,8 +113,9 @@
     auto encoder = [backend commandEncoder];
     [encoder setTexture:input_buffer_->image() atIndex:(0)];
     [encoder setTexture:output_buffer_->image() atIndex:(1)];
-    if (function_name_ == "relu" || function_name_ == "leaky_relu" ||
-        function_name_ == "hard_swish" || function_name_ == "hard_sigmoid") {
+    if (function_name_ == "leaky_relu" || function_name_ == "relu6" ||
+        function_name_ == "hard_swish" || function_name_ == "hard_sigmoid" ||
+        function_name_ == "swish") {
         [encoder setBuffer:param_buffer_->buffer() offset:(0) atIndex:(0)];
     }
 

lite/kernels/metal/image_op/metal_params.h

@@ -184,6 +184,10 @@ struct HardSigmoidMetalParam {
     float offset;
 };
 
+struct SwishMetalParam {
+    float beta;
+};
+
 struct HardSwishMetalParam {
     float offset;
     float threshold;