diff --git a/mace/kernels/opencl/cl/conv_2d_1x1.cl b/mace/kernels/opencl/cl/conv_2d_1x1.cl index 8a07486c243b280a68317538f19fc6aaf8c85287..d371698862460087fa67812282cc72b181e431df 100644 --- a/mace/kernels/opencl/cl/conv_2d_1x1.cl +++ b/mace/kernels/opencl/cl/conv_2d_1x1.cl @@ -49,13 +49,11 @@ __kernel void conv_2d_1x1_v2(__global const float *input, /* n, c, h, w */ int out_chan_len = out_chan_end - out_chan_begin; int pixel_len = out_pixel_end - out_pixel_begin; - if (bias != NULL) { - for (int out_chan = out_chan_begin; out_chan < out_chan_end; ++out_chan) { - float *output_ptr = output_base + out_chan * pixel_num; - float bias_value = bias[out_chan]; - for (int p = 0; p < pixel_len; ++p) { - output_ptr[p] = bias_value; - } + for (int out_chan = out_chan_begin; out_chan < out_chan_end; ++out_chan) { + float *output_ptr = output_base + out_chan * pixel_num; + float bias_value = bias == NULL ? 0 : bias[out_chan]; + for (int p = 0; p < pixel_len; ++p) { + output_ptr[p] = bias_value; } } diff --git a/tools/validate_icnet.py b/tools/validate_icnet.py new file mode 100644 index 0000000000000000000000000000000000000000..319011e7241dfb8007f47f959bb531a1b6302c59 --- /dev/null +++ b/tools/validate_icnet.py @@ -0,0 +1,124 @@ +import argparse +import sys +import tensorflow as tf +import numpy as np + +from tensorflow import gfile + +# Validation Flow: +# 1. Generate input data +# python validate_icnet.py --generate_data 1 \ +# --random_seed 1 +# 2. Use mace_run to run icnet on phone. +# 3. adb pull the result. +# 4. Compare output data of mace and tf +# python validate_icnet.py --model_file opt_icnet.pb \ +# --tf_input_file input_file \ +# --mace_out_file icnet.out + + +def generate_data(shape): + np.random.seed(FLAGS.random_seed) + data = np.random.random(shape) + print FLAGS.tf_input_file + data.astype(np.float32).tofile(FLAGS.tf_input_file) + mace_data = np.transpose(data, axes=(2, 0, 1)) + mace_data.astype(np.float32).tofile(FLAGS.mace_input_file) + print "Generate input file done." + +def load_data(file): + return np.fromfile(file=file, dtype=np.float32) + +def valid_output(out_shape, mace_out_file, tf_out_value): + mace_out_value = load_data(mace_out_file) + mace_out_value = mace_out_value.reshape(out_shape) + tf_out_data_t = np.transpose(tf_out_value, axes=(0, 3, 1, 2)) + res = np.allclose(mace_out_value, tf_out_data_t, rtol=0, atol=1e-5) + print 'Passed! Haha' if res else 'Failed! Oops' + + +def run_model(input_shape): + if not gfile.Exists(FLAGS.model_file): + print("Input graph file '" + FLAGS.model_file + "' does not exist!") + return -1 + + input_graph_def = tf.GraphDef() + with gfile.Open(FLAGS.model_file, "rb") as f: + data = f.read() + input_graph_def.ParseFromString(data) + tf.import_graph_def(input_graph_def, name="") + + with tf.Session() as session: + with session.graph.as_default() as graph: + tf.import_graph_def(input_graph_def, name="") + input_node = graph.get_tensor_by_name('input_node:0') + output_node = graph.get_tensor_by_name('output_node:0') + + input_value = load_data(FLAGS.tf_input_file) + input_value = input_value.reshape(input_shape) + + output_value = session.run(output_node, feed_dict={input_node: [input_value]}) + return output_value + +def main(unused_args): + input_shape = [int(x) for x in FLAGS.input_shape.split(',')] + output_shape = [int(x) for x in FLAGS.output_shape.split(',')] + if FLAGS.generate_data: + generate_data(input_shape) + else: + output_value = run_model(input_shape) + valid_output(output_shape, FLAGS.mace_out_file, output_value) + + +def parse_args(): + """Parses command line arguments.""" + parser = argparse.ArgumentParser() + parser.register("type", "bool", lambda v: v.lower() == "true") + parser.add_argument( + "--model_file", + type=str, + default="", + help="TensorFlow \'GraphDef\' file to load.") + parser.add_argument( + "--tf_input_file", + type=str, + default="", + help="tensorflow input data to load.") + parser.add_argument( + "--mace_input_file", + type=str, + default="", + help="mace input data to load.") + parser.add_argument( + "--mace_out_file", + type=str, + default="", + help="mace output file to load.") + parser.add_argument( + "--input_shape", + type=str, + default="480,480,3", + help="input shape.") + parser.add_argument( + "--output_shape", + type=str, + default="1,2,480,480", + help="output shape.") + parser.add_argument( + "--generate_data", + type='bool', + default="false", + help="Random seed for generate test case.") + parser.add_argument( + "--random_seed", + type=int, + default="0", + help="Random seed for generate test case.") + + return parser.parse_known_args() + + +if __name__ == '__main__': + FLAGS, unparsed = parse_args() + main(unused_args=[sys.argv[0]] + unparsed) + diff --git a/tools/validate_icnet.sh b/tools/validate_icnet.sh new file mode 100644 index 0000000000000000000000000000000000000000..8b5501f5bd66669d3a7d74b46cfb6b965c946fd8 --- /dev/null +++ b/tools/validate_icnet.sh @@ -0,0 +1,71 @@ +#!/bin/bash +# Must run at root dir of mace project. +set -e + +Usage() { + echo 'Usage: bash tools/validate_icnet.sh tf_model_file' +} + +if [ $# != 1 ];then + Usage + exit -1 +fi + +TF_MODEL_FILE_PATH=$1 +MODEL_DIR=$(dirname ${TF_MODEL_FILE_PATH}) +MACE_MODEL_NAME='mace_model.pb' +TF_INPUT_FILE_NAME='tf_model_input' +MACE_INPUT_FILE_NAME='mace_model_input' +OUTPUT_FILE_NAME='icnet.out' +PHONE_DATA_DIR="/data/local/tmp/${MACE_MODEL_NAME}" +KERNEL_DIR="${PHONE_DATA_DIR}/cl/" + +# Step 1: convert tf model to mace model +echo "Step 1: convert tf model to mace model" +bazel build //mace/python/tools:tf_converter +bazel-bin/mace/python/tools/tf_converter --input=${TF_MODEL_FILE_PATH} --output=${MODEL_DIR}/${MACE_MODEL_NAME} + +# Step 2: Generate input data +echo "Step 2: Generate input data" +python tools/validate_icnet.py --generate_data true --random_seed 1 \ + --mace_input_file ${MODEL_DIR}/${MACE_INPUT_FILE_NAME} --tf_input_file ${MODEL_DIR}/${TF_INPUT_FILE_NAME} + +# Step 3: Run model on the phone +echo "Step 3: Run model on the phone" +bazel build -c opt --strip always mace/examples:mace_run \ + --crosstool_top=//external:android/crosstool \ + --host_crosstool_top=@bazel_tools//tools/cpp:toolchain \ + --cpu=arm64-v8a + +adb shell "mkdir -p ${PHONE_DATA_DIR}" +adb shell "mkdir -p ${KERNEL_DIR}" +adb push mace/kernels/opencl/cl/* ${KERNEL_DIR} +adb push ${MODEL_DIR}/${MACE_MODEL_NAME} ${PHONE_DATA_DIR} +adb push ${MODEL_DIR}/${MACE_INPUT_FILE_NAME} ${PHONE_DATA_DIR} +adb push bazel-bin/mace/examples/mace_run ${PHONE_DATA_DIR} + +num_threads=${1:-1} + +adb shell MACE_RUN_PARAMETER_PATH=${PHONE_DATA_DIR}/mace_run.config \ + MACE_KERNEL_PATH=$KERNEL_DIR \ + OMP_NUM_THREADS=$num_threads \ + ${PHONE_DATA_DIR}/mace_run \ + --model=${PHONE_DATA_DIR}/${MACE_MODEL_NAME} \ + --input=input_node \ + --output=icnet/Conv_11/BatchNorm/batchnorm/add_1 \ + --input_shape=1,3,480,480\ + --input_file=${PHONE_DATA_DIR}/${MACE_INPUT_FILE_NAME} \ + --output_file=${PHONE_DATA_DIR}/${OUTPUT_FILE_NAME} \ + --device=OPENCL + +# Step 4: pull the mace run result. +echo "Step 4: pull the mace run result." +adb pull ${PHONE_DATA_DIR}/${OUTPUT_FILE_NAME} ${MODEL_DIR} + +# Step 5: validate the result +echo "Step 5: validate the result" +python tools/validate_icnet.py --model_file ${TF_MODEL_FILE_PATH} \ + --tf_input_file ${MODEL_DIR}/${TF_INPUT_FILE_NAME} \ + --mace_out_file ${MODEL_DIR}/${OUTPUT_FILE_NAME} + +