diff --git a/mace/core/workspace.cc b/mace/core/workspace.cc index a0eab8bf3d4678815d0815157d494773cc44004d..2a172f3e93161c1e8fe8cc3c9f4105b5bd88a3fb 100644 --- a/mace/core/workspace.cc +++ b/mace/core/workspace.cc @@ -55,7 +55,9 @@ Tensor *Workspace::GetTensor(const string &name) { void Workspace::LoadModelTensor(const NetDef &net_def, DeviceType type) { Serializer serializer; for (auto &tensor_proto : net_def.tensors()) { - VLOG(1) << "Load tensor: " << tensor_proto.name() << " has shape: " + VLOG(1) << "Load tensor: " << tensor_proto.name() + << ", with data type: " << tensor_proto.data_type() + << ", has shape: " << internal::MakeString(vector(tensor_proto.dims().begin(), tensor_proto.dims().end())); tensor_map_[tensor_proto.name()] = diff --git a/mace/examples/mace_run.cc b/mace/examples/mace_run.cc index d75b7b9483883f5dc192068c5831d7bb3b61d6ee..c13d6a95529afda817e34c92aa43799c8e55a959 100644 --- a/mace/examples/mace_run.cc +++ b/mace/examples/mace_run.cc @@ -81,30 +81,36 @@ int main(int argc, char **argv) { net_def.ParseFromIstream(&file_stream); file_stream.close(); + DeviceType device_type; + DeviceType_Parse(device, &device_type); + VLOG(0) << device_type; Workspace ws; - ws.LoadModelTensor(net_def, DeviceType::CPU); + ws.LoadModelTensor(net_def, device_type); Tensor *input_tensor = - ws.CreateTensor(input_node + ":0", GetDeviceAllocator(DeviceType::CPU), DT_FLOAT); + ws.CreateTensor(input_node + ":0", GetDeviceAllocator(device_type), DT_FLOAT); input_tensor->Resize(shape); - float *input_data = input_tensor->mutable_data(); + { + Tensor::MappingGuard input_guard(input_tensor); + float *input_data = input_tensor->mutable_data(); + + // load input + ifstream in_file(input_file, ios::in | ios::binary); + in_file.read(reinterpret_cast(input_data), + input_tensor->size() * sizeof(float)); + in_file.close(); + } - // load input - ifstream in_file(input_file, ios::in | ios::binary); - in_file.read(reinterpret_cast(input_data), - input_tensor->size() * sizeof(float)); - in_file.close(); // run model - DeviceType device_type; - DeviceType_Parse(device, &device_type); - VLOG(0) << device_type; auto net = CreateNet(net_def, &ws, device_type); + VLOG(0) << "warm up"; // warm up - for (int i = 0; i < 2; ++i) { + for (int i = 0; i < 1; ++i) { net->Run(); } + VLOG(0) << "run"; timeval tv1, tv2; gettimeofday(&tv1, NULL); for (int i = 0; i < round; ++i) { @@ -120,9 +126,15 @@ int main(int argc, char **argv) { // save output const Tensor *output = ws.GetTensor(output_node + ":0"); + Tensor::MappingGuard output_guard(output); ofstream out_file(output_file, ios::binary); out_file.write((const char *)(output->data()), output->size() * sizeof(float)); out_file.flush(); out_file.close(); + VLOG(0) << "Output shape: [" + << output->dim(0) << ", " + << output->dim(1) << ", " + << output->dim(2) << ", " + << output->dim(3) << "]"; } \ No newline at end of file diff --git a/mace/kernels/conv_pool_2d_util.cc b/mace/kernels/conv_pool_2d_util.cc index f3fe94c8a82223d1658de1cd4f781847686af233..580a124b21d2d313a8f8a42358de64834cc6e76c 100644 --- a/mace/kernels/conv_pool_2d_util.cc +++ b/mace/kernels/conv_pool_2d_util.cc @@ -17,7 +17,7 @@ void CalcPaddingAndOutputSize(const index_t *input_shape, // NCHW MACE_CHECK(dilations[0] > 0 && dilations[1] > 0, "Invalid dilations, must >= 1"); MACE_CHECK((dilations[0] == 1 || strides[0] == 1) && - (dilations[1] == 1 || strides[1] == 1), + (dilations[1] == 1 || strides[1] == 1), "If dilations > 1, strides should be 1"); MACE_CHECK_NOTNULL(output_shape); MACE_CHECK_NOTNULL(padding_size); @@ -51,8 +51,7 @@ void CalcPaddingAndOutputSize(const index_t *input_shape, // NCHW output_height = (input_shape[2] + k_extent_height - 2) / strides[0] + 1; output_width = (input_shape[3] + k_extent_width - 2) / strides[1] + 1; break; - default: - MACE_CHECK(false, "Unsupported padding type: ", padding); + default:MACE_CHECK(false, "Unsupported padding type: ", padding); } // Note: TensorFlow may padded one more on the right/bottom side @@ -61,10 +60,10 @@ void CalcPaddingAndOutputSize(const index_t *input_shape, // NCHW // based on the model accuracy. padding_size[0] = - std::max(0, (output_height - 1) * strides[0] + std::max(0, (output_height - 1) * strides[0] + k_extent_height - input_shape[2]); padding_size[1] = - std::max(0, (output_width - 1) * strides[1] + std::max(0, (output_width - 1) * strides[1] + k_extent_width - input_shape[3]); output_shape[0] = input_shape[0]; @@ -82,7 +81,7 @@ void CalPaddingSize(const index_t *input_shape, // NCHW MACE_CHECK(dilations[0] > 0 && dilations[1] > 0, "Invalid dilations, must >= 1"); MACE_CHECK((dilations[0] == 1 || strides[0] == 1) && - (dilations[1] == 1 || strides[1] == 1), + (dilations[1] == 1 || strides[1] == 1), "If dilations > 1, strides should be 1"); MACE_CHECK_NOTNULL(padding_size); @@ -103,8 +102,7 @@ void CalPaddingSize(const index_t *input_shape, // NCHW output_height = (input_shape[2] + k_extent_height - 2) / strides[0] + 1; output_width = (input_shape[3] + k_extent_width - 2) / strides[1] + 1; break; - default: - MACE_CHECK(false, "Unsupported padding type: ", padding); + default:MACE_CHECK(false, "Unsupported padding type: ", padding); } // Note: TensorFlow may padded one more on the right/bottom side @@ -112,10 +110,10 @@ void CalPaddingSize(const index_t *input_shape, // NCHW // utilize the more centered features. We need to benchmark // based on the model accuracy. padding_size[0] = - std::max(0, (output_height - 1) * strides[0] + std::max(0, (output_height - 1) * strides[0] + k_extent_height - input_shape[2]); padding_size[1] = - std::max(0, (output_width - 1) * strides[1] + std::max(0, (output_width - 1) * strides[1] + k_extent_width - input_shape[3]); } @@ -123,6 +121,7 @@ void ConstructInputWithPadding(const Tensor *input_tensor, const int *paddings, Tensor *output_tensor, bool padding_same_value) { + VLOG(1) << "input: " << input_tensor->NumElements(); Tensor::MappingGuard input_mapper(input_tensor); const float *input = input_tensor->data(); const index_t *input_shape = input_tensor->shape().data(); diff --git a/mace/kernels/opencl/cl/conv_2d_1x1.cl b/mace/kernels/opencl/cl/conv_2d_1x1.cl index ae004658f707f5a9d1c74a76afd69ef324d4ce18..d371698862460087fa67812282cc72b181e431df 100644 --- a/mace/kernels/opencl/cl/conv_2d_1x1.cl +++ b/mace/kernels/opencl/cl/conv_2d_1x1.cl @@ -51,7 +51,7 @@ __kernel void conv_2d_1x1_v2(__global const float *input, /* n, c, h, w */ 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]; + 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/mace/kernels/opencl/cl/conv_2d_3x3.cl b/mace/kernels/opencl/cl/conv_2d_3x3.cl index c51cc2b6f03993538cd300b33688bb23fd447104..b3f7735d5f6ac78e465d3bfefd2ab6aeed903250 100644 --- a/mace/kernels/opencl/cl/conv_2d_3x3.cl +++ b/mace/kernels/opencl/cl/conv_2d_3x3.cl @@ -39,7 +39,8 @@ void kernel conv_2d_3x3(global const float *input, float *output_ptr = output_base + i * out_pixel; const float *filter_base = filter + i * in_chan_num * 9; if (pixels == 4) { - float4 res = (float4)bias[i]; + + float4 res = bias == NULL ? 0 : (float4)bias[i]; for (int in_chan_idx = 0; in_chan_idx < in_chan_num; ++in_chan_idx) { const float* input_ptr = input_base + in_chan_idx * in_pixel; const float* filter_ptr = filter_base + in_chan_idx * 9; @@ -56,7 +57,7 @@ void kernel conv_2d_3x3(global const float *input, vstore4(res, 0, output_ptr); } else { for (int p = 0; p < pixels; ++p) { - float res = bias[i]; + float res = bias == NULL ? 0 : bias[i]; for (uint in_chan_idx = 0; in_chan_idx < in_chan_num; ++in_chan_idx) { const float* input_ptr = input_base + in_chan_idx * in_pixel + p * stride_w; const float* filter_ptr = filter_base + in_chan_idx * 9; diff --git a/mace/kernels/opencl/conv_2d_opencl_1x1.cc b/mace/kernels/opencl/conv_2d_opencl_1x1.cc index 3e988d8304a1fc78a9fa39ee652b4d1e52e084cd..ba784d0552bd3f5a67558ab1392905db35ae2c4a 100644 --- a/mace/kernels/opencl/conv_2d_opencl_1x1.cc +++ b/mace/kernels/opencl/conv_2d_opencl_1x1.cc @@ -68,8 +68,12 @@ void Conv1x1V2(const Tensor *input, *(static_cast(input->buffer()))); conv_2d_kernel.setArg(idx++, *(static_cast(filter->buffer()))); - conv_2d_kernel.setArg(idx++, - *(static_cast(bias->buffer()))); + if (bias == NULL) { + conv_2d_kernel.setArg(idx++, NULL); + } else { + conv_2d_kernel.setArg(idx++, + *(static_cast(bias->buffer()))); + } conv_2d_kernel.setArg(idx++, *(static_cast(output->buffer()))); conv_2d_kernel.setArg(idx++, static_cast(input_channels)); conv_2d_kernel.setArg(idx++, static_cast(channels)); diff --git a/mace/kernels/opencl/conv_2d_opencl_3x3.cc b/mace/kernels/opencl/conv_2d_opencl_3x3.cc index 537204f74a7f07d258f4e8afd7dac9c76b53f48b..a0224484c6acd647778caf421fc534b02db0b8a2 100644 --- a/mace/kernels/opencl/conv_2d_opencl_3x3.cc +++ b/mace/kernels/opencl/conv_2d_opencl_3x3.cc @@ -27,7 +27,11 @@ static void InnerConv2dK3x3S12(const Tensor *input, const Tensor *filter, uint32_t idx = 0; conv_kernel.setArg(idx++, *(static_cast(input->buffer()))); conv_kernel.setArg(idx++, *(static_cast(filter->buffer()))); - conv_kernel.setArg(idx++, *(static_cast(bias->buffer()))); + if (bias == nullptr) { + conv_kernel.setArg(idx++, NULL); + } else { + conv_kernel.setArg(idx++, *(static_cast(bias->buffer()))); + } conv_kernel.setArg(idx++, *(static_cast(output->buffer()))); conv_kernel.setArg(idx++, static_cast(input->dim(1))); conv_kernel.setArg(idx++, static_cast(channels)); diff --git a/mace/kernels/opencl/resize_bilinear_opencl.cc b/mace/kernels/opencl/resize_bilinear_opencl.cc index 6ad4a2d661ca092dfeb8198dda6519aa5c50bb44..50d717ef7bd225ab4aee4222042d9ee0ac7770f2 100644 --- a/mace/kernels/opencl/resize_bilinear_opencl.cc +++ b/mace/kernels/opencl/resize_bilinear_opencl.cc @@ -36,12 +36,13 @@ void ResizeBilinearFunctor::operator()( uint32_t idx = 0; rb_kernel.setArg(idx++, *(static_cast(input->buffer()))); rb_kernel.setArg(idx++, *(static_cast(output->buffer()))); - rb_kernel.setArg(idx++, static_cast(height_scale)); - rb_kernel.setArg(idx++, static_cast(width_scale)); + rb_kernel.setArg(idx++, height_scale); + rb_kernel.setArg(idx++, width_scale); rb_kernel.setArg(idx++, static_cast(in_height)); rb_kernel.setArg(idx++, static_cast(in_width)); auto command_queue = runtime->command_queue(); + cl_int error = command_queue.enqueueNDRangeKernel( rb_kernel, cl::NullRange, cl::NDRange(static_cast(batch * channels), diff --git a/mace/kernels/resize_bilinear.h b/mace/kernels/resize_bilinear.h index 59986cf6d8ca4ba5a16887545e7e019f6588dda8..59bb2505c9c379c1b0700d7a515a880a704d72db 100644 --- a/mace/kernels/resize_bilinear.h +++ b/mace/kernels/resize_bilinear.h @@ -154,7 +154,7 @@ class ResizeBilinearFunctor { if (size_[0] < 0 || size_[1] < 0) { MACE_CHECK(resize_dims != nullptr && resize_dims->dim_size() == 1); Tensor::MappingGuard resize_dims_mapper(resize_dims); - auto dims_data = resize_dims->data(); + auto dims_data = resize_dims->data(); *out_height = dims_data[0]; *out_width = dims_data[1]; } else { diff --git a/mace/ops/conv_2d_test.cc b/mace/ops/conv_2d_test.cc index 31d3130a88796bd21776e2a67ea12be1ed9a7fdb..8eb805d3987dde93a640468918f663ad9da5377e 100644 --- a/mace/ops/conv_2d_test.cc +++ b/mace/ops/conv_2d_test.cc @@ -9,7 +9,7 @@ using namespace mace; class Conv2dOpTest : public OpsTestBase {}; -template +template void TestSimple3x3VALID() { OpsTestNet net; OpDefBuilder("Conv2D", "Conv2dTest") @@ -44,7 +44,7 @@ void TestSimple3x3VALID() { } -template +template void TestSimple3x3SAME() { OpsTestNet net; OpDefBuilder("Conv2D", "Conv2dTest") @@ -93,7 +93,51 @@ TEST_F(Conv2dOpTest, OPENCLSimple) { TestSimple3x3SAME(); } -template +template +void TestSimple3x3WithoutBias() { + OpsTestNet net; + OpDefBuilder("Conv2D", "Conv2dTest") + .Input("Input") + .Input("Filter") + .Output("Output") + .AddIntsArg("strides", {1, 1}) + .AddIntArg("padding", Padding::VALID) + .AddIntsArg("dilations", {1, 1}) + .Finalize(net.NewOperatorDef()); + + // Add args + + // Add input data + net.AddInputFromArray( + "Input", {1, 2, 3, 3}, + {1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1}); + net.AddInputFromArray( + "Filter", {1, 2, 3, 3}, + {1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, + 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f}); + + // Run + net.RunOp(D); + + // Check + auto expected = CreateTensor({1, 1, 1, 1}, {18.0f}); + + ExpectTensorNear(*expected, *net.GetOutput("Output"), 0.001); +} + +TEST_F(Conv2dOpTest, CPUWithoutBias) { + TestSimple3x3WithoutBias(); +} + +TEST_F(Conv2dOpTest, NEONWithouBias) { + TestSimple3x3WithoutBias(); +} + +TEST_F(Conv2dOpTest, OPENCLWithoutBias) { + TestSimple3x3WithoutBias(); +} + +template static void TestCombined3x3() { // Construct graph OpsTestNet net; @@ -143,7 +187,7 @@ TEST_F(Conv2dOpTest, OPENCLCombined) { TestCombined3x3(); } -template +template void TestConv1x1() { // Construct graph OpsTestNet net; @@ -178,9 +222,9 @@ void TestConv1x1() { // Check auto expected = CreateTensor( {1, 2, 3, 10}, - {5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, - 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, - 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, + {5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, + 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, + 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 5.1f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f, 10.2f}); @@ -196,7 +240,7 @@ TEST_F(Conv2dOpTest, OPENCLConv1x1) { TestConv1x1(); } -template +template static void TestAlignedConvNxNS12() { testing::internal::LogToStderr(); auto func = [&](int kernel_h, int kernel_w, int stride_h, int stride_w, @@ -254,7 +298,7 @@ TEST_F(Conv2dOpTest, OPENCLAlignedConvNxNS12) { TestAlignedConvNxNS12(); } -template +template static void TestUnalignedConvNxNS12() { testing::internal::LogToStderr(); auto func = [&](int kernel_h, int kernel_w, int stride_h, int stride_w, diff --git a/mace/ops/resize_bilinear_test.cc b/mace/ops/resize_bilinear_test.cc index c64a38151658eb1f9a84953c1113e83bcd9f32a5..7b7cee9d97da3afd98e80ff710815f06cf1d8eef 100644 --- a/mace/ops/resize_bilinear_test.cc +++ b/mace/ops/resize_bilinear_test.cc @@ -24,7 +24,7 @@ TEST_F(ResizeBilinearTest, CPUResizeBilinearWOAlignCorners) { vector input(24); std::iota(begin(input), end(input), 0); net.AddInputFromArray("Input", {1, 3, 2, 4}, input); - net.AddInputFromArray("OutSize", {2}, {1, 2}); + net.AddInputFromArray("OutSize", {2}, {1, 2}); // Run net.RunOp(); @@ -50,7 +50,7 @@ TEST_F(ResizeBilinearTest, ResizeBilinearWAlignCorners) { vector input(24); std::iota(begin(input), end(input), 0); net.AddInputFromArray("Input", {1, 3, 2, 4}, input); - net.AddInputFromArray("OutSize", {2}, {1, 2}); + net.AddInputFromArray("OutSize", {2}, {1, 2}); // Run net.RunOp(); @@ -86,7 +86,7 @@ void TestRandomResizeBilinear() { // Add input data net.AddRandomInput("Input", {batch, channels, in_height, in_width}); - net.AddInputFromArray("OutSize", {2}, {height, width}); + net.AddInputFromArray("OutSize", {2}, {height, width}); // Run net.RunOp(D); diff --git a/mace/python/tools/tf_converter_lib.py b/mace/python/tools/tf_converter_lib.py index 0568bfcd82839c31c9ffcf2b56d1f1bde490089e..97575bf23ce9583f1db75ce37d5bc699d0f0189e 100644 --- a/mace/python/tools/tf_converter_lib.py +++ b/mace/python/tools/tf_converter_lib.py @@ -2,6 +2,7 @@ from mace.proto import mace_pb2 import tensorflow as tf import numpy as np +# TODO: support NCHW formt, now only support NHWC. padding_mode = { 'VALID': 0, 'SAME': 1, @@ -22,7 +23,7 @@ def convert_tensor(op, tensor): op.name.endswith('weights') or op.name.endswith('kernel')) \ and op.outputs[0].consumers()[0].type.find('Conv') != -1: - if op.outputs[0].consumers()[0].get_attr('data_format') == 'NCHW': + if op.outputs[0].consumers()[0].get_attr('data_format') == 'NHWC': tf_tensor = np.transpose(tf_tensor, axes=(3, 2, 0, 1)) shape = [shape[3], shape[2], shape[0], shape[1]] # print (tensor.name, shape) @@ -70,7 +71,7 @@ def convert_ops(unresolved_ops, net_def): padding_arg.i = padding_mode[first_op.get_attr('padding')] strides_arg = op_def.arg.add() strides_arg.name = 'strides' - strides_arg.ints.extend(first_op.get_attr('strides')[2:]) + strides_arg.ints.extend(first_op.get_attr('strides')[1:3]) data_format_arg = op_def.arg.add() data_format_arg.name = 'data_format' data_format_arg.s = 'NCHW' @@ -129,10 +130,10 @@ def convert_ops(unresolved_ops, net_def): padding_arg.i = padding_mode[first_op.get_attr('padding')] strides_arg = op_def.arg.add() strides_arg.name = 'strides' - strides_arg.ints.extend(first_op.get_attr('strides')[2:]) + strides_arg.ints.extend(first_op.get_attr('strides')[1:3]) kernels_arg = op_def.arg.add() kernels_arg.name = 'kernels' - kernels_arg.ints.extend(first_op.get_attr('ksize')[2:]) + kernels_arg.ints.extend(first_op.get_attr('ksize')[1:3]) data_format_arg = op_def.arg.add() data_format_arg.name = 'data_format' data_format_arg.s = 'NCHW' 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} + +