From c5d834952f222dfce76362c2ed93b2203088d5b5 Mon Sep 17 00:00:00 2001 From: liuqi Date: Wed, 31 Jan 2018 16:47:49 +0800 Subject: [PATCH] Support to run on local PC. --- mace/core/BUILD | 3 +- mace/core/operator.h | 1 - .../hexagon/hexagon_control_wrapper.cc | 2 - mace/kernels/conv_2d.h | 2 - mace/kernels/depthwise_conv2d.h | 2 - mace/kernels/opencl/cl/conv_2d_3x3.cl | 14 +- mace/kernels/pooling.h | 1 - mace/utils/utils.h | 4 +- tools/export_local_lib.sh | 124 ++++++++++++++++++ 9 files changed, 135 insertions(+), 18 deletions(-) create mode 100755 tools/export_local_lib.sh diff --git a/mace/core/BUILD b/mace/core/BUILD index 81b731d7..afaed252 100644 --- a/mace/core/BUILD +++ b/mace/core/BUILD @@ -42,9 +42,8 @@ cc_library( "runtime/opencl/*.h", "runtime/hexagon/*.h", ]), - linkopts = if_android([ + linkopts = ["-ldl",] + if_android([ "-pie", - "-ldl", "-lm", ]), deps = [ diff --git a/mace/core/operator.h b/mace/core/operator.h index d673ca81..185bbc76 100644 --- a/mace/core/operator.h +++ b/mace/core/operator.h @@ -94,7 +94,6 @@ class Operator : public OperatorBase { for (const string &output_str : operator_def.output()) { if (ws->HasTensor(output_str)) { - Tensor *found_tensor = ws->GetTensor(output_str); outputs_.push_back(ws->GetTensor(output_str)); } else { outputs_.push_back(MACE_CHECK_NOTNULL(ws->CreateTensor( diff --git a/mace/core/runtime/hexagon/hexagon_control_wrapper.cc b/mace/core/runtime/hexagon/hexagon_control_wrapper.cc index b7d2bfda..eab3a902 100644 --- a/mace/core/runtime/hexagon/hexagon_control_wrapper.cc +++ b/mace/core/runtime/hexagon/hexagon_control_wrapper.cc @@ -158,7 +158,6 @@ bool HexagonControlWrapper::TeardownGraph() { void HexagonControlWrapper::PrintLog() { char *buf; - unsigned char *p; if ((buf = new char[PRINT_BUFSIZE]) == NULL) return; hexagon_nn_getlog(nn_id_, reinterpret_cast(buf), PRINT_BUFSIZE); LOG(INFO) << string(buf); @@ -168,7 +167,6 @@ void HexagonControlWrapper::PrintLog() { void HexagonControlWrapper::PrintGraph() { LOG(INFO) << "Print Graph"; char *buf; - unsigned char *p; if ((buf = new char[PRINT_BUFSIZE]) == NULL) return; hexagon_nn_snpprint(nn_id_, reinterpret_cast(buf), PRINT_BUFSIZE); LOG(INFO) << string(buf); diff --git a/mace/kernels/conv_2d.h b/mace/kernels/conv_2d.h index 01c55434..8da36579 100644 --- a/mace/kernels/conv_2d.h +++ b/mace/kernels/conv_2d.h @@ -94,8 +94,6 @@ struct Conv2dFunctor : Conv2dFunctorBase { index_t padded_h_stop = input_height + paddings[0] - paddings[0] / 2; index_t padded_w_stop = input_width + paddings[1] - paddings[1] / 2; - index_t kernel_size = input_channels * kernel_h * kernel_w; - Tensor::MappingGuard input_mapper(input); Tensor::MappingGuard filter_mapper(filter); Tensor::MappingGuard bias_mapper(bias); diff --git a/mace/kernels/depthwise_conv2d.h b/mace/kernels/depthwise_conv2d.h index caff1893..cdf915c0 100644 --- a/mace/kernels/depthwise_conv2d.h +++ b/mace/kernels/depthwise_conv2d.h @@ -105,8 +105,6 @@ struct DepthwiseConv2dFunctor : public DepthwiseConv2dFunctorBase { index_t padded_h_stop = input_height + paddings[0] - paddings[0] / 2; index_t padded_w_stop = input_width + paddings[1] - paddings[1] / 2; - const index_t kernel_size = kernel_h * kernel_w; - Tensor::MappingGuard input_mapper(input); Tensor::MappingGuard filter_mapper(filter); Tensor::MappingGuard bias_mapper(bias); diff --git a/mace/kernels/opencl/cl/conv_2d_3x3.cl b/mace/kernels/opencl/cl/conv_2d_3x3.cl index 8d0b4d1a..9403c905 100644 --- a/mace/kernels/opencl/cl/conv_2d_3x3.cl +++ b/mace/kernels/opencl/cl/conv_2d_3x3.cl @@ -62,17 +62,17 @@ __kernel void conv_2d_3x3(__read_only image2d_t input, /* [c%4 * w * c/4, h * b] for (short in_ch_blk = 0; in_ch_blk < in_ch_blks; ++in_ch_blk) { const int in_idx = mul24(in_ch_blk, in_width); int filter_x_part0 = in_ch_blk << 2; + int in_hb_idx = height_idx; for (short hb_idx = 0; hb_idx < 3; ++hb_idx) { - // TODO (heliangliang) optimize out these muls - int in_hb_value = height_idx + mul24(hb_idx, dilation_h); - in_hb_value = select(in_hb_value + batch_idx, - -1, - (in_hb_value < 0 || in_hb_value >= in_height)); + int in_hb_value = select(in_hb_idx + batch_idx, + -1, + (in_hb_idx < 0 || in_hb_idx >= in_height)); int filter_x_part1 = 0; + int in_width_idx = 0; for (short width_idx = 0; width_idx < 3; ++width_idx) { int in_width_value; #define READ_INPUT(i) \ - in_width_value = in_width##i + mul24(width_idx, dilation_w); \ + in_width_value = in_width##i + in_width_idx; \ in_width_value = select(in_idx + in_width_value, \ -1, \ (in_width_value < 0 || in_width_value >= in_width)); \ @@ -120,8 +120,10 @@ __kernel void conv_2d_3x3(__read_only image2d_t input, /* [c%4 * w * c/4, h * b] out4 = mad(in4.w, weights3, out4); filter_x_part1 += rounded_in_ch; + in_width_idx += dilation_w; } filter_x_part0 += rounded_in_ch_x_3; + in_hb_idx += dilation_h; } } diff --git a/mace/kernels/pooling.h b/mace/kernels/pooling.h index 99d2363f..dcf1c53a 100644 --- a/mace/kernels/pooling.h +++ b/mace/kernels/pooling.h @@ -76,7 +76,6 @@ struct PoolingFunctor : PoolingFunctorBase { index_t height = output_shape[1]; index_t width = output_shape[2]; index_t channels = output_shape[3]; - index_t out_image_size = height * width; index_t input_height = input_shape[1]; index_t input_width = input_shape[2]; diff --git a/mace/utils/utils.h b/mace/utils/utils.h index 73e1c6db..ce682f06 100644 --- a/mace/utils/utils.h +++ b/mace/utils/utils.h @@ -52,7 +52,7 @@ inline std::string ObfuscateString(const std::string &src, const std::string &lookup_table) { std::string dest; dest.resize(src.size()); - for (int i = 0; i < src.size(); i++) { + for (size_t i = 0; i < src.size(); i++) { dest[i] = src[i] ^ lookup_table[i % lookup_table.size()]; } return std::move(dest); @@ -73,7 +73,7 @@ inline std::string ObfuscateSymbol(const std::string &src) { dest[0] = src[0]; // avoid invalid symbol which starts from 0-9 const std::string encode_dict = "ABCDEFGHIJKLMNOPQRSTUVWXYZabcdefghijklmnopqrstuvwxyz0123456789_"; - for (int i = 1; i < src.size(); i++) { + for (size_t i = 1; i < src.size(); i++) { char ch = src[i]; int idx; if (ch >= '0' && ch <= '9') { diff --git a/tools/export_local_lib.sh b/tools/export_local_lib.sh new file mode 100755 index 00000000..26aa3fa4 --- /dev/null +++ b/tools/export_local_lib.sh @@ -0,0 +1,124 @@ +#!/bin/bash + +set -e + +Usage() { + echo "Usage: ./tools/export_local_lib.sh export_include_dir export_lib_dir" + echo "eg: ./tools/export_local_lib.sh ../include ../lib/libmace_v7" +} + +if [ $# -lt 2 ]; then + Usage + exit -1 +fi + +EXPORT_INCLUDE_DIR=$1 +EXPORT_LIB_DIR=$2 + +MACE_SOURCE_DIR=`/bin/pwd` +CODEGEN_DIR=${MACE_SOURCE_DIR}/mace/codegen +VERSION_CODEGEN_DIR=${CODEGEN_DIR}/version +STRIP="--strip always" + +# TO REMOVE +LIBMACE_TEMP_DIR=`mktemp -d -t libmace.XXXX` + +LIBMACE_NAME="libmace" +LIBMACE_DEV_NAME="libmace_dev" +LIBMACE_PROD_NAME="libmace_prod" + +libmace_targets=( + "//mace/ops:ops" + "//mace/kernels:kernels" + "//mace/codegen:generated_version" + "//mace/core:core" + "//mace/utils:logging" +) + +libmace_dev_targets=( + "//mace/codegen:generated_opencl_dev" + "//mace/core:opencl_dev" + "//mace/utils:tuner_dev" +) + +libmace_prod_targets=( + "//mace/core:opencl_prod" + "//mace/utils:tuner_prod" +) + +all_targets=(${libmace_targets[*]} ${libmace_dev_targets[*]} ${libmace_prod_targets[*]}) + +build_target() +{ + BAZEL_TARGET=$1 + bazel build --verbose_failures -c opt --strip always $BAZEL_TARGET \ + --copt="-std=c++11" \ + --copt="-D_GLIBCXX_USE_C99_MATH_TR1" \ + --copt="-Werror=return-type" \ + --copt="-DMACE_OBFUSCATE_LITERALS" || exit -1 +} + +merge_libs() +{ + CREATE_LIB_NAME=$1 + LIBS_LIST=$2 + echo "create ${LIBMACE_TEMP_DIR}/${CREATE_LIB_NAME}.a" > ${LIBMACE_TEMP_DIR}/${CREATE_LIB_NAME}.mri || exit -1 + + for lib_target in ${LIBS_LIST[*]} + do + lib_dir=`echo ${lib_target} | cut -d: -f1` + lib_dir=${lib_dir#//} + lib_name_prefix=lib`echo ${lib_target} | cut -d: -f2` + bin_path="${MACE_SOURCE_DIR}/bazel-bin/${lib_dir}/${lib_name_prefix}" + if [ -f "${bin_path}.a" ]; then + bin_path="${bin_path}.a" + else + bin_path="${bin_path}.lo" + fi + echo "addlib ${bin_path}" >> ${LIBMACE_TEMP_DIR}/${CREATE_LIB_NAME}.mri || exit -1 + done + + echo "save" >> ${LIBMACE_TEMP_DIR}/${CREATE_LIB_NAME}.mri || exit -1 + echo "end" >> ${LIBMACE_TEMP_DIR}/${CREATE_LIB_NAME}.mri || exit -1 + + $ANDROID_NDK_HOME/toolchains/aarch64-linux-android-4.9/prebuilt/linux-x86_64/bin/aarch64-linux-android-ar \ + -M < ${LIBMACE_TEMP_DIR}/${CREATE_LIB_NAME}.mri || exit -1 +} + +echo "Step 1: Generate encrypted opencl source" +python mace/python/tools/encrypt_opencl_codegen.py \ + --cl_kernel_dir=./mace/kernels/opencl/cl/ \ + --output_path=${CODEGEN_DIR}/opencl/opencl_encrypt_program.cc || exit -1 + + +echo "Step 2: Generate version source" +rm -rf ${VERSION_CODEGEN_DIR} +mkdir ${VERSION_CODEGEN_DIR} +bash mace/tools/git/gen_version_source.sh ${CODEGEN_DIR}/version/version.cc || exit -1 + + +echo "Step 3: Build libmace targets" +bazel clean +for target in ${all_targets[*]} +do + build_target ${target} +done + +echo "Step 4: Create mri files and generate merged libs" +merge_libs "libmace" "${libmace_targets[*]}" +merge_libs "libmace_dev" "${libmace_dev_targets[*]}" +merge_libs "libmace_prod" "${libmace_prod_targets[*]}" + +echo "Step 5: Export lib" +rm -rf ${EXPORT_INCLUDE_DIR} +mkdir -p ${EXPORT_INCLUDE_DIR}/mace/core/public +rm -rf ${EXPORT_LIB_DIR} +mkdir -p ${EXPORT_LIB_DIR} + +cp ${MACE_SOURCE_DIR}/mace/core/public/* ${EXPORT_INCLUDE_DIR}/mace/core/public || exit -1 +cp ${LIBMACE_TEMP_DIR}/libmace.a ${LIBMACE_TEMP_DIR}/libmace_dev.a ${LIBMACE_TEMP_DIR}/libmace_prod.a ${EXPORT_LIB_DIR}/ || exit -1 + +echo "Step 6: Remove temporary file" +rm -rf ${LIBMACE_TEMP_DIR} + +echo "Done!" -- GitLab