From 736e8c4ccb16718d11cf7c8e1fac843bf6e388a7 Mon Sep 17 00:00:00 2001 From: ManHyuk Date: Wed, 14 Feb 2018 18:26:20 +0900 Subject: [PATCH 001/528] fix typo --- tensorflow/core/lib/io/record_writer.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/core/lib/io/record_writer.cc b/tensorflow/core/lib/io/record_writer.cc index 3657243c5d..ebc5648269 100644 --- a/tensorflow/core/lib/io/record_writer.cc +++ b/tensorflow/core/lib/io/record_writer.cc @@ -49,7 +49,7 @@ RecordWriterOptions RecordWriterOptions::CreateRecordWriterOptions( #endif // IS_SLIM_BUILD } else if (compression_type != compression::kNone) { LOG(ERROR) << "Unsupported compression_type:" << compression_type - << ". No comprression will be used."; + << ". No compression will be used."; } return options; } -- GitLab From 617fa4e5fa634270c36a2a8762e6ce96bd38f2f8 Mon Sep 17 00:00:00 2001 From: ManHyuk Date: Wed, 14 Feb 2018 18:35:31 +0900 Subject: [PATCH 002/528] fix typo --- tensorflow/contrib/makefile/README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/contrib/makefile/README.md b/tensorflow/contrib/makefile/README.md index b0228c5435..995230dfa8 100644 --- a/tensorflow/contrib/makefile/README.md +++ b/tensorflow/contrib/makefile/README.md @@ -155,7 +155,7 @@ CC_PREFIX=ccache tensorflow/contrib/makefile/build_all_android.sh -s tensorflow/ (add -T on subsequent builds to skip protobuf downloading/building) -#### Testing the the CUDA-enabled benchmark via adb: +#### Testing the CUDA-enabled benchmark via adb: Build binaries first as above, then run: ```bash -- GitLab From 672ec270f96144bca5e1d75d002421c1e9b49921 Mon Sep 17 00:00:00 2001 From: Hovhannes Harutyunyan Date: Mon, 19 Feb 2018 12:56:40 +0400 Subject: [PATCH 003/528] Add broadcasting functionality fro Div and Sub ops. --- tensorflow/contrib/lite/kernels/div.cc | 117 ++++++-- tensorflow/contrib/lite/kernels/div_test.cc | 174 ++++++++++++ .../internal/optimized/optimized_ops.h | 268 +++++++++++++++++- .../internal/reference/reference_ops.h | 257 +++++++++++++++++ tensorflow/contrib/lite/kernels/sub.cc | 135 +++++++-- tensorflow/contrib/lite/kernels/sub_test.cc | 213 ++++++++++++++ .../testing/generated_examples_zip_test.cc | 15 +- 7 files changed, 1122 insertions(+), 57 deletions(-) create mode 100644 tensorflow/contrib/lite/kernels/div_test.cc create mode 100644 tensorflow/contrib/lite/kernels/sub_test.cc diff --git a/tensorflow/contrib/lite/kernels/div.cc b/tensorflow/contrib/lite/kernels/div.cc index 44bd0dc85d..c77a0de9b7 100644 --- a/tensorflow/contrib/lite/kernels/div.cc +++ b/tensorflow/contrib/lite/kernels/div.cc @@ -37,7 +37,23 @@ constexpr int kInputTensor1 = 0; constexpr int kInputTensor2 = 1; constexpr int kOutputTensor = 0; +struct OpData { + bool requires_broadcast; +}; + +void* Init(TfLiteContext* context, const char* buffer, size_t length) { + auto* data = new OpData; + data->requires_broadcast = false; + return data; +} + +void Free(TfLiteContext* context, void* buffer) { + delete reinterpret_cast(buffer); +} + TfLiteStatus Prepare(TfLiteContext* context, TfLiteNode* node) { + OpData* data = reinterpret_cast(node->user_data); + TF_LITE_ENSURE_EQ(context, NumInputs(node), 2); TF_LITE_ENSURE_EQ(context, NumOutputs(node), 1); @@ -45,35 +61,85 @@ TfLiteStatus Prepare(TfLiteContext* context, TfLiteNode* node) { TfLiteTensor* input2 = GetInput(context, node, kInputTensor2); TfLiteTensor* output = GetOutput(context, node, kOutputTensor); - TF_LITE_ENSURE_EQ(context, NumDimensions(input1), NumDimensions(input2)); - for (int i = 0; i < NumDimensions(input1); ++i) { - TF_LITE_ENSURE_EQ(context, SizeOfDimension(input1, i), - SizeOfDimension(input2, i)); - } + TF_LITE_ENSURE_EQ(context, input1->type, input2->type); + output->type = input2->type; + + data->requires_broadcast = !HaveSameShapes(input1, input2); - TF_LITE_ENSURE_EQ(context, input1->type, output->type); - TF_LITE_ENSURE_EQ(context, input2->type, output->type); + TfLiteIntArray* output_size = nullptr; + if (data->requires_broadcast) { + TF_LITE_ENSURE_OK(context, CalculateShapeForBroadcast( + context, input1, input2, &output_size)); + } else { + output_size = TfLiteIntArrayCopy(input1->dims); + } - TfLiteIntArray* output_size = TfLiteIntArrayCopy(input1->dims); return context->ResizeTensor(context, output, output_size); } template -void EvalDivFloat(TfLiteContext* context, TfLiteNode* node, - TfLiteDivParams* params, TfLiteTensor* input1, - TfLiteTensor* input2, TfLiteTensor* output) { +void EvalFloat(TfLiteContext* context, TfLiteNode* node, + TfLiteDivParams* params, const OpData* data, + TfLiteTensor* input1, TfLiteTensor* input2, + TfLiteTensor* output) { float output_activation_min, output_activation_max; CalculateActivationRangeFloat(params->activation, &output_activation_min, &output_activation_max); -#define TF_LITE_DIV(type) \ - type::Div(GetTensorData(input1), GetTensorDims(input1), \ - GetTensorData(input2), GetTensorDims(input2), \ - output_activation_min, output_activation_max, \ - GetTensorData(output), GetTensorDims(output)) +#define TF_LITE_DIV(type, opname) \ + type::opname(GetTensorData(input1), GetTensorDims(input1), \ + GetTensorData(input2), GetTensorDims(input2), \ + output_activation_min, output_activation_max, \ + GetTensorData(output), GetTensorDims(output)) + if (kernel_type == kReference) { + if (data->requires_broadcast) { + TF_LITE_DIV(reference_ops, BroadcastDiv); + } else { + TF_LITE_DIV(reference_ops, Div); + } + } else { + if (data->requires_broadcast) { + TF_LITE_DIV(optimized_ops, BroadcastDiv); + } else { + TF_LITE_DIV(optimized_ops, Div); + } + } +#undef TF_LITE_DIV +} + +template +void EvalQuantized(TfLiteContext* context, TfLiteNode* node, + TfLiteDivParams* params, const OpData* data, + TfLiteTensor* input1, TfLiteTensor* input2, + TfLiteTensor* output) { + auto input1_offset = -input1->params.zero_point; + auto input2_offset = -input2->params.zero_point; + auto output_offset = output->params.zero_point; + + int32_t output_multiplier; + int output_shift; + + double real_multiplier = + input1->params.scale * input2->params.scale / output->params.scale; + QuantizeMultiplierSmallerThanOne(real_multiplier, &output_multiplier, + &output_shift); + + int32 output_activation_min, output_activation_max; + CalculateActivationRangeUint8(params->activation, output, + &output_activation_min, &output_activation_max); + +#define TF_LITE_DIV(type, opname) \ + type::opname(GetTensorData(input1), GetTensorDims(input1), \ + input1_offset, GetTensorData(input2), \ + GetTensorDims(input2), input2_offset, output_offset, \ + output_multiplier, output_shift, output_activation_min, \ + output_activation_max, GetTensorData(output), \ + GetTensorDims(output)); + // The quantized version of Div doesn't support activations, so we + // always use BroadcastDiv. if (kernel_type == kReference) { - TF_LITE_DIV(reference_ops); + TF_LITE_DIV(reference_ops, BroadcastDiv); } else { - TF_LITE_DIV(optimized_ops); + TF_LITE_DIV(optimized_ops, BroadcastDiv); } #undef TF_LITE_DIV } @@ -81,15 +147,20 @@ void EvalDivFloat(TfLiteContext* context, TfLiteNode* node, template TfLiteStatus Eval(TfLiteContext* context, TfLiteNode* node) { auto* params = reinterpret_cast(node->builtin_data); + OpData* data = reinterpret_cast(node->user_data); TfLiteTensor* input1 = GetInput(context, node, kInputTensor1); TfLiteTensor* input2 = GetInput(context, node, kInputTensor2); TfLiteTensor* output = GetOutput(context, node, kOutputTensor); if (output->type == kTfLiteFloat32) { - EvalDivFloat(context, node, params, input1, input2, output); + EvalFloat(context, node, params, data, input1, input2, output); + } else if (output->type == kTfLiteUInt8) { + EvalQuantized(context, node, params, data, input1, input2, + output); } else { - context->ReportError(context, "Inputs and outputs not all float types."); + context->ReportError(context, + "Div only supports FLOAT32 and quantized UINT8 now."); return kTfLiteError; } @@ -99,19 +170,19 @@ TfLiteStatus Eval(TfLiteContext* context, TfLiteNode* node) { } // namespace div TfLiteRegistration* Register_DIV_REF() { - static TfLiteRegistration r = {nullptr, nullptr, div::Prepare, + static TfLiteRegistration r = {div::Init, div::Free, div::Prepare, div::Eval}; return &r; } TfLiteRegistration* Register_DIV_GENERIC_OPT() { - static TfLiteRegistration r = {nullptr, nullptr, div::Prepare, + static TfLiteRegistration r = {div::Init, div::Free, div::Prepare, div::Eval}; return &r; } TfLiteRegistration* Register_DIV_NEON_OPT() { - static TfLiteRegistration r = {nullptr, nullptr, div::Prepare, + static TfLiteRegistration r = {div::Init, div::Free, div::Prepare, div::Eval}; return &r; } diff --git a/tensorflow/contrib/lite/kernels/div_test.cc b/tensorflow/contrib/lite/kernels/div_test.cc new file mode 100644 index 0000000000..78918a0d79 --- /dev/null +++ b/tensorflow/contrib/lite/kernels/div_test.cc @@ -0,0 +1,174 @@ +/* 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. +==============================================================================*/ +#include +#include "tensorflow/contrib/lite/interpreter.h" +#include "tensorflow/contrib/lite/kernels/register.h" +#include "tensorflow/contrib/lite/kernels/test_util.h" +#include "tensorflow/contrib/lite/model.h" + +namespace tflite { +namespace { + +using ::testing::ElementsAreArray; + +class BaseDivOpModel : public SingleOpModel { + public: + BaseDivOpModel(const TensorData& input1, const TensorData& input2, + const TensorData& output, + ActivationFunctionType activation_type) { + input1_ = AddInput(input1); + input2_ = AddInput(input2); + output_ = AddOutput(output); + SetBuiltinOp(BuiltinOperator_DIV, BuiltinOptions_DivOptions, + CreateDivOptions(builder_, activation_type).Union()); + BuildInterpreter({GetShape(input1_), GetShape(input2_)}); + } + + int input1() { return input1_; } + int input2() { return input2_; } + + protected: + int input1_; + int input2_; + int output_; +}; + +class FloatDivOpModel : public BaseDivOpModel { + public: + using BaseDivOpModel::BaseDivOpModel; + + std::vector GetOutput() { return ExtractVector(output_); } +}; + +// For quantized Div, the error shouldn't exceed (2*step + step^2). +// The param min=-1.0 & max=1.0 is used in the following tests. +// The tolerance value is ~0.0157. +const float kQuantizedStep = 2.0 / 255.0; +const float kQuantizedTolerance = + 2.0 * kQuantizedStep + kQuantizedStep * kQuantizedStep; + +class QuantizedDivOpModel : public BaseDivOpModel { + public: + using BaseDivOpModel::BaseDivOpModel; + + std::vector GetDequantizedOutput() { + return Dequantize(ExtractVector(output_), + GetScale(output_), GetZeroPoint(output_)); + } +}; + +TEST(FloatDivOpTest, NoActivation) { + FloatDivOpModel m({TensorType_FLOAT32, {1, 2, 2, 1}}, + {TensorType_FLOAT32, {1, 2, 2, 1}}, + {TensorType_FLOAT32, {}}, ActivationFunctionType_NONE); + m.PopulateTensor(m.input1(), {-0.2, 0.2, -1.2, 0.8}); + m.PopulateTensor(m.input2(), {0.5, 0.2, -1.5, 0.5}); + m.Invoke(); + EXPECT_THAT(m.GetOutput(), + ElementsAreArray(ArrayFloatNear({-0.4, 1.0, 0.8, 1.6}))); +} + +TEST(FloatDivOpTest, ActivationRELU_N1_TO_1) { + FloatDivOpModel m( + {TensorType_FLOAT32, {1, 2, 2, 1}}, {TensorType_FLOAT32, {1, 2, 2, 1}}, + {TensorType_FLOAT32, {}}, ActivationFunctionType_RELU_N1_TO_1); + m.PopulateTensor(m.input1(), {-0.2, 0.2, -1.2, 0.8}); + m.PopulateTensor(m.input2(), {0.1, 0.2, -1.5, 0.5}); + m.Invoke(); + EXPECT_THAT(m.GetOutput(), + ElementsAreArray(ArrayFloatNear({-1.0, 1.0, 0.8, 1.0}))); +} + +TEST(FloatDivOpTest, VariousInputShapes) { + std::vector> test_shapes = { + {6}, {2, 3}, {2, 1, 3}, {1, 3, 1, 2}}; + for (int i = 0; i < test_shapes.size(); ++i) { + FloatDivOpModel m({TensorType_FLOAT32, test_shapes[i]}, + {TensorType_FLOAT32, test_shapes[i]}, + {TensorType_FLOAT32, {}}, ActivationFunctionType_NONE); + m.PopulateTensor(m.input1(), {-2.0, 0.2, 0.3, 0.8, 1.1, -2.0}); + m.PopulateTensor(m.input2(), {0.1, 0.2, 0.6, 0.5, -1.1, -0.1}); + m.Invoke(); + EXPECT_THAT( + m.GetOutput(), + ElementsAreArray(ArrayFloatNear({-20.0, 1.0, 0.5, 1.6, -1.0, 20.0}))) + << "With shape number " << i; + } +} + +TEST(FloatDivOpTest, WithBroadcast) { + std::vector> test_shapes = { + {6}, {2, 3}, {2, 1, 3}, {1, 3, 1, 2}}; + for (int i = 0; i < test_shapes.size(); ++i) { + FloatDivOpModel m({TensorType_FLOAT32, test_shapes[i]}, + {TensorType_FLOAT32, {}}, // always a scalar + {TensorType_FLOAT32, {}}, ActivationFunctionType_NONE); + m.PopulateTensor(m.input1(), {-0.2, 0.2, 0.07, 0.08, 0.11, -0.123}); + m.PopulateTensor(m.input2(), {0.1}); + m.Invoke(); + EXPECT_THAT( + m.GetOutput(), + ElementsAreArray(ArrayFloatNear({-2.0, 2.0, 0.7, 0.8, 1.1, -1.23}))) + << "With shape number " << i; + } +} + +TEST(QuantizedDivOpTest, NoActivation) { + QuantizedDivOpModel m({TensorType_UINT8, {1, 2, 2, 1}, -1.0, 1.0}, + {TensorType_UINT8, {1, 2, 2, 1}, -1.0, 1.0}, + {TensorType_UINT8, {}, -1.0, 1.0}, + ActivationFunctionType_NONE); + m.QuantizeAndPopulate(m.input1(), {-0.6, 0.2, 0.9, -0.7}); + m.QuantizeAndPopulate(m.input2(), {0.8, 0.4, 0.9, -0.8}); + m.Invoke(); + EXPECT_THAT(m.GetDequantizedOutput(), + ElementsAreArray(ArrayFloatNear({-0.75, 0.5, 1.0, 0.875}, + kQuantizedTolerance))); +} + +// for quantized Div, the error shouldn't exceed 2*step +float GetTolerance(int min, int max) { + float kQuantizedStep = (max - min) / 255.0; + float kQuantizedTolerance = 2.0 * kQuantizedStep; + return kQuantizedTolerance; +} + +TEST(QuantizedDivOpTest, WithBroadcast) { + float kQuantizedTolerance = GetTolerance(-3.0, 3.0); + std::vector> test_shapes = { + {6}, {2, 3}, {2, 1, 3}, {1, 3, 1, 2}}; + for (int i = 0; i < test_shapes.size(); ++i) { + QuantizedDivOpModel m({TensorType_UINT8, test_shapes[i], -3.0, 3.0}, + {TensorType_UINT8, {}, -3.0, 3.0}, // always a scalar + {TensorType_UINT8, {}, -3.0, 3.0}, + ActivationFunctionType_NONE); + m.QuantizeAndPopulate(m.input1(), {-0.2, 0.2, 0.07, 0.08, 0.11, -0.123}); + m.QuantizeAndPopulate(m.input2(), {0.1}); + m.Invoke(); + EXPECT_THAT(m.GetDequantizedOutput(), + ElementsAreArray(ArrayFloatNear( + {-2.0, 2.0, 0.7, 0.8, 1.1, -1.23}, kQuantizedTolerance))) + << "With shape number " << i; + } +} + +} // namespace +} // namespace tflite + +int main(int argc, char** argv) { + ::tflite::LogToStderr(); + ::testing::InitGoogleTest(&argc, argv); + return RUN_ALL_TESTS(); +} diff --git a/tensorflow/contrib/lite/kernels/internal/optimized/optimized_ops.h b/tensorflow/contrib/lite/kernels/internal/optimized/optimized_ops.h index dec58fea4f..d12a3eca1d 100644 --- a/tensorflow/contrib/lite/kernels/internal/optimized/optimized_ops.h +++ b/tensorflow/contrib/lite/kernels/internal/optimized/optimized_ops.h @@ -1928,6 +1928,126 @@ inline void Div(const float* input1_data, const Dims<4>& input1_dims, } } +// TODO(jiawen): We can implement BroadcastDiv on buffers of arbitrary +// dimensionality if the runtime code does a single loop over one dimension +// that handles broadcasting as the base case. The code generator would then +// generate max(D1, D2) nested for loops. +// TODO(benoitjacob): BroadcastDiv is intentionally duplicated from +// reference_ops.h. Once an optimized version is implemented and NdArrayDesc +// is no longer referenced in this file, move NdArrayDesc from types.h to +// reference_ops.h. +template +void BroadcastDiv(const T* input1_data, const Dims<4>& input1_dims, + const T* input2_data, const Dims<4>& input2_dims, + T output_activation_min, T output_activation_max, + T* output_data, const Dims<4>& output_dims) { + gemmlowp::ScopedProfilingLabel label("BroadcastDiv"); + + NdArrayDesc<4> desc1; + NdArrayDesc<4> desc2; + NdArrayDescsForElementwiseBroadcast(input1_dims, input2_dims, &desc1, &desc2); + + // In Tensorflow, the dimensions are canonically named (batch_number, row, + // col, channel), with extents (batches, height, width, depth), with the + // trailing dimension changing most rapidly (channels has the smallest stride, + // typically 1 element). + // + // In generated C code, we store arrays with the dimensions reversed. The + // first dimension has smallest stride. + // + // We name our variables by their Tensorflow convention, but generate C code + // nesting loops such that the innermost loop has the smallest stride for the + // best cache behavior. + for (int b = 0; b < ArraySize(output_dims, 3); ++b) { + for (int y = 0; y < ArraySize(output_dims, 2); ++y) { + for (int x = 0; x < ArraySize(output_dims, 1); ++x) { + for (int c = 0; c < ArraySize(output_dims, 0); ++c) { + output_data[Offset(output_dims, c, x, y, b)] = + ActivationFunctionWithMinMax( + input1_data[SubscriptToIndex(desc1, c, x, y, b)] / + input2_data[SubscriptToIndex(desc2, c, x, y, b)], + output_activation_min, output_activation_max); + } + } + } + } +} + +// legacy, for compatibility with old checked-in code +template +void BroadcastDiv(const T* input1_data, const Dims<4>& input1_dims, + const T* input2_data, const Dims<4>& input2_dims, + T* output_data, const Dims<4>& output_dims) { + T output_activation_min, output_activation_max; + GetActivationMinMax(Ac, &output_activation_min, &output_activation_max); + + BroadcastDiv(input1_data, input1_dims, input2_data, input2_dims, + output_activation_min, output_activation_max, output_data, + output_dims); +} + +inline void BroadcastDiv(const uint8* input1_data, const Dims<4>& input1_dims, + int32 input1_offset, const uint8* input2_data, + const Dims<4>& input2_dims, int32 input2_offset, + int32 output_offset, int32 output_multiplier, + int output_shift, int32 output_activation_min, + int32 output_activation_max, uint8* output_data, + const Dims<4>& output_dims) { + gemmlowp::ScopedProfilingLabel label("BroadcastDiv/8bit"); + + NdArrayDesc<4> desc1; + NdArrayDesc<4> desc2; + NdArrayDescsForElementwiseBroadcast(input1_dims, input2_dims, &desc1, &desc2); + + // In Tensorflow, the dimensions are canonically named (batch_number, row, + // col, channel), with extents (batches, height, width, depth), with the + // trailing dimension changing most rapidly (channels has the smallest stride, + // typically 1 element). + // + // In generated C code, we store arrays with the dimensions reversed. The + // first dimension has smallest stride. + // + // We name our variables by their Tensorflow convention, but generate C code + // nesting loops such that the innermost loop has the smallest stride for the + // best cache behavior. + for (int b = 0; b < ArraySize(output_dims, 3); ++b) { + for (int y = 0; y < ArraySize(output_dims, 2); ++y) { + for (int x = 0; x < ArraySize(output_dims, 1); ++x) { + for (int c = 0; c < ArraySize(output_dims, 0); ++c) { + const int32 input1_val = + input1_offset + input1_data[SubscriptToIndex(desc1, c, x, y, b)]; + const int32 input2_val = + input2_offset + input2_data[SubscriptToIndex(desc2, c, x, y, b)]; + const int32 unclamped_result = + output_offset + + MultiplyByQuantizedMultiplierSmallerThanOne( + input1_val / input2_val, output_multiplier, output_shift); + const int32 clamped_output = + std::min(output_activation_max, + std::max(output_activation_min, unclamped_result)); + output_data[Offset(output_dims, c, x, y, b)] = + static_cast(clamped_output); + } + } + } + } +} + +// legacy, for compatibility with old checked-in code +template +inline void BroadcastDiv(const uint8* input1_data, const Dims<4>& input1_dims, + int32 input1_offset, const uint8* input2_data, + const Dims<4>& input2_dims, int32 input2_offset, + int32 output_offset, int32 output_multiplier, + int output_shift, int32 output_activation_min, + int32 output_activation_max, uint8* output_data, + const Dims<4>& output_dims) { + BroadcastDiv(input1_data, input1_dims, input1_offset, input2_data, + input2_dims, input2_offset, output_offset, output_multiplier, + output_shift, output_activation_min, output_activation_max, + output_data, output_dims); +} + // TODO(aselle): This is not actually optimized yet. inline void Sub(const float* input1_data, const Dims<4>& input1_dims, const float* input2_data, const Dims<4>& input2_dims, @@ -1955,6 +2075,152 @@ inline void Sub(const float* input1_data, const Dims<4>& input1_dims, } } } + +// TODO(jiawen): We can implement BroadcastSub on buffers of arbitrary +// dimensionality if the runtime code does a single loop over one dimension +// that handles broadcasting as the base case. The code generator would then +// generate max(D1, D2) nested for loops. +// TODO(benoitjacob): BroadcastSub is intentionally duplicated from +// reference_ops.h. Once an optimized version is implemented and NdArrayDesc +// is no longer referenced in this file, move NdArrayDesc from types.h to +// reference_ops.h. +template +void BroadcastSub(const T* input1_data, const Dims<4>& input1_dims, + const T* input2_data, const Dims<4>& input2_dims, + T output_activation_min, T output_activation_max, + T* output_data, const Dims<4>& output_dims) { + gemmlowp::ScopedProfilingLabel label("BroadcastSub"); + + NdArrayDesc<4> desc1; + NdArrayDesc<4> desc2; + NdArrayDescsForElementwiseBroadcast(input1_dims, input2_dims, &desc1, &desc2); + + // In Tensorflow, the dimensions are canonically named (batch_number, row, + // col, channel), with extents (batches, height, width, depth), with the + // trailing dimension changing most rapidly (channels has the smallest stride, + // typically 1 element). + // + // In generated C code, we store arrays with the dimensions reversed. The + // first dimension has smallest stride. + // + // We name our variables by their Tensorflow convention, but generate C code + // nesting loops such that the innermost loop has the smallest stride for the + // best cache behavior. + for (int b = 0; b < ArraySize(output_dims, 3); ++b) { + for (int y = 0; y < ArraySize(output_dims, 2); ++y) { + for (int x = 0; x < ArraySize(output_dims, 1); ++x) { + for (int c = 0; c < ArraySize(output_dims, 0); ++c) { + output_data[Offset(output_dims, c, x, y, b)] = + ActivationFunctionWithMinMax( + input1_data[SubscriptToIndex(desc1, c, x, y, b)] - + input2_data[SubscriptToIndex(desc2, c, x, y, b)], + output_activation_min, output_activation_max); + } + } + } + } +} + +// legacy, for compatibility with old checked-in code +template +void BroadcastSub(const T* input1_data, const Dims<4>& input1_dims, + const T* input2_data, const Dims<4>& input2_dims, + T* output_data, const Dims<4>& output_dims) { + T output_activation_min, output_activation_max; + GetActivationMinMax(Ac, &output_activation_min, &output_activation_max); + + BroadcastSub(input1_data, input1_dims, input2_data, input2_dims, + output_activation_min, output_activation_max, output_data, + output_dims); +} + +inline void BroadcastSub(int left_shift, const uint8* input1_data, + const Dims<4>& input1_dims, int32 input1_offset, + int32 input1_multiplier, int input1_shift, + const uint8* input2_data, const Dims<4>& input2_dims, + int32 input2_offset, int32 input2_multiplier, + int input2_shift, int32 output_offset, + int32 output_multiplier, int output_shift, + int32 output_activation_min, + int32 output_activation_max, uint8* output_data, + const Dims<4>& output_dims) { + gemmlowp::ScopedProfilingLabel label("BroadcastSub/8bit"); + + NdArrayDesc<4> desc1; + NdArrayDesc<4> desc2; + NdArrayDescsForElementwiseBroadcast(input1_dims, input2_dims, &desc1, &desc2); + + // In Tensorflow, the dimensions are canonically named (batch_number, row, + // col, channel), with extents (batches, height, width, depth), with the + // trailing dimension changing most rapidly (channels has the smallest stride, + // typically 1 element). + // + // In generated C code, we store arrays with the dimensions reversed. The + // first dimension has smallest stride. + // + // We name our variables by their Tensorflow convention, but generate C code + // nesting loops such that the innermost loop has the smallest stride for the + // best cache behavior. + for (int b = 0; b < ArraySize(output_dims, 3); ++b) { + for (int y = 0; y < ArraySize(output_dims, 2); ++y) { + for (int x = 0; x < ArraySize(output_dims, 1); ++x) { + for (int c = 0; c < ArraySize(output_dims, 0); ++c) { + const int32 input1_val = + input1_offset + input1_data[SubscriptToIndex(desc1, c, x, y, b)]; + const int32 input2_val = + input2_offset + input2_data[SubscriptToIndex(desc2, c, x, y, b)]; + const int32 shifted_input1_val = input1_val * (1 << left_shift); + const int32 shifted_input2_val = input2_val * (1 << left_shift); + const int32 scaled_input1_val = + MultiplyByQuantizedMultiplierSmallerThanOne( + shifted_input1_val, input1_multiplier, input1_shift); + const int32 scaled_input2_val = + MultiplyByQuantizedMultiplierSmallerThanOne( + shifted_input2_val, input2_multiplier, input2_shift); + const int32 raw_sum = scaled_input1_val - scaled_input2_val; + const int32 raw_output = + MultiplyByQuantizedMultiplierSmallerThanOne( + raw_sum, output_multiplier, output_shift) + + output_offset; + const int32 clamped_output = + std::min(output_activation_max, + std::max(output_activation_min, raw_output)); + output_data[Offset(output_dims, c, x, y, b)] = + static_cast(clamped_output); + } + } + } + } +} + +template +inline void BroadcastSub(int left_shift, const uint8* input1_data, + const Dims<4>& input1_dims, int32 input1_offset, + int32 input1_multiplier, int input1_shift, + const uint8* input2_data, const Dims<4>& input2_dims, + int32 input2_offset, int32 input2_multiplier, + int input2_shift, int32 output_offset, + int32 output_multiplier, int output_shift, + int32 output_activation_min, + int32 output_activation_max, uint8* output_data, + const Dims<4>& output_dims) { + static_assert(Ac == FusedActivationFunctionType::kNone || + Ac == FusedActivationFunctionType::kRelu || + Ac == FusedActivationFunctionType::kRelu6 || + Ac == FusedActivationFunctionType::kRelu1, + ""); + TFLITE_DCHECK_LE(output_activation_min, output_activation_max); + if (Ac == FusedActivationFunctionType::kNone) { + TFLITE_DCHECK_EQ(output_activation_min, 0); + TFLITE_DCHECK_EQ(output_activation_max, 255); + } + BroadcastSub(left_shift, input1_data, input1_dims, input1_offset, + input1_multiplier, input1_shift, input2_data, input2_dims, + input2_offset, input2_multiplier, input2_shift, output_offset, + output_multiplier, output_shift, output_activation_min, + output_activation_max, output_data, output_dims); +} + template void Concatenation(int concat_dim, const Scalar* const* input_data, const Dims<4>* const* input_dims, int inputs_count, @@ -2866,7 +3132,7 @@ inline void Softmax(const uint8* input_data, const Dims<4>& input_dims, using FixedPointAccum = gemmlowp::FixedPoint; using FixedPoint0 = gemmlowp::FixedPoint; - gemmlowp::ScopedProfilingLabel label("Softmax/8bit"); +gemmlowp::ScopedProfilingLabel label("Softmax/8bit"); const int batches = MatchingArraySize(input_dims, 3, output_dims, 3); const int height = MatchingArraySize(input_dims, 2, output_dims, 2); const int width = MatchingArraySize(input_dims, 1, output_dims, 1); diff --git a/tensorflow/contrib/lite/kernels/internal/reference/reference_ops.h b/tensorflow/contrib/lite/kernels/internal/reference/reference_ops.h index 5f4d5be323..c7b7687622 100644 --- a/tensorflow/contrib/lite/kernels/internal/reference/reference_ops.h +++ b/tensorflow/contrib/lite/kernels/internal/reference/reference_ops.h @@ -1208,6 +1208,122 @@ inline void Div(const float* input1_data, const Dims<4>& input1_dims, } } +// TODO(jiawen): We can implement BroadcastDiv on buffers of arbitrary +// dimensionality if the runtime code does a single loop over one dimension +// that handles broadcasting as the base case. The code generator would then +// generate max(D1, D2) nested for loops. +template +void BroadcastDiv(const T* input1_data, const Dims<4>& input1_dims, + const T* input2_data, const Dims<4>& input2_dims, + T output_activation_min, T output_activation_max, + T* output_data, const Dims<4>& output_dims) { + gemmlowp::ScopedProfilingLabel label("BroadcastDiv"); + + NdArrayDesc<4> desc1; + NdArrayDesc<4> desc2; + NdArrayDescsForElementwiseBroadcast(input1_dims, input2_dims, &desc1, &desc2); + + // In Tensorflow, the dimensions are canonically named (batch_number, row, + // col, channel), with extents (batches, height, width, depth), with the + // trailing dimension changing most rapidly (channels has the smallest + // stride, typically 1 element). + // + // In generated C code, we store arrays with the dimensions reversed. The + // first dimension has smallest stride. + // + // We name our variables by their Tensorflow convention, but generate C code + // nesting loops such that the innermost loop has the smallest stride for + // the best cache behavior. + for (int b = 0; b < ArraySize(output_dims, 3); ++b) { + for (int y = 0; y < ArraySize(output_dims, 2); ++y) { + for (int x = 0; x < ArraySize(output_dims, 1); ++x) { + for (int c = 0; c < ArraySize(output_dims, 0); ++c) { + output_data[Offset(output_dims, c, x, y, b)] = + ActivationFunctionWithMinMax( + input1_data[SubscriptToIndex(desc1, c, x, y, b)] / + input2_data[SubscriptToIndex(desc2, c, x, y, b)], + output_activation_min, output_activation_max); + } + } + } + } +} + +// legacy, for compatibility with old checked-in code +template +void BroadcastDiv(const T* input1_data, const Dims<4>& input1_dims, + const T* input2_data, const Dims<4>& input2_dims, + T* output_data, const Dims<4>& output_dims) { + T output_activation_min, output_activation_max; + GetActivationMinMax(Ac, &output_activation_min, &output_activation_max); + + BroadcastDiv(input1_data, input1_dims, input2_data, input2_dims, + output_activation_min, output_activation_max, output_data, + output_dims); +} + +inline void BroadcastDiv(const uint8* input1_data, const Dims<4>& input1_dims, + int32 input1_offset, const uint8* input2_data, + const Dims<4>& input2_dims, int32 input2_offset, + int32 output_offset, int32 output_multiplier, + int output_shift, int32 output_activation_min, + int32 output_activation_max, uint8* output_data, + const Dims<4>& output_dims) { + gemmlowp::ScopedProfilingLabel label("BroadcastDiv/8bit"); + + NdArrayDesc<4> desc1; + NdArrayDesc<4> desc2; + NdArrayDescsForElementwiseBroadcast(input1_dims, input2_dims, &desc1, &desc2); + + // In Tensorflow, the dimensions are canonically named (batch_number, row, + // col, channel), with extents (batches, height, width, depth), with the + // trailing dimension changing most rapidly (channels has the smallest + // stride, typically 1 element). + // + // In generated C code, we store arrays with the dimensions reversed. The + // first dimension has smallest stride. + // + // We name our variables by their Tensorflow convention, but generate C code + // nesting loops such that the innermost loop has the smallest stride for + // the best cache behavior. + for (int b = 0; b < ArraySize(output_dims, 3); ++b) { + for (int y = 0; y < ArraySize(output_dims, 2); ++y) { + for (int x = 0; x < ArraySize(output_dims, 1); ++x) { + for (int c = 0; c < ArraySize(output_dims, 0); ++c) { + const int32 input1_val = + input1_offset + input1_data[SubscriptToIndex(desc1, c, x, y, b)]; + const int32 input2_val = + input2_offset + input2_data[SubscriptToIndex(desc2, c, x, y, b)]; + const int32 unclamped_result = + output_offset + + MultiplyByQuantizedMultiplierSmallerThanOne( + input1_val / input2_val, output_multiplier, output_shift); + const int32 clamped_output = + std::min(output_activation_max, + std::max(output_activation_min, unclamped_result)); + output_data[Offset(output_dims, c, x, y, b)] = + static_cast(clamped_output); + } + } + } + } +} + +// legacy, for compatibility with old checked-in code +template +inline void BroadcastDiv(const uint8* input1_data, const Dims<4>& input1_dims, + int32 input1_offset, const uint8* input2_data, + const Dims<4>& input2_dims, int32 input2_offset, + int32 output_offset, int32 output_multiplier, + int output_shift, int32 output_activation_min, + int32 output_activation_max, uint8* output_data, + const Dims<4>& output_dims) { + BroadcastDiv(input1_data, input1_dims, input1_offset, input2_data, + input2_dims, input2_offset, output_offset, output_multiplier, + output_shift, output_activation_min, output_activation_max, + output_data, output_dims); +} + inline void Sub(const float* input1_data, const Dims<4>& input1_dims, const float* input2_data, const Dims<4>& input2_dims, float output_activation_min, float output_activation_max, @@ -1235,6 +1351,147 @@ inline void Sub(const float* input1_data, const Dims<4>& input1_dims, } } +// TODO(jiawen): We can implement BroadcastSub on buffers of arbitrary +// dimensionality if the runtime code does a single loop over one dimension +// that handles broadcasting as the base case. The code generator would then +// generate max(D1, D2) nested for loops. +template +void BroadcastSub(const T* input1_data, const Dims<4>& input1_dims, + const T* input2_data, const Dims<4>& input2_dims, + T output_activation_min, T output_activation_max, + T* output_data, const Dims<4>& output_dims) { + gemmlowp::ScopedProfilingLabel label("BroadcastSub"); + + NdArrayDesc<4> desc1; + NdArrayDesc<4> desc2; + NdArrayDescsForElementwiseBroadcast(input1_dims, input2_dims, &desc1, &desc2); + + // In Tensorflow, the dimensions are canonically named (batch_number, row, + // col, channel), with extents (batches, height, width, depth), with the + // trailing dimension changing most rapidly (channels has the smallest stride, + // typically 1 element). + // + // In generated C code, we store arrays with the dimensions reversed. The + // first dimension has smallest stride. + // + // We name our variables by their Tensorflow convention, but generate C code + // nesting loops such that the innermost loop has the smallest stride for the + // best cache behavior. + for (int b = 0; b < ArraySize(output_dims, 3); ++b) { + for (int y = 0; y < ArraySize(output_dims, 2); ++y) { + for (int x = 0; x < ArraySize(output_dims, 1); ++x) { + for (int c = 0; c < ArraySize(output_dims, 0); ++c) { + output_data[Offset(output_dims, c, x, y, b)] = + ActivationFunctionWithMinMax( + input1_data[SubscriptToIndex(desc1, c, x, y, b)] - + input2_data[SubscriptToIndex(desc2, c, x, y, b)], + output_activation_min, output_activation_max); + } + } + } + } +} + +// legacy, for compatibility with old checked-in code +template +void BroadcastSub(const T* input1_data, const Dims<4>& input1_dims, + const T* input2_data, const Dims<4>& input2_dims, + T* output_data, const Dims<4>& output_dims) { + T output_activation_min, output_activation_max; + GetActivationMinMax(Ac, &output_activation_min, &output_activation_max); + + BroadcastSub(input1_data, input1_dims, input2_data, input2_dims, + output_activation_min, output_activation_max, output_data, + output_dims); +} + +inline void BroadcastSub(int left_shift, const uint8* input1_data, + const Dims<4>& input1_dims, int32 input1_offset, + int32 input1_multiplier, int input1_shift, + const uint8* input2_data, const Dims<4>& input2_dims, + int32 input2_offset, int32 input2_multiplier, + int input2_shift, int32 output_offset, + int32 output_multiplier, int output_shift, + int32 output_activation_min, + int32 output_activation_max, uint8* output_data, + const Dims<4>& output_dims) { + gemmlowp::ScopedProfilingLabel label("BroadcastSub/8bit"); + + NdArrayDesc<4> desc1; + NdArrayDesc<4> desc2; + NdArrayDescsForElementwiseBroadcast(input1_dims, input2_dims, &desc1, &desc2); + + // In Tensorflow, the dimensions are canonically named (batch_number, row, + // col, channel), with extents (batches, height, width, depth), with the + // trailing dimension changing most rapidly (channels has the smallest stride, + // typically 1 element). + // + // In generated C code, we store arrays with the dimensions reversed. The + // first dimension has smallest stride. + // + // We name our variables by their Tensorflow convention, but generate C code + // nesting loops such that the innermost loop has the smallest stride for the + // best cache behavior. + for (int b = 0; b < ArraySize(output_dims, 3); ++b) { + for (int y = 0; y < ArraySize(output_dims, 2); ++y) { + for (int x = 0; x < ArraySize(output_dims, 1); ++x) { + for (int c = 0; c < ArraySize(output_dims, 0); ++c) { + const int32 input1_val = + input1_offset + input1_data[SubscriptToIndex(desc1, c, x, y, b)]; + const int32 input2_val = + input2_offset + input2_data[SubscriptToIndex(desc2, c, x, y, b)]; + const int32 shifted_input1_val = input1_val * (1 << left_shift); + const int32 shifted_input2_val = input2_val * (1 << left_shift); + const int32 scaled_input1_val = + MultiplyByQuantizedMultiplierSmallerThanOne( + shifted_input1_val, input1_multiplier, input1_shift); + const int32 scaled_input2_val = + MultiplyByQuantizedMultiplierSmallerThanOne( + shifted_input2_val, input2_multiplier, input2_shift); + const int32 raw_sum = scaled_input1_val - scaled_input2_val; + const int32 raw_output = + MultiplyByQuantizedMultiplierSmallerThanOne( + raw_sum, output_multiplier, output_shift) + + output_offset; + const int32 clamped_output = + std::min(output_activation_max, + std::max(output_activation_min, raw_output)); + output_data[Offset(output_dims, c, x, y, b)] = + static_cast(clamped_output); + } + } + } + } +} + +template +inline void BroadcastSub(int left_shift, const uint8* input1_data, + const Dims<4>& input1_dims, int32 input1_offset, + int32 input1_multiplier, int input1_shift, + const uint8* input2_data, const Dims<4>& input2_dims, + int32 input2_offset, int32 input2_multiplier, + int input2_shift, int32 output_offset, + int32 output_multiplier, int output_shift, + int32 output_activation_min, + int32 output_activation_max, uint8* output_data, + const Dims<4>& output_dims) { + static_assert(Ac == FusedActivationFunctionType::kNone || + Ac == FusedActivationFunctionType::kRelu || + Ac == FusedActivationFunctionType::kRelu6 || + Ac == FusedActivationFunctionType::kRelu1, + ""); + TFLITE_DCHECK_LE(output_activation_min, output_activation_max); + if (Ac == FusedActivationFunctionType::kNone) { + TFLITE_DCHECK_EQ(output_activation_min, 0); + TFLITE_DCHECK_EQ(output_activation_max, 255); + } + BroadcastSub(left_shift, input1_data, input1_dims, input1_offset, + input1_multiplier, input1_shift, input2_data, input2_dims, + input2_offset, input2_multiplier, input2_shift, output_offset, + output_multiplier, output_shift, output_activation_min, + output_activation_max, output_data, output_dims); +} + template void Concatenation(int concat_dim, const Scalar* const* input_data, const Dims<4>* const* input_dims, int inputs_count, diff --git a/tensorflow/contrib/lite/kernels/sub.cc b/tensorflow/contrib/lite/kernels/sub.cc index ddaf498d5b..410585a293 100644 --- a/tensorflow/contrib/lite/kernels/sub.cc +++ b/tensorflow/contrib/lite/kernels/sub.cc @@ -26,7 +26,7 @@ namespace ops { namespace builtin { namespace sub { -// This file has three implementation of Div. +// This file has three implementation of Sub. enum KernelType { kReference, kGenericOptimized, // Neon-free @@ -37,7 +37,23 @@ constexpr int kInputTensor1 = 0; constexpr int kInputTensor2 = 1; constexpr int kOutputTensor = 0; +struct OpData { + bool requires_broadcast; +}; + +void* Init(TfLiteContext* context, const char* buffer, size_t length) { + auto* data = new OpData; + data->requires_broadcast = false; + return data; +} + +void Free(TfLiteContext* context, void* buffer) { + delete reinterpret_cast(buffer); +} + TfLiteStatus Prepare(TfLiteContext* context, TfLiteNode* node) { + OpData* data = reinterpret_cast(node->user_data); + TF_LITE_ENSURE_EQ(context, NumInputs(node), 2); TF_LITE_ENSURE_EQ(context, NumOutputs(node), 1); @@ -45,51 +61,122 @@ TfLiteStatus Prepare(TfLiteContext* context, TfLiteNode* node) { TfLiteTensor* input2 = GetInput(context, node, kInputTensor2); TfLiteTensor* output = GetOutput(context, node, kOutputTensor); - TF_LITE_ENSURE_EQ(context, NumDimensions(input1), NumDimensions(input2)); - for (int i = 0; i < NumDimensions(input1); ++i) { - TF_LITE_ENSURE_EQ(context, SizeOfDimension(input1, i), - SizeOfDimension(input2, i)); - } + TF_LITE_ENSURE_EQ(context, input1->type, input2->type); + output->type = input2->type; - TF_LITE_ENSURE_EQ(context, input1->type, output->type); - TF_LITE_ENSURE_EQ(context, input2->type, output->type); + data->requires_broadcast = !HaveSameShapes(input1, input2); + + TfLiteIntArray* output_size = nullptr; + if (data->requires_broadcast) { + TF_LITE_ENSURE_OK(context, CalculateShapeForBroadcast( + context, input1, input2, &output_size)); + } else { + output_size = TfLiteIntArrayCopy(input1->dims); + } - TfLiteIntArray* output_size = TfLiteIntArrayCopy(input1->dims); return context->ResizeTensor(context, output, output_size); } template void EvalSubFloat(TfLiteContext* context, TfLiteNode* node, - TfLiteSubParams* params, TfLiteTensor* input1, - TfLiteTensor* input2, TfLiteTensor* output) { + TfLiteSubParams* params, const OpData* data, + TfLiteTensor* input1, TfLiteTensor* input2, + TfLiteTensor* output) { float output_activation_min, output_activation_max; CalculateActivationRangeFloat(params->activation, &output_activation_min, &output_activation_max); -#define TF_LITE_Sub(type) \ - type::Sub(GetTensorData(input1), GetTensorDims(input1), \ - GetTensorData(input2), GetTensorDims(input2), \ - output_activation_min, output_activation_max, \ - GetTensorData(output), GetTensorDims(output)) +#define TF_LITE_SUB(type, opname) \ + type::opname(GetTensorData(input1), GetTensorDims(input1), \ + GetTensorData(input2), GetTensorDims(input2), \ + output_activation_min, output_activation_max, \ + GetTensorData(output), GetTensorDims(output)) + if (kernel_type == kReference) { + if (data->requires_broadcast) { + TF_LITE_SUB(reference_ops, BroadcastSub); + } else { + TF_LITE_SUB(reference_ops, Sub); + } + } else { + if (data->requires_broadcast) { + TF_LITE_SUB(optimized_ops, BroadcastSub); + } else { + TF_LITE_SUB(optimized_ops, Sub); + } + } +#undef TF_LITE_SUB +} + +template +void EvalSubQuantized(TfLiteContext* context, TfLiteNode* node, + TfLiteSubParams* params, const OpData* data, + TfLiteTensor* input1, TfLiteTensor* input2, + TfLiteTensor* output) { + auto input1_offset = -input1->params.zero_point; + auto input2_offset = -input2->params.zero_point; + auto output_offset = output->params.zero_point; + const int left_shift = 20; + const double twice_max_input_scale = + 2 * std::max(input1->params.scale, input2->params.scale); + const double real_input1_multiplier = + input1->params.scale / twice_max_input_scale; + const double real_input2_multiplier = + input2->params.scale / twice_max_input_scale; + const double real_output_multiplier = + twice_max_input_scale / ((1 << left_shift) * output->params.scale); + + int32 input1_multiplier; + int input1_shift; + QuantizeMultiplierSmallerThanOne(real_input1_multiplier, &input1_multiplier, + &input1_shift); + int32 input2_multiplier; + int input2_shift; + QuantizeMultiplierSmallerThanOne(real_input2_multiplier, &input2_multiplier, + &input2_shift); + int32 output_multiplier; + int output_shift; + QuantizeMultiplierSmallerThanOne(real_output_multiplier, &output_multiplier, + &output_shift); + + int32 output_activation_min, output_activation_max; + CalculateActivationRangeUint8(params->activation, output, + &output_activation_min, &output_activation_max); + +#define TF_LITE_SUB(type, opname) \ + type::opname(left_shift, GetTensorData(input1), \ + GetTensorDims(input1), input1_offset, input1_multiplier, \ + input1_shift, GetTensorData(input2), \ + GetTensorDims(input2), input2_offset, input2_multiplier, \ + input2_shift, output_offset, output_multiplier, output_shift, \ + output_activation_min, output_activation_max, \ + GetTensorData(output), GetTensorDims(output)); + // The quantized version of Sub doesn't support activations, so we + // always use BroadcastSub. if (kernel_type == kReference) { - TF_LITE_Sub(reference_ops); + TF_LITE_SUB(reference_ops, BroadcastSub); } else { - TF_LITE_Sub(optimized_ops); + TF_LITE_SUB(optimized_ops, BroadcastSub); } -#undef TF_LITE_Sub +#undef TF_LITE_SUB } template TfLiteStatus Eval(TfLiteContext* context, TfLiteNode* node) { auto* params = reinterpret_cast(node->builtin_data); + OpData* data = reinterpret_cast(node->user_data); TfLiteTensor* input1 = GetInput(context, node, kInputTensor1); TfLiteTensor* input2 = GetInput(context, node, kInputTensor2); TfLiteTensor* output = GetOutput(context, node, kOutputTensor); if (output->type == kTfLiteFloat32) { - EvalSubFloat(context, node, params, input1, input2, output); + EvalSubFloat(context, node, params, data, input1, input2, + output); + } else if (output->type == kTfLiteUInt8) { + EvalSubQuantized(context, node, params, data, input1, input2, + output); } else { - context->ReportError(context, "Inputs and outputs not all float types."); + context->ReportError(context, + "Inputs and outputs not all float|unit8 types."); return kTfLiteError; } @@ -99,19 +186,19 @@ TfLiteStatus Eval(TfLiteContext* context, TfLiteNode* node) { } // namespace sub TfLiteRegistration* Register_SUB_REF() { - static TfLiteRegistration r = {nullptr, nullptr, sub::Prepare, + static TfLiteRegistration r = {sub::Init, sub::Free, sub::Prepare, sub::Eval}; return &r; } TfLiteRegistration* Register_SUB_GENERIC_OPT() { - static TfLiteRegistration r = {nullptr, nullptr, sub::Prepare, + static TfLiteRegistration r = {sub::Init, sub::Free, sub::Prepare, sub::Eval}; return &r; } TfLiteRegistration* Register_SUB_NEON_OPT() { - static TfLiteRegistration r = {nullptr, nullptr, sub::Prepare, + static TfLiteRegistration r = {sub::Init, sub::Free, sub::Prepare, sub::Eval}; return &r; } diff --git a/tensorflow/contrib/lite/kernels/sub_test.cc b/tensorflow/contrib/lite/kernels/sub_test.cc new file mode 100644 index 0000000000..b2c6d05f62 --- /dev/null +++ b/tensorflow/contrib/lite/kernels/sub_test.cc @@ -0,0 +1,213 @@ +/* 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. +==============================================================================*/ +#include +#include "tensorflow/contrib/lite/interpreter.h" +#include "tensorflow/contrib/lite/kernels/register.h" +#include "tensorflow/contrib/lite/kernels/test_util.h" +#include "tensorflow/contrib/lite/model.h" + +namespace tflite { +namespace { + +using ::testing::ElementsAreArray; + +class BaseSubOpModel : public SingleOpModel { + public: + BaseSubOpModel(const TensorData& input1, const TensorData& input2, + const TensorData& output, + ActivationFunctionType activation_type) { + input1_ = AddInput(input1); + input2_ = AddInput(input2); + output_ = AddOutput(output); + SetBuiltinOp(BuiltinOperator_Sub, BuiltinOptions_SubOptions, + CreateSubOptions(builder_, activation_type).Union()); + BuildInterpreter({GetShape(input1_), GetShape(input2_)}); + } + + int input1() { return input1_; } + int input2() { return input2_; } + + protected: + int input1_; + int input2_; + int output_; +}; + +class FloatSubOpModel : public BaseSubOpModel { + public: + using BaseSubOpModel::BaseSubOpModel; + + std::vector GetOutput() { return ExtractVector(output_); } +}; + +class QuantizedSubOpModel : public BaseSubOpModel { + public: + using BaseSubOpModel::BaseSubOpModel; + + std::vector GetDequantizedOutput() { + return Dequantize(ExtractVector(output_), + GetScale(output_), GetZeroPoint(output_)); + } +}; + +// for quantized Sub, the error shouldn't exceed 2*step +float GetTolerance(int min, int max) { + float kQuantizedStep = (max - min) / 255.0; + float kQuantizedTolerance = 2.0 * kQuantizedStep; + return kQuantizedTolerance; +} + +TEST(FloatSubOpModel, NoActivation) { + FloatSubOpModel m({TensorType_FLOAT32, {1, 2, 2, 1}}, + {TensorType_FLOAT32, {1, 2, 2, 1}}, + {TensorType_FLOAT32, {}}, ActivationFunctionType_NONE); + m.PopulateTensor(m.input1(), {-2.0, 0.2, 1.7, 0.5}); + m.PopulateTensor(m.input2(), {0.1, 0.2, 0.3, 0.8}); + m.Invoke(); + EXPECT_THAT(m.GetOutput(), ElementsAreArray({-2.1, 0.0, 1.4, -0.3})); +} + +TEST(FloatSubOpModel, ActivationRELU_N1_TO_1) { + FloatSubOpModel m( + {TensorType_FLOAT32, {1, 2, 2, 1}}, {TensorType_FLOAT32, {1, 2, 2, 1}}, + {TensorType_FLOAT32, {}}, ActivationFunctionType_RELU_N1_TO_1); + m.PopulateTensor(m.input1(), {-2.0, 0.2, 1.7, 0.5}); + m.PopulateTensor(m.input2(), {0.1, 0.2, 0.3, 0.8}); + m.Invoke(); + EXPECT_THAT(m.GetOutput(), ElementsAreArray({-1.0, 0.0, 1.0, -0.3})); +} + +TEST(FloatSubOpModel, VariousInputShapes) { + std::vector> test_shapes = { + {6}, {2, 3}, {2, 1, 3}, {1, 3, 1, 2}}; + for (int i = 0; i < test_shapes.size(); ++i) { + FloatSubOpModel m({TensorType_FLOAT32, test_shapes[i]}, + {TensorType_FLOAT32, test_shapes[i]}, + {TensorType_FLOAT32, {}}, ActivationFunctionType_NONE); + m.PopulateTensor(m.input1(), {-2.0, 0.2, 1.7, 0.5, -1.1, 2.0}); + m.PopulateTensor(m.input2(), {0.1, 0.2, 0.3, 0.8, -1.1, 0.1}); + m.Invoke(); + EXPECT_THAT(m.GetOutput(), + ElementsAreArray({-2.1, 0.0, 1.4, -0.3, 0.0, 1.9})) + << "With shape number " << i; + } +} + +TEST(FloatSubOpModel, WithBroadcast) { + std::vector> test_shapes = { + {6}, {2, 3}, {2, 1, 3}, {1, 3, 1, 2}}; + for (int i = 0; i < test_shapes.size(); ++i) { + FloatSubOpModel m({TensorType_FLOAT32, test_shapes[i]}, + {TensorType_FLOAT32, {}}, // always a scalar + {TensorType_FLOAT32, {}}, ActivationFunctionType_NONE); + m.PopulateTensor(m.input1(), {-2.0, 0.2, 1.7, 0.5, -1.1, 2.0}); + m.PopulateTensor(m.input2(), {0.5}); + m.Invoke(); + EXPECT_THAT( + m.GetOutput(), + ElementsAreArray(ArrayFloatNear({-2.5, -0.3, 1.2, 0.0, -1.6, 1.5}))) + << "With shape number " << i; + } +} + +TEST(QuantizedSubOpModel, QuantizedTestsNoActivation) { + float kQuantizedTolerance = GetTolerance(-1.0, 1.0); + std::vector> inputs1 = { + {0.1, 0.2, 0.3, 0.4}, {-0.2, 0.2, 0.4, 0.7}, {-0.01, 0.2, 0.7, 0.3}}; + std::vector> inputs2 = { + {0.6, 0.4, 0.3, 0.1}, {0.6, 0.4, 0.5, -0.2}, {0.6, 0.4, -0.18, 0.5}}; + std::vector> results = { + {-0.5, -0.2, 0.0, 0.3}, {-0.8, -0.2, -0.1, 0.9}, {-0.61, -0.2, 0.88, -0.2}}; + for (int i = 0; i < inputs1.size(); ++i) { + QuantizedSubOpModel m({TensorType_UINT8, {1, 2, 2, 1}, -1.0, 1.0}, + {TensorType_UINT8, {1, 2, 2, 1}, -1.0, 1.0}, + {TensorType_UINT8, {}, -1.0, 1.0}, + ActivationFunctionType_NONE); + m.QuantizeAndPopulate(m.input1(), inputs1[i]); + m.QuantizeAndPopulate(m.input2(), inputs2[i]); + m.Invoke(); + EXPECT_THAT(m.GetDequantizedOutput(), ElementsAreArray(ArrayFloatNear( + results[i], kQuantizedTolerance))) + << "With test number " << i; + } +} + +TEST(QuantizedSubOpModel, QuantizedTestsActivationRELU_N1_TO_1) { + float kQuantizedTolerance = GetTolerance(-1.0, 1.0); + std::vector> inputs1 = {{-0.8, 0.2, 0.9, 0.7}, + {-0.8, 0.2, 0.7, 0.5}}; + std::vector> inputs2 = {{0.6, 0.4, 0.9, -0.8}, + {0.6, 0.4, -0.8, 0.3}}; + std::vector> results = {{-1.0, -0.2, 0.0, 1.0}, + {-1.0, -0.2, 1.0, 0.2}}; + for (int i = 0; i < inputs1.size(); ++i) { + QuantizedSubOpModel m({TensorType_UINT8, {1, 2, 2, 1}, -1.0, 1.0}, + {TensorType_UINT8, {1, 2, 2, 1}, -1.0, 1.0}, + {TensorType_UINT8, {}, -1.0, 1.0}, + ActivationFunctionType_RELU_N1_TO_1); + m.QuantizeAndPopulate(m.input1(), inputs1[i]); + m.QuantizeAndPopulate(m.input2(), inputs2[i]); + m.Invoke(); + EXPECT_THAT(m.GetDequantizedOutput(), ElementsAreArray(ArrayFloatNear( + results[i], kQuantizedTolerance))) + << "With test number " << i; + } +} + +TEST(QuantizedSubOpModel, QuantizedVariousInputShapes) { + float kQuantizedTolerance = GetTolerance(-3.0, 3.0); + std::vector> test_shapes = { + {6}, {2, 3}, {2, 1, 3}, {1, 3, 1, 2}}; + for (int i = 0; i < test_shapes.size(); ++i) { + QuantizedSubOpModel m({TensorType_UINT8, test_shapes[i], -3.0, 3.0}, + {TensorType_UINT8, test_shapes[i], -3.0, 3.0}, + {TensorType_UINT8, {}, -3.0, 3.0}, + ActivationFunctionType_NONE); + m.QuantizeAndPopulate(m.input1(), {-2.0, 0.2, 0.7, 0.8, 1.1, 2.0}); + m.QuantizeAndPopulate(m.input2(), {0.1, 0.3, 0.3, 0.5, 1.1, 0.1}); + m.Invoke(); + EXPECT_THAT(m.GetDequantizedOutput(), + ElementsAreArray(ArrayFloatNear({-2.1, -0.1, 0.4, 0.3, 0.0, 1.9}, + kQuantizedTolerance))) + << "With shape number " << i; + } +} + +TEST(QuantizedSubOpModel, QuantizedWithBroadcast) { + float kQuantizedTolerance = GetTolerance(-3.0, 3.0); + std::vector> test_shapes = { + {6}, {2, 3}, {2, 1, 3}, {1, 3, 1, 2}}; + for (int i = 0; i < test_shapes.size(); ++i) { + QuantizedSubOpModel m({TensorType_UINT8, test_shapes[i], -3.0, 3.0}, + {TensorType_UINT8, {}, -3.0, 3.0}, + {TensorType_UINT8, {}, -3.0, 3.0}, + ActivationFunctionType_NONE); + m.QuantizeAndPopulate(m.input1(), {-2.0, 0.2, 0.7, 0.8, 1.1, 2.0}); + m.QuantizeAndPopulate(m.input2(), {0.7}); + m.Invoke(); + EXPECT_THAT(m.GetDequantizedOutput(), + ElementsAreArray(ArrayFloatNear({-2.7, -0.5, 0.0, 0.1, 0.4, 1.3}, + kQuantizedTolerance))) + << "With shape number " << i; + } +} + +} // namespace +} // namespace tflite +int main(int argc, char** argv) { + ::tflite::LogToStderr(); + ::testing::InitGoogleTest(&argc, argv); + return RUN_ALL_TESTS(); +} diff --git a/tensorflow/contrib/lite/testing/generated_examples_zip_test.cc b/tensorflow/contrib/lite/testing/generated_examples_zip_test.cc index 49766cedac..1e177d5f6e 100644 --- a/tensorflow/contrib/lite/testing/generated_examples_zip_test.cc +++ b/tensorflow/contrib/lite/testing/generated_examples_zip_test.cc @@ -47,9 +47,6 @@ tensorflow::Env* env = tensorflow::Env::Default(); // Key is a substring of the test name and value is a bug number. // TODO(ahentz): make sure we clean this list up frequently. std::map kBrokenTests = { - // Sub and Div don't support broadcasting. - {R"(^\/diva.*input_shape_1=\[1,3,4,3\],input_shape_2=\[3\])", "68500195"}, - {R"(^\/suba.*input_shape_1=\[1,3,4,3\],input_shape_2=\[3\])", "68500195"}, // Add only supports float32. (and "constant" tests use Add) {R"(^\/adda.*int32)", "68808744"}, @@ -235,22 +232,23 @@ TEST_P(OpsTest, RunStuff) { INSTANTIATE_TESTS(add) INSTANTIATE_TESTS(avg_pool) -INSTANTIATE_TESTS(space_to_batch_nd) INSTANTIATE_TESTS(batch_to_space_nd) INSTANTIATE_TESTS(concat) INSTANTIATE_TESTS(constant) INSTANTIATE_TESTS(control_dep) INSTANTIATE_TESTS(conv) INSTANTIATE_TESTS(depthwiseconv) +INSTANTIATE_TESTS(div) INSTANTIATE_TESTS(exp) INSTANTIATE_TESTS(fully_connected) INSTANTIATE_TESTS(fused_batch_norm) INSTANTIATE_TESTS(gather) INSTANTIATE_TESTS(global_batch_norm) -INSTANTIATE_TESTS(l2norm) INSTANTIATE_TESTS(l2_pool) +INSTANTIATE_TESTS(l2norm) INSTANTIATE_TESTS(local_response_norm) INSTANTIATE_TESTS(max_pool) +INSTANTIATE_TESTS(mean) INSTANTIATE_TESTS(mul) INSTANTIATE_TESTS(pad) INSTANTIATE_TESTS(relu) @@ -260,14 +258,13 @@ INSTANTIATE_TESTS(reshape) INSTANTIATE_TESTS(resize_bilinear) INSTANTIATE_TESTS(sigmoid) INSTANTIATE_TESTS(softmax) +INSTANTIATE_TESTS(space_to_batch_nd) INSTANTIATE_TESTS(space_to_depth) -INSTANTIATE_TESTS(sub) INSTANTIATE_TESTS(split) -INSTANTIATE_TESTS(div) -INSTANTIATE_TESTS(transpose) -INSTANTIATE_TESTS(mean) INSTANTIATE_TESTS(squeeze) INSTANTIATE_TESTS(strided_slice) +INSTANTIATE_TESTS(sub) +INSTANTIATE_TESTS(transpose) } // namespace testing } // namespace tflite -- GitLab From 779d457008ab7ea2c11f4d73370099a1e56c0652 Mon Sep 17 00:00:00 2001 From: ManHyuk Date: Sun, 25 Feb 2018 21:39:52 +0900 Subject: [PATCH 004/528] fix typo --- .../python/kernel_tests/linalg/linear_operator_diag_test.py | 2 +- tensorflow/python/ops/linalg/linear_operator_diag.py | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/tensorflow/python/kernel_tests/linalg/linear_operator_diag_test.py b/tensorflow/python/kernel_tests/linalg/linear_operator_diag_test.py index 343d158498..8cb9f9e621 100644 --- a/tensorflow/python/kernel_tests/linalg/linear_operator_diag_test.py +++ b/tensorflow/python/kernel_tests/linalg/linear_operator_diag_test.py @@ -129,7 +129,7 @@ class LinearOperatorDiagTest( with self.test_session() as sess: x = random_ops.random_normal(shape=(2, 2, 3, 4)) - # This LinearOperatorDiag will be brodacast to (2, 2, 3, 3) during solve + # This LinearOperatorDiag will be broadcast to (2, 2, 3, 3) during solve # and matmul with 'x' as the argument. diag = random_ops.random_uniform(shape=(2, 1, 3)) operator = linalg.LinearOperatorDiag(diag, is_self_adjoint=True) diff --git a/tensorflow/python/ops/linalg/linear_operator_diag.py b/tensorflow/python/ops/linalg/linear_operator_diag.py index b3ec3d5b7c..e180e83026 100644 --- a/tensorflow/python/ops/linalg/linear_operator_diag.py +++ b/tensorflow/python/ops/linalg/linear_operator_diag.py @@ -67,7 +67,7 @@ class LinearOperatorDiag(linear_operator.LinearOperator): operator = LinearOperatorDiag(diag) # Create a shape [2, 1, 4, 2] vector. Note that this shape is compatible - # since the batch dimensions, [2, 1], are brodcast to + # since the batch dimensions, [2, 1], are broadcast to # operator.batch_shape = [2, 3]. y = tf.random_normal(shape=[2, 1, 4, 2]) x = operator.solve(y) -- GitLab From b569035378ef4a8595c64e5f398d74244cac376e Mon Sep 17 00:00:00 2001 From: ManHyuk Date: Sun, 25 Feb 2018 21:44:12 +0900 Subject: [PATCH 005/528] fix typo --- tensorflow/contrib/slim/python/slim/data/parallel_reader.py | 2 +- tensorflow/python/ops/distributions/special_math.py | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/tensorflow/contrib/slim/python/slim/data/parallel_reader.py b/tensorflow/contrib/slim/python/slim/data/parallel_reader.py index ad5e985487..b3343aef47 100644 --- a/tensorflow/contrib/slim/python/slim/data/parallel_reader.py +++ b/tensorflow/contrib/slim/python/slim/data/parallel_reader.py @@ -221,7 +221,7 @@ def parallel_read(data_sources, the data will be cycled through indefinitely. num_readers: a integer, number of Readers to create. reader_kwargs: an optional dict, of kwargs for the reader. - shuffle: boolean, wether should shuffle the files and the records by using + shuffle: boolean, whether should shuffle the files and the records by using RandomShuffleQueue as common_queue. dtypes: A list of types. The length of dtypes must equal the number of elements in each record. If it is None it will default to diff --git a/tensorflow/python/ops/distributions/special_math.py b/tensorflow/python/ops/distributions/special_math.py index bed4cbb2c1..1d605c5dfc 100644 --- a/tensorflow/python/ops/distributions/special_math.py +++ b/tensorflow/python/ops/distributions/special_math.py @@ -213,7 +213,7 @@ def _ndtri(p): # Compute x for p <= exp(-2): x = z - log(z)/z - (1/z) P(1/z) / Q(1/z), # where z = sqrt(-2. * log(p)), and P/Q are chosen between two different - # arrays based on wether p < exp(-32). + # arrays based on whether p < exp(-32). z = math_ops.sqrt(-2. * math_ops.log(sanitized_mcp)) first_term = z - math_ops.log(z) / z second_term_small_p = (_create_polynomial(1. / z, p2) -- GitLab From f1f70ef5c268d6ce41bdab4867ed0f2e19d6f924 Mon Sep 17 00:00:00 2001 From: Hovhannes Harutyunyan Date: Mon, 26 Feb 2018 10:52:11 +0400 Subject: [PATCH 006/528] Remove code that was written for compatibility with old checked-in code. Update code to have 80 characters per line. --- tensorflow/contrib/lite/kernels/div_test.cc | 3 +- .../internal/optimized/optimized_ops.h | 41 ------------------- .../internal/reference/reference_ops.h | 41 ------------------- tensorflow/contrib/lite/kernels/sub_test.cc | 18 +++++--- 4 files changed, 15 insertions(+), 88 deletions(-) diff --git a/tensorflow/contrib/lite/kernels/div_test.cc b/tensorflow/contrib/lite/kernels/div_test.cc index 78918a0d79..e67e0ec034 100644 --- a/tensorflow/contrib/lite/kernels/div_test.cc +++ b/tensorflow/contrib/lite/kernels/div_test.cc @@ -154,7 +154,8 @@ TEST(QuantizedDivOpTest, WithBroadcast) { {TensorType_UINT8, {}, -3.0, 3.0}, // always a scalar {TensorType_UINT8, {}, -3.0, 3.0}, ActivationFunctionType_NONE); - m.QuantizeAndPopulate(m.input1(), {-0.2, 0.2, 0.07, 0.08, 0.11, -0.123}); + m.QuantizeAndPopulate(m.input1(), {-0.2, 0.2, 0.07, + 0.08, 0.11, -0.123}); m.QuantizeAndPopulate(m.input2(), {0.1}); m.Invoke(); EXPECT_THAT(m.GetDequantizedOutput(), diff --git a/tensorflow/contrib/lite/kernels/internal/optimized/optimized_ops.h b/tensorflow/contrib/lite/kernels/internal/optimized/optimized_ops.h index d12a3eca1d..b19f46beaa 100644 --- a/tensorflow/contrib/lite/kernels/internal/optimized/optimized_ops.h +++ b/tensorflow/contrib/lite/kernels/internal/optimized/optimized_ops.h @@ -1973,19 +1973,6 @@ void BroadcastDiv(const T* input1_data, const Dims<4>& input1_dims, } } -// legacy, for compatibility with old checked-in code -template -void BroadcastDiv(const T* input1_data, const Dims<4>& input1_dims, - const T* input2_data, const Dims<4>& input2_dims, - T* output_data, const Dims<4>& output_dims) { - T output_activation_min, output_activation_max; - GetActivationMinMax(Ac, &output_activation_min, &output_activation_max); - - BroadcastDiv(input1_data, input1_dims, input2_data, input2_dims, - output_activation_min, output_activation_max, output_data, - output_dims); -} - inline void BroadcastDiv(const uint8* input1_data, const Dims<4>& input1_dims, int32 input1_offset, const uint8* input2_data, const Dims<4>& input2_dims, int32 input2_offset, @@ -2033,21 +2020,6 @@ inline void BroadcastDiv(const uint8* input1_data, const Dims<4>& input1_dims, } } -// legacy, for compatibility with old checked-in code -template -inline void BroadcastDiv(const uint8* input1_data, const Dims<4>& input1_dims, - int32 input1_offset, const uint8* input2_data, - const Dims<4>& input2_dims, int32 input2_offset, - int32 output_offset, int32 output_multiplier, - int output_shift, int32 output_activation_min, - int32 output_activation_max, uint8* output_data, - const Dims<4>& output_dims) { - BroadcastDiv(input1_data, input1_dims, input1_offset, input2_data, - input2_dims, input2_offset, output_offset, output_multiplier, - output_shift, output_activation_min, output_activation_max, - output_data, output_dims); -} - // TODO(aselle): This is not actually optimized yet. inline void Sub(const float* input1_data, const Dims<4>& input1_dims, const float* input2_data, const Dims<4>& input2_dims, @@ -2121,19 +2093,6 @@ void BroadcastSub(const T* input1_data, const Dims<4>& input1_dims, } } -// legacy, for compatibility with old checked-in code -template -void BroadcastSub(const T* input1_data, const Dims<4>& input1_dims, - const T* input2_data, const Dims<4>& input2_dims, - T* output_data, const Dims<4>& output_dims) { - T output_activation_min, output_activation_max; - GetActivationMinMax(Ac, &output_activation_min, &output_activation_max); - - BroadcastSub(input1_data, input1_dims, input2_data, input2_dims, - output_activation_min, output_activation_max, output_data, - output_dims); -} - inline void BroadcastSub(int left_shift, const uint8* input1_data, const Dims<4>& input1_dims, int32 input1_offset, int32 input1_multiplier, int input1_shift, diff --git a/tensorflow/contrib/lite/kernels/internal/reference/reference_ops.h b/tensorflow/contrib/lite/kernels/internal/reference/reference_ops.h index c7b7687622..847075e207 100644 --- a/tensorflow/contrib/lite/kernels/internal/reference/reference_ops.h +++ b/tensorflow/contrib/lite/kernels/internal/reference/reference_ops.h @@ -1249,19 +1249,6 @@ void BroadcastDiv(const T* input1_data, const Dims<4>& input1_dims, } } -// legacy, for compatibility with old checked-in code -template -void BroadcastDiv(const T* input1_data, const Dims<4>& input1_dims, - const T* input2_data, const Dims<4>& input2_dims, - T* output_data, const Dims<4>& output_dims) { - T output_activation_min, output_activation_max; - GetActivationMinMax(Ac, &output_activation_min, &output_activation_max); - - BroadcastDiv(input1_data, input1_dims, input2_data, input2_dims, - output_activation_min, output_activation_max, output_data, - output_dims); -} - inline void BroadcastDiv(const uint8* input1_data, const Dims<4>& input1_dims, int32 input1_offset, const uint8* input2_data, const Dims<4>& input2_dims, int32 input2_offset, @@ -1309,21 +1296,6 @@ inline void BroadcastDiv(const uint8* input1_data, const Dims<4>& input1_dims, } } -// legacy, for compatibility with old checked-in code -template -inline void BroadcastDiv(const uint8* input1_data, const Dims<4>& input1_dims, - int32 input1_offset, const uint8* input2_data, - const Dims<4>& input2_dims, int32 input2_offset, - int32 output_offset, int32 output_multiplier, - int output_shift, int32 output_activation_min, - int32 output_activation_max, uint8* output_data, - const Dims<4>& output_dims) { - BroadcastDiv(input1_data, input1_dims, input1_offset, input2_data, - input2_dims, input2_offset, output_offset, output_multiplier, - output_shift, output_activation_min, output_activation_max, - output_data, output_dims); -} - inline void Sub(const float* input1_data, const Dims<4>& input1_dims, const float* input2_data, const Dims<4>& input2_dims, float output_activation_min, float output_activation_max, @@ -1392,19 +1364,6 @@ void BroadcastSub(const T* input1_data, const Dims<4>& input1_dims, } } -// legacy, for compatibility with old checked-in code -template -void BroadcastSub(const T* input1_data, const Dims<4>& input1_dims, - const T* input2_data, const Dims<4>& input2_dims, - T* output_data, const Dims<4>& output_dims) { - T output_activation_min, output_activation_max; - GetActivationMinMax(Ac, &output_activation_min, &output_activation_max); - - BroadcastSub(input1_data, input1_dims, input2_data, input2_dims, - output_activation_min, output_activation_max, output_data, - output_dims); -} - inline void BroadcastSub(int left_shift, const uint8* input1_data, const Dims<4>& input1_dims, int32 input1_offset, int32 input1_multiplier, int input1_shift, diff --git a/tensorflow/contrib/lite/kernels/sub_test.cc b/tensorflow/contrib/lite/kernels/sub_test.cc index b2c6d05f62..1fd0ee2a0e 100644 --- a/tensorflow/contrib/lite/kernels/sub_test.cc +++ b/tensorflow/contrib/lite/kernels/sub_test.cc @@ -125,11 +125,17 @@ TEST(FloatSubOpModel, WithBroadcast) { TEST(QuantizedSubOpModel, QuantizedTestsNoActivation) { float kQuantizedTolerance = GetTolerance(-1.0, 1.0); std::vector> inputs1 = { - {0.1, 0.2, 0.3, 0.4}, {-0.2, 0.2, 0.4, 0.7}, {-0.01, 0.2, 0.7, 0.3}}; + {0.1, 0.2, 0.3, 0.4}, + {-0.2, 0.2, 0.4, 0.7}, + {-0.01, 0.2, 0.7, 0.3}}; std::vector> inputs2 = { - {0.6, 0.4, 0.3, 0.1}, {0.6, 0.4, 0.5, -0.2}, {0.6, 0.4, -0.18, 0.5}}; + {0.6, 0.4, 0.3, 0.1}, + {0.6, 0.4, 0.5, -0.2}, + {0.6, 0.4, -0.18, 0.5}}; std::vector> results = { - {-0.5, -0.2, 0.0, 0.3}, {-0.8, -0.2, -0.1, 0.9}, {-0.61, -0.2, 0.88, -0.2}}; + {-0.5, -0.2, 0.0, 0.3}, + {-0.8, -0.2, -0.1, 0.9}, + {-0.61, -0.2, 0.88, -0.2}}; for (int i = 0; i < inputs1.size(); ++i) { QuantizedSubOpModel m({TensorType_UINT8, {1, 2, 2, 1}, -1.0, 1.0}, {TensorType_UINT8, {1, 2, 2, 1}, -1.0, 1.0}, @@ -179,7 +185,8 @@ TEST(QuantizedSubOpModel, QuantizedVariousInputShapes) { m.QuantizeAndPopulate(m.input2(), {0.1, 0.3, 0.3, 0.5, 1.1, 0.1}); m.Invoke(); EXPECT_THAT(m.GetDequantizedOutput(), - ElementsAreArray(ArrayFloatNear({-2.1, -0.1, 0.4, 0.3, 0.0, 1.9}, + ElementsAreArray(ArrayFloatNear({-2.1, -0.1, 0.4, + 0.3, 0.0, 1.9}, kQuantizedTolerance))) << "With shape number " << i; } @@ -198,7 +205,8 @@ TEST(QuantizedSubOpModel, QuantizedWithBroadcast) { m.QuantizeAndPopulate(m.input2(), {0.7}); m.Invoke(); EXPECT_THAT(m.GetDequantizedOutput(), - ElementsAreArray(ArrayFloatNear({-2.7, -0.5, 0.0, 0.1, 0.4, 1.3}, + ElementsAreArray(ArrayFloatNear({-2.7, -0.5, 0.0, + 0.1, 0.4, 1.3}, kQuantizedTolerance))) << "With shape number " << i; } -- GitLab From 0489bf25930ea0dc4b7d8ffc792b0390bfbc06bc Mon Sep 17 00:00:00 2001 From: Jingwen Date: Tue, 27 Feb 2018 18:30:09 -0500 Subject: [PATCH 007/528] Include cstring in logging.cc for use of strrchr() --- tensorflow/core/platform/default/logging.cc | 1 + 1 file changed, 1 insertion(+) diff --git a/tensorflow/core/platform/default/logging.cc b/tensorflow/core/platform/default/logging.cc index 2b874da198..c6e5777c26 100644 --- a/tensorflow/core/platform/default/logging.cc +++ b/tensorflow/core/platform/default/logging.cc @@ -21,6 +21,7 @@ limitations under the License. #include #include #include +#include #endif #include -- GitLab From 4b7db48218799ef172c7c9794d9d98e56d838ecb Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Mon, 5 Mar 2018 17:41:00 +0000 Subject: [PATCH 008/528] Update the documentation of `softmax_cross_entropy` This fix updates the documentation of `softmax_cross_entropy`, and removed the shape restrictions of `onehot_labels` and `logits`. They only needs to be of the same shape, not necessary `[batch_size, num_classes]`. This fix fixes 16263. Signed-off-by: Yong Tang --- tensorflow/python/ops/losses/losses_impl.py | 13 +++++++++---- 1 file changed, 9 insertions(+), 4 deletions(-) diff --git a/tensorflow/python/ops/losses/losses_impl.py b/tensorflow/python/ops/losses/losses_impl.py index 7386976e93..04c13cb6c6 100644 --- a/tensorflow/python/ops/losses/losses_impl.py +++ b/tensorflow/python/ops/losses/losses_impl.py @@ -710,11 +710,16 @@ def softmax_cross_entropy( new_onehot_labels = onehot_labels * (1 - label_smoothing) + label_smoothing / num_classes + Note that `onehot_labels` and `logits` must have the same shape, + e.g. `[batch_size, num_classes]`. The shape of `weights` must be + broadcastable to loss, whose shape is decided by the shape of `logits`. + In case the shape of `logits` is `[batch_size, num_classes]`, loss is + a `Tensor` of shape `[batch_size]`. + Args: - onehot_labels: `[batch_size, num_classes]` target one-hot-encoded labels. - logits: `[batch_size, num_classes]` logits outputs of the network . - weights: Optional `Tensor` whose rank is either 0, or rank 1 and is - broadcastable to the loss which is a `Tensor` of shape `[batch_size]`. + onehot_labels: One-hot-encoded labels. + logits: Logits outputs of the network. + weights: Optional `Tensor` that is broadcastable to loss. label_smoothing: If greater than 0 then smooth the labels. scope: the scope for the operations performed in computing the loss. loss_collection: collection to which the loss will be added. -- GitLab From 82e34cd19f554509113d438ca98ad76e42fdf4e9 Mon Sep 17 00:00:00 2001 From: Hovhannes Harutyunyan Date: Wed, 7 Mar 2018 09:14:53 +0400 Subject: [PATCH 009/528] Remove quantized versiaon of Div till fixing it. --- .../internal/optimized/optimized_ops.h | 47 ------------------- 1 file changed, 47 deletions(-) diff --git a/tensorflow/contrib/lite/kernels/internal/optimized/optimized_ops.h b/tensorflow/contrib/lite/kernels/internal/optimized/optimized_ops.h index b19f46beaa..9c181fddad 100644 --- a/tensorflow/contrib/lite/kernels/internal/optimized/optimized_ops.h +++ b/tensorflow/contrib/lite/kernels/internal/optimized/optimized_ops.h @@ -1973,53 +1973,6 @@ void BroadcastDiv(const T* input1_data, const Dims<4>& input1_dims, } } -inline void BroadcastDiv(const uint8* input1_data, const Dims<4>& input1_dims, - int32 input1_offset, const uint8* input2_data, - const Dims<4>& input2_dims, int32 input2_offset, - int32 output_offset, int32 output_multiplier, - int output_shift, int32 output_activation_min, - int32 output_activation_max, uint8* output_data, - const Dims<4>& output_dims) { - gemmlowp::ScopedProfilingLabel label("BroadcastDiv/8bit"); - - NdArrayDesc<4> desc1; - NdArrayDesc<4> desc2; - NdArrayDescsForElementwiseBroadcast(input1_dims, input2_dims, &desc1, &desc2); - - // In Tensorflow, the dimensions are canonically named (batch_number, row, - // col, channel), with extents (batches, height, width, depth), with the - // trailing dimension changing most rapidly (channels has the smallest stride, - // typically 1 element). - // - // In generated C code, we store arrays with the dimensions reversed. The - // first dimension has smallest stride. - // - // We name our variables by their Tensorflow convention, but generate C code - // nesting loops such that the innermost loop has the smallest stride for the - // best cache behavior. - for (int b = 0; b < ArraySize(output_dims, 3); ++b) { - for (int y = 0; y < ArraySize(output_dims, 2); ++y) { - for (int x = 0; x < ArraySize(output_dims, 1); ++x) { - for (int c = 0; c < ArraySize(output_dims, 0); ++c) { - const int32 input1_val = - input1_offset + input1_data[SubscriptToIndex(desc1, c, x, y, b)]; - const int32 input2_val = - input2_offset + input2_data[SubscriptToIndex(desc2, c, x, y, b)]; - const int32 unclamped_result = - output_offset + - MultiplyByQuantizedMultiplierSmallerThanOne( - input1_val / input2_val, output_multiplier, output_shift); - const int32 clamped_output = - std::min(output_activation_max, - std::max(output_activation_min, unclamped_result)); - output_data[Offset(output_dims, c, x, y, b)] = - static_cast(clamped_output); - } - } - } - } -} - // TODO(aselle): This is not actually optimized yet. inline void Sub(const float* input1_data, const Dims<4>& input1_dims, const float* input2_data, const Dims<4>& input2_dims, -- GitLab From b21ceeb518ca9462a247d8be05870f12bebad201 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Thu, 15 Mar 2018 23:13:25 -0700 Subject: [PATCH 010/528] Enhancement with deprecated_argument_lookup for argmax This fix makes some enhancement for argmax, using deprecated_argument_lookup instread of customerized logic. Signed-off-by: Yong Tang --- tensorflow/python/ops/math_ops.py | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/tensorflow/python/ops/math_ops.py b/tensorflow/python/ops/math_ops.py index e18d0e9501..9a88b71398 100644 --- a/tensorflow/python/ops/math_ops.py +++ b/tensorflow/python/ops/math_ops.py @@ -208,11 +208,9 @@ def argmax(input, name=None, dimension=None, output_type=dtypes.int64): - if dimension is not None: - if axis is not None: - raise ValueError("Cannot specify both 'axis' and 'dimension'") - axis = dimension - elif axis is None: + axis = deprecation.deprecated_argument_lookup( + "axis", axis, "dimension", dimension) + if axis is None: axis = 0 return gen_math_ops.arg_max(input, axis, name=name, output_type=output_type) -- GitLab From 82571ca199869f60fe2036d15d0071031d997b47 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Thu, 15 Mar 2018 23:15:37 -0700 Subject: [PATCH 011/528] Enhancement with deprecated_argument_lookup for argmin This fix makes some enhancement for argmin, using deprecated_argument_lookup instread of customerized logic. Signed-off-by: Yong Tang --- tensorflow/python/ops/math_ops.py | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/tensorflow/python/ops/math_ops.py b/tensorflow/python/ops/math_ops.py index 9a88b71398..a2892d206d 100644 --- a/tensorflow/python/ops/math_ops.py +++ b/tensorflow/python/ops/math_ops.py @@ -226,11 +226,9 @@ def argmin(input, name=None, dimension=None, output_type=dtypes.int64): - if dimension is not None: - if axis is not None: - raise ValueError("Cannot specify both 'axis' and 'dimension'") - axis = dimension - elif axis is None: + axis = deprecation.deprecated_argument_lookup( + "axis", axis, "dimension", dimension) + if axis is None: axis = 0 return gen_math_ops.arg_min(input, axis, name=name, output_type=output_type) -- GitLab From 52fef7f6b8b41d4fffa92bddcb78d96eb6333051 Mon Sep 17 00:00:00 2001 From: ManHyuk Date: Fri, 16 Mar 2018 16:03:26 +0900 Subject: [PATCH 012/528] fix typo --- tensorflow/tools/graph_transforms/fold_old_batch_norms_test.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/tools/graph_transforms/fold_old_batch_norms_test.cc b/tensorflow/tools/graph_transforms/fold_old_batch_norms_test.cc index 272410c693..7651a03fe5 100644 --- a/tensorflow/tools/graph_transforms/fold_old_batch_norms_test.cc +++ b/tensorflow/tools/graph_transforms/fold_old_batch_norms_test.cc @@ -398,7 +398,7 @@ TEST_F(FoldOldBatchNormsTest, TestFoldFusedBatchNorms) { } TEST_F(FoldOldBatchNormsTest, TestFoldFusedBatchNormsWithConcat) { - // Test axis is not 3, so all weigths and offsets are fused to each of inputs + // Test axis is not 3, so all weights and offsets are fused to each of inputs // of conv2d. TestFoldFusedBatchNormsWithConcat(/*split=*/true); // Test axis = 3, BatchNorm weights and offsets will be split before fused -- GitLab From 083cf6b91a380641933457a4301f9b1efa13af92 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Wed, 25 Oct 2017 17:03:15 +0000 Subject: [PATCH 013/528] Add customerized kernel implementation for clip_by_value This fix tries to address the issue raised in 7225 where `tf.clip_by_value` does not have a custom kernel and reused `tf.maximum` and `tf.mimimum`. In case scalar values are passed to `tf.clip_by_value`, unnecessary memory might incur. This fix adds the customerized kernel implementation for `tf.clip_by_value`. This fix fixes 7225. Signed-off-by: Yong Tang --- tensorflow/core/kernels/cwise_op_clip.cc | 150 +++++++++++++++++++++++ tensorflow/core/ops/math_ops.cc | 23 ++++ 2 files changed, 173 insertions(+) create mode 100644 tensorflow/core/kernels/cwise_op_clip.cc diff --git a/tensorflow/core/kernels/cwise_op_clip.cc b/tensorflow/core/kernels/cwise_op_clip.cc new file mode 100644 index 0000000000..6ce062b08f --- /dev/null +++ b/tensorflow/core/kernels/cwise_op_clip.cc @@ -0,0 +1,150 @@ +/* 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/core/kernels/cwise_ops_common.h" + +//#include "third_party/eigen3/Eigen/Core/CwiseTernaryOp.h" + +namespace tensorflow { + +// Unary functor for clip +template +struct UnaryClipOp { + UnaryClipOp(const T& value_min, const T& value_max) + : value_min_(value_min), value_max_(value_max) {} + const T operator()(const T& value) const { + return std::max(std::min(value, value_max_), value_min_); + } + T value_min_; + T value_max_; +}; + +// Binary functor for clip +template +struct BinaryClipMinOp { + BinaryClipMinOp(const T& value_min) : value_min_(value_min) {} + const T operator()(const T& value, const T& value_max) const { + return std::max(std::min(value, value_max), value_min_); + } + T value_min_; +}; + +// Binary functor for clip +template +struct BinaryClipMaxOp { + BinaryClipMaxOp(const T& value_max) : value_max_(value_max) {} + const T operator()(const T& value, const T& value_min) const { + return std::max(std::min(value, value_max_), value_min); + } + T value_max_; +}; + +// Basic coefficient-wise tenary operations. +// This is the case for example of the clip_by_value. +// Device: E.g., CPUDevice, GPUDevice. +// Functor: defined above. E.g., functor::clip. +template +class TenaryOp : public OpKernel { + public: + explicit TenaryOp(OpKernelConstruction* ctx) : OpKernel(ctx) {} + + void Compute(OpKernelContext* ctx) override { + const Tensor& in0 = ctx->input(0); + const Tensor& in1 = ctx->input(1); + const Tensor& in2 = ctx->input(2); + + auto in0_flat = in0.flat(); + auto in1_flat = in1.flat(); + auto in2_flat = in2.flat(); + const Device& d = ctx->eigen_device(); + + Tensor* out = nullptr; + OP_REQUIRES_OK( + ctx, ctx->forward_input_or_allocate_output({0}, 0, in0.shape(), &out)); + auto out_flat = out->flat(); + if (in1.shape() == in2.shape()) { + if (in0.shape() == in1.shape()) { + out_flat = in0_flat.cwiseMin(in2_flat).cwiseMax(in1_flat); + } else { + OP_REQUIRES(ctx, TensorShapeUtils::IsScalar(in1.shape()), + errors::InvalidArgument( + "clip_value_min and clip_value_max must be either of " + "the same shape as input, or a scalar. ", + "input shape: ", in0.shape().DebugString(), + "clip_value_min shape: ", in1.shape().DebugString(), + "clip_value_max shape: ", in2.shape().DebugString())); + out_flat = in0_flat.unaryExpr(UnaryClipOp(in1_flat(0), in2_flat(0))); + } + } else { + if (in0.shape() == in1.shape()) { + OP_REQUIRES(ctx, TensorShapeUtils::IsScalar(in2.shape()), + errors::InvalidArgument( + "clip_value_min and clip_value_max must be either of " + "the same shape as input, or a scalar. ", + "input shape: ", in0.shape().DebugString(), + "clip_value_min shape: ", in1.shape().DebugString(), + "clip_value_max shape: ", in2.shape().DebugString())); + out_flat = + in0_flat.binaryExpr(in1_flat, BinaryClipMaxOp(in2_flat(0))); + + } else { + OP_REQUIRES(ctx, (in0.shape() == in2.shape() && + TensorShapeUtils::IsScalar(in1.shape())), + errors::InvalidArgument( + "clip_value_min and clip_value_max must be either of " + "the same shape as input, or a scalar. ", + "input shape: ", in0.shape().DebugString(), + "clip_value_min shape: ", in1.shape().DebugString(), + "clip_value_max shape: ", in2.shape().DebugString())); + out_flat = + in0_flat.binaryExpr(in2_flat, BinaryClipMinOp(in1_flat(0))); + } + } + } +}; + +#define REGISTER_CPU_KERNEL(type) \ + REGISTER_KERNEL_BUILDER( \ + Name("ClipByValue").Device(DEVICE_CPU).TypeConstraint("T"), \ + TenaryOp); + +REGISTER_CPU_KERNEL(Eigen::half); +REGISTER_CPU_KERNEL(float); +REGISTER_CPU_KERNEL(double); +REGISTER_CPU_KERNEL(int8); +REGISTER_CPU_KERNEL(int16); +REGISTER_CPU_KERNEL(int32); +REGISTER_CPU_KERNEL(int64); +REGISTER_CPU_KERNEL(uint8); +REGISTER_CPU_KERNEL(uint16); + +#undef REGISTER_CPU_KERNEL + +#if GOOGLE_CUDA +// REGISTER3(BinaryOp, GPU, "Add", functor::add, float, Eigen::half, double); + +// A special GPU kernel for int32. +// TODO(b/25387198): Also enable int32 in device memory. This kernel +// registration requires all int32 inputs and outputs to be in host memory. +REGISTER_KERNEL_BUILDER(Name("ClipByValue") + .Device(DEVICE_GPU) + .HostMemory("t") + .HostMemory("clip_value_min") + .HostMemory("clip_value_min") + .TypeConstraint("T"), + TenaryOp); +#endif + +} // namespace tensorflow diff --git a/tensorflow/core/ops/math_ops.cc b/tensorflow/core/ops/math_ops.cc index 8f33d51d5a..602a6ec115 100644 --- a/tensorflow/core/ops/math_ops.cc +++ b/tensorflow/core/ops/math_ops.cc @@ -1558,6 +1558,29 @@ REGISTER_OP("Bucketize") .Attr("boundaries: list(float)") .SetShapeFn(shape_inference::UnchangedShape); +REGISTER_OP("ClipByValue") + .Input("t: T") + .Input("clip_value_min: T") + .Input("clip_value_max: T") + .Output("output: T") + .Attr("T: numbertype") + .SetShapeFn(shape_inference::UnchangedShape) + .Doc(R"doc( +Clips tensor values to a specified min and max. + +Given a tensor `t`, this operation returns a tensor of the same type and +shape as `t` with its values clipped to `clip_value_min` and `clip_value_max`. +Any values less than `clip_value_min` are set to `clip_value_min`. Any values +greater than `clip_value_max` are set to `clip_value_max`. + +t: A `Tensor`. +clip_value_min: A 0-D (scalar) `Tensor`, or a `Tensor` with the same shape + as `t`. The minimum value to clip by. +clip_value_max: A 0-D (scalar) `Tensor`, or a `Tensor` with the same shape + as `t`. The maximum value to clip by. +output: A clipped `Tensor` with the same shape as input 't'. +)doc"); + #ifdef INTEL_MKL REGISTER_OP("_MklAddN") .Input("inputs: N * T") -- GitLab From daf0b206b5afde875a19270136ad22d9d2bb138c Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Wed, 25 Oct 2017 17:08:32 +0000 Subject: [PATCH 014/528] Add python wrapper for tf.clip_by_value Signed-off-by: Yong Tang --- tensorflow/python/ops/clip_ops.py | 17 +- tensorflow/python/ops/hidden_ops.txt | 395 +++++++++++++++++++++++++++ 2 files changed, 400 insertions(+), 12 deletions(-) create mode 100644 tensorflow/python/ops/hidden_ops.txt diff --git a/tensorflow/python/ops/clip_ops.py b/tensorflow/python/ops/clip_ops.py index 49f8c66531..a5baebb3f6 100644 --- a/tensorflow/python/ops/clip_ops.py +++ b/tensorflow/python/ops/clip_ops.py @@ -26,6 +26,7 @@ from tensorflow.python.framework import constant_op from tensorflow.python.framework import dtypes from tensorflow.python.framework import ops from tensorflow.python.ops import array_ops +from tensorflow.python.ops import gen_math_ops from tensorflow.python.ops import gen_nn_ops from tensorflow.python.ops import math_ops from tensorflow.python.util.tf_export import tf_export @@ -58,18 +59,10 @@ def clip_by_value(t, clip_value_min, clip_value_max, """ with ops.name_scope(name, "clip_by_value", [t, clip_value_min, clip_value_max]) as name: - t = ops.convert_to_tensor(t, name="t") - - # Go through list of tensors, for each value in each tensor clip - t_min = math_ops.minimum(t, clip_value_max) - # Assert that the shape is compatible with the initial shape, - # to prevent unintentional broadcasting. - _ = t.shape.merge_with(t_min.shape) - - t_max = math_ops.maximum(t_min, clip_value_min, name=name) - _ = t.shape.merge_with(t_max.shape) - - return t_max + return gen_math_ops._clip_by_value(t, + clip_value_min, + clip_value_max, + name=name) @tf_export("clip_by_norm") diff --git a/tensorflow/python/ops/hidden_ops.txt b/tensorflow/python/ops/hidden_ops.txt new file mode 100644 index 0000000000..e1217e984c --- /dev/null +++ b/tensorflow/python/ops/hidden_ops.txt @@ -0,0 +1,395 @@ +# array_ops +BatchToSpace +BroadcastArgs +BroadcastGradientArgs +ConcatOffset +Concat +ConcatV2 +ConjugateTranspose +Const +DebugGradientIdentity +DebugGradientRefIdentity +EditDistance +ExpandDims +ListDiff +MirrorPad +MirrorPadGrad +OneHot +Pack +Pad +PadV2 +ParallelConcat +Placeholder +RefIdentity +Reverse +Snapshot +SpaceToBatch +Split +SplitV +Squeeze +Slice +TileGrad # Exported through array_grad instead of array_ops. +ZerosLike # TODO(josh11b): Use this instead of the Python version. +Unique +UniqueV2 +UniqueWithCounts +UniqueWithCountsV2 +Unpack + +# candidate_sampling_ops +AllCandidateSampler +ComputeAccidentalHits +FixedUnigramCandidateSampler +LearnedUnigramCandidateSampler +LogUniformCandidateSampler +ThreadUnsafeUnigramCandidateSampler +UniformCandidateSampler + +# checkpoint_ops +GenerateVocabRemapping +LoadAndRemapMatrix + + +# control_flow_ops +Switch +Merge +RefMerge +Exit +RefExit + +# ctc_ops +CTCLoss +CTCGreedyDecoder +CTCBeamSearchDecoder + +# data_flow_ops +Barrier +BarrierClose +BarrierIncompleteSize +BarrierInsertMany +BarrierReadySize +BarrierTakeMany +DeleteSessionTensor +FakeQueue +FIFOQueue +FIFOQueueV2 +GetSessionHandle +GetSessionHandleV2 +GetSessionTensor +HashTable +HashTableV2 +InitializeTable +InitializeTableV2 +InitializeTableFromTextFile +InitializeTableFromTextFileV2 +LookupTableExport +LookupTableExportV2 +LookupTableFind +LookupTableFindV2 +LookupTableImport +LookupTableImportV2 +LookupTableInsert +LookupTableInsertV2 +LookupTableSize +LookupTableSizeV2 +MutableDenseHashTable +MutableDenseHashTableV2 +MutableHashTable +MutableHashTableV2 +MutableHashTableOfTensors +MutableHashTableOfTensorsV2 +Mutex +MutexAcquire +MutexRelease +PaddingFIFOQueue +PaddingFIFOQueueV2 +PriorityQueue +PriorityQueueV2 +QueueClose +QueueCloseV2 +QueueDequeue +QueueDequeueV2 +QueueDequeueMany +QueueDequeueManyV2 +QueueDequeueUpTo +QueueDequeueUpToV2 +QueueEnqueue +QueueEnqueueV2 +QueueEnqueueMany +QueueEnqueueManyV2 +QueueSize +QueueSizeV2 +RandomShuffleQueue +RandomShuffleQueueV2 +Stack +StackClose +StackPop +StackPush +StackV2 +StackCloseV2 +StackPopV2 +StackPushV2 +TensorArray +TensorArrayClose +TensorArrayCloseV2 +TensorArrayConcat +TensorArrayConcatV2 +TensorArrayGather +TensorArrayGatherV2 +TensorArrayGrad +TensorArrayGradV2 +TensorArrayPack +TensorArrayPackV2 +TensorArrayRead +TensorArrayReadV2 +TensorArrayScatter +TensorArrayScatterV2 +TensorArraySize +TensorArraySizeV2 +TensorArraySplit +TensorArraySplitV2 +TensorArrayUnpack +TensorArrayUnpackV2 +TensorArrayV2 +TensorArrayWrite +TensorArrayWriteV2 +TensorArrayV3 +TensorArrayCloseV3 +TensorArrayConcatV3 +TensorArrayGatherV3 +TensorArrayGradV3 +TensorArrayReadV3 +TensorArrayPackV3 +TensorArrayScatterV3 +TensorArraySizeV3 +TensorArraySplitV3 +TensorArrayUnpackV3 +TensorArrayWriteV3 + +# functional_ops +SymbolicGradient + +# image_ops +AdjustContrastv2 +NonMaxSuppression +NonMaxSuppressionV2 +RandomCrop +ResizeBilinearGrad +ResizeBicubicGrad +ResizeNearestNeighborGrad +SampleDistortedBoundingBox +SampleDistortedBoundingBoxV2 +ScaleImageGrad + +# io_ops +FixedLengthRecordReader +IdentityReader +ReaderNumRecordsProduced +ReaderNumWorkUnitsCompleted +ReaderRead +ReaderReadUpTo +ReaderReset +ReaderRestoreState +ReaderSerializeState +ReaderWorkQueueLength +FixedLengthRecordReaderV2 +IdentityReaderV2 +ReaderNumRecordsProducedV2 +ReaderNumWorkUnitsCompletedV2 +ReaderReadV2 +ReaderReadUpToV2 +ReaderResetV2 +ReaderRestoreStateV2 +ReaderSerializeStateV2 +ReaderWorkQueueLengthV2 +Restore +RestoreSlice +Save +SaveSlices +ShardedFilename +ShardedFilespec +TextLineReader +TFRecordReader +WholeFileReader +TextLineReaderV2 +TFRecordReaderV2 +WholeFileReaderV2 +LMDBReader +DecodeCSV + +# linalg_ops +BatchCholesky +BatchCholeskyGrad +BatchMatrixDeterminant +BatchMatrixInverse +BatchMatrixSolve +BatchMatrixSolveLs +BatchMatrixTriangularSolve +BatchSelfAdjointEig +BatchSelfAdjointEigV2 +BatchSvd +LogMatrixDeterminant +MatrixExponential +MatrixLogarithm +MatrixSolveLs +SelfAdjointEig +SelfAdjointEigV2 +Svd + +# logging_ops +Assert +AudioSummary +AudioSummaryV2 +HistogramSummary +ImageSummary +MergeSummary +Print +ScalarSummary +TensorSummary +TensorSummaryV2 + +# math_ops +Abs +AccumulateNV2 +AddN +AddV2 +All +Any +BatchMatMul +BatchFFT +BatchFFT2D +BatchFFT3D +BatchIFFT +BatchIFFT2D +BatchIFFT3D +Bucketize +ClipByValue +Complex +ComplexAbs +Conj +FloorDiv +FloorMod +HistogramFixedWidth +Max +Mean +Min +Mul +Neg +Pow +Prod +Range +RealDiv +Select +SparseMatMul +Sub +Sum +MatMul +Sigmoid +Tanh +SigmoidGrad +TanhGrad +InvGrad +ReciprocalGrad +SqrtGrad +RsqrtGrad +TruncateDiv +TruncateMod + +# nn_ops +AvgPoolGrad # "*Grad" accessible through nn_grad instead of nn_ops. +AvgPool3DGrad +BatchNormWithGlobalNormalization +BatchNormWithGlobalNormalizationGrad +FusedBatchNorm +FusedBatchNormV2 +SoftmaxCrossEntropyWithLogits +SparseSoftmaxCrossEntropyWithLogits +LRNGrad +MaxPoolGrad +MaxPoolGradWithArgmax +MaxPoolGradGrad +MaxPoolGradGradWithArgmax +MaxPool3DGrad +MaxPool3DGradGrad +ReluGrad +Relu6Grad +EluGrad +SeluGrad +SoftplusGrad +SoftsignGrad +TopK +TopKV2 +BiasAdd +BiasAddV1 +Relu6 +AvgPool +MaxPool +MaxPoolV2 +Softmax +LogSoftmax +FractionalAvgPoolGrad +FractionalMaxPoolGrad +InTopK +InTopKV2 + +# parsing_ops +ParseExample +ParseSingleSequenceExample + +# random_ops +RandomGamma +RandomPoisson +RandomUniform +RandomUniformInt +RandomShuffle +RandomStandardNormal +ParameterizedTruncatedNormal +TruncatedNormal + +# script_ops +PyFunc +PyFuncStateless +EagerPyFunc + +# sdca_ops + +# state_ops +Variable +VariableV2 +TemporaryVariable +DestroyTemporaryVariable + +# sparse_ops +AddSparseToTensorsMap +AddManySparseToTensorsMap +TakeManySparseFromTensorsMap +DeserializeManySparse +DeserializeSparse +SerializeManySparse +SerializeSparse +SparseAdd +SparseAddGrad +SparseConcat +SparseCross +SparseFillEmptyRows +SparseFillEmptyRowsGrad +SparseSplit +SparseSelectLastK +SparseReorder +SparseReshape +SparseToDense +SparseTensorDenseAdd +SparseTensorDenseMatMul + +# string_ops +StringSplit + +# user_ops +Fact + +# training_ops +# (None) + +# word2vec deprecated ops +NegTrain +Skipgram -- GitLab From 90a271e7a37574fc1c90fd6042c3b3972645d114 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Wed, 25 Oct 2017 17:09:05 +0000 Subject: [PATCH 015/528] Update tests for `tf.clip_by_value` Signed-off-by: Yong Tang --- tensorflow/python/kernel_tests/clip_ops_test.py | 11 +++++++---- 1 file changed, 7 insertions(+), 4 deletions(-) diff --git a/tensorflow/python/kernel_tests/clip_ops_test.py b/tensorflow/python/kernel_tests/clip_ops_test.py index 5c8b71da17..d47930350e 100644 --- a/tensorflow/python/kernel_tests/clip_ops_test.py +++ b/tensorflow/python/kernel_tests/clip_ops_test.py @@ -19,6 +19,7 @@ from __future__ import division from __future__ import print_function from tensorflow.python.framework import constant_op +from tensorflow.python.framework import errors_impl from tensorflow.python.framework import ops from tensorflow.python.ops import clip_ops from tensorflow.python.platform import test @@ -42,10 +43,12 @@ class ClipTest(test.TestCase): x = constant_op.constant([-5.0, 2.0, 3.0, 4.0, 5.0, 6.0], shape=[2, 3, 1]) # Use a nonsensical shape. clip = constant_op.constant([1.0, 2.0]) - with self.assertRaises(ValueError): - _ = clip_ops.clip_by_value(x, -clip, clip) - with self.assertRaises(ValueError): - _ = clip_ops.clip_by_value(x, 1.0, clip) + with self.assertRaises(errors_impl.InvalidArgumentError): + ans = clip_ops.clip_by_value(x, -clip, clip) + tf_ans = ans.eval() + with self.assertRaises(errors_impl.InvalidArgumentError): + ans = clip_ops.clip_by_value(x, 1.0, clip) + tf_ans = ans.eval() def testClipByValueNonFinite(self): with self.test_session(): -- GitLab From cff8abcb1a9305491637dc44559316aa1d8184e6 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Thu, 26 Oct 2017 04:37:55 +0000 Subject: [PATCH 016/528] Add GPU kernel for tf.clip_by_value Signed-off-by: Yong Tang --- tensorflow/core/kernels/cwise_op_clip.cc | 162 +++++++++++++----- tensorflow/core/kernels/cwise_op_clip.h | 61 +++++++ .../core/kernels/cwise_op_clip_gpu.cu.cc | 134 +++++++++++++++ 3 files changed, 313 insertions(+), 44 deletions(-) create mode 100644 tensorflow/core/kernels/cwise_op_clip.h create mode 100644 tensorflow/core/kernels/cwise_op_clip_gpu.cu.cc diff --git a/tensorflow/core/kernels/cwise_op_clip.cc b/tensorflow/core/kernels/cwise_op_clip.cc index 6ce062b08f..c2980acdd8 100644 --- a/tensorflow/core/kernels/cwise_op_clip.cc +++ b/tensorflow/core/kernels/cwise_op_clip.cc @@ -13,43 +13,12 @@ See the License for the specific language governing permissions and limitations under the License. ==============================================================================*/ -#include "tensorflow/core/kernels/cwise_ops_common.h" - -//#include "third_party/eigen3/Eigen/Core/CwiseTernaryOp.h" +#include "tensorflow/core/kernels/cwise_op_clip.h" namespace tensorflow { -// Unary functor for clip -template -struct UnaryClipOp { - UnaryClipOp(const T& value_min, const T& value_max) - : value_min_(value_min), value_max_(value_max) {} - const T operator()(const T& value) const { - return std::max(std::min(value, value_max_), value_min_); - } - T value_min_; - T value_max_; -}; - -// Binary functor for clip -template -struct BinaryClipMinOp { - BinaryClipMinOp(const T& value_min) : value_min_(value_min) {} - const T operator()(const T& value, const T& value_max) const { - return std::max(std::min(value, value_max), value_min_); - } - T value_min_; -}; - -// Binary functor for clip -template -struct BinaryClipMaxOp { - BinaryClipMaxOp(const T& value_max) : value_max_(value_max) {} - const T operator()(const T& value, const T& value_min) const { - return std::max(std::min(value, value_max_), value_min); - } - T value_max_; -}; +typedef Eigen::ThreadPoolDevice CPUDevice; +typedef Eigen::GpuDevice GPUDevice; // Basic coefficient-wise tenary operations. // This is the case for example of the clip_by_value. @@ -76,7 +45,8 @@ class TenaryOp : public OpKernel { auto out_flat = out->flat(); if (in1.shape() == in2.shape()) { if (in0.shape() == in1.shape()) { - out_flat = in0_flat.cwiseMin(in2_flat).cwiseMax(in1_flat); + functor::TernaryClipOp()(d, in0_flat, in1_flat, in2_flat, + out_flat); } else { OP_REQUIRES(ctx, TensorShapeUtils::IsScalar(in1.shape()), errors::InvalidArgument( @@ -85,7 +55,8 @@ class TenaryOp : public OpKernel { "input shape: ", in0.shape().DebugString(), "clip_value_min shape: ", in1.shape().DebugString(), "clip_value_max shape: ", in2.shape().DebugString())); - out_flat = in0_flat.unaryExpr(UnaryClipOp(in1_flat(0), in2_flat(0))); + functor::UnaryClipOp()(d, in0_flat, in1_flat, in2_flat, + out_flat); } } else { if (in0.shape() == in1.shape()) { @@ -96,9 +67,8 @@ class TenaryOp : public OpKernel { "input shape: ", in0.shape().DebugString(), "clip_value_min shape: ", in1.shape().DebugString(), "clip_value_max shape: ", in2.shape().DebugString())); - out_flat = - in0_flat.binaryExpr(in1_flat, BinaryClipMaxOp(in2_flat(0))); - + functor::BinaryLeftClipOp()(d, in0_flat, in1_flat, in2_flat, + out_flat); } else { OP_REQUIRES(ctx, (in0.shape() == in2.shape() && TensorShapeUtils::IsScalar(in1.shape())), @@ -108,13 +78,103 @@ class TenaryOp : public OpKernel { "input shape: ", in0.shape().DebugString(), "clip_value_min shape: ", in1.shape().DebugString(), "clip_value_max shape: ", in2.shape().DebugString())); - out_flat = - in0_flat.binaryExpr(in2_flat, BinaryClipMinOp(in1_flat(0))); + functor::BinaryRightClipOp()(d, in0_flat, in1_flat, in2_flat, + out_flat); } } } }; +namespace functor { +// Unary functor for clip [Tensor, Scalar, Scalar] +template +struct UnaryClipFunc { + UnaryClipFunc(const T& value_min, const T& value_max) + : value_min_(value_min), value_max_(value_max) {} + const T operator()(const T& value) const { + return std::max(std::min(value, value_max_), value_min_); + } + T value_min_; + T value_max_; +}; +template +struct UnaryClipOp { + void operator()(const CPUDevice& d, typename TTypes::ConstFlat& in0_flat, + typename TTypes::ConstFlat& in1_flat, + typename TTypes::ConstFlat& in2_flat, + typename TTypes::Flat& out_flat) const { + out_flat = in0_flat.unaryExpr(UnaryClipFunc(in1_flat(0), in2_flat(0))); + } +}; + +// Binary functor for clip [Tensor, Scalar, Tensor] +template +struct BinaryRightClipFunc { + BinaryRightClipFunc(const T& value_min) : value_min_(value_min) {} + const T operator()(const T& value, const T& value_max) const { + return std::max(std::min(value, value_max), value_min_); + } + T value_min_; +}; +template +struct BinaryRightClipOp { + void operator()(const CPUDevice& d, typename TTypes::ConstFlat& in0_flat, + typename TTypes::ConstFlat& in1_flat, + typename TTypes::ConstFlat& in2_flat, + typename TTypes::Flat& out_flat) const { + out_flat = + in0_flat.binaryExpr(in2_flat, BinaryRightClipFunc(in1_flat(0))); + } +}; + +// Binary functor for clip [Tensor, Tensor, Scalar] +template +struct BinaryLeftClipFunc { + BinaryLeftClipFunc(const T& value_max) : value_max_(value_max) {} + const T operator()(const T& value, const T& value_min) const { + return std::max(std::min(value, value_max_), value_min); + } + T value_max_; +}; +template +struct BinaryLeftClipOp { + void operator()(const CPUDevice& d, typename TTypes::ConstFlat& in0_flat, + typename TTypes::ConstFlat& in1_flat, + typename TTypes::ConstFlat& in2_flat, + typename TTypes::Flat& out_flat) const { + out_flat = + in0_flat.binaryExpr(in1_flat, BinaryLeftClipFunc(in2_flat(0))); + } +}; + +// Ternary functor for clip [Tensor, Tensor, Tensor] +template +struct TernaryClipOp { + void operator()(const CPUDevice& d, typename TTypes::ConstFlat& in0_flat, + typename TTypes::ConstFlat& in1_flat, + typename TTypes::ConstFlat& in2_flat, + typename TTypes::Flat& out_flat) const { + out_flat.device(d) = in0_flat.cwiseMin(in2_flat).cwiseMax(in1_flat); + } +}; + +#define INSTANTIATE_CPU(T) \ + template struct UnaryClipOp; \ + template struct BinaryRightClipOp; \ + template struct BinaryLeftClipOp; \ + template struct TernaryClipOp; +INSTANTIATE_CPU(Eigen::half); +INSTANTIATE_CPU(float); +INSTANTIATE_CPU(double); +INSTANTIATE_CPU(int8); +INSTANTIATE_CPU(int16); +INSTANTIATE_CPU(int32); +INSTANTIATE_CPU(int64); +INSTANTIATE_CPU(uint8); +INSTANTIATE_CPU(uint16); +#undef INSTANTIATE_CPU +} // namespace functor + #define REGISTER_CPU_KERNEL(type) \ REGISTER_KERNEL_BUILDER( \ Name("ClipByValue").Device(DEVICE_CPU).TypeConstraint("T"), \ @@ -129,11 +189,22 @@ REGISTER_CPU_KERNEL(int32); REGISTER_CPU_KERNEL(int64); REGISTER_CPU_KERNEL(uint8); REGISTER_CPU_KERNEL(uint16); - #undef REGISTER_CPU_KERNEL #if GOOGLE_CUDA -// REGISTER3(BinaryOp, GPU, "Add", functor::add, float, Eigen::half, double); + +#define REGISTER_GPU_KERNEL(type) \ + REGISTER_KERNEL_BUILDER( \ + Name("ClipByValue").Device(DEVICE_GPU).TypeConstraint("T"), \ + TenaryOp); +REGISTER_GPU_KERNEL(Eigen::half); +REGISTER_GPU_KERNEL(float); +REGISTER_GPU_KERNEL(double); +REGISTER_GPU_KERNEL(int8); +REGISTER_GPU_KERNEL(int16); +REGISTER_GPU_KERNEL(int64); +REGISTER_GPU_KERNEL(uint8); +REGISTER_GPU_KERNEL(uint16); // A special GPU kernel for int32. // TODO(b/25387198): Also enable int32 in device memory. This kernel @@ -142,9 +213,12 @@ REGISTER_KERNEL_BUILDER(Name("ClipByValue") .Device(DEVICE_GPU) .HostMemory("t") .HostMemory("clip_value_min") - .HostMemory("clip_value_min") + .HostMemory("clip_value_max") + .HostMemory("output") .TypeConstraint("T"), TenaryOp); + +#undef REGISTER_GPU_KERNEL #endif } // namespace tensorflow diff --git a/tensorflow/core/kernels/cwise_op_clip.h b/tensorflow/core/kernels/cwise_op_clip.h new file mode 100644 index 0000000000..1a4bf8cf1d --- /dev/null +++ b/tensorflow/core/kernels/cwise_op_clip.h @@ -0,0 +1,61 @@ +/* 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. +==============================================================================*/ + +#ifndef TENSORFLOW_KERNELS_CWISE_OP_CLIP_H_ +#define TENSORFLOW_KERNELS_CWISE_OP_CLIP_H_ + +#include "tensorflow/core/kernels/cwise_ops_common.h" + +namespace tensorflow { +namespace functor { +// Unary functor for clip [Tensor, Scalar, Scalar] +template +struct UnaryClipOp { + void operator()(const Device &d, typename TTypes::ConstFlat &in0_flat, + typename TTypes::ConstFlat &in1_flat, + typename TTypes::ConstFlat &in2_flat, + typename TTypes::Flat &out_flat) const; +}; + +// Binary functor for clip [Tensor, Scalar, Tensor] +template +struct BinaryRightClipOp { + void operator()(const Device &d, typename TTypes::ConstFlat &in0_flat, + typename TTypes::ConstFlat &in1_flat, + typename TTypes::ConstFlat &in2_flat, + typename TTypes::Flat &out_flat) const; +}; + +// Binary functor for clip [Tensor, Tensor, Scalar] +template +struct BinaryLeftClipOp { + void operator()(const Device &d, typename TTypes::ConstFlat &in0_flat, + typename TTypes::ConstFlat &in1_flat, + typename TTypes::ConstFlat &in2_flat, + typename TTypes::Flat &out_flat) const; +}; + +// Ternary functor for clip [Tensor, Tensor, Tensor] +template +struct TernaryClipOp { + void operator()(const Device &d, typename TTypes::ConstFlat &in0_flat, + typename TTypes::ConstFlat &in1_flat, + typename TTypes::ConstFlat &in2_flat, + typename TTypes::Flat &out_flat) const; +}; +} +} // namespace tensorflow + +#endif // TENSORFLOW_KERNELS_CWISE_OP_CLIP_H_ diff --git a/tensorflow/core/kernels/cwise_op_clip_gpu.cu.cc b/tensorflow/core/kernels/cwise_op_clip_gpu.cu.cc new file mode 100644 index 0000000000..5c07847548 --- /dev/null +++ b/tensorflow/core/kernels/cwise_op_clip_gpu.cu.cc @@ -0,0 +1,134 @@ +/* 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. +==============================================================================*/ + +#if GOOGLE_CUDA + +#define EIGEN_USE_GPU + +#include "tensorflow/core/kernels/cwise_op_clip.h" +#include "tensorflow/core/kernels/cwise_ops_gpu_common.cu.h" +#include "tensorflow/core/util/cuda_kernel_helper.h" + +namespace tensorflow { + +template +__global__ void UnaryClipCustomKernel(const int32 size_in, const T *in0, + const T *in1, const T *in2, T *out) { + CUDA_1D_KERNEL_LOOP(i, size_in) { + T value = in2[0] < in0[i] ? in2[0] : in0[i]; + out[i] = value < in1[0] ? in1[0] : value; + } +} + +template +__global__ void BinaryRightClipCustomKernel(const int32 size_in, const T *in0, + const T *in1, const T *in2, + T *out) { + CUDA_1D_KERNEL_LOOP(i, size_in) { + T value = in2[i] < in0[i] ? in2[i] : in0[i]; + out[i] = value < in1[0] ? in1[0] : value; + } +} + +template +__global__ void BinaryLeftClipCustomKernel(const int32 size_in, const T *in0, + const T *in1, const T *in2, T *out) { + CUDA_1D_KERNEL_LOOP(i, size_in) { + T value = in2[0] < in0[i] ? in2[0] : in0[i]; + out[i] = value < in1[i] ? in1[i] : value; + } +} + +namespace functor { + +// Unary functor for clip [Tensor, Scalar, Scalar] +template +struct UnaryClipOp { + void operator()(const GPUDevice &d, typename TTypes::ConstFlat &in0_flat, + typename TTypes::ConstFlat &in1_flat, + typename TTypes::ConstFlat &in2_flat, + typename TTypes::Flat &out_flat) const { + CudaLaunchConfig config = GetCudaLaunchConfig(in0_flat.size(), d); + + UnaryClipCustomKernel< + T><<>>( + in0_flat.size(), in0_flat.data(), in1_flat.data(), in2_flat.data(), + out_flat.data()); + } +}; + +// Binary functor for clip [Tensor, Scalar, Tensor] +template +struct BinaryRightClipOp { + void operator()(const GPUDevice &d, typename TTypes::ConstFlat &in0_flat, + typename TTypes::ConstFlat &in1_flat, + typename TTypes::ConstFlat &in2_flat, + typename TTypes::Flat &out_flat) const { + CudaLaunchConfig config = GetCudaLaunchConfig(in0_flat.size(), d); + + BinaryRightClipCustomKernel< + T><<>>( + in0_flat.size(), in0_flat.data(), in1_flat.data(), in2_flat.data(), + out_flat.data()); + } +}; + +// Binary functor for clip [Tensor, Tensor, Scalar] +template +struct BinaryLeftClipOp { + void operator()(const GPUDevice &d, typename TTypes::ConstFlat &in0_flat, + typename TTypes::ConstFlat &in1_flat, + typename TTypes::ConstFlat &in2_flat, + typename TTypes::Flat &out_flat) const { + CudaLaunchConfig config = GetCudaLaunchConfig(in0_flat.size(), d); + + BinaryLeftClipCustomKernel< + T><<>>( + in0_flat.size(), in0_flat.data(), in1_flat.data(), in2_flat.data(), + out_flat.data()); + } +}; + +// Ternary functor for clip [Tensor, Tensor, Tensor] +template +struct TernaryClipOp { + void operator()(const GPUDevice &d, typename TTypes::ConstFlat &in0_flat, + typename TTypes::ConstFlat &in1_flat, + typename TTypes::ConstFlat &in2_flat, + typename TTypes::Flat &out_flat) const { + out_flat.device(d) = in0_flat.cwiseMin(in2_flat).cwiseMax(in1_flat); + } +}; + +#define INSTANTIATE_GPU(T) \ + template struct UnaryClipOp; \ + template struct BinaryRightClipOp; \ + template struct BinaryLeftClipOp; \ + template struct TernaryClipOp; +INSTANTIATE_GPU(Eigen::half); +INSTANTIATE_GPU(float); +INSTANTIATE_GPU(double); +INSTANTIATE_GPU(int8); +INSTANTIATE_GPU(int16); +INSTANTIATE_GPU(int32); +INSTANTIATE_GPU(int64); +INSTANTIATE_GPU(uint8); +INSTANTIATE_GPU(uint16); +#undef INSTANTIATE_GPU + +} // namespace functor +} // namespace tensorflow + +#endif // GOOGLE_CUDA -- GitLab From a3553d45b63fba1cd4eb8d1d5b6dd0d565c94879 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Thu, 26 Oct 2017 04:38:38 +0000 Subject: [PATCH 017/528] Update test cases for tf.clip_by_value Signed-off-by: Yong Tang --- .../python/kernel_tests/clip_ops_test.py | 105 ++++++++++++++---- 1 file changed, 85 insertions(+), 20 deletions(-) diff --git a/tensorflow/python/kernel_tests/clip_ops_test.py b/tensorflow/python/kernel_tests/clip_ops_test.py index d47930350e..2d03fb99e4 100644 --- a/tensorflow/python/kernel_tests/clip_ops_test.py +++ b/tensorflow/python/kernel_tests/clip_ops_test.py @@ -19,6 +19,7 @@ from __future__ import division from __future__ import print_function from tensorflow.python.framework import constant_op +from tensorflow.python.framework import dtypes from tensorflow.python.framework import errors_impl from tensorflow.python.framework import ops from tensorflow.python.ops import clip_ops @@ -29,7 +30,7 @@ class ClipTest(test.TestCase): # ClipByValue test def testClipByValue(self): - with self.test_session(): + with self.test_session(use_gpu=True): x = constant_op.constant([-5.0, 2.0, 3.0, 4.0, 5.0, 6.0], shape=[2, 3]) np_ans = [[-4.4, 2.0, 3.0], [4.0, 4.4, 4.4]] clip_value = 4.4 @@ -38,8 +39,72 @@ class ClipTest(test.TestCase): self.assertAllClose(np_ans, tf_ans) + # [Tensor, Scalar, Scalar] + def testClipByValue0Type(self): + for dtype in [dtypes.float16, dtypes.float32, dtypes.float64, + dtypes.int8, dtypes.int16, dtypes.int32, dtypes.int64, + dtypes.uint8, dtypes.uint16]: + with self.test_session(use_gpu=True): + x = constant_op.constant([1, 2, 3, 4, 5, 6], shape=[2, 3], dtype=dtype) + np_ans = [[2, 2, 3], [4, 4, 4]] + clip_value_min = 2 + clip_value_max = 4 + ans = clip_ops.clip_by_value(x, clip_value_min, clip_value_max) + tf_ans = ans.eval() + + self.assertAllClose(np_ans, tf_ans) + + # [Tensor, Tensor, Scalar] + def testClipByValue1Type(self): + for dtype in [dtypes.float16, dtypes.float32, dtypes.float64, + dtypes.int8, dtypes.int16, dtypes.int32, dtypes.int64, + dtypes.uint8, dtypes.uint16]: + with self.test_session(use_gpu=True): + x = constant_op.constant([1, 2, 3, 4, 5, 6], shape=[2, 3], dtype=dtype) + np_ans = [[2, 2, 3], [4, 4, 4]] + clip_value_min = constant_op.constant([2, 2, 2, 3, 3, 3], shape=[2, 3], + dtype=dtype) + clip_value_max = 4 + ans = clip_ops.clip_by_value(x, clip_value_min, clip_value_max) + tf_ans = ans.eval() + + self.assertAllClose(np_ans, tf_ans) + + # [Tensor, Scalar, Tensor] + def testClipByValue2Type(self): + for dtype in [dtypes.float16, dtypes.float32, dtypes.float64, + dtypes.int8, dtypes.int16, dtypes.int32, dtypes.int64, + dtypes.uint8, dtypes.uint16]: + with self.test_session(use_gpu=True): + x = constant_op.constant([1, 2, 3, 4, 5, 6], shape=[2, 3], dtype=dtype) + np_ans = [[4, 4, 4], [4, 5, 6]] + clip_value_min = 4 + clip_value_max = constant_op.constant([6, 6, 6, 6, 6, 6], shape=[2, 3], + dtype=dtype) + ans = clip_ops.clip_by_value(x, clip_value_min, clip_value_max) + tf_ans = ans.eval() + + self.assertAllClose(np_ans, tf_ans) + + # [Tensor, Tensor, Tensor] + def testClipByValue3Type(self): + for dtype in [dtypes.float16, dtypes.float32, dtypes.float64, + dtypes.int8, dtypes.int16, dtypes.int32, dtypes.int64, + dtypes.uint8, dtypes.uint16]: + with self.test_session(use_gpu=True): + x = constant_op.constant([1, 2, 3, 4, 5, 6], shape=[2, 3], dtype=dtype) + np_ans = [[2, 2, 3], [5, 5, 6]] + clip_value_min = constant_op.constant([2, 2, 2, 5, 5, 5], shape=[2, 3], + dtype=dtype) + clip_value_max = constant_op.constant([5, 5, 5, 7, 7, 7], shape=[2, 3], + dtype=dtype) + ans = clip_ops.clip_by_value(x, clip_value_min, clip_value_max) + tf_ans = ans.eval() + + self.assertAllClose(np_ans, tf_ans) + def testClipByValueBadShape(self): - with self.test_session(): + with self.test_session(use_gpu=True): x = constant_op.constant([-5.0, 2.0, 3.0, 4.0, 5.0, 6.0], shape=[2, 3, 1]) # Use a nonsensical shape. clip = constant_op.constant([1.0, 2.0]) @@ -51,7 +116,7 @@ class ClipTest(test.TestCase): tf_ans = ans.eval() def testClipByValueNonFinite(self): - with self.test_session(): + with self.test_session(use_gpu=True): x = constant_op.constant([float('NaN'), float('Inf'), -float('Inf')]) np_ans = [float('NaN'), 4.0, -4.0] clip_value = 4.0 @@ -63,7 +128,7 @@ class ClipTest(test.TestCase): # ClipByNorm tests def testClipByNormClipped(self): # Norm clipping when clip_norm < 5 - with self.test_session(): + with self.test_session(use_gpu=True): x = constant_op.constant([-3.0, 0.0, 0.0, 4.0, 0.0, 0.0], shape=[2, 3]) # Norm of x = sqrt(3^2 + 4^2) = 5 np_ans = [[-2.4, 0.0, 0.0], [3.2, 0.0, 0.0]] @@ -79,7 +144,7 @@ class ClipTest(test.TestCase): self.assertAllClose(np_ans, tf_ans_tensor) def testClipByNormBadShape(self): - with self.test_session(): + with self.test_session(use_gpu=True): x = constant_op.constant([-3.0, 0.0, 0.0, 4.0, 0.0, 0.0], shape=[2, 3, 1]) # Use a nonsensical shape. clip = constant_op.constant([1.0, 2.0]) @@ -88,7 +153,7 @@ class ClipTest(test.TestCase): def testClipByNormNotClipped(self): # No norm clipping when clip_norm >= 5 - with self.test_session(): + with self.test_session(use_gpu=True): x = constant_op.constant([-3.0, 0.0, 0.0, 4.0, 0.0, 0.0], shape=[2, 3]) # Norm of x = sqrt(3^2 + 4^2) = 5 np_ans = [[-3.0, 0.0, 0.0], [4.0, 0.0, 0.0]] @@ -100,7 +165,7 @@ class ClipTest(test.TestCase): def testClipByNormZero(self): # No norm clipping when norm = 0 - with self.test_session(): + with self.test_session(use_gpu=True): x = constant_op.constant([0.0, 0.0, 0.0, 0.0, 0.0, 0.0], shape=[2, 3]) # Norm = 0, no changes np_ans = [[0.0, 0.0, 0.0], [0.0, 0.0, 0.0]] @@ -112,7 +177,7 @@ class ClipTest(test.TestCase): def testClipByNormClippedWithDim0(self): # Norm clipping when clip_norm < 5 - with self.test_session(): + with self.test_session(use_gpu=True): x = constant_op.constant([-3.0, 0.0, 0.0, 4.0, 0.0, 3.0], shape=[2, 3]) # Norm of x[:, 0] = sqrt(3^2 + 4^2) = 5, x[:, 2] = 3 np_ans = [[-2.4, 0.0, 0.0], [3.2, 0.0, 3.0]] @@ -124,7 +189,7 @@ class ClipTest(test.TestCase): def testClipByNormClippedWithDim1(self): # Norm clipping when clip_norm < 5 - with self.test_session(): + with self.test_session(use_gpu=True): x = constant_op.constant([-3.0, 0.0, 0.0, 4.0, 0.0, 3.0], shape=[2, 3]) # Norm of x[0, :] = 3, x[1, :] = sqrt(3^2 + 4^2) = 5 np_ans = [[-3.0, 0.0, 0.0], [3.2, 0.0, 2.4]] @@ -136,7 +201,7 @@ class ClipTest(test.TestCase): def testClipByNormNotClippedWithAxes(self): # No norm clipping when clip_norm >= 5 - with self.test_session(): + with self.test_session(use_gpu=True): x = constant_op.constant([-3.0, 0.0, 0.0, 4.0, 0.0, 3.0], shape=[2, 3]) # Norm of x[0, :] = 3, x[1, :] = sqrt(3^2 + 4^2) = 5 np_ans = [[-3.0, 0.0, 0.0], [4.0, 0.0, 3.0]] @@ -149,7 +214,7 @@ class ClipTest(test.TestCase): # ClipByGlobalNorm tests def testClipByGlobalNormClipped(self): # Norm clipping when clip_norm < 5 - with self.test_session(): + with self.test_session(use_gpu=True): x0 = constant_op.constant([-2.0, 0.0, 0.0, 4.0, 0.0, 0.0], shape=[2, 3]) x1 = constant_op.constant([1.0, -2.0]) # Global norm of x0 and x1 = sqrt(1 + 4^2 + 2^2 + 2^2) = 5 @@ -170,7 +235,7 @@ class ClipTest(test.TestCase): def testClipByGlobalNormClippedTensor(self): # Norm clipping when clip_norm < 5 - with self.test_session(): + with self.test_session(use_gpu=True): x0 = constant_op.constant([-2.0, 0.0, 0.0, 4.0, 0.0, 0.0], shape=[2, 3]) x1 = constant_op.constant([1.0, -2.0]) # Global norm of x0 and x1 = sqrt(1 + 4^2 + 2^2 + 2^2) = 5 @@ -191,7 +256,7 @@ class ClipTest(test.TestCase): def testClipByGlobalNormSupportsNone(self): # Norm clipping when clip_norm < 5 - with self.test_session(): + with self.test_session(use_gpu=True): x0 = constant_op.constant([-2.0, 0.0, 0.0, 4.0, 0.0, 0.0], shape=[2, 3]) x1 = constant_op.constant([1.0, -2.0]) # Global norm of x0 and x1 = sqrt(1 + 4^2 + 2^2 + 2^2) = 5 @@ -214,7 +279,7 @@ class ClipTest(test.TestCase): def testClipByGlobalNormWithIndexedSlicesClipped(self): # Norm clipping when clip_norm < 5 - with self.test_session(): + with self.test_session(use_gpu=True): x0 = constant_op.constant([-2.0, 0.0, 0.0, 4.0, 0.0, 0.0], shape=[2, 3]) x1 = ops.IndexedSlices( constant_op.constant([1.0, -2.0]), constant_op.constant([3, 4])) @@ -247,7 +312,7 @@ class ClipTest(test.TestCase): def testClipByGlobalNormNotClipped(self): # No norm clipping when clip_norm >= 5 - with self.test_session(): + with self.test_session(use_gpu=True): x0 = constant_op.constant([-2.0, 0.0, 0.0, 4.0, 0.0, 0.0], shape=[2, 3]) x1 = constant_op.constant([1.0, -2.0]) # Global norm of x0 and x1 = sqrt(1 + 4^2 + 2^2 + 2^2) = 5 @@ -266,7 +331,7 @@ class ClipTest(test.TestCase): def testClipByGlobalNormZero(self): # No norm clipping when norm = 0 - with self.test_session(): + with self.test_session(use_gpu=True): x0 = constant_op.constant([0.0, 0.0, 0.0, 0.0, 0.0, 0.0], shape=[2, 3]) x1 = constant_op.constant([0.0, 0.0]) # Norm = 0, no changes @@ -285,7 +350,7 @@ class ClipTest(test.TestCase): def testClipByAverageNormClipped(self): # Norm clipping when average clip_norm < 0.83333333 - with self.test_session(): + with self.test_session(use_gpu=True): x = constant_op.constant([-3.0, 0.0, 0.0, 4.0, 0.0, 0.0], shape=[2, 3]) # Average norm of x = sqrt(3^2 + 4^2) / 6 = 0.83333333 np_ans = [[-2.88, 0.0, 0.0], [3.84, 0.0, 0.0]] @@ -297,7 +362,7 @@ class ClipTest(test.TestCase): def testClipByAverageNormClippedTensor(self): # Norm clipping when average clip_norm < 0.83333333 - with self.test_session(): + with self.test_session(use_gpu=True): x = constant_op.constant([-3.0, 0.0, 0.0, 4.0, 0.0, 0.0], shape=[2, 3]) # Average norm of x = sqrt(3^2 + 4^2) / 6 = 0.83333333 np_ans = [[-2.88, 0.0, 0.0], [3.84, 0.0, 0.0]] @@ -309,7 +374,7 @@ class ClipTest(test.TestCase): def testClipByAverageNormNotClipped(self): # No norm clipping when average clip_norm >= 0.83333333 - with self.test_session(): + with self.test_session(use_gpu=True): x = constant_op.constant([-3.0, 0.0, 0.0, 4.0, 0.0, 0.0], shape=[2, 3]) # Average norm of x = sqrt(3^2 + 4^2) / 6 = 0.83333333 np_ans = [[-3.0, 0.0, 0.0], [4.0, 0.0, 0.0]] @@ -321,7 +386,7 @@ class ClipTest(test.TestCase): def testClipByAverageNormZero(self): # No norm clipping when average clip_norm = 0 - with self.test_session(): + with self.test_session(use_gpu=True): x = constant_op.constant([0.0, 0.0, 0.0, 0.0, 0.0, 0.0], shape=[2, 3]) # Average norm = 0, no changes np_ans = [[0.0, 0.0, 0.0], [0.0, 0.0, 0.0]] -- GitLab From a5e9d9a387680b0b1d7d8ed08fc9c07477a7efe7 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Mon, 30 Oct 2017 23:42:08 +0000 Subject: [PATCH 018/528] Add grad registration for clip_by_value and address review feedbacks. Signed-off-by: Yong Tang --- tensorflow/core/kernels/cwise_op_clip.cc | 2 +- .../python/kernel_tests/clip_ops_test.py | 16 ++++++++++++ tensorflow/python/ops/clip_ops.py | 25 +++++++++++++++++++ 3 files changed, 42 insertions(+), 1 deletion(-) diff --git a/tensorflow/core/kernels/cwise_op_clip.cc b/tensorflow/core/kernels/cwise_op_clip.cc index c2980acdd8..f30c49fdf8 100644 --- a/tensorflow/core/kernels/cwise_op_clip.cc +++ b/tensorflow/core/kernels/cwise_op_clip.cc @@ -1,4 +1,4 @@ -/* Copyright 2015 The TensorFlow Authors. All Rights Reserved. +/* 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. diff --git a/tensorflow/python/kernel_tests/clip_ops_test.py b/tensorflow/python/kernel_tests/clip_ops_test.py index 2d03fb99e4..cb1359be15 100644 --- a/tensorflow/python/kernel_tests/clip_ops_test.py +++ b/tensorflow/python/kernel_tests/clip_ops_test.py @@ -23,11 +23,27 @@ from tensorflow.python.framework import dtypes from tensorflow.python.framework import errors_impl from tensorflow.python.framework import ops from tensorflow.python.ops import clip_ops +from tensorflow.python.ops import gradient_checker from tensorflow.python.platform import test class ClipTest(test.TestCase): + def testClipByValueGradient(self): + inputs = constant_op.constant([1.0, 2.0, 3.0, 4.0], dtype=dtypes.float32) + outputs_1 = clip_ops.clip_by_value(inputs, 0.5, 3.5) + min_val = constant_op.constant([0.5, 0.5, 0.5, 0.5], dtype=dtypes.float32) + max_val = constant_op.constant([3.5, 3.5, 3.5, 3.5], dtype=dtypes.float32) + outputs_2 = clip_ops.clip_by_value(inputs, min_val, max_val) + with self.test_session(): + error_1 = gradient_checker.compute_gradient_error(inputs, [4], + outputs_1, [4]) + self.assertLess(error_1, 1e-4) + + error_2 = gradient_checker.compute_gradient_error(inputs, [4], + outputs_2, [4]) + self.assertLess(error_2, 1e-4) + # ClipByValue test def testClipByValue(self): with self.test_session(use_gpu=True): diff --git a/tensorflow/python/ops/clip_ops.py b/tensorflow/python/ops/clip_ops.py index a5baebb3f6..e84cfc6944 100644 --- a/tensorflow/python/ops/clip_ops.py +++ b/tensorflow/python/ops/clip_ops.py @@ -26,6 +26,7 @@ from tensorflow.python.framework import constant_op from tensorflow.python.framework import dtypes from tensorflow.python.framework import ops from tensorflow.python.ops import array_ops +from tensorflow.python.ops import gen_array_ops from tensorflow.python.ops import gen_math_ops from tensorflow.python.ops import gen_nn_ops from tensorflow.python.ops import math_ops @@ -64,6 +65,30 @@ def clip_by_value(t, clip_value_min, clip_value_max, clip_value_max, name=name) +@ops.RegisterGradient("ClipByValue") +def _ClipByValueGrad(op, grad): + """Returns grad of clip_by_value.""" + x = op.inputs[0] + y = op.inputs[1] + z = op.inputs[2] + gdtype = grad.dtype + sx = array_ops.shape(x) + sy = array_ops.shape(y) + sz = array_ops.shape(z) + gradshape = array_ops.shape(grad) + zeros = array_ops.zeros(gradshape, gdtype) + xymask = math_ops.less(x, y) + xzmask = math_ops.greater(x, z) + rx, ry = gen_array_ops._broadcast_gradient_args(sx, sy) + rx, rz = gen_array_ops._broadcast_gradient_args(sx, sz) + xgrad = array_ops.where(math_ops.logical_or(xymask, xzmask), zeros, grad) + ygrad = array_ops.where(xymask, grad, zeros) + zgrad = array_ops.where(xzmask, grad, zeros) + gx = array_ops.reshape(math_ops.reduce_sum(xgrad, rx), sx) + gy = array_ops.reshape(math_ops.reduce_sum(ygrad, ry), sy) + gz = array_ops.reshape(math_ops.reduce_sum(zgrad, rz), sz) + return (gx, gy, gz) + @tf_export("clip_by_norm") def clip_by_norm(t, clip_norm, axes=None, name=None): -- GitLab From 71ddf90d3c8c49d4401c0d298bf63b92150dadaa Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Thu, 14 Dec 2017 04:06:58 +0000 Subject: [PATCH 019/528] Update with `TenaryOp` -> `ClipOp` Signed-off-by: Yong Tang --- tensorflow/core/kernels/cwise_op_clip.cc | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/tensorflow/core/kernels/cwise_op_clip.cc b/tensorflow/core/kernels/cwise_op_clip.cc index f30c49fdf8..bd22f5777c 100644 --- a/tensorflow/core/kernels/cwise_op_clip.cc +++ b/tensorflow/core/kernels/cwise_op_clip.cc @@ -25,9 +25,9 @@ typedef Eigen::GpuDevice GPUDevice; // Device: E.g., CPUDevice, GPUDevice. // Functor: defined above. E.g., functor::clip. template -class TenaryOp : public OpKernel { +class ClipOp : public OpKernel { public: - explicit TenaryOp(OpKernelConstruction* ctx) : OpKernel(ctx) {} + explicit ClipOp(OpKernelConstruction* ctx) : OpKernel(ctx) {} void Compute(OpKernelContext* ctx) override { const Tensor& in0 = ctx->input(0); @@ -178,7 +178,7 @@ INSTANTIATE_CPU(uint16); #define REGISTER_CPU_KERNEL(type) \ REGISTER_KERNEL_BUILDER( \ Name("ClipByValue").Device(DEVICE_CPU).TypeConstraint("T"), \ - TenaryOp); + ClipOp); REGISTER_CPU_KERNEL(Eigen::half); REGISTER_CPU_KERNEL(float); @@ -196,7 +196,7 @@ REGISTER_CPU_KERNEL(uint16); #define REGISTER_GPU_KERNEL(type) \ REGISTER_KERNEL_BUILDER( \ Name("ClipByValue").Device(DEVICE_GPU).TypeConstraint("T"), \ - TenaryOp); + ClipOp); REGISTER_GPU_KERNEL(Eigen::half); REGISTER_GPU_KERNEL(float); REGISTER_GPU_KERNEL(double); @@ -216,7 +216,7 @@ REGISTER_KERNEL_BUILDER(Name("ClipByValue") .HostMemory("clip_value_max") .HostMemory("output") .TypeConstraint("T"), - TenaryOp); + ClipOp); #undef REGISTER_GPU_KERNEL #endif -- GitLab From d1078b562532e2de60bc16fc544a94823149ae77 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Mon, 18 Dec 2017 17:42:37 +0000 Subject: [PATCH 020/528] Fix failing test //tensorflow/python:function_test Signed-off-by: Yong Tang --- tensorflow/python/framework/function_test.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/python/framework/function_test.py b/tensorflow/python/framework/function_test.py index 65ca801cbe..24aaff3748 100644 --- a/tensorflow/python/framework/function_test.py +++ b/tensorflow/python/framework/function_test.py @@ -1333,7 +1333,7 @@ class UnrollLSTMTest(test.TestCase): value=math_ops.matmul(xm, weights), num_or_size_splits=4, axis=1) new_c = math_ops.sigmoid(f_g) * cprev + math_ops.sigmoid( i_g) * math_ops.tanh(i_i) - new_c = clip_ops.clip_by_value(new_c, -50.0, 50.0) + new_c = math_ops.maximum(math_ops.minimum(new_c, 50.0), -50.0) new_m = math_ops.sigmoid(o_g) * math_ops.tanh(new_c) return new_m, new_c -- GitLab From 14e9c14ecdb9e9ddb283c5ec9cf27b3c5dbb900e Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Mon, 18 Dec 2017 18:58:42 +0000 Subject: [PATCH 021/528] Fix api_compatibility_test with `--update_goldens True` Signed-off-by: Yong Tang --- .../base_api/api_def_ClipByValue.pbtxt | 36 +++++++++++++++++++ .../python_api/api_def_ClipByValue.pbtxt | 4 +++ 2 files changed, 40 insertions(+) create mode 100644 tensorflow/core/api_def/base_api/api_def_ClipByValue.pbtxt create mode 100644 tensorflow/core/api_def/python_api/api_def_ClipByValue.pbtxt diff --git a/tensorflow/core/api_def/base_api/api_def_ClipByValue.pbtxt b/tensorflow/core/api_def/base_api/api_def_ClipByValue.pbtxt new file mode 100644 index 0000000000..803d8970ab --- /dev/null +++ b/tensorflow/core/api_def/base_api/api_def_ClipByValue.pbtxt @@ -0,0 +1,36 @@ +op { + graph_op_name: "ClipByValue" + in_arg { + name: "t" + description: <