Rename fc to fully_connected to keep consistency.

mace/kernels/fully_connected.h

@@ -26,12 +26,12 @@ struct FullyConnectedBase {
   const float prelu_alpha_;
 };
 
-template <DeviceType D, typename T>
-struct FullyConnectedFunctor : FullyConnectedBase{
+template<DeviceType D, typename T>
+struct FullyConnectedFunctor : FullyConnectedBase {
   FullyConnectedFunctor(const ActivationType activation,
-                     const float relux_max_limit,
-                     const float prelu_alpha) :
-      FullyConnectedBase(activation, relux_max_limit, prelu_alpha){}
+                        const float relux_max_limit,
+                        const float prelu_alpha) :
+      FullyConnectedBase(activation, relux_max_limit, prelu_alpha) {}
 
   void operator()(const Tensor *input,
                   const Tensor *weight,
@@ -74,13 +74,12 @@ struct FullyConnectedFunctor : FullyConnectedBase{
   }
 };
 
-
-template <typename T>
-struct FullyConnectedFunctor<DeviceType::OPENCL, T> : FullyConnectedBase{
+template<typename T>
+struct FullyConnectedFunctor<DeviceType::OPENCL, T> : FullyConnectedBase {
   FullyConnectedFunctor(const ActivationType activation,
                         const float relux_max_limit,
                         const float prelu_alpha) :
-      FullyConnectedBase(activation, relux_max_limit, prelu_alpha){}
+      FullyConnectedBase(activation, relux_max_limit, prelu_alpha) {}
 
   void operator()(const Tensor *input,
                   const Tensor *weight,

mace/kernels/opencl/cl/fc.cl → mace/kernels/opencl/cl/fully_connected.cl

 #include <common.h>
 
 // output = weight * input + bias
-__kernel void fc(__read_only image2d_t input,
-                 __read_only image2d_t weight,
+__kernel void fully_connected(__read_only image2d_t input,
+                              __read_only image2d_t weight,
 #ifdef BIAS
-                 __read_only image2d_t bias,
+                              __read_only image2d_t bias,
 #endif
-                 __write_only image2d_t output,
-                 __private const int input_height,
-                 __private const int input_width,
-                 __private const int input_channel,
-                 __private const float relux_max_limit,
-                 __private const float prelu_alpha) {
+                              __write_only image2d_t output,
+                              __private const int input_height,
+                              __private const int input_width,
+                              __private const int input_channel,
+                              __private const float relux_max_limit,
+                              __private const float prelu_alpha) {
   const int batch_idx = get_global_id(0);
   const int out_blk_idx = get_global_id(1);
   const int input_chan_blk = (input_channel + 3) >> 2;

mace/kernels/opencl/fully_connected_opencl.cc

@@ -32,8 +32,8 @@ void FullyConnectedFunctor<DeviceType::OPENCL, T>::operator()(
     auto runtime = OpenCLRuntime::Global();
     std::set<std::string> built_options;
     auto dt = DataTypeToEnum<T>::value;
-    std::string kernel_name = MACE_OBFUSCATE_SYMBOL("fc");
-    built_options.emplace("-Dfc=" + kernel_name);
+    std::string kernel_name = MACE_OBFUSCATE_SYMBOL("fully_connected");
+    built_options.emplace("-Dfully_connected=" + kernel_name);
     built_options.emplace("-DDATA_TYPE=" + DtToUpstreamCLDt(dt));
     built_options.emplace("-DCMD_DATA_TYPE=" + DtToUpstreamCLCMDDt(dt));
     if (bias != nullptr) {
@@ -60,7 +60,7 @@ void FullyConnectedFunctor<DeviceType::OPENCL, T>::operator()(
       default:
         LOG(FATAL) << "Unknown activation type: " << activation_;
     }
-    kernel_ = runtime->BuildKernel("fc", kernel_name, built_options);
+    kernel_ = runtime->BuildKernel("fully_connected", kernel_name, built_options);
 
     uint32_t idx = 0;
     kernel_.setArg(idx++,