From 397e0ec2cc1bcde3d73b4e884de01e3fb54e0207 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Fri, 24 Nov 2017 17:36:01 -0800 Subject: [PATCH 0001/1867] Add DT_HALF support for SpaceToDepth on GPU This fix tries to address the issue raised in 14871 where there were no DT_HALF support for SpaceToDepth on GPU. This fix adds DT_HALF support on GPU and adds aditional test cases. This fix fixes 14871. Signed-off-by: Yong Tang --- tensorflow/core/kernels/spacetodepth_op.cc | 3 +++ tensorflow/core/kernels/spacetodepth_op_gpu.cu.cc | 4 ++++ 2 files changed, 7 insertions(+) diff --git a/tensorflow/core/kernels/spacetodepth_op.cc b/tensorflow/core/kernels/spacetodepth_op.cc index 23df1c35e5..d93a2a9bad 100644 --- a/tensorflow/core/kernels/spacetodepth_op.cc +++ b/tensorflow/core/kernels/spacetodepth_op.cc @@ -187,6 +187,9 @@ TF_CALL_ALL_TYPES(REGISTER); REGISTER_KERNEL_BUILDER( Name("SpaceToDepth").Device(DEVICE_GPU).TypeConstraint("T"), SpaceToDepthOp); +REGISTER_KERNEL_BUILDER( + Name("SpaceToDepth").Device(DEVICE_GPU).TypeConstraint("T"), + SpaceToDepthOp); REGISTER_KERNEL_BUILDER( Name("SpaceToDepth").Device(DEVICE_GPU).TypeConstraint("T"), SpaceToDepthOp); diff --git a/tensorflow/core/kernels/spacetodepth_op_gpu.cu.cc b/tensorflow/core/kernels/spacetodepth_op_gpu.cu.cc index a1a01e8813..e841472972 100644 --- a/tensorflow/core/kernels/spacetodepth_op_gpu.cu.cc +++ b/tensorflow/core/kernels/spacetodepth_op_gpu.cu.cc @@ -225,6 +225,10 @@ struct SpaceToDepthOpFunctor { template struct functor::SpaceToDepthOpFunctor; template struct functor::SpaceToDepthOpFunctor; +// Instantiate the GPU implementations for Eigen::Half. +template struct functor::SpaceToDepthOpFunctor; +template struct functor::SpaceToDepthOpFunctor; + // NCHW_VECT_C with 4 x qint8 can be treated as NCHW int32. template struct functor::SpaceToDepthOpFunctor; -- GitLab From 1d77785e9e13241cb318edce4661e0bdc2dd3095 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Fri, 24 Nov 2017 17:37:27 -0800 Subject: [PATCH 0002/1867] Add test cases for DT_HALF support for SpaceToDepth on GPU. Signed-off-by: Yong Tang --- tensorflow/python/kernel_tests/spacetodepth_op_test.py | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/tensorflow/python/kernel_tests/spacetodepth_op_test.py b/tensorflow/python/kernel_tests/spacetodepth_op_test.py index 3c98a685e0..4af0e6f9db 100644 --- a/tensorflow/python/kernel_tests/spacetodepth_op_test.py +++ b/tensorflow/python/kernel_tests/spacetodepth_op_test.py @@ -34,8 +34,8 @@ from tensorflow.python.platform import tf_logging class SpaceToDepthTest(test.TestCase): - def _testOne(self, inputs, block_size, outputs): - input_nhwc = math_ops.to_float(inputs) + def _testOne(self, inputs, block_size, outputs, dtype=dtypes.float32): + input_nhwc = math_ops.cast(inputs, dtype) with self.test_session(use_gpu=False): # test NHWC (default) on CPU x_tf = array_ops.space_to_depth(input_nhwc, block_size) @@ -58,6 +58,12 @@ class SpaceToDepthTest(test.TestCase): x_out = [[[[1, 2, 3, 4]]]] self._testOne(x_np, block_size, x_out) + def testBasicFloat16(self): + x_np = [[[[1], [2]], [[3], [4]]]] + block_size = 2 + x_out = [[[[1, 2, 3, 4]]]] + self._testOne(x_np, block_size, x_out, dtype=dtypes.float16) + # Tests for larger input dimensions. To make sure elements are # correctly ordered spatially. def testLargerInput2x2(self): -- GitLab From 3e6edce1f41a79ca83358b14af9230826e871b66 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Fri, 24 Nov 2017 17:50:04 -0800 Subject: [PATCH 0003/1867] Address `Eigen::Half` -> `Eigen::half` Signed-off-by: Yong Tang --- tensorflow/core/kernels/spacetodepth_op.cc | 4 ++-- tensorflow/core/kernels/spacetodepth_op_gpu.cu.cc | 6 +++--- 2 files changed, 5 insertions(+), 5 deletions(-) diff --git a/tensorflow/core/kernels/spacetodepth_op.cc b/tensorflow/core/kernels/spacetodepth_op.cc index d93a2a9bad..e59adfc6ac 100644 --- a/tensorflow/core/kernels/spacetodepth_op.cc +++ b/tensorflow/core/kernels/spacetodepth_op.cc @@ -188,8 +188,8 @@ REGISTER_KERNEL_BUILDER( Name("SpaceToDepth").Device(DEVICE_GPU).TypeConstraint("T"), SpaceToDepthOp); REGISTER_KERNEL_BUILDER( - Name("SpaceToDepth").Device(DEVICE_GPU).TypeConstraint("T"), - SpaceToDepthOp); + Name("SpaceToDepth").Device(DEVICE_GPU).TypeConstraint("T"), + SpaceToDepthOp); REGISTER_KERNEL_BUILDER( Name("SpaceToDepth").Device(DEVICE_GPU).TypeConstraint("T"), SpaceToDepthOp); diff --git a/tensorflow/core/kernels/spacetodepth_op_gpu.cu.cc b/tensorflow/core/kernels/spacetodepth_op_gpu.cu.cc index e841472972..8466fa192f 100644 --- a/tensorflow/core/kernels/spacetodepth_op_gpu.cu.cc +++ b/tensorflow/core/kernels/spacetodepth_op_gpu.cu.cc @@ -225,9 +225,9 @@ struct SpaceToDepthOpFunctor { template struct functor::SpaceToDepthOpFunctor; template struct functor::SpaceToDepthOpFunctor; -// Instantiate the GPU implementations for Eigen::Half. -template struct functor::SpaceToDepthOpFunctor; -template struct functor::SpaceToDepthOpFunctor; +// Instantiate the GPU implementations for Eigen::half. +template struct functor::SpaceToDepthOpFunctor; +template struct functor::SpaceToDepthOpFunctor; // NCHW_VECT_C with 4 x qint8 can be treated as NCHW int32. template struct functor::SpaceToDepthOpFunctor; -- GitLab From 17b982cad07799feeb00614b0faeba4cf95474c2 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Sat, 25 Nov 2017 17:33:43 -0800 Subject: [PATCH 0004/1867] Add DT_HALF support for DepthToSpace on GPU Signed-off-by: Yong Tang --- tensorflow/core/kernels/depthtospace_op.cc | 3 +++ tensorflow/core/kernels/depthtospace_op_gpu.cu.cc | 4 ++++ 2 files changed, 7 insertions(+) diff --git a/tensorflow/core/kernels/depthtospace_op.cc b/tensorflow/core/kernels/depthtospace_op.cc index 39aa3e9eb0..b74a09e2cb 100644 --- a/tensorflow/core/kernels/depthtospace_op.cc +++ b/tensorflow/core/kernels/depthtospace_op.cc @@ -187,6 +187,9 @@ TF_CALL_ALL_TYPES(REGISTER); REGISTER_KERNEL_BUILDER( Name("DepthToSpace").Device(DEVICE_GPU).TypeConstraint("T"), DepthToSpaceOp); +REGISTER_KERNEL_BUILDER( + Name("DepthToSpace").Device(DEVICE_GPU).TypeConstraint("T"), + DepthToSpaceOp); REGISTER_KERNEL_BUILDER( Name("DepthToSpace").Device(DEVICE_GPU).TypeConstraint("T"), DepthToSpaceOp); diff --git a/tensorflow/core/kernels/depthtospace_op_gpu.cu.cc b/tensorflow/core/kernels/depthtospace_op_gpu.cu.cc index 7a66285383..2d39abce16 100644 --- a/tensorflow/core/kernels/depthtospace_op_gpu.cu.cc +++ b/tensorflow/core/kernels/depthtospace_op_gpu.cu.cc @@ -229,6 +229,10 @@ struct DepthToSpaceOpFunctor { template struct functor::DepthToSpaceOpFunctor; template struct functor::DepthToSpaceOpFunctor; +// Instantiate the GPU implementations for Eigen::half. +template struct functor::DepthToSpaceOpFunctor; +template struct functor::DepthToSpaceOpFunctor; + // NCHW_VECT_C with 4 x qint8 can be treated as NCHW int32. template struct functor::DepthToSpaceOpFunctor; -- GitLab From 1100256692a2b130f3ef2b4e36cd5b63241672ce Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Sat, 25 Nov 2017 17:34:14 -0800 Subject: [PATCH 0005/1867] Add test cases for DT_HALF support with DepthToSpace on GPU. Signed-off-by: Yong Tang --- tensorflow/python/kernel_tests/depthtospace_op_test.py | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/tensorflow/python/kernel_tests/depthtospace_op_test.py b/tensorflow/python/kernel_tests/depthtospace_op_test.py index 7df2366954..f03ad85f17 100644 --- a/tensorflow/python/kernel_tests/depthtospace_op_test.py +++ b/tensorflow/python/kernel_tests/depthtospace_op_test.py @@ -35,8 +35,8 @@ from tensorflow.python.platform import tf_logging class DepthToSpaceTest(test.TestCase): - def _testOne(self, inputs, block_size, outputs): - input_nhwc = math_ops.to_float(inputs) + def _testOne(self, inputs, block_size, outputs, dtype=dtypes.float32): + input_nhwc = math_ops.cast(inputs, dtype) with self.test_session(use_gpu=False): # test NHWC (default) on CPU x_tf = array_ops.depth_to_space(input_nhwc, block_size) @@ -59,6 +59,12 @@ class DepthToSpaceTest(test.TestCase): x_out = [[[[1], [2]], [[3], [4]]]] self._testOne(x_np, block_size, x_out) + def testBasicFloat16(self): + x_np = [[[[1, 2, 3, 4]]]] + block_size = 2 + x_out = [[[[1], [2]], [[3], [4]]]] + self._testOne(x_np, block_size, x_out, dtype=dtypes.float16) + # Tests for larger input dimensions. To make sure elements are # correctly ordered spatially. def testBlockSize2(self): -- GitLab From 7a590cd8ea21ae085845efc6d9b1724d42800659 Mon Sep 17 00:00:00 2001 From: Benoit Steiner Date: Fri, 19 Jan 2018 19:13:43 -0800 Subject: [PATCH 0006/1867] Turn the op_performance_data proto lib into a header only library by default PiperOrigin-RevId: 182621348 Signed-off-by: Jie --- tensorflow/core/BUILD | 6 +++-- tensorflow/core/grappler/costs/BUILD | 24 +++++++++---------- .../core/platform/default/build_config.bzl | 8 +++++++ tensorflow/python/BUILD | 4 ++-- 4 files changed, 26 insertions(+), 16 deletions(-) diff --git a/tensorflow/core/BUILD b/tensorflow/core/BUILD index 579174efa3..f2f66fc567 100644 --- a/tensorflow/core/BUILD +++ b/tensorflow/core/BUILD @@ -136,6 +136,8 @@ load( "tf_nano_proto_library", "tf_protos_all", "tf_protos_all_impl", + "tf_protos_grappler", + "tf_protos_grappler_impl", ) load( "//tensorflow/core:platform/default/build_config_root.bzl", @@ -1529,7 +1531,7 @@ cc_library( "@snappy", "@zlib_archive//:zlib", "@protobuf_archive//:protobuf", - ] + tf_protos_all_impl(), + ] + tf_protos_all_impl() + tf_protos_grappler_impl(), ) # File compiled with extra flags to get cpu-specific acceleration. @@ -2094,7 +2096,7 @@ tf_cuda_library( ":core_cpu_base", ":proto_text", "//tensorflow/core/grappler:grappler_item", - ] + if_static([":core_cpu_impl"]) + tf_protos_all(), + ] + if_static([":core_cpu_impl"]) + tf_protos_all() + tf_protos_grappler(), ) tf_cuda_library( diff --git a/tensorflow/core/grappler/costs/BUILD b/tensorflow/core/grappler/costs/BUILD index 7abc155c19..0fe01e9c9e 100644 --- a/tensorflow/core/grappler/costs/BUILD +++ b/tensorflow/core/grappler/costs/BUILD @@ -1,6 +1,10 @@ licenses(["notice"]) # Apache 2.0 load("//tensorflow:tensorflow.bzl", "tf_cuda_library", "tf_cc_test") +load( + "//tensorflow/core:platform/default/build_config.bzl", + "tf_protos_grappler", +) filegroup( name = "all_files", @@ -37,6 +41,7 @@ tf_proto_library( name = "op_performance_data", srcs = ["op_performance_data.proto"], cc_api_version = 2, + default_header = True, protodeps = tf_additional_all_protos(), visibility = ["//visibility:public"], ) @@ -47,7 +52,6 @@ cc_library( hdrs = ["graph_properties.h"], visibility = ["//visibility:public"], deps = [ - ":op_performance_data_cc", ":utils", "//tensorflow/core:core_cpu_base", "//tensorflow/core:framework", @@ -55,7 +59,7 @@ cc_library( "//tensorflow/core/grappler:grappler_item", "//tensorflow/core/grappler:utils", "//tensorflow/core/grappler/clusters:cluster", - ], + ] + tf_protos_grappler(), ) tf_cc_test( @@ -135,7 +139,7 @@ tf_cuda_library( hdrs = ["utils.h"], visibility = ["//visibility:public"], deps = [ - ":op_performance_data_cc", + "//third_party/eigen3", "//tensorflow/core:framework", "//tensorflow/core:graph", "//tensorflow/core:lib", @@ -143,8 +147,7 @@ tf_cuda_library( "//tensorflow/core:protos_all_cc", "//tensorflow/core/grappler:utils", "//tensorflow/core/grappler/clusters:utils", - "//third_party/eigen3", - ], + ] + tf_protos_grappler(), ) tf_cc_test( @@ -207,9 +210,8 @@ cc_library( hdrs = ["op_context.h"], visibility = ["//visibility:public"], deps = [ - ":op_performance_data_cc", "//tensorflow/core:protos_all_cc", - ], + ] + tf_protos_grappler(), ) cc_library( @@ -276,12 +278,11 @@ cc_library( deps = [ ":cost_estimator", ":op_context", - ":op_performance_data_cc", + "//third_party/eigen3", "//tensorflow/core:framework", "//tensorflow/core:protos_all_cc", "//tensorflow/core/grappler/clusters:utils", - "//third_party/eigen3", - ], + ] + tf_protos_grappler(), ) tf_cc_test( @@ -305,7 +306,6 @@ cc_library( ":cost_estimator", ":graph_properties", ":op_level_cost_estimator", - ":op_performance_data_cc", ":utils", ":virtual_placer", ":virtual_scheduler", @@ -314,7 +314,7 @@ cc_library( "//tensorflow/core:lib", "//tensorflow/core:protos_all_cc", "//tensorflow/core/grappler:grappler_item", - ], + ] + tf_protos_grappler(), ) tf_cc_test( diff --git a/tensorflow/core/platform/default/build_config.bzl b/tensorflow/core/platform/default/build_config.bzl index e9c510c93c..2102c5cca3 100644 --- a/tensorflow/core/platform/default/build_config.bzl +++ b/tensorflow/core/platform/default/build_config.bzl @@ -378,6 +378,14 @@ def tf_protos_all(): extra_deps=tf_protos_all_impl(), otherwise=["//tensorflow/core:protos_all_cc"]) +def tf_protos_grappler_impl(): + return ["//tensorflow/core/grappler/costs:op_performance_data_cc_impl"] + +def tf_protos_grappler(): + return if_static( + extra_deps=tf_protos_grappler_impl(), + otherwise=["//tensorflow/core/grappler/costs:op_performance_data_cc"]) + def tf_env_time_hdrs(): return [ "platform/env_time.h", diff --git a/tensorflow/python/BUILD b/tensorflow/python/BUILD index 3493ed76f3..dbb29d9878 100644 --- a/tensorflow/python/BUILD +++ b/tensorflow/python/BUILD @@ -32,6 +32,7 @@ load("//tensorflow/core:platform/default/build_config.bzl", "tf_proto_library") load("//tensorflow/core:platform/default/build_config.bzl", "tf_proto_library_py") load("//tensorflow/core:platform/default/build_config.bzl", "tf_additional_lib_deps") load("//tensorflow/core:platform/default/build_config.bzl", "tf_additional_all_protos") +load("//tensorflow/core:platform/default/build_config.bzl", "tf_protos_grappler") load("//tensorflow/core:platform/default/build_config_root.bzl", "tf_additional_plugin_deps") load("//tensorflow/python:build_defs.bzl", "tf_gen_op_wrapper_private_py") load("//tensorflow/core:platform/default/build_config_root.bzl", "tf_additional_verbs_deps") @@ -209,9 +210,8 @@ cc_library( "//tensorflow/core/grappler/costs:analytical_cost_estimator", "//tensorflow/core/grappler/costs:cost_estimator", "//tensorflow/core/grappler/costs:measuring_cost_estimator", - "//tensorflow/core/grappler/costs:op_performance_data_cc", "//tensorflow/core/grappler/costs:utils", - ], + ] + tf_protos_grappler(), ) cc_library( -- GitLab From 550a8fa4e9a29bde527730eb45bcbfb7e9067436 Mon Sep 17 00:00:00 2001 From: Jie Date: Mon, 22 Jan 2018 18:07:49 -0800 Subject: [PATCH 0007/1867] [Update] Refactor optimization pass through grappler tensorflow fixed dependency issues in core/grappler/constant_folding removed python calls for optimization(layout/constfold), moved optimization to convert_graph.cc bug: dependency issue with //tensorflow/core/grappler/clusters:single_machine TODO: shape inference through grappler. cluster for optimization pass. --- tensorflow/contrib/tensorrt/BUILD | 6 +- .../contrib/tensorrt/convert/convert_graph.cc | 56 +++++++++++++++++-- .../contrib/tensorrt/python/trt_convert.py | 36 ++++++------ 3 files changed, 76 insertions(+), 22 deletions(-) diff --git a/tensorflow/contrib/tensorrt/BUILD b/tensorflow/contrib/tensorrt/BUILD index 723c9f5434..1cb916e4c3 100644 --- a/tensorflow/contrib/tensorrt/BUILD +++ b/tensorflow/contrib/tensorrt/BUILD @@ -192,7 +192,11 @@ cc_library( "//tensorflow/core:protos_all_cc", "//tensorflow/core:framework_headers_lib", "//tensorflow/core:core_cpu_base", - #"//third_party/eigen3", + "//tensorflow/core/grappler/optimizers:constant_folding", + "//tensorflow/core/grappler/optimizers:layout_optimizer", + "//tensorflow/core/grappler/clusters:virtual_cluster", + "//tensorflow/core/grappler:devices", + #"//tensorflow/core/grappler/clusters:single_machine", ], ) diff --git a/tensorflow/contrib/tensorrt/convert/convert_graph.cc b/tensorflow/contrib/tensorrt/convert/convert_graph.cc index 29aa555467..c1948c8144 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_graph.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_graph.cc @@ -40,6 +40,15 @@ limitations under the License. #include "tensorflow/core/platform/logging.h" #define _TF_LOG_DEBUG ::tensorflow::internal::LogMessage(__FILE__, __LINE__, -1) +#include "tensorflow/core/grappler/optimizers/constant_folding.h" +#include "tensorflow/core/grappler/optimizers/layout_optimizer.h" +#include "tensorflow/core/grappler/devices.h" +//#include "tensorflow/core/grappler/clusters/single_machine.h" +#include "tensorflow/core/grappler/clusters/virtual_cluster.h" +#include "tensorflow/core/protobuf/device_properties.pb.h" +#include "tensorflow/core/grappler/grappler_item.h" +#include "tensorflow/core/grappler/utils.h" + //------------------------------------------------------------------------------ namespace tensorrt { namespace convert { @@ -199,9 +208,48 @@ tensorflow::Status ConvertGraphDefToTensorRT( const tensorflow::GraphDef& graph_def, const std::vector& output_names, size_t max_batch_size, size_t max_workspace_size, tensorflow::GraphDef* new_graph_def) { + + // optimization pass + tensorflow::grappler::GrapplerItem item; + item.fetch = output_names; + tensorflow::GraphDef gdef; + + // layout optimization + item.graph = graph_def; + tensorflow::grappler::LayoutOptimizer optimizer; + tensorflow::grappler::Cluster* gCluster; + + // virtual cluster + tensorflow::DeviceProperties device_properties; + device_properties.set_type("GPU"); + device_properties.mutable_environment()->insert({"architecture", "6"}); + gCluster = + new tensorflow::grappler::VirtualCluster({{"/GPU:0", device_properties}}); + + // single machine + int num_cpu_cores = tensorflow::grappler::GetNumAvailableLogicalCPUCores(); + int num_gpus = tensorflow::grappler::GetNumAvailableGPUs(); + LOG(DEBUG) << "cpu_cores: " << num_cpu_cores; + LOG(DEBUG) << "gpus: " << num_gpus; + // int timeout_s = 60 * 10; + // gCluster = new tensorflow::grappler::SingleMachine( + // timeout_s, num_cpu_cores, num_gpus); + + tensorflow::Status status = optimizer.Optimize(gCluster, item, &gdef); + + if (status !=tensorflow::Status::OK()) + return status; + + // constant folding + item.graph = gdef; + tensorflow::grappler::ConstantFolding fold(nullptr); + status = fold.Optimize(nullptr, item, &gdef); + if (status !=tensorflow::Status::OK()) + return status; + ShapeMap shape_map; TF_RETURN_IF_ERROR( - tensorflow::trt::inferShapes(graph_def, output_names, shape_map)); + tensorflow::trt::inferShapes(gdef, output_names, shape_map)); std::stringstream oss; for (auto& n : shape_map) { // nodes oss << " Node= " << n.first << ", "; @@ -213,10 +261,10 @@ tensorflow::Status ConvertGraphDefToTensorRT( } // Build full graph tensorflow::FunctionLibraryDefinition flib(tensorflow::OpRegistry::Global(), - graph_def.library()); + gdef.library()); tensorflow::Graph graph(flib); TF_RETURN_IF_ERROR(tensorflow::ConvertGraphDefToGraph( - tensorflow::GraphConstructorOptions(), graph_def, &graph)); + tensorflow::GraphConstructorOptions(), gdef, &graph)); // Segment the graph into subgraphs that can be converted to TensorRT tensorrt::segment::SegmentOptions segment_options; @@ -227,7 +275,7 @@ tensorflow::Status ConvertGraphDefToTensorRT( segment_options.minimum_segment_size = 2; tensorrt::segment::SegmentNodesVector segments; TF_RETURN_IF_ERROR(tensorrt::segment::SegmentGraph( - graph_def, IsTensorRTCandidate, segment_options, &segments)); + gdef, IsTensorRTCandidate, segment_options, &segments)); if (segments.size() > 1) { // LOG(WARNING) << "Multiple TensorRT candidate subgraphs were found, " //<< "but only the first can be converted."; diff --git a/tensorflow/contrib/tensorrt/python/trt_convert.py b/tensorflow/contrib/tensorrt/python/trt_convert.py index a66afa8d05..354f0c8b42 100644 --- a/tensorflow/contrib/tensorrt/python/trt_convert.py +++ b/tensorflow/contrib/tensorrt/python/trt_convert.py @@ -48,25 +48,27 @@ def CreateInferenceGraph(input_graph_def, outputs,max_batch_size=1,max_workspace # output_graph_def_string = trt_convert( # input_graph_def_string,outputs, # max_batch_size,max_workspace_size, status) - g = tf.Graph() - with g.as_default(): - tf.import_graph_def(input_graph_def, name="") - rewriter_config = rewriter_config_pb2.RewriterConfig() - rewriter_config.optimizers.append('layout') - rewriter_config.optimizers.append('constfold') + # g = tf.Graph() + # with g.as_default(): + # tf.import_graph_def(input_graph_def, name="") + # rewriter_config = rewriter_config_pb2.RewriterConfig() + # rewriter_config.optimizers.append('layout') + # rewriter_config.optimizers.append('constfold') - # mark output nodes as fetch - train_op = ops.get_collection_ref(ops.GraphKeys.TRAIN_OP) - for node_name in outputs: - out_node = g.get_operation_by_name(node_name) - for i in range(0,len(out_node.outputs)): - train_op.append(out_node.outputs[0]) + # # mark output nodes as fetch + # train_op = ops.get_collection_ref(ops.GraphKeys.TRAIN_OP) + # for node_name in outputs: + # out_node = g.get_operation_by_name(node_name) + # for i in range(0,len(out_node.outputs)): + # train_op.append(out_node.outputs[0]) - # constant folding - mg = meta_graph.create_meta_graph_def(graph=g) - meta_graph.add_collection_def(mg, ops.GraphKeys.TRAIN_OP) - optimized_graph_def_str = \ - tf_optimizer.OptimizeGraph(rewriter_config, mg).SerializeToString() + # # constant folding + # mg = meta_graph.create_meta_graph_def(graph=g) + # meta_graph.add_collection_def(mg, ops.GraphKeys.TRAIN_OP) + # optimized_graph_def_str = \ + # tf_optimizer.OptimizeGraph(rewriter_config, mg).SerializeToString() + + optimized_graph_def_str = input_graph_def.SerializeToString() # TODO(sami): Fix this when we can return status from C++ library # There is a problem with the TF internal library setup that doesn't allow us to return a status object from C++. -- GitLab From da188d378bc6826a8f182b42aa8175a932a0c2f8 Mon Sep 17 00:00:00 2001 From: Jie Date: Tue, 23 Jan 2018 17:23:00 -0800 Subject: [PATCH 0008/1867] [UPDATE] Refactoring shape inference Removed shape refiner and apply shape inference through grappler/costs/graph_properties Currently using static shape inference --- tensorflow/contrib/tensorrt/BUILD | 3 +- .../contrib/tensorrt/convert/convert_graph.cc | 39 +++--- .../contrib/tensorrt/convert/convert_nodes.cc | 24 ++-- .../contrib/tensorrt/convert/convert_nodes.h | 5 +- .../contrib/tensorrt/convert/inferShapes.cc | 125 ------------------ .../contrib/tensorrt/convert/inferShapes.h | 39 ------ 6 files changed, 40 insertions(+), 195 deletions(-) delete mode 100644 tensorflow/contrib/tensorrt/convert/inferShapes.cc delete mode 100644 tensorflow/contrib/tensorrt/convert/inferShapes.h diff --git a/tensorflow/contrib/tensorrt/BUILD b/tensorflow/contrib/tensorrt/BUILD index 1cb916e4c3..f92b60b03a 100644 --- a/tensorflow/contrib/tensorrt/BUILD +++ b/tensorflow/contrib/tensorrt/BUILD @@ -174,12 +174,10 @@ cc_library( "convert/convert_nodes.cc", "convert/convert_graph.cc", "segment/segment.cc", - "convert/inferShapes.cc", ], hdrs=[ "convert/convert_nodes.h", "convert/convert_graph.h", - "convert/inferShapes.h", "segment/segment.h", "segment/union_find.h", ], @@ -196,6 +194,7 @@ cc_library( "//tensorflow/core/grappler/optimizers:layout_optimizer", "//tensorflow/core/grappler/clusters:virtual_cluster", "//tensorflow/core/grappler:devices", + "//tensorflow/core/grappler/costs:graph_properties", #"//tensorflow/core/grappler/clusters:single_machine", ], ) diff --git a/tensorflow/contrib/tensorrt/convert/convert_graph.cc b/tensorflow/contrib/tensorrt/convert/convert_graph.cc index c1948c8144..e90790716c 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_graph.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_graph.cc @@ -28,7 +28,6 @@ limitations under the License. #include "NvInfer.h" #include "tensorflow/contrib/tensorrt/convert/convert_nodes.h" -#include "tensorflow/contrib/tensorrt/convert/inferShapes.h" #include "tensorflow/contrib/tensorrt/segment/segment.h" #include "tensorflow/core/framework/graph.pb.h" #include "tensorflow/core/framework/node_def.pb.h" @@ -49,6 +48,8 @@ limitations under the License. #include "tensorflow/core/grappler/grappler_item.h" #include "tensorflow/core/grappler/utils.h" +#include "tensorflow/core/grappler/costs/graph_properties.h" + //------------------------------------------------------------------------------ namespace tensorrt { namespace convert { @@ -123,7 +124,8 @@ std::unordered_map> BuildTensorNameMap( tensorflow::Status ConvertSubGraphToTensorRT( tensorflow::Graph& graph, const std::vector& output_names, const std::set& subgraph_node_ids, size_t max_batch_size, - size_t max_workspace_size, const ShapeMap& shape_map) { + size_t max_workspace_size, + const tensorflow::grappler::GraphProperties& graph_properties) { tensorflow::EdgeSet subgraph_incoming_edges; GetSubGraphIncomingEdges(graph, subgraph_node_ids, &subgraph_incoming_edges); @@ -161,7 +163,7 @@ tensorflow::Status ConvertSubGraphToTensorRT( tensorflow::NodeDef trt_node_def; TF_RETURN_IF_ERROR(ConvertSubGraphToTensorRTNodeDef( graph, subgraph_node_ids, subgraph_inputs, subgraph_outputs, - max_batch_size, max_workspace_size, shape_map, &trt_node_def)); + max_batch_size, max_workspace_size, graph_properties, &trt_node_def)); tensorflow::Status status; tensorflow::Node* trt_node = graph.AddNode(trt_node_def, &status); @@ -246,19 +248,24 @@ tensorflow::Status ConvertGraphDefToTensorRT( status = fold.Optimize(nullptr, item, &gdef); if (status !=tensorflow::Status::OK()) return status; + + // AJ refactoring shape inference through grappler/GraphProperties. + tensorflow::grappler::GraphProperties static_graph_properties(item); + static_graph_properties.InferStatically(false); + // TF_CHECK_OK(static_graph_prop.InferStatically(false)); + // ShapeMap shape_map; + // TF_RETURN_IF_ERROR( + // tensorflow::trt::inferShapes(gdef, output_names, shape_map)); + // std::stringstream oss; + // for (auto& n : shape_map) { // nodes + // oss << " Node= " << n.first << ", "; + // for (auto o : n.second) { // outputs + // oss << o.first.DebugString() << " T= " << o.second << ", "; + // } + // LOG(DEBUG) << oss.str(); + // oss.str(""); + // } - ShapeMap shape_map; - TF_RETURN_IF_ERROR( - tensorflow::trt::inferShapes(gdef, output_names, shape_map)); - std::stringstream oss; - for (auto& n : shape_map) { // nodes - oss << " Node= " << n.first << ", "; - for (auto o : n.second) { // outputs - oss << o.first.DebugString() << " T= " << o.second << ", "; - } - LOG(DEBUG) << oss.str(); - oss.str(""); - } // Build full graph tensorflow::FunctionLibraryDefinition flib(tensorflow::OpRegistry::Global(), gdef.library()); @@ -291,7 +298,7 @@ tensorflow::Status ConvertGraphDefToTensorRT( } TF_RETURN_IF_ERROR(ConvertSubGraphToTensorRT( graph, output_names, subgraph_node_ids, max_batch_size, - max_workspace_size, shape_map)); + max_workspace_size, static_graph_properties)); } graph.ToGraphDef(new_graph_def); return tensorflow::Status::OK(); diff --git a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc index 83f78d7eff..6c77cdc0b6 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc @@ -1548,7 +1548,8 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( const tensorflow::Graph& graph, const std::set& subgraph_node_ids, const std::vector>& input_inds, const std::vector>& output_inds, size_t max_batch_size, - size_t max_workspace_size, const ShapeMap& shape_map, + size_t max_workspace_size, + const tensorflow::grappler::GraphProperties& graph_properties, tensorflow::NodeDef* trt_node) { // Visit nodes in reverse topological order and construct the TRT network. @@ -1605,20 +1606,20 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( input_names.push_back(node_name); // insert original node name without port // TODO(jie): alternative :) // tensorflow::DataType tf_dtype = node->output_type(output_idx); - if (shape_map.count(node_name) == 0) + if (!graph_properties.HasOutputProperties(node_name)) return tensorflow::errors::Internal("failed to find input node: " + node_name); - auto input_entry_vec = shape_map.at(node_name); - if (static_cast(input_entry_vec.size()) < output_idx) + auto op_info_vec = graph_properties.GetOutputProperties(node_name); + if (static_cast(op_info_vec.size()) < output_idx) return tensorflow::errors::Internal( "accessing output index of: " + std::to_string(output_idx) + ", at node: " + node_name + "with output entry from shape_map: " + - std::to_string(input_entry_vec.size())); + std::to_string(op_info_vec.size())); - auto input_entry = input_entry_vec.at(output_idx); + auto op_info = op_info_vec.at(output_idx); - tensorflow::DataType tf_dtype = input_entry.second; + tensorflow::DataType tf_dtype = op_info.dtype(); input_dtypes.push_back(tf_dtype); nvinfer1::DataType dtype(nvinfer1::DataType::kFLOAT); @@ -1627,15 +1628,16 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( LOG(DEBUG) << "accessing output index of: " << std::to_string(output_idx) << ", at node: " << node_name << "with output entry from shape_map: " - << std::to_string(input_entry_vec.size()); + << std::to_string(op_info_vec.size()); + // TODO(ben,jie): update TRT input format/dimension nvinfer1::DimsCHW input_dim_psuedo_chw; for (int i = 0; i < 3; i++) input_dim_psuedo_chw.d[i] = 1; - for (int i = 1; i < input_entry.first.dims(); i++) { + for (int i = 1; i < op_info.shape().dim_size(); i++) { LOG(DEBUG) << "dimension: " << i - << " , size: " << input_entry.first.dim_size(i); - input_dim_psuedo_chw.d[i - 1] = input_entry.first.dim_size(i); + << " , size: " << op_info.shape().dim(i).size(); + input_dim_psuedo_chw.d[i - 1] = op_info.shape().dim(i).size(); } // TODO(ben,jie): proper way to restore input tensor name? diff --git a/tensorflow/contrib/tensorrt/convert/convert_nodes.h b/tensorflow/contrib/tensorrt/convert/convert_nodes.h index a624582dec..dc59c37892 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_nodes.h +++ b/tensorflow/contrib/tensorrt/convert/convert_nodes.h @@ -20,10 +20,10 @@ limitations under the License. #include #include -#include "tensorflow/contrib/tensorrt/convert/inferShapes.h" #include "tensorflow/core/framework/graph.pb.h" #include "tensorflow/core/graph/graph.h" #include "tensorflow/core/lib/core/status.h" +#include "tensorflow/core/grappler/costs/graph_properties.h" namespace tensorrt { namespace convert { @@ -34,7 +34,8 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( input_inds, // {node_id, output_idx} const std::vector>& output_inds, // {node_id, output_idx} - size_t max_batch_size, size_t max_workspace_size, const ShapeMap& shape_map, + size_t max_batch_size, size_t max_workspace_size, + const tensorflow::grappler::GraphProperties& graph_prop, tensorflow::NodeDef* trt_node); } // namespace convert } // namespace tensorrt diff --git a/tensorflow/contrib/tensorrt/convert/inferShapes.cc b/tensorflow/contrib/tensorrt/convert/inferShapes.cc deleted file mode 100644 index c7f0f0023d..0000000000 --- a/tensorflow/contrib/tensorrt/convert/inferShapes.cc +++ /dev/null @@ -1,125 +0,0 @@ -/* Copyright 2015 The TensorFlow Authors. All Rights Reserved. -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at - http://www.apache.org/licenses/LICENSE-2.0 -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. -==============================================================================*/ -#include "tensorflow/contrib/tensorrt/convert/inferShapes.h" -#include -#include "tensorflow/core/common_runtime/shape_refiner.h" -#include "tensorflow/core/framework/node_def.pb.h" -#include "tensorflow/core/framework/shape_inference.h" -#include "tensorflow/core/framework/tensor.h" -#include "tensorflow/core/framework/types.pb_text.h" -#include "tensorflow/core/graph/algorithm.h" -#include "tensorflow/core/graph/graph.h" -#include "tensorflow/core/graph/graph_constructor.h" -#include "tensorflow/core/lib/core/errors.h" -#include "tensorflow/core/lib/core/status.h" -#include "tensorflow/core/platform/logging.h" - -#define _TF_LOG_DEBUG ::tensorflow::internal::LogMessage(__FILE__, __LINE__, -1) - -namespace tensorflow { -namespace trt { -std::vector getTypes(const tensorflow::OpDef& op, - const tensorflow::NodeDef& nd, - bool inp = true) { - const auto& attrMap = nd.attr(); - auto getType = [&attrMap](decltype( - op.input_arg(0)) a) -> std::vector { - std::vector tvec; - if (!a.type_list_attr().empty()) { // get the list types - const auto& tl = attrMap.at(a.type_list_attr()).list(); - int tsize = tl.type_size(); - tvec.reserve(tsize); - for (int t = 0; t < tsize; t++) { - tvec.push_back(tl.type(t)); - } - return tvec; - } - tensorflow::DataType cType = tensorflow::DT_INVALID; - if (a.type() != tensorflow::DT_INVALID) { // get defined types - cType = a.type(); - } else if (!a.type_attr().empty()) { - cType = attrMap.at(a.type_attr()).type(); - } - if (!a.number_attr().empty()) { // numbertypes - int64 nTensors = attrMap.at(a.number_attr()).i(); - tvec = std::vector(nTensors, cType); - return tvec; - } - tvec.push_back(cType); - return tvec; - }; - std::vector types; - if (inp) { - int n_inputs = op.input_arg_size(); - for (int i = 0; i < n_inputs; i++) { - auto tout = getType(op.input_arg(i)); - LOG(DEBUG) << "Node= " << nd.name() << " #inputs" << tout.size(); - types.insert(types.end(), tout.begin(), tout.end()); - } - } else { - int n_outputs = op.output_arg_size(); - // types.resize(n_outputs); - for (int i = 0; i < n_outputs; i++) { - auto tout = getType(op.output_arg(i)); - LOG(DEBUG) << "Node= " << nd.name() << " #outputs" << tout.size(); - types.insert(types.end(), tout.begin(), tout.end()); - } - } - return types; -} - -tensorflow::Status inferShapes(const tensorflow::GraphDef& graph_def, - const std::vector& output_names, - ShapeMap& shapes) { - tensorflow::Graph g(OpRegistry::Global()); - TF_RETURN_IF_ERROR(tensorflow::ConvertGraphDefToGraph( - tensorflow::GraphConstructorOptions(), graph_def, &g)); - std::vector POnodes; - tensorflow::GetPostOrder(g, &POnodes); - tensorflow::ShapeRefiner refiner(graph_def.versions().producer(), - OpRegistry::Global()); - for (auto n = POnodes.rbegin(); n != POnodes.rend(); ++n) { - TF_CHECK_OK(refiner.AddNode(*n)); - } - - auto shape2PTS = [](tensorflow::shape_inference::InferenceContext* ic, - const tensorflow::shape_inference::ShapeHandle& sh) - -> tensorflow::PartialTensorShape { - std::vector dims; - int64 rank = ic->Rank(sh); - for (int64 i = 0; i < rank; i++) { - auto dh = ic->Dim(sh, i); - dims.push_back(ic->Value(dh)); - } - return tensorflow::PartialTensorShape(dims); - }; - for (const auto& n : POnodes) { - auto ic = refiner.GetContext(n); - if (ic) { - int nOuts = ic->num_outputs(); - auto types = getTypes(n->op_def(), n->def(), false); - std::vector< - std::pair> - SAT; - for (int i = 0; i < nOuts; i++) { - auto PTS = shape2PTS(ic, ic->output(i)); - SAT.push_back({PTS, types.at(i)}); - } - shapes[n->name()] = SAT; - } else { - LOG(WARNING) << "Node " << n->name() << " doesn't have InferenceContext!"; - } - } - return tensorflow::Status::OK(); -} -} // namespace trt -} // namespace tensorflow diff --git a/tensorflow/contrib/tensorrt/convert/inferShapes.h b/tensorflow/contrib/tensorrt/convert/inferShapes.h deleted file mode 100644 index b94f1ee893..0000000000 --- a/tensorflow/contrib/tensorrt/convert/inferShapes.h +++ /dev/null @@ -1,39 +0,0 @@ -/* Copyright 2017 The TensorFlow Authors. All Rights Reserved. - -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at - - http://www.apache.org/licenses/LICENSE-2.0 - -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. -==============================================================================*/ -#ifndef TENSORFLOW_CONTRIB_TENSORRT_CONVERT_INFERSHAPES_H_ -#define TENSORFLOW_CONTRIB_TENSORRT_CONVERT_INFERSHAPES_H_ - -#include -#include -#include -#include - -#include "tensorflow/core/framework/graph.pb.h" -#include "tensorflow/core/framework/tensor_shape.h" -#include "tensorflow/core/lib/core/status.h" - -typedef std::unordered_map>> - ShapeMap; -namespace tensorflow { -namespace trt { -tensorflow::Status inferShapes(const tensorflow::GraphDef& graph_def, - const std::vector& output_names, - ShapeMap& shapes); -} -} // namespace tensorflow - -#endif // TENSORFLOW_CONTRIB_TENSORRT_CONVERT_INFERSHAPES_H_ -- GitLab From ccb555f1e7947785763cf65a6713634a85c72607 Mon Sep 17 00:00:00 2001 From: Jie Date: Wed, 24 Jan 2018 16:32:02 -0800 Subject: [PATCH 0009/1867] [BUG_FIX] 'Mean' converter ConvertReduce fixed 1. permutation index 2. output tensor pushed back into map --- tensorflow/contrib/tensorrt/convert/convert_graph.cc | 2 +- tensorflow/contrib/tensorrt/convert/convert_nodes.cc | 9 ++++++--- 2 files changed, 7 insertions(+), 4 deletions(-) diff --git a/tensorflow/contrib/tensorrt/convert/convert_graph.cc b/tensorflow/contrib/tensorrt/convert/convert_graph.cc index e90790716c..16d6e6ec7d 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_graph.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_graph.cc @@ -60,7 +60,7 @@ static std::unordered_set output_nodes; bool IsTensorRTCandidate(const tensorflow::NodeDef& node_def) { static const std::set candidate_ops = { "Identity", "Const", "Conv2D", "MaxPool", "BiasAdd", "Relu", - "Add", "Mul", "Sub", "Rsqrt", "Pad" // "Placeholder" ,"Mean" + "Add", "Mul", "Sub", "Rsqrt", "Pad" , "Mean" // TODO(ben,jie): ... }; if (output_nodes.count(node_def.name())) return false; diff --git a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc index 6c77cdc0b6..6a93edfb47 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc @@ -1334,7 +1334,7 @@ tensorflow::Status ConvertReduce(Converter& ctx, int nbDims = dims.nbDims + 1; TRT_ShapedWeights index_list = inputs.at(1).weights(); - + TFAttrs attrs(node_def); // TODO(jie): handle data type // auto data_type = attrs.get("T"); @@ -1372,7 +1372,9 @@ tensorflow::Status ConvertReduce(Converter& ctx, if (index_list_data[i] == 0) return tensorflow::errors::InvalidArgument("TRT cannot reduce at 0, at" + node_def.name()); - if (index_list_data[i] == 1) permuted_index = 1; + if (index_list_data[i] == 1) + permuted_index = 1; + idx_set.emplace(index_list_data[i]); } @@ -1380,7 +1382,7 @@ tensorflow::Status ConvertReduce(Converter& ctx, nvinfer1::DimsHW pool_kernel; if (permuted_index == 1) { for (int i = 2; i < nbDims; i++) { - if (idx_set.count(i)) { + if (idx_set.count(i)==0) { permuted_index = i; break; } @@ -1415,6 +1417,7 @@ tensorflow::Status ConvertReduce(Converter& ctx, output_tensor = ctx.transposeTensor( const_cast(output_tensor), permutation_order); } + outputs->push_back(TRT_TensorOrWeights(output_tensor)); return tensorflow::Status::OK(); } -- GitLab From e1eb01e5edf1b5814d7f50e8bcdf910c02a49256 Mon Sep 17 00:00:00 2001 From: Sami Kama Date: Wed, 24 Jan 2018 19:29:22 -0800 Subject: [PATCH 0010/1867] Adding Resources for calibration and execution --- tensorflow/contrib/tensorrt/BUILD | 21 ++++++ .../contrib/tensorrt/convert/convert_nodes.cc | 1 + .../tensorrt/resources/TRTInt8Calibrator.cc | 65 +++++++++++++++++++ .../tensorrt/resources/TRTInt8Calibrator.h | 40 ++++++++++++ .../tensorrt/resources/TRTResourceManager.cc | 18 +++++ .../tensorrt/resources/TRTResourceManager.h | 37 +++++++++++ .../contrib/tensorrt/resources/TRTResources.h | 32 +++++++++ 7 files changed, 214 insertions(+) create mode 100644 tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.cc create mode 100644 tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.h create mode 100644 tensorflow/contrib/tensorrt/resources/TRTResourceManager.cc create mode 100644 tensorflow/contrib/tensorrt/resources/TRTResourceManager.h create mode 100644 tensorflow/contrib/tensorrt/resources/TRTResources.h diff --git a/tensorflow/contrib/tensorrt/BUILD b/tensorflow/contrib/tensorrt/BUILD index 1cb916e4c3..37aa573cdb 100644 --- a/tensorflow/contrib/tensorrt/BUILD +++ b/tensorflow/contrib/tensorrt/BUILD @@ -168,6 +168,26 @@ tf_py_wrap_cc( ], ) +cc_library( + name = "trt_resources", + srcs = [ + "resources/TRTInt8Calibrator.cc", + "resources/TRTResourceManager.cc", + ], + hdrs = [ + "resources/TRTInt8Calibrator.h", + "resources/TRTResourceManager.h", + "resources/TRTResources.h", + ], + deps = [ + "@local_config_tensorrt//:tensorrt", + "//tensorflow/core:framework_headers_lib", + "//tensorflow/core:framework_lite", + "//tensorflow/core:core_cpu_base", + + ], +) + cc_library( name= "trt_conversion", srcs=[ @@ -188,6 +208,7 @@ cc_library( "@protobuf_archive//:protobuf_headers", "@nsync//:nsync_headers", ":trt_logging", + ":trt_resources", "//tensorflow/core:framework_lite", "//tensorflow/core:protos_all_cc", "//tensorflow/core:framework_headers_lib", diff --git a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc index 83f78d7eff..3684ac8e78 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc @@ -39,6 +39,7 @@ limitations under the License. #include "tensorflow/core/lib/core/errors.h" #include "tensorflow/core/lib/core/status.h" #include "tensorflow/core/platform/logging.h" +#include "tensorflow/contrib/tensorrt/resources/TRTResourceManager.h" #define _TF_LOG_DEBUG ::tensorflow::internal::LogMessage(__FILE__, __LINE__, -1) // Check if the types are equal. Cast to int first so that failure log message diff --git a/tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.cc b/tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.cc new file mode 100644 index 0000000000..3c94b52ea6 --- /dev/null +++ b/tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.cc @@ -0,0 +1,65 @@ +// +// Created by skama on 1/24/18. +// + +#include "tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.h" + +#include +#include "tensorflow/core/platform/logging.h" + +namespace tensorflow { +namespace trt { + +int TRTInt8Calibrator::getBatchSize() const { return batch_size_; } + +bool TRTInt8Calibrator::setBatch( + const std::unordered_map& data) { + while (calib_running_.load( + std::memory_order_acquire)) { // wait while calibration is running + tensorflow::mutex_lock l(cond_mtx_); + cond_.wait_for(l, std::chrono::milliseconds(50)); + } + for (const auto it : data) { + auto devptr = dev_buffers_.find(it.first); + if (devptr == dev_buffers_.end()) { + LOG(FATAL) << "FATAL input name '" << it.first + << "' does not match with the buffer names"; + } + const auto& d = devptr->second; + auto status = + cudaMemcpy(d.first, it.second, d.second, cudaMemcpyHostToDevice); + if (status != 0) { + LOG(FATAL) << "cudaMemcpy for '" << it.first << "' failed with " + << status; + } + } + calib_running_.store(true, std::memory_order_release); // release builder + cond_.notify_all(); + return true; +} + +bool TRTInt8Calibrator::getBatch(void** bindings, const char** names, + int nbBindings) { + calib_running_.store(false, std::memory_order_release); // wait for new batch + cond_.notify_all(); + while (!calib_running_.load( + std::memory_order_acquire)) { // wait until new batch arrives + tensorflow::mutex_lock l(cond_mtx_); + cond_.wait_for(l, std::chrono::milliseconds(50)); + } + if (done_) { + return false; + } + for (int i = 0; i < nbBindings; i++) { + auto it = dev_buffers_.find(names[i]); + if (it == dev_buffers_.end()) { + LOG(FATAL) << "Calibration engine asked for unknown tensor name '" + << names[i] << "' at position " << i; + } + bindings[i] = it->second.first; + } + return true; +} + +} // namespace trt +} // namespace tensorflow \ No newline at end of file diff --git a/tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.h b/tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.h new file mode 100644 index 0000000000..b0e904b666 --- /dev/null +++ b/tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.h @@ -0,0 +1,40 @@ +// +// Created by skama on 1/24/18. +// + +#ifndef TENSORFLOW_CONTRIB_TENSORRT_RESOURCES_TRTINT8CALIBRATOR_H_ +#define TENSORFLOW_CONTRIB_TENSORRT_RESOURCES_TRTINT8CALIBRATOR_H_ + +#include +#include +#include +#include +#include +#include "tensorflow/core/platform/mutex.h" +namespace tensorflow { +namespace trt { + +struct TRTInt8Calibrator : public nvinfer1::IInt8Calibrator { + public: + TRTInt8Calibrator(const std::unordered_map< + std::string, std::pair>& dev_buffers, + int batch_size) + : batch_size_(batch_size), + done_(false), + dev_buffers_(dev_buffers), + calib_running_(false){}; + int getBatchSize() const; + bool getBatch(void* bindings[], const char* names[], int nbBindings) override; + bool setBatch(const std::unordered_map &data); + void setDone(){done_=true;} + private: + int batch_size_; + tensorflow::mutex cond_mtx_; + tensorflow::condition_variable cond_; + bool done_; + std::unordered_map> dev_buffers_; + std::atomic_bool calib_running_; +}; +} // namespace trt +} // namespace tensorflow +#endif // TENSORFLOW_CONTRIB_TENSORRT_RESOURCES_TRTINT8CALIBRATOR_H_ diff --git a/tensorflow/contrib/tensorrt/resources/TRTResourceManager.cc b/tensorflow/contrib/tensorrt/resources/TRTResourceManager.cc new file mode 100644 index 0000000000..b060295301 --- /dev/null +++ b/tensorflow/contrib/tensorrt/resources/TRTResourceManager.cc @@ -0,0 +1,18 @@ +// +// Created by skama on 1/23/18. +// + +#include "tensorflow/contrib/tensorrt/resources/TRTResourceManager.h" + + +std::shared_ptr tensorflow::trt::TRTResourceManager::getManager(const std::string &mgr_name) { + // mutex is held for lookup only. Most instantiations where mutex will be held longer + // will be during op creation and should be ok. + tensorflow::mutex_lock lock(map_mutex_); + auto s=managers_.find(mgr_name); + if(s==managers_.end()){ + auto it=managers_.emplace(mgr_name,std::make_shared(mgr_name)); + return it.first->second; + } + return s->second; +} diff --git a/tensorflow/contrib/tensorrt/resources/TRTResourceManager.h b/tensorflow/contrib/tensorrt/resources/TRTResourceManager.h new file mode 100644 index 0000000000..5ec66ab582 --- /dev/null +++ b/tensorflow/contrib/tensorrt/resources/TRTResourceManager.h @@ -0,0 +1,37 @@ +// +// Created by skama on 1/23/18. +// + +#ifndef TENSORFLOW_CONTRIB_TENSORRT_RESOURCES_TRTRESOURCEMANAGER_H_ + +#define TENSORFLOW_CONTRIB_TENSORRT_RESOURCE_TRTRESOURCEMANAGER_H_ +#include + +#include +#include +#include "tensorflow/core/framework/resource_mgr.h" +#include "tensorflow/core/platform/mutex.h" + +namespace tensorflow { +namespace trt { +class TRTResourceManager { + TRTResourceManager() = default; + + public: + static std::shared_ptr instance() { + static std::shared_ptr instance_( + new TRTResourceManager); + return instance_; + } + // returns a manager for given op, if it doesn't exists it creates one + std::shared_ptr getManager( + const std::string& op_name); + + private: + std::unordered_map> + managers_; + tensorflow::mutex map_mutex_; +}; +} // namespace trt +} // namespace tensorflow +#endif // TENSORFLOW_CONTRIB_TENSORRT_RESOURCES_TRTRESOURCEMANAGER_H_ diff --git a/tensorflow/contrib/tensorrt/resources/TRTResources.h b/tensorflow/contrib/tensorrt/resources/TRTResources.h new file mode 100644 index 0000000000..2b65017943 --- /dev/null +++ b/tensorflow/contrib/tensorrt/resources/TRTResources.h @@ -0,0 +1,32 @@ +// +// Created by skama on 1/23/18. +// + +#ifndef TENSORFLOW_CONTRIB_TENSORRT_RESOURCES_TRTRESOURCES_H_ + +#define TENSORFLOW_CONTRIB_TENSORRT_RESOURCES_TRTRESOURCES_H_ + +#include +#include +#include "tensorflow/contrib/tensorrt/resourcemgr/TRTInt8Calibrator.h" +#include "tensorflow/core/framework/resource_mgr.h" + +namespace tensorflow { +namespace trt { + +struct TRTCalibrationResource : public tensorflow::ResourceBase { + TRTCalibrationResource():calibrator(nullptr), builder(nullptr), thr(nullptr){}; + TRTInt8Calibrator* calibrator; + nvinfer1::IBuilder* builder; + std::thread *thr; +}; + +struct TRTEngineResource:public tensorflow::ResourceBase{ + TRTEngineResource():runtime(nullptr), ctx(nullptr){}; + nvinfer1::IRuntime *runtime; + nvinfer1::IExecutionContext *ctx; +}; + +} +} +#endif // TENSORFLOW_CONTRIB_TENSORRT_RESOURCEMGR_TRTRESOURCES_H_ -- GitLab From 6ea7a24c615e7cd9445395539a37e67cb74eede2 Mon Sep 17 00:00:00 2001 From: Jie Date: Thu, 25 Jan 2018 15:14:50 -0800 Subject: [PATCH 0011/1867] [UPDATE] Converter update ConcatV2 AvgPool inception_v1 passed --- .../contrib/tensorrt/convert/convert_graph.cc | 3 +- .../contrib/tensorrt/convert/convert_nodes.cc | 122 +++++++++++++++++- 2 files changed, 123 insertions(+), 2 deletions(-) diff --git a/tensorflow/contrib/tensorrt/convert/convert_graph.cc b/tensorflow/contrib/tensorrt/convert/convert_graph.cc index 16d6e6ec7d..2b6a26491b 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_graph.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_graph.cc @@ -60,7 +60,8 @@ static std::unordered_set output_nodes; bool IsTensorRTCandidate(const tensorflow::NodeDef& node_def) { static const std::set candidate_ops = { "Identity", "Const", "Conv2D", "MaxPool", "BiasAdd", "Relu", - "Add", "Mul", "Sub", "Rsqrt", "Pad" , "Mean" + "Add", "Mul", "Sub", "Rsqrt", "Pad" , "Mean", + "AvgPool", "ConcatV2" // TODO(ben,jie): ... }; if (output_nodes.count(node_def.name())) return false; diff --git a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc index 86c43d960a..ff2e37b7da 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc @@ -1093,6 +1093,8 @@ tensorflow::Status ConvertPool(Converter& ctx, // TODO(jie): support other pooling type if (node_def.op() == "MaxPool") type = nvinfer1::PoolingType::kMAX; + else if (node_def.op() == "AvgPool") + type = nvinfer1::PoolingType::kAVERAGE; else return tensorflow::errors::Unimplemented("only supports Max pool"); @@ -1253,6 +1255,25 @@ tensorflow::Status ConvertConst(Converter& ctx, // weights = ctx.get_temp_weights(dtype, scalar_shape); // std::memcpy(const_cast(weights.values), // weights_tensor.float_val().data(), weights.size_bytes()); + } else if (!weights_tensor.int_val().empty()) { + LOG(DEBUG) << "int!!!" << node_def.name(); + nvinfer1::Dims scalar_shape; + if (tensor.dims() > 0) { + LOG(DEBUG) << "dimensions: " << tensor.dims(); + weights = TRT_ShapedWeights(dtype, weights_tensor.int_val().data(), + get_tensor_shape(tensor)); + } else { + LOG(DEBUG) << "dimensions: " << tensor.dims(); + scalar_shape.nbDims = 1; + scalar_shape.d[0] = 1; + scalar_shape.type[0] = nvinfer1::DimensionType::kSPATIAL; + for (int i = 1; i < nvinfer1::Dims::MAX_DIMS; i++) { + scalar_shape.d[i] = 0; + scalar_shape.type[i] = nvinfer1::DimensionType::kSPATIAL; + } + weights = TRT_ShapedWeights(dtype, weights_tensor.int_val().data(), + scalar_shape); + } } else if (!weights_tensor.tensor_content().empty()) { LOG(DEBUG) << "TENSOR!!!" << node_def.name(); weights = TRT_ShapedWeights(dtype, weights_tensor.tensor_content().data(), @@ -1261,6 +1282,7 @@ tensorflow::Status ConvertConst(Converter& ctx, return tensorflow::errors::Unimplemented( "not supported constant type, at " + node_def.name()); } + // pass the output outputs->push_back(TRT_TensorOrWeights(weights)); return tensorflow::Status::OK(); @@ -1522,19 +1544,115 @@ tensorflow::Status ConvertPad(Converter& ctx, return tensorflow::Status::OK(); } +tensorflow::Status ConvertConcat( + Converter& ctx, tensorflow::NodeDef const& node_def, + std::vector const& inputs, + std::vector* outputs) { + + // not including the last input (axis) here + int input_size = static_cast(inputs.size()) - 1; + + if (!inputs.at(0).is_tensor()) + return tensorflow::errors::InvalidArgument( + "Concat in TRT support only Tensor input, at " + node_def.name()); + + // We are retrieving the axis + TRT_ShapedWeights axis = inputs.at(input_size).weights(); + + TFAttrs attrs(node_def); + auto attr_size = attrs.at("N")->i(); + auto data_type = attrs.get("T"); + auto index_type = attrs.get("Tidx"); + + // TODO(jie): handle data type + // Only expect to handle INT32 as index attributes for now + if (index_type != tensorflow::DataType::DT_INT32) + return tensorflow::errors::Unimplemented("Tidx supports only DT_INT32, at " + + node_def.name()); + + int index = + *(static_cast(const_cast(axis.values_))); + + // TODO(jie): early termination with no-op (attr_size==1) + + auto dim = inputs.at(0).tensor()->getDimensions(); + // dimension check + if (index > dim.nbDims + 1) + return tensorflow::errors::InvalidArgument( + "Concatenate on axis out of dimension range, at " + + node_def.name()); + + if (index == 0) + return tensorflow::errors::InvalidArgument( + "Concatenate on batch dimension not supported, at " + + node_def.name()); + + // incase we need permutation; + std::vector permutation_order(dim.nbDims+1); + + for (int i=0; i inputs_vec; + // Shap chack (all input tensor should have same shape) + // starting from 0 since we are probably also doing transpose here; + for (int i=0; i < input_size; i++) { + auto tensor_i = inputs.at(i).tensor(); + auto dim_i = tensor_i->getDimensions(); + if ( dim_i.nbDims != dim.nbDims ) + return tensorflow::errors::InvalidArgument( + "Concatenate receives inputs with inconsistent dimensions, at " + + node_def.name()); + + for (int j=0; j < dim.nbDims; j++) { + // check dimension consistency on non-concatenate axis + if (j != index-1 && dim_i.d[j] != dim.d[j]) + return tensorflow::errors::InvalidArgument( + "Concatenate receives inputs with inconsistent shape, at" + + node_def.name()); + } + + // TRT does concatenation only on channel! + if (index != 1) + tensor_i = ctx.transposeTensor(const_cast(tensor_i), + permutation_order); + + inputs_vec.push_back(tensor_i); + } + + // nvinfer1::ITensor const* tensor = inputs.at(0).tensor(); + nvinfer1::IConcatenationLayer* layer = ctx.network()->addConcatenation( + const_cast(inputs_vec.data()), + inputs_vec.size()); + nvinfer1::ITensor* output_tensor = layer->getOutput(0); + + if (index != 1) + { + output_tensor= ctx.transposeTensor(output_tensor, permutation_order); + } + outputs->push_back(TRT_TensorOrWeights(output_tensor)); + return tensorflow::Status::OK(); +} + void Converter::register_op_converters() { // vgg_16 slim implementation _op_registry["Placeholder"] = ConvertPlaceholder; _op_registry["Conv2D"] = ConvertConv2D; _op_registry["Relu"] = ConvertActivation; _op_registry["MaxPool"] = ConvertPool; + _op_registry["AvgPool"] = ConvertPool; // This could be really handled as ConvertBinary _op_registry["BiasAdd"] = ConvertScale; _op_registry["Const"] = ConvertConst; // _op_registry["MatMul"] = ConvertFullyConnected; // not used in vgg // TODO(ben,jie): this is a temp hack. _op_registry["Identity"] = ConvertIdentity; // Identity should be removed - // _op_registry["AvgPool"] = ConvertPool; // resnet_50_v1 slim implementation _op_registry["Add"] = ConvertBinary; @@ -1544,6 +1662,8 @@ void Converter::register_op_converters() { _op_registry["Mean"] = ConvertReduce; _op_registry["Pad"] = ConvertPad; // TODO(ben,jie): Add more ops + + _op_registry["ConcatV2"] = ConvertConcat; } } // namespace -- GitLab From cf30a7549e026d5c50117ae011af2b0148a81a89 Mon Sep 17 00:00:00 2001 From: Jie Date: Thu, 25 Jan 2018 17:21:07 -0800 Subject: [PATCH 0012/1867] [UPDATE] Converter update Grouped convolution support added (depthwise as a special case) inception_v2 passed --- .../contrib/tensorrt/convert/convert_graph.cc | 2 +- .../contrib/tensorrt/convert/convert_nodes.cc | 220 +++++++++++------- 2 files changed, 140 insertions(+), 82 deletions(-) diff --git a/tensorflow/contrib/tensorrt/convert/convert_graph.cc b/tensorflow/contrib/tensorrt/convert/convert_graph.cc index 2b6a26491b..c7fa4144b1 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_graph.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_graph.cc @@ -61,7 +61,7 @@ bool IsTensorRTCandidate(const tensorflow::NodeDef& node_def) { static const std::set candidate_ops = { "Identity", "Const", "Conv2D", "MaxPool", "BiasAdd", "Relu", "Add", "Mul", "Sub", "Rsqrt", "Pad" , "Mean", - "AvgPool", "ConcatV2" + "AvgPool", "ConcatV2", "DepthwiseConv2dNative" // TODO(ben,jie): ... }; if (output_nodes.count(node_def.name())) return false; diff --git a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc index ff2e37b7da..ff47cdfe4a 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc @@ -366,15 +366,20 @@ void reorder4(nvinfer1::DimsNCHW shape, T const* idata, } void reorder_rsck_to_kcrs(TRT_ShapedWeights const& iweights, - TRT_ShapedWeights* oweights) { + TRT_ShapedWeights* oweights, int nbGroups) { CHECK_EQ(iweights.type_, oweights->type_); CHECK_EQ(iweights.size_bytes(), oweights->size_bytes()); int r = iweights.shape_.d[0]; int s = iweights.shape_.d[1]; - int c = iweights.shape_.d[2]; - int k = iweights.shape_.d[3]; - oweights->shape_.d[0] = k; - oweights->shape_.d[1] = c; + // TRT requires GKcRS, while TF depthwise has RSCK + // where c=1, C=G + LOG(DEBUG) << "nbGroups: " << nbGroups; + int c = iweights.shape_.d[2]/nbGroups; + LOG(DEBUG) << "c" << iweights.shape_.d[2] << " then " << c; + int k = iweights.shape_.d[3]*nbGroups; + LOG(DEBUG) << "k" << iweights.shape_.d[3] << " then " << k; + oweights->shape_.d[0] = k/nbGroups; + oweights->shape_.d[1] = c*nbGroups; oweights->shape_.d[2] = r; oweights->shape_.d[3] = s; // nvinfer1::DimsNCHW istrides = {1, s, c*r*s, r*s}; @@ -911,87 +916,23 @@ tensorflow::Status BinaryTensorOpWeight( return tensorflow::Status::OK(); } -tensorflow::Status BinaryTensorOpTensor( - Converter& ctx, tensorflow::NodeDef const& node_def, - const nvinfer1::ITensor* tensor_l, const nvinfer1::ITensor* tensor_r, - std::vector* outputs) { - static const std::unordered_map - ops{ - {"Add", nvinfer1::ElementWiseOperation::kSUM}, - {"Mul", nvinfer1::ElementWiseOperation::kPROD}, - // {"max", nvinfer1::ElementWiseOperation::kMAX}, - // {"min", nvinfer1::ElementWiseOperation::kMIN}, - {"Sub", nvinfer1::ElementWiseOperation::kSUB}, - {"Div", nvinfer1::ElementWiseOperation::kDIV}, - }; - - // FIXME assume type matches input weights - // get trt type & shape - TFAttrs attrs(node_def); - // maybe this part has to be moved into the block of rsqrt later - nvinfer1::DataType dtype = attrs.get("T"); - - // check type consistency - CHECK_EQ_TYPE(tensor_l->getType(), dtype); - CHECK_EQ_TYPE(tensor_r->getType(), dtype); - auto op_pair = ops.find(node_def.op()); - if (op_pair == ops.end()) - return tensorflow::errors::Unimplemented( - "binary op: " + node_def.op() + - " not supported at: " + node_def.name()); - - nvinfer1::IElementWiseLayer* layer = ctx.network()->addElementWise( - *const_cast(tensor_l), - *const_cast(tensor_r), op_pair->second); - - nvinfer1::ITensor* output_tensor = layer->getOutput(0); - - // pass the output - outputs->push_back(TRT_TensorOrWeights(output_tensor)); - return tensorflow::Status::OK(); -} - -tensorflow::Status ConvertPlaceholder( - Converter& ctx, tensorflow::NodeDef const& node_def, - std::vector const& inputs, - std::vector* outputs) { - LOG(DEBUG) << "Placeholder should have been replace already"; - return tensorflow::errors::Unimplemented("cannot convert Placeholder op"); - // OK this make sense since we are supposed to replace it with input - TFAttrs attrs(node_def); - nvinfer1::DataType dtype = attrs.get("dtype"); - nvinfer1::Dims dims = attrs.get("shape"); - - dims.nbDims--; - for (int i = 0; i < dims.nbDims; i++) dims.d[i] = dims.d[i + 1]; - - nvinfer1::ITensor* output = - ctx.network()->addInput(node_def.name().c_str(), dtype, dims); - if (!output) { - return tensorflow::errors::InvalidArgument("Failed to create Input layer"); - } - outputs->push_back(TRT_TensorOrWeights(output)); - return tensorflow::Status::OK(); -} +enum class ConvolutionType { + DEFAULT, + DEPTHWISE_CONV +}; -tensorflow::Status ConvertConv2D(Converter& ctx, +tensorflow::Status ConvertConv2DHelper( + Converter& ctx, tensorflow::NodeDef const& node_def, std::vector const& inputs, - std::vector* outputs) { + std::vector* outputs, + int group // group ==0 specifies depthwise conv + ) { nvinfer1::ITensor const* tensor = inputs.at(0).tensor(); - // nvinfer1::ITensor* tensor = inputs.at(0).tensor(); - // TODO(jie): handle NHWC/NCHW transpose; - TRT_ShapedWeights weights_rsck = inputs.at(1).weights(); - TRT_ShapedWeights weights = ctx.get_temp_weights_like(weights_rsck); - reorder_rsck_to_kcrs(weights_rsck, &weights); - TRT_ShapedWeights biases(weights.type_); - int noutput = weights.shape_.d[0]; - nvinfer1::DimsHW kernel_size; - kernel_size.h() = weights.shape_.d[2]; - kernel_size.w() = weights.shape_.d[3]; - LOG(DEBUG) << "kernel size: " << kernel_size.h() << ", " << kernel_size.w(); + TFAttrs attrs(node_def); + int c_index = 1; int h_index = 2; int w_index = 3; auto data_format = attrs.get("data_format"); @@ -1000,17 +941,36 @@ tensorflow::Status ConvertConv2D(Converter& ctx, {0, 3, 1, 2}); h_index = 1; w_index = 2; + c_index = 3; // TODO(jie): transpose it } else { LOG(DEBUG) << "NCHW !!!!"; } + + // tensor after transpose (NCHW) + auto tensor_dim = tensor->getDimensions(); + + int nbGroups = group; + if (nbGroups == 0) // depthwise convolution + nbGroups = tensor_dim.d[0]; + LOG(DEBUG) << "groups count: " << nbGroups; + + TRT_ShapedWeights weights_rsck = inputs.at(1).weights(); + TRT_ShapedWeights weights = ctx.get_temp_weights_like(weights_rsck); + reorder_rsck_to_kcrs(weights_rsck, &weights, nbGroups); + TRT_ShapedWeights biases(weights.type_); + int noutput = weights.shape_.d[0] * nbGroups; + nvinfer1::DimsHW kernel_size; + kernel_size.h() = weights.shape_.d[2]; + kernel_size.w() = weights.shape_.d[3]; + LOG(DEBUG) << "kernel size: " << kernel_size.h() << ", " << kernel_size.w(); + // TODO(jie): stride. (NHWC/NCHW) auto tf_stride = attrs.get>("strides"); LOG(DEBUG) << "h_INDEX" << h_index << ", w_index " << w_index; LOG(DEBUG) << "stride!!!: " << tf_stride[0] << tf_stride[1] << tf_stride[2] << tf_stride[3]; nvinfer1::DimsHW stride(tf_stride[h_index], tf_stride[w_index]); - auto tensor_dim = tensor->getDimensions(); std::vector> padding; // TODO(jie): padding. if (attrs.get("padding") == "SAME") { @@ -1055,6 +1015,7 @@ tensorflow::Status ConvertConv2D(Converter& ctx, layer->setStride(stride); layer->setPadding({padding[0].first, padding[1].first}); layer->setName(node_def.name().c_str()); + layer->setNbGroups(nbGroups); nvinfer1::ITensor* output_tensor = layer->getOutput(0); auto dim_after = output_tensor->getDimensions(); @@ -1071,6 +1032,102 @@ tensorflow::Status ConvertConv2D(Converter& ctx, return tensorflow::Status::OK(); } +tensorflow::Status ConvertConv2DHelper( + Converter& ctx, + tensorflow::NodeDef const& node_def, + std::vector const& inputs, + std::vector* outputs, + ConvolutionType type) { + switch(type) { + case ConvolutionType::DEFAULT: + return ConvertConv2DHelper(ctx, node_def, inputs, outputs, 1); + case ConvolutionType::DEPTHWISE_CONV: + return ConvertConv2DHelper(ctx, node_def, inputs, outputs, 0); + } + return tensorflow::errors::Unimplemented( + "unsupported convolution type at, " + node_def.name()); +} + +tensorflow::Status BinaryTensorOpTensor( + Converter& ctx, tensorflow::NodeDef const& node_def, + const nvinfer1::ITensor* tensor_l, const nvinfer1::ITensor* tensor_r, + std::vector* outputs) { + static const std::unordered_map + ops{ + {"Add", nvinfer1::ElementWiseOperation::kSUM}, + {"Mul", nvinfer1::ElementWiseOperation::kPROD}, + // {"max", nvinfer1::ElementWiseOperation::kMAX}, + // {"min", nvinfer1::ElementWiseOperation::kMIN}, + {"Sub", nvinfer1::ElementWiseOperation::kSUB}, + {"Div", nvinfer1::ElementWiseOperation::kDIV}, + }; + + // FIXME assume type matches input weights + // get trt type & shape + TFAttrs attrs(node_def); + // maybe this part has to be moved into the block of rsqrt later + nvinfer1::DataType dtype = attrs.get("T"); + + // check type consistency + CHECK_EQ_TYPE(tensor_l->getType(), dtype); + CHECK_EQ_TYPE(tensor_r->getType(), dtype); + auto op_pair = ops.find(node_def.op()); + if (op_pair == ops.end()) + return tensorflow::errors::Unimplemented( + "binary op: " + node_def.op() + + " not supported at: " + node_def.name()); + + nvinfer1::IElementWiseLayer* layer = ctx.network()->addElementWise( + *const_cast(tensor_l), + *const_cast(tensor_r), op_pair->second); + + nvinfer1::ITensor* output_tensor = layer->getOutput(0); + + // pass the output + outputs->push_back(TRT_TensorOrWeights(output_tensor)); + return tensorflow::Status::OK(); +} + +tensorflow::Status ConvertPlaceholder( + Converter& ctx, tensorflow::NodeDef const& node_def, + std::vector const& inputs, + std::vector* outputs) { + LOG(DEBUG) << "Placeholder should have been replace already"; + return tensorflow::errors::Unimplemented("cannot convert Placeholder op"); + // OK this make sense since we are supposed to replace it with input + TFAttrs attrs(node_def); + nvinfer1::DataType dtype = attrs.get("dtype"); + nvinfer1::Dims dims = attrs.get("shape"); + + dims.nbDims--; + for (int i = 0; i < dims.nbDims; i++) dims.d[i] = dims.d[i + 1]; + + nvinfer1::ITensor* output = + ctx.network()->addInput(node_def.name().c_str(), dtype, dims); + if (!output) { + return tensorflow::errors::InvalidArgument("Failed to create Input layer"); + } + outputs->push_back(TRT_TensorOrWeights(output)); + return tensorflow::Status::OK(); +} + +tensorflow::Status ConvertConv2D(Converter& ctx, + tensorflow::NodeDef const& node_def, + std::vector const& inputs, + std::vector* outputs) { + return ConvertConv2DHelper(ctx, node_def, inputs, outputs, + ConvolutionType::DEFAULT); +} + +tensorflow::Status ConvertConv2DDepthwise( + Converter& ctx, + tensorflow::NodeDef const& node_def, + std::vector const& inputs, + std::vector* outputs) { + return ConvertConv2DHelper(ctx, node_def, inputs, outputs, + ConvolutionType::DEPTHWISE_CONV); +} + tensorflow::Status ConvertPool(Converter& ctx, tensorflow::NodeDef const& node_def, std::vector const& inputs, @@ -1644,6 +1701,7 @@ void Converter::register_op_converters() { // vgg_16 slim implementation _op_registry["Placeholder"] = ConvertPlaceholder; _op_registry["Conv2D"] = ConvertConv2D; + _op_registry["DepthwiseConv2dNative"] = ConvertConv2DDepthwise; _op_registry["Relu"] = ConvertActivation; _op_registry["MaxPool"] = ConvertPool; _op_registry["AvgPool"] = ConvertPool; -- GitLab From 51ce6cf02c0a445e1a7c89225353ff20fdb538cb Mon Sep 17 00:00:00 2001 From: Jie Date: Tue, 30 Jan 2018 10:43:21 -0800 Subject: [PATCH 0013/1867] [DEBUG] Converter update 1. ConvertConst float length doesn't match tensor shape. handling default broadcast. -> fixed resnet_200 2. Control dependency edge normalizing (remove '^' prefix) -> fixed inception_resnet_v2 --- .../contrib/tensorrt/convert/convert_graph.cc | 2 +- .../contrib/tensorrt/convert/convert_nodes.cc | 39 +++++++++++++------ 2 files changed, 28 insertions(+), 13 deletions(-) diff --git a/tensorflow/contrib/tensorrt/convert/convert_graph.cc b/tensorflow/contrib/tensorrt/convert/convert_graph.cc index c7fa4144b1..185451e28b 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_graph.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_graph.cc @@ -280,7 +280,7 @@ tensorflow::Status ConvertGraphDefToTensorRT( for (auto node : output_names) output_nodes.insert(node); // TODO(sami): this should be passed as a knob!!!! - segment_options.minimum_segment_size = 2; + segment_options.minimum_segment_size = 10; tensorrt::segment::SegmentNodesVector segments; TF_RETURN_IF_ERROR(tensorrt::segment::SegmentGraph( gdef, IsTensorRTCandidate, segment_options, &segments)); diff --git a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc index ff47cdfe4a..6cdfc837fc 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc @@ -437,8 +437,14 @@ class Converter { tensorflow::NodeDef const& node_def) { std::vector inputs; for (auto const& input_name : node_def.input()) { - LOG(DEBUG) << "retrieve input: " << input_name; - inputs.push_back(_trt_tensors.at(input_name)); + std::string name = input_name[0] == '^'? input_name.substr(1) : input_name; + LOG(DEBUG) << "retrieve input: " << name; + if (_trt_tensors.count(name)) { + inputs.push_back(_trt_tensors.at(name)); + } else { + LOG(FATAL) << "input: " << name << "not availabled for node at, " + << node_def.name(); + } } return inputs; } @@ -462,6 +468,7 @@ class Converter { } tensorflow::Status convert_node(tensorflow::NodeDef const& node_def) { + //LOG(DEBUG) << node_def.DebugString(); std::vector inputs = this->get_inputs(node_def); std::string op = node_def.op(); if (!_op_registry.count(op)) { @@ -1292,20 +1299,24 @@ tensorflow::Status ConvertConst(Converter& ctx, nvinfer1::Dims scalar_shape; if (tensor.dims() > 0) { LOG(DEBUG) << "dimensions: " << tensor.dims(); - weights = TRT_ShapedWeights(dtype, weights_tensor.float_val().data(), - get_tensor_shape(tensor)); + scalar_shape = get_tensor_shape(tensor); + if (get_shape_size(scalar_shape) != weights_tensor.float_val_size()) { + LOG(FATAL) << "Broadcast on weights not supported, at: " + << node_def.name(); + } } else { LOG(DEBUG) << "dimensions: " << tensor.dims(); scalar_shape.nbDims = 1; - scalar_shape.d[0] = 1; + // no dimension provided. flatten it + scalar_shape.d[0] = weights_tensor.float_val_size(); scalar_shape.type[0] = nvinfer1::DimensionType::kSPATIAL; for (int i = 1; i < nvinfer1::Dims::MAX_DIMS; i++) { scalar_shape.d[i] = 0; scalar_shape.type[i] = nvinfer1::DimensionType::kSPATIAL; } - weights = TRT_ShapedWeights(dtype, weights_tensor.float_val().data(), - scalar_shape); } + weights = TRT_ShapedWeights(dtype, weights_tensor.float_val().data(), + scalar_shape); // LOG(INFO) << " add: " << weights_tensor.float_val().data(); // LOG(INFO) << " value: " << (*weights_tensor.float_val().data()); @@ -1317,20 +1328,24 @@ tensorflow::Status ConvertConst(Converter& ctx, nvinfer1::Dims scalar_shape; if (tensor.dims() > 0) { LOG(DEBUG) << "dimensions: " << tensor.dims(); - weights = TRT_ShapedWeights(dtype, weights_tensor.int_val().data(), - get_tensor_shape(tensor)); + scalar_shape = get_tensor_shape(tensor); + if (get_shape_size(scalar_shape) != weights_tensor.int_val_size()) { + LOG(FATAL) << "Broadcast on weights not supported, at: " + << node_def.name(); + } } else { LOG(DEBUG) << "dimensions: " << tensor.dims(); scalar_shape.nbDims = 1; - scalar_shape.d[0] = 1; + // no dimension provided. flatten it + scalar_shape.d[0] = weights_tensor.int_val_size(); scalar_shape.type[0] = nvinfer1::DimensionType::kSPATIAL; for (int i = 1; i < nvinfer1::Dims::MAX_DIMS; i++) { scalar_shape.d[i] = 0; scalar_shape.type[i] = nvinfer1::DimensionType::kSPATIAL; } - weights = TRT_ShapedWeights(dtype, weights_tensor.int_val().data(), - scalar_shape); } + weights = TRT_ShapedWeights(dtype, weights_tensor.int_val().data(), + scalar_shape); } else if (!weights_tensor.tensor_content().empty()) { LOG(DEBUG) << "TENSOR!!!" << node_def.name(); weights = TRT_ShapedWeights(dtype, weights_tensor.tensor_content().data(), -- GitLab From 359329893e9db38d08be605bad85c3d3eef1a4cd Mon Sep 17 00:00:00 2001 From: Jie Date: Tue, 30 Jan 2018 21:31:10 -0800 Subject: [PATCH 0014/1867] [Debug + Feature] Feature: input tensor shape inference passing output_edge_map to allow ops absorbed by TRT subgraph to infer shape without running another shape infer Debug: fixed BiasAdd broadcasting Debug: fixed rewiring input edge to TRT_ENGINE_OP TODO: incoming edge check (shape / dimension) TRT dimension requirement for 3.1 makes input tensor with 2 dimension (NC) tricky to interpret. --- .../contrib/tensorrt/convert/convert_graph.cc | 38 +++++- .../contrib/tensorrt/convert/convert_nodes.cc | 118 ++++++++++++++---- .../contrib/tensorrt/convert/convert_nodes.h | 1 + .../contrib/tensorrt/kernels/trt_engine_op.cc | 3 +- 4 files changed, 134 insertions(+), 26 deletions(-) diff --git a/tensorflow/contrib/tensorrt/convert/convert_graph.cc b/tensorflow/contrib/tensorrt/convert/convert_graph.cc index 185451e28b..258a850b21 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_graph.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_graph.cc @@ -77,8 +77,10 @@ void GetSubGraphIncomingEdges(tensorflow::Graph const& graph, for (tensorflow::Edge const* edge : node->in_edges()) { if (!subgraph_node_ids.count(edge->src()->id()) && !edge->src()->IsSource()) { - LOG(DEBUG) << edge->src()->name() << ", "; + LOG(DEBUG) << edge->src()->name() << " Y, "; incoming_edges->insert(edge); + } else { + LOG(DEBUG) << edge->src()->name() << " N, "; } } } @@ -93,7 +95,10 @@ void GetSubGraphOutgoingEdges(tensorflow::Graph const& graph, for (tensorflow::Edge const* edge : node->out_edges()) { if (!subgraph_node_ids.count(edge->dst()->id()) && !edge->dst()->IsSink()) { + LOG(DEBUG) << edge->dst()->name() << " Y, "; outgoing_edges->insert(edge); + } else { + LOG(DEBUG) << edge->dst()->name() << " N, "; } } } @@ -126,6 +131,7 @@ tensorflow::Status ConvertSubGraphToTensorRT( tensorflow::Graph& graph, const std::vector& output_names, const std::set& subgraph_node_ids, size_t max_batch_size, size_t max_workspace_size, + std::unordered_map>* output_edge_map, const tensorflow::grappler::GraphProperties& graph_properties) { tensorflow::EdgeSet subgraph_incoming_edges; GetSubGraphIncomingEdges(graph, subgraph_node_ids, &subgraph_incoming_edges); @@ -164,10 +170,32 @@ tensorflow::Status ConvertSubGraphToTensorRT( tensorflow::NodeDef trt_node_def; TF_RETURN_IF_ERROR(ConvertSubGraphToTensorRTNodeDef( graph, subgraph_node_ids, subgraph_inputs, subgraph_outputs, - max_batch_size, max_workspace_size, graph_properties, &trt_node_def)); + max_batch_size, max_workspace_size, graph_properties, output_edge_map, + &trt_node_def)); tensorflow::Status status; tensorflow::Node* trt_node = graph.AddNode(trt_node_def, &status); + // AddNode does not wire edges. + // Re-map incoming edges to use the new TRT node instead of the orig subgraph + std::map, int> subgraph_edge_to_input_map; + for (size_t i=0; i old_src = {edge->src()->id(), edge->src_output()}; + int new_src_output = subgraph_edge_to_input_map.at(old_src); + graph.AddEdge( + edge->src(), edge->src_output(), trt_node, new_src_output); + graph.RemoveEdge(edge); + } + + + LOG(DEBUG) << "new wiring edges: " << trt_node->in_edges().size(); + for (tensorflow::Edge const* edge : trt_node->in_edges()) { + LOG(DEBUG) << edge->src()->name() << " port: " << edge->src_output(); + } + TF_RETURN_IF_ERROR(status); // Re-map outgoing edges to use the new TRT node instead of the orig subgraph @@ -176,6 +204,7 @@ tensorflow::Status ConvertSubGraphToTensorRT( subgraph_edge_to_output_map.insert({subgraph_outputs.at(i), i}); } TF_RETURN_IF_ERROR(status); + LOG(DEBUG) << "OUT going edge size: " << subgraph_outgoing_edges.size(); for (tensorflow::Edge const* edge : subgraph_outgoing_edges) { std::pair old_src = {edge->src()->id(), edge->src_output()}; int new_src_output = subgraph_edge_to_output_map.at(old_src); @@ -280,7 +309,7 @@ tensorflow::Status ConvertGraphDefToTensorRT( for (auto node : output_names) output_nodes.insert(node); // TODO(sami): this should be passed as a knob!!!! - segment_options.minimum_segment_size = 10; + segment_options.minimum_segment_size = 2; tensorrt::segment::SegmentNodesVector segments; TF_RETURN_IF_ERROR(tensorrt::segment::SegmentGraph( gdef, IsTensorRTCandidate, segment_options, &segments)); @@ -292,6 +321,7 @@ tensorflow::Status ConvertGraphDefToTensorRT( } std::unordered_map node_map; TF_RETURN_IF_ERROR(BuildNodeMap(graph, &node_map)); + std::unordered_map> output_edge_map; for (std::set const& subgraph_node_names : segments) { std::set subgraph_node_ids; for (std::string const& node_name : subgraph_node_names) { @@ -299,7 +329,7 @@ tensorflow::Status ConvertGraphDefToTensorRT( } TF_RETURN_IF_ERROR(ConvertSubGraphToTensorRT( graph, output_names, subgraph_node_ids, max_batch_size, - max_workspace_size, static_graph_properties)); + max_workspace_size, &output_edge_map, static_graph_properties)); } graph.ToGraphDef(new_graph_def); return tensorflow::Status::OK(); diff --git a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc index 6cdfc837fc..bf6a9be8be 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc @@ -437,6 +437,17 @@ class Converter { tensorflow::NodeDef const& node_def) { std::vector inputs; for (auto const& input_name : node_def.input()) { + /************************************************************************* + * TODO(jie) handle case 1) here + * Normalizes the inputs and extracts associated metadata: + * 1) Inputs can contain a colon followed by a suffix of characters. + * That suffix may be a single number (e.g. inputName:1) or several + * word characters separated from a number by a colon + * (e.g. inputName:foo:1). The + * latter case is used to denote inputs and outputs of functions. + * 2) Control dependency inputs contain caret at the beginning and we + * remove this and annotate the edge as a control dependency. + ************************************************************************/ std::string name = input_name[0] == '^'? input_name.substr(1) : input_name; LOG(DEBUG) << "retrieve input: " << name; if (_trt_tensors.count(name)) { @@ -1261,9 +1272,26 @@ tensorflow::Status ConvertScale(Converter& ctx, } else { LOG(DEBUG) << "NCHW !!!!"; } + + auto dims = tensor->getDimensions(); + LOG(DEBUG) << "tensor dimensions: " << dims.nbDims; + for (int i = 0; i < dims.nbDims; i++) { + LOG(DEBUG) << "i: " << dims.d[i]; + } + dims = weights.shape_; + LOG(DEBUG) << "tensor dimensions: " << dims.nbDims; + for (int i = 0; i < dims.nbDims; i++) { + LOG(DEBUG) << "i: " << dims.d[i]; + } + + nvinfer1::ScaleMode mode = nvinfer1::ScaleMode::kCHANNEL; + if (weights.shape_.d[0] == 1) { + mode = nvinfer1::ScaleMode::kUNIFORM; + } + nvinfer1::IScaleLayer* layer = ctx.network()->addScale( - *const_cast(tensor), nvinfer1::ScaleMode::kCHANNEL, - weights, empty_weights, empty_weights); + *const_cast(tensor), mode, weights, + empty_weights, empty_weights); nvinfer1::ITensor* output_tensor = layer->getOutput(0); if (data_format == "NHWC") { @@ -1299,11 +1327,21 @@ tensorflow::Status ConvertConst(Converter& ctx, nvinfer1::Dims scalar_shape; if (tensor.dims() > 0) { LOG(DEBUG) << "dimensions: " << tensor.dims(); + LOG(DEBUG) << "size: " << weights_tensor.float_val_size(); scalar_shape = get_tensor_shape(tensor); + for (int i=0; i < scalar_shape.nbDims; i++) LOG(DEBUG) << scalar_shape.d[i]; if (get_shape_size(scalar_shape) != weights_tensor.float_val_size()) { - LOG(FATAL) << "Broadcast on weights not supported, at: " - << node_def.name(); + if (weights_tensor.float_val_size() == 1 || + scalar_shape.d[0] == weights_tensor.float_val_size()) { + scalar_shape.nbDims = 1; + // no dimension provided. flatten it + scalar_shape.d[0] = weights_tensor.float_val_size(); + scalar_shape.type[0] = nvinfer1::DimensionType::kSPATIAL; + } else { + LOG(FATAL) << "Broadcast on weights only supports kCHANNEL and" + << " kUNIFORM, at: " << node_def.name(); } + } } else { LOG(DEBUG) << "dimensions: " << tensor.dims(); scalar_shape.nbDims = 1; @@ -1330,9 +1368,17 @@ tensorflow::Status ConvertConst(Converter& ctx, LOG(DEBUG) << "dimensions: " << tensor.dims(); scalar_shape = get_tensor_shape(tensor); if (get_shape_size(scalar_shape) != weights_tensor.int_val_size()) { - LOG(FATAL) << "Broadcast on weights not supported, at: " - << node_def.name(); + if (weights_tensor.int_val_size() == 1 || + scalar_shape.d[0] == weights_tensor.int_val_size()) { + scalar_shape.nbDims = 1; + // no dimension provided. flatten it + scalar_shape.d[0] = weights_tensor.int_val_size(); + scalar_shape.type[0] = nvinfer1::DimensionType::kSPATIAL; + } else { + LOG(FATAL) << "Broadcast on weights only supports kCHANNEL and" + << " kUNIFORM, at: " << node_def.name(); } + } } else { LOG(DEBUG) << "dimensions: " << tensor.dims(); scalar_shape.nbDims = 1; @@ -1747,6 +1793,7 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( const std::vector>& output_inds, size_t max_batch_size, size_t max_workspace_size, const tensorflow::grappler::GraphProperties& graph_properties, + std::unordered_map>* output_edge_map, tensorflow::NodeDef* trt_node) { // Visit nodes in reverse topological order and construct the TRT network. @@ -1800,21 +1847,39 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( int output_idx = input.second; tensorflow::Node* node = graph.FindNodeId(node_id); auto node_name = node->name(); - input_names.push_back(node_name); // insert original node name without port + // input_names should use the node name in the graph + // insert original node name without port + input_names.push_back(node_name); + + auto tensor_name = node_name; + if (output_idx != 0) + tensor_name = tensor_name + ":" + std::to_string(output_idx); + + LOG(DEBUG) << "input name: " << node_name << " tensor_name: " << tensor_name << " idx: " << output_idx; + + auto shape_inference_node_name = node_name; + auto shape_inference_output_idx = output_idx; + // rewire the shape inference to original node in the graph + if (output_edge_map->count(tensor_name)) { + shape_inference_node_name = output_edge_map->at(tensor_name).second; + shape_inference_output_idx = output_edge_map->at(tensor_name).first; + } + LOG(DEBUG) << "shapeinference name: " << shape_inference_node_name << " idx: " << shape_inference_output_idx; + // TODO(jie): alternative :) - // tensorflow::DataType tf_dtype = node->output_type(output_idx); - if (!graph_properties.HasOutputProperties(node_name)) + // tensorflow::DataType tf_dtype = node->output_type(); + if (!graph_properties.HasOutputProperties(shape_inference_node_name)) return tensorflow::errors::Internal("failed to find input node: " + - node_name); + shape_inference_node_name); - auto op_info_vec = graph_properties.GetOutputProperties(node_name); - if (static_cast(op_info_vec.size()) < output_idx) + auto op_info_vec = graph_properties.GetOutputProperties(shape_inference_node_name); + if (static_cast(op_info_vec.size()) <= shape_inference_output_idx) return tensorflow::errors::Internal( - "accessing output index of: " + std::to_string(output_idx) + - ", at node: " + node_name + "with output entry from shape_map: " + + "accessing output index of: " + std::to_string(shape_inference_output_idx) + + ", at node: " + shape_inference_node_name + " with output entry from shape_map: " + std::to_string(op_info_vec.size())); - auto op_info = op_info_vec.at(output_idx); + auto op_info = op_info_vec.at(shape_inference_output_idx); tensorflow::DataType tf_dtype = op_info.dtype(); input_dtypes.push_back(tf_dtype); @@ -1822,9 +1887,9 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( nvinfer1::DataType dtype(nvinfer1::DataType::kFLOAT); TF_CHECK_OK(convert_dtype(tf_dtype, &dtype)); - LOG(DEBUG) << "accessing output index of: " << std::to_string(output_idx) - << ", at node: " << node_name - << "with output entry from shape_map: " + LOG(DEBUG) << "accessing output index of: " << std::to_string(shape_inference_output_idx) + << ", at node: " << shape_inference_node_name + << " with output entry from shape_map: " << std::to_string(op_info_vec.size()); // TODO(ben,jie): update TRT input format/dimension @@ -1866,15 +1931,26 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( LOG(DEBUG) << "finished conversion"; + // TODO(sami,ben,jie): proper naming! + static int static_id = 0; + std::string engine_name = "my_trt_op" + std::to_string(static_id++); + // Gather output metadata std::vector output_names; std::vector output_dtypes; + int trt_engine_op_output_idx = 0; for (std::pair const& output : output_inds) { int node_id = output.first; int output_idx = output.second; tensorflow::Node* node = graph.FindNodeId(node_id); std::string op_name = node->name(); std::string tensor_name = op_name; + + output_edge_map->insert( + {trt_engine_op_output_idx == 0 ? + engine_name : engine_name + std::to_string(trt_engine_op_output_idx), + {output_idx, tensor_name}}); + if (output_idx != 0) tensor_name = tensor_name + ":" + std::to_string(output_idx); LOG(DEBUG) << "output tensor name: " << tensor_name; @@ -1923,12 +1999,12 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( LOG(INFO) << "finished engine"; // Build the TRT op - // TODO(sami,ben,jie): proper naming! - static int static_id = 0; tensorflow::NodeDefBuilder op_builder( - "my_trt_op" + std::to_string(static_id++), "TRTEngineOp"); + engine_name, "TRTEngineOp"); std::vector income_edges; + LOG(DEBUG) << "input edge size: " << input_names.size(); for (size_t i = 0; i < input_names.size(); ++i) { + LOG(DEBUG) << "input edges: " << std::to_string(i) << " " << input_names.at(i); int output_idx = input_inds.at(i).second; // we wired up the input here already, it is redundant to do it again in // ConvertSubGraphToTensorRT(convert_graph.cc) diff --git a/tensorflow/contrib/tensorrt/convert/convert_nodes.h b/tensorflow/contrib/tensorrt/convert/convert_nodes.h index dc59c37892..23ca9fcc82 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_nodes.h +++ b/tensorflow/contrib/tensorrt/convert/convert_nodes.h @@ -36,6 +36,7 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( output_inds, // {node_id, output_idx} size_t max_batch_size, size_t max_workspace_size, const tensorflow::grappler::GraphProperties& graph_prop, + std::unordered_map>* output_edge_map, tensorflow::NodeDef* trt_node); } // namespace convert } // namespace tensorrt diff --git a/tensorflow/contrib/tensorrt/kernels/trt_engine_op.cc b/tensorflow/contrib/tensorrt/kernels/trt_engine_op.cc index a1524a592a..445900f08c 100644 --- a/tensorflow/contrib/tensorrt/kernels/trt_engine_op.cc +++ b/tensorflow/contrib/tensorrt/kernels/trt_engine_op.cc @@ -175,7 +175,8 @@ void TRTEngineOp::Compute(OpKernelContext* context) { ->CudaStreamMemberHack())); trt_context_ptr_->enqueue(nbBatch, &buffers[0], *stream, nullptr); - cudaStreamSynchronize(*stream); + // sync should be done by TF. + //cudaStreamSynchronize(*stream); } REGISTER_KERNEL_BUILDER(Name("TRTEngineOp").Device(DEVICE_GPU), TRTEngineOp); -- GitLab From c91050a97b9816627865dd367c93c3ef88ca212f Mon Sep 17 00:00:00 2001 From: Jie Date: Wed, 31 Jan 2018 14:35:49 -0800 Subject: [PATCH 0015/1867] [Feature] subgraph conversion graceful failure conversion failure would result in skipping current subgraph. incoming edge check. require subgraph with incoming edge passing 4 dimensional tensor. TODO binary op -> still needs transpose (since current layout optimization is not working properly --- .../contrib/tensorrt/convert/convert_graph.cc | 17 +++++++++++++---- .../contrib/tensorrt/convert/convert_nodes.cc | 9 ++++++--- 2 files changed, 19 insertions(+), 7 deletions(-) diff --git a/tensorflow/contrib/tensorrt/convert/convert_graph.cc b/tensorflow/contrib/tensorrt/convert/convert_graph.cc index 258a850b21..34a2e9ce6a 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_graph.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_graph.cc @@ -143,6 +143,7 @@ tensorflow::Status ConvertSubGraphToTensorRT( for (tensorflow::Edge const* edge : subgraph_incoming_edges) { subgraph_inputs.push_back({edge->src()->id(), edge->src_output()}); } + std::set> subgraph_outputs_set; // Collect outputs referenced from output_names auto output_name_to_index_map = BuildTensorNameMap(output_names); @@ -168,11 +169,11 @@ tensorflow::Status ConvertSubGraphToTensorRT( subgraph_outputs_set.begin(), subgraph_outputs_set.end()); // Build TensorRT node and add it to the graph tensorflow::NodeDef trt_node_def; + tensorflow::Status status; TF_RETURN_IF_ERROR(ConvertSubGraphToTensorRTNodeDef( graph, subgraph_node_ids, subgraph_inputs, subgraph_outputs, max_batch_size, max_workspace_size, graph_properties, output_edge_map, &trt_node_def)); - tensorflow::Status status; tensorflow::Node* trt_node = graph.AddNode(trt_node_def, &status); // AddNode does not wire edges. @@ -253,6 +254,7 @@ tensorflow::Status ConvertGraphDefToTensorRT( // virtual cluster tensorflow::DeviceProperties device_properties; + device_properties.set_type("GPU"); device_properties.mutable_environment()->insert({"architecture", "6"}); gCluster = @@ -322,14 +324,21 @@ tensorflow::Status ConvertGraphDefToTensorRT( std::unordered_map node_map; TF_RETURN_IF_ERROR(BuildNodeMap(graph, &node_map)); std::unordered_map> output_edge_map; + int count = 0; for (std::set const& subgraph_node_names : segments) { std::set subgraph_node_ids; for (std::string const& node_name : subgraph_node_names) { subgraph_node_ids.insert(node_map.at(node_name)->id()); } - TF_RETURN_IF_ERROR(ConvertSubGraphToTensorRT( - graph, output_names, subgraph_node_ids, max_batch_size, - max_workspace_size, &output_edge_map, static_graph_properties)); + tensorflow::Status status = + ConvertSubGraphToTensorRT(graph, output_names, subgraph_node_ids, + max_batch_size, max_workspace_size, &output_edge_map, + static_graph_properties); + if ( status != tensorflow::Status::OK()) { + LOG(WARNING) << "subgraph conversion error for subgraph_index:" << count + << " due to: \n" << status.ToString() << "SKIPPING......"; + } + count++; } graph.ToGraphDef(new_graph_def); return tensorflow::Status::OK(); diff --git a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc index bf6a9be8be..da6252b25d 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc @@ -1866,8 +1866,6 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( } LOG(DEBUG) << "shapeinference name: " << shape_inference_node_name << " idx: " << shape_inference_output_idx; - // TODO(jie): alternative :) - // tensorflow::DataType tf_dtype = node->output_type(); if (!graph_properties.HasOutputProperties(shape_inference_node_name)) return tensorflow::errors::Internal("failed to find input node: " + shape_inference_node_name); @@ -1885,7 +1883,7 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( input_dtypes.push_back(tf_dtype); nvinfer1::DataType dtype(nvinfer1::DataType::kFLOAT); - TF_CHECK_OK(convert_dtype(tf_dtype, &dtype)); + TF_RETURN_IF_ERROR(convert_dtype(tf_dtype, &dtype)); LOG(DEBUG) << "accessing output index of: " << std::to_string(shape_inference_output_idx) << ", at node: " << shape_inference_node_name @@ -1896,6 +1894,11 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( nvinfer1::DimsCHW input_dim_psuedo_chw; for (int i = 0; i < 3; i++) input_dim_psuedo_chw.d[i] = 1; + // TODO(jie): TRT 3.x only support 4 dimensional input tensor. + // update the code once TRT 4.0 comes out. + if (op_info.shape().dim_size() != 4) + return tensorflow::errors::Unimplemented("require 4 dimensional input"); + for (int i = 1; i < op_info.shape().dim_size(); i++) { LOG(DEBUG) << "dimension: " << i << " , size: " << op_info.shape().dim(i).size(); -- GitLab From 45487b143f890eac31844bfdea171954ddae9e38 Mon Sep 17 00:00:00 2001 From: Jie Date: Wed, 31 Jan 2018 21:13:07 -0800 Subject: [PATCH 0016/1867] [UPDATE] 1. debug binary ops: transpose added again since TF layout optimization is not sufficient 2. debug consecutive trt_engine_op binding names TODO: binding names + input wiring needs refactoring Also change the trt_engine_op attrs (input/output nodes might not be necessary --- .../contrib/tensorrt/convert/convert_nodes.cc | 99 ++++++++++++------- 1 file changed, 63 insertions(+), 36 deletions(-) diff --git a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc index da6252b25d..5df1132f01 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc @@ -449,6 +449,10 @@ class Converter { * remove this and annotate the edge as a control dependency. ************************************************************************/ std::string name = input_name[0] == '^'? input_name.substr(1) : input_name; + auto first = name.find_first_of(':'); + if (first != std::string::npos && first+2 == name.size() && name[first+1]=='0') + name.erase(first); + LOG(DEBUG) << "retrieve input: " << name; if (_trt_tensors.count(name)) { inputs.push_back(_trt_tensors.at(name)); @@ -833,9 +837,12 @@ tensorflow::Status BinaryTensorOpWeight( auto dims_w = weights.shape_; auto dims_t = tensor->getDimensions(); - // default to channel-wise + // default to element-wise auto scale_mode = nvinfer1::ScaleMode::kELEMENTWISE; + // TODO(jie): maybe use a permuatation instead to support more cases; + bool permutation_flag = false; + /* if (weights.count() == 1) { LOG(DEBUG) << "UNIFORM"; @@ -857,44 +864,63 @@ tensorflow::Status BinaryTensorOpWeight( scale_mode = nvinfer1::ScaleMode::kUNIFORM; } else { // no broadcasting on Batch dimension; - assert(dims_w.d[0]==1); - - // broadcasting on Channel dimension only allowed in kUNIFORM - assert(dims_w.d[1]==dims_t.d[0]); - assert(dims_w.nbDims==dims_t.nbDims); - - // default is element; - for (int i=2; i permutation(dims_t.nbDims + 1); - if (scale_mode == nvinfer1::ScaleMode::kCHANNEL && dims_t.nbDims > 1) { - // we swap the last dimension into channel for trt. - // because of tensorflow default broadcasting rules. - for (int i = 0; i < static_cast(permutation.size()); i++) { - permutation[i] = i; + if (permutation_flag) { + if (scale_mode == nvinfer1::ScaleMode::kCHANNEL && dims_t.nbDims > 1) { + // we swap the last dimension into channel for trt. + // because of tensorflow default broadcasting rules. + for (int i = 0; i < static_cast(permutation.size()); i++) { + permutation[i] = i; + } + permutation[1] = dims_t.nbDims; + permutation[dims_t.nbDims] = 1; + tensor = ctx.transposeTensor(const_cast(tensor), + permutation); + } else { + return tensorflow::errors::InvalidArgument( + "Transpose cannot be applied, " + node_def.name()); } - permutation[1] = dims_t.nbDims; - permutation[dims_t.nbDims] = 1; - tensor = ctx.transposeTensor(const_cast(tensor), - permutation); } - */ // prepare weights TRT_ShapedWeights shiftWeights(weights.type_); @@ -923,11 +949,9 @@ tensorflow::Status BinaryTensorOpWeight( nvinfer1::ITensor* output_tensor = layer->getOutput(0); // transpose back dimension - /* - if (scale_mode == nvinfer1::ScaleMode::kCHANNEL && dims_t.nbDims > 1) { + if (permutation_flag) { output_tensor = ctx.transposeTensor(output_tensor, permutation); } - */ // pass the output outputs->push_back(TRT_TensorOrWeights(output_tensor)); @@ -1847,9 +1871,11 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( int output_idx = input.second; tensorflow::Node* node = graph.FindNodeId(node_id); auto node_name = node->name(); + // input_names should use the node name in the graph + // here it should be the input tensor name -> matching the binding // insert original node name without port - input_names.push_back(node_name); + // input_names.push_back(node_name); auto tensor_name = node_name; if (output_idx != 0) @@ -1910,6 +1936,7 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( if (output_idx != 0) input_tensor_name = node_name + ":" + std::to_string(output_idx); + input_names.push_back(input_tensor_name); nvinfer1::ITensor* input_tensor = converter.network()->addInput( input_tensor_name.c_str(), dtype, input_dim_psuedo_chw); @@ -1951,9 +1978,9 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( output_edge_map->insert( {trt_engine_op_output_idx == 0 ? - engine_name : engine_name + std::to_string(trt_engine_op_output_idx), + engine_name : engine_name + ":" + std::to_string(trt_engine_op_output_idx), {output_idx, tensor_name}}); - + trt_engine_op_output_idx++; if (output_idx != 0) tensor_name = tensor_name + ":" + std::to_string(output_idx); LOG(DEBUG) << "output tensor name: " << tensor_name; @@ -1999,7 +2026,7 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( // engine_out << engine_plan_string; // engine_out.close(); - LOG(INFO) << "finished engine"; + LOG(INFO) << "finished engine" << engine_name; // Build the TRT op tensorflow::NodeDefBuilder op_builder( -- GitLab From 10a642da150356d1072e9a5197967f3f3a2bcd7b Mon Sep 17 00:00:00 2001 From: Jie Date: Thu, 1 Feb 2018 07:13:40 -0800 Subject: [PATCH 0017/1867] [UPDATE] converter update: MatMul added TODO: reshape --- .../contrib/tensorrt/convert/convert_graph.cc | 2 +- .../contrib/tensorrt/convert/convert_nodes.cc | 67 ++++++++++++++++++- 2 files changed, 66 insertions(+), 3 deletions(-) diff --git a/tensorflow/contrib/tensorrt/convert/convert_graph.cc b/tensorflow/contrib/tensorrt/convert/convert_graph.cc index 34a2e9ce6a..254a428104 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_graph.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_graph.cc @@ -61,7 +61,7 @@ bool IsTensorRTCandidate(const tensorflow::NodeDef& node_def) { static const std::set candidate_ops = { "Identity", "Const", "Conv2D", "MaxPool", "BiasAdd", "Relu", "Add", "Mul", "Sub", "Rsqrt", "Pad" , "Mean", - "AvgPool", "ConcatV2", "DepthwiseConv2dNative" + "AvgPool", "ConcatV2", "DepthwiseConv2dNative" , "MatMul" // TODO(ben,jie): ... }; if (output_nodes.count(node_def.name())) return false; diff --git a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc index 5df1132f01..6c0ee5e527 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc @@ -347,7 +347,7 @@ template <> tensorflow::DataType TFAttrs::get(std::string key) const { return this->at(key)->type(); } - +// TODO(jie): reorder4 & reorder2 should be merged? template void reorder4(nvinfer1::DimsNCHW shape, T const* idata, nvinfer1::DimsNCHW istrides, T* odata, @@ -365,6 +365,38 @@ void reorder4(nvinfer1::DimsNCHW shape, T const* idata, } } +template +void reorder2(nvinfer1::DimsHW shape, T const* idata, + nvinfer1::DimsHW istrides, T* odata, + nvinfer1::DimsHW ostrides) { + for (int h = 0; h < shape.h(); ++h) { + for (int w = 0; w < shape.w(); ++w) { + odata[h * ostrides.h() + w * ostrides.w()] + = idata[h * ostrides.h() + w * ostrides.w()]; + } + } +} + +// TODO(jie): fail to tensorflow!! +void reorder_ck_to_kc(TRT_ShapedWeights const& iweights, + TRT_ShapedWeights* oweights) { + int c = iweights.shape_.d[0]; + int k = iweights.shape_.d[1]; + oweights->shape_.d[0] = k; + oweights->shape_.d[1] = c; + nvinfer1::DimsHW istrides = {1, k}; + nvinfer1::DimsHW ostrides = {c, 1}; + switch (iweights.type_) { + case tensorflow::DataType::DT_FLOAT: + reorder2( + {k, c}, static_cast(iweights.values_), istrides, + static_cast(const_cast(oweights->values_)), ostrides); + break; + default: + LOG(FATAL) << "!!!!!!!!!!!!!!!!!!!!!!!!broke!!!!!!!!!!!!"; + } +} + void reorder_rsck_to_kcrs(TRT_ShapedWeights const& iweights, TRT_ShapedWeights* oweights, int nbGroups) { CHECK_EQ(iweights.type_, oweights->type_); @@ -382,7 +414,6 @@ void reorder_rsck_to_kcrs(TRT_ShapedWeights const& iweights, oweights->shape_.d[1] = c*nbGroups; oweights->shape_.d[2] = r; oweights->shape_.d[3] = s; - // nvinfer1::DimsNCHW istrides = {1, s, c*r*s, r*s}; nvinfer1::DimsNCHW istrides = {1, k, s * k * c, c * k}; nvinfer1::DimsNCHW ostrides = {c * r * s, r * s, s, 1}; switch (iweights.type_) { @@ -1782,6 +1813,37 @@ tensorflow::Status ConvertConcat( return tensorflow::Status::OK(); } +tensorflow::Status ConvertMatMul( + Converter& ctx, + tensorflow::NodeDef const& node_def, + std::vector const& inputs, + std::vector* outputs) { + nvinfer1::ITensor const* tensor = inputs.at(0).tensor(); + + // TODO(jie): transpose! + TFAttrs attrs(node_def); + //bool transpose_w = bool(attrs->at("transpose_b")->i()); + + // tensor after transpose (NCHW) + auto tensor_dim = tensor->getDimensions(); + + TRT_ShapedWeights weights_ck = inputs.at(1).weights(); + TRT_ShapedWeights weights = ctx.get_temp_weights_like(weights_ck); + reorder_ck_to_kc(weights_ck, &weights); + TRT_ShapedWeights biases(weights.type_); + + int noutput = weights.shape_.d[0]; + + nvinfer1::IFullyConnectedLayer* layer = + ctx.network()->addFullyConnected(*const_cast(tensor), + noutput, weights, biases); + + nvinfer1::ITensor* output_tensor = layer->getOutput(0); + outputs->push_back(TRT_TensorOrWeights(output_tensor)); + return tensorflow::Status::OK(); + +} + void Converter::register_op_converters() { // vgg_16 slim implementation _op_registry["Placeholder"] = ConvertPlaceholder; @@ -1804,6 +1866,7 @@ void Converter::register_op_converters() { _op_registry["Rsqrt"] = ConvertUnary; _op_registry["Mean"] = ConvertReduce; _op_registry["Pad"] = ConvertPad; + _op_registry["MatMul"] = ConvertMatMul; // TODO(ben,jie): Add more ops _op_registry["ConcatV2"] = ConvertConcat; -- GitLab From c5d9369831bfcb66ea54f06349ebae5979c4912d Mon Sep 17 00:00:00 2001 From: Jie Date: Thu, 1 Feb 2018 09:43:24 -0800 Subject: [PATCH 0018/1867] [debug] binary op mode/dimension bug fixed TODO: reshape / debug Matmul --- .../contrib/tensorrt/convert/convert_graph.cc | 3 ++- .../contrib/tensorrt/convert/convert_nodes.cc | 24 ++++++++++++------- 2 files changed, 17 insertions(+), 10 deletions(-) diff --git a/tensorflow/contrib/tensorrt/convert/convert_graph.cc b/tensorflow/contrib/tensorrt/convert/convert_graph.cc index 254a428104..e9ab542f31 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_graph.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_graph.cc @@ -61,7 +61,8 @@ bool IsTensorRTCandidate(const tensorflow::NodeDef& node_def) { static const std::set candidate_ops = { "Identity", "Const", "Conv2D", "MaxPool", "BiasAdd", "Relu", "Add", "Mul", "Sub", "Rsqrt", "Pad" , "Mean", - "AvgPool", "ConcatV2", "DepthwiseConv2dNative" , "MatMul" + "AvgPool", "ConcatV2", "DepthwiseConv2dNative" //, "MatMul", + //"Reshape" // TODO(ben,jie): ... }; if (output_nodes.count(node_def.name())) return false; diff --git a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc index 6c0ee5e527..c697093d12 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc @@ -897,17 +897,22 @@ tensorflow::Status BinaryTensorOpWeight( // no broadcasting on Batch dimension; LOG(DEBUG) << "WEIGHTS DIM: " << dims_w.nbDims << " tensor DIM: " << dims_t.nbDims; - if (dims_w.nbDims==dims_t.nbDims && dims_w.d[0]==1) { - for (int i=1; i Date: Thu, 1 Feb 2018 11:07:53 -0800 Subject: [PATCH 0019/1867] [update] converter update: reshape implemented. I cannot support reshape or MatMul at this moment because of the backend. TODO: wait until TRT 4.0 for backend support on reshape. --- .../contrib/tensorrt/convert/convert_graph.cc | 2 +- .../contrib/tensorrt/convert/convert_nodes.cc | 72 ++++++++++++++++++- 2 files changed, 72 insertions(+), 2 deletions(-) diff --git a/tensorflow/contrib/tensorrt/convert/convert_graph.cc b/tensorflow/contrib/tensorrt/convert/convert_graph.cc index e9ab542f31..573394f309 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_graph.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_graph.cc @@ -61,7 +61,7 @@ bool IsTensorRTCandidate(const tensorflow::NodeDef& node_def) { static const std::set candidate_ops = { "Identity", "Const", "Conv2D", "MaxPool", "BiasAdd", "Relu", "Add", "Mul", "Sub", "Rsqrt", "Pad" , "Mean", - "AvgPool", "ConcatV2", "DepthwiseConv2dNative" //, "MatMul", + "AvgPool", "ConcatV2", "DepthwiseConv2dNative" //, "MatMul", //"Reshape" // TODO(ben,jie): ... }; diff --git a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc index c697093d12..09c1b959ce 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc @@ -1849,6 +1849,76 @@ tensorflow::Status ConvertMatMul( } +tensorflow::Status ConvertReshape(Converter& ctx, + tensorflow::NodeDef const& node_def, + std::vector const& inputs, + std::vector* outputs) { + if (inputs.size() != 2 || !inputs.at(0).is_tensor() || + !inputs.at(1).is_weights()) + return tensorflow::errors::InvalidArgument( + "Input expects tensor and weights, at" + node_def.name()); + + // implement tensor binaryOp weight [channel wise] for now; + nvinfer1::ITensor const* tensor = inputs.at(0).tensor(); + auto dims = tensor->getDimensions(); + // restore implicit batch dimension + int nbDims = dims.nbDims + 1; + + TRT_ShapedWeights shape = inputs.at(1).weights(); + + TFAttrs attrs(node_def); + + auto padding_type = attrs.get("Tshape"); + + if (shape.shape_.nbDims != 1) + return tensorflow::errors::InvalidArgument( + "reshape new shape is not 1 dimensional, at " + node_def.name()); + + // Only expect to handle INT32 as attributes for now + if (padding_type != tensorflow::DataType::DT_INT32) + return tensorflow::errors::Unimplemented( + "reshape new shape supports only DT_INT32, at "+ node_def.name()); + + auto shape_data = static_cast(const_cast(shape.values_)); + + if (shape_data[0] != -1) + return tensorflow::errors::InvalidArgument( + "reshape new shape first dimension is not -1, at "+ node_def.name()); + + auto shape_num_dims = shape.shape_.d[0]; + LOG(DEBUG) << "shape dimensions: " << shape_num_dims; + int volume_w = 1; + for (int i = 1; i < shape.shape_.d[0]; i++) + volume_w *= shape_data[i]; + + int volume_t = 1; + for (int i = 0; i < dims.nbDims; i++) + volume_t *= dims.d[i]; + + LOG(DEBUG) << "volume: " << volume_t << " volume weights: " << volume_w; + if (volume_w != volume_t) + return tensorflow::errors::InvalidArgument( + "volume does not agree between tensor and new shape, at "+ node_def.name()); + + nvinfer1::IShuffleLayer* layer = + ctx.network()->addShuffle(*const_cast(tensor)); + + nvinfer1::Dims reshapeDims; + LOG(DEBUG) << "new dimension: " << shape_num_dims-1; + reshapeDims.nbDims = shape_num_dims-1; + for (int32_t i = 0; i < reshapeDims.nbDims; ++i) { + reshapeDims.d[i] = shape_data[i+1]; + } + layer->setReshapeDimensions(reshapeDims); + LOG(DEBUG) << "new dimension: " << shape_num_dims-1; + + nvinfer1::ITensor* output_tensor = layer->getOutput(0); + auto dims_output = output_tensor->getDimensions(); + LOG(DEBUG) << "output tensor dimension:" << dims_output.nbDims; + outputs->push_back(TRT_TensorOrWeights(output_tensor)); + return tensorflow::Status::OK(); +} + void Converter::register_op_converters() { // vgg_16 slim implementation _op_registry["Placeholder"] = ConvertPlaceholder; @@ -1875,7 +1945,7 @@ void Converter::register_op_converters() { _op_registry["ConcatV2"] = ConvertConcat; _op_registry["MatMul"] = ConvertMatMul; - //_op_registry["Reshape"] = ConvertReshape; + _op_registry["Reshape"] = ConvertReshape; } } // namespace -- GitLab From 28c52d14afb5a54930bcca0db60c9d5068a2c63e Mon Sep 17 00:00:00 2001 From: Sami Kama Date: Mon, 5 Feb 2018 09:55:39 -0800 Subject: [PATCH 0020/1867] WIP: adding int8 calibration --- tensorflow/contrib/tensorrt/BUILD | 4 + .../contrib/tensorrt/convert/convert_graph.cc | 167 +++++++++---- .../contrib/tensorrt/convert/convert_graph.h | 3 +- .../contrib/tensorrt/convert/convert_nodes.cc | 230 ++++++++++++++++-- .../contrib/tensorrt/convert/convert_nodes.h | 38 ++- .../contrib/tensorrt/kernels/trt_calib_op.cc | 68 ++++++ .../contrib/tensorrt/kernels/trt_calib_op.h | 35 +++ .../contrib/tensorrt/kernels/trt_engine_op.cc | 6 +- .../contrib/tensorrt/ops/trt_calib_op.cc | 34 +++ .../contrib/tensorrt/python/trt_convert.py | 4 +- .../tensorrt/resources/TRTInt8Calibrator.cc | 2 +- .../contrib/tensorrt/resources/TRTResources.h | 35 ++- tensorflow/contrib/tensorrt/trt_conversion.i | 9 +- 13 files changed, 543 insertions(+), 92 deletions(-) create mode 100644 tensorflow/contrib/tensorrt/kernels/trt_calib_op.cc create mode 100644 tensorflow/contrib/tensorrt/kernels/trt_calib_op.h create mode 100644 tensorflow/contrib/tensorrt/ops/trt_calib_op.cc diff --git a/tensorflow/contrib/tensorrt/BUILD b/tensorflow/contrib/tensorrt/BUILD index c10e85cffa..bcb8573045 100644 --- a/tensorflow/contrib/tensorrt/BUILD +++ b/tensorflow/contrib/tensorrt/BUILD @@ -60,9 +60,11 @@ tf_kernel_library( name = "trt_engine_op_kernel", srcs = [ "kernels/trt_engine_op.cc", + "kernels/trt_calib_op.cc", ], hdrs=[ "kernels/trt_engine_op.h", + "kernels/trt_calib_op.h", ], gpu_srcs = [ ], @@ -82,6 +84,7 @@ tf_kernel_library( tf_gen_op_libs( op_lib_names = [ "trt_engine_op", + "trt_calib_op", ], deps=[ "@local_config_tensorrt//:tensorrt", @@ -108,6 +111,7 @@ tf_gen_op_wrapper_py( name = "trt_engine_op", deps = [ ":trt_engine_op_op_lib", + ":trt_calib_op_op_lib", ":trt_shape_function", ], ) diff --git a/tensorflow/contrib/tensorrt/convert/convert_graph.cc b/tensorflow/contrib/tensorrt/convert/convert_graph.cc index 16d6e6ec7d..d14abf14dd 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_graph.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_graph.cc @@ -40,9 +40,8 @@ limitations under the License. #define _TF_LOG_DEBUG ::tensorflow::internal::LogMessage(__FILE__, __LINE__, -1) #include "tensorflow/core/grappler/optimizers/constant_folding.h" -#include "tensorflow/core/grappler/optimizers/layout_optimizer.h" +#include "tensorflow/core/grappler/optimizers/layout_optimizer.h" #include "tensorflow/core/grappler/devices.h" -//#include "tensorflow/core/grappler/clusters/single_machine.h" #include "tensorflow/core/grappler/clusters/virtual_cluster.h" #include "tensorflow/core/protobuf/device_properties.pb.h" #include "tensorflow/core/grappler/grappler_item.h" @@ -121,73 +120,146 @@ std::unordered_map> BuildTensorNameMap( return result; } -tensorflow::Status ConvertSubGraphToTensorRT( - tensorflow::Graph& graph, const std::vector& output_names, - const std::set& subgraph_node_ids, size_t max_batch_size, - size_t max_workspace_size, - const tensorflow::grappler::GraphProperties& graph_properties) { - tensorflow::EdgeSet subgraph_incoming_edges; - GetSubGraphIncomingEdges(graph, subgraph_node_ids, &subgraph_incoming_edges); +struct ConvertGraphParams{ + ConvertGraphParams(tensorflow::Graph &graph_, + const std::vector &output_names_, + const std::set& subgraph_node_ids_, + size_t max_batch_size_, + size_t max_workspace_size_, + const tensorflow::grappler::GraphProperties &graph_properties_, + bool int8_ + ):graph(graph_),output_names(output_names_),subgraph_node_ids(subgraph_node_ids_), + max_batch_size(max_batch_size_),max_workspace_size(max_workspace_size_), + graph_properties(graph_properties_),int8(int8_){ + + } - std::vector> subgraph_inputs; + tensorflow::Graph& graph; + const std::vector& output_names; + const std::set& subgraph_node_ids; + size_t max_batch_size; + size_t max_workspace_size; + const tensorflow::grappler::GraphProperties& graph_properties; + bool int8; + std::vector> subgraph_inputs; + std::vector> subgraph_outputs; + tensorflow::EdgeSet subgraph_incoming_edges; + tensorflow::EdgeSet subgraph_outgoing_edges; +}; +tensorflow::Status FillSubGraphEdgeSets(ConvertGraphParams &p){ - // Collect inputs by looking for incoming edges - for (tensorflow::Edge const* edge : subgraph_incoming_edges) { - subgraph_inputs.push_back({edge->src()->id(), edge->src_output()}); + GetSubGraphIncomingEdges(p.graph, p.subgraph_node_ids, &p.subgraph_incoming_edges); + for (tensorflow::Edge const* edge : p.subgraph_incoming_edges) { + p.subgraph_inputs.push_back({edge->src()->id(), edge->src_output()}); } + auto output_name_to_index_map = BuildTensorNameMap(p.output_names); std::set> subgraph_outputs_set; - // Collect outputs referenced from output_names - auto output_name_to_index_map = BuildTensorNameMap(output_names); - // for (int node_id : subgraph_node_ids_no_placeholder) { - for (int node_id : subgraph_node_ids) { - tensorflow::Node* node = graph.FindNodeId(node_id); + + for (int node_id : p.subgraph_node_ids) { + tensorflow::Node* node = p.graph.FindNodeId(node_id); if (output_name_to_index_map.count(node->name())) { for (int index : output_name_to_index_map.at(node->name())) { subgraph_outputs_set.insert({node_id, index}); } } } - // Collect outputs referenced from outgoing edges - tensorflow::EdgeSet subgraph_outgoing_edges; - // GetSubGraphOutgoingEdges(graph, subgraph_node_ids_no_placeholder, - // &subgraph_outgoing_edges); - GetSubGraphOutgoingEdges(graph, subgraph_node_ids, &subgraph_outgoing_edges); - for (tensorflow::Edge const* edge : subgraph_outgoing_edges) { + + GetSubGraphOutgoingEdges(p.graph, p.subgraph_node_ids, &p.subgraph_outgoing_edges); + for (tensorflow::Edge const* edge : p.subgraph_outgoing_edges) { subgraph_outputs_set.insert({edge->src()->id(), edge->src_output()}); } - // Impose an ordering on the outputs - std::vector> subgraph_outputs( + p.subgraph_outputs.reserve(subgraph_outputs_set.size()); + p.subgraph_outputs.insert(p.subgraph_outputs.begin(), subgraph_outputs_set.begin(), subgraph_outputs_set.end()); - // Build TensorRT node and add it to the graph + return tensorflow::Status::OK(); + +}; + +tensorflow::Status GetCalibNode(ConvertGraphParams *params){ + + FillSubGraphEdgeSets(*params); tensorflow::NodeDef trt_node_def; - TF_RETURN_IF_ERROR(ConvertSubGraphToTensorRTNodeDef( - graph, subgraph_node_ids, subgraph_inputs, subgraph_outputs, - max_batch_size, max_workspace_size, graph_properties, &trt_node_def)); + + SubGraphParams s(params->graph, params->subgraph_node_ids, params->subgraph_inputs, params->subgraph_outputs, + params->max_batch_size, params->max_workspace_size, params->graph_properties, &trt_node_def); + TF_RETURN_IF_ERROR(InjectCalibrationNode(s)); tensorflow::Status status; - tensorflow::Node* trt_node = graph.AddNode(trt_node_def, &status); + tensorflow::Node* trt_node = params->graph.AddNode(trt_node_def, &status); + + TF_RETURN_IF_ERROR(status); + + for (auto inp_port: params->subgraph_inputs) { // loop over incoming edges and attach them to calib node + tensorflow::Node * in_node =params->graph.FindNodeId(inp_port.first); + params->graph.UpdateEdge(trt_node, inp_port.second, in_node, inp_port.second); + } + return tensorflow::Status::OK(); +} + +tensorflow::Status ConvertSubGraphToTensorRT(ConvertGraphParams* params ) { + +// tensorflow::EdgeSet subgraph_incoming_edges; +// +// std::vector> subgraph_inputs; +// +// +// // Collect inputs by looking for incoming edges +// for (tensorflow::Edge const* edge : subgraph_incoming_edges) { +// subgraph_inputs.push_back({edge->src()->id(), edge->src_output()}); +// } +// std::set> subgraph_outputs_set; +// // Collect outputs referenced from output_names +// auto output_name_to_index_map = BuildTensorNameMap(output_names); +// for (int node_id : subgraph_node_ids) { +// tensorflow::Node* node = graph.FindNodeId(node_id); +// if (output_name_to_index_map.count(node->name())) { +// for (int index : output_name_to_index_map.at(node->name())) { +// subgraph_outputs_set.insert({node_id, index}); +// } +// } +// } +// // Collect outputs referenced from outgoing edges +// tensorflow::EdgeSet subgraph_outgoing_edges; +// // GetSubGraphOutgoingEdges(graph, subgraph_node_ids_no_placeholder, +// // &subgraph_outgoing_edges); +// GetSubGraphOutgoingEdges(graph, subgraph_node_ids, &subgraph_outgoing_edges); +// for (tensorflow::Edge const* edge : subgraph_outgoing_edges) { +// subgraph_outputs_set.insert({edge->src()->id(), edge->src_output()}); +// } +// // Impose an ordering on the outputs +// std::vector> subgraph_outputs( +// subgraph_outputs_set.begin(), subgraph_outputs_set.end()); +// // Build TensorRT node and add it to the graph + FillSubGraphEdgeSets(*params); + tensorflow::NodeDef trt_node_def; + + SubGraphParams s(params->graph, params->subgraph_node_ids, params->subgraph_inputs, params->subgraph_outputs, + params->max_batch_size, params->max_workspace_size, params->graph_properties, &trt_node_def); + TF_RETURN_IF_ERROR(ConvertSubGraphToTensorRTNodeDef(s)); + tensorflow::Status status; + tensorflow::Node* trt_node = params->graph.AddNode(trt_node_def, &status); TF_RETURN_IF_ERROR(status); // Re-map outgoing edges to use the new TRT node instead of the orig subgraph std::map, int> subgraph_edge_to_output_map; - for (size_t i = 0; i < subgraph_outputs.size(); ++i) { - subgraph_edge_to_output_map.insert({subgraph_outputs.at(i), i}); + for (size_t i = 0; i < params->subgraph_outputs.size(); ++i) { + subgraph_edge_to_output_map.insert({params->subgraph_outputs.at(i), i}); } TF_RETURN_IF_ERROR(status); - for (tensorflow::Edge const* edge : subgraph_outgoing_edges) { + for (tensorflow::Edge const* edge : params->subgraph_outgoing_edges) { std::pair old_src = {edge->src()->id(), edge->src_output()}; int new_src_output = subgraph_edge_to_output_map.at(old_src); - graph.UpdateEdge(trt_node, new_src_output, edge->dst(), edge->dst_input()); + params->graph.UpdateEdge(trt_node, new_src_output, edge->dst(), edge->dst_input()); } // Remove the original subgraph - for (int node_id : subgraph_node_ids) { - tensorflow::Node* node = graph.FindNodeId(node_id); + for (int node_id : params->subgraph_node_ids) { + tensorflow::Node* node = params->graph.FindNodeId(node_id); // Don't remove the input placeholders if (node->type_string() == "Placeholder") { continue; } - graph.RemoveNode(node); + params->graph.RemoveNode(node); } return tensorflow::Status::OK(); } @@ -209,7 +281,9 @@ tensorflow::Status BuildNodeMap( tensorflow::Status ConvertGraphDefToTensorRT( const tensorflow::GraphDef& graph_def, const std::vector& output_names, size_t max_batch_size, - size_t max_workspace_size, tensorflow::GraphDef* new_graph_def) { + size_t max_workspace_size, + tensorflow::GraphDef* new_graph_def, + bool int8=false) { // optimization pass tensorflow::grappler::GrapplerItem item; @@ -246,9 +320,9 @@ tensorflow::Status ConvertGraphDefToTensorRT( item.graph = gdef; tensorflow::grappler::ConstantFolding fold(nullptr); status = fold.Optimize(nullptr, item, &gdef); - if (status !=tensorflow::Status::OK()) + if (status !=tensorflow::Status::OK()) { return status; - + } // AJ refactoring shape inference through grappler/GraphProperties. tensorflow::grappler::GraphProperties static_graph_properties(item); static_graph_properties.InferStatically(false); @@ -296,9 +370,14 @@ tensorflow::Status ConvertGraphDefToTensorRT( for (std::string const& node_name : subgraph_node_names) { subgraph_node_ids.insert(node_map.at(node_name)->id()); } - TF_RETURN_IF_ERROR(ConvertSubGraphToTensorRT( - graph, output_names, subgraph_node_ids, max_batch_size, - max_workspace_size, static_graph_properties)); + + ConvertGraphParams p(graph,output_names,subgraph_node_ids,max_batch_size,max_workspace_size, + static_graph_properties,int8); + if(int8) { + TF_RETURN_IF_ERROR(GetCalibNode(&p)); + } else{ + TF_RETURN_IF_ERROR(ConvertSubGraphToTensorRT(&p)); + } } graph.ToGraphDef(new_graph_def); return tensorflow::Status::OK(); diff --git a/tensorflow/contrib/tensorrt/convert/convert_graph.h b/tensorflow/contrib/tensorrt/convert/convert_graph.h index cd713de888..4ac33cf128 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_graph.h +++ b/tensorflow/contrib/tensorrt/convert/convert_graph.h @@ -27,7 +27,8 @@ namespace convert { tensorflow::Status ConvertGraphDefToTensorRT( const tensorflow::GraphDef& graph_def, const std::vector& output_names, size_t max_batch_size, - size_t max_workspace_size, tensorflow::GraphDef* new_graph_def); + size_t max_workspace_size, + tensorflow::GraphDef* new_graph_def,bool int8); } } // namespace tensorrt diff --git a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc index 86c43d960a..d54c88d9f3 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc @@ -40,6 +40,7 @@ limitations under the License. #include "tensorflow/core/lib/core/status.h" #include "tensorflow/core/platform/logging.h" #include "tensorflow/contrib/tensorrt/resources/TRTResourceManager.h" +#include "tensorflow/contrib/tensorrt/resources/TRTResources.h" #define _TF_LOG_DEBUG ::tensorflow::internal::LogMessage(__FILE__, __LINE__, -1) // Check if the types are equal. Cast to int first so that failure log message @@ -1547,23 +1548,216 @@ void Converter::register_op_converters() { } } // namespace +tensorflow::Status GetTensorRTGraph(tensorrt::convert::SubGraphParams &s ){ + return tensorflow::errors::Unimplemented("Not implemented yet"); +} + +tensorflow::Status InjectCalibrationNode(tensorrt::convert::SubGraphParams &s){ + // Visit nodes in reverse topological order and construct the TRT network. + + // Toposort + std::vector order_vec; + tensorflow::GetPostOrder(s.graph, &order_vec); + // Select just the subgraph + std::list order; + for (tensorflow::Node* node : order_vec) { + if (s.subgraph_node_ids.count(node->id())) { + // order.push_back(node); + order.push_front(node); // we want topological order to contstruct the + // network layer by layer + } + } + // topological order is needed to build TRT network + LOG(DEBUG) << "BUILDING 1"; + static int static_id = 0; + std::string calib_op_name=std::string("my_trt_calib_op_") + std::to_string(static_id++); + + + LOG(DEBUG) << "BUILDING 2"; + auto trt_rmgr=tensorflow::trt::TRTResourceManager::instance(); + auto op_rmgr=trt_rmgr->getManager("TRTCalibOps"); + auto op_res=new tensorflow::trt::TRTCalibrationResource(); + TF_CHECK_OK(op_rmgr->Create(calib_op_name,calib_op_name,op_res)); + op_res->logger=new tensorflow::tensorrt::Logger(); + op_res->builder = nvinfer1::createInferBuilder(*(op_res->logger)); + + if (!op_res->builder) { + return tensorflow::errors::Internal( + "failed to create TensorRT builder object"); + } + + LOG(DEBUG) << "BUILDING 3"; + + op_res->network = op_res->builder->createNetwork(); + if (!op_res->network) { + return tensorflow::errors::Internal( + "failed to create TensorRT network object"); + } + + LOG(DEBUG) << "BUILDING 4"; + + // Build the network + Converter converter(op_res->network); + + LOG(DEBUG) << "BUILDING 5"; + std::vector input_names; + std::vector input_dtypes; + for (std::pair const& input : s.input_inds) { + LOG(DEBUG) << "parsing input!!!!!"; + int node_id = input.first; + int output_idx = input.second; + tensorflow::Node* node = s.graph.FindNodeId(node_id); + auto node_name = node->name(); + input_names.push_back(node_name); // insert original node name without port + // TODO(jie): alternative :) + // tensorflow::DataType tf_dtype = node->output_type(output_idx); + if (!s.graph_properties.HasOutputProperties(node_name)) + return tensorflow::errors::Internal("failed to find input node: " + + node_name); + + auto op_info_vec = s.graph_properties.GetOutputProperties(node_name); + if (static_cast(op_info_vec.size()) < output_idx) + return tensorflow::errors::Internal( + "accessing output index of: " + std::to_string(output_idx) + + ", at node: " + node_name + "with output entry from shape_map: " + + std::to_string(op_info_vec.size())); + + auto op_info = op_info_vec.at(output_idx); + + tensorflow::DataType tf_dtype = op_info.dtype(); + input_dtypes.push_back(tf_dtype); + + nvinfer1::DataType dtype(nvinfer1::DataType::kFLOAT); + TF_CHECK_OK(convert_dtype(tf_dtype, &dtype)); + + LOG(DEBUG) << "accessing output index of: " << std::to_string(output_idx) + << ", at node: " << node_name + << "with output entry from shape_map: " + << std::to_string(op_info_vec.size()); + + // TODO(ben,jie): update TRT input format/dimension + nvinfer1::DimsCHW input_dim_psuedo_chw; + for (int i = 0; i < 3; i++) input_dim_psuedo_chw.d[i] = 1; + + for (int i = 1; i < op_info.shape().dim_size(); i++) { + LOG(DEBUG) << "dimension: " << i + << " , size: " << op_info.shape().dim(i).size(); + input_dim_psuedo_chw.d[i - 1] = op_info.shape().dim(i).size(); + } + + // TODO(ben,jie): proper way to restore input tensor name? + auto input_tensor_name = node_name; + if (output_idx != 0) + input_tensor_name = node_name + ":" + std::to_string(output_idx); + + nvinfer1::ITensor* input_tensor = converter.network()->addInput( + input_tensor_name.c_str(), dtype, input_dim_psuedo_chw); + + if (!input_tensor) + return tensorflow::errors::InvalidArgument( + "Failed to create Input layer"); + LOG(DEBUG) << "input tensor name :" << input_tensor_name; + + if (!converter.insert_input_tensor(input_tensor_name, input_tensor)) + return tensorflow::errors::AlreadyExists( + "output tensor already exists for op: " + input_tensor_name); + } + + LOG(DEBUG) << "finished sorting"; + + for (const tensorflow::Node* node : order) { + tensorflow::NodeDef const& node_def = node->def(); + LOG(DEBUG) << "converting node: " << node_def.name() << " , " + << node_def.op(); + TF_RETURN_IF_ERROR(converter.convert_node(node_def)); + } + + LOG(DEBUG) << "finished conversion"; + + // Gather output metadata + std::vector output_names; + std::vector output_dtypes; + for (std::pair const& output : s.output_inds) { + int node_id = output.first; + int output_idx = output.second; + tensorflow::Node* node = s.graph.FindNodeId(node_id); + std::string op_name = node->name(); + std::string tensor_name = op_name; + if (output_idx != 0) + tensor_name = tensor_name + ":" + std::to_string(output_idx); + LOG(DEBUG) << "output tensor name: " << tensor_name; + output_names.push_back(tensor_name); + auto tensor_or_weights = converter.get_tensor(tensor_name); + if (!tensor_or_weights.is_tensor()) { + return tensorflow::errors::InvalidArgument( + "Output node is weights not tensor"); + } + nvinfer1::ITensor* tensor = tensor_or_weights.tensor(); + if (!tensor) { + return tensorflow::errors::NotFound("Output tensor not found: " + + tensor_name); + } + converter.network()->markOutput(*tensor); + tensorflow::DataType tf_dtype = node->output_type(output_idx); + output_dtypes.push_back(tf_dtype); + nvinfer1::DataType trt_dtype = nvinfer1::DataType::kFLOAT; + TF_RETURN_IF_ERROR(convert_dtype(tf_dtype, &trt_dtype)); + tensor->setType(trt_dtype); + } + + LOG(DEBUG) << "finished output"; + + // Build the engine + op_res->builder->setMaxBatchSize(s.max_batch_size); + op_res->builder->setMaxWorkspaceSize(s.max_workspace_size); + + // Build the TRT op + // TODO(sami,ben,jie): proper naming! + tensorflow::NodeDefBuilder op_builder( + calib_op_name, "TRTCalibOp"); + std::vector income_edges; + for (size_t i = 0; i < input_names.size(); ++i) { + int output_idx = s.input_inds.at(i).second; + // we wired up the input here already, it is redundant to do it again in + // ConvertSubGraphToTensorRT(convert_graph.cc) + auto incoming_edge = tensorflow::NodeDefBuilder::NodeOut(input_names.at(i), + output_idx, input_dtypes.at(i)); + income_edges.push_back(incoming_edge); + } + tensorflow::gtl::ArraySlice + input_list(income_edges); + op_builder.Input(input_list); + std::vector segment_names; + segment_names.reserve(s.subgraph_node_ids.size()); + for(int i : s.subgraph_node_ids){ + auto node=s.graph.FindNodeId(i); + segment_names.push_back(node->name()); + } + LOG(INFO) << "finished op preparation"; + + auto status = op_builder.Attr("segment_names", segment_names ) + .Attr("segment_output_names", output_names) + .Finalize(s.trt_node); + + LOG(INFO) << status.ToString(); + LOG(INFO) << "finished op building"; + + return tensorflow::Status::OK(); + +} tensorflow::Status ConvertSubGraphToTensorRTNodeDef( - const tensorflow::Graph& graph, const std::set& subgraph_node_ids, - const std::vector>& input_inds, - const std::vector>& output_inds, size_t max_batch_size, - size_t max_workspace_size, - const tensorflow::grappler::GraphProperties& graph_properties, - tensorflow::NodeDef* trt_node) { + tensorrt::convert::SubGraphParams &s +) { // Visit nodes in reverse topological order and construct the TRT network. // Toposort std::vector order_vec; - tensorflow::GetPostOrder(graph, &order_vec); + tensorflow::GetPostOrder(s.graph, &order_vec); // Select just the subgraph std::list order; for (tensorflow::Node* node : order_vec) { - if (subgraph_node_ids.count(node->id())) { + if (s.subgraph_node_ids.count(node->id())) { // order.push_back(node); order.push_front(node); // we want topological order to contstruct the // network layer by layer @@ -1601,20 +1795,20 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( LOG(DEBUG) << "BUILDING 5"; std::vector input_names; std::vector input_dtypes; - for (std::pair const& input : input_inds) { + for (std::pair const& input : s.input_inds) { LOG(DEBUG) << "parsing input!!!!!"; int node_id = input.first; int output_idx = input.second; - tensorflow::Node* node = graph.FindNodeId(node_id); + tensorflow::Node* node = s.graph.FindNodeId(node_id); auto node_name = node->name(); input_names.push_back(node_name); // insert original node name without port // TODO(jie): alternative :) // tensorflow::DataType tf_dtype = node->output_type(output_idx); - if (!graph_properties.HasOutputProperties(node_name)) + if (!s.graph_properties.HasOutputProperties(node_name)) return tensorflow::errors::Internal("failed to find input node: " + node_name); - auto op_info_vec = graph_properties.GetOutputProperties(node_name); + auto op_info_vec = s.graph_properties.GetOutputProperties(node_name); if (static_cast(op_info_vec.size()) < output_idx) return tensorflow::errors::Internal( "accessing output index of: " + std::to_string(output_idx) + @@ -1676,10 +1870,10 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( // Gather output metadata std::vector output_names; std::vector output_dtypes; - for (std::pair const& output : output_inds) { + for (std::pair const& output : s.output_inds) { int node_id = output.first; int output_idx = output.second; - tensorflow::Node* node = graph.FindNodeId(node_id); + tensorflow::Node* node = s.graph.FindNodeId(node_id); std::string op_name = node->name(); std::string tensor_name = op_name; if (output_idx != 0) @@ -1707,8 +1901,8 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( LOG(DEBUG) << "finished output"; // Build the engine - trt_builder->setMaxBatchSize(max_batch_size); - trt_builder->setMaxWorkspaceSize(max_workspace_size); + trt_builder->setMaxBatchSize(s.max_batch_size); + trt_builder->setMaxWorkspaceSize(s.max_workspace_size); LOG(INFO) << "starting build engine"; // TODO(ben,jie): half2 and int8 mode support std::string engine_plan_string; @@ -1736,7 +1930,7 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( "my_trt_op" + std::to_string(static_id++), "TRTEngineOp"); std::vector income_edges; for (size_t i = 0; i < input_names.size(); ++i) { - int output_idx = input_inds.at(i).second; + int output_idx = s.input_inds.at(i).second; // we wired up the input here already, it is redundant to do it again in // ConvertSubGraphToTensorRT(convert_graph.cc) auto incoming_edge = tensorflow::NodeDefBuilder::NodeOut(input_names.at(i), @@ -1753,7 +1947,7 @@ tensorflow::Status ConvertSubGraphToTensorRTNodeDef( .Attr("input_nodes", input_names) .Attr("output_nodes", output_names) .Attr("OutT", output_dtypes) - .Finalize(trt_node); + .Finalize(s.trt_node); LOG(INFO) << status.ToString(); LOG(INFO) << "finished op building"; diff --git a/tensorflow/contrib/tensorrt/convert/convert_nodes.h b/tensorflow/contrib/tensorrt/convert/convert_nodes.h index dc59c37892..9f552d0990 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_nodes.h +++ b/tensorflow/contrib/tensorrt/convert/convert_nodes.h @@ -28,15 +28,37 @@ limitations under the License. namespace tensorrt { namespace convert { +struct SubGraphParams{ + SubGraphParams(const tensorflow::Graph &graph_, + const std::set &subgraph_node_ids_, + const std::vector> &input_inds_, + const std::vector> &output_inds_, + size_t max_batch_size_, + size_t max_workspace_size_, + const tensorflow::grappler::GraphProperties &graph_properties_, + tensorflow::NodeDef* trt_node_, + bool int8_=false):graph(graph_), subgraph_node_ids(subgraph_node_ids_), + input_inds(input_inds_),output_inds(output_inds_), + max_batch_size(max_batch_size_), + max_workspace_size(max_workspace_size_), + graph_properties(graph_properties_), + trt_node(trt_node_),int8(int8_){} + + const tensorflow::Graph &graph; + const std::set& subgraph_node_ids; + const std::vector>& input_inds; // {node_id, output_idx} + const std::vector>& output_inds; // {node_id, output_idx} + size_t max_batch_size; + size_t max_workspace_size; + const tensorflow::grappler::GraphProperties& graph_properties; + tensorflow::NodeDef* trt_node; + const bool int8; +}; + tensorflow::Status ConvertSubGraphToTensorRTNodeDef( - const tensorflow::Graph& graph, const std::set& subgraph_node_ids, - const std::vector>& - input_inds, // {node_id, output_idx} - const std::vector>& - output_inds, // {node_id, output_idx} - size_t max_batch_size, size_t max_workspace_size, - const tensorflow::grappler::GraphProperties& graph_prop, - tensorflow::NodeDef* trt_node); + SubGraphParams & params + ); +tensorflow::Status InjectCalibrationNode(SubGraphParams ¶ms); } // namespace convert } // namespace tensorrt diff --git a/tensorflow/contrib/tensorrt/kernels/trt_calib_op.cc b/tensorflow/contrib/tensorrt/kernels/trt_calib_op.cc new file mode 100644 index 0000000000..6fdb583b9a --- /dev/null +++ b/tensorflow/contrib/tensorrt/kernels/trt_calib_op.cc @@ -0,0 +1,68 @@ +// +// Created by skama on 1/25/18. +// + +#include "tensorflow/contrib/tensorrt/kernels/trt_calib_op.h" +#include +#include +#include "tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.h" +#include "tensorflow/contrib/tensorrt/resources/TRTResourceManager.h" +#include "tensorflow/contrib/tensorrt/resources/TRTResources.h" +#include "tensorflow/core/framework/tensor.h" +#include "tensorflow/core/framework/tensor_shape.h" +#include "tensorflow/core/framework/tensor_types.h" +#include "tensorflow/core/framework/types.h" +namespace tensorflow{ +namespace trt{ +TRTCalibOp::TRTCalibOp(OpKernelConstruction* context) : OpKernel(context){ + OP_REQUIRES_OK(context, + context->GetAttr("segment_nodes", &segment_nodes_)); + OP_REQUIRES_OK(context, context->GetAttr("input_names", &input_names_)); + dev_tensors_.resize(segment_nodes_.size()); + +}; + +void TRTCalibOp::Compute(OpKernelContext *ctx) { + auto trt_rm = tensorflow::trt::TRTResourceManager::instance(); + auto resmgr = trt_rm->getManager(name()); + TRTCalibrationResource *calibRes= nullptr; + auto status=resmgr->Lookup(name(), name(), &calibRes); + if (status.ok()){ + int batchSize=ctx->input(0).dim_size(0); + int numInputs=ctx->num_inputs(); + if ( calibRes->calibrator == nullptr){// first run + for(int i = 0 ; i < numInputs; i++){ + const Tensor& t=ctx->input(i); + OP_REQUIRES_OK(ctx, ctx->allocate_persistent(t.dtype(), t.shape(),&dev_tensors_.at(i), nullptr)); + const auto dTensor=dev_tensors_.at(i).AccessTensor(ctx); + CHECK_EQ(t.TotalBytes(),dTensor->TotalBytes()); + auto dType=t.dtype(); + void* devAddr=(void*)dTensor->flat::Type>().data(); + device_buffers_.emplace({input_names_.at(i),std::make_pair(devAddr,dTensor->TotalBytes())}); + } + calibRes->calibrator=new TRTInt8Calibrator(device_buffers_,batchSize); + auto builder=calibRes->builder; + calibRes->thr=new std::thread([calibRes](){ + calibRes->engine=calibRes->builder->buildCudaEngine(*calibRes->network); // will loop until we terminate calibrator + }); + } + std::unordered_map input_data; + for(int i = 0; i < numInputs; i++){ + const Tensor& t = ctx->input(i); + auto dType = t.dtype(); + void* data_address = (void*)t.flat::Type>().data(); + const auto dTensor = dev_tensors_.at(i).AccessTensor(ctx); + CHECK_EQ(t.TotalBytes(), dTensor->TotalBytes()); // use the tensor so FW keeps it + input_data.emplace(input_names_.at(i), data_address); + ctx->set_output(i,t); + } + calibRes->calibrator->setBatch(input_data); + }else{ + ctx->SetStatus(status); + return; + } + +}; + +} +} \ No newline at end of file diff --git a/tensorflow/contrib/tensorrt/kernels/trt_calib_op.h b/tensorflow/contrib/tensorrt/kernels/trt_calib_op.h new file mode 100644 index 0000000000..aefafb29d5 --- /dev/null +++ b/tensorflow/contrib/tensorrt/kernels/trt_calib_op.h @@ -0,0 +1,35 @@ +// +// Created by skama on 1/25/18. +// + +#ifndef TFGITHUB_TRT_CALIB_OP_H +#define TFGITHUB_TRT_CALIB_OP_H + +#include +#include +#include +#include +#include +#include "tensorflow/core/framework/op.h" +#include "tensorflow/core/framework/op_kernel.h" +#include "tensorflow/core/framework/tensor_shape.h" + +namespace tensorflow { +namespace trt { +class TRTCalibOp: public OpKernel { +public: + explicit TRTCalibOp(OpKernelConstruction* context); + + void Compute(OpKernelContext* context) override; + + private: + std::vector segment_nodes_; + std::vector input_names_; + std::vector shapes_; + std::unordered_map> device_buffers_; + std::vector dev_tensors_; + +}; +} +} +#endif //TFGITHUB_TRT_CALIB_OP_H diff --git a/tensorflow/contrib/tensorrt/kernels/trt_engine_op.cc b/tensorflow/contrib/tensorrt/kernels/trt_engine_op.cc index a1524a592a..54b8d0d431 100644 --- a/tensorflow/contrib/tensorrt/kernels/trt_engine_op.cc +++ b/tensorflow/contrib/tensorrt/kernels/trt_engine_op.cc @@ -24,8 +24,8 @@ limitations under the License. namespace tensorflow { static ::tensorflow::tensorrt::Logger gLogger; -using namespace nvinfer1; - +using IRuntime=nvinfer1::IRuntime; +using Dims=nvinfer1::Dims; namespace tensorrt { TRTEngineOp::TRTEngineOp(OpKernelConstruction* context) : OpKernel(context) { @@ -44,7 +44,7 @@ TRTEngineOp::TRTEngineOp(OpKernelConstruction* context) : OpKernel(context) { // TODO(samikama) runtime should be taken from a resourcemanager as well. // Only engine should be in the op and context and runtime should be taken // from resourcemanager - IRuntime* infer = createInferRuntime(gLogger); + IRuntime* infer = nvinfer1::createInferRuntime(gLogger); trt_engine_ptr_.reset(infer->deserializeCudaEngine( serialized_engine.c_str(), serialized_engine.size(), nullptr)); diff --git a/tensorflow/contrib/tensorrt/ops/trt_calib_op.cc b/tensorflow/contrib/tensorrt/ops/trt_calib_op.cc new file mode 100644 index 0000000000..ddf2baa526 --- /dev/null +++ b/tensorflow/contrib/tensorrt/ops/trt_calib_op.cc @@ -0,0 +1,34 @@ +/* Copyright 2018 The TensorFlow Authors. All Rights Reserved. + +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + +Unless required by applicable law or agreed to in writing, software +distributed under the License is distributed on an "AS IS" BASIS, +WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +See the License for the specific language governing permissions and +limitations under the License. +==============================================================================*/ + +#include "tensorflow/core/framework/op.h" +#include "tensorflow/core/framework/shape_inference.h" +namespace tensorflow { + + +REGISTER_OP("TRTCalibOp") + .Attr("segment_nodes: list(string)") // names of the ops in segment + .Attr("segment_output_names: list(string)") // names of the output ops in segment + .Attr("InT: list({int8, float16, float32})") + .Input("in_tensor: InT") + .Output("out_tensor: InT") + .SetShapeFn([](tensorflow::shape_inference::InferenceContext* c) { + for (int i = 0; i < c->num_inputs(); i++){ + c->set_output(i, c->input(i)); + } + return Status::OK(); + }); + +} // namespace tensorflow diff --git a/tensorflow/contrib/tensorrt/python/trt_convert.py b/tensorflow/contrib/tensorrt/python/trt_convert.py index 354f0c8b42..5aba371a03 100644 --- a/tensorflow/contrib/tensorrt/python/trt_convert.py +++ b/tensorflow/contrib/tensorrt/python/trt_convert.py @@ -30,7 +30,7 @@ from tensorflow.python.framework import meta_graph from tensorflow.python.framework import ops -def CreateInferenceGraph(input_graph_def, outputs,max_batch_size=1,max_workspace_size=2<<20): +def CreateInferenceGraph(input_graph_def, outputs,max_batch_size=1,max_workspace_size=2<<20, int8=False): """Python wrapper for the TRT transormation. @@ -76,7 +76,7 @@ def CreateInferenceGraph(input_graph_def, outputs,max_batch_size=1,max_workspace # transformed graphs protobuf string. out = trt_convert( optimized_graph_def_str ,outputs, - max_batch_size,max_workspace_size) + max_batch_size,max_workspace_size,int8) status = out[0] output_graph_def_string = out[1] del optimized_graph_def_str #save some memory diff --git a/tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.cc b/tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.cc index 3c94b52ea6..fe414c45ce 100644 --- a/tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.cc +++ b/tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.cc @@ -9,7 +9,7 @@ namespace tensorflow { namespace trt { - +// set the batch size before constructing the thread to execute engine int TRTInt8Calibrator::getBatchSize() const { return batch_size_; } bool TRTInt8Calibrator::setBatch( diff --git a/tensorflow/contrib/tensorrt/resources/TRTResources.h b/tensorflow/contrib/tensorrt/resources/TRTResources.h index 2b65017943..2fe78b882d 100644 --- a/tensorflow/contrib/tensorrt/resources/TRTResources.h +++ b/tensorflow/contrib/tensorrt/resources/TRTResources.h @@ -6,27 +6,40 @@ #define TENSORFLOW_CONTRIB_TENSORRT_RESOURCES_TRTRESOURCES_H_ -#include #include -#include "tensorflow/contrib/tensorrt/resourcemgr/TRTInt8Calibrator.h" +#include +#include "tensorflow/contrib/tensorrt/log/trt_logger.h" +#include "tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.h" #include "tensorflow/core/framework/resource_mgr.h" namespace tensorflow { namespace trt { struct TRTCalibrationResource : public tensorflow::ResourceBase { - TRTCalibrationResource():calibrator(nullptr), builder(nullptr), thr(nullptr){}; + TRTCalibrationResource() + : calibrator(nullptr), + builder(nullptr), + network(nullptr), + engine(nullptr), + logger(nullptr), + thr(nullptr) {} + string DebugString() override { + return ""; + } TRTInt8Calibrator* calibrator; nvinfer1::IBuilder* builder; - std::thread *thr; + nvinfer1::INetworkDefinition* network; + nvinfer1::ICudaEngine* engine; + tensorflow::tensorrt::Logger* logger; + std::thread* thr; }; -struct TRTEngineResource:public tensorflow::ResourceBase{ - TRTEngineResource():runtime(nullptr), ctx(nullptr){}; - nvinfer1::IRuntime *runtime; - nvinfer1::IExecutionContext *ctx; +struct TRTEngineResource : public tensorflow::ResourceBase { + TRTEngineResource() : runtime(nullptr), ctx(nullptr){}; + nvinfer1::IRuntime* runtime; + nvinfer1::IExecutionContext* ctx; }; -} -} -#endif // TENSORFLOW_CONTRIB_TENSORRT_RESOURCEMGR_TRTRESOURCES_H_ +} // namespace trt +} // namespace tensorflow +#endif // TENSORFLOW_CONTRIB_TENSORRT_RESOURCEMGR_TRTRESOURCES_H_ diff --git a/tensorflow/contrib/tensorrt/trt_conversion.i b/tensorflow/contrib/tensorrt/trt_conversion.i index 5f8e73a59f..3e8baf91ae 100644 --- a/tensorflow/contrib/tensorrt/trt_conversion.i +++ b/tensorflow/contrib/tensorrt/trt_conversion.i @@ -28,7 +28,8 @@ std::pair trt_convert(string graph_def_string,//const tensorflow::GraphDef& std::vector output_names, size_t max_batch_size, - size_t max_workspace_size + size_t max_workspace_size_bytes, + bool int8 // unfortunately we can't use TF_Status here since it // is in c/c_api and brings in a lot of other libraries // which in turn declare ops. These ops are included @@ -57,8 +58,8 @@ tensorrt::convert::ConvertGraphDefToTensorRT(graph_def, output_names, max_batch_size, - max_workspace_size, - &outGraph); + max_workspace_size_bytes, + &outGraph,int8); if (!conversion_status.ok()) { auto retCode=(int)conversion_status.code(); char buff[2000]; @@ -79,6 +80,6 @@ std::pair trt_convert(string graph_def_string, std::vector output_names, size_t max_batch_size, - size_t max_workspace_size); + size_t max_workspace_size,bool int8); %unignoreall -- GitLab From adaabc11680fa2823d029cf67214b23fa6652a4b Mon Sep 17 00:00:00 2001 From: Jie Date: Mon, 5 Feb 2018 18:56:48 -0800 Subject: [PATCH 0021/1867] [DEBUG] multiple GPU crash with [cuda_illigal_memory_address] added cudaSetDevice before ICudaEngine::createExecutionContext() To make sure TRT engine gets allocated on the same GPU (to access IO memory) --- .../contrib/tensorrt/kernels/trt_engine_op.cc | 26 ++++++++++++++++--- .../contrib/tensorrt/segment/segment.cc | 10 ------- 2 files changed, 22 insertions(+), 14 deletions(-) diff --git a/tensorflow/contrib/tensorrt/kernels/trt_engine_op.cc b/tensorflow/contrib/tensorrt/kernels/trt_engine_op.cc index 445900f08c..81fd4c9747 100644 --- a/tensorflow/contrib/tensorrt/kernels/trt_engine_op.cc +++ b/tensorflow/contrib/tensorrt/kernels/trt_engine_op.cc @@ -44,11 +44,22 @@ TRTEngineOp::TRTEngineOp(OpKernelConstruction* context) : OpKernel(context) { // TODO(samikama) runtime should be taken from a resourcemanager as well. // Only engine should be in the op and context and runtime should be taken // from resourcemanager + // TODO(jie): cudaSetDevice make sure trt engine is allocated on the same + // gpu where the input/output is also located. + int gpu_id = context->device()->tensorflow_gpu_device_info()->gpu_id; + cudaSetDevice(gpu_id); + int device; + cudaGetDevice(&device); + if (gpu_id != device) + LOG(FATAL) << "set device failed!"; + IRuntime* infer = createInferRuntime(gLogger); trt_engine_ptr_.reset(infer->deserializeCudaEngine( serialized_engine.c_str(), serialized_engine.size(), nullptr)); trt_context_ptr_.reset(trt_engine_ptr_->createExecutionContext()); + + // trt_context_ptr_.reset(nullptr); // runtime is safe to delete after engine creation infer->destroy(); std::stringstream oss; @@ -103,12 +114,16 @@ void TRTEngineOp::Compute(OpKernelContext* context) { const TensorShape& input_shape = input_tensor.shape(); if (i == 0) { nbBatch = input_shape.dim_size(0); + if (nbBatch > trt_engine_ptr_->getMaxBatchSize()) + LOG(FATAL) << "input tensor batch larger than max_batch_size: " + << trt_engine_ptr_->getMaxBatchSize(); } else if (nbBatch != input_shape.dim_size(0)) { valid = false; break; } // int64 input_shape.dim_size(int d) // int input_shape.dims() + LOG(INFO) << "INPUT BINDING index: " << bindingIndex << " with name: " << input_nodes_[i]; switch (trt_engine_ptr_->getBindingDataType(bindingIndex)) { case nvinfer1::DataType::kFLOAT: LOG(INFO) << "float"; @@ -125,7 +140,7 @@ void TRTEngineOp::Compute(OpKernelContext* context) { } } - if (!valid) LOG(WARNING) << "input data inconsistent batch size"; + if (!valid) LOG(FATAL) << "input data inconsistent batch size"; for (int i = 0; i < static_cast(output_nodes_.size()); i++) { // This is bad that we have to reallocate output buffer every run. @@ -135,7 +150,7 @@ void TRTEngineOp::Compute(OpKernelContext* context) { TensorShape output_shape; if (bindingIndex != -1) { - LOG(INFO) << "got binding " << bindingIndex; + LOG(INFO) << "got binding " << bindingIndex << " with name: " << output_nodes_[i]; auto dims = trt_engine_ptr_->getBindingDimensions(bindingIndex); std::vector trt_shape(dims.nbDims + 1); trt_shape[0] = nbBatch; @@ -167,6 +182,7 @@ void TRTEngineOp::Compute(OpKernelContext* context) { break; } } + LOG(INFO) << "getting stream"; // copied from cuda_kernel_helper since it seems only valid in *.cu.cc files const cudaStream_t* stream = CHECK_NOTNULL( reinterpret_cast(context->op_device_context() @@ -174,9 +190,11 @@ void TRTEngineOp::Compute(OpKernelContext* context) { ->implementation() ->CudaStreamMemberHack())); - trt_context_ptr_->enqueue(nbBatch, &buffers[0], *stream, nullptr); + // TODO(jie): trt enqueue does not return error + LOG(INFO) << "enqueue returns: " << trt_context_ptr_->enqueue(nbBatch, &buffers[0], *stream, nullptr); + LOG(INFO) << "all good"; // sync should be done by TF. - //cudaStreamSynchronize(*stream); + // cudaStreamSynchronize(*stream); } REGISTER_KERNEL_BUILDER(Name("TRTEngineOp").Device(DEVICE_GPU), TRTEngineOp); diff --git a/tensorflow/contrib/tensorrt/segment/segment.cc b/tensorflow/contrib/tensorrt/segment/segment.cc index 41da528247..d749d0d0e8 100644 --- a/tensorflow/contrib/tensorrt/segment/segment.cc +++ b/tensorflow/contrib/tensorrt/segment/segment.cc @@ -220,16 +220,6 @@ tensorflow::Status SegmentGraph( } } - // Cleanup the graph to remove disconnected nodes before outputting - if (VLOG_IS_ON(2)) { - for (tensorflow::Node* node : graph.nodes()) { - if ((node->in_edges().size() == 0) && (node->out_edges().size() == 0)) { - graph.RemoveNode(node); - } - } - // tensorflow::DumpGraph("Post-Segment", &graph); - } - // Convert the segments into the expected return format for (const auto& itr : sg_map) { const auto& segment_node_names = itr.second; -- GitLab From 0b8492b612eef6057440c4d1fe5dca41cacf5d6d Mon Sep 17 00:00:00 2001 From: Sami Kama Date: Mon, 12 Feb 2018 18:40:07 -0800 Subject: [PATCH 0022/1867] Debugging calibration --- .../contrib/tensorrt/convert/convert_graph.cc | 28 +++- .../contrib/tensorrt/convert/convert_graph.h | 5 +- .../contrib/tensorrt/convert/convert_nodes.cc | 147 +++++++++++++++++- .../contrib/tensorrt/convert/convert_nodes.h | 8 +- .../contrib/tensorrt/kernels/trt_calib_op.cc | 22 +-- .../contrib/tensorrt/python/__init__.py | 1 + .../contrib/tensorrt/python/trt_convert.py | 20 ++- .../tensorrt/resources/TRTInt8Calibrator.cc | 65 +++++++- .../tensorrt/resources/TRTInt8Calibrator.h | 9 +- tensorflow/contrib/tensorrt/trt_conversion.i | 135 ++++++++++------ 10 files changed, 363 insertions(+), 77 deletions(-) diff --git a/tensorflow/contrib/tensorrt/convert/convert_graph.cc b/tensorflow/contrib/tensorrt/convert/convert_graph.cc index 494920fb7c..8aa4e42fa6 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_graph.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_graph.cc @@ -216,11 +216,11 @@ tensorflow::Status GetCalibNode(ConvertGraphParams* params) { TF_RETURN_IF_ERROR(status); for (auto in_edge: params->subgraph_incoming_edges) { // loop over incoming edges and attach them to calib node - tensorflow::Node* src_node = in_edge->src(); + // tensorflow::Node* src_node = in_edge->src(); auto src_output=in_edge->src_output(); auto dst_node=in_edge->dst(); auto dst_input=in_edge->dst_input(); - VLOG(0)<<" update edge "<name()<<":"< "<name()<<":"<name()<<":"< "<name()<<":"<graph.UpdateEdge(trt_node, src_output, dst_node, dst_input); } @@ -330,6 +330,30 @@ tensorflow::Status BuildNodeMap( } } // namespace +tensorflow::Status ConvertCalibGraphToInferGraph( + const tensorflow::GraphDef& graph_def, + tensorflow::GraphDef* infer_graph){ + VLOG(0)<<"Starting Calib Conversion"; + tensorflow::Graph graph(tensorflow::OpRegistry::Global()); + TF_RETURN_IF_ERROR(tensorflow::ConvertGraphDefToGraph( + tensorflow::GraphConstructorOptions(), graph_def, &graph)); + // get calib nodes + std::vector calibNodes; + for(auto node : graph.op_nodes()){ + if(node->type_string()=="TRTCalibOp"){ + VLOG(1)<<"Found Calib Node"; + calibNodes.push_back(node); + } + } + VLOG(0)<<"Num Calib nodes in graph= "<& output_names, size_t max_batch_size, diff --git a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc index 4e70fb00f9..588cecf8dd 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc @@ -40,6 +40,7 @@ limitations under the License. #include "tensorflow/core/graph/graph_constructor.h" #include "tensorflow/core/lib/core/errors.h" #include "tensorflow/core/lib/core/status.h" +#include "tensorflow/core/lib/strings/str_util.h" #include "tensorflow/core/platform/logging.h" #define _TF_LOG_DEBUG ::tensorflow::internal::LogMessage(__FILE__, __LINE__, -1) @@ -299,6 +300,11 @@ std::vector TFAttrs::get>(std::string key) const { return std::vector(attr.begin(), attr.end()); } template <> +std::vector TFAttrs::get>(std::string key) const { + auto attr = this->at(key)->list().s(); + return std::vector(attr.begin(), attr.end()); +} +template <> nvinfer1::Dims TFAttrs::get(std::string key) const { auto values = this->get>(key); nvinfer1::Dims dims; @@ -1938,6 +1944,125 @@ void Converter::register_op_converters() { tensorflow::Status GetTensorRTGraph(tensorrt::convert::SubGraphParams& s) { return tensorflow::errors::Unimplemented("Not implemented yet"); } +tensorflow::Status ConvertCalibrationNodeToEngineNode(tensorflow::Graph &graph, + tensorflow::Node *c_node) { + const auto ndef=c_node->def(); + + TFAttrs attrs(ndef); + std::vector segment_nodes(attrs.get>("segment_nodes")); + std::vector output_nodes(attrs.get>("segment_output_names")); + std::vector input_names(attrs.get>("input_names")); + std::string res_name = attrs.get("resource_name"); + VLOG(1) << "Node name " << c_node->name() << " res_name " << res_name; + std::string engine_name="my_trt_op"; + { + const auto node_id=tensorflow::str_util::Split(res_name,"_"); + engine_name+=node_id.back(); + } + std::map nodeMaps; + + for(auto n: graph.op_nodes()){ + nodeMaps.insert({n->name(),n}); + } + VLOG(1)<<"Output Nodes:"; + std::vector out_types; + std::vector out_edges; + for(auto &i : output_nodes ){ + auto node_port=tensorflow::str_util::Split(i,":"); + VLOG(1) << " " << i << " in graph " << nodeMaps.count(i); + auto out_node_name = node_port.at(0); + if(node_port.size()>1){ + VLOG(1) << "Multi port output" << node_port.at(0) << + " " << node_port.at(1) << " size=" << node_port.size(); + } + auto nodeIt=nodeMaps.find(out_node_name); + if(nodeIt!=nodeMaps.end()){ + tensorflow::Node* outNode=nodeIt->second; + int port=0; + if(node_port.size()==2){ + port=std::strtoul(node_port.at(1).c_str(),nullptr,10); + out_types.push_back(outNode->output_type(port)); + }else{ + out_types.push_back(outNode->output_type(0)); + } + for(auto outEdge : outNode->out_edges()){ + if(outEdge->src_output()==port){ + out_edges.push_back(outEdge); + break; + } + } + }else{ + LOG(WARNING)<<" couldn't find output node "<getManager("TRTCalibOps"); + tensorflow::trt::TRTCalibrationResource* calibRes = nullptr; + auto status = resmgr->Lookup(res_name, res_name, &calibRes); + if(!status.ok() || !calibRes->calibrator){ + return tensorflow::errors::FailedPrecondition("You must run calibration"\ + " and inference conversion in the same proces"); + } + + calibRes->calibrator->setDone(); + VLOG(1)<<"Waiting for calibration thread to join"; + calibRes->thr->join(); + delete calibRes->thr; + if(!calibRes->engine){ + LOG(FATAL)<<"Calibration failed!, engine is nullptr"; + } + auto engine_plan_string=calibRes->engine->serialize(); + calibRes->engine->destroy(); + calibRes->network->destroy(); + calibRes->builder->destroy(); + calibRes->thr= nullptr; + calibRes->engine= nullptr; + calibRes->builder= nullptr; + tensorflow::NodeDefBuilder op_builder(engine_name, "TRTEngineOp"); + std::vector income_edges; + for(const auto in_edge : c_node->in_edges()){ + auto src=in_edge->src(); + int dest_port=in_edge->dst_input(); + income_edges.emplace_back(src->name(),in_edge->src_output(),c_node->input_type(dest_port)); + } + tensorflow::gtl::ArraySlice input_list( + income_edges); + op_builder.Input(input_list); + tensorflow::NodeDef engine_node; + status = op_builder.Attr("serialized_engine", engine_plan_string) + .Attr("input_nodes", input_names) + .Attr("output_nodes", output_nodes) + .Attr("OutT", out_types) + .Finalize(&engine_node); + if(!status.ok()){ + LOG(ERROR)<<"Engine Node creation failed"; + return status; + } + auto trt_engine_node=graph.AddNode(engine_node,&status); + TF_CHECK_OK(status); + for(size_t i=0;idst()->name() << " port " + << out_edges.at(i)->dst_input(); + TF_RETURN_IF_ERROR(graph.UpdateEdge(trt_engine_node, i, + out_edges.at(i)->dst(), + out_edges.at(i)->dst_input())); + } + VLOG(1) << "Segment nodes:"; + for (auto &i : segment_nodes){ + VLOG(1) << " " << i << " in graph " << nodeMaps.count(i); + auto it=nodeMaps.find(i); + if(it!=nodeMaps.end()){ + graph.RemoveNode(it->second); + } + } + return tensorflow::Status::OK(); +} tensorflow::Status InjectCalibrationNode(tensorrt::convert::SubGraphParams& s) { // Visit nodes in reverse topological order and construct the TRT network. @@ -1958,13 +2083,15 @@ tensorflow::Status InjectCalibrationNode(tensorrt::convert::SubGraphParams& s) { LOG(DEBUG) << "BUILDING 1"; static int static_id = 0; std::string calib_op_name = - std::string("my_trt_calib_op_") + std::to_string(static_id++); - + std::string("my_trt_calib_op_") + std::to_string(static_id); + std::string engine_name = + std::string("my_trt_op") + std::to_string(static_id); + static_id++; LOG(DEBUG) << "BUILDING 2"; auto trt_rmgr = tensorflow::trt::TRTResourceManager::instance(); auto op_rmgr = trt_rmgr->getManager("TRTCalibOps"); auto op_res = new tensorflow::trt::TRTCalibrationResource(); - VLOG(0)<<"SAMI Creating calibresource "<Create(calib_op_name, calib_op_name, op_res)); op_res->logger = new tensorflow::tensorrt::Logger(); op_res->builder = nvinfer1::createInferBuilder(*(op_res->logger)); @@ -2065,15 +2192,23 @@ tensorflow::Status InjectCalibrationNode(tensorrt::convert::SubGraphParams& s) { // Gather output metadata std::vector output_names; std::vector output_dtypes; + int trt_engine_op_output_idx = 0; for (std::pair const& output : s.output_inds) { int node_id = output.first; int output_idx = output.second; tensorflow::Node* node = s.graph.FindNodeId(node_id); std::string op_name = node->name(); std::string tensor_name = op_name; + + s.output_edge_map->insert( + {trt_engine_op_output_idx == 0 + ? engine_name + : engine_name + ":" + std::to_string(trt_engine_op_output_idx), + {output_idx, tensor_name}}); + trt_engine_op_output_idx++; if (output_idx != 0) tensor_name = tensor_name + ":" + std::to_string(output_idx); - LOG(DEBUG) << "output tensor name: " << tensor_name; + VLOG(1) << "output tensor name: " << tensor_name; output_names.push_back(tensor_name); auto tensor_or_weights = converter.get_tensor(tensor_name); if (!tensor_or_weights.is_tensor()) { @@ -2083,7 +2218,7 @@ tensorflow::Status InjectCalibrationNode(tensorrt::convert::SubGraphParams& s) { nvinfer1::ITensor* tensor = tensor_or_weights.tensor(); if (!tensor) { return tensorflow::errors::NotFound("Output tensor not found: " + - tensor_name); + tensor_name); } converter.network()->markOutput(*tensor); tensorflow::DataType tf_dtype = node->output_type(output_idx); @@ -2109,7 +2244,7 @@ tensorflow::Status InjectCalibrationNode(tensorrt::convert::SubGraphParams& s) { // ConvertSubGraphToTensorRT(convert_graph.cc) auto incoming_edge = tensorflow::NodeDefBuilder::NodeOut( input_names.at(i), output_idx, input_dtypes.at(i)); - VLOG(0) << calib_op_name << " input " << i << " = " << input_names.at(i) + VLOG(1) << calib_op_name << " input " << i << " = " << input_names.at(i) << ":" << output_idx <<" dType= "<< tensorflow::DataTypeString(input_dtypes.at(i)); income_edges.push_back(incoming_edge); diff --git a/tensorflow/contrib/tensorrt/convert/convert_nodes.h b/tensorflow/contrib/tensorrt/convert/convert_nodes.h index 2f754968dc..71f61e2dc4 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_nodes.h +++ b/tensorflow/contrib/tensorrt/convert/convert_nodes.h @@ -31,7 +31,7 @@ namespace tensorrt { namespace convert { struct SubGraphParams { - SubGraphParams(const tensorflow::Graph& graph_, + SubGraphParams(tensorflow::Graph& graph_, const std::set& subgraph_node_ids_, const std::vector>& input_inds_, const std::vector>& output_inds_, @@ -52,7 +52,7 @@ struct SubGraphParams { trt_node(trt_node_), int8(int8_) {} - const tensorflow::Graph& graph; + tensorflow::Graph& graph; const std::set& subgraph_node_ids; const std::vector>& input_inds; // {node_id, output_idx} const std::vector>& output_inds; // {node_id, output_idx} @@ -64,8 +64,10 @@ struct SubGraphParams { const bool int8; }; -tensorflow::Status ConvertSubGraphToTensorRTNodeDef(SubGraphParams& params); +tensorflow::Status ConvertSubGraphToTensorRTNodeDef(SubGraphParams ¶ms); tensorflow::Status InjectCalibrationNode(SubGraphParams& params); +tensorflow::Status ConvertCalibrationNodeToEngineNode(tensorflow::Graph& graph, + tensorflow::Node* c_node); } // namespace convert } // namespace tensorrt diff --git a/tensorflow/contrib/tensorrt/kernels/trt_calib_op.cc b/tensorflow/contrib/tensorrt/kernels/trt_calib_op.cc index 4996b3cd40..41906b6090 100644 --- a/tensorflow/contrib/tensorrt/kernels/trt_calib_op.cc +++ b/tensorflow/contrib/tensorrt/kernels/trt_calib_op.cc @@ -43,23 +43,22 @@ TRTCalibOp::TRTCalibOp(OpKernelConstruction* context) : OpKernel(context) { } void TRTCalibOp::Compute(tensorflow::OpKernelContext* ctx) { auto trt_rm = tensorflow::trt::TRTResourceManager::instance(); - VLOG(0) << "Op Name= " << name() << " nodedef name= " << repo_name; + VLOG(2) << "Op Name= " << name() << " nodedef name= " << repo_name; auto resmgr = trt_rm->getManager("TRTCalibOps"); tensorflow::trt::TRTCalibrationResource* calibRes = nullptr; auto status = resmgr->Lookup(repo_name, repo_name, &calibRes); - VLOG(0) << "SAMI status " << status.ToString(); if (status.ok()) { int batchSize = ctx->input(0).dim_size(0); - VLOG(0) << "SAMI Batchsize= " << batchSize; + VLOG(2) << "SAMI Batchsize= " << batchSize; int numInputs = ctx->num_inputs(); - VLOG(0) << "SAMI numInputs= " << numInputs; + VLOG(2) << "SAMI numInputs= " << numInputs; dev_tensors_.resize(numInputs); if (calibRes->calibrator == nullptr) { - VLOG(0) << " Constructing calibrator"; + VLOG(1) << " Constructing calibrator"; // first run for (int i = 0; i < numInputs; i++) { const tensorflow::Tensor& t = ctx->input(i); - VLOG(0) << "Tensor " << i << " " << t.shape().DebugString(); + VLOG(1) << "Tensor " << i << " " << t.shape().DebugString(); OP_REQUIRES_OK(ctx, ctx->allocate_persistent(t.dtype(), t.shape(), &dev_tensors_.at(i), nullptr)); @@ -73,11 +72,14 @@ void TRTCalibOp::Compute(tensorflow::OpKernelContext* ctx) { } calibRes->calibrator = new TRTInt8Calibrator(device_buffers_, batchSize); calibRes->thr = new std::thread([calibRes]() { + VLOG(0)<<"Starting calibration thread, Calibration Resource @ "<builder->setInt8Calibrator(calibRes->calibrator); + calibRes->builder->setInt8Mode(true); calibRes->engine = calibRes->builder->buildCudaEngine( *calibRes->network); // will loop until we terminate calibrator - VLOG(1) << "Calibration loop terminated"; + VLOG(0) << "SAMI Calibration loop terminated"; }); - VLOG(0) << "SAMI intialized calibrator resource"; + VLOG(0) << "SAMI initialized calibrator resource"; } std::unordered_map input_data; @@ -92,9 +94,9 @@ void TRTCalibOp::Compute(tensorflow::OpKernelContext* ctx) { input_data.emplace(input_names_.at(i), data_address); ctx->set_output(i, t); } - VLOG(0) << "Filled map"; + VLOG(1) << "Filled map for sending"; calibRes->calibrator->setBatch(input_data); - VLOG(0) << "Passed calibration data"; + VLOG(1) << "Passed calibration data"; } else { ctx->SetStatus(status); return; diff --git a/tensorflow/contrib/tensorrt/python/__init__.py b/tensorflow/contrib/tensorrt/python/__init__.py index 4aeea48515..9eb589664c 100644 --- a/tensorflow/contrib/tensorrt/python/__init__.py +++ b/tensorflow/contrib/tensorrt/python/__init__.py @@ -5,4 +5,5 @@ from __future__ import print_function # pylint: disable=unused-import,wildcard-import from tensorflow.contrib.tensorrt.python.ops import trt_engine_op from tensorflow.contrib.tensorrt.python.trt_convert import CreateInferenceGraph +from tensorflow.contrib.tensorrt.python.trt_convert import CalibGraphToInferGraph # pylint: enable=unused-import,wildcard-import diff --git a/tensorflow/contrib/tensorrt/python/trt_convert.py b/tensorflow/contrib/tensorrt/python/trt_convert.py index 5aba371a03..18ea6c83cc 100644 --- a/tensorflow/contrib/tensorrt/python/trt_convert.py +++ b/tensorflow/contrib/tensorrt/python/trt_convert.py @@ -21,7 +21,7 @@ from __future__ import print_function from tensorflow.core.framework import graph_pb2 from tensorflow.python.framework import errors from tensorflow.python.framework import errors_impl as _impl -from tensorflow.contrib.tensorrt.wrap_conversion import trt_convert +from tensorflow.contrib.tensorrt.wrap_conversion import trt_convert,calib_convert from tensorflow.python.util import compat import tensorflow as tf from tensorflow.python.grappler import tf_optimizer @@ -91,3 +91,21 @@ def CreateInferenceGraph(input_graph_def, outputs,max_batch_size=1,max_workspace output_graph_def.ParseFromString(output_graph_def_string) del output_graph_def_string #save some memory return output_graph_def + +def CalibGraphToInferGraph(calibration_graph_def): + graph_str=calibration_graph_def.SerializeToString() + out=calib_convert(graph_str) + status=out[0] + output_graph_def_string = out[1] + del graph_str #save some memory + if len(status) < 2: + raise _impl.UnknownError(None,None,status) + if status[:2] != "OK": + msg=status.split(";") + if len(msg) == 1: + raise RuntimeError("Status message is malformed {}".format(status)) + raise _impl._make_specific_exception(None,None,";".join(msg[1:]), int(msg[0])) + output_graph_def = graph_pb2.GraphDef() + output_graph_def.ParseFromString(output_graph_def_string) + del output_graph_def_string #save some memory + return output_graph_def diff --git a/tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.cc b/tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.cc index 10d9350d7a..e1ab243b07 100644 --- a/tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.cc +++ b/tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.cc @@ -5,6 +5,10 @@ #include "tensorflow/contrib/tensorrt/resources/TRTInt8Calibrator.h" #include +#include +#include +#include + #include "tensorflow/core/platform/logging.h" namespace tensorflow { @@ -12,26 +16,67 @@ namespace trt { // set the batch size before constructing the thread to execute engine int TRTInt8Calibrator::getBatchSize() const { return batch_size_; } +TRTInt8Calibrator::TRTInt8Calibrator(const std::unordered_map< + std::string, std::pair>& dev_buffers, + int batch_size) + : batch_size_(batch_size), + done_(false), + dev_buffers_(dev_buffers), + calib_running_(false){ + cudaPointerAttributes pa; + int devid=-1; + cudaGetDevice(&devid); + VLOG(0)<<"Constructing calibrator with batch size "<& data) { + VLOG(1)<<"SAMI SAMI Waiting to set new batch"; + if(done_)return false; while (calib_running_.load( std::memory_order_acquire)) { // wait while calibration is running tensorflow::mutex_lock l(cond_mtx_); cond_.wait_for(l, std::chrono::milliseconds(50)); + if(done_)return false; } + VLOG(1)<<"Set Batch Waiting finished"; for (const auto it : data) { + auto devptr = dev_buffers_.find(it.first); if (devptr == dev_buffers_.end()) { LOG(FATAL) << "FATAL input name '" << it.first << "' does not match with the buffer names"; } + cudaPointerAttributes pa; const auto& d = devptr->second; + VLOG(1)<<"cuda memcopy buff name= "<second.first; bindings[i] = it->second.first; + float f[2]; + f[0]=3.; + f[1]=0.14159; + auto status=cudaMemcpy(f,bindings[i],sizeof(float)*2,cudaMemcpyDeviceToHost); + int devid=-1; + cudaGetDevice(&devid); + VLOG(0)<<"SAMI ORDER GETTING, Data in perm storage [0]="<>& dev_buffers, - int batch_size) - : batch_size_(batch_size), - done_(false), - dev_buffers_(dev_buffers), - calib_running_(false){}; + int batch_size); int getBatchSize() const; bool getBatch(void* bindings[], const char* names[], int nbBindings) override; bool setBatch(const std::unordered_map &data); void setDone(){done_=true;} const void *readCalibrationCache(std::size_t &length) override; void writeCalibrationCache(const void *ptr, std::size_t length) override; + ~TRTInt8Calibrator(); private: int batch_size_; tensorflow::mutex cond_mtx_; tensorflow::condition_variable cond_; bool done_; - std::unordered_map> dev_buffers_; + const std::unordered_map> dev_buffers_; std::atomic_bool calib_running_; }; } // namespace trt diff --git a/tensorflow/contrib/tensorrt/trt_conversion.i b/tensorflow/contrib/tensorrt/trt_conversion.i index 3e8baf91ae..ee87d7fae1 100644 --- a/tensorflow/contrib/tensorrt/trt_conversion.i +++ b/tensorflow/contrib/tensorrt/trt_conversion.i @@ -23,58 +23,98 @@ %ignoreall %unignore tensorflow; %unignore trt_convert; +%unignore calib_convert; %{ - std::pair trt_convert(string graph_def_string,//const tensorflow::GraphDef& - std::vector output_names, - size_t max_batch_size, - size_t max_workspace_size_bytes, - bool int8 - // unfortunately we can't use TF_Status here since it - // is in c/c_api and brings in a lot of other libraries - // which in turn declare ops. These ops are included - // statically in our library and cause an abort when - // module is loaded due to double registration - // until Tensorflow properly exposes these headers - // we have to work around this by returning a string - // and converting it to exception on python side. - //,TF_Status* out_status) { - ) { - string out_status; +std::pair trt_convert(string graph_def_string,//const tensorflow::GraphDef& + std::vector output_names, + size_t max_batch_size, + size_t max_workspace_size_bytes, + bool int8 + // unfortunately we can't use TF_Status here since it + // is in c/c_api and brings in a lot of other libraries + // which in turn declare ops. These ops are included + // statically in our library and cause an abort when + // module is loaded due to double registration + // until Tensorflow properly exposes these headers + // we have to work around this by returning a string + // and converting it to exception on python side. + //,TF_Status* out_status) { +) { + string out_status; - tensorflow::GraphDef graph_def; - if (!graph_def.ParseFromString(graph_def_string)) { - out_status="InvalidArgument;Couldn't interpret input as a GraphDef"; - return std::pair{out_status,""}; - } + tensorflow::GraphDef graph_def; + if (!graph_def.ParseFromString(graph_def_string)) { + out_status="InvalidArgument;Couldn't interpret input as a GraphDef"; + return std::pair{out_status,""}; + } - if (!output_names.size()) { - out_status="InvalidArgument;Size of the output_names vector is 0"; - return std::pair{out_status,""}; - //return ""; - } - tensorflow::GraphDef outGraph; - tensorflow::Status conversion_status = + if (!output_names.size()) { + out_status="InvalidArgument;Size of the output_names vector is 0"; + return std::pair{out_status,""}; + //return ""; + } + tensorflow::GraphDef outGraph; + tensorflow::Status conversion_status = tensorrt::convert::ConvertGraphDefToTensorRT(graph_def, - output_names, - max_batch_size, - max_workspace_size_bytes, - &outGraph,int8); - if (!conversion_status.ok()) { - auto retCode=(int)conversion_status.code(); - char buff[2000]; - snprintf(buff,2000,"%d;%s",retCode,conversion_status.error_message().c_str()); - out_status=buff; - return std::pair{out_status,""}; - } - string result; - if (!outGraph.SerializeToString(&result)) { - out_status="InvalidArgument;Couldn't serialize output as a GraphDef"; - return std::pair{out_status,""}; - } - out_status="OK;All good!"; - return std::pair{out_status,result}; + output_names, + max_batch_size, + max_workspace_size_bytes, + &outGraph,int8); + if (!conversion_status.ok()) { + auto retCode=(int)conversion_status.code(); + char buff[2000]; + snprintf(buff,2000,"%d;%s",retCode,conversion_status.error_message().c_str()); + out_status=buff; + return std::pair{out_status,""}; + } + string result; + if (!outGraph.SerializeToString(&result)) { + out_status="InvalidArgument;Couldn't serialize output as a GraphDef"; + return std::pair{out_status,""}; + } + out_status="OK;All good!"; + return std::pair{out_status,result}; +} + +std::pair calib_convert(string graph_def_string // const tensorflow::GraphDef& + // unfortunately we can't use TF_Status here since it + // is in c/c_api and brings in a lot of other libraries + // which in turn declare ops. These ops are included + // statically in our library and cause an abort when + // module is loaded due to double registration + // until Tensorflow properly exposes these headers + // we have to work around this by returning a string + // and converting it to exception on python side. + //,TF_Status* out_status) { +) { + string out_status; + + tensorflow::GraphDef graph_def; + if (!graph_def.ParseFromString(graph_def_string)) { + out_status="InvalidArgument;Couldn't interpret input as a GraphDef"; + return std::pair{out_status,""}; } + + tensorflow::GraphDef outGraph; + tensorflow::Status conversion_status = + tensorrt::convert::ConvertCalibGraphToInferGraph(graph_def, + &outGraph); + if (!conversion_status.ok()) { + auto retCode=(int)conversion_status.code(); + char buff[2000]; + snprintf(buff,2000,"%d;%s",retCode,conversion_status.error_message().c_str()); + out_status=buff; + return std::pair{out_status,""}; + } + string result; + if (!outGraph.SerializeToString(&result)) { + out_status="InvalidArgument;Couldn't serialize output as a GraphDef"; + return std::pair{out_status,""}; + } + out_status="OK;All good!"; + return std::pair{out_status,result}; +} %} std::pair trt_convert(string graph_def_string, @@ -82,4 +122,7 @@ std::pair trt_convert(string graph_def_string, size_t max_batch_size, size_t max_workspace_size,bool int8); +std::pair calib_convert(string graph_def_string); + + %unignoreall -- GitLab From ca19b32e4d1574ad29e36dbc164c320aeca80d47 Mon Sep 17 00:00:00 2001 From: Guozhong Zhuang Date: Wed, 14 Feb 2018 00:13:00 -0800 Subject: [PATCH 0023/1867] cifar 10 divergance fix and batchnorm unit test fix --- .../core/kernels/mkl_fused_batch_norm_op.cc | 96 +++++++++++++------ tensorflow/core/kernels/mkl_relu_op.cc | 20 +++- 2 files changed, 81 insertions(+), 35 deletions(-) diff --git a/tensorflow/core/kernels/mkl_fused_batch_norm_op.cc b/tensorflow/core/kernels/mkl_fused_batch_norm_op.cc index 8313224d7f..b7dee3fb3e 100644 --- a/tensorflow/core/kernels/mkl_fused_batch_norm_op.cc +++ b/tensorflow/core/kernels/mkl_fused_batch_norm_op.cc @@ -1110,19 +1110,12 @@ class MklFusedBatchNormGradOp : public OpKernel { return; } - if (dnn_shape_src.IsMklTensor()) - depth_ = dnn_shape_src.DimSize(MklDnnDims::Dim_C); - else - ExtractParams(context); - - memory::format format_m; if (dnn_shape_src.IsMklTensor()) { - if (dnn_shape_src.IsTensorInNCHWFormat()) - format_m = memory::format::nchw; - else - format_m = memory::format::nhwc; + depth_ = dnn_shape_src.DimSize(MklDnnDims::Dim_C); + } else if (dnn_shape_diff_dst.IsMklTensor()) { + depth_ = dnn_shape_diff_dst.DimSize(MklDnnDims::Dim_C); } else { - format_m = TFDataFormatToMklDnnDataFormat(tensor_format_); + ExtractParams(context); } MklDnnData src(&cpu_engine); @@ -1146,20 +1139,20 @@ class MklFusedBatchNormGradOp : public OpKernel { diff_dst_dims = TFShapeToMklDnnDimsInNCHW(diff_dst_tensor.shape(), tensor_format_); - // set src and diff_dst primitives + // set src and diff_dst primitives according to input layout memory::desc src_md({}, memory::data_undef, memory::format_undef); memory::desc diff_dst_md({}, memory::data_undef, memory::format_undef); - if (dnn_shape_src.IsMklTensor() || dnn_shape_diff_dst.IsMklTensor()) { - if (dnn_shape_src.IsMklTensor()) { - src_md = dnn_shape_src.GetMklLayout(); - diff_dst_md = src_md; - } else { - diff_dst_md = dnn_shape_diff_dst.GetMklLayout(); - src_md = diff_dst_md; - } + if (dnn_shape_src.IsMklTensor()) { + src_md = dnn_shape_src.GetMklLayout(); } else { - src_md = memory::desc(src_dims, MklDnnType(), format_m); - diff_dst_md = src_md; + src_md = memory::desc(src_dims, MklDnnType(), + TFDataFormatToMklDnnDataFormat(tensor_format_)); + } + if (dnn_shape_diff_dst.IsMklTensor()) { + diff_dst_md = dnn_shape_diff_dst.GetMklLayout(); + } else { + diff_dst_md = memory::desc(diff_dst_dims, MklDnnType(), + TFDataFormatToMklDnnDataFormat(tensor_format_)); } src.SetUsrMem(src_md, &src_tensor); diff_dst.SetUsrMem(diff_dst_md, &diff_dst_tensor); @@ -1211,28 +1204,64 @@ class MklFusedBatchNormGradOp : public OpKernel { // allocate diff_src tensor MklDnnShape dnn_shape_diff_src; TensorShape tf_shape_diff_src; - if (dnn_shape_src.IsMklTensor()) { + + // MKL-DNN's BN primitive not provide API to fetch internal format + // set common_md as OpMem + // src and diff_dst will reorder to common_md + // diff_src will set as common_md + memory::desc common_md({}, memory::data_undef, memory::format_undef); + if (dnn_shape_src.IsMklTensor() || dnn_shape_diff_dst.IsMklTensor()) { + if (dnn_shape_src.IsMklTensor()) { + common_md = dnn_shape_src.GetMklLayout(); + } else { + common_md = dnn_shape_diff_dst.GetMklLayout(); + } + } else { + common_md = memory::desc(src_dims, MklDnnType(), + TFDataFormatToMklDnnDataFormat(tensor_format_)); + } + // if any of src and diff_dst as mkl layout, + // then we set diff_src as mkl layout + if (dnn_shape_src.IsMklTensor() || + dnn_shape_diff_dst.IsMklTensor()) { dnn_shape_diff_src.SetMklTensor(true); - auto diff_src_pd = bnrm_fwd_pd.dst_primitive_desc(); + // set diff_src's mkl layout as common_md + auto diff_src_pd = memory::primitive_desc(common_md, cpu_engine); dnn_shape_diff_src.SetMklLayout(&diff_src_pd); dnn_shape_diff_src.SetElemType(MklDnnType()); - dnn_shape_diff_src.SetTfLayout(dnn_shape_src.GetDimension(), src_dims, - format_m); - dnn_shape_diff_src.SetTfDimOrder(dnn_shape_src.GetDimension(), - tensor_format_); + if (dnn_shape_src.IsMklTensor()) { + dnn_shape_diff_src.SetTfLayout( + dnn_shape_src.GetDimension(), + src_dims, + dnn_shape_src.GetTfDataFormat()); + dnn_shape_diff_src.SetTfDimOrder( + dnn_shape_src.GetDimension(), + tensor_format_); + } else { + dnn_shape_diff_src.SetTfLayout( + dnn_shape_diff_dst.GetDimension(), + src_dims, + dnn_shape_diff_dst.GetTfDataFormat()); + dnn_shape_diff_src.SetTfDimOrder( + dnn_shape_diff_dst.GetDimension(), + tensor_format_); + } tf_shape_diff_src.AddDim(diff_src_pd.get_size() / sizeof(T)); } else { dnn_shape_diff_src.SetMklTensor(false); + // both src and diff_dst are tf layout, + // so get tf shape from anyont should be ok tf_shape_diff_src = src_tensor.shape(); } AllocateOutputSetMklShape(context, kDiffSrcIndex, &diff_src_tensor, tf_shape_diff_src, dnn_shape_diff_src); - diff_src.SetUsrMem(src_md, diff_src_tensor); + // set diff_src + diff_src.SetUsrMem(common_md, diff_src_tensor); prop_kind pk = prop_kind::backward; auto bnrm_bwd_desc = batch_normalization_backward::desc( - pk, diff_src.GetUsrMemDesc(), src.GetUsrMemDesc(), epsilon_, + pk, common_md, common_md, epsilon_, /* for inference, specify use_global_stats 1. on fwd prop, use mean and variance provided as inputs @@ -1245,11 +1274,16 @@ class MklFusedBatchNormGradOp : public OpKernel { auto bnrm_bwd_pd = batch_normalization_backward::primitive_desc( bnrm_bwd_desc, cpu_engine, bnrm_fwd_pd); + std::vector net; + src.CheckReorderToOpMem(memory::primitive_desc(common_md, + cpu_engine), &net); + diff_dst.CheckReorderToOpMem(memory::primitive_desc(common_md, + cpu_engine), &net); + auto bnrm_bwd_op = batch_normalization_backward( bnrm_bwd_pd, src.GetOpMem(), mean.GetOpMem(), variance.GetOpMem(), diff_dst.GetOpMem(), weights_m, diff_src.GetOpMem(), diff_weights_m); - std::vector net; net.push_back(bnrm_bwd_op); stream(stream::kind::eager).submit(net).wait(); diff --git a/tensorflow/core/kernels/mkl_relu_op.cc b/tensorflow/core/kernels/mkl_relu_op.cc index 51db3991e2..924b9da7e0 100644 --- a/tensorflow/core/kernels/mkl_relu_op.cc +++ b/tensorflow/core/kernels/mkl_relu_op.cc @@ -368,8 +368,11 @@ void MklReluGradOp::Compute(OpKernelContext* context) { mkl_context.MklCleanup(); } + + #else // INTEL_MKL_ML + template class MklReluOpBase : public OpKernel { public: @@ -579,17 +582,26 @@ class MklReluGradOpBase : public OpKernel { // allocate diff_src tensor MklDnnShape dnn_shape_diff_src; TensorShape tf_shape_diff_src; - if (dnn_shape_src.IsMklTensor()) { + if (dnn_shape_src.IsMklTensor() || + dnn_shape_diff_dst.IsMklTensor()) { dnn_shape_diff_src.SetMklTensor(true); auto diff_src_pd = relu_bwd_pd.diff_src_primitive_desc(); dnn_shape_diff_src.SetMklLayout(&diff_src_pd); dnn_shape_diff_src.SetElemType(MklDnnType()); - dnn_shape_diff_src.SetTfLayout(dnn_shape_src.GetDimension(), - dnn_shape_src.GetSizesAsMklDnnDims(), - dnn_shape_src.GetTfDataFormat()); + if (dnn_shape_src.IsMklTensor()) { + dnn_shape_diff_src.SetTfLayout(dnn_shape_src.GetDimension(), + dnn_shape_src.GetSizesAsMklDnnDims(), + dnn_shape_src.GetTfDataFormat()); + } else { + dnn_shape_diff_src.SetTfLayout(dnn_shape_diff_dst.GetDimension(), + dnn_shape_diff_dst.GetSizesAsMklDnnDims(), + dnn_shape_diff_dst.GetTfDataFormat()); + } tf_shape_diff_src.AddDim(diff_src_pd.get_size() / sizeof(T)); } else { dnn_shape_diff_src.SetMklTensor(false); + // both src and diff_dst are tf layout, + // so get tf shape from anyone should be ok tf_shape_diff_src = src_tensor.shape(); } AllocateOutputSetMklShape(context, diff_src_index, &diff_src_tensor, -- GitLab From c6cd20dbcaaa601977d1b63ab17e04d137de5133 Mon Sep 17 00:00:00 2001 From: Ben Barsdell Date: Thu, 15 Feb 2018 19:01:57 -0800 Subject: [PATCH 0024/1867] Add node converter for FusedBatchNorm op --- .../contrib/tensorrt/convert/convert_graph.cc | 9 ++- .../contrib/tensorrt/convert/convert_nodes.cc | 67 +++++++++++++++++++ 2 files changed, 73 insertions(+), 3 deletions(-) diff --git a/tensorflow/contrib/tensorrt/convert/convert_graph.cc b/tensorflow/contrib/tensorrt/convert/convert_graph.cc index 31ba30b2d9..8c0aada355 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_graph.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_graph.cc @@ -68,9 +68,12 @@ bool IsTensorRTCandidate(const tensorflow::NodeDef& node_def) { "Mean", "AvgPool", "ConcatV2", - "DepthwiseConv2dNative" //, "MatMul", - //"Reshape" - // TODO(ben,jie): ... + "DepthwiseConv2dNative", + "FusedBatchNorm", + "FusedBatchNormV2", + //, "MatMul", + //"Reshape" + // TODO(ben,jie): ... }; // LINT.ThenChange(//tensorflow/contrib/tensorrt/convert/convert_nodes.h) return candidate_ops.count(node_def.op()); diff --git a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc index ea0eb480f2..e3b16126f1 100644 --- a/tensorflow/contrib/tensorrt/convert/convert_nodes.cc +++ b/tensorflow/contrib/tensorrt/convert/convert_nodes.cc @@ -276,6 +276,17 @@ template <> tensorflow::DataType TFAttrs::get(string key) const { return this->at(key)->type(); } + +template <> +float TFAttrs::get(string key) const { + return this->at(key)->f(); +} + +template <> +bool TFAttrs::get(string key) const { + return this->at(key)->b(); +} + // TODO(jie): reorder4 & reorder2 should be merged? template void Reorder4(nvinfer1::DimsNCHW shape, const T* idata, @@ -1703,6 +1714,60 @@ tensorflow::Status ConvertConcat(Converter& ctx, return tensorflow::Status::OK(); } +tensorflow::Status ConvertFusedBatchNorm(Converter& ctx, + tensorflow::NodeDef const& node_def, + std::vector const& inputs, + std::vector* outputs) { + TFAttrs attrs(node_def); + float epsilon = attrs.get("epsilon"); + auto data_format = attrs.get("data_format"); + if (data_format != "NCHW" ) { + return tensorflow::errors::Unimplemented( + "only data_format=NCHW is supported, at " + node_def.name()); + } + bool is_training = attrs.get("is_training"); + if (is_training) { + return tensorflow::errors::Unimplemented( + "only is_training=false is supported, at " + node_def.name()); + } + nvinfer1::ITensor const* tensor = inputs.at(0).tensor(); + TRT_ShapedWeights scale_weights = inputs.at(1).weights(); + TRT_ShapedWeights offset_weights = inputs.at(2).weights(); + TRT_ShapedWeights mean_weights = inputs.at(3).weights(); + TRT_ShapedWeights variance_weights = inputs.at(4).weights(); + TRT_ShapedWeights dummy_power_weights(scale_weights.type_); + TRT_ShapedWeights combined_scale_weights = + ctx.get_temp_weights_like(scale_weights); + TRT_ShapedWeights combined_offset_weights = + ctx.get_temp_weights_like(offset_weights); + size_t nweight = scale_weights.count(); + if (scale_weights.type_ != tensorflow::DataType::DT_FLOAT || + offset_weights.type_ != tensorflow::DataType::DT_FLOAT || + mean_weights.type_ != tensorflow::DataType::DT_FLOAT || + variance_weights.type_ != tensorflow::DataType::DT_FLOAT) { + return tensorflow::errors::Unimplemented( + "only float32 weights data type is supported, at " + node_def.name()); + } + for (size_t i=0; i(scale_weights.GetValues()))[i]; + float offset = (static_cast(offset_weights.GetValues()))[i]; + float mean = (static_cast(mean_weights.GetValues()))[i]; + float variance = (static_cast(variance_weights.GetValues()))[i]; + float& combined_scale_ref = const_cast( + static_cast(combined_scale_weights.GetValues()))[i]; + float& combined_offset_ref = const_cast( + static_cast(combined_offset_weights.GetValues()))[i]; + combined_scale_ref = scale / sqrtf(variance + epsilon); + combined_offset_ref = offset - mean * combined_scale_ref; + } + nvinfer1::IScaleLayer* layer = ctx.network()->addScale( + *const_cast(tensor), nvinfer1::ScaleMode::kCHANNEL, + combined_offset_weights, combined_scale_weights, dummy_power_weights); + nvinfer1::ITensor* output_tensor = layer->getOutput(0); + outputs->push_back(TRT_TensorOrWeights(output_tensor)); + return tensorflow::Status::OK(); +} + tensorflow::Status ConvertMatMul(Converter& ctx, tensorflow::NodeDef const& node_def, std::vector const& inputs, @@ -1827,6 +1892,8 @@ void Converter::register_op_converters() { op_registry_["ConcatV2"] = ConvertConcat; op_registry_["MatMul"] = ConvertMatMul; op_registry_["Reshape"] = ConvertReshape; + op_registry_["FusedBatchNorm"] = ConvertFusedBatchNorm; + op_registry_["FusedBatchNormV2"] = ConvertFusedBatchNorm; } } // namespace -- GitLab From 203caffbee9470109e3f750ba847e0aa4894a1e6 Mon Sep 17 00:00:00 2001 From: Rajendra arora Date: Fri, 16 Feb 2018 10:56:25 +0530 Subject: [PATCH 0025/1867] Documentation api reference badge added in Readme.md --- README.md | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/README.md b/README.md index 916e5200b2..efacf063e3 100644 --- a/README.md +++ b/README.md @@ -4,9 +4,10 @@ ----------------- -| **`Linux CPU`** | **`Linux GPU`** | **`Mac OS CPU`** | **`Windows CPU`** | **`Android`** | -|-----------------|---------------------|------------------|-------------------|---------------| -| [![Build Status](https://ci.tensorflow.org/buildStatus/icon?job=tensorflow-master-cpu)](https://ci.tensorflow.org/job/tensorflow-master-cpu) | [![Build Status](https://ci.tensorflow.org/buildStatus/icon?job=tensorflow-master-linux-gpu)](https://ci.tensorflow.org/job/tensorflow-master-linux-gpu) | [![Build Status](https://ci.tensorflow.org/buildStatus/icon?job=tensorflow-master-mac)](https://ci.tensorflow.org/job/tensorflow-master-mac) | [![Build Status](https://ci.tensorflow.org/buildStatus/icon?job=tensorflow-master-win-cmake-py)](https://ci.tensorflow.org/job/tensorflow-master-win-cmake-py) | [![Build Status](https://ci.tensorflow.org/buildStatus/icon?job=tensorflow-master-android)](https://ci.tensorflow.org/job/tensorflow-master-android) [ ![Download](https://api.bintray.com/packages/google/tensorflow/tensorflow/images/download.svg) ](https://bintray.com/google/tensorflow/tensorflow/_latestVersion) | + +| **`Documentation`** | **`Linux CPU`** | **`Linux GPU`** | **`Mac OS CPU`** | **`Windows CPU`** | **`Android`** | +|-----------------|---------------------|------------------|-------------------|---------------|---------------| +| [![Documentation](https://img.shields.io/badge/api-reference-blue.svg)](https://www.tensorflow.org/api_docs/) | [![Build Status](https://ci.tensorflow.org/buildStatus/icon?job=tensorflow-master-cpu)](https://ci.tensorflow.org/job/tensorflow-master-cpu) | [![Build Status](https://ci.tensorflow.org/buildStatus/icon?job=tensorflow-master-linux-gpu)](https://ci.tensorflow.org/job/tensorflow-master-linux-gpu) | [![Build Status](https://ci.tensorflow.org/buildStatus/icon?job=tensorflow-master-mac)](https://ci.tensorflow.org/job/tensorflow-master-mac) | [![Build Status](https://ci.tensorflow.org/buildStatus/icon?job=tensorflow-master-win-cmake-py)](https://ci.tensorflow.org/job/tensorflow-master-win-cmake-py) | [![Build Status](https://ci.tensorflow.org/buildStatus/icon?job=tensorflow-master-android)](https://ci.tensorflow.org/job/tensorflow-master-android) [ ![Download](https://api.bintray.com/packages/google/tensorflow/tensorflow/images/download.svg) ](https://bintray.com/google/tensorflow/tensorflow/_latestVersion) **TensorFlow** is an open source software library for numerical computation using data flow graphs. The graph nodes represent mathematical operations, while -- GitLab From 4e7772e0c74a663809f9fcf39545032eb8277e6a Mon Sep 17 00:00:00 2001 From: Rajendra arora Date: Fri, 16 Feb 2018 11:48:10 +0530 Subject: [PATCH 0026/1867] Added a contribution guideline header in readme.md --- README.md | 30 ++++++++++++++++-------------- 1 file changed, 16 insertions(+), 14 deletions(-) diff --git a/README.md b/README.md index efacf063e3..ef5bdc66ef 100644 --- a/README.md +++ b/README.md @@ -22,20 +22,6 @@ organization for the purposes of conducting machine learning and deep neural networks research. The system is general enough to be applicable in a wide variety of other domains, as well. -**If you want to contribute to TensorFlow, be sure to review the [contribution -guidelines](CONTRIBUTING.md). This project adheres to TensorFlow's -[code of conduct](CODE_OF_CONDUCT.md). By participating, you are expected to -uphold this code.** - -**We use [GitHub issues](https://github.com/tensorflow/tensorflow/issues) for -tracking requests and bugs. So please see -[TensorFlow Discuss](https://groups.google.com/a/tensorflow.org/forum/#!forum/discuss) for general questions -and discussion, and please direct specific questions to [Stack Overflow](https://stackoverflow.com/questions/tagged/tensorflow).** - -The TensorFlow project strives to abide by generally accepted best practices in open-source software development: - -[![CII Best Practices](https://bestpractices.coreinfrastructure.org/projects/1486/badge)](https://bestpractices.coreinfrastructure.org/projects/1486) - ## Installation *See [Installing TensorFlow](https://www.tensorflow.org/get_started/os_setup.html) for instructions on how to install our release binaries or how to build from source.* @@ -76,6 +62,22 @@ $ python >>> sess.close() ``` +## Contribution guidelines + +**If you want to contribute to TensorFlow, be sure to review the [contribution +guidelines](CONTRIBUTING.md). This project adheres to TensorFlow's +[code of conduct](CODE_OF_CONDUCT.md). By participating, you are expected to +uphold this code.** + +**We use [GitHub issues](https://github.com/tensorflow/tensorflow/issues) for +tracking requests and bugs. So please see +[TensorFlow Discuss](https://groups.google.com/a/tensorflow.org/forum/#!forum/discuss) for general questions +and discussion, and please direct specific questions to [Stack Overflow](https://stackoverflow.com/questions/tagged/tensorflow).** + +The TensorFlow project strives to abide by generally accepted best practices in open-source software development: + +[![CII Best Practices](https://bestpractices.coreinfrastructure.org/projects/1486/badge)](https://bestpractices.coreinfrastructure.org/projects/1486) + ## For more information * [TensorFlow Website](https://www.tensorflow.org) -- GitLab From f4d95b4abc45645ff5ed1670abc73fe0ffe49a82 Mon Sep 17 00:00:00 2001 From: kdavis-mozilla Date: Fri, 16 Feb 2018 11:41:19 +0100 Subject: [PATCH 0027/1867] Added Deep Speech use --- tensorflow/docs_src/about/uses.md | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/tensorflow/docs_src/about/uses.md b/tensorflow/docs_src/about/uses.md index 8818177a28..d646880bd3 100644 --- a/tensorflow/docs_src/about/uses.md +++ b/tensorflow/docs_src/about/uses.md @@ -22,6 +22,14 @@ This section describes some of the current uses of the TensorFlow system. > TensorFlow, or even better, send us a pull request to add an entry to this > file. +* **Deep Speech** +
    +
  • **Organization**: Mozilla
    • +
  • **Domain**: Speech Recognition
    • +
  • **Description**: A TensorFlow implementation motivated by Baidu's Deep Speech architecture.
    • +
  • **More info**: [GitHub Repo](https://github.com/mozilla/deepspeech)
    • +
+ * **RankBrain**