From 2389d3f1a3d0d169fe6371efe922808df0e8b698 Mon Sep 17 00:00:00 2001 From: liuqi Date: Thu, 1 Mar 2018 09:37:31 +0800 Subject: [PATCH] Rename fc to fully_connected to keep consistency. --- mace/kernels/fully_connected.h | 17 ++++++++--------- .../opencl/cl/{fc.cl => fully_connected.cl} | 18 +++++++++--------- mace/kernels/opencl/fully_connected_opencl.cc | 6 +++--- 3 files changed, 20 insertions(+), 21 deletions(-) rename mace/kernels/opencl/cl/{fc.cl => fully_connected.cl} (74%) diff --git a/mace/kernels/fully_connected.h b/mace/kernels/fully_connected.h index d95d8a48..c3b66222 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 ced7adf7..021012ff 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 06dc74e2..f719bb01 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++, -- GitLab