COMMON_FLAGS += -DDEBUG -g -O0
NVCCFLAGS += -G
else
- COMMON_FLAGS += -DNDEBUG -O2
+ COMMON_FLAGS += -DNDEBUG -O3
endif
# cuDNN acceleration configuration.
COMMON_FLAGS += -DUSE_ACL -std=c++11
endif
+ifeq ($(USE_OPENCL), 1)
+ COMMON_FLAGS += -DUSE_OPENCL
+endif
+
#USE_PROFILING -- get profiling informations, is controled by LOGACL
#LAYER_PERF_STAT -- haitao's net profiling information
ifeq ($(USE_PROFILING), 1)
## Refer to http://caffe.berkeleyvision.org/installation.html
# Contributions simplifying and improving our build system are welcome!
-AIDDIR=/usr/local/AID
-export PKG_CONFIG_PATH=${AIDDIR}/opencv3.3.0/lib/pkgconfig
-
# cuDNN acceleration switch (uncomment to build with cuDNN).
# USE_CUDNN := 1
+AIDDIR=/usr/local/AID
+export PKG_CONFIG_PATH=${AIDDIR}/opencv3.3.0/lib/pkgconfig
+
# CPU-only switch (uncomment to build without GPU support).
CPU_ONLY := 1
USE_PROFILING := 0
USE_ACL :=1
+USE_OPENCL:=1
ACL_ROOT :=$(AIDDIR)/ComputeLibrary
ACL_INCS :=$(ACL_ROOT)/include
ACL_INCS +=$(ACL_ROOT)
# enable pretty build (comment to see full commands)
Q ?= @
+
+ifeq ($(wildcard $(AIDDIR)),)
+ACL_ROOT :=$(shell pwd)/../ComputeLibrary
+ACL_INCS :=$(ACL_ROOT)/include
+ACL_INCS +=$(ACL_ROOT)
+ACL_LIBS_DIR :=$(ACL_ROOT)/build
+OPENCV_VERSION :=
+endif
-
# CaffeOnACL
[](LICENSE)
* Caffe is a fast open framework for deep learning. See also [Caffe](https://github.com/BVLC/caffe).
### Documents
-* [Installation instructions](https://github.com/OAID/CaffeOnACL/blob/master/acl_openailab/installation.md)
-* [User Manuals PDF](https://github.com/OAID/CaffeOnACL/blob/master/acl_openailab/user_manual.pdf)
-* [Performance Report PDF](https://github.com/OAID/CaffeOnACL/blob/master/acl_openailab/performance_report.pdf)
+* [Installation instructions](acl_openailab/installation.md)
+* [User Manuals PDF](acl_openailab/user_manual.pdf)
+* [Performance Report PDF](acl_openailab/performance_report.pdf)
+* [Accuracy Report PDF](acl_openailab/accuracy_report.pdf)
### Arm Compute Library Compatibility Issues :
There are some compatibility issues between ACL and Caffe Layers, we bypass it to Caffe's original layer class as the workaround solution for the below issues
wget --no-check-certificate https://github.com/opencv/opencv/archive/3.3.0.tar.gz
tar -xvf 3.3.0.tar.gz
#### Download "gen-pkg-config-pc"
- wget https://github.com/OAID/AID-tools/raw/master/script/gen-pkg-config-pc.sh
+ wget ftp://ftp.openailab.net/tools/script/gen-pkg-config-pc.sh
+ chmod +x ./gen-pkg-config-pc.sh
#### Download "ACL"
git clone https://github.com/ARM-software/ComputeLibrary.git
- git checkout bf8b01d
+ git checkout 48bc34e
#### Download "CaffeOnACL" :
git clone https://github.com/OAID/CaffeOnACL.git
#### Download "Googletest" :
mkdir build
aarch64-linux-gnu-gcc opencl-1.2-stubs/opencl_stubs.c -Iinclude -shared -o build/libOpenCL.so
scons Werror=1 -j4 debug=0 asserts=1 neon=1 opencl=1 embed_kernels=1 os=linux arch=arm64-v8a
- wget https://github.com/OAID/AID-tools/raw/master/script/Computelibrary/Makefile
+ wget ftp://ftp.openailab.net/tools/script/Computelibrary/Makefile
sudo make install
+ sudo ~/gen-pkg-config-pc.sh /usr/local/AID
## 3.3 Build Caffe :
cd ~/CaffeOnACL
make all
make distribute
sudo make install
+ sudo ~/gen-pkg-config-pc.sh /usr/local/AID
## 3.4 Build Unit tests
##### Build the gtest libraries
tar -xf caffe_ilsvrc12.tar.gz && rm -f caffe_ilsvrc12.tar.gz
-echo "Done."
+echo "Done. "
\ No newline at end of file
--- /dev/null
+#include <caffe/caffe.hpp>
+#ifdef USE_OPENCV
+#include <opencv2/core/core.hpp>
+#include <opencv2/highgui/highgui.hpp>
+#include <opencv2/imgproc/imgproc.hpp>
+#endif // USE_OPENCV
+#include <algorithm>
+#include <iosfwd>
+#include <memory>
+#include <string>
+#include <utility>
+#include <vector>
+
+#ifdef USE_PROFILING
+
+#include <iostream>
+
+#include <time.h>
+
+#define REPEAT_TEST
+
+unsigned long get_cur_time(void)
+{
+ struct timespec tm;
+
+ clock_gettime(CLOCK_MONOTONIC_COARSE, &tm);
+
+ return (tm.tv_sec*1000+tm.tv_nsec/1000000);
+}
+
+#endif //USE_PROFILING
+
+#ifdef USE_OPENCV
+using namespace caffe; // NOLINT(build/namespaces)
+using std::string;
+
+/* Pair (label, confidence) representing a prediction. */
+typedef std::pair<string, float> Prediction;
+
+class Classifier {
+ public:
+ Classifier(const string& model_file,
+ const string& trained_file,
+ const string& mean_file,
+ const string& label_file);
+
+ std::vector<Prediction> Classify(const cv::Mat& img, int N = 5);
+
+#ifdef USE_PROFILING
+
+#ifdef LAYER_PERF_STAT
+ void dump_perf_stat(void);
+ void dump_single_layer_io(int idx, Layer<float> * p_layer);
+ void dump_single_layer_perf(int idx, Layer<float> * p_layer,uint64_t total_net_time);
+#ifdef REPEAT_TEST
+ void collect_layer_stat(vector<vector<perf_stat> * > & all_stat);
+ void dump_all_stat(vector <vector<perf_stat>*>& all_stat);
+ void reset_layer_stat();
+#endif
+#endif
+
+#endif //USE_PROFILING
+
+ private:
+ void SetMean(const string& mean_file);
+
+ std::vector<float> Predict(const cv::Mat& img);
+
+ void WrapInputLayer(std::vector<cv::Mat>* input_channels);
+
+ void Preprocess(const cv::Mat& img,
+ std::vector<cv::Mat>* input_channels);
+
+ private:
+ shared_ptr<Net<float> > net_;
+ cv::Size input_geometry_;
+ int num_channels_;
+ cv::Mat mean_;
+ std::vector<string> labels_;
+};
+
+Classifier::Classifier(const string& model_file,
+ const string& trained_file,
+ const string& mean_file,
+ const string& label_file) {
+#ifdef CPU_ONLY
+ Caffe::set_mode(Caffe::CPU);
+#else
+ Caffe::set_mode(Caffe::GPU);
+#endif
+
+ AclEnableSchedule();
+ /* Load the network. */
+ net_.reset(new Net<float>(model_file, TEST));
+ net_->CopyTrainedLayersFrom(trained_file);
+
+ CHECK_EQ(net_->num_inputs(), 1) << "Network should have exactly one input.";
+ CHECK_EQ(net_->num_outputs(), 1) << "Network should have exactly one output.";
+
+ Blob<float>* input_layer = net_->input_blobs()[0];
+ num_channels_ = input_layer->channels();
+ CHECK(num_channels_ == 3 || num_channels_ == 1)
+ << "Input layer should have 1 or 3 channels.";
+ input_geometry_ = cv::Size(input_layer->width(), input_layer->height());
+
+ /* Load the binaryproto mean file. */
+ SetMean(mean_file);
+
+ /* Load labels. */
+ std::ifstream labels(label_file.c_str());
+ CHECK(labels) << "Unable to open labels file " << label_file;
+ string line;
+ while (std::getline(labels, line))
+ labels_.push_back(string(line));
+
+ Blob<float>* output_layer = net_->output_blobs()[0];
+ CHECK_EQ(labels_.size(), output_layer->channels())
+ << "Number of labels is different from the output layer dimension.";
+}
+
+static bool PairCompare(const std::pair<float, int>& lhs,
+ const std::pair<float, int>& rhs) {
+ return lhs.first > rhs.first;
+}
+
+/* Return the indices of the top N values of vector v. */
+static std::vector<int> Argmax(const std::vector<float>& v, int N) {
+ std::vector<std::pair<float, int> > pairs;
+ for (size_t i = 0; i < v.size(); ++i)
+ pairs.push_back(std::make_pair(v[i], i));
+ std::partial_sort(pairs.begin(), pairs.begin() + N, pairs.end(), PairCompare);
+
+ std::vector<int> result;
+ for (int i = 0; i < N; ++i)
+ result.push_back(pairs[i].second);
+ return result;
+}
+
+/* Return the top N predictions. */
+std::vector<Prediction> Classifier::Classify(const cv::Mat& img, int N) {
+ std::vector<float> output = Predict(img);
+
+ N = std::min<int>(labels_.size(), N);
+ std::vector<int> maxN = Argmax(output, N);
+ std::vector<Prediction> predictions;
+ for (int i = 0; i < N; ++i) {
+ int idx = maxN[i];
+ predictions.push_back(std::make_pair(labels_[idx], output[idx]));
+ }
+
+ return predictions;
+}
+
+/* Load the mean file in binaryproto format. */
+void Classifier::SetMean(const string& mean_file) {
+ BlobProto blob_proto;
+ ReadProtoFromBinaryFileOrDie(mean_file.c_str(), &blob_proto);
+
+ /* Convert from BlobProto to Blob<float> */
+ Blob<float> mean_blob;
+ mean_blob.FromProto(blob_proto);
+ CHECK_EQ(mean_blob.channels(), num_channels_)
+ << "Number of channels of mean file doesn't match input layer.";
+
+ /* The format of the mean file is planar 32-bit float BGR or grayscale. */
+ std::vector<cv::Mat> channels;
+ float* data = mean_blob.mutable_cpu_data();
+ for (int i = 0; i < num_channels_; ++i) {
+ /* Extract an individual channel. */
+ cv::Mat channel(mean_blob.height(), mean_blob.width(), CV_32FC1, data);
+ channels.push_back(channel);
+ data += mean_blob.height() * mean_blob.width();
+ }
+
+ /* Merge the separate channels into a single image. */
+ cv::Mat mean;
+ cv::merge(channels, mean);
+
+ /* Compute the global mean pixel value and create a mean image
+ * filled with this value. */
+ cv::Scalar channel_mean = cv::mean(mean);
+ mean_ = cv::Mat(input_geometry_, mean.type(), channel_mean);
+}
+
+std::vector<float> Classifier::Predict(const cv::Mat& img) {
+ Blob<float>* input_layer = net_->input_blobs()[0];
+ input_layer->Reshape(1, num_channels_,
+ input_geometry_.height, input_geometry_.width);
+ /* Forward dimension change to all layers. */
+ net_->Reshape();
+
+ std::vector<cv::Mat> input_channels;
+ WrapInputLayer(&input_channels);
+
+ Preprocess(img, &input_channels);
+
+#ifdef USE_PROFILING
+ unsigned long tstart=get_cur_time();
+#endif //USE_PROFILING
+
+ net_->Forward();
+
+#ifdef USE_PROFILING
+
+ unsigned long tend=get_cur_time();
+
+ std::cout<<"used time: "<<tend-tstart<<std::endl;
+
+#ifdef LAYER_PERF_STAT
+ dump_perf_stat();
+#ifdef REPEAT_TEST
+
+ reset_layer_stat();
+
+ vector<vector<perf_stat>* > all_stat;
+ int rep_number=10;
+
+ for(int i=0;i<rep_number;i++)
+ {
+ net_->Forward();
+ collect_layer_stat(all_stat);
+ reset_layer_stat();
+ }
+
+ //dump stats
+ dump_all_stat(all_stat);
+
+ for(int i=0;i<all_stat.size();i++)
+ delete all_stat[i];
+
+#endif //REPEAT_TEST
+#endif //LAYER_PERF_STAT
+#endif //USE_PROFILING
+
+ /* Copy the output layer to a std::vector */
+ Blob<float>* output_layer = net_->output_blobs()[0];
+ const float* begin = output_layer->cpu_data();
+ const float* end = begin + output_layer->channels();
+ return std::vector<float>(begin, end);
+}
+
+#ifdef USE_PROFILING
+
+#ifdef LAYER_PERF_STAT
+
+#ifdef REPEAT_TEST
+void Classifier::collect_layer_stat(vector<vector<perf_stat>*>& all_stat)
+{
+ vector<perf_stat > * p_stat;
+ perf_stat * p_time_stat;
+ const vector<shared_ptr<Layer<float> > >& layers=net_->layers();
+
+
+ p_stat=new vector<perf_stat>;
+
+ for (int i =0;i< layers.size(); i++) {
+ p_time_stat=layers[i]->get_time_stat();
+ p_stat->push_back(*p_time_stat);
+
+ }
+
+ all_stat.push_back(p_stat);
+}
+
+void Classifier::reset_layer_stat(void)
+{
+ const vector<shared_ptr<Layer<float> > >& layers=net_->layers();
+ perf_stat * p_time_stat;
+
+ for (int i =0;i< layers.size(); i++) {
+ p_time_stat=layers[i]->get_time_stat();
+
+ p_time_stat->count=0;
+ p_time_stat->total=0;
+ p_time_stat->used=p_time_stat->start=p_time_stat->end=0;
+ }
+}
+
+void Classifier::dump_all_stat(vector<vector<perf_stat>*>& all_stat)
+{
+
+ struct new_perf_stat {
+ perf_stat stat;
+ int idx;
+ };
+
+ vector<new_perf_stat > layer_stat;
+ perf_stat * p_stat;
+
+ uint64_t total_time=0;
+
+ layer_stat.resize(all_stat[0]->size());
+
+ for(int i=0;i<all_stat.size();i++)
+ {
+ for(int j=0;j<layer_stat.size();j++)
+ {
+ p_stat=&layer_stat[j].stat;
+
+ p_stat->total+=(*all_stat[i])[j].total;
+ p_stat->count+=(*all_stat[i])[j].count;
+ total_time+=(*all_stat[i])[j].total;
+ }
+ }
+
+ total_time=total_time/all_stat.size();
+
+ std::cout<<std::endl<<"----------------------------------"<<std::endl;
+ std::cout<<"STATS for "<<all_stat.size()<<" reptitions: ..."<<std::endl;
+ std::cout<<"Total time: "<<total_time<<" per forward"<<std::endl;
+ std::cout<<"Each layer stats: ..."<<std::endl;
+
+
+ for(int i=layer_stat.size()-1;i>=0;i--)
+ {
+ p_stat=&layer_stat[i].stat;
+
+ layer_stat[i].idx=i;
+
+ std::cout<<" "<<i<<": used time: "<<p_stat->total/all_stat.size();
+ std::cout<<" ratio: "<<((float)p_stat->total)/all_stat.size()/total_time*100;
+ std::cout<<" enter count: "<<p_stat->count/all_stat.size()<<std::endl;
+ }
+
+ std::cout<<std::endl;
+
+ std::cout<<"time cost top 10 layers are: ..."<<std::endl;
+
+ std::sort(layer_stat.begin(),layer_stat.end(),[](const new_perf_stat& a, const new_perf_stat& b)
+ {
+ if(a.stat.total>b.stat.total)
+ return true;
+ else
+ return false;
+ });
+
+ uint64_t top_total_time=0;
+
+ for(int i=0; i<10; i++)
+ {
+ p_stat=&layer_stat[i].stat;
+
+ std::cout<<" "<<layer_stat[i].idx<<": used time: "<<p_stat->total/all_stat.size();
+ std::cout<<" ratio: "<<((float)p_stat->total)/all_stat.size()/total_time*100;
+ std::cout<<" enter count: "<<p_stat->count/all_stat.size()<<std::endl;
+ top_total_time+=p_stat->total;
+ }
+
+ std::cout<<"Top cost layers occupied: "<<(float)top_total_time/all_stat.size()/total_time*100<<std::endl;
+
+ std::cout<<std::endl;
+}
+
+#endif
+
+void Classifier::dump_single_layer_io(int idx, Layer<float> * p_layer)
+{
+ const LayerParameter& layer_param=p_layer->layer_param();
+
+ std::cout<<std::endl<<"LAYER IDX: "<<idx<<" name: "<<layer_param.name();
+ std::cout<<" type: "<<layer_param.type()<<std::endl;
+
+ const vector<Blob<float>*> *p_bottom_vec=p_layer->saved_bottom;
+
+ for(int i=0;i<layer_param.bottom_size(); i++)
+ {
+ std::cout<<"bottom "<<layer_param.bottom(i)<<": ";
+
+ Blob<float> * p_blob=(*p_bottom_vec)[i];
+
+ for(int j=0;j<p_blob->num_axes();j++)
+ {
+ std::cout<<p_blob->shape(j)<<" ";
+ }
+ std::cout<<std::endl;
+ }
+
+ const vector<Blob<float>*> *p_top_vec=p_layer->saved_top;
+ for(int i=0;i<layer_param.top_size(); i++)
+ {
+ std::cout<<"top "<<layer_param.top(i)<<": ";
+ Blob<float> * p_blob=(*p_top_vec)[i];
+
+ for(int j=0;j<p_blob->num_axes();j++)
+ {
+ std::cout<<p_blob->shape(j)<<" ";
+ }
+ std::cout<<std::endl;
+ }
+}
+
+void Classifier::dump_single_layer_perf(int idx, Layer<float> * p_layer, uint64_t total_net_time)
+{
+ const LayerParameter& layer_param=p_layer->layer_param();
+ perf_stat * p_time_stat;
+
+ p_time_stat=p_layer->get_time_stat();
+
+ std::cout<<std::endl<<"LAYER IDX: "<<idx<<" name: "<<layer_param.name();
+ std::cout<<" type: "<<layer_param.type();
+ std::cout<<" ratio: "<<(float)p_time_stat->total/total_net_time*100<<std::endl;
+
+
+ std::cout<<"time stat: total: "<<p_time_stat->total<<" count: "<<p_time_stat->count;
+ if(p_time_stat->count)
+ {
+ std::cout<<" average: "<<((float)p_time_stat->total)/p_time_stat->count;
+ }
+
+ std::cout<<" start: "<<p_time_stat->start<<" end: "<<p_time_stat->end;
+ std::cout<<std::endl;
+
+
+}
+
+void Classifier::dump_perf_stat(void)
+{
+ uint64_t total_net_time=0;
+
+ const vector<shared_ptr<Layer<float> > >& layers=net_->layers();
+
+ std::cout<<"Input/output shape for each layer ... total: "<<layers.size()<<std::endl;
+
+ for (int i = layers.size() - 1; i >= 0; --i) {
+ dump_single_layer_io(i,layers[i].get());
+ }
+
+
+ for (int i = layers.size() - 1; i >= 0; --i) {
+
+ perf_stat * p_time_stat;
+
+ p_time_stat=layers[i]->get_time_stat();
+
+ total_net_time+=p_time_stat->total;
+
+ }
+
+ std::cout<<"Time for each layer ... sum of all layers is : ";
+ std::cout<<total_net_time<<std::endl;
+
+ for (int i = layers.size() - 1; i >= 0; --i) {
+
+ dump_single_layer_perf(i,layers[i].get(),total_net_time);
+ }
+
+}
+
+#endif
+
+#endif //USE_PROFILING
+
+/* Wrap the input layer of the network in separate cv::Mat objects
+ * (one per channel). This way we save one memcpy operation and we
+ * don't need to rely on cudaMemcpy2D. The last preprocessing
+ * operation will write the separate channels directly to the input
+ * layer. */
+void Classifier::WrapInputLayer(std::vector<cv::Mat>* input_channels) {
+ Blob<float>* input_layer = net_->input_blobs()[0];
+
+ int width = input_layer->width();
+ int height = input_layer->height();
+ float* input_data = input_layer->mutable_cpu_data();
+ for (int i = 0; i < input_layer->channels(); ++i) {
+ cv::Mat channel(height, width, CV_32FC1, input_data);
+ input_channels->push_back(channel);
+ input_data += width * height;
+ }
+}
+
+void Classifier::Preprocess(const cv::Mat& img,
+ std::vector<cv::Mat>* input_channels) {
+ /* Convert the input image to the input image format of the network. */
+ cv::Mat sample;
+ if (img.channels() == 3 && num_channels_ == 1)
+ cv::cvtColor(img, sample, cv::COLOR_BGR2GRAY);
+ else if (img.channels() == 4 && num_channels_ == 1)
+ cv::cvtColor(img, sample, cv::COLOR_BGRA2GRAY);
+ else if (img.channels() == 4 && num_channels_ == 3)
+ cv::cvtColor(img, sample, cv::COLOR_BGRA2BGR);
+ else if (img.channels() == 1 && num_channels_ == 3)
+ cv::cvtColor(img, sample, cv::COLOR_GRAY2BGR);
+ else
+ sample = img;
+
+ cv::Mat sample_resized;
+ if (sample.size() != input_geometry_)
+ cv::resize(sample, sample_resized, input_geometry_);
+ else
+ sample_resized = sample;
+
+ cv::Mat sample_float;
+ if (num_channels_ == 3)
+ sample_resized.convertTo(sample_float, CV_32FC3);
+ else
+ sample_resized.convertTo(sample_float, CV_32FC1);
+
+ cv::Mat sample_normalized;
+ cv::subtract(sample_float, mean_, sample_normalized);
+
+ /* This operation will write the separate BGR planes directly to the
+ * input layer of the network because it is wrapped by the cv::Mat
+ * objects in input_channels. */
+ cv::split(sample_normalized, *input_channels);
+
+ CHECK(reinterpret_cast<float*>(input_channels->at(0).data)
+ == net_->input_blobs()[0]->cpu_data())
+ << "Input channels are not wrapping the input layer of the network.";
+}
+
+int main(int argc, char** argv) {
+ if (argc != 6) {
+ std::cerr << "Usage: " << argv[0]
+ << " deploy.prototxt network.caffemodel"
+ << " mean.binaryproto labels.txt img.jpg" << std::endl;
+ return 1;
+ }
+
+ ::google::InitGoogleLogging(argv[0]);
+
+ string model_file = argv[1];
+ string trained_file = argv[2];
+ string mean_file = argv[3];
+ string label_file = argv[4];
+ Classifier classifier(model_file, trained_file, mean_file, label_file);
+
+ string file = argv[5];
+
+ std::cout << "---------- Prediction for "
+ << file << " ----------" << std::endl;
+
+ cv::Mat img = cv::imread(file, -1);
+ CHECK(!img.empty()) << "Unable to decode image " << file;
+ std::vector<Prediction> predictions = classifier.Classify(img);
+
+ /* Print the top N predictions. */
+ for (size_t i = 0; i < predictions.size(); ++i) {
+ Prediction p = predictions[i];
+ std::cout << std::fixed << std::setprecision(4) << p.second << " - \""
+ << p.first << "\"" << std::endl;
+ }
+}
+#else
+int main(int argc, char** argv) {
+ LOG(FATAL) << "This example requires OpenCV; compile with USE_OPENCV.";
+}
+#endif // USE_OPENCV
+++ /dev/null
-#ifndef CAFFE_ACL_LAYER_HPP_
-#define CAFFE_ACL_LAYER_HPP_
-
-#ifdef USE_ACL
-#include "arm_compute/runtime/NEON/functions/NEConvolutionLayer.h"
-#include "arm_compute/runtime/NEON/functions/NEDirectConvolutionLayer.h"
-#include "arm_compute/runtime/CL/functions/CLConvolutionLayer.h"
-#include "arm_compute/runtime/NEON/functions/NEActivationLayer.h"
-#include "arm_compute/runtime/CL/functions/CLActivationLayer.h"
-#include "arm_compute/runtime/NEON/functions/NENormalizationLayer.h"
-#include "arm_compute/runtime/CL/functions/CLNormalizationLayer.h"
-#include "arm_compute/runtime/NEON/functions/NEPoolingLayer.h"
-#include "arm_compute/runtime/CL/functions/CLPoolingLayer.h"
-#include "arm_compute/runtime/NEON/functions/NESoftmaxLayer.h"
-#include "arm_compute/runtime/CL/functions/CLSoftmaxLayer.h"
-#include "arm_compute/runtime/NEON/functions/NEFullyConnectedLayer.h"
-#include "arm_compute/runtime/CL/functions/CLFullyConnectedLayer.h"
-#include "arm_compute/runtime/NEON/functions/NELocallyConnectedLayer.h"
-#include "arm_compute/runtime/CL/functions/CLLocallyConnectedLayer.h"
-#include "arm_compute/runtime/NEON/functions/NEBatchNormalizationLayer.h"
-#include "arm_compute/runtime/CL/functions/CLBatchNormalizationLayer.h"
-#include "arm_compute/core/NEON/kernels/NEDepthConcatenateKernel.h"
-#include "arm_compute/runtime/NEON/functions/NEDepthConcatenate.h"
-#include "arm_compute/core/CL/kernels/CLDepthConcatenateKernel.h"
-#include "arm_compute/runtime/CL/functions/CLDepthConcatenate.h"
-#include "arm_compute/runtime/CL/CLTensor.h"
-#include "arm_compute/runtime/Tensor.h"
-#include "arm_compute/runtime/CL/CLScheduler.h"
-using namespace arm_compute;
-#define FLAGS_ENABLE_ACL_ABSVAL 0x00000001
-#define FLAGS_ENABLE_ACL_BNLL 0x00000002
-#define FLAGS_ENABLE_ACL_CONV 0x00000004
-#define FLAGS_ENABLE_ACL_FC 0x00000008
-#define FLAGS_ENABLE_ACL_LRN 0x00000010
-#define FLAGS_ENABLE_ACL_POOLING 0x00000020
-#define FLAGS_ENABLE_ACL_RELU 0x00000040
-#define FLAGS_ENABLE_ACL_SIGMOID 0x00000080
-#define FLAGS_ENABLE_ACL_SOFTMAX 0x00000100
-#define FLAGS_ENABLE_ACL_TANH 0x00000200
-#define FLAGS_ENABLE_ACL_LC 0x00000400
-#define FLAGS_ENABLE_ACL_BN 0x00000800
-#define FLAGS_ENABLE_ACL_CONCAT 0x00001000
-extern unsigned int bypass_acl_class_layer;
-#endif
-#ifdef USE_PROFILING
-#include "layer.hpp"
-
-#define MASK_LOG_APP_TIME 0x00000001
-#define MASK_LOG_ALLOCATE 0x00000002
-#define MASK_LOG_RUN 0x00000004
-#define MASK_LOG_CONFIG 0x00000008
-#define MASK_LOG_COPY 0x00000010
-#define MASK_LOG_ABSVAL 0x00000020
-#define MASK_LOG_BNLL 0x00000040
-#define MASK_LOG_CONV 0x00000080
-#define MASK_LOG_FC 0x00000100
-#define MASK_LOG_LRN 0x00000200
-#define MASK_LOG_POOLING 0x00000400
-#define MASK_LOG_RELU 0x00000800
-#define MASK_LOG_SIGMOID 0x00001000
-#define MASK_LOG_SOFTMAX 0x00002000
-#define MASK_LOG_TANH 0x00004000
-#define MASK_LOG_LC 0x00008000
-#define MASK_LOG_BN 0x00010000
-#define MASK_LOG_CONCAT 0x00020000
-#define APP_TIME_INFO MASK_LOG_APP_TIME,"time: \t"
-#define ACL_ALLOCATE_INFO MASK_LOG_ALLOCATE,"allocate: \t\t"
-#define ACL_RUN_INFO MASK_LOG_RUN, "run: \t\t\t"
-#define ACL_CONFIG_INFO MASK_LOG_CONFIG, "configure: \t\t\t\t"
-#define ACL_COPY_INFO MASK_LOG_COPY, "tensor_copy:\t\t\t\t\t"
-#define ACL_ABSVAL_INFO MASK_LOG_ABSVAL, "ACL_ABSVAL :\t\t\t\t\t\t"
-#define ACL_BNLL_INFO MASK_LOG_BNLL, "ACL_BNLL :\t\t\t\t\t\t\t"
-#define ACL_CONV_INFO MASK_LOG_CONV, "ACL_CONV :\t\t\t\t\t\t\t\t"
-#define ACL_FC_INFO MASK_LOG_FC, "ACL_FC :\t\t\t\t\t\t\t\t\t"
-#define ACL_LRN_INFO MASK_LOG_LRN, "ACL_LRN :\t\t\t\t\t\t\t\t\t\t"
-#define ACL_POOLING_INFO MASK_LOG_POOLING, "ACL_POOLING:\t\t\t\t\t\t\t\t\t\t\t"
-#define ACL_RELU_INFO MASK_LOG_RELU, "ACL_RELU :\t\t\t\t\t\t\t\t\t\t\t\t"
-#define ACL_SIGMOID_INFO MASK_LOG_SIGMOID, "ACL_SIGMOID:\t\t\t\t\t\t\t\t\t\t\t\t\t"
-#define ACL_SOFTMAX_INFO MASK_LOG_SOFTMAX, "ACL_SOFTMAX:\t\t\t\t\t\t\t\t\t\t\t\t\t\t"
-#define ACL_TANH_INFO MASK_LOG_TANH, "ACL_TANH :\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t"
-#define ACL_LC_INFO MASK_LOG_LC, "ACL_LC :\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t"
-#define ACL_BN_INFO MASK_LOG_BN, "ACL_BN :\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t"
-#define ACL_CONCAT_INFO MASK_LOG_CONCAT, "ACL_CONCAT :\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t"
-extern unsigned int acl_log_flags;
-#endif //USE_PROFILING
-namespace caffe {
-#ifdef USE_ACL
-enum TensorType{
- tensor_input,
- tensor_output,
- tensor_weights,
- tensor_biases,
- tensor_data,
-};
-template <typename ACLTensor>
-class BaseTensor:public ACLTensor{
-public:
- BaseTensor(bool share)
- :share_(share),type_(tensor_input),allocate_(false){
- }
- virtual void bindmem(void *mem,bool share){
- mem_=mem;
- share_=share;
- }
- virtual void settensortype(TensorType type){
- type_=type;
- };
- virtual void map(bool blocking = true){}
- virtual void unmap(){}
- virtual void commit(TensorType type=tensor_data);
- int tensor_copy(void * mem, bool toTensor=true);
-protected:
- void* mem_;
- bool share_;
- TensorType type_;
- bool allocate_;
-};
-class GPUTensor:public BaseTensor<CLTensor>{
-public:
- explicit GPUTensor(bool share)
- :BaseTensor(share){}
- virtual void map(bool blocking = true){
- if (!allocate_){
- CLTensor::allocator()->allocate();
- allocate_=true;
- }
- CLTensor::map(blocking);
- }
- virtual void unmap(){
- CLTensor::unmap();
- }
-};
-class CPUTensor:public BaseTensor<Tensor>{
-public:
- explicit CPUTensor(bool share)
- :BaseTensor(share){}
- virtual void map(bool blocking = true){
- if (!allocate_){
- Tensor::allocator()->allocate();
- allocate_=true;
- }
- }
- virtual void unmap(){
- }
-};
-template <typename ACLLayer,typename ACLTensor>
-class ACLXPUBaseLayer{
-public:
- virtual void commit(){
- if (input) {
- input->commit(tensor_input);
- }
- if (output){
- output->commit(tensor_output);
- }
- if (weights){
- weights->commit(tensor_weights);
- }
- if (biases){
- biases->commit(tensor_biases);
- }
- }
- virtual void run(bool gpu){
- commit();
-#ifdef USE_PROFILING
- logtime_util log_time(ACL_RUN_INFO);
-#endif //USE_PROFILING
- layer->run();
- if (gpu) {
- // Make sure all the OpenCL jobs are done executing:
- CLScheduler::get().sync();
- }
- }
- virtual bool reshape(TensorShape &shape,TensorType type);
- explicit ACLXPUBaseLayer(){
- layer=nullptr;
- input=nullptr;
- output=nullptr;
- weights=nullptr;
- biases=nullptr;
- mean=nullptr;
- var=nullptr;
- beta=nullptr;
- gamma=nullptr;
-#ifdef USE_CONV_CACHE
- for(int i = 0; i < 16; ++i){
- cache.layer[i] = nullptr;
- cache.input[i] = nullptr;
- cache.output[i] = nullptr;
- cache.weights[i] = nullptr;
- cache.biases[i] = nullptr;
- }
-#endif //USE_CONV_CACHE
- }
- virtual void freelayer(){
-#ifndef USE_CONV_CACHE
- if (layer) delete layer;
- if (input) delete input;
- if (output) delete output;
- if (weights) delete weights;
- if (biases) delete biases;
- if (mean) delete mean;
- if (var) delete var;
- if (beta) delete beta;
- if (gamma) delete gamma;
-#endif //USE_CONV_CACHE
- layer=nullptr;
- input=nullptr;
- output=nullptr;
- weights=nullptr;
- biases=nullptr;
- mean=nullptr;
- var=nullptr;
- beta=nullptr;
- gamma=nullptr;
- }
- virtual ~ACLXPUBaseLayer(){
- freelayer();
- }
- ACLLayer *layer;
- ACLTensor *input;
- ACLTensor *output;
- ACLTensor *weights;
- ACLTensor *biases;
- //for BN
- ACLTensor *mean;
- ACLTensor *var;
- ACLTensor *beta;
- ACLTensor *gamma;
-#ifdef USE_CONV_CACHE
- struct{
- ACLLayer *layer[16];
- ACLTensor *input[16];
- ACLTensor *output[16];
- ACLTensor *weights[16];
- ACLTensor *biases[16];
- }cache;
-#endif //USE_CONV_CACHE
-};
-template <typename GPULayer, typename CPULayer>
-class ACLBaseLayer {
-public:
- explicit ACLBaseLayer();
- virtual void gpu_run();
- virtual void cpu_run();
- virtual ~ACLBaseLayer();
- virtual GPULayer * new_gpulayer();
- virtual CPULayer * new_cpulayer();
- ACLXPUBaseLayer<GPULayer,GPUTensor>& gpu(){
- return gpu_;
- }
- ACLXPUBaseLayer<CPULayer,CPUTensor>& cpu(){
- return cpu_;
- }
- bool checkreshape(TensorShape shape,bool gpu=false, TensorType type=tensor_input);
- template <typename ACLTensor> bool tensor_mem(ACLTensor *tensor,void *mem,bool share=false);
- template <typename ACLTensor> bool tensor_mem(void *mem,ACLTensor *tensor,bool share=false);
- template <typename ACLTensor> bool new_tensor(ACLTensor *&tensor,TensorShape shape,void *mem=nullptr,bool share=false);
-protected:
- ACLXPUBaseLayer<GPULayer,GPUTensor> gpu_;
- ACLXPUBaseLayer<CPULayer,CPUTensor> cpu_;
- bool init_layer_;
- bool force_bypass_acl_path_;
-
-};
-
-#endif
-}
-#define INSTANTIATE_ACLBASECLASS(GPULayer,CPULayer) \
- template class ACLBaseLayer<GPULayer,CPULayer>;
-
-#define INSTANTIATE_ACLBASE_FUNCTION(GPULayer,CPULayer,ACLTensor) \
- template bool ACLBaseLayer<GPULayer,CPULayer>::tensor_mem(ACLTensor *tensor,void *mem,bool share); \
- template bool ACLBaseLayer<GPULayer,CPULayer>::tensor_mem(void *mem,ACLTensor *tensor,bool share); \
- template bool ACLBaseLayer<GPULayer,CPULayer>::new_tensor(ACLTensor *&tensor,TensorShape shape,void *mem,bool share); \
-
-
-#endif
--- /dev/null
+#ifndef CAFFE_ACL_LAYER_HPP_
+#define CAFFE_ACL_LAYER_HPP_
+
+#ifdef USE_ACL
+#include "arm_compute/runtime/NEON/functions/NEDepthConcatenateLayer.h"
+#include "arm_compute/runtime/NEON/functions/NEConvolutionLayer.h"
+#include "arm_compute/runtime/NEON/functions/NEDirectConvolutionLayer.h"
+#include "arm_compute/runtime/NEON/functions/NEActivationLayer.h"
+#include "arm_compute/runtime/NEON/functions/NENormalizationLayer.h"
+#include "arm_compute/runtime/NEON/functions/NEPoolingLayer.h"
+#include "arm_compute/runtime/NEON/functions/NESoftmaxLayer.h"
+#include "arm_compute/runtime/NEON/functions/NEFullyConnectedLayer.h"
+#include "arm_compute/runtime/NEON/functions/NELocallyConnectedLayer.h"
+#include "arm_compute/runtime/NEON/functions/NEBatchNormalizationLayer.h"
+#include "arm_compute/runtime/Tensor.h"
+
+#include "arm_compute/core/CL/OpenCL.h"
+#include "arm_compute/runtime/CL/functions/CLDepthConcatenateLayer.h"
+#include "arm_compute/runtime/CL/functions/CLDirectConvolutionLayer.h"
+#include "arm_compute/runtime/CL/functions/CLConvolutionLayer.h"
+#include "arm_compute/runtime/CL/functions/CLActivationLayer.h"
+#include "arm_compute/runtime/CL/functions/CLNormalizationLayer.h"
+#include "arm_compute/runtime/CL/functions/CLPoolingLayer.h"
+#include "arm_compute/runtime/CL/functions/CLSoftmaxLayer.h"
+#include "arm_compute/runtime/CL/functions/CLFullyConnectedLayer.h"
+#include "arm_compute/runtime/CL/functions/CLLocallyConnectedLayer.h"
+#include "arm_compute/runtime/CL/functions/CLBatchNormalizationLayer.h"
+#include "arm_compute/runtime/CL/CLTensor.h"
+#include "arm_compute/runtime/CL/CLScheduler.h"
+#include "acl_tensor.hpp"
+#include "caffe/common.hpp"
+#include "caffe/layer.hpp"
+#define FLAGS_ENABLE_ACL_ABSVAL 0x00000001
+#define FLAGS_ENABLE_ACL_BNLL 0x00000002
+#define FLAGS_ENABLE_ACL_CONV 0x00000004
+#define FLAGS_ENABLE_ACL_FC 0x00000008
+#define FLAGS_ENABLE_ACL_LRN 0x00000010
+#define FLAGS_ENABLE_ACL_POOLING 0x00000020
+#define FLAGS_ENABLE_ACL_RELU 0x00000040
+#define FLAGS_ENABLE_ACL_SIGMOID 0x00000080
+#define FLAGS_ENABLE_ACL_SOFTMAX 0x00000100
+#define FLAGS_ENABLE_ACL_TANH 0x00000200
+#define FLAGS_ENABLE_ACL_LC 0x00000400
+#define FLAGS_ENABLE_ACL_BN 0x00000800
+#define FLAGS_ENABLE_ACL_CONCAT 0x00001000
+extern unsigned int bypass_acl_class_layer;
+extern unsigned int openailab_intfp;
+#endif
+#ifdef USE_PROFILING
+#include "layer.hpp"
+
+#define MASK_LOG_APP_TIME 0x00000001
+#define MASK_LOG_ALLOCATE 0x00000002
+#define MASK_LOG_RUN 0x00000004
+#define MASK_LOG_CONFIG 0x00000008
+#define MASK_LOG_COPY 0x00000010
+#define MASK_LOG_ABSVAL 0x00000020
+#define MASK_LOG_BNLL 0x00000040
+#define MASK_LOG_CONV 0x00000080
+#define MASK_LOG_FC 0x00000100
+#define MASK_LOG_LRN 0x00000200
+#define MASK_LOG_POOLING 0x00000400
+#define MASK_LOG_RELU 0x00000800
+#define MASK_LOG_SIGMOID 0x00001000
+#define MASK_LOG_SOFTMAX 0x00002000
+#define MASK_LOG_TANH 0x00004000
+#define MASK_LOG_LC 0x00008000
+#define MASK_LOG_BN 0x00010000
+#define MASK_LOG_CONCAT 0x00020000
+#define APP_TIME_INFO MASK_LOG_APP_TIME,"time: \t"
+#define ACL_ALLOCATE_INFO MASK_LOG_ALLOCATE,"allocate: \t\t"
+#define ACL_RUN_INFO MASK_LOG_RUN, "run: \t\t\t"
+#define ACL_CONFIG_INFO MASK_LOG_CONFIG, "configure: \t\t\t\t"
+#define ACL_COPY_INFO MASK_LOG_COPY, "tensor_copy:\t\t\t\t\t"
+#define ACL_ABSVAL_INFO MASK_LOG_ABSVAL, "ACL_ABSVAL :\t\t\t\t\t\t"
+#define ACL_BNLL_INFO MASK_LOG_BNLL, "ACL_BNLL :\t\t\t\t\t\t\t"
+#define ACL_CONV_INFO MASK_LOG_CONV, "ACL_CONV :\t\t\t\t\t\t\t\t"
+#define ACL_FC_INFO MASK_LOG_FC, "ACL_FC :\t\t\t\t\t\t\t\t\t"
+#define ACL_LRN_INFO MASK_LOG_LRN, "ACL_LRN :\t\t\t\t\t\t\t\t\t\t"
+#define ACL_POOLING_INFO MASK_LOG_POOLING, "ACL_POOLING:\t\t\t\t\t\t\t\t\t\t\t"
+#define ACL_RELU_INFO MASK_LOG_RELU, "ACL_RELU :\t\t\t\t\t\t\t\t\t\t\t\t"
+#define ACL_SIGMOID_INFO MASK_LOG_SIGMOID, "ACL_SIGMOID:\t\t\t\t\t\t\t\t\t\t\t\t\t"
+#define ACL_SOFTMAX_INFO MASK_LOG_SOFTMAX, "ACL_SOFTMAX:\t\t\t\t\t\t\t\t\t\t\t\t\t\t"
+#define ACL_TANH_INFO MASK_LOG_TANH, "ACL_TANH :\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t"
+#define ACL_LC_INFO MASK_LOG_LC, "ACL_LC :\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t"
+#define ACL_BN_INFO MASK_LOG_BN, "ACL_BN :\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t"
+#define ACL_CONCAT_INFO MASK_LOG_CONCAT, "ACL_CONCAT :\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t"
+extern unsigned int acl_log_flags;
+#endif //USE_PROFILING
+namespace caffe {
+#ifdef USE_ACL
+enum TensorType{
+ tensor_input,
+ tensor_output,
+ tensor_weights,
+ tensor_biases,
+ tensor_mean,
+ tensor_var,
+ tensor_beta,
+ tensor_gamma,
+ tensor_concat,
+ tensor_data,
+};
+enum OperatorState{
+ operator_not_init,
+ operator_init_done,
+ operator_reinit,
+};
+enum OperateType{
+ operate_type_pooling,
+ operate_type_activation,
+ operate_type_lrn,
+ operate_type_conv,
+ operate_type_lc,
+ operate_type_fc,
+ operate_type_bn,
+ operate_type_softmax,
+ operate_type_concat,
+};
+class BaseACLTensor{
+public:
+ BaseACLTensor()
+ :type_(tensor_input),allocate_(false){
+ }
+ virtual void bindmem(void *mem){
+ mem_=mem;
+ }
+ virtual void settensortype(TensorType type){
+ type_=type;
+ };
+ virtual void map(bool blocking = true){
+ }
+ virtual void unmap(){}
+ virtual void commit(TensorType type=tensor_data){}
+ int tensor_copy(arm_compute::ITensor* tensor,void * mem, bool toTensor=true);
+protected:
+ void* mem_;
+ TensorType type_;
+ bool allocate_;
+};
+class ACLTensor:public BaseACLTensor,public Tensor{
+public:
+ ACLTensor(arm_compute::TensorInfo &&info)
+ :Tensor(info){
+ }
+ virtual void map(bool blocking = true){
+ if (!allocate_){
+ Tensor::allocate();
+ allocate_=true;
+ }
+ Tensor::map(blocking);
+ }
+ virtual int tensor_copy(void * mem, bool toTensor=true){
+ auto acl_tensor=this;
+ arm_compute::ITensor* tensor=acl_tensor->tensor();
+ BaseACLTensor::tensor_copy(tensor,mem,toTensor);
+ return 0;
+ }
+ virtual void unmap(){Tensor::unmap();}
+ virtual void commit(TensorType type=tensor_data);
+};
+class ACLSubTensor:public BaseACLTensor,public SubTensor{
+public:
+ ACLSubTensor(std::unique_ptr<ACLTensor> &parent,arm_compute::TensorShape &shape,arm_compute::Coordinates& coord)
+ :SubTensor(parent.get(),shape,coord){
+ }
+ virtual int tensor_copy(void * mem, bool toTensor=true){
+ return 0;
+ }
+};
+
+template <typename T>
+class TensorPair{
+public:
+ TensorPair(){}
+ ~TensorPair(){}
+ TensorType type;
+ std::unique_ptr<T> tensor;
+};
+template <typename T>
+std::unique_ptr<T> &tensor_item(std::vector<std::unique_ptr<TensorPair<T>>>& pool,TensorType type,int idx){
+ int count=0;
+ for (auto &item: pool) {
+ if(item.get()->type==type){
+ ++count;
+ }
+ if(item.get()->type==type && idx==count-1){
+ return item.get()->tensor;
+ }
+ }
+ pool.push_back((std::unique_ptr<TensorPair<T>>)std::move(new TensorPair<T>));
+ auto item=pool[pool.size()-1].get();
+ item->type=type;
+ item->tensor=NULL;
+ return item->tensor;
+}
+class ACLOperator {
+public:
+ virtual void commit(){
+ for (auto & item: tensor_pool_) {
+ if(item.get()->tensor)item.get()->tensor->commit(item.get()->type);
+ }
+ }
+ inline void run(){
+ commit();
+ #ifdef USE_PROFILING
+ logtime_util log_time(ACL_RUN_INFO);
+ #endif //USE_PROFILING
+ for(auto &c : funcs_)
+ {
+ c->run();
+ }
+ }
+
+ inline std::vector<std::unique_ptr<arm_compute::IFunction>> &funcs(){return funcs_;}
+
+ inline std::unique_ptr<ACLSubTensor> &sinput(int idx=0){return tensor_item(subtensor_pool_,tensor_input,idx);}
+ inline std::unique_ptr<ACLSubTensor> &soutput(int idx=0){return tensor_item(subtensor_pool_,tensor_output,idx);}
+ inline std::unique_ptr<ACLSubTensor> &sweights(int idx=0){return tensor_item(subtensor_pool_,tensor_weights,idx);}
+ inline std::unique_ptr<ACLSubTensor> &sbiases(int idx=0){return tensor_item(subtensor_pool_,tensor_biases,idx);}
+
+ inline std::unique_ptr<ACLTensor> &cinput(int idx=0){return tensor_item(tensor_pool_,tensor_concat,idx);}
+ inline std::unique_ptr<ACLTensor> &input(int idx=0){return tensor_item(tensor_pool_,tensor_input,idx);}
+ inline std::unique_ptr<ACLTensor> &output(int idx=0){return tensor_item(tensor_pool_,tensor_output,idx);}
+ inline std::unique_ptr<ACLTensor> &weights(int idx=0){return tensor_item(tensor_pool_,tensor_weights,idx);}
+ inline std::unique_ptr<ACLTensor> &biases(int idx=0){return tensor_item(tensor_pool_,tensor_biases,idx);}
+ inline std::unique_ptr<ACLTensor> &mean(int idx=0){return tensor_item(tensor_pool_,tensor_mean,idx);}
+ inline std::unique_ptr<ACLTensor> &var(int idx=0){return tensor_item(tensor_pool_,tensor_var,idx);}
+ inline std::unique_ptr<ACLTensor> &beta(int idx=0){return tensor_item(tensor_pool_,tensor_beta,idx);}
+ inline std::unique_ptr<ACLTensor> &gamma(int idx=0){return tensor_item(tensor_pool_,tensor_gamma,idx);}
+ inline std::unique_ptr<ACLTensor> &tensor(TensorType type){
+ switch (type) {
+ case tensor_biases:
+ return biases();
+ break;
+ case tensor_weights:
+ return weights();
+ break;
+ case tensor_output:
+ return output();
+ break;
+ default:
+ case tensor_input:
+ return input();
+ break;
+ }
+ return input();
+ }
+
+
+ explicit ACLOperator(const LayerParameter& param);
+ virtual ~ACLOperator();
+ inline TargetHint getTargetHint(){
+#ifdef USE_OPENCL
+ if (target_hint_==TargetHint::DONT_CARE) {
+ if (Caffe::arm_gpu_mode()) {
+ return TargetHint::OPENCL;
+ }
+ return TargetHint::NEON;
+ }
+ return target_hint_;
+#else
+ return TargetHint::NEON;
+#endif
+ }
+ inline void setTargetHint(TargetHint hint){
+ target_hint_=hint;
+ }
+ inline ConvolutionMethodHint & getConvMethod(){ return convolution_method_hint_;}
+ inline bool tensor_mem(std::unique_ptr<ACLTensor> &tensor,void *mem){
+ tensor->bindmem(mem);
+ return true;
+ }
+ inline bool tensor_mem(void *mem,std::unique_ptr<ACLTensor> &tensor){
+ tensor->tensor_copy(mem,false);
+ return true;
+ }
+ bool new_tensor(std::unique_ptr<ACLTensor> &tensor,arm_compute::TensorShape &shape,void *mem=nullptr,bool commit=false);
+ bool new_tensor(std::unique_ptr<ACLSubTensor> &tensor,std::unique_ptr<ACLTensor> &parent,arm_compute::TensorShape &shape,arm_compute::Coordinates& coord);
+ inline int & group(){return _group;}
+ inline void set_operator_property(OperateType type,const char*name){
+ name_=name;
+ type_=type;
+ }
+ inline void acl_run(void *input_data, void *output_data){
+ if(input_data)tensor_mem(input(),input_data);
+ run();
+ tensor_mem(output_data,output());
+ }
+
+
+protected:
+ inline bool isGPUMode(){
+ if (!support_opencl_) return false;
+ return getTargetHint()==TargetHint::OPENCL;
+ }
+ inline OperatorState & opstate(){return operator_state_;}
+ inline bool is_operator_init_done(arm_compute::TensorShape shape,TensorType type=tensor_input){
+ checkreshape(shape,type);
+ return operator_state_==operator_init_done;
+ }
+ inline void set_operator_init_done(){
+ opstate()=operator_init_done;
+ set_bypass_state(false);
+ }
+ inline void set_bypass_state(bool state=false){
+ force_bypass_acl_path_=state;
+ }
+ inline OperatorState checkreshape(arm_compute::TensorShape shape,TensorType type=tensor_input){
+ opstate()=reshape(shape,type);
+ if (opstate()==operator_reinit) {
+ freeres();
+ }
+ return opstate();
+ }
+ inline OperatorState reshape(arm_compute::TensorShape &shape,TensorType type){
+ arm_compute::TensorShape _shape;
+ std::unique_ptr<ACLTensor> &acl_tensor=tensor(type);
+ if (!acl_tensor.get()) return operator_not_init;
+ _shape = acl_tensor->info().tensor_shape();
+ if (_shape.total_size()==shape.total_size() && _shape[0]==shape[0] && _shape[1]==shape[1]) {
+ return operator_init_done;
+ }
+ return operator_reinit;
+ }
+ inline void freeres(){
+ tensor_pool_.clear();
+ subtensor_pool_.clear();
+ funcs_.clear();
+ }
+ inline const char* &name(){return name_;}
+
+protected:
+ std::vector<std::unique_ptr<TensorPair<ACLTensor>>>tensor_pool_;
+ std::vector<std::unique_ptr<TensorPair<ACLSubTensor>>>subtensor_pool_;
+ std::vector<std::unique_ptr<arm_compute::IFunction>> funcs_;
+ OperatorState operator_state_;
+ bool force_bypass_acl_path_;
+ TargetHint target_hint_;
+ ConvolutionMethodHint convolution_method_hint_;
+ static bool support_opencl_;
+ static bool init_cl_env;
+ int _group;
+ const char* name_;
+ OperateType type_;
+};
+
+int isScheduleEnable();
+
+template <typename OperatorType, typename TensorType>
+std::unique_ptr<arm_compute::IFunction> instantiate_function(arm_compute::ITensor *input, arm_compute::ITensor *output){
+ auto op = cpp14::make_unique<OperatorType>();
+ op->configure(
+ dynamic_cast<TensorType *>(input),
+ dynamic_cast<TensorType *>(output)
+ );
+
+ return std::move(op);
+}
+
+template <typename OperatorType, typename TensorType>
+std::unique_ptr<arm_compute::IFunction> instantiate(arm_compute::ITensor *input, arm_compute::ITensor *output)
+{
+ return instantiate_function<OperatorType, TensorType>(input, output);
+}
+
+template <typename GPUOpType,typename GPUTensor,typename CPUOpType,typename CPUTensor>
+std::unique_ptr<arm_compute::IFunction> instantiate_op_func(std::unique_ptr<ACLTensor>& input, std::unique_ptr<ACLTensor>& output,TargetHint& hint){
+ std::unique_ptr<arm_compute::IFunction> func;
+#ifdef USE_OPENCL
+ if(hint == TargetHint::OPENCL)
+ {
+ func = instantiate<GPUOpType, GPUTensor>(input->tensor(), output->tensor());
+ }
+ else
+#endif
+ {
+ func = instantiate<CPUOpType, CPUTensor>(input->tensor(), output->tensor());
+ }
+ return func;
+}
+
+
+template <typename OperatorType, typename TensorType,typename VectorTensor>
+std::unique_ptr<arm_compute::IFunction> instantiate_function(VectorTensor inputs, arm_compute::ITensor *output){
+ auto op = cpp14::make_unique<OperatorType>();
+ op->configure(
+ inputs,
+ dynamic_cast<TensorType *>(output)
+ );
+
+ return std::move(op);
+}
+
+template <typename OperatorType, typename TensorType,typename VectorTensor>
+std::unique_ptr<arm_compute::IFunction> instantiate(VectorTensor inputs, arm_compute::ITensor *output)
+{
+ return instantiate_function<OperatorType, TensorType,VectorTensor>(inputs, output);
+}
+
+template <typename GPUOpType,typename GPUTensor,typename CPUOpType,typename CPUTensor>
+std::unique_ptr<arm_compute::IFunction> instantiate_op_func_lists(ACLOperator*& acl_op, std::unique_ptr<ACLTensor>& output,int num,TargetHint& hint){
+ std::unique_ptr<arm_compute::IFunction> func;
+#ifdef USE_OPENCL
+ if(hint == TargetHint::OPENCL)
+ {
+ static std::vector<arm_compute::ICLTensor*> tensors;
+ tensors.clear();
+ for (int i=0;i<num;++i) {
+ tensors.push_back(dynamic_cast<arm_compute::ICLTensor*>(acl_op->cinput(i).get()->tensor()));
+ }
+ func = instantiate<GPUOpType, GPUTensor, std::vector<arm_compute::ICLTensor *>>(tensors, output->tensor());
+ }
+ else
+#endif
+ {
+ static std::vector<arm_compute::ITensor*> tensors;
+ tensors.clear();
+ for (int i=0;i<num;++i) {
+ tensors.push_back(dynamic_cast<arm_compute::ITensor*>(acl_op->cinput(i).get()->tensor()));
+ }
+ func = instantiate<CPUOpType, CPUTensor,std::vector<arm_compute::ITensor*>>(tensors, output->tensor());
+ }
+ return func;
+}
+
+template <typename OperatorType, typename TensorType,typename OperatorInfo>
+std::unique_ptr<arm_compute::IFunction> instantiate_function(arm_compute::ITensor *input, arm_compute::ITensor *output, const OperatorInfo &info){
+ auto op = cpp14::make_unique<OperatorType>();
+ op->configure(
+ dynamic_cast<TensorType *>(input),
+ dynamic_cast<TensorType *>(output),
+ info);
+
+ return std::move(op);
+}
+
+template <typename OperatorType, typename TensorType,typename OperatorInfo>
+std::unique_ptr<arm_compute::IFunction> instantiate(arm_compute::ITensor *input, arm_compute::ITensor *output, const OperatorInfo &info)
+{
+ return instantiate_function<OperatorType, TensorType, OperatorInfo>(input, output, info);
+}
+
+template <typename GPUOpType,typename GPUTensor,typename CPUOpType,typename CPUTensor, typename OperatorInfo>
+std::unique_ptr<arm_compute::IFunction> instantiate_op_func(std::unique_ptr<ACLTensor>& input, std::unique_ptr<ACLTensor>& output, const OperatorInfo &info,TargetHint& hint){
+ std::unique_ptr<arm_compute::IFunction> func;
+#ifdef USE_OPENCL
+ if(hint == TargetHint::OPENCL)
+ {
+ func = instantiate<GPUOpType, GPUTensor,OperatorInfo>(input->tensor(), output->tensor(), info);
+ }
+ else
+#endif
+ {
+ func = instantiate<CPUOpType, CPUTensor,OperatorInfo>(input->tensor(), output->tensor(), info);
+ }
+ return func;
+}
+
+
+template <typename OperatorType, typename TensorType,typename OperatorInfo>
+std::unique_ptr<arm_compute::IFunction> instantiate_function(arm_compute::ITensor *input,arm_compute::ITensor *weights,arm_compute::ITensor *biases, arm_compute::ITensor *output, const OperatorInfo &info){
+ auto op = cpp14::make_unique<OperatorType>();
+ op->configure(
+ dynamic_cast<TensorType *>(input),
+ dynamic_cast<TensorType *>(weights),
+ dynamic_cast<TensorType *>(biases),
+ dynamic_cast<TensorType *>(output),
+ info);
+ return std::move(op);
+}
+
+template <typename OperatorType, typename TensorType,typename OperatorInfo>
+std::unique_ptr<arm_compute::IFunction> instantiate(arm_compute::ITensor *input,arm_compute::ITensor *weights,arm_compute::ITensor *biases, arm_compute::ITensor *output, const OperatorInfo &info)
+{
+ return instantiate_function<OperatorType, TensorType, OperatorInfo>(input,weights,biases,output, info);
+}
+
+template <typename GPUOpType,typename GPUTensor,typename CPUOpType,typename CPUTensor, typename OperatorInfo,typename ACLTensor>
+std::unique_ptr<arm_compute::IFunction> instantiate_op_func(std::unique_ptr<ACLTensor>& input,std::unique_ptr<ACLTensor>& weights,std::unique_ptr<ACLTensor>& biases, std::unique_ptr<ACLTensor>& output, const OperatorInfo &info,TargetHint& hint){
+ std::unique_ptr<arm_compute::IFunction> func;
+ arm_compute::ITensor * biases_tensor=NULL;
+
+ if (biases.get()) {
+ biases_tensor=biases->tensor();
+ }
+#ifdef USE_OPENCL
+ if (hint == TargetHint::OPENCL)
+ {
+ func = instantiate<GPUOpType, GPUTensor,OperatorInfo>(input->tensor(), weights->tensor(),biases_tensor,output->tensor(), info);
+ }
+ else
+#endif
+ {
+ func = instantiate<CPUOpType, CPUTensor,OperatorInfo>(input->tensor(), weights->tensor(),biases_tensor, output->tensor(), info);
+ }
+ return func;
+}
+
+
+
+template <typename Dtype,typename OperatorType, typename TensorType>
+std::unique_ptr<arm_compute::IFunction> instantiate_function(arm_compute::ITensor *input, arm_compute::ITensor *output,
+ arm_compute::ITensor *mean,arm_compute::ITensor *var,arm_compute::ITensor *beta,arm_compute::ITensor *gamma,Dtype & eps){
+ auto op = cpp14::make_unique<OperatorType>();
+ op->configure(
+ dynamic_cast<TensorType *>(input),
+ dynamic_cast<TensorType *>(output),
+ dynamic_cast<TensorType *>(mean),
+ dynamic_cast<TensorType *>(var),
+ dynamic_cast<TensorType *>(beta),
+ dynamic_cast<TensorType *>(gamma),
+ eps);
+
+ return std::move(op);
+}
+
+template <typename Dtype,typename OperatorType, typename TensorType>
+std::unique_ptr<arm_compute::IFunction> instantiate(arm_compute::ITensor * input,arm_compute::ITensor * output,
+ arm_compute::ITensor * mean,arm_compute::ITensor * var,arm_compute::ITensor * beta,arm_compute::ITensor * gamma,Dtype eps){
+ return instantiate_function<Dtype,OperatorType, TensorType>(input,output, mean,var,beta,gamma,eps);
+}
+
+template <typename Dtype,typename GPUOpType,typename GPUTensor,typename CPUOpType,typename CPUTensor>
+std::unique_ptr<arm_compute::IFunction> instantiate_op_func(std::unique_ptr<ACLTensor>& input,std::unique_ptr<ACLTensor>& output,
+ std::unique_ptr<ACLTensor>& mean,std::unique_ptr<ACLTensor>& var,std::unique_ptr<ACLTensor>& beta,std::unique_ptr<ACLTensor>& gamma,Dtype eps,TargetHint hint){
+ std::unique_ptr<arm_compute::IFunction> func;
+#ifdef USE_OPENCL
+ if(hint == TargetHint::OPENCL)
+ {
+ func = instantiate<Dtype,GPUOpType, GPUTensor>(input->tensor(),output->tensor(), mean->tensor(),var->tensor(),beta->tensor(),gamma->tensor(),eps);
+ }
+ else
+#endif
+ {
+ func = instantiate<Dtype,CPUOpType, CPUTensor>(input->tensor(),output->tensor(), mean->tensor(),var->tensor(),beta->tensor(),gamma->tensor(),eps);
+ }
+ return func;
+}
+
+
+template <typename OperatorInfo>
+bool instantiate_op_pooling(ACLOperator* acl_op,std::vector<std::unique_ptr<arm_compute::IFunction>> & func,std::unique_ptr<ACLTensor> & input, std::unique_ptr<ACLTensor> & output,TargetHint hint, const OperatorInfo &info){
+ func.push_back(instantiate_op_func<arm_compute::CLPoolingLayer, arm_compute::ICLTensor, arm_compute::NEPoolingLayer, arm_compute::ITensor, arm_compute::PoolingLayerInfo>(input, output, info, hint));
+ return true;
+}
+template <typename OperatorInfo>
+bool instantiate_op_activation(ACLOperator* acl_op,std::vector<std::unique_ptr<arm_compute::IFunction>> & func,std::unique_ptr<ACLTensor> & input,std::unique_ptr<ACLTensor> & output,TargetHint hint, const OperatorInfo &info){
+ func.push_back(instantiate_op_func<arm_compute::CLActivationLayer,arm_compute::ICLTensor,arm_compute::NEActivationLayer,arm_compute::ITensor, arm_compute::ActivationLayerInfo>(input, output, info, hint));
+ return true;
+}
+template <typename OperatorInfo>
+bool instantiate_op_lrn(ACLOperator* acl_op,std::vector<std::unique_ptr<arm_compute::IFunction>> & func,std::unique_ptr<ACLTensor> & input,std::unique_ptr<ACLTensor> & output,TargetHint hint, const OperatorInfo &info){
+ func.push_back(instantiate_op_func<arm_compute::CLNormalizationLayer,arm_compute::ICLTensor,arm_compute::NENormalizationLayer,arm_compute::ITensor, arm_compute::NormalizationLayerInfo>(input, output, info, hint));
+ return true;
+}
+template <typename OperatorInfo>
+bool instantiate_op_conv(ACLOperator* acl_op,std::vector<std::unique_ptr<arm_compute::IFunction>> & func,std::unique_ptr<ACLTensor> & input,std::unique_ptr<ACLTensor> & output,TargetHint hint,const OperatorInfo &info){
+ std::unique_ptr<ACLTensor> & weights=acl_op->weights();
+ std::unique_ptr<ACLTensor> & biases=acl_op->biases();
+ ConvolutionMethodHint& conv_method=acl_op->getConvMethod();
+ bool has_biases=biases.get()?true:false;
+ int& groups=acl_op->group();
+ arm_compute::TensorShape input_shape=input->info().tensor_shape();
+ arm_compute::TensorShape weights_shape=weights->info().tensor_shape();
+ arm_compute::TensorShape biases_shape;
+ if (has_biases) {
+ biases_shape = biases->info().tensor_shape();
+ }
+ arm_compute::TensorShape output_shape=output->info().tensor_shape();
+
+ if (groups==1) {
+ if (conv_method == ConvolutionMethodHint::GEMM) {
+ func.push_back(instantiate_op_func<arm_compute::CLConvolutionLayer, arm_compute::ICLTensor, arm_compute::NEConvolutionLayer, arm_compute::ITensor, arm_compute::PadStrideInfo>(acl_op->input(), acl_op->weights(), acl_op->biases(), acl_op->output(), info, hint));
+ }else{
+ func.push_back(instantiate_op_func<arm_compute::CLDirectConvolutionLayer, arm_compute::ICLTensor, arm_compute::NEDirectConvolutionLayer, arm_compute::ITensor, arm_compute::PadStrideInfo>(acl_op->input(), acl_op->weights(), acl_op->biases(), acl_op->output(), info, hint));
+ }
+ return true;
+ }
+
+ // Calculate sub-tensor splits
+ const int input_split = input_shape.z() / groups;
+ const int output_split = output_shape.z() / groups;
+ const int weights_split = weights_shape[3] / groups;
+ const int biases_split = biases_shape.x() / groups;
+
+ // Calculate sub-tensor shapes
+ input_shape.set(2, input_split);
+ output_shape.set(2, output_split);
+ weights_shape.set(3, weights_split);
+ biases_shape.set(0, biases_split);
+
+ for (auto i = 0; i < groups; ++i) {
+ // Calculate sub-tensors starting coordinates
+ arm_compute::Coordinates input_coord(0, 0, input_split * i);
+ arm_compute::Coordinates output_coord(0, 0, output_split * i);
+ arm_compute::Coordinates weights_coord(0, 0, 0, weights_split * i);
+ arm_compute::Coordinates biases_coord(biases_split * i);
+
+ // Create sub-tensors for input, output, weights and bias
+ acl_op->new_tensor(acl_op->sinput(i), acl_op->input(), input_shape, input_coord);
+ acl_op->new_tensor(acl_op->soutput(i),acl_op->output(),output_shape, output_coord);
+ acl_op->new_tensor(acl_op->sweights(i),acl_op->weights(), weights_shape, weights_coord);
+ if (has_biases) {
+ acl_op->new_tensor(acl_op->sbiases(i),acl_op->biases(), biases_shape, biases_coord);
+ }
+
+ if (conv_method == ConvolutionMethodHint::GEMM) {
+ func.push_back(instantiate_op_func<arm_compute::CLConvolutionLayer, arm_compute::ICLTensor, arm_compute::NEConvolutionLayer, arm_compute::ITensor, arm_compute::PadStrideInfo,ACLSubTensor>(acl_op->sinput(i), acl_op->sweights(i), acl_op->sbiases(i), acl_op->soutput(i), info, hint));
+ }else{
+ func.push_back(instantiate_op_func<arm_compute::CLDirectConvolutionLayer, arm_compute::ICLTensor, arm_compute::NEDirectConvolutionLayer, arm_compute::ITensor, arm_compute::PadStrideInfo,ACLSubTensor>(acl_op->sinput(i), acl_op->sweights(i), acl_op->sbiases(i), acl_op->soutput(i), info, hint));
+ }
+ }
+ return true;
+}
+template <typename OperatorInfo>
+bool instantiate_op_lc(ACLOperator* acl_op,std::vector<std::unique_ptr<arm_compute::IFunction>> & func,std::unique_ptr<ACLTensor> & input,std::unique_ptr<ACLTensor> & output,TargetHint hint, const OperatorInfo &info){
+ std::unique_ptr<ACLTensor> & weights=acl_op->weights();
+ std::unique_ptr<ACLTensor> & biases=acl_op->biases();
+ func.push_back(instantiate_op_func<arm_compute::CLLocallyConnectedLayer,arm_compute::ICLTensor,arm_compute::NELocallyConnectedLayer,arm_compute::ITensor, arm_compute::PadStrideInfo>(input, weights,biases,output,info, hint));
+ return true;
+}
+template <typename OperatorInfo>
+bool instantiate_op_fc(ACLOperator* acl_op,std::vector<std::unique_ptr<arm_compute::IFunction>> & func,std::unique_ptr<ACLTensor> & input,std::unique_ptr<ACLTensor> & output,TargetHint hint, const OperatorInfo &info){
+ std::unique_ptr<ACLTensor> & weights=acl_op->weights();
+ std::unique_ptr<ACLTensor> & biases=acl_op->biases();
+ func.push_back(instantiate_op_func<arm_compute::CLFullyConnectedLayer,arm_compute::ICLTensor,arm_compute::NEFullyConnectedLayer,arm_compute::ITensor, bool>(input, weights,biases,output,info, hint));
+ return true;
+}
+template <typename Dtype>
+bool instantiate_op_bn(ACLOperator* acl_op,std::vector<std::unique_ptr<arm_compute::IFunction>> & func,std::unique_ptr<ACLTensor> & input,std::unique_ptr<ACLTensor> & output,TargetHint hint, Dtype eps){
+ std::unique_ptr<ACLTensor> & mean=acl_op->mean();
+ std::unique_ptr<ACLTensor> & var=acl_op->var();
+ std::unique_ptr<ACLTensor> & beta=acl_op->beta();
+ std::unique_ptr<ACLTensor> & gamma=acl_op->gamma();
+ func.push_back(instantiate_op_func<Dtype,arm_compute::CLBatchNormalizationLayer,arm_compute::ICLTensor,arm_compute::NEBatchNormalizationLayer,arm_compute::ITensor>(input, output, mean,var,beta,gamma,eps, hint));
+ return true;
+}
+inline bool instantiate_op_softmax(ACLOperator* acl_op,std::vector<std::unique_ptr<arm_compute::IFunction>> & func,std::unique_ptr<ACLTensor> & input,std::unique_ptr<ACLTensor> & output,TargetHint hint,void *data){
+ func.push_back(instantiate_op_func<arm_compute::CLSoftmaxLayer,arm_compute::ICLTensor,arm_compute::NESoftmaxLayer,arm_compute::ITensor>(input, output, hint));
+ return true;
+}
+inline bool instantiate_op_concat(ACLOperator* acl_op,std::vector<std::unique_ptr<arm_compute::IFunction>> & func,std::unique_ptr<ACLTensor> & input,std::unique_ptr<ACLTensor> & output,TargetHint hint,int num){
+ func.push_back(instantiate_op_func_lists<arm_compute::CLDepthConcatenateLayer,arm_compute::ICLTensor,arm_compute::NEDepthConcatenateLayer,arm_compute::ITensor>(acl_op, output, num,hint));
+ return true;
+}
+template <typename Dtype>
+Dtype* GetDataPtr(ACLOperator* op,Blob<Dtype>* const &blob,bool isconst=false){
+ if (!isconst) {
+ if (op->getTargetHint() == TargetHint::NEON) {
+ return blob->mutable_cpu_data();
+ }
+ return blob->mutable_gpu_data();
+ }
+ if (op->getTargetHint()==TargetHint::NEON) {
+ return (Dtype*)blob->cpu_data();
+ }
+ return (Dtype*)blob->gpu_data();
+}
+
+template <typename Dtype>
+Dtype* InputdataPtr(ACLOperator* op,const vector<Blob<Dtype>*>& bottom,int index=-1){
+ if (index==-1) index=0;
+ return GetDataPtr(op, bottom[index], true);
+}
+template <typename Dtype>
+Dtype* OutputdataPtr(ACLOperator* op,const vector<Blob<Dtype>*>& top){
+ return GetDataPtr(op,top[0]);
+}
+
+template <typename Dtype>
+void acl_run(ACLOperator* op,const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top,bool multi_input_run=true){
+ if (multi_input_run) {
+ for (int i = 0; i < bottom.size(); ++i) {
+ const Dtype* bottom_data = bottom[i]->cpu_data();
+ Dtype* top_data = top[i]->mutable_cpu_data();
+ op->acl_run((void*)bottom_data,(void*)top_data);
+ }
+ return ;
+ }
+ for (int i = 0; i < bottom.size(); ++i) {
+ op->tensor_mem(op->cinput(i),InputdataPtr(op,bottom,i));
+ }
+ op->acl_run(NULL,OutputdataPtr(op,top));
+}
+}
+
+#define INIT_GLOBAL_FUNCS_TYPE(Dtype) \
+template <> \
+Dtype* InputdataPtr(ACLOperator* op,const vector<Blob<Dtype>*>& bottom,int index); \
+template <> \
+Dtype* OutputdataPtr(ACLOperator* op,const vector<Blob<Dtype>*>& top); \
+template <> \
+Dtype* GetDataPtr(ACLOperator* op,Blob<Dtype>* const & blob,bool isconst); \
+
+#define INIT_GLOBAL_FUNCS() \
+INIT_GLOBAL_FUNCS_TYPE(double); \
+INIT_GLOBAL_FUNCS_TYPE(float); \
+
+
+#ifdef USE_PROFILING
+#define acl_configure(opname,acl_op,args...)\
+{\
+ set_operator_property(operate_type_##opname,#opname); \
+ logtime_util log_time(ACL_CONFIG_INFO); \
+ instantiate_op_##opname(acl_op,acl_op->funcs(),acl_op->input(),acl_op->output(),acl_op->getTargetHint(),args);\
+}
+#else
+#define acl_configure(opname,acl_op,args...)\
+{\
+ set_operator_property(operate_type_##opname,#opname); \
+ instantiate_op_##opname(acl_op,acl_op->funcs(),acl_op->input(),acl_op->output(),acl_op->getTargetHint(),args);\
+}
+#endif
+
+#endif
+
+#endif
--- /dev/null
+#ifndef __TENSOR_H__
+#define __TENSOR_H__
+
+#ifdef USE_ACL
+#include "arm_compute/runtime/CL/CLSubTensor.h"
+#include "arm_compute/runtime/SubTensor.h"
+#include "arm_compute/runtime/CL/CLTensor.h"
+#include "arm_compute/runtime/Tensor.h"
+
+#include <memory>
+
+namespace caffe{
+enum class TargetHint{
+ DONT_CARE,
+ OPENCL,
+ NEON,
+};
+
+enum class ConvolutionMethodHint{
+ GEMM,
+ DIRECT,
+};
+namespace cpp14{
+template <class T>
+struct _Unique_if{
+ typedef std::unique_ptr<T> _Single_object;
+};
+
+template <class T>
+struct _Unique_if<T[]>{
+ typedef std::unique_ptr<T[]> _Unknown_bound;
+};
+
+template <class T, size_t N>
+struct _Unique_if<T[N]>{
+ typedef void _Known_bound;
+};
+
+template <class T, class... Args>
+typename _Unique_if<T>::_Single_object
+make_unique(Args &&... args){
+ return std::unique_ptr<T>(new T(std::forward<Args>(args)...));
+}
+
+template <class T>
+typename _Unique_if<T>::_Unknown_bound
+make_unique(size_t n){
+ typedef typename std::remove_extent<T>::type U;
+ return std::unique_ptr<T>(new U[n]());
+}
+
+template <class T, class... Args>
+typename _Unique_if<T>::_Known_bound
+make_unique(Args &&...) ;
+}
+
+class Tensor {
+public:
+ Tensor(arm_compute::TensorInfo &info) noexcept;
+ ~Tensor(){
+ }
+ Tensor(Tensor &&src) noexcept ;
+ void set_info(arm_compute::TensorInfo &&info){
+ _info = info;
+ }
+ arm_compute::ITensor *set_target(TargetHint target);
+ const arm_compute::TensorInfo &info() const{
+ return _info;
+ }
+ arm_compute::ITensor * tensor(){
+ return _tensor.get();
+ }
+ void allocate();
+ void init(){
+
+ }
+ TargetHint target() const{
+ return _target;
+ }
+ virtual void map(bool blocking = true);
+ virtual void unmap();
+
+private:
+ TargetHint _target;
+ arm_compute::TensorInfo _info;
+ std::unique_ptr<arm_compute::ITensor> _tensor;
+};
+
+class SubTensor
+{
+public:
+ SubTensor(Tensor* parent, arm_compute::TensorShape& tensor_shape, arm_compute::Coordinates& coords)noexcept;
+ ~SubTensor(){}
+ arm_compute::ITensor *tensor() ;
+ const arm_compute::ITensor *tensor() const ;
+ TargetHint target() const ;
+ void allocate() ;
+ arm_compute::ITensor *set_target(TargetHint target);
+
+private:
+ /** Instantiates a sub-tensor */
+ void instantiate_subtensor();
+
+private:
+ TargetHint _target; /**< Target that this tensor is pinned on */
+ arm_compute::TensorShape _tensor_shape; /**< SubTensor shape */
+ arm_compute::Coordinates _coords; /**< SubTensor Coordinates */
+ arm_compute::ITensor *_parent; /**< Parent tensor */
+ std::unique_ptr<arm_compute::ITensor> _subtensor; /**< SubTensor */
+};
+
+}
+#endif
+#endif //__TENSOR_H__
namespace boost { class mutex; }
namespace caffe {
+bool AclEnableSchedule(int enable=1);
#ifdef USE_PROFILING
class logtime_util
{
#include "caffe/layers/absval_layer.hpp"
#ifdef USE_ACL
-#include "caffe/acl_layer.hpp"
+#include "caffe/acl_operator.hpp"
#include "caffe/layers/acl_base_activation_layer.hpp"
#endif
const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom){
NOT_IMPLEMENTED;
}
- virtual void SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
- const vector<Blob<Dtype>*>& top, ActivationLayerInfo::ActivationFunction type);
+ virtual void SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top, arm_compute::ActivationLayerInfo::ActivationFunction type);
+ virtual bool Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top);
};
#endif
#include "caffe/layers/neuron_layer.hpp"
#ifdef USE_ACL
-#include "caffe/acl_layer.hpp"
+#include "caffe/acl_operator.hpp"
#include "caffe/layers/acl_base_activation_layer.hpp"
#endif
* Fallback to BNLLLayer for some corner cases.
*/
template <typename Dtype>
-class ACLBaseActivationLayer : public ACLBaseLayer<CLActivationLayer,NEActivationLayer> {
+class ACLBaseActivationLayer : public ACLOperator {
public:
- explicit ACLBaseActivationLayer(const LayerParameter& param)
+ explicit ACLBaseActivationLayer(const LayerParameter& param):ACLOperator(param)
{}
virtual void LayerSetUp(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top);
const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom){
NOT_IMPLEMENTED;
}
- virtual void SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
- const vector<Blob<Dtype>*>& top,ActivationLayerInfo::ActivationFunction type=ActivationLayerInfo::ActivationFunction::RELU);
+ virtual void SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top,arm_compute::ActivationLayerInfo::ActivationFunction type=arm_compute::ActivationLayerInfo::ActivationFunction::RELU);
};
#endif
+++ /dev/null
-#ifndef CAFFE_ACL_BASE_CONV_LAYER_HPP_
-#define CAFFE_ACL_BASE_CONV_LAYER_HPP_
-
-#include <vector>
-
-#include "caffe/blob.hpp"
-#include "caffe/layer.hpp"
-#include "caffe/proto/caffe.pb.h"
-
-#include "caffe/layers/conv_layer.hpp"
-
-#ifdef USE_ACL
-#include "caffe/acl_layer.hpp"
-#endif
-
-namespace caffe {
-
-#ifdef USE_ACL
-/*
- * @brief ACL implementation of ConvolutionLayer.
- * Fallback to ConvolutionLayer for some corner cases.
- *
-*/
-template <typename Dtype,typename GPUConvLayer,typename CPUConvLayer>
-class ACLConvolutionLayer : public ACLBaseLayer<GPUConvLayer,CPUConvLayer>,public ConvolutionLayer<Dtype> {
- public:
- explicit ACLConvolutionLayer(const LayerParameter& param)
- : ConvolutionLayer<Dtype>(param) {}
- virtual void LayerSetUp(const vector<Blob<Dtype>*>& bottom,
- const vector<Blob<Dtype>*>& top);
- virtual void Reshape(const vector<Blob<Dtype>*>& bottom,
- const vector<Blob<Dtype>*>& top);
- virtual ~ACLConvolutionLayer();
-
- protected:
- virtual void Forward_gpu(const vector<Blob<Dtype>*>& bottom,
- const vector<Blob<Dtype>*>& top);
- virtual void Forward_cpu(const vector<Blob<Dtype>*>& bottom,
- const vector<Blob<Dtype>*>& top);
- virtual void Backward_gpu(const vector<Blob<Dtype>*>& top,
- const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom){
- NOT_IMPLEMENTED;
- }
- virtual void Backward_cpu(const vector<Blob<Dtype>*>& top,
- const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom){
- NOT_IMPLEMENTED;
- }
- virtual void SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
- const vector<Blob<Dtype>*>& top);
-
-};
-#endif
-
-} // namespace caffe
-
-// Instantiate a class with float and double specifications.
-#define INSTANTIATE_CONV_CLASS(classname,GPUConvLayer,CPUConvLayer) \
- template class classname<float,GPUConvLayer,CPUConvLayer>; \
- template class classname<double,GPUConvLayer,CPUConvLayer>
-
-#endif // CAFFE_ACL_BASE_CONV_LAYER_HPP_
#include "caffe/layers/batch_norm_layer.hpp"
#ifdef USE_ACL
-#include "caffe/acl_layer.hpp"
+#include "caffe/acl_operator.hpp"
#endif
namespace caffe {
* Fallback to BatchNormLayer for some corner cases.
*/
template <typename Dtype>
-class ACLBatchNormLayer : public ACLBaseLayer<CLBatchNormalizationLayer,NEBatchNormalizationLayer>,public BatchNormLayer<Dtype> {
+class ACLBatchNormLayer : public ACLOperator,public BatchNormLayer<Dtype> {
public:
explicit ACLBatchNormLayer(const LayerParameter& param)
- : BatchNormLayer<Dtype>(param) {}
+ : ACLOperator(param),BatchNormLayer<Dtype>(param) {}
virtual void LayerSetUp(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top);
virtual void Reshape(const vector<Blob<Dtype>*>& bottom,
const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom){
NOT_IMPLEMENTED;
}
- virtual void SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
+ virtual void SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top);
+ virtual bool Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top);
};
#endif
#include "caffe/layers/bnll_layer.hpp"
#ifdef USE_ACL
-#include "caffe/acl_layer.hpp"
+#include "caffe/acl_operator.hpp"
#include "caffe/layers/acl_base_activation_layer.hpp"
#endif
const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom){
NOT_IMPLEMENTED;
}
- virtual void SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
- const vector<Blob<Dtype>*>& top, ActivationLayerInfo::ActivationFunction type);
+ virtual void SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top, arm_compute::ActivationLayerInfo::ActivationFunction type);
+ virtual bool Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top);
};
#endif
#include "caffe/layers/concat_layer.hpp"
#ifdef USE_ACL
-#include "caffe/acl_layer.hpp"
+#include "caffe/acl_operator.hpp"
#endif
namespace caffe {
* Fallback to ConcatLayer for some corner cases.
*/
template <typename Dtype>
-class ACLConcatLayer : public ACLBaseLayer<CLDepthConcatenate,NEDepthConcatenate>,public ConcatLayer<Dtype> {
+class ACLConcatLayer : public ACLOperator,public ConcatLayer<Dtype> {
public:
explicit ACLConcatLayer(const LayerParameter& param)
- : ConcatLayer<Dtype>(param) {}
+ : ACLOperator(param),ConcatLayer<Dtype>(param) {}
virtual void LayerSetUp(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top);
virtual void Reshape(const vector<Blob<Dtype>*>& bottom,
const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom){
NOT_IMPLEMENTED;
}
- virtual void SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
+ virtual void SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top);
- private:
- std::vector<ITensor *> cpu_vectors;
- std::vector<ICLTensor *> gpu_vectors;
+ virtual bool Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top);
};
#endif
#ifndef CAFFE_ACL_CONV_LAYER_HPP_
#define CAFFE_ACL_CONV_LAYER_HPP_
+#include <vector>
+
+#include "caffe/blob.hpp"
+#include "caffe/layer.hpp"
+#include "caffe/proto/caffe.pb.h"
+
+#include "caffe/layers/conv_layer.hpp"
+
#ifdef USE_ACL
-#include "caffe/layers/acl_base_conv_layer.hpp"
+#include "caffe/acl_operator.hpp"
#endif
namespace caffe {
-extern bool use_direct_conv_;
#ifdef USE_ACL
+/*
+ * @brief ACL implementation of ConvolutionLayer.
+ * Fallback to ConvolutionLayer for some corner cases.
+ *
+*/
template <typename Dtype>
-inline shared_ptr<Layer<Dtype> > GetACLConvolutionLayer(
- const LayerParameter& param) {
- ConvolutionParameter conv_param = param.convolution_param();
- const char* pDirectConv;
- pDirectConv = getenv ("DIRECTCONV");
- if (pDirectConv){
- unsigned int bdirectconv;
- sscanf(pDirectConv,"%i", &bdirectconv);
- if(bdirectconv != use_direct_conv_){
- use_direct_conv_ = bdirectconv;
- printf("DIRECTCONV<%s>\n", pDirectConv);
- printf("DIRECTCONV: %x\n", use_direct_conv_);
+class ACLConvolutionLayer : public ACLOperator,public ConvolutionLayer<Dtype> {
+ public:
+ explicit ACLConvolutionLayer(const LayerParameter& param)
+ : ACLOperator(param),ConvolutionLayer<Dtype>(param) {
+ }
+ virtual void LayerSetUp(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top);
+ virtual void Reshape(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top);
+ virtual ~ACLConvolutionLayer();
+
+ protected:
+ virtual void Forward_gpu(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top);
+ virtual void Forward_cpu(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top);
+ virtual void Backward_gpu(const vector<Blob<Dtype>*>& top,
+ const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom){
+ NOT_IMPLEMENTED;
}
- }
- int pad_data[3];
- if (conv_param.has_pad_h() || conv_param.has_pad_w()) {
- pad_data[0] = conv_param.pad_h();
- pad_data[1] = conv_param.pad_w();
- } else {
- const int kDefaultPad = 0;
- const int num_pad_dims = conv_param.pad_size();
- for (int i = 0; i < 2; ++i) {
- pad_data[i] = (num_pad_dims == 0) ? kDefaultPad :
- conv_param.pad((num_pad_dims == 1) ? 0 : i);
+ virtual void Backward_cpu(const vector<Blob<Dtype>*>& top,
+ const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom){
+ NOT_IMPLEMENTED;
}
- }
- if (use_direct_conv_ && ( (conv_param.kernel_size(0)==1 &&pad_data[0]==0 && pad_data[1]==0) || (conv_param.kernel_size(0)==3 && pad_data[0]<=1 && pad_data[1] <=1 ) )) {
- return shared_ptr<Layer<Dtype> >(new ACLConvolutionLayer<Dtype, CLConvolutionLayer, NEDirectConvolutionLayer>(param)); //NEDirectConvolutionLayer only for 1x1 and 3x3
- }
- return shared_ptr<Layer<Dtype> >(new ACLConvolutionLayer<Dtype, CLConvolutionLayer, NEConvolutionLayer>(param));
-}
+ virtual void SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top);
+ virtual bool Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top);
+};
+
#endif
} // namespace caffe
#include "caffe/layers/inner_product_layer.hpp"
#ifdef USE_ACL
-#include "caffe/acl_layer.hpp"
+#include "caffe/acl_operator.hpp"
#endif
namespace caffe {
* Fallback to InnerProductLayer for some corner cases.
*/
template <typename Dtype>
-class ACLInnerProductLayer : public ACLBaseLayer<CLFullyConnectedLayer,NEFullyConnectedLayer>,public InnerProductLayer<Dtype> {
+class ACLInnerProductLayer : public ACLOperator,public InnerProductLayer<Dtype> {
public:
explicit ACLInnerProductLayer(const LayerParameter& param)
- : InnerProductLayer<Dtype>(param) {}
+ : ACLOperator(param),InnerProductLayer<Dtype>(param) {
+ }
virtual void LayerSetUp(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top);
virtual void Reshape(const vector<Blob<Dtype>*>& bottom,
const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom){
NOT_IMPLEMENTED;
}
- virtual void SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
+ virtual void SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top);
+ virtual bool Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top);
};
#endif
#include "caffe/layers/local_connect_layer.hpp"
#ifdef USE_ACL
-#include "caffe/acl_layer.hpp"
+#include "caffe/acl_operator.hpp"
#endif
namespace caffe {
*
*/
template <typename Dtype>
-class ACLLocalConnectLayer : public ACLBaseLayer<CLLocallyConnectedLayer,NELocallyConnectedLayer>,public LocalConnectLayer<Dtype> {
+class ACLLocalConnectLayer : public ACLOperator,public LocalConnectLayer<Dtype> {
public:
explicit ACLLocalConnectLayer(const LayerParameter& param)
- : LocalConnectLayer<Dtype>(param) {}
+ : ACLOperator(param),LocalConnectLayer<Dtype>(param) {}
virtual void LayerSetUp(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top);
virtual void Reshape(const vector<Blob<Dtype>*>& bottom,
const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom){
NOT_IMPLEMENTED;
}
- virtual void SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
+ virtual void SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top);
-
+ virtual bool Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top);
};
#endif
#include "caffe/layers/lrn_layer.hpp"
#ifdef USE_ACL
-#include "caffe/acl_layer.hpp"
+#include "caffe/acl_operator.hpp"
#endif
namespace caffe {
* Fallback to LRNLayer for some corner cases.
*/
template <typename Dtype>
-class ACLLRNLayer : public ACLBaseLayer<CLNormalizationLayer,NENormalizationLayer>,public LRNLayer<Dtype> {
+class ACLLRNLayer : public ACLOperator,public LRNLayer<Dtype> {
public:
explicit ACLLRNLayer(const LayerParameter& param)
- : LRNLayer<Dtype>(param) {}
+ : ACLOperator(param),LRNLayer<Dtype>(param) {}
virtual void LayerSetUp(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top);
virtual void Reshape(const vector<Blob<Dtype>*>& bottom,
const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom){
NOT_IMPLEMENTED;
}
- virtual void SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
+ virtual void SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top);
+ virtual bool Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top);
};
#endif
#include "caffe/layers/pooling_layer.hpp"
#ifdef USE_ACL
-#include "caffe/acl_layer.hpp"
+#include "caffe/acl_operator.hpp"
#endif
namespace caffe {
* Fallback to PoolingLayer for some corner cases.
*/
template <typename Dtype>
-class ACLPoolingLayer : public ACLBaseLayer<CLPoolingLayer,NEPoolingLayer>,public PoolingLayer<Dtype> {
+class ACLPoolingLayer : public ACLOperator,public PoolingLayer<Dtype> {
public:
explicit ACLPoolingLayer(const LayerParameter& param)
- : PoolingLayer<Dtype>(param) {}
+ : ACLOperator(param),PoolingLayer<Dtype>(param) {}
virtual void LayerSetUp(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top);
virtual void Reshape(const vector<Blob<Dtype>*>& bottom,
const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom){
NOT_IMPLEMENTED;
}
- virtual void SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
+ virtual void SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top);
+ virtual bool Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top);
};
#endif
#include "caffe/layers/relu_layer.hpp"
#ifdef USE_ACL
-#include "caffe/acl_layer.hpp"
+#include "caffe/acl_operator.hpp"
#include "caffe/layers/acl_base_activation_layer.hpp"
#endif
const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom){
NOT_IMPLEMENTED;
}
- virtual void SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
+ virtual void SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top);
+ virtual bool Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top);
};
#endif
#include "caffe/layers/sigmoid_layer.hpp"
#ifdef USE_ACL
-#include "caffe/acl_layer.hpp"
+#include "caffe/acl_operator.hpp"
#include "caffe/layers/acl_base_activation_layer.hpp"
#endif
const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom){
NOT_IMPLEMENTED;
}
- virtual void SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
- const vector<Blob<Dtype>*>& top, ActivationLayerInfo::ActivationFunction type);
+ virtual void SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top, arm_compute::ActivationLayerInfo::ActivationFunction type);
+ virtual bool Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top);
};
#endif
#include "caffe/layers/softmax_layer.hpp"
#ifdef USE_ACL
-#include "caffe/acl_layer.hpp"
+#include "caffe/acl_operator.hpp"
#endif
namespace caffe {
* Fallback to SoftmaxLayer for some corner cases.
*/
template <typename Dtype>
-class ACLSoftmaxLayer : public ACLBaseLayer<CLSoftmaxLayer,NESoftmaxLayer>,public SoftmaxLayer<Dtype> {
+class ACLSoftmaxLayer : public ACLOperator,public SoftmaxLayer<Dtype> {
public:
explicit ACLSoftmaxLayer(const LayerParameter& param)
- : SoftmaxLayer<Dtype>(param) {}
+ : ACLOperator(param),SoftmaxLayer<Dtype>(param) {}
virtual void LayerSetUp(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top);
virtual void Reshape(const vector<Blob<Dtype>*>& bottom,
const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom){
NOT_IMPLEMENTED;
}
- virtual void SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
+ virtual void SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top);
+ virtual bool Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top);
};
#endif
#include "caffe/layers/tanh_layer.hpp"
#ifdef USE_ACL
-#include "caffe/acl_layer.hpp"
+#include "caffe/acl_operator.hpp"
#include "caffe/layers/acl_base_activation_layer.hpp"
#endif
const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom){
NOT_IMPLEMENTED;
}
- virtual void SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
- const vector<Blob<Dtype>*>& top, ActivationLayerInfo::ActivationFunction type);
+ virtual void SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top, arm_compute::ActivationLayerInfo::ActivationFunction type);
+ virtual bool Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top);
};
#endif
+++ /dev/null
-#ifdef USE_ACL
-#include "caffe/acl_layer.hpp"
-
-unsigned int bypass_acl_class_layer = (0 | \
- /*0xffffffff |*/ \
- /*FLAGS_ENABLE_ACL_FC |*/ \
- /*FLAGS_ENABLE_ACL_LRN |*/ \
- 0 );
-
-#ifdef USE_PROFILING
-
-#include "arm_neon.h"
-
-unsigned int acl_log_flags = (0 | \
- MASK_LOG_APP_TIME | \
- /*MASK_LOG_ALLOCATE | */\
- /*MASK_LOG_ALLOCATE | */\
- /*MASK_LOG_RUN | */\
- /*MASK_LOG_CONFIG | */\
- /*MASK_LOG_COPY | */\
- MASK_LOG_ABSVAL | \
- MASK_LOG_BNLL | \
- MASK_LOG_CONV | \
- MASK_LOG_FC | \
- MASK_LOG_LRN | \
- MASK_LOG_POOLING | \
- MASK_LOG_RELU | \
- MASK_LOG_SIGMOID | \
- MASK_LOG_SOFTMAX | \
- MASK_LOG_TANH | \
- MASK_LOG_LC | \
- MASK_LOG_BN | \
- MASK_LOG_CONCAT | \
- 0);
-#include <stdio.h> /* printf */
-#include <stdlib.h> /* getenv */
-#endif //USE_PROFILING
-
-namespace caffe {
-template <typename GPULayer, typename CPULayer>
-ACLBaseLayer<GPULayer,CPULayer>::ACLBaseLayer()
- :init_layer_(true),force_bypass_acl_path_(false){
- const char* pBypassACL;
- pBypassACL = getenv ("BYPASSACL");
- if (pBypassACL){
- unsigned int bacl;
- sscanf(pBypassACL,"%i", &bacl);
- if(bacl != bypass_acl_class_layer){
- bypass_acl_class_layer = bacl;
- printf("BYPASSACL<%s>\n", pBypassACL);
- printf("BYPASSACL: %x\n", bypass_acl_class_layer);
- }
- }
-#ifdef USE_PROFILING
- const char* pLogACL;
- pLogACL = getenv("LOGACL");
- if (pLogACL){
- unsigned int alf;
- sscanf(pLogACL,"%i", &alf);
- if (alf != acl_log_flags){
- acl_log_flags = alf;
- printf("LOGACL<%s>\n", pLogACL);
- printf("LOGACL: %x\n", acl_log_flags);
- }
- }
-#endif //USE_PROFILING
-}
-template <typename GPULayer, typename CPULayer>
-void ACLBaseLayer<GPULayer,CPULayer>::gpu_run() {
- gpu_.run(true);
-}
-template <typename GPULayer, typename CPULayer>
-void ACLBaseLayer<GPULayer,CPULayer>::cpu_run() {
- cpu_.run(false);
-}
-
-template <typename GPULayer, typename CPULayer>
-ACLBaseLayer<GPULayer,CPULayer>::~ACLBaseLayer(){
-}
-template <typename GPULayer, typename CPULayer>
-template <typename ACLTensor> bool ACLBaseLayer<GPULayer,CPULayer>::new_tensor(ACLTensor *&tensor,TensorShape shape,void *mem,bool share)
-{
- tensor=new ACLTensor(share);
-#if 1 //F32
- tensor->allocator()->init(TensorInfo(shape, Format::F32));
-#else //F16
- tensor->allocator()->init(TensorInfo(shape, Format::F16));
-#endif
- tensor->bindmem(mem,share);
- return true;
-}
-
-template <typename ACLTensor>
-void BaseTensor<ACLTensor>::commit(TensorType type){
- settensortype(type);
- if (!share_&&mem_) {
- if (!allocate_){
-#ifdef USE_PROFILING
- logtime_util log_time(ACL_ALLOCATE_INFO);
-#endif //USE_PROFILING
- ACLTensor::allocator()->allocate();
- allocate_=true;
- }
- if (type_!= tensor_output) {
- tensor_copy(mem_);
- }
- mem_=nullptr;
- }
-}
-
-template <typename ACLTensor>
-int BaseTensor<ACLTensor>::tensor_copy(void * mem,bool toTensor)
-{
-#ifdef USE_PROFILING
- logtime_util log_time(ACL_COPY_INFO);
-#endif //USE_PROFILING
- arm_compute::Window window;
- ACLTensor* tensor=this;
- window.use_tensor_dimensions(tensor->info()->tensor_shape(), /* first_dimension =*/Window::DimY); // Iterate through the rows (not each element)
- int width = tensor->info()->tensor_shape()[0]; //->dimension(0); //window.x().end() - window.x().start(); // + 1;
- int height = tensor->info()->tensor_shape()[1]; //->dimension(1); //window.y().end() - window.y().start(); // + 1;
- int deepth = tensor->info()->tensor_shape()[2];
- map();
- // Create an iterator:
- arm_compute::Iterator it(tensor, window);
- // Except it works for an arbitrary number of dimensions
- if (toTensor) { //mem->tensor
- arm_compute::execute_window_loop(window, [&](const arm_compute::Coordinates & id)
- {
-#if 0 //F16
- if (tensor->info()->element_size() ==2)
- {
- for(int i = 0; i < width; i+= 4){
- auto pa = (float32x4_t*)((char*)mem) + ((id[3] * (width * height * deepth) + id.z() * (width * height) + id.y() * width + id.x() + i) * 4);
- *(float16x4_t*)(((char*)it.ptr()) + i*2) = vcvt_f16_f32(*pa);
- }
- }
- else{
-#endif
- memcpy(it.ptr(), ((char*)mem) + ((id[3] * (width * height * deepth) + id.z() * (width * height) + id.y() * width + id.x()) * tensor->info()->element_size()), width * tensor->info()->element_size());
-#if 0 //F16
- }
-#endif
- },
- it);
- }else{ //tensor-->mem
- arm_compute::execute_window_loop(window, [&](const arm_compute::Coordinates & id)
- {
-#if 0 //F16
- if (tensor->info()->element_size() ==2)
- {
- for(int i = 0; i < width; i+= 4){
- auto pa = (float32x4_t*)(((char*)mem) + ((id[3] * (width * height * deepth) + id.z() * (width * height) + id.y() * width + id.x() + i) * 4));
- *pa = vcvt_f32_f16(*(float16x4_t*)(((char*)it.ptr()) + i*2));
- }
- }
- else{
-#endif
- memcpy(((char*)mem) + ((id[3] * (width * height * deepth) + id.z() * (width * height) + id.y() * width) * tensor->info()->element_size()), it.ptr(), width * tensor->info()->element_size());
-#if 0 //F16
- }
-#endif
- },
- it);
- }
- unmap();
-
- return 0;
-}
-
-template <typename GPULayer, typename CPULayer>
-template <typename ACLTensor> bool ACLBaseLayer<GPULayer,CPULayer>::tensor_mem(ACLTensor *tensor,void *mem,bool share)
-{
- tensor->bindmem(mem,share);
- return true;
-}
-
-template <typename GPULayer, typename CPULayer>
-template <typename ACLTensor> bool ACLBaseLayer<GPULayer,CPULayer>::tensor_mem(void *mem,ACLTensor *tensor,bool share)
-{
- if (mem==tensor->buffer()) return true;
- if (!share) {
- tensor->tensor_copy(mem,false);
- }
- return true;
-}
-
-
-template <typename GPULayer, typename CPULayer>
-bool ACLBaseLayer<GPULayer,CPULayer>::checkreshape(TensorShape shape,bool gpu, TensorType type)
-{
- if (gpu) {
- init_layer_ = gpu_.reshape(shape,type);
- }else{
- init_layer_ = cpu_.reshape(shape,type);
- }
- return init_layer_;
-}
-
-template <typename GPULayer, typename CPULayer>
-GPULayer * ACLBaseLayer<GPULayer,CPULayer>::new_gpulayer(){
- gpu_.layer= new GPULayer;
- return gpu_.layer;
-}
-template <typename GPULayer, typename CPULayer>
-CPULayer * ACLBaseLayer<GPULayer,CPULayer>::new_cpulayer(){
- cpu_.layer= new CPULayer;
- return cpu_.layer;
-}
-template <typename ACLLayer,typename ACLTensor>
-bool ACLXPUBaseLayer<ACLLayer,ACLTensor>::reshape(TensorShape &shape,TensorType type)
-{
- TensorShape _shape;
- if (!layer) return true;
-#ifdef USE_CONV_CACHE
- if (tensor_input == type){
- _shape = input->info()->tensor_shape();
- if (_shape.total_size()==shape.total_size() && _shape[0]==shape[0] && _shape[1]==shape[1]) {
- return false;
- }
- for(int i = 0; i < 16; ++i){
- if(cache.input[i] == nullptr) break;
- _shape = cache.input[i]->info()->tensor_shape();
- if (_shape.total_size()==shape.total_size() && _shape[0]==shape[0] && _shape[1]==shape[1]) {
- this->layer = cache.layer[i];
- this->input = cache.input[i];
- this->output = cache.output[i];
- this->weights = cache.weights[i];
- this->biases = cache.biases[i];
- return false;
- }
- }
- }
-#endif //USE_CONV_CACHE
- switch (type) {
- case tensor_biases:
- _shape = biases->info()->tensor_shape();
- break;
- case tensor_weights:
- _shape = weights->info()->tensor_shape();
- break;
- case tensor_output:
- _shape = output->info()->tensor_shape();
- break;
- case tensor_input:
- default:
- _shape = input->info()->tensor_shape();
- break;
- }
- if (_shape.total_size()==shape.total_size() && _shape[0]==shape[0] && _shape[1]==shape[1]) {
- return false;
- }
- freelayer();
- return true;
-}
-
-INSTANTIATE_ACLBASECLASS(CLNormalizationLayer,NENormalizationLayer);
- INSTANTIATE_ACLBASE_FUNCTION(CLNormalizationLayer,NENormalizationLayer,GPUTensor);
- INSTANTIATE_ACLBASE_FUNCTION(CLNormalizationLayer,NENormalizationLayer,CPUTensor);
-INSTANTIATE_ACLBASECLASS(CLActivationLayer,NEActivationLayer);
- INSTANTIATE_ACLBASE_FUNCTION(CLActivationLayer,NEActivationLayer,GPUTensor);
- INSTANTIATE_ACLBASE_FUNCTION(CLActivationLayer,NEActivationLayer,CPUTensor);
-INSTANTIATE_ACLBASECLASS(CLPoolingLayer,NEPoolingLayer);
- INSTANTIATE_ACLBASE_FUNCTION(CLPoolingLayer,NEPoolingLayer,GPUTensor);
- INSTANTIATE_ACLBASE_FUNCTION(CLPoolingLayer,NEPoolingLayer,CPUTensor);
-INSTANTIATE_ACLBASECLASS(CLSoftmaxLayer,NESoftmaxLayer);
- INSTANTIATE_ACLBASE_FUNCTION(CLSoftmaxLayer,NESoftmaxLayer,GPUTensor);
- INSTANTIATE_ACLBASE_FUNCTION(CLSoftmaxLayer,NESoftmaxLayer,CPUTensor);
-INSTANTIATE_ACLBASECLASS(CLFullyConnectedLayer,NEFullyConnectedLayer);
- INSTANTIATE_ACLBASE_FUNCTION(CLFullyConnectedLayer,NEFullyConnectedLayer,GPUTensor);
- INSTANTIATE_ACLBASE_FUNCTION(CLFullyConnectedLayer,NEFullyConnectedLayer,CPUTensor);
-INSTANTIATE_ACLBASECLASS(CLConvolutionLayer,NEConvolutionLayer);
- INSTANTIATE_ACLBASE_FUNCTION(CLConvolutionLayer,NEConvolutionLayer,GPUTensor);
- INSTANTIATE_ACLBASE_FUNCTION(CLConvolutionLayer,NEConvolutionLayer,CPUTensor);
-INSTANTIATE_ACLBASECLASS(CLConvolutionLayer,NEDirectConvolutionLayer);
- INSTANTIATE_ACLBASE_FUNCTION(CLConvolutionLayer,NEDirectConvolutionLayer,GPUTensor);
- INSTANTIATE_ACLBASE_FUNCTION(CLConvolutionLayer,NEDirectConvolutionLayer,CPUTensor);
-INSTANTIATE_ACLBASECLASS(CLBatchNormalizationLayer,NEBatchNormalizationLayer);
- INSTANTIATE_ACLBASE_FUNCTION(CLBatchNormalizationLayer,NEBatchNormalizationLayer,GPUTensor);
- INSTANTIATE_ACLBASE_FUNCTION(CLBatchNormalizationLayer,NEBatchNormalizationLayer,CPUTensor);
-INSTANTIATE_ACLBASECLASS(CLLocallyConnectedLayer,NELocallyConnectedLayer);
- INSTANTIATE_ACLBASE_FUNCTION(CLLocallyConnectedLayer,NELocallyConnectedLayer,GPUTensor);
- INSTANTIATE_ACLBASE_FUNCTION(CLLocallyConnectedLayer,NELocallyConnectedLayer,CPUTensor);
-INSTANTIATE_ACLBASECLASS(CLDepthConcatenate,NEDepthConcatenate);
- INSTANTIATE_ACLBASE_FUNCTION(CLDepthConcatenate,NEDepthConcatenate,GPUTensor);
- INSTANTIATE_ACLBASE_FUNCTION(CLDepthConcatenate,NEDepthConcatenate,CPUTensor);
-}
-
-#endif
--- /dev/null
+#ifdef USE_ACL
+#include "caffe/acl_operator.hpp"
+#include "caffe/common.hpp"
+
+unsigned int bypass_acl_class_layer = (0 | \
+ FLAGS_ENABLE_ACL_CONCAT | \
+ /*0xffffffff |*/ \
+ /*FLAGS_ENABLE_ACL_FC |*/ \
+ /*FLAGS_ENABLE_ACL_LRN |*/ \
+ 0 );
+
+unsigned int openailab_intfp = 0;
+int enable_schedule=0;
+
+#ifdef USE_PROFILING
+
+#include "arm_neon.h"
+
+unsigned int acl_log_flags = (0 | \
+ MASK_LOG_APP_TIME | \
+ /*MASK_LOG_ALLOCATE | */\
+ /*MASK_LOG_ALLOCATE | */\
+ /*MASK_LOG_RUN | */\
+ /*MASK_LOG_CONFIG | */\
+ /*MASK_LOG_COPY | */\
+ MASK_LOG_ABSVAL | \
+ MASK_LOG_BNLL | \
+ MASK_LOG_CONV | \
+ MASK_LOG_FC | \
+ MASK_LOG_LRN | \
+ MASK_LOG_POOLING | \
+ MASK_LOG_RELU | \
+ MASK_LOG_SIGMOID | \
+ MASK_LOG_SOFTMAX | \
+ MASK_LOG_TANH | \
+ MASK_LOG_LC | \
+ MASK_LOG_BN | \
+ MASK_LOG_CONCAT | \
+ 0);
+#include <stdio.h> /* printf */
+#include <stdlib.h> /* getenv */
+#endif //USE_PROFILING
+
+namespace caffe {
+bool AclEnableSchedule(int enable){
+ enable_schedule=enable;
+ if (enable) {
+ Caffe::set_mode(Caffe::GPU);
+ }
+ return true;
+}
+int isScheduleEnable()
+{
+ return enable_schedule;
+}
+bool ACLOperator::init_cl_env=true;
+bool ACLOperator::support_opencl_=false;
+bool opencl_is_available()
+{
+ return arm_compute::opencl_is_available();
+}
+ACLOperator::ACLOperator(const LayerParameter& param)
+ :operator_state_(operator_not_init),force_bypass_acl_path_(false),
+ target_hint_(TargetHint::DONT_CARE),
+ convolution_method_hint_(ConvolutionMethodHint::GEMM),
+ _group(1),name_(""){
+ const char* pBypassACL;
+ if(init_cl_env){
+#ifdef USE_OPENCL
+ try {
+ if (opencl_is_available()) {
+ arm_compute::CLScheduler::get().default_init();
+ support_opencl_=true;
+ }
+ }catch(std::exception& e){
+ support_opencl_=false;
+ }
+#endif
+ init_cl_env=false;
+ }
+ pBypassACL = getenv ("BYPASSACL");
+ if (pBypassACL){
+ unsigned int bacl;
+ sscanf(pBypassACL,"%i", &bacl);
+ if(bacl != bypass_acl_class_layer){
+ bypass_acl_class_layer = bacl;
+ printf("BYPASSACL<%s>\n", pBypassACL);
+ printf("BYPASSACL: %x\n", bypass_acl_class_layer);
+ }
+ }
+
+ const string& layer_type = param.type();
+ if (layer_type=="Convolution") {
+ ConvolutionParameter conv_param = param.convolution_param();
+ const char* pDirectConv;
+ unsigned int use_direct_conv=0;
+ pDirectConv = getenv ("DIRECTCONV");
+ if (pDirectConv){
+ unsigned int bdirectconv;
+ sscanf(pDirectConv,"%i", &bdirectconv);
+ if(bdirectconv != use_direct_conv){
+ use_direct_conv = bdirectconv;
+ printf("DIRECTCONV<%s>\n", pDirectConv);
+ printf("DIRECTCONV: %x\n", use_direct_conv);
+ }
+ }
+ int pad_data[3];
+ if (conv_param.has_pad_h() || conv_param.has_pad_w()) {
+ pad_data[0] = conv_param.pad_h();
+ pad_data[1] = conv_param.pad_w();
+ } else {
+ const int kDefaultPad = 0;
+ const int num_pad_dims = conv_param.pad_size();
+ for (int i = 0; i < 2; ++i) {
+ pad_data[i] = (num_pad_dims == 0) ? kDefaultPad :
+ conv_param.pad((num_pad_dims == 1) ? 0 : i);
+ }
+ }
+ if (use_direct_conv && ( (conv_param.kernel_size(0)==1 &&pad_data[0]==0 && pad_data[1]==0) || (conv_param.kernel_size(0)==3 && pad_data[0]<=1 && pad_data[1] <=1 ) )) {
+ convolution_method_hint_=ConvolutionMethodHint::DIRECT; //NEDirectConvolutionLayer only for 1x1 and 3x3
+ }
+ }
+
+#ifdef USE_PROFILING
+ const char* pLogACL;
+ pLogACL = getenv("LOGACL");
+ if (pLogACL){
+ unsigned int alf;
+ sscanf(pLogACL,"%i", &alf);
+ if (alf != acl_log_flags){
+ acl_log_flags = alf;
+ printf("LOGACL<%s>\n", pLogACL);
+ printf("LOGACL: %x\n", acl_log_flags);
+ }
+ }
+#endif //USE_PROFILING
+ const char* pEnableSchedule;
+ pEnableSchedule = getenv ("ENABLESCHEDULE");
+ if (pEnableSchedule){
+ unsigned int bshedule;
+ sscanf(pEnableSchedule,"%i", &bshedule);
+ if(bshedule != enable_schedule){
+ enable_schedule = bshedule;
+ printf("ENABLESCHEDULE<%s>\n", pEnableSchedule);
+ printf("ENABLESCHEDULE: %x\n", enable_schedule);
+ }
+ if (enable_schedule) {
+ AclEnableSchedule(1);
+ }
+ }
+}
+ACLOperator::~ACLOperator() {
+}
+
+bool ACLOperator::new_tensor(std::unique_ptr<ACLTensor> &tensor,arm_compute::TensorShape &shape,void *mem,bool commit)
+{
+ auto acl_tensor=new ACLTensor(arm_compute::TensorInfo(shape, arm_compute::Format::F32));
+ acl_tensor->set_target(getTargetHint());
+ acl_tensor->bindmem(mem);
+ if (commit) acl_tensor->commit();
+ tensor=(std::unique_ptr<ACLTensor>) std::move(acl_tensor);
+ return true;
+}
+bool ACLOperator::new_tensor(std::unique_ptr<ACLSubTensor> &tensor,std::unique_ptr<ACLTensor> &parent,arm_compute::TensorShape &shape,arm_compute::Coordinates& coord)
+{
+ auto acl_tensor=new ACLSubTensor(parent,shape, coord);
+ acl_tensor->set_target(getTargetHint());
+ tensor=(std::unique_ptr<ACLSubTensor>) std::move(acl_tensor);
+ return true;
+}
+
+void ACLTensor::commit(TensorType type)
+{
+ settensortype(type);
+ if (mem_) {
+ if (!allocate_){
+#ifdef USE_PROFILING
+ logtime_util log_time(ACL_ALLOCATE_INFO);
+#endif //USE_PROFILING
+ allocate();
+ allocate_=true;
+ }
+ if (type_!= tensor_output) {
+ tensor_copy(mem_);
+ }
+ mem_=nullptr;
+ }
+}
+
+int BaseACLTensor::tensor_copy(arm_compute::ITensor* tensor,void * mem,bool toTensor)
+{
+#ifdef USE_PROFILING
+ logtime_util log_time(ACL_COPY_INFO);
+#endif //USE_PROFILING
+ arm_compute::Window window;
+ window.use_tensor_dimensions(tensor->info()->tensor_shape(), /* first_dimension =*/arm_compute::Window::DimY); // Iterate through the rows (not each element)
+ int width = tensor->info()->tensor_shape()[0];
+ int height = tensor->info()->tensor_shape()[1];
+ int deepth = tensor->info()->tensor_shape()[2];
+ map();
+ // Create an iterator:
+ arm_compute::Iterator it(tensor, window);
+ // Except it works for an arbitrary number of dimensions
+ if (toTensor) { //mem->tensor
+ arm_compute::execute_window_loop(window, [&](const arm_compute::Coordinates & id)
+ {
+ memcpy(it.ptr(), ((char*)mem) + ((id[3] * (width * height * deepth) + id.z() * (width * height) + id.y() * width + id.x()) * tensor->info()->element_size()), width * tensor->info()->element_size());
+ },
+ it);
+ }else{ //tensor-->mem
+ arm_compute::execute_window_loop(window, [&](const arm_compute::Coordinates & id)
+ {
+ memcpy(((char*)mem) + ((id[3] * (width * height * deepth) + id.z() * (width * height) + id.y() * width) * tensor->info()->element_size()), it.ptr(), width * tensor->info()->element_size());
+ },
+ it);
+ }
+ unmap();
+
+ return 0;
+}
+
+INIT_GLOBAL_FUNCS();
+
+}
+
+
+#endif
--- /dev/null
+#include "caffe/acl_tensor.hpp"
+
+namespace caffe {
+
+#ifdef USE_ACL
+template <typename TensorType>
+std::unique_ptr<arm_compute::ITensor> initialise_tensor(arm_compute::TensorInfo &info)
+{
+ auto tensor = cpp14::make_unique<TensorType>();
+ tensor->allocator()->init(info);
+ return std::move(tensor);
+}
+
+template <typename TensorType>
+void tensor_allocate(arm_compute::ITensor &tensor)
+{
+ auto itensor = dynamic_cast<TensorType *>(&tensor);
+ itensor->allocator()->allocate();
+}
+
+Tensor::Tensor(arm_compute::TensorInfo &info) noexcept
+ : _target(TargetHint::DONT_CARE), _info(info), _tensor(nullptr)
+{
+}
+
+Tensor::Tensor(Tensor &&src) noexcept
+ : _target(src._target),
+ _info(std::move(src._info)),
+ _tensor(std::move(src._tensor))
+{
+}
+
+arm_compute::ITensor *Tensor::set_target(TargetHint target)
+{
+ switch(target)
+ {
+#ifdef USE_OPENCL
+ case TargetHint::OPENCL:
+ _tensor = initialise_tensor<arm_compute::CLTensor>(_info);
+ break;
+#endif
+ case TargetHint::NEON:
+ _tensor = initialise_tensor<arm_compute::Tensor>(_info);
+ break;
+ default:
+ break;
+ }
+ _target = target;
+ return _tensor.get();
+}
+
+void Tensor::allocate()
+{
+ switch(_target)
+ {
+#ifdef USE_OPENCL
+ case TargetHint::OPENCL:
+ tensor_allocate<arm_compute::CLTensor>(*_tensor);
+ break;
+#endif
+ case TargetHint::NEON:
+ tensor_allocate<arm_compute::Tensor>(*_tensor);
+ break;
+ default:
+ break;
+ }
+}
+void Tensor::map(bool blocking){
+#ifdef USE_OPENCL
+ if (_target==TargetHint::OPENCL)
+ dynamic_cast<arm_compute::CLTensor *>(tensor())->map(blocking);
+#endif
+}
+void Tensor::unmap(){
+#ifdef USE_OPENCL
+ if (_target==TargetHint::OPENCL)
+ dynamic_cast<arm_compute::CLTensor *>(tensor())->unmap();
+#endif
+}
+
+template <typename SubTensorType, typename ParentTensorType>
+std::unique_ptr<arm_compute::ITensor> initialise_subtensor(arm_compute::ITensor *parent, arm_compute::TensorShape shape, arm_compute::Coordinates coords)
+{
+ auto ptensor = dynamic_cast<ParentTensorType *>(parent);
+ auto subtensor = cpp14::make_unique<SubTensorType>(ptensor, shape, coords);
+ return std::move(subtensor);
+}
+SubTensor::SubTensor(Tensor* parent, arm_compute::TensorShape& tensor_shape, arm_compute::Coordinates& coords) noexcept
+ : _target(TargetHint::DONT_CARE), _tensor_shape(tensor_shape), _coords(coords), _parent(nullptr), _subtensor(nullptr)
+{
+ _parent = parent->tensor();
+ _target = parent->target();
+
+ instantiate_subtensor();
+}
+arm_compute::ITensor *SubTensor::set_target(TargetHint target)
+{
+ return (target == _target) ? _subtensor.get() : nullptr;
+}
+
+arm_compute::ITensor *SubTensor::tensor()
+{
+ return _subtensor.get();
+}
+
+const arm_compute::ITensor *SubTensor::tensor() const
+{
+ return _subtensor.get();
+}
+
+TargetHint SubTensor::target() const
+{
+ return _target;
+}
+
+void SubTensor::allocate()
+{
+ // NOP for sub-tensors
+}
+
+void SubTensor::instantiate_subtensor()
+{
+ switch(_target)
+ {
+#ifdef USE_OPENCL
+ case TargetHint::OPENCL:
+ _subtensor = initialise_subtensor<arm_compute::CLSubTensor, arm_compute::ICLTensor>(_parent, _tensor_shape, _coords);
+ break;
+#endif
+ default:
+ case TargetHint::NEON:
+ _subtensor = initialise_subtensor<arm_compute::SubTensor, arm_compute::ITensor>(_parent, _tensor_shape, _coords);
+ break;
+ }
+}
+
+#endif
+}
#include "caffe/common.hpp"
#include "caffe/util/rng.hpp"
-#ifdef USE_ACL
-#include "arm_compute/runtime/CL/CLScheduler.h"
-using namespace arm_compute;
-#endif
namespace caffe {
Caffe::Caffe()
: random_generator_(), mode_(Caffe::CPU),use_mali_gpu_(false),
- solver_count_(1), solver_rank_(0), multiprocess_(false) {
-#ifdef USE_ACL
-
- try {
- CLScheduler::get().default_init();
- }
- catch(std::exception& e)
+ solver_count_(1), solver_rank_(0), multiprocess_(false)
{
- std::cout << "OPENCL initialization failed"<< std::endl;
- }
-
-#endif
}
Caffe::~Caffe() { }
ConvolutionParameter conv_param = param.convolution_param();
ConvolutionParameter_Engine engine = conv_param.engine();
#ifdef USE_ACL
- return GetACLConvolutionLayer<Dtype>(param);
+ return shared_ptr<Layer<Dtype> >(new ACLConvolutionLayer<Dtype>(param));
#endif
#ifdef USE_CUDNN
bool use_dilation = false;
}
template <typename Dtype>
-void ACLAbsValLayer<Dtype>::SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
- const vector<Blob<Dtype>*>& top,ActivationLayerInfo::ActivationFunction type){
- ACLBaseActivationLayer<Dtype>::SetupACLLayer(bottom, top,ActivationLayerInfo::ActivationFunction::ABS);
+void ACLAbsValLayer<Dtype>::SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top,arm_compute::ActivationLayerInfo::ActivationFunction type){
+ ACLBaseActivationLayer<Dtype>::SetupACLOperator(bottom, top,arm_compute::ActivationLayerInfo::ActivationFunction::ABS);
}
template <typename Dtype>
}
template <typename Dtype>
+bool ACLAbsValLayer<Dtype>::Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top){
+ bool bypass_acl=false;
+ if (this->force_bypass_acl_path_) {
+ bypass_acl=true;
+ }
+ return bypass_acl;
+}
+
+template <typename Dtype>
void ACLAbsValLayer<Dtype>::Forward_cpu(
const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) {
#ifdef USE_PROFILING
logtime_util log_time(ACL_ABSVAL_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_) {
+ if (Bypass_acl(bottom,top)) {
AbsValLayer<Dtype>::Forward_cpu(bottom,top);
return;
}
#ifdef USE_PROFILING
logtime_util log_time(ACL_ABSVAL_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_) {
+ if (Bypass_acl(bottom,top)) {
AbsValLayer<Dtype>::Forward_cpu(bottom,top);
return;
}
const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) {
}
template <typename Dtype>
-void ACLBaseActivationLayer<Dtype>::SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
- const vector<Blob<Dtype>*>& top,ActivationLayerInfo::ActivationFunction type){
+void ACLBaseActivationLayer<Dtype>::SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top,arm_compute::ActivationLayerInfo::ActivationFunction type){
const unsigned int count = bottom[0]->count();
const unsigned int count_ = top[0]->count();
- TensorShape input_shape(count);
- TensorShape output_shape(count_);
- checkreshape(input_shape,Caffe::arm_gpu_mode());
- if (!this->init_layer_) return;
- this->init_layer_=false;
- // Initialize ACL.
- if (Caffe::arm_gpu_mode()) {
- new_gpulayer();
- }else{
- new_cpulayer();
- }
+ arm_compute::TensorShape input_shape(count);
+ arm_compute::TensorShape output_shape(count_);
+ if (is_operator_init_done(input_shape)) return;
+ set_operator_init_done();
- this->force_bypass_acl_path_=false;
- ActivationLayerInfo act_info(type);
+ // Initialize ACL.
+ arm_compute::ActivationLayerInfo act_info(type);
- if(type== ActivationLayerInfo::ActivationFunction::TANH)
- act_info=ActivationLayerInfo(type,1.0,1.0);
-
-
+ if(type== arm_compute::ActivationLayerInfo::ActivationFunction::TANH)
+ act_info=arm_compute::ActivationLayerInfo(type,1.0,1.0);
- if (Caffe::arm_gpu_mode()) {
- Dtype *top_data = top[0]->mutable_gpu_data();
- const Dtype* bottom_data = bottom[0]->gpu_data();
- new_tensor(this->gpu().input,input_shape,(void*)bottom_data);
- new_tensor(this->gpu().output,output_shape,(void*)top_data);
-#ifdef USE_PROFILING
- logtime_util log_time(ACL_CONFIG_INFO);
-#endif //USE_PROFILING
- this->gpu().layer->configure(this->gpu().input,this->gpu().output,act_info);
- }else{
- Dtype *top_data = top[0]->mutable_cpu_data();
- const Dtype* bottom_data = bottom[0]->cpu_data();
- new_tensor(this->cpu().input,input_shape,(void*)bottom_data);
- new_tensor(this->cpu().output,output_shape,(void*)top_data);
-#ifdef USE_PROFILING
- logtime_util log_time(ACL_CONFIG_INFO);
-#endif //USE_PROFILING
- this->cpu().layer->configure(this->cpu().input,this->cpu().output,act_info);
- }
+ new_tensor(input(),input_shape,(void*)InputdataPtr(this,bottom));
+ new_tensor(output(),output_shape,(void*)OutputdataPtr(this,top));
+ acl_configure(activation,this,act_info);
}
template <typename Dtype>
void ACLBaseActivationLayer<Dtype>::Reshape(
template <typename Dtype>
void ACLBaseActivationLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top) {
- if(Caffe::arm_gpu_mode()){
- Forward_gpu(bottom, top);
+ if(isGPUMode()){
+ ACLBaseActivationLayer<Dtype>::Forward_gpu(bottom, top);
return;
}
- Dtype* top_data = top[0]->mutable_cpu_data();
- const Dtype* bottom_data = bottom[0]->cpu_data();
- SetupACLLayer(bottom,top);
- tensor_mem(this->cpu().input,(void*)(bottom_data));
- cpu_run();
- tensor_mem((void*)(top_data),this->cpu().output);
+ SetupACLOperator(bottom,top);
+ caffe::acl_run(this,bottom,top);
}
template <typename Dtype>
void ACLBaseActivationLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top) {
- Dtype* top_data = top[0]->mutable_gpu_data();
- const Dtype* bottom_data = bottom[0]->gpu_data();
- SetupACLLayer(bottom,top);
- tensor_mem(this->gpu().input,(void*)(bottom_data));
- gpu_run();
- tensor_mem((void*)(top_data),this->gpu().output);
+ SetupACLOperator(bottom,top);
+ caffe::acl_run(this,bottom,top);
}
template <typename Dtype>
+++ /dev/null
-#ifdef USE_ACL
-#include <algorithm>
-#include <vector>
-
-#include "caffe/filler.hpp"
-#include "caffe/layers/acl_conv_layer.hpp"
-
-namespace caffe {
-
-bool use_direct_conv_=false;
-template <typename Dtype,typename GPUConvLayer,typename CPUConvLayer>
-void ACLConvolutionLayer<Dtype,GPUConvLayer,CPUConvLayer>::LayerSetUp(
- const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) {
- ConvolutionLayer<Dtype>::LayerSetUp(bottom, top);
- this->force_bypass_acl_path_= bypass_acl_class_layer & FLAGS_ENABLE_ACL_CONV;
-}
-
-template <typename Dtype,typename GPUConvLayer,typename CPUConvLayer>
-void ACLConvolutionLayer<Dtype,GPUConvLayer,CPUConvLayer>::SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
- const vector<Blob<Dtype>*>& top){
-
- TensorShape input_shape((unsigned int)bottom[0]->width(), (unsigned int)bottom[0]->height(),(unsigned int)bottom[0]->channels(),(unsigned int)bottom[0]->num());
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::checkreshape(input_shape,Caffe::arm_gpu_mode());
- if (!this->init_layer_) return;
- this->init_layer_=false;
- // Initialize ACL.
- if (Caffe::arm_gpu_mode()) {
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::new_gpulayer();
- }else{
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::new_cpulayer();
- }
- this->force_bypass_acl_path_=false;
- ConvolutionParameter conv_param = this->layer_param_.convolution_param();
- int stride_x =this->stride_.mutable_cpu_data()[1];
- int stride_y =this->stride_.mutable_cpu_data()[0];
- int pad_x=this->pad_.mutable_cpu_data()[1];
- int pad_y=this->pad_.mutable_cpu_data()[0];
- unsigned int kernel_x=this->kernel_shape_.mutable_cpu_data()[1];
- unsigned int kernel_y=this->kernel_shape_.mutable_cpu_data()[0];
- PadStrideInfo conv_info(stride_x,stride_y,pad_x,pad_y);
- TensorShape weights_shape(kernel_x,kernel_y,(unsigned int)this->channels_, (unsigned int)this->num_output_);
- TensorShape biases_shape ((unsigned int)this->num_output_);
- TensorShape output_shape((unsigned int)top[0]->width(), (unsigned int)top[0]->height(),(unsigned int)top[0]->channels(),(unsigned int)top[0]->num());
-
- if (Caffe::arm_gpu_mode()) {
- Dtype *top_data = top[0]->mutable_gpu_data();
- const Dtype* bottom_data = bottom[0]->gpu_data();
- //[kernel_x, kernel_y, IFM, OFM]
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::new_tensor(this->gpu().weights,weights_shape,(void*)(this->blobs_[0].get()->mutable_gpu_data()));
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::tensor_mem(this->gpu().weights,(void*)(this->blobs_[0].get()->mutable_gpu_data()));
- //[OFM]
- if (this->bias_term_) {
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::new_tensor(this->gpu().biases,biases_shape,(void*)(this->blobs_[1].get()->mutable_gpu_data()));
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::tensor_mem(this->gpu().biases,(void*)(this->blobs_[1].get()->mutable_gpu_data()));
- }
-
- //[width, height, IFM]
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::new_tensor(this->gpu().input,input_shape,(void*)bottom_data);
- //[width, height, OFM]
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::new_tensor(this->gpu().output,output_shape,(void*)top_data);
-#ifdef USE_PROFILING
- {
- logtime_util log_time(ACL_CONFIG_INFO);
-#endif //USE_PROFILING
- this->gpu().layer->configure(this->gpu().input,this->gpu().weights,this->gpu().biases,this->gpu().output,conv_info);
-#ifdef USE_PROFILING
- }
-#endif //USE_PROFILING
-#ifdef USE_CONV_CACHE
- for(int i = 0; i < 16; ++i){
- fprintf(stderr, "<GPU>check cache[%d]\n", i);
- if(this->gpu().cache.layer[i] == nullptr){
- this->gpu().cache.layer[i] = this->gpu().layer;
- this->gpu().cache.input[i] = this->gpu().input;
- this->gpu().cache.output[i] = this->gpu().output;
- this->gpu().cache.weights[i] = this->gpu().weights;
- this->gpu().cache.biases[i] = this->gpu().biases;
- break;
- }
- }
-#endif //USE_CONV_CACHE
- }else{
- Dtype *top_data = top[0]->mutable_cpu_data();
- const Dtype* bottom_data = bottom[0]->cpu_data();
- //[kernel_x, kernel_y, IFM, OFM]
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::new_tensor(this->cpu().weights,weights_shape,(void*)(this->blobs_[0].get()->mutable_cpu_data()));
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::tensor_mem(this->cpu().weights,(void*)(this->blobs_[0].get()->mutable_cpu_data()));
- //[OFM]
- if (this->bias_term_) {
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::new_tensor(this->cpu().biases,biases_shape,(void*)(this->blobs_[1].get()->mutable_cpu_data()));
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::tensor_mem(this->cpu().biases,(void*)(this->blobs_[1].get()->mutable_cpu_data()));
- }
-
- //[width, height, IFM]
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::new_tensor(this->cpu().input,input_shape,(void*)bottom_data);
- //[width, height, OFM]
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::new_tensor(this->cpu().output,output_shape,(void*)top_data);
-#ifdef USE_PROFILING
- {
- logtime_util log_time(ACL_CONFIG_INFO);
-#endif //USE_PROFILING
- this->cpu().layer->configure(this->cpu().input,this->cpu().weights,this->cpu().biases,this->cpu().output,conv_info);
-#ifdef USE_PROFILING
- }
-#endif //USE_PROFILING
-#ifdef USE_CONV_CACHE
- for(int i = 0; i < 16; ++i){
- fprintf(stderr, "<CPU>check cache[%d]\n", i);
- if(this->cpu().cache.layer[i] == nullptr){
- this->cpu().cache.layer[i] = this->cpu().layer;
- this->cpu().cache.input[i] = this->cpu().input;
- this->cpu().cache.output[i] = this->cpu().output;
- this->cpu().cache.weights[i] = this->cpu().weights;
- this->cpu().cache.biases[i] = this->cpu().biases;
- break;
- }
- }
-#endif //USE_CONV_CACHE
- }
-}
-template <typename Dtype,typename GPUConvLayer,typename CPUConvLayer>
-void ACLConvolutionLayer<Dtype,GPUConvLayer,CPUConvLayer>::Reshape(
- const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) {
- ConvolutionLayer<Dtype>::Reshape(bottom, top);
-}
-
-template <typename Dtype,typename GPUConvLayer,typename CPUConvLayer>
-void ACLConvolutionLayer<Dtype,GPUConvLayer,CPUConvLayer>::Forward_cpu(
- const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) {
- if(Caffe::arm_gpu_mode()){
- Forward_gpu(bottom, top);
- return;
- }
-#ifdef USE_PROFILING
- logtime_util log_time(ACL_CONV_INFO);
-#endif //USE_PROFILING
- if (this->force_bypass_acl_path_|| this->group_!=1) {
- ConvolutionLayer<Dtype>::Forward_cpu(bottom,top);
- return;
- }
-
- ConvolutionParameter conv_param = this->layer_param_.convolution_param();
- if (conv_param.kernel_size_size()>2 || this->num_spatial_axes_>2 || this->num_spatial_axes_==0) {
- ConvolutionLayer<Dtype>::Forward_cpu(bottom,top);
- return;
- }
- /* check dilation */
- int dilated=0;
-
- for(int i=0;i<this->num_spatial_axes_;i++)
- {
- const int *p=this->dilation_.cpu_data();
-
- if(p[i]!=1)
- dilated=1;
- }
- if(dilated) {
- ConvolutionLayer<Dtype>::Forward_cpu(bottom,top);
- return;
- }
-
- SetupACLLayer(bottom,top);
- for (int i = 0; i < bottom.size(); ++i) {
- const Dtype* bottom_data = bottom[i]->cpu_data();
- Dtype* top_data = top[i]->mutable_cpu_data();
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::tensor_mem(this->cpu().input,(void*)bottom_data);
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::cpu_run();
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::tensor_mem((void*)top_data,this->cpu().output);
- }
-}
-
-template <typename Dtype,typename GPUConvLayer,typename CPUConvLayer>
-void ACLConvolutionLayer<Dtype,GPUConvLayer,CPUConvLayer>::Forward_gpu(
- const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) {
-#ifdef USE_PROFILING
- logtime_util log_time(ACL_CONV_INFO);
-#endif //USE_PROFILING
- ConvolutionParameter conv_param = this->layer_param_.convolution_param();
- if (this->force_bypass_acl_path_|| this->group_!=1) {
- ConvolutionLayer<Dtype>::Forward_cpu(bottom,top);
- return;
- }
- if (conv_param.kernel_size_size()>2 || this->num_spatial_axes_>2 ) {
- ConvolutionLayer<Dtype>::Forward_cpu(bottom,top);
- return;
- }
- /* check dilation */
- int dilated=0;
-
- for(int i=0;i<this->num_spatial_axes_;i++)
- {
- const int *p=this->dilation_.gpu_data();
-
- if(p[i]!=1)
- dilated=1;
- }
-
- if(dilated) {
- ConvolutionLayer<Dtype>::Forward_cpu(bottom,top);
- return;
- }
- SetupACLLayer(bottom,top);
- for (int i = 0; i < bottom.size(); ++i) {
- const Dtype* bottom_data = bottom[i]->gpu_data();
- Dtype* top_data = top[i]->mutable_gpu_data();
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::tensor_mem(this->gpu().input,(void*)bottom_data);
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::gpu_run();
- ACLBaseLayer<GPUConvLayer,CPUConvLayer>::tensor_mem((void*)top_data,this->gpu().output);
- }
-}
-
-template <typename Dtype,typename GPUConvLayer,typename CPUConvLayer>
-ACLConvolutionLayer<Dtype,GPUConvLayer,CPUConvLayer>::~ACLConvolutionLayer() {
-}
-
-#ifdef USE_ACL
-INSTANTIATE_CONV_CLASS(ACLConvolutionLayer,CLConvolutionLayer,NEDirectConvolutionLayer);
-INSTANTIATE_CONV_CLASS(ACLConvolutionLayer,CLConvolutionLayer,NEConvolutionLayer);
-#endif
-
-} // namespace caffe
-#endif // USE_ACL
this->force_bypass_acl_path_= bypass_acl_class_layer & FLAGS_ENABLE_ACL_BN;
}
template <typename Dtype>
-void ACLBatchNormLayer<Dtype>::SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
+void ACLBatchNormLayer<Dtype>::SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top){
- if (!this->init_layer_) return;
- this->init_layer_=false;
- // Initialize ACL.
- if (Caffe::arm_gpu_mode()) {
- new_gpulayer();
- }else{
- new_cpulayer();
- }
-
- this->force_bypass_acl_path_=false;
+ arm_compute::TensorShape in_shape ((unsigned int)bottom[0]->width(), (unsigned int)bottom[0]->height(),(unsigned int)bottom[0]->channels(),(unsigned int)bottom[0]->num());
+ if (is_operator_init_done(in_shape)) return;
+ set_operator_init_done();
- TensorShape in_shape ((unsigned int)bottom[0]->width(), (unsigned int)bottom[0]->height(),(unsigned int)bottom[0]->channels(),(unsigned int)bottom[0]->num());
- TensorShape out_shape((unsigned int)top[0]->width(), (unsigned int)top[0]->height(),(unsigned int)top[0]->channels(),(unsigned int)top[0]->num());
- TensorShape mean_shape((unsigned int)this->channels_);
- TensorShape var_shape=mean_shape;
- TensorShape beta_shape=mean_shape;
- TensorShape gamma_shape=mean_shape;
+ // Initialize ACL.
+ arm_compute::TensorShape out_shape((unsigned int)top[0]->width(), (unsigned int)top[0]->height(),(unsigned int)top[0]->channels(),(unsigned int)top[0]->num());
+ arm_compute::TensorShape mean_shape((unsigned int)this->channels_);
+ arm_compute::TensorShape var_shape=mean_shape;
+ arm_compute::TensorShape beta_shape=mean_shape;
+ arm_compute::TensorShape gamma_shape=mean_shape;
Dtype beta_val[beta_shape.total_size()];
Dtype gamma_val[gamma_shape.total_size()];
-
for (int i=0;i<beta_shape.total_size();++i) {
beta_val[i]=0.0;
}
for (int i=0;i<gamma_shape.total_size();++i) {
gamma_val[i]=1.0;
}
- if (Caffe::arm_gpu_mode()) {
- Dtype *top_data = top[0]->mutable_gpu_data();
- const Dtype* bottom_data = bottom[0]->gpu_data();
- // use the stored mean/variance estimates.
- const Dtype scale_factor = this->blobs_[2]->cpu_data()[0] == 0 ?
- 0 : 1 / this->blobs_[2]->cpu_data()[0];
- caffe_cpu_scale(this->variance_.count(), scale_factor,
- this->blobs_[0]->gpu_data(), this->mean_.mutable_gpu_data());
- caffe_cpu_scale(this->variance_.count(), scale_factor,
- this->blobs_[1]->gpu_data(), this->variance_.mutable_gpu_data());
- new_tensor(this->gpu().input,in_shape,(void*)bottom_data);
- new_tensor(this->gpu().output,out_shape,(void*)top_data);
- new_tensor(this->gpu().mean,mean_shape);
- new_tensor(this->gpu().var,var_shape);
- new_tensor(this->gpu().beta,beta_shape);
- new_tensor(this->gpu().gamma,gamma_shape);
- tensor_mem(this->gpu().mean,(void*)this->mean_.mutable_gpu_data());
- tensor_mem(this->gpu().var,(void*)this->variance_.mutable_gpu_data());
- tensor_mem(this->gpu().beta,(void*)beta_val);
- tensor_mem(this->gpu().gamma,(void*)gamma_val);
- this->gpu().mean->commit();
- this->gpu().var->commit();
- this->gpu().beta->commit();
- this->gpu().gamma->commit();
-#ifdef USE_PROFILING
- logtime_util log_time(ACL_CONFIG_INFO);
-#endif //USE_PROFILING
- this->gpu().layer->configure(this->gpu().input,this->gpu().output,this->gpu().mean,this->gpu().var,this->gpu().beta,this->gpu().gamma,this->eps_);
- }else{
- Dtype *top_data = top[0]->mutable_cpu_data();
- const Dtype* bottom_data = bottom[0]->cpu_data();
- // use the stored mean/variance estimates.
- const Dtype scale_factor = this->blobs_[2]->cpu_data()[0] == 0 ?
- 0 : 1 / this->blobs_[2]->cpu_data()[0];
- caffe_cpu_scale(this->variance_.count(), scale_factor,
- this->blobs_[0]->cpu_data(), this->mean_.mutable_cpu_data());
- caffe_cpu_scale(this->variance_.count(), scale_factor,
- this->blobs_[1]->cpu_data(), this->variance_.mutable_cpu_data());
- new_tensor(this->cpu().input,in_shape,(void*)bottom_data);
- new_tensor(this->cpu().output,out_shape,(void*)top_data);
- new_tensor(this->cpu().mean,mean_shape);
- new_tensor(this->cpu().var,var_shape);
- new_tensor(this->cpu().beta,beta_shape);
- new_tensor(this->cpu().gamma,gamma_shape);
- tensor_mem(this->cpu().mean,(void*)this->mean_.mutable_cpu_data());
- tensor_mem(this->cpu().var,(void*)this->variance_.mutable_cpu_data());
- tensor_mem(this->cpu().beta,(void*)beta_val);
- tensor_mem(this->cpu().gamma,(void*)gamma_val);
- this->cpu().mean->commit();
- this->cpu().var->commit();
- this->cpu().beta->commit();
- this->cpu().gamma->commit();
+ new_tensor(input(),in_shape,InputdataPtr(this,bottom));
+ new_tensor(output(),out_shape,OutputdataPtr(this,top));
+ // use the stored mean/variance estimates.
+ const Dtype scale_factor = this->blobs_[2]->cpu_data()[0] == 0 ?
+ 0 : 1 / this->blobs_[2]->cpu_data()[0];
+ caffe_cpu_scale(this->variance_.count(), scale_factor,
+ this->blobs_[0]->cpu_data(), GetDataPtr(this,&this->mean_));
+ caffe_cpu_scale(this->variance_.count(), scale_factor,
+ this->blobs_[1]->cpu_data(), GetDataPtr(this,&this->variance_));
-#ifdef USE_PROFILING
- logtime_util log_time(ACL_CONFIG_INFO);
-#endif //USE_PROFILING
- this->cpu().layer->configure(this->cpu().input,this->cpu().output,this->cpu().mean,this->cpu().var,this->cpu().beta,this->cpu().gamma,this->eps_);
- }
+ new_tensor(mean(),mean_shape,GetDataPtr(this,&this->mean_));
+ new_tensor(var(),var_shape,GetDataPtr(this,&this->variance_));
+ new_tensor(beta(),beta_shape,(void*)beta_val,true);
+ new_tensor(gamma(),gamma_shape,(void*)gamma_val,true);
+ acl_configure(bn,this,this->eps_);
}
template <typename Dtype>
void ACLBatchNormLayer<Dtype>::Reshape(const vector<Blob<Dtype>*>& bottom,
}
template <typename Dtype>
+bool ACLBatchNormLayer<Dtype>::Bypass_acl(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top){
+ bool bypass_acl=false;
+ if (this->force_bypass_acl_path_||!this->use_global_stats_) {
+ bypass_acl=true;
+ }
+ if (isScheduleEnable()) {
+ bypass_acl=true;
+ }
+ return bypass_acl;
+}
+
+template <typename Dtype>
void ACLBatchNormLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top) {
- if(Caffe::arm_gpu_mode()){
+ if(isGPUMode()){
Forward_gpu(bottom, top);
return;
}
#ifdef USE_PROFILING
logtime_util log_time(ACL_BN_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_||!this->use_global_stats_) {
+ if (Bypass_acl(bottom,top)) {
BatchNormLayer<Dtype>::Forward_cpu(bottom,top);
return;
}
- const Dtype* bottom_data = bottom[0]->cpu_data();
- Dtype* top_data = top[0]->mutable_cpu_data();
- SetupACLLayer(bottom,top);
- tensor_mem(this->cpu().input,(void*)(bottom_data));
- cpu_run();
- tensor_mem((void*)(top_data),this->cpu().output);
+ SetupACLOperator(bottom,top);
+ caffe::acl_run(this,bottom,top);
}
template <typename Dtype>
void ACLBatchNormLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top) {
- if (this->force_bypass_acl_path_||!this->use_global_stats_) {
- BatchNormLayer<Dtype>::Forward_cpu(bottom,top);
- return;
- }
#ifdef USE_PROFILING
logtime_util log_time(ACL_BN_INFO);
#endif //USE_PROFILING
- const Dtype* bottom_data = bottom[0]->gpu_data();
- Dtype* top_data = top[0]->mutable_gpu_data();
- SetupACLLayer(bottom,top);
- tensor_mem(this->gpu().input,(void*)(bottom_data));
- gpu_run();
- tensor_mem((void*)(top_data),this->gpu().output);
+ if (Bypass_acl(bottom,top)) {
+ BatchNormLayer<Dtype>::Forward_cpu(bottom,top);
+ return;
+ }
+ SetupACLOperator(bottom,top);
+ caffe::acl_run(this,bottom,top);
}
template <typename Dtype>
this->force_bypass_acl_path_= bypass_acl_class_layer & FLAGS_ENABLE_ACL_BNLL;
}
template <typename Dtype>
-void ACLBNLLLayer<Dtype>::SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
- const vector<Blob<Dtype>*>& top, ActivationLayerInfo::ActivationFunction type){
- ACLBaseActivationLayer<Dtype>::SetupACLLayer(bottom, top,ActivationLayerInfo::ActivationFunction::SOFT_RELU);
+void ACLBNLLLayer<Dtype>::SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top, arm_compute::ActivationLayerInfo::ActivationFunction type){
+ ACLBaseActivationLayer<Dtype>::SetupACLOperator(bottom, top,arm_compute::ActivationLayerInfo::ActivationFunction::SOFT_RELU);
}
template <typename Dtype>
void ACLBNLLLayer<Dtype>::Reshape(
}
template <typename Dtype>
+bool ACLBNLLLayer<Dtype>::Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top){
+ bool bypass_acl=false;
+ if (this->force_bypass_acl_path_) {
+ bypass_acl=true;
+ }
+ return bypass_acl;
+}
+
+template <typename Dtype>
void ACLBNLLLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top) {
#ifdef USE_PROFILING
logtime_util log_time(ACL_BNLL_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_) {
+ if (Bypass_acl(bottom,top)) {
BNLLLayer<Dtype>::Forward_cpu(bottom,top);
return;
}
#ifdef USE_PROFILING
logtime_util log_time(ACL_BNLL_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_) {
+ if (Bypass_acl(bottom,top)) {
BNLLLayer<Dtype>::Forward_cpu(bottom,top);
return;
}
void ACLConcatLayer<Dtype>::LayerSetUp(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top) {
ConcatLayer<Dtype>::LayerSetUp(bottom, top);
- //this->force_bypass_acl_path_= bypass_acl_class_layer & FLAGS_ENABLE_ACL_CONCAT;
- this->force_bypass_acl_path_= true;
+ this->force_bypass_acl_path_= bypass_acl_class_layer & FLAGS_ENABLE_ACL_CONCAT;
}
template <typename Dtype>
-void ACLConcatLayer<Dtype>::SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
+void ACLConcatLayer<Dtype>::SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top){
unsigned int channels=0;
for (int i = 0; i < bottom.size(); ++i) {
channels+=bottom[i]->channels();
}
- TensorShape out_shape((unsigned int)top[0]->width(), (unsigned int)top[0]->height(),channels);
+ arm_compute::TensorShape out_shape((unsigned int)top[0]->width(), (unsigned int)top[0]->height(),channels);
+
+ if (is_operator_init_done(out_shape,tensor_output)) return;
+ set_operator_init_done();
- if (!this->init_layer_) return;
- this->init_layer_=false;
// Initialize ACL.
- if (Caffe::arm_gpu_mode()) {
- new_gpulayer();
- }else{
- new_cpulayer();
+ std::vector<arm_compute::TensorShape> shapes;
+ for (int i = 0; i < bottom.size(); ++i) {
+ arm_compute::TensorShape in_shape((unsigned int)bottom[i]->width(), (unsigned int)bottom[i]->height(),(unsigned int)bottom[i]->channels());
+ new_tensor(cinput(i),in_shape,InputdataPtr(this,bottom,i));
}
+ new_tensor(output(),out_shape,OutputdataPtr(this,top));
+ acl_configure(concat,this,bottom.size());
- this->force_bypass_acl_path_=false;
-
- if (Caffe::arm_gpu_mode()) {
- Dtype *top_data = top[0]->mutable_gpu_data();
- for (int i = 0; i < bottom.size(); ++i) {
- const Dtype* bottom_data = bottom[i]->gpu_data();
- TensorShape vec_shape((unsigned int)bottom[i]->width(), (unsigned int)bottom[i]->height(),(unsigned int)bottom[0]->channels());
- GPUTensor *vector;
- new_tensor(vector,vec_shape,(void*)bottom_data);
- tensor_mem(vector,(void*)bottom_data);
- vector->commit();
- gpu_vectors.push_back(vector);
- }
- new_tensor(this->gpu().output,out_shape,(void*)top_data);
-#ifdef USE_PROFILING
- logtime_util log_time(ACL_CONFIG_INFO);
-#endif //USE_PROFILING
- this->gpu().layer->configure(gpu_vectors,this->gpu().output);
- }else{
- Dtype *top_data = top[0]->mutable_cpu_data();
- for (int i = 0; i < bottom.size(); ++i) {
- const Dtype* bottom_data = bottom[i]->cpu_data();
- TensorShape vec_shape((unsigned int)bottom[i]->width(), (unsigned int)bottom[i]->height(),(unsigned int)bottom[0]->channels());
- CPUTensor *vector;
- new_tensor(vector,vec_shape,(void*)bottom_data);
- tensor_mem(vector,(void*)bottom_data);
- vector->commit();
- cpu_vectors.push_back(vector);
- }
- new_tensor(this->cpu().output,out_shape,(void*)top_data);
-#ifdef USE_PROFILING
- logtime_util log_time(ACL_CONFIG_INFO);
-#endif //USE_PROFILING
- this->cpu().layer->configure(cpu_vectors,this->cpu().output);
- }
}
template <typename Dtype>
void ACLConcatLayer<Dtype>::Reshape(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top) {
- ConcatLayer<Dtype>::Reshape(bottom, top);
+ ConcatLayer<Dtype>::Reshape(bottom, top);
+}
+template <typename Dtype>
+bool ACLConcatLayer<Dtype>::Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top){
+ bool bypass_acl=false;
+ if (this->force_bypass_acl_path_||this->concat_axis_==0) {
+ bypass_acl=true;
+ }
+ if(isScheduleEnable()){
+ bypass_acl=true;
+ }
+ return bypass_acl;
}
template <typename Dtype>
void ACLConcatLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top) {
- if(Caffe::arm_gpu_mode()){
+ if(isGPUMode()){
Forward_gpu(bottom, top);
return;
}
#ifdef USE_PROFILING
logtime_util log_time(ACL_CONCAT_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_||this->concat_axis_==0) {
+ if (Bypass_acl(bottom,top)) {
ConcatLayer<Dtype>::Forward_cpu(bottom,top);
return;
}
- Dtype* top_data = top[0]->mutable_cpu_data();
- SetupACLLayer(bottom,top);
- cpu_run();
- tensor_mem((void*)(top_data),this->cpu().output);
+ SetupACLOperator(bottom,top);
+ caffe::acl_run(this,bottom,top,false);
}
template <typename Dtype>
#ifdef USE_PROFILING
logtime_util log_time(ACL_CONCAT_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_||this->concat_axis_==0) {
+ if (Bypass_acl(bottom,top)) {
ConcatLayer<Dtype>::Forward_cpu(bottom,top);
return;
}
- Dtype* top_data = top[0]->mutable_gpu_data();
- SetupACLLayer(bottom,top);
- gpu_run();
- tensor_mem((void*)(top_data),this->gpu().output);
+ SetupACLOperator(bottom,top);
+ caffe::acl_run(this,bottom,top,false);
}
template <typename Dtype>
ACLConcatLayer<Dtype>::~ACLConcatLayer() {
- if(this->force_bypass_acl_path_)return;
- for (int i =0; i < cpu_vectors.size(); i ++) {
- delete cpu_vectors[i];
- }
- for (int i =0; i < gpu_vectors.size(); i ++) {
- delete gpu_vectors[i];
- }
- cpu_vectors.erase(cpu_vectors.begin());
- gpu_vectors.erase(gpu_vectors.begin());
}
INSTANTIATE_CLASS(ACLConcatLayer);
--- /dev/null
+#ifdef USE_ACL
+#include <algorithm>
+#include <vector>
+#include "caffe/filler.hpp"
+#include "caffe/layers/acl_conv_layer.hpp"
+
+namespace caffe {
+
+bool use_direct_conv_=false;
+template <typename Dtype>
+void ACLConvolutionLayer<Dtype>::LayerSetUp(
+ const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) {
+ ConvolutionLayer<Dtype>::LayerSetUp(bottom, top);
+ this->force_bypass_acl_path_= bypass_acl_class_layer & FLAGS_ENABLE_ACL_CONV;
+}
+
+template <typename Dtype>
+void ACLConvolutionLayer<Dtype>::SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top){
+
+ arm_compute::TensorShape input_shape((unsigned int)bottom[0]->width(), (unsigned int)bottom[0]->height(),(unsigned int)bottom[0]->channels(),(unsigned int)bottom[0]->num());
+ if (is_operator_init_done(input_shape)) return;
+ set_operator_init_done();
+
+ // Initialize ACL.
+ ConvolutionParameter conv_param = this->layer_param_.convolution_param();
+ int stride_x =this->stride_.mutable_cpu_data()[1];
+ int stride_y =this->stride_.mutable_cpu_data()[0];
+ int pad_x=this->pad_.mutable_cpu_data()[1];
+ int pad_y=this->pad_.mutable_cpu_data()[0];
+ unsigned int kernel_x=this->kernel_shape_.mutable_cpu_data()[1];
+ unsigned int kernel_y=this->kernel_shape_.mutable_cpu_data()[0];
+ arm_compute::PadStrideInfo conv_info(stride_x,stride_y,pad_x,pad_y);
+ arm_compute::TensorShape weights_shape(kernel_x,kernel_y,(unsigned int)this->channels_/this->group_, (unsigned int)this->num_output_);
+ arm_compute::TensorShape biases_shape ((unsigned int)this->num_output_);
+ arm_compute::TensorShape output_shape((unsigned int)top[0]->width(), (unsigned int)top[0]->height(),(unsigned int)top[0]->channels(),(unsigned int)top[0]->num());
+ group()=this->group_;
+
+ //[kernel_x, kernel_y, IFM, OFM]
+ new_tensor(weights(),weights_shape,GetDataPtr(this,this->blobs_[0].get()));
+ //[OFM]
+ if (this->bias_term_) {
+ new_tensor(biases(),biases_shape,GetDataPtr(this,this->blobs_[1].get()));
+ }
+
+ //[width, height, IFM]
+ new_tensor(input(),input_shape,InputdataPtr(this,bottom));
+ //[width, height, OFM]
+ new_tensor(output(),output_shape,OutputdataPtr(this,top));
+
+ acl_configure(conv,this,conv_info);
+}
+template <typename Dtype>
+void ACLConvolutionLayer<Dtype>::Reshape(
+ const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) {
+ ConvolutionLayer<Dtype>::Reshape(bottom, top);
+}
+
+template <typename Dtype>
+bool ACLConvolutionLayer<Dtype>::Bypass_acl(
+ const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) {
+ bool bypass_acl=false;
+ if (this->force_bypass_acl_path_|| ((openailab_intfp==0) && (this->group_>=5)) //for performance, more groups impact GPU performance
+ || ((openailab_intfp != 0 && (top[0]->channels() / this->group_ == 1)))) {
+ bypass_acl=true;
+ }
+
+ ConvolutionParameter conv_param = this->layer_param_.convolution_param();
+ if (conv_param.kernel_size_size()>2 || this->num_spatial_axes_>2 || this->num_spatial_axes_==0) {
+ bypass_acl=true;
+ }
+ /* check dilation */
+ int dilated=0;
+
+ for(int i=0;i<this->num_spatial_axes_;i++)
+ {
+ const int *p=this->dilation_.cpu_data();
+
+ if(p[i]!=1)
+ dilated=1;
+ }
+ if(dilated) {
+ bypass_acl=true;
+ }
+
+
+ if((this->kernel_shape_.mutable_cpu_data()[1]==1||this->kernel_shape_.mutable_cpu_data()[0]==1) &&
+ isScheduleEnable()){
+ bypass_acl=true;
+ }
+ if((this->kernel_shape_.mutable_cpu_data()[1]==3||this->kernel_shape_.mutable_cpu_data()[0]==3) &&
+ (bottom[0]->channels()<150) && isScheduleEnable()){
+ bypass_acl=true;
+ }
+
+ return bypass_acl;
+}
+
+template <typename Dtype>
+void ACLConvolutionLayer<Dtype>::Forward_cpu(
+ const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) {
+ if(isGPUMode()){
+ Forward_gpu(bottom, top);
+ return;
+ }
+#ifdef USE_PROFILING
+ logtime_util log_time(ACL_CONV_INFO);
+#endif //USE_PROFILING
+
+ if (Bypass_acl(bottom,top)) {
+ ConvolutionLayer<Dtype>::Forward_cpu(bottom,top);
+ return;
+ }
+
+ SetupACLOperator(bottom,top);
+
+ // acl fp
+ if (openailab_intfp==0){
+ caffe::acl_run(this,bottom,top);
+ }
+ return;
+}
+
+template <typename Dtype>
+void ACLConvolutionLayer<Dtype>::Forward_gpu(
+ const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) {
+#ifdef USE_PROFILING
+ logtime_util log_time(ACL_CONV_INFO);
+#endif //USE_PROFILING
+ if (Bypass_acl(bottom,top)) {
+ ConvolutionLayer<Dtype>::Forward_cpu(bottom,top);
+ return;
+ }
+ SetupACLOperator(bottom,top);
+ caffe::acl_run(this,bottom,top);
+}
+
+template <typename Dtype>
+ACLConvolutionLayer<Dtype>::~ACLConvolutionLayer() {
+}
+
+#ifdef USE_ACL
+INSTANTIATE_CLASS(ACLConvolutionLayer);
+#endif
+
+} // namespace caffe
+#endif // USE_ACL
#ifdef USE_ACL
#include <vector>
-
#include "caffe/filler.hpp"
#include "caffe/layers/acl_inner_product_layer.hpp"
#include "caffe/util/math_functions.hpp"
this->force_bypass_acl_path_= bypass_acl_class_layer & FLAGS_ENABLE_ACL_FC;
}
template <typename Dtype>
-void ACLInnerProductLayer<Dtype>::SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
+void ACLInnerProductLayer<Dtype>::SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top){
- TensorShape weights_shape_t((unsigned int)this->K_, (unsigned int)this->N_);
- TensorShape weights_shape((unsigned int)this->N_, (unsigned int)this->K_);
- TensorShape biases_shape((unsigned int)this->N_);
- TensorShape input_shape((unsigned int)this->K_, (unsigned int)this->M_);
- TensorShape output_shape((unsigned int)this->N_, (unsigned int)this->M_);
- checkreshape(input_shape,Caffe::arm_gpu_mode());
- if (!this->init_layer_) return;
- this->init_layer_=false;
+ arm_compute::TensorShape weights_shape_t((unsigned int)this->K_, (unsigned int)this->N_);
+ arm_compute::TensorShape weights_shape((unsigned int)this->N_, (unsigned int)this->K_);
+ arm_compute::TensorShape biases_shape((unsigned int)this->N_);
+ arm_compute::TensorShape input_shape((unsigned int)this->K_, (unsigned int)this->M_);
+ arm_compute::TensorShape output_shape((unsigned int)this->N_, (unsigned int)this->M_);
+ if (is_operator_init_done(input_shape)) return;
+ set_operator_init_done();
// Initialize ACL.
- if (Caffe::arm_gpu_mode()) {
- new_gpulayer();
- }else{
- new_cpulayer();
- }
bool transpose = !this->layer_param_.inner_product_param().transpose();
- this->force_bypass_acl_path_ = false;
- if (Caffe::arm_gpu_mode()) {
- Dtype *top_data = top[0]->mutable_gpu_data();
- const Dtype* bottom_data = bottom[0]->gpu_data();
- if (transpose) {
- new_tensor(this->gpu().weights,weights_shape_t,(void*)(this->blobs_[0].get()->mutable_gpu_data()));
- }else{
- new_tensor(this->gpu().weights,weights_shape,(void*)(this->blobs_[0].get()->mutable_gpu_data()));
- }
- tensor_mem(this->gpu().weights,(void*)(this->blobs_[0].get()->mutable_gpu_data()));
- if (this->bias_term_) {
- new_tensor(this->gpu().biases,biases_shape,(void*)(this->blobs_[1].get()->mutable_gpu_data()));
- tensor_mem(this->gpu().biases,(void*)(this->blobs_[1].get()->mutable_gpu_data()));
- }
- new_tensor(this->gpu().input,input_shape,(void*)bottom_data);
- new_tensor(this->gpu().output,output_shape,(void*)top_data);
-#ifdef USE_PROFILING
- logtime_util log_time(ACL_CONFIG_INFO);
-#endif //USE_PROFILING
- this->gpu().layer->configure(this->gpu().input,this->gpu().weights,this->gpu().biases,this->gpu().output,transpose);
+ if (transpose) {
+ new_tensor(weights(),weights_shape_t,GetDataPtr(this,this->blobs_[0].get()));
}else{
- Dtype *top_data = top[0]->mutable_cpu_data();
- const Dtype* bottom_data = bottom[0]->cpu_data();
- if (transpose) {
- new_tensor(this->cpu().weights,weights_shape_t,(void*)(this->blobs_[0].get()->mutable_cpu_data()));
- }else{
- new_tensor(this->cpu().weights,weights_shape,(void*)(this->blobs_[0].get()->mutable_cpu_data()));
- }
- tensor_mem(this->cpu().weights,(void*)(this->blobs_[0].get()->mutable_cpu_data()));
- if (this->bias_term_) {
- new_tensor(this->cpu().biases,biases_shape,(void*)(this->blobs_[1].get()->mutable_cpu_data()));
- tensor_mem(this->cpu().biases,(void*)(this->blobs_[1].get()->mutable_cpu_data()));
- }
- new_tensor(this->cpu().input,input_shape,(void*)bottom_data);
- new_tensor(this->cpu().output,output_shape,(void*)top_data);
-#ifdef USE_PROFILING
- logtime_util log_time(ACL_CONFIG_INFO);
-#endif //USE_PROFILING
- this->cpu().layer->configure(this->cpu().input,this->cpu().weights,this->cpu().biases,this->cpu().output,transpose);
+ new_tensor(weights(),weights_shape,GetDataPtr(this,this->blobs_[0].get()));
+ }
+ if (this->bias_term_) {
+ new_tensor(biases(),biases_shape,GetDataPtr(this,this->blobs_[1].get()));
}
+ new_tensor(input(),input_shape,InputdataPtr(this,bottom));
+ new_tensor(output(),output_shape,OutputdataPtr(this,top));
+ acl_configure(fc,this,transpose);
}
template <typename Dtype>
void ACLInnerProductLayer<Dtype>::Reshape(const vector<Blob<Dtype>*>& bottom,
}
template <typename Dtype>
+bool ACLInnerProductLayer<Dtype>::Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top){
+ bool bypass_acl=false;
+ if (this->force_bypass_acl_path_) {
+ bypass_acl=true;
+ }
+ return bypass_acl;
+}
+
+template <typename Dtype>
void ACLInnerProductLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top) {
- if(Caffe::arm_gpu_mode()){
+ if(isGPUMode()){
Forward_gpu(bottom, top);
return;
}
#ifdef USE_PROFILING
logtime_util log_time(ACL_FC_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_) {
+ if (Bypass_acl(bottom, top)) {
InnerProductLayer<Dtype>::Forward_cpu(bottom,top);
return;
}
- Dtype* top_data = top[0]->mutable_cpu_data();
- const Dtype* bottom_data = bottom[0]->cpu_data();
- SetupACLLayer(bottom,top);
- tensor_mem(this->cpu().input,(void*)(bottom_data));
- cpu_run();
- tensor_mem((void*)(top_data),this->cpu().output);
+ SetupACLOperator(bottom,top);
+
+ if (this->M_ != 1 && openailab_intfp != 0){
+ InnerProductLayer<Dtype>::Forward_cpu(bottom,top);
+ return;
+ }
+
+ // ACL FP
+ if(openailab_intfp == 0){
+ caffe::acl_run(this,bottom,top);
+ }
+ return;
}
template <typename Dtype>
#ifdef USE_PROFILING
logtime_util log_time(ACL_FC_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_) {
+ if (Bypass_acl(bottom, top)) {
InnerProductLayer<Dtype>::Forward_cpu(bottom,top);
return;
}
- Dtype* top_data = top[0]->mutable_gpu_data();
- const Dtype* bottom_data = bottom[0]->gpu_data();
- SetupACLLayer(bottom,top);
- tensor_mem(this->gpu().input,(void*)(bottom_data));
- gpu_run();
- tensor_mem((void*)(top_data),this->gpu().output);
+ SetupACLOperator(bottom,top);
+ caffe::acl_run(this,bottom,top);
}
template <typename Dtype>
}
template <typename Dtype>
-void ACLLocalConnectLayer<Dtype>::SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
+void ACLLocalConnectLayer<Dtype>::SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top){
- TensorShape input_shape((unsigned int)bottom[0]->width(), (unsigned int)bottom[0]->height(),(unsigned int)bottom[0]->channels(),(unsigned int)bottom[0]->num());
- checkreshape(input_shape,Caffe::arm_gpu_mode());
- if (!this->init_layer_) return;
- this->init_layer_=false;
- // Initialize ACL.
- if (Caffe::arm_gpu_mode()) {
- new_gpulayer();
- }else{
- new_cpulayer();
- }
- this->force_bypass_acl_path_=false;
+ arm_compute::TensorShape input_shape((unsigned int)bottom[0]->width(), (unsigned int)bottom[0]->height(),(unsigned int)bottom[0]->channels(),(unsigned int)bottom[0]->num());
+ if (is_operator_init_done(input_shape)) return;
+ set_operator_init_done();
+
+ // Initialize ACL.
ConvolutionParameter conv_param = this->layer_param_.convolution_param();
int stride_x =this->stride_;
int stride_y =this->stride_;
int pad_y=this->pad_;
unsigned int kernel_x=this->kernel_size_;
unsigned int kernel_y=this->kernel_size_;
- PadStrideInfo conv_info(stride_x,stride_y,pad_x,pad_y);
- TensorShape weights_shape(kernel_x,kernel_y,(unsigned int)this->channels_, (unsigned int)this->num_output_);
- TensorShape biases_shape ((unsigned int)this->num_output_);
- TensorShape output_shape((unsigned int)top[0]->width(), (unsigned int)top[0]->height(),(unsigned int)top[0]->channels(),(unsigned int)top[0]->num());
-
- if (Caffe::arm_gpu_mode()) {
- Dtype *top_data = top[0]->mutable_gpu_data();
- const Dtype* bottom_data = bottom[0]->gpu_data();
- //[kernel_x, kernel_y, IFM, OFM]
- new_tensor(this->gpu().weights,weights_shape,(void*)(this->blobs_[0].get()->mutable_gpu_data()));
- tensor_mem(this->gpu().weights,(void*)(this->blobs_[0].get()->mutable_gpu_data()));
- //[OFM]
- if (this->bias_term_) {
- new_tensor(this->gpu().biases,biases_shape,(void*)(this->blobs_[1].get()->mutable_gpu_data()));
- tensor_mem(this->gpu().biases,(void*)(this->blobs_[1].get()->mutable_gpu_data()));
- }
+ arm_compute::PadStrideInfo conv_info(stride_x,stride_y,pad_x,pad_y);
+ arm_compute::TensorShape weights_shape(kernel_x,kernel_y,(unsigned int)this->channels_, (unsigned int)this->num_output_);
+ arm_compute::TensorShape biases_shape ((unsigned int)this->num_output_);
+ arm_compute::TensorShape output_shape((unsigned int)top[0]->width(), (unsigned int)top[0]->height(),(unsigned int)top[0]->channels(),(unsigned int)top[0]->num());
- //[width, height, IFM]
- new_tensor(this->gpu().input,input_shape,(void*)bottom_data);
- //[width, height, OFM]
- new_tensor(this->gpu().output,output_shape,(void*)top_data);
-#ifdef USE_PROFILING
- {
- logtime_util log_time(ACL_CONFIG_INFO);
-#endif //USE_PROFILING
- this->gpu().layer->configure(this->gpu().input,this->gpu().weights,this->gpu().biases,this->gpu().output,conv_info);
-#ifdef USE_PROFILING
- }
-#endif //USE_PROFILING
-#ifdef USE_CONV_CACHE
- for(int i = 0; i < 16; ++i){
- fprintf(stderr, "<GPU>check cache[%d]\n", i);
- if(this->gpu().cache.layer[i] == nullptr){
- this->gpu().cache.layer[i] = this->gpu().layer;
- this->gpu().cache.input[i] = this->gpu().input;
- this->gpu().cache.output[i] = this->gpu().output;
- this->gpu().cache.weights[i] = this->gpu().weights;
- this->gpu().cache.biases[i] = this->gpu().biases;
- break;
- }
- }
-#endif //USE_CONV_CACHE
- }else{
- Dtype *top_data = top[0]->mutable_cpu_data();
- const Dtype* bottom_data = bottom[0]->cpu_data();
- //[kernel_x, kernel_y, IFM, OFM]
- new_tensor(this->cpu().weights,weights_shape,(void*)(this->blobs_[0].get()->mutable_cpu_data()));
- tensor_mem(this->cpu().weights,(void*)(this->blobs_[0].get()->mutable_cpu_data()));
- //[OFM]
- if (this->bias_term_) {
- new_tensor(this->cpu().biases,biases_shape,(void*)(this->blobs_[1].get()->mutable_cpu_data()));
- tensor_mem(this->cpu().biases,(void*)(this->blobs_[1].get()->mutable_cpu_data()));
- }
-
- //[width, height, IFM]
- new_tensor(this->cpu().input,input_shape,(void*)bottom_data);
- //[width, height, OFM]
- new_tensor(this->cpu().output,output_shape,(void*)top_data);
-#ifdef USE_PROFILING
- {
- logtime_util log_time(ACL_CONFIG_INFO);
-#endif //USE_PROFILING
- this->cpu().layer->configure(this->cpu().input,this->cpu().weights,this->cpu().biases,this->cpu().output,conv_info);
-#ifdef USE_PROFILING
- }
-#endif //USE_PROFILING
+ //[kernel_x, kernel_y, IFM, OFM]
+ new_tensor(weights(),weights_shape,GetDataPtr(this,this->blobs_[0].get()));
+ //[OFM]
+ if (this->bias_term_) {
+ new_tensor(biases(),biases_shape,GetDataPtr(this,this->blobs_[1].get()));
}
+
+ //[width, height, IFM]
+ new_tensor(input(),input_shape,InputdataPtr(this,bottom));
+ //[width, height, OFM]
+ new_tensor(output(),output_shape,OutputdataPtr(this,top));
+ acl_configure(lc,this,conv_info);
}
template <typename Dtype>
void ACLLocalConnectLayer<Dtype>::Reshape(
}
template <typename Dtype>
+bool ACLLocalConnectLayer<Dtype>::Bypass_acl(
+ const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top){
+ bool bypass_acl=false;
+ if (this->force_bypass_acl_path_) {
+ bypass_acl=true;
+ }
+
+ ConvolutionParameter conv_param = this->layer_param_.convolution_param();
+ if (conv_param.kernel_size_size()>2 ) {
+ bypass_acl=true;
+ }
+ return bypass_acl;
+}
+
+template <typename Dtype>
void ACLLocalConnectLayer<Dtype>::Forward_cpu(
const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) {
- if(Caffe::arm_gpu_mode()){
+ if(isGPUMode()){
Forward_gpu(bottom, top);
return;
}
#ifdef USE_PROFILING
logtime_util log_time(ACL_LC_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_) {
- LocalConnectLayer<Dtype>::Forward_cpu(bottom,top);
- return;
- }
-
- ConvolutionParameter conv_param = this->layer_param_.convolution_param();
- if (conv_param.kernel_size_size()>2 ) {
+ if (Bypass_acl(bottom,top)) {
LocalConnectLayer<Dtype>::Forward_cpu(bottom,top);
return;
}
- SetupACLLayer(bottom,top);
- for (int i = 0; i < bottom.size(); ++i) {
- const Dtype* bottom_data = bottom[i]->cpu_data();
- Dtype* top_data = top[i]->mutable_cpu_data();
- tensor_mem(this->cpu().input,(void*)bottom_data);
- cpu_run();
- tensor_mem((void*)top_data,this->cpu().output);
- }
+ SetupACLOperator(bottom,top);
+ caffe::acl_run(this,bottom,top);
}
template <typename Dtype>
logtime_util log_time(ACL_LC_INFO);
#endif //USE_PROFILING
ConvolutionParameter conv_param = this->layer_param_.convolution_param();
- if (this->force_bypass_acl_path_) {
- LocalConnectLayer<Dtype>::Forward_cpu(bottom,top);
- return;
- }
- if (conv_param.kernel_size_size()>2 ) {
+ if (Bypass_acl(bottom,top)) {
LocalConnectLayer<Dtype>::Forward_cpu(bottom,top);
return;
}
- SetupACLLayer(bottom,top);
- for (int i = 0; i < bottom.size(); ++i) {
- const Dtype* bottom_data = bottom[i]->gpu_data();
- Dtype* top_data = top[i]->mutable_gpu_data();
- tensor_mem(this->gpu().input,(void*)bottom_data);
- gpu_run();
- tensor_mem((void*)top_data,this->gpu().output);
- }
+ SetupACLOperator(bottom,top);
+ caffe::acl_run(this,bottom,top);
}
template <typename Dtype>
namespace caffe {
-const NormType IN_MAP=(arm_compute::NormType)0;
+const arm_compute::NormType IN_MAP=(arm_compute::NormType)0;
template <typename Dtype>
void ACLLRNLayer<Dtype>::LayerSetUp(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top) {
this->force_bypass_acl_path_= bypass_acl_class_layer & FLAGS_ENABLE_ACL_LRN;
}
template <typename Dtype>
-void ACLLRNLayer<Dtype>::SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
+void ACLLRNLayer<Dtype>::SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top){
- TensorShape shape((unsigned int)this->width_,(unsigned int)this->height_, (unsigned int)this->channels_);
- checkreshape(shape,Caffe::arm_gpu_mode());
- if (!this->init_layer_) return;
- // Initialize ACL.
- if (Caffe::arm_gpu_mode()) {
- new_gpulayer();
- }else{
- new_cpulayer();
- }
+ arm_compute::TensorShape shape((unsigned int)this->width_,(unsigned int)this->height_, (unsigned int)this->channels_);
+ if (is_operator_init_done(shape)) return;
+ set_operator_init_done();
- //this->force_bypass_acl_path_=false;
- NormalizationLayerInfo *norm_info;
+ // Initialize ACL.
+ arm_compute::NormalizationLayerInfo norm_info(IN_MAP, this->size_, this->alpha_, this->beta_, this->k_);
if(this->layer_param_.lrn_param().norm_region() == LRNParameter_NormRegion_WITHIN_CHANNEL)
- norm_info=new NormalizationLayerInfo(IN_MAP, this->size_, this->alpha_, this->beta_, this->k_);
+ norm_info=arm_compute::NormalizationLayerInfo(IN_MAP, this->size_, this->alpha_, this->beta_, this->k_);
else
- norm_info=new NormalizationLayerInfo(NormType::CROSS_MAP, this->size_, this->alpha_, this->beta_, this->k_);
+ norm_info=arm_compute::NormalizationLayerInfo(arm_compute::NormType::CROSS_MAP, this->size_, this->alpha_, this->beta_, this->k_);
+
+ new_tensor(input(),shape,InputdataPtr(this,bottom));
+ new_tensor(output(),shape,OutputdataPtr(this,top));
+ acl_configure(lrn,this,norm_info);
- if (Caffe::arm_gpu_mode()) {
- Dtype *top_data = top[0]->mutable_gpu_data();
- const Dtype* bottom_data = bottom[0]->gpu_data();
- new_tensor(this->gpu().input,shape,(void*)bottom_data);
- new_tensor(this->gpu().output,shape,(void*)top_data);
-#ifdef USE_PROFILING
- logtime_util log_time(ACL_CONFIG_INFO);
-#endif //USE_PROFILING
- this->gpu().layer->configure(this->gpu().input,this->gpu().output,*norm_info);
- }else{
- Dtype *top_data = top[0]->mutable_cpu_data();
- const Dtype* bottom_data = bottom[0]->cpu_data();
- new_tensor(this->cpu().input,shape,(void*)bottom_data);
- new_tensor(this->cpu().output,shape,(void*)top_data);
-#ifdef USE_PROFILING
- logtime_util log_time(ACL_CONFIG_INFO);
-#endif //USE_PROFILING
- this->cpu().layer->configure(this->cpu().input,this->cpu().output,*norm_info);
- }
- delete norm_info;
}
template <typename Dtype>
void ACLLRNLayer<Dtype>::Reshape(const vector<Blob<Dtype>*>& bottom,
}
template <typename Dtype>
+bool ACLLRNLayer<Dtype>::Bypass_acl(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top){
+ bool bypass_acl=false;
+ if (this->force_bypass_acl_path_ || this->layer_param_.lrn_param().norm_region() == LRNParameter_NormRegion_WITHIN_CHANNEL) {
+ bypass_acl=true;
+ }
+ return bypass_acl;
+}
+
+template <typename Dtype>
void ACLLRNLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top) {
- if(Caffe::arm_gpu_mode()){
+ if(isGPUMode()){
Forward_gpu(bottom, top);
return;
}
#ifdef USE_PROFILING
logtime_util log_time(ACL_LRN_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_ || this->layer_param_.lrn_param().norm_region() == LRNParameter_NormRegion_WITHIN_CHANNEL) {
+ if (Bypass_acl(bottom, top)) {
LRNLayer<Dtype>::Forward_cpu(bottom,top);
return;
}
const Dtype* bottom_data = bottom[0]->cpu_data();
Dtype* top_data = top[0]->mutable_cpu_data();
- SetupACLLayer(bottom,top);
+ SetupACLOperator(bottom,top);
switch (this->layer_param_.lrn_param().norm_region()) {
case LRNParameter_NormRegion_ACROSS_CHANNELS:
for (int n = 0; n < this->num_; ++n) {
- tensor_mem(this->cpu().input,(void*)(bottom_data+ bottom[0]->offset(n)));
- cpu_run();
- tensor_mem((void*)(top_data + top[0]->offset(n)),this->cpu().output);
+ acl_run((void*)(bottom_data+ bottom[0]->offset(n)),(void*)(top_data + top[0]->offset(n)));
}
break;
case LRNParameter_NormRegion_WITHIN_CHANNEL:
for (int n = 0; n < bottom[0]->num(); ++n) {
- tensor_mem(this->cpu().input,(void*)(bottom_data));
- cpu_run();
- tensor_mem((void*)(top_data),this->cpu().output);
+ acl_run((void*)bottom_data,(void*)top_data);
bottom_data += bottom[0]->offset(0, 1);
top_data += top[0]->offset(0, 1);
}
#ifdef USE_PROFILING
logtime_util log_time(ACL_LRN_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_) {
+ if (Bypass_acl(bottom, top)) {
LRNLayer<Dtype>::Forward_cpu(bottom,top);
return;
}
const Dtype* bottom_data = bottom[0]->gpu_data();
Dtype* top_data = top[0]->mutable_gpu_data();
- SetupACLLayer(bottom,top);
+ SetupACLOperator(bottom,top);
switch (this->layer_param_.lrn_param().norm_region()) {
case LRNParameter_NormRegion_ACROSS_CHANNELS:
for (int n = 0; n < this->num_; ++n) {
- tensor_mem(this->gpu().input,(void*)(bottom_data+ bottom[0]->offset(n)));
- gpu_run();
- tensor_mem((void*)(top_data + top[0]->offset(n)),this->gpu().output);
+ acl_run((void*)(bottom_data+ bottom[0]->offset(n)),(void*)(top_data + top[0]->offset(n)));
}
break;
case LRNParameter_NormRegion_WITHIN_CHANNEL:
for (int n = 0; n < bottom[0]->num(); ++n) {
- tensor_mem(this->gpu().input,(void*)(bottom_data));
- gpu_run();
- tensor_mem((void*)(top_data),this->gpu().output);
+ acl_run((void*)bottom_data,(void*)top_data);
bottom_data += bottom[0]->offset(0, 1);
top_data += top[0]->offset(0, 1);
}
this->force_bypass_acl_path_= bypass_acl_class_layer & FLAGS_ENABLE_ACL_POOLING;
}
template <typename Dtype>
-void ACLPoolingLayer<Dtype>::SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
+void ACLPoolingLayer<Dtype>::SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top){
- TensorShape in_shape ((unsigned int)this->width_, (unsigned int)this->height_,(unsigned int)this->channels_);
- TensorShape out_shape((unsigned int)this->pooled_width_, (unsigned int)this->pooled_height_,(unsigned int)this->channels_);
- checkreshape(in_shape,Caffe::arm_gpu_mode());
- if (!this->init_layer_) return;
- this->init_layer_=false;
- // Initialize ACL.
- if (Caffe::arm_gpu_mode()) {
- new_gpulayer();
- }else{
- new_cpulayer();
- }
+ arm_compute::TensorShape in_shape ((unsigned int)this->width_, (unsigned int)this->height_,(unsigned int)this->channels_);
+ arm_compute::TensorShape out_shape((unsigned int)this->pooled_width_, (unsigned int)this->pooled_height_,(unsigned int)this->channels_);
+ if (is_operator_init_done(in_shape)) return;
+ set_operator_init_done();
- this->force_bypass_acl_path_=false;
- PoolingLayerInfo *pool_info;
+ // Initialize ACL.
+ arm_compute::PoolingLayerInfo pool_info;
if(this->layer_param_.pooling_param().pool()==PoolingParameter_PoolMethod_MAX)
- pool_info=new PoolingLayerInfo(PoolingType::MAX, this->kernel_w_, PadStrideInfo(this->stride_w_,this->stride_h_,this->pad_w_,this->pad_h_,DimensionRoundingType::CEIL));
+ pool_info=arm_compute::PoolingLayerInfo(arm_compute::PoolingType::MAX, this->kernel_w_, arm_compute::PadStrideInfo(this->stride_w_,this->stride_h_,this->pad_w_,this->pad_h_,arm_compute::DimensionRoundingType::CEIL));
else
- pool_info=new PoolingLayerInfo(PoolingType::AVG, this->kernel_w_, PadStrideInfo(this->stride_w_,this->stride_h_,this->pad_w_,this->pad_h_,DimensionRoundingType::CEIL));
+ pool_info=arm_compute::PoolingLayerInfo(arm_compute::PoolingType::AVG, this->kernel_w_, arm_compute::PadStrideInfo(this->stride_w_,this->stride_h_,this->pad_w_,this->pad_h_,arm_compute::DimensionRoundingType::CEIL));
+
+ new_tensor(input(),in_shape,InputdataPtr(this,bottom));
+ new_tensor(output(),out_shape,OutputdataPtr(this,top));
+ acl_configure(pooling,this,pool_info);
- if (Caffe::arm_gpu_mode()) {
- Dtype *top_data = top[0]->mutable_gpu_data();
- const Dtype* bottom_data = bottom[0]->gpu_data();
- new_tensor(this->gpu().input,in_shape,(void*)bottom_data);
- new_tensor(this->gpu().output,out_shape,(void*)top_data);
-#ifdef USE_PROFILING
- logtime_util log_time(ACL_CONFIG_INFO);
-#endif //USE_PROFILING
- this->gpu().layer->configure(this->gpu().input,this->gpu().output,*pool_info);
- }else{
- Dtype *top_data = top[0]->mutable_cpu_data();
- const Dtype* bottom_data = bottom[0]->cpu_data();
- new_tensor(this->cpu().input,in_shape,(void*)bottom_data);
- new_tensor(this->cpu().output,out_shape,(void*)top_data);
-#ifdef USE_PROFILING
- logtime_util log_time(ACL_CONFIG_INFO);
-#endif //USE_PROFILING
- this->cpu().layer->configure(this->cpu().input,this->cpu().output,*pool_info);
- }
- delete pool_info;
}
template <typename Dtype>
void ACLPoolingLayer<Dtype>::Reshape(const vector<Blob<Dtype>*>& bottom,
}
template <typename Dtype>
+bool ACLPoolingLayer<Dtype>::Bypass_acl(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top){
+ bool bypass_acl=false;
+ if (this->force_bypass_acl_path_|| this->layer_param_.pooling_param().global_pooling()) {
+ bypass_acl=true;
+ }
+ if (this->layer_param_.pooling_param().pool()!=PoolingParameter_PoolMethod_MAX &&
+ this->layer_param_.pooling_param().pool()!=PoolingParameter_PoolMethod_AVE) {
+ bypass_acl=true;
+ }
+ if (this->kernel_h_!=this->kernel_w_) {
+ bypass_acl=true;
+ }
+ if (this->kernel_h_!=2 && this->kernel_h_!=3) {
+ bypass_acl=true;
+ }
+ return bypass_acl;
+}
+
+template <typename Dtype>
void ACLPoolingLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top) {
- if(Caffe::arm_gpu_mode()){
+ if(isGPUMode()){
Forward_gpu(bottom, top);
return;
}
#ifdef USE_PROFILING
logtime_util log_time(ACL_POOLING_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_|| this->layer_param_.pooling_param().global_pooling()) {
+ if (Bypass_acl(bottom,top)) {
PoolingLayer<Dtype>::Forward_cpu(bottom,top);
return;
}
const Dtype* bottom_data = bottom[0]->cpu_data();
Dtype* top_data = top[0]->mutable_cpu_data();
- if (this->layer_param_.pooling_param().pool()!=PoolingParameter_PoolMethod_MAX &&
- this->layer_param_.pooling_param().pool()!=PoolingParameter_PoolMethod_AVE) {
- PoolingLayer<Dtype>::Forward_cpu(bottom,top);
- return ;
- }
- if (this->kernel_h_!=this->kernel_w_ || top.size()>1) {
- PoolingLayer<Dtype>::Forward_cpu(bottom,top);
- return ;
- }
- if (this->kernel_h_!=2 && this->kernel_h_!=3) {
- PoolingLayer<Dtype>::Forward_cpu(bottom,top);
- return ;
- }
- SetupACLLayer(bottom,top);
+ SetupACLOperator(bottom,top);
for (int n = 0; n < bottom[0]->num(); ++n) {
- tensor_mem(this->cpu().input,(void*)(bottom_data));
- cpu_run();
- tensor_mem((void*)(top_data),this->cpu().output);
+ acl_run((void*)bottom_data,(void*)top_data);
bottom_data += bottom[0]->offset(1);
top_data += top[0]->offset(1);
}
#ifdef USE_PROFILING
logtime_util log_time(ACL_POOLING_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_|| this->layer_param_.pooling_param().global_pooling()) {
+ if (Bypass_acl(bottom,top)) {
PoolingLayer<Dtype>::Forward_cpu(bottom,top);
- return;
+ return ;
}
const Dtype* bottom_data = bottom[0]->gpu_data();
Dtype* top_data = top[0]->mutable_gpu_data();
- if (this->layer_param_.pooling_param().pool()!=PoolingParameter_PoolMethod_MAX &&
- this->layer_param_.pooling_param().pool()!=PoolingParameter_PoolMethod_AVE) {
- PoolingLayer<Dtype>::Forward_cpu(bottom,top);
- return ;
- }
- if (this->kernel_h_!=this->kernel_w_) {
- PoolingLayer<Dtype>::Forward_cpu(bottom,top);
- return ;
- }
- if (this->kernel_h_!=2 && this->kernel_h_!=3) {
- PoolingLayer<Dtype>::Forward_cpu(bottom,top);
- return ;
- }
- SetupACLLayer(bottom,top);
+ SetupACLOperator(bottom,top);
for (int n = 0; n < bottom[0]->num(); ++n) {
- tensor_mem(this->gpu().input,(void*)(bottom_data));
- gpu_run();
- tensor_mem((void*)(top_data),this->gpu().output);
+ acl_run((void*)bottom_data,(void*)top_data);
bottom_data += bottom[0]->offset(1);
top_data += top[0]->offset(1);
}
this->force_bypass_acl_path_= bypass_acl_class_layer & FLAGS_ENABLE_ACL_RELU;
}
template <typename Dtype>
-void ACLReLULayer<Dtype>::SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
+void ACLReLULayer<Dtype>::SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top){
- ACLBaseActivationLayer<Dtype>::SetupACLLayer(bottom, top,ActivationLayerInfo::ActivationFunction::RELU);
+ ACLBaseActivationLayer<Dtype>::SetupACLOperator(bottom, top,arm_compute::ActivationLayerInfo::ActivationFunction::RELU);
}
template <typename Dtype>
void ACLReLULayer<Dtype>::Reshape(const vector<Blob<Dtype>*>& bottom,
}
template <typename Dtype>
+bool ACLReLULayer<Dtype>::Bypass_acl(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top) {
+ bool bypass_acl=false;
+ if (this->force_bypass_acl_path_) {
+ bypass_acl=true;
+ }
+ // Fallback to standard Caffe for leaky ReLU.
+ if (ReLULayer<Dtype>::layer_param_.relu_param().negative_slope() != 0) {
+ bypass_acl=true;
+ }
+ if (isScheduleEnable()) {
+ bypass_acl=true;
+ }
+ return bypass_acl;
+}
+
+template <typename Dtype>
void ACLReLULayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top) {
#ifdef USE_PROFILING
logtime_util log_time(ACL_RELU_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_) {
+ if (Bypass_acl(bottom,top)) {
ReLULayer<Dtype>::Forward_cpu(bottom,top);
return;
}
- // Fallback to standard Caffe for leaky ReLU.
- if (ReLULayer<Dtype>::layer_param_.relu_param().negative_slope() != 0) {
- ReLULayer<Dtype>::Forward_cpu(bottom, top);
- return;
- }
ACLBaseActivationLayer<Dtype>::Forward_cpu(bottom,top);
}
#ifdef USE_PROFILING
logtime_util log_time(ACL_RELU_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_) {
- ReLULayer<Dtype>::Forward_cpu(bottom,top);
- return;
- }
- // Fallback to standard Caffe for leaky ReLU.
- if (ReLULayer<Dtype>::layer_param_.relu_param().negative_slope() != 0) {
+ if (Bypass_acl(bottom,top)) {
ReLULayer<Dtype>::Forward_cpu(bottom, top);
return;
}
}
template <typename Dtype>
-void ACLSigmoidLayer<Dtype>::SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
- const vector<Blob<Dtype>*>& top,ActivationLayerInfo::ActivationFunction type){
- ACLBaseActivationLayer<Dtype>::SetupACLLayer(bottom, top,ActivationLayerInfo::ActivationFunction::LOGISTIC);
+void ACLSigmoidLayer<Dtype>::SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top,arm_compute::ActivationLayerInfo::ActivationFunction type){
+ ACLBaseActivationLayer<Dtype>::SetupACLOperator(bottom, top,arm_compute::ActivationLayerInfo::ActivationFunction::LOGISTIC);
}
template <typename Dtype>
void ACLSigmoidLayer<Dtype>::Reshape(const vector<Blob<Dtype>*>& bottom,
}
template <typename Dtype>
+bool ACLSigmoidLayer<Dtype>::Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top){
+ bool bypass_acl=false;
+ if (this->force_bypass_acl_path_) {
+ bypass_acl=true;
+ }
+ return bypass_acl;
+}
+
+template <typename Dtype>
void ACLSigmoidLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top) {
#ifdef USE_PROFILING
logtime_util log_time(ACL_SIGMOID_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_) {
+ if (Bypass_acl(bottom,top)) {
SigmoidLayer<Dtype>::Forward_cpu(bottom,top);
return;
}
#ifdef USE_PROFILING
logtime_util log_time(ACL_SIGMOID_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_) {
+ if (Bypass_acl(bottom,top)) {
SigmoidLayer<Dtype>::Forward_cpu(bottom,top);
return;
}
this->force_bypass_acl_path_= bypass_acl_class_layer & FLAGS_ENABLE_ACL_SOFTMAX;
}
template <typename Dtype>
-void ACLSoftmaxLayer<Dtype>::SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
+void ACLSoftmaxLayer<Dtype>::SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top){
unsigned int channels = bottom[0]->shape(this->softmax_axis_);
- TensorShape shape(channels*this->inner_num_);
- checkreshape(shape,Caffe::arm_gpu_mode());
- if (!this->init_layer_) return;
- this->init_layer_=false;
+ arm_compute::TensorShape shape(channels*this->inner_num_);
+ if (is_operator_init_done(shape)) return;
+ set_operator_init_done();
// Initialize ACL.
- if (Caffe::arm_gpu_mode()) {
- new_gpulayer();
- }else{
- new_cpulayer();
- }
+ new_tensor(input(),shape,InputdataPtr(this,bottom));
+ new_tensor(output(),shape,OutputdataPtr(this,top));
+ acl_configure(softmax,this,NULL);
- //this->force_bypass_acl_path_=false;
- if (Caffe::arm_gpu_mode()) {
- Dtype *top_data = top[0]->mutable_gpu_data();
- const Dtype* bottom_data = bottom[0]->gpu_data();
- new_tensor(this->gpu().input,shape,(void*)bottom_data);
- new_tensor(this->gpu().output,shape,(void*)top_data);
-#ifdef USE_PROFILING
- logtime_util log_time(ACL_CONFIG_INFO);
-#endif //USE_PROFILING
- this->gpu().layer->configure(this->gpu().input,this->gpu().output);
- }else{
- Dtype *top_data = top[0]->mutable_cpu_data();
- const Dtype* bottom_data = bottom[0]->cpu_data();
- new_tensor(this->cpu().input,shape,(void*)bottom_data);
- new_tensor(this->cpu().output,shape,(void*)top_data);
-#ifdef USE_PROFILING
- logtime_util log_time(ACL_CONFIG_INFO);
-#endif //USE_PROFILING
- this->cpu().layer->configure(this->cpu().input,this->cpu().output);
- }
}
template <typename Dtype>
void ACLSoftmaxLayer<Dtype>::Reshape(const vector<Blob<Dtype>*>& bottom,
}
template <typename Dtype>
+bool ACLSoftmaxLayer<Dtype>::Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top){
+ bool bypass_acl=false;
+ if (this->force_bypass_acl_path_ || this->inner_num_>1) {
+ bypass_acl=true;
+ }
+ return bypass_acl;
+}
+
+template <typename Dtype>
void ACLSoftmaxLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top) {
- if(Caffe::arm_gpu_mode()){
+ if(isGPUMode()){
Forward_gpu(bottom, top);
return;
}
#ifdef USE_PROFILING
logtime_util log_time(ACL_SOFTMAX_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_ || this->inner_num_>1) {
+ if (Bypass_acl(bottom,top)) {
SoftmaxLayer<Dtype>::Forward_cpu(bottom,top);
return ;
}
const Dtype* bottom_data = bottom[0]->cpu_data();
Dtype* top_data = top[0]->mutable_cpu_data();
- SetupACLLayer(bottom,top);
+ SetupACLOperator(bottom,top);
int channels = bottom[0]->shape(this->softmax_axis_);
for (int i = 0; i < this->outer_num_; ++i) {
- tensor_mem(this->cpu().input,(void*)(bottom_data));
- cpu_run();
- tensor_mem((void*)(top_data),this->cpu().output);
+ acl_run((void*)bottom_data,(void*)top_data);
top_data += channels;
bottom_data += channels;
}
#ifdef USE_PROFILING
logtime_util log_time(ACL_SOFTMAX_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_|| this->inner_num_>1) {
+ if (Bypass_acl(bottom,top)) {
SoftmaxLayer<Dtype>::Forward_cpu(bottom,top);
return;
}
const Dtype* bottom_data = bottom[0]->gpu_data();
Dtype* top_data = top[0]->mutable_gpu_data();
- SetupACLLayer(bottom,top);
+ SetupACLOperator(bottom,top);
for (int i = 0; i < this->outer_num_; ++i) {
- tensor_mem(this->gpu().input,(void*)(bottom_data));
- gpu_run();
- tensor_mem((void*)(top_data),this->gpu().output);
+ acl_run((void*)bottom_data,(void*)top_data);
top_data += this->inner_num_;
bottom_data += this->inner_num_;
}
}
template <typename Dtype>
-void ACLTanHLayer<Dtype>::SetupACLLayer(const vector<Blob<Dtype>*>& bottom,
- const vector<Blob<Dtype>*>& top, ActivationLayerInfo::ActivationFunction type){
- ACLBaseActivationLayer<Dtype>::SetupACLLayer(bottom, top,ActivationLayerInfo::ActivationFunction::TANH);
+void ACLTanHLayer<Dtype>::SetupACLOperator(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top, arm_compute::ActivationLayerInfo::ActivationFunction type){
+ ACLBaseActivationLayer<Dtype>::SetupACLOperator(bottom, top,arm_compute::ActivationLayerInfo::ActivationFunction::TANH);
}
template <typename Dtype>
}
template <typename Dtype>
+bool ACLTanHLayer<Dtype>::Bypass_acl(const vector<Blob<Dtype>*>& bottom,const vector<Blob<Dtype>*>& top){
+ bool bypass_acl=false;
+ if (this->force_bypass_acl_path_) {
+ bypass_acl=true;
+ }
+ return bypass_acl;
+}
+template <typename Dtype>
void ACLTanHLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top) {
#ifdef USE_PROFILING
logtime_util log_time(ACL_TANH_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_) {
+ if (Bypass_acl(bottom,top)) {
TanHLayer<Dtype>::Forward_cpu(bottom,top);
return;
}
#ifdef USE_PROFILING
logtime_util log_time(ACL_TANH_INFO);
#endif //USE_PROFILING
- if (this->force_bypass_acl_path_) {
+ if (Bypass_acl(bottom,top)) {
TanHLayer<Dtype>::Forward_cpu(bottom,top);
return;
}
--- /dev/null
+import sys
+import os
+import re
+import pdb
+import xlwt
+
+help_ = '''
+Usage:
+ python tpi.py log.txt
+'''
+
+#data_list= {}
+data_list1= []
+data_list2= []
+cnt=0
+times=0.0
+table_val=''
+name_list1= ['allocate','run','configure','tensor_copy','ACL_CONV','ACL_FC','ACL_LRN','ACL_POOLING','ACL_RELU','ACL_SOFTMAX']
+
+
+
+
+def getvalpairs(words):
+ name=''
+ val=''
+ for word in words:
+ if word=='':
+ continue
+ if name=='':
+ name=word
+ else:
+ val=word
+ break;
+ #print word,
+ #print ''
+ return (name,val)
+
+def addpairstolist(db,name,val,idx):
+ #pdb.set_trace()
+ #if idx in db:
+ # db[idx]['val'] += val
+ #else:
+ # db[idx] = {'val':val,'name':name}
+
+ #pdb.set_trace()
+
+ for i in db:
+ if i['name']==name:
+ i['val'] += val
+ return
+ db.append({'idx':idx,'val':val,'name':name})
+
+def gettabnum(line):
+ start=line.find(':')
+ if start==-1:
+ start=0
+ else:
+ start+=1
+ #pdb.set_trace()
+ str=line[start:-1].lstrip(' ')
+ words=re.split('\t',str)
+ idx=0
+ for word in words:
+ idx+=1
+ if word=='':
+ continue
+ break
+ return idx
+
+def decodefile(logfile):
+ data_list=data_list1
+ for line in open(logfile):
+ if line.find(':')==-1:
+ continue
+ #pdb.set_trace()
+ #print line,
+ idx=gettabnum(line)
+ words=re.split('\t|:| |\r|\n',line)
+ #print(words)
+ (name,val)=getvalpairs(words)
+ #print (name,float(val),eval(val))
+ if name == 'used' and val == 'time':
+ data_list=data_list2
+ try:
+ addpairstolist(data_list,name,float(val),idx)
+ except ValueError as e:
+ #print(line)
+ continue
+
+def printresult(db):
+ #for i in db:
+ # print i, db[i]['idx'],db[i]['val']
+ #pdb.set_trace()
+ db.sort(key=lambda obj:obj.get('idx'), reverse=False)
+ tpi_start=0
+ conv_str='ACL_CONV'
+ find_acl = 0
+ name_index=0
+ global trow
+ global tcol
+ for i in db:
+ if i['name']==conv_str:
+ tpi_start=i['idx']
+
+ tpi=0
+ for i in db:
+ if i['idx']>=tpi_start:
+ tpi+=i['val']
+
+ start=len('ACL_')
+
+ table_head='TPI'+'\t'
+ table_val='%.4f' % (tpi/times)+'\t'
+
+ for i in db:
+ #print i
+ if i['idx']<tpi_start:
+ if i['name'].find('ACL_')==0:
+ table_head+=i['name'][start:]+'\t'
+ else:
+ table_head+=i['name']+'\t'
+ table_val+='%.4f' % (i['val']/times)+'\t'
+
+ print(table_head)
+ print(table_val)
+
+ table_head='TPI'+'\t'
+ table_val='%.4f' % (tpi/times)+'\t'
+
+ for i in db:
+ if i['idx']>=tpi_start:
+ if i['name'].find('ACL_')==0:
+ #pdb.set_trace()
+ table_head+=i['name'][start:]+'\t'
+ else:
+ table_head+=i['name']+'\t'
+ table_val+='%.4f' % (i['val']/times)+'\t'
+
+ print(table_head)
+ print(table_val)
+
+ ws.write(trow, tcol, 'TPI')
+ ws.write(trow+1,tcol,'%.4f' % (tpi/times))
+ tcol+=1
+
+ temp_row=trow
+ temp_col=tcol
+ for i in name_list1:
+ if i.find('ACL_')==0 and find_acl==0:
+ temp_row+=2
+ temp_col=2
+ find_acl=1
+ ws.write(temp_row,temp_col,i)
+ ws.write(temp_row+1,temp_col,'0')
+ temp_col+=1
+ find_acl=0
+
+ for i in db:
+ curname=i['name']
+ curvalue='%.4f' % (i['val']/times)
+ if curname == 'ACL_BN':
+ ws.write(trow+2,7,curname)
+ ws.write(trow+3,7,curvalue)
+
+ if curname in name_list1:
+ val_col=name_list1.index(curname)+2
+ val_row=trow
+ # print ('name found'+ curname + curvalue)
+ # print(val_col)
+ # print (val_row)
+ if val_col>5:
+ val_col-=4
+ val_row+=2
+ ws.write(val_row,val_col,curname)
+ ws.write(val_row+1,val_col,curvalue)
+
+ tcol=0
+ trow+=4
+
+
+if __name__ == '__main__' :
+ if len(sys.argv) < 2:
+ print(help_)
+ sys.exit()
+ else:
+ logfile = sys.argv[1]
+
+ filename = os.path.basename(logfile)
+ decodefile(logfile)
+
+ wb = xlwt.Workbook()
+ ws = wb.add_sheet('testsheet',True)
+ trow = 0
+ tcol = 0
+ cnt=1
+ times=1.0
+ table_val=''
+ print('1st time:')
+ ws.write(trow,tcol,'1st time')
+ tcol+=1
+ printresult(data_list1)
+
+ cnt=2
+ times=10.0
+ table_val=''
+ print('\nAverage of 2-11 times:')
+ ws.write(trow, tcol, '2-11 times')
+ tcol+=1
+ printresult(data_list2)
+ wb.save(filename+'.xls')
+ print ('Xls file generated')
+
+
projects / platform / upstream / caffeonacl.git / commitdiff