diff --git a/mace/kernels/fully_connected.h b/mace/kernels/fully_connected.h index d95d8a48ce3e8e8a7f105cd99250d364b59c38b1..c3b662222a58d8dea1180195e2bdb8aae0c56eca 100644 --- a/mace/kernels/fully_connected.h +++ b/mace/kernels/fully_connected.h @@ -26,12 +26,12 @@ struct FullyConnectedBase { const float prelu_alpha_; }; -template -struct FullyConnectedFunctor : FullyConnectedBase{ +template +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 -struct FullyConnectedFunctor : FullyConnectedBase{ +template +struct FullyConnectedFunctor : 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, diff --git a/mace/kernels/opencl/cl/fc.cl b/mace/kernels/opencl/cl/fully_connected.cl similarity index 74% rename from mace/kernels/opencl/cl/fc.cl rename to mace/kernels/opencl/cl/fully_connected.cl index ced7adf70f89ba8703d5553f5a545ab7c9686e4c..021012ffe39eb5f9a3744267bdb088001123f1aa 100644 --- a/mace/kernels/opencl/cl/fc.cl +++ b/mace/kernels/opencl/cl/fully_connected.cl @@ -1,17 +1,17 @@ #include // 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; diff --git a/mace/kernels/opencl/fully_connected_opencl.cc b/mace/kernels/opencl/fully_connected_opencl.cc index 06dc74e2d83a19aa655665f175577c90f6ddca29..f719bb0100e763514360d4bed837ba12e2ee9e0d 100644 --- a/mace/kernels/opencl/fully_connected_opencl.cc +++ b/mace/kernels/opencl/fully_connected_opencl.cc @@ -32,8 +32,8 @@ void FullyConnectedFunctor::operator()( auto runtime = OpenCLRuntime::Global(); std::set built_options; auto dt = DataTypeToEnum::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::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++,