experiments/STT-tensorflow
1
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

Merge changes from github.

Browse Source 
END_PUBLIC

I also integrated #13073 by hand to make TAP happy.

---
Commit 92362d0f0 authored by Skye Wanderman-Milne<skyewm@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Add WhileContext class and add plumbing for creating them.

This change introduces WhileContext, which stores information about a
while loop and will be used in future changes to generate while loop
gradient graphs. Exit nodes in a while loop now have a pointer to
their associated WhileContext. This will be used to retrieve the
context for a given loop.

This change adds an optional parameter to BuildWhileLoop() to create a
WhileContext for the while loop (currently this is always true, but
gradients will generate while loops without associated contexts). This
change also adds a as-yet-unused option to BuildWhileLoop() to return
the predicate output.

PiperOrigin-RevId: 168562303

---
Commit a4f6e7c1a authored by RJ Ryan<rjryan@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Add mel-scale conversion matrix support to tf.contrib.signal.

PiperOrigin-RevId: 168560255

---
Commit b00b6d23c authored by Henry Tan<henrytan@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Fix a segmentation fault caused by invalid log directory in InternalFlush().

PiperOrigin-RevId: 168557063

---
Commit 2bc7a155a authored by Yong Tang<yong.tang.github@outlook.com>
Committed by Rasmus Munk Larsen<rmlarsen@google.com>:
Add uint16 support for tf.decode_raw (#12719)

* Add uint16 support for tf.decode_raw

This fix tries to address the request raised in 10124 where
uint16 support for tf.decode_raw is needed. tf.decode_raw
already support half, float32, float64, int8, int16, int32, int64,
uint8. And uint16 was not supported.

This fix adds uint16 support for tf.decode_raw.

This fix fixes 10124.

Signed-off-by: Yong Tang <yong.tang.github@outlook.com>

* Fix test failure caused by uint16 support of decode_raw and add unit tests.

Signed-off-by: Yong Tang <yong.tang.github@outlook.com>

---
Commit 009285c09 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Remove benchmark for TensorShapeOld.

PiperOrigin-RevId: 168551108

---
Commit dc1eda8a6 authored by Peter Hawkins<phawkins@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
[XLA] Fix CHECK-failure crash if a non-tuple was passed to GetTupleElement.

PiperOrigin-RevId: 168550703

---
Commit 010922ed9 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Go: Update generated wrapper functions for TensorFlow ops.

PiperOrigin-RevId: 168549989

---
Commit c8a6131e9 authored by Mark Daoust<markdaoust@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
make `tf.sets` examples executable

Fixes #12969

PiperOrigin-RevId: 168549712

---
Commit bece65c6f authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Use a map instead of a vector of Children() in the BeamEntry.

The assumption is that since the entries are sparse (they are all populated, but most are never Active()), using the map will save memory and make iterating over the Children() more efficient.

PiperOrigin-RevId: 168548814

---
Commit 0d5ab82ce authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Update ops-related pbtxt files.

PiperOrigin-RevId: 168548642

---
Commit 3331c574b authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Implementing gradients for tf.image.resize_bicubic.

PiperOrigin-RevId: 168547412

---
Commit 4982ef0fa authored by Martin Wicke<wicke@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Add the ability to warn only once if deprecated functionality is used, and make that the default.

PiperOrigin-RevId: 168545655

---
Commit 99423416a authored by Peter Hawkins<phawkins@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
[XLA] Make shape inference error messages for the While HLO more readable. Build the error lazily.

PiperOrigin-RevId: 168531083

---
Commit d10374e45 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Discard some unneccessary logging commands.

PiperOrigin-RevId: 168500721

---
Commit 83cbabb85 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Fix wrong format of logging message.

PiperOrigin-RevId: 168497373

---
Commit eec4f1b3a authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Go: Update generated wrapper functions for TensorFlow ops.

PiperOrigin-RevId: 168494944

---
Commit 69301f352 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Update ops-related pbtxt files.

PiperOrigin-RevId: 168494220

---
Commit 9d56f419c authored by Mingxing Tan<tanmingxing@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Add crop_and_decode_jpeg_op that combines the crop and decode for better
performance.

PiperOrigin-RevId: 168493125

---
Commit 48ddf64d0 authored by Chris Leary<leary@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
[XLA] Make large params test only run in opt builds.

PiperOrigin-RevId: 168491913

---
Commit 11d3ac29d authored by Chris Leary<leary@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
[XLA] Add tests for large numbers of parameter / return values and while loops.

PiperOrigin-RevId: 168487225

---
Commit 3cd6bdef5 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Added test cases on R4 slice.

PiperOrigin-RevId: 168482049

---
Commit 46a81b5c3 authored by Jacques Pienaar<jpienaar@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Add cast S64 to F32 test.

PiperOrigin-RevId: 168473650

---
Commit 59bdf598d authored by Derek Murray<mrry@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Add an automatically-generated "tensorflow.python.platform.build_info" script.

The motivation for this script is to provide better tools for
diagnosing load-time errors (such as the ones that plague the Windows
build due to DLL issues). Note that the script is intended to be
self-contained, so that it is possible to import it without loading
the entire TensorFlow runtime.

This generated script currently contains a single symbol,
`is_cuda_build`, which records whether the build has GPU support or not.

PiperOrigin-RevId: 168471034

---
Commit c3b86347f authored by Olivia Nordquist<nolivia@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
reenabling tests that are passing

PiperOrigin-RevId: 168466361

---
Commit c728665ec authored by Henry Tan<henrytan@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Add const qualifiers whenever appropriate.

PiperOrigin-RevId: 168465926

---
Commit bf96fcd13 authored by Alexandre Passos<apassos@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Use the scalar cache in MeanGrad.

PiperOrigin-RevId: 168462267

---
Commit 1cada9ea2 authored by Olivia Nordquist<nolivia@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
reenabling test that passed after 100 runs w/o timing out

PiperOrigin-RevId: 168458634

---
Commit 00c865566 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Generate error (instead of segfault) when trying to copy string tensor
to GPU in EagerTensor constructor.

PiperOrigin-RevId: 168457320

---
Commit 655f26fc7 authored by Alexandre Passos<apassos@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Resurrects autograd-free eager gradients.

PiperOrigin-RevId: 168448557

---
Commit 8f37f3002 authored by Peter Hawkins<phawkins@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
[TF:XLA] Cleanups to handling of arguments during XLA compilation:
* combine resource kinds in XlaCompiler::Argument::Kind, use a separate XlaResource::Kind field to distinguish different kinds of resource.
* merge XlaContext::HandleOrConstant and XlaExpression, which were almost identical.
* remove XlaContext::Argument; instead, build XlaExpressions directly from XlaCompiler and add them to the XlaContext.

PiperOrigin-RevId: 168439341

---
Commit 7f5346a80 authored by Gunhan Gulsoy<gunan@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Reduce cmake log mess.

* Echo off for the .bat scripts.
* TF cmake: disable warnings in some of the patched projects (gif,jpeg,lmdb).

PiperOrigin-RevId: 168432070

---
Commit 2ad85aa4d authored by Mark Heffernan<meheff@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Use xla/tests:xla_internal_test_main for all tests under tf/compiler/xla
and remove any main() definitions in tests. This enables use of flags
in all tests.

PiperOrigin-RevId: 168424796

---
Commit cd377811d authored by Henry Tan<henrytan@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Comment and error message consistency cleanup.

PiperOrigin-RevId: 168422582

---
Commit 7c19b82af authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Update tf.sparse_reset_shape so that when shrinking the shape of an empty
sparse tensor, the result has a shape of all zeros.

PiperOrigin-RevId: 168419639

---
Commit fcacb40d4 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
FirstReadyManager for scheduling nodes in VirtualScheduler.
The current FIFOManager may yield inefficient scheduling; _Recv pushed to the
FIFO blocks other nodes that can run before _Recv due to the node order in FIFO.
FirstReadyManager picks a node with the earliest time_ready in the queue,
avoiding this problem.

Also, fixed VirtualPlacer to properly set device when Node's device name does not
include job name and to set GPU:0 as default device.

PiperOrigin-RevId: 168418455

---
Commit 7e47624f5 authored by Asim Shankar<ashankar@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
eager: Initial support for iteration over tf.contrib.data.Dataset objects.

TODO:
- Support function-valued operation attributes in eager
  (Required for MapDataset, FilterDataset etc. which encode the
  per-element computation in a TensorFlow function)
PiperOrigin-RevId: 168418250

---
Commit b0a397fce authored by Asim Shankar<ashankar@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
eager: Remove unnecessary TFE_Context argument to TFE_OpSetDevice.

PiperOrigin-RevId: 168417999

---
Commit 86211d554 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Graph transform to flatten atrous (dilated) convolutions (i.e., a sequence of SpaceToBatchND-Conv-BatchToSpaceND ops) to a regular Conv op with upsampled filters.

PiperOrigin-RevId: 168414124

---
Commit 3438981ca authored by David G. Andersen<dga@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Apply exported symbol filtering to the c++ API analogously to
what is filtered for the C API.
Fixes bug reported in comments on #1924

PiperOrigin-RevId: 168413719

---
Commit 7e023d865 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
[XLA:CPU] Remove code from parallel CPU backend outlining that was causing unnecessary copies to be inserted, and which is no longer necessary since we added co-located buffer support for kCall.
*) All bitcast copy is no longer necessary as CopyInsertion will insert copies
at the root of the computation for a parameter which is live-out.
*) Copy if root does not define buffer no longer necessary because colocated
assignment looks at points-to set of root instruction.

PiperOrigin-RevId: 168412076

---
Commit 5da4df92c authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Simplify some code in grappler_item_builder.cc, no change in logic.

PiperOrigin-RevId: 168409110

---
Commit 82ec6241a authored by drpngx<drpngx@users.noreply.github.com>
Committed by GitHub<noreply@github.com>:
Add six and numpy imports
---
Commit 9c4ce2452 authored by Mark Heffernan<meheff@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Add flag parsing to more tests in xla/service specifically those which build
HLO graphs. This enables, for example, dumping of the graphs with
--xla_generate_hlo_graph. Also remove some superfluous tensorflow test_main
dependencies.

PiperOrigin-RevId: 168406746

---
Commit d4efa695c authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Relax the feed_nodes collection check, which triggers a false positive in some modes where the feed node collection is auto-generated. Keep it as a warning to help correct user-provided feed node lists.

PiperOrigin-RevId: 168396408

---
Commit cbc46a856 authored by Changming Sun<chasun@microsoft.com>
Committed by gunan<gunan@google.com>:
Add a missing template explicit instantiation of SetZeroFunctor (#12791)

---
Commit 7bb08f5bf authored by Kevin Slagle<kjslag@gmail.com>
Committed by drpngx<drpngx@users.noreply.github.com>:
fix ExponentialMovingAverage documentation so that ExponentialMovingAverage.apply is evaluated within control_dependencies (#12987)

---
Commit e6b011763 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Extend c++ gradient_checker to complex types.

PiperOrigin-RevId: 168392949

---
Commit 4086219a4 authored by Lyndon White<oxinabox@ucc.asn.au>
Committed by drpngx<drpngx@users.noreply.github.com>:
Correct minor typo in substr docs example (#12991)

---
Commit f63aa7f49 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Migrate core TFGAN functions to opensource.

PiperOrigin-RevId: 168391923

---
Commit bc6b60f1b authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Fix tuple_losses bug caused by Python bug.

PiperOrigin-RevId: 168386341

---
Commit 7a8c63da3 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Migrate `leaky_relu` to `nn_ops.py`. Will be used for TFGAN.

PiperOrigin-RevId: 168386268

---
Commit f7ba16fdf authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Do not export from eval on train data steps.

PiperOrigin-RevId: 168374021

---
Commit 9b9e54b34 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Adding NCCL sum op, register all_sum gradient.
Streamlining nccl test.

PiperOrigin-RevId: 168347428

---
Commit bc300318e authored by Gunhan Gulsoy<gunan@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Update gemmlowp hash as the commit history seems to have changed in the
repository.

PiperOrigin-RevId: 168343607

---
Commit 1e96d54d9 authored by gunan<gunan@google.com>
Committed by GitHub<noreply@github.com>:
Also accept non-k8 CPU types in build pip package. (#12975)

* Also accept non-k8 CPU types in build pip package.
Fixes #12735

* Make the script work with `set -e`.

---
Commit c0a4c7ffc authored by Chris Leary<leary@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
[XLA] Fix bug in ShapeUtil::ShapeIs that would lead to type inference errors.

PiperOrigin-RevId: 168323589

---
Commit 4af9be964 authored by Amy<amy@infosleuth.net>
Committed by drpngx<drpngx@users.noreply.github.com>:
support passing in a source url to the mnist read_data_sets function, to make it easier to use 'fashion mnist' etc. (#12983)

---
Commit 9f848734f authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Tweak layer a bit to be eager friendly.

PiperOrigin-RevId: 168312865

---
Commit 60f15462b authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Change conv_input_scale and side_input_scale from attributes to inputs for improved flexibility, in fused_conv2d_bias_activation op.

PiperOrigin-RevId: 168311988

---
Commit 4b4e10f9c authored by Jianwei Xie<xiejw@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Adds dict support of eval metrics.

PiperOrigin-RevId: 168310444

---
Commit ab7f22de6 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Move FusedConvBiasActivationShape out of common_shape_fns.cc to a lambda inside the op.

PiperOrigin-RevId: 168300911

---
Commit 3a98035fa authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
[XLA] Augment metadata output with source-line info, as before.

PiperOrigin-RevId: 168292527

---
Commit 349188152 authored by Yao Zhang<yaozhang@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Enable fused batch norm, which is 15-20% faster for training and inference.

PiperOrigin-RevId: 168288154

---
Commit 08587d45b authored by Yuefeng Zhou<yuefengz@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Added back persistent memory tracking in queue op. The new tracking logic has avoided the crash in previous implementation:  the queue_ passed to CreateTypedQueue may be unreffed if the resource is already created by another resource op that shares the same resource name and type.

PiperOrigin-RevId: 168284509

---
Commit 733063d55 authored by Amit Patankar<amitpatankar@google.com>
Committed by Amit Patankar<amitpatankar@google.com>:
Fixing awkward wording.

---
Commit c7ad6bfef authored by Amit Patankar<amitpatankar@google.com>
Committed by Amit Patankar<amitpatankar@google.com>:
Removing accidental hash.

---
Commit 53dbc761a authored by Amit Patankar<amitpatankar@google.com>
Committed by Amit Patankar<amitpatankar@google.com>:
Adding Windows self check script to docs.

---
Commit ed1135994 authored by Andrew Harp<andrewharp@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Add -latomic flag to benchmark_model target to fix Android x86 build.

PiperOrigin-RevId: 168281337

---
Commit c0348bb55 authored by Anna R<annarev@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Update tf_export.py to take constant name as an argument instead of a constant.

PiperOrigin-RevId: 168280613

---
Commit c3d19e40a authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Cleanup training_ops to reduce code redudancy.

PiperOrigin-RevId: 168280069

---
Commit 123fb01ee authored by Yao Zhang<yaozhang@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Set fused=False for batch norm, because the test assumes no bessel's
correction. Fused=True would add bessel's correction to variance.

PiperOrigin-RevId: 168274392

---
Commit f0e8c545e authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Switch resource variables from copy-on-read to copy-on-write.

RELNOTES: Change the signature of (C++) GetInputTensorFromVariable in
training_op_helpers to support new copy-on-write semenatics of resource
variables.
PiperOrigin-RevId: 168273249

---
Commit 495cc8e47 authored by Yuan (Terry) Tang<terrytangyuan@users.noreply.github.com>
Committed by drpngx<drpngx@users.noreply.github.com>:
Minor wording change in timeseries module's README (#12938)

* Minor wording change in timeseries module's README

* Address comments

---
Commit f13b876ed authored by Amit Patankar<amitpatankar@google.com>
Committed by Amit Patankar<amitpatankar@google.com>:
Making the default build from source version 1.4.0dev. The whl files that are built will be 1.3.0devDDMMYYYY.

---
Commit 2356c0ff4 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Delete ScopedTFStatus to avoid leaking it for long running trainers(1+day).

PiperOrigin-RevId: 168259652

---
Commit e15f4cae2 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Don't remove all aliases from linalg namespace.
Get rid of redundant aliases.

PiperOrigin-RevId: 168257658

---
Commit c58082642 authored by postBG<profile2697@gmail.com>
Committed by drpngx<drpngx@users.noreply.github.com>:
Fix minor typo in Programmers guide (#12965)

* Fix minor typo in Programmers guide

* change to "this"

---
Commit 509372c2e authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Add a lot of operations' flops calculations

PiperOrigin-RevId: 168256746

---
Commit 80ed8afc0 authored by Francois Chollet<fchollet@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Add Flatten to core layers.

PiperOrigin-RevId: 168254118

---
Commit a6223c01a authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Fix locking of variables in SparseProximalGradientDescent,
AdagradDA, SparseAdagradDA.

PiperOrigin-RevId: 168252530

---
Commit abde00830 authored by Olivia Nordquist<nolivia@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
adding InputTensor class for symmetry with OutputTensor

PiperOrigin-RevId: 168250085

---
Commit 0451032ca authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
[XLA] Fix variable naming style guide violation.

PiperOrigin-RevId: 168245542

---
Commit a202a5a94 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Update ops-related pbtxt files.

PiperOrigin-RevId: 168245371

---
Commit f93e354cb authored by Derek Murray<mrry@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
[tf.contrib.data] Switch backend Dataset representation to DT_VARIANT.

This change introduces a new `DatasetWrapper` type that wraps a
`DatasetBase*` and can be stored in a DT_VARIANT tensor. All Dataset
ops now consume and produce DT_VARIANT instead of DT_RESOURCE, and the
underlying implementation is simplified because the `DatasetWrapper`
can be passed directly by value without using the `ResourceMgr`.

PiperOrigin-RevId: 168240571

---
Commit a4042cd2a authored by Jianwei Xie<xiejw@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Introduces the placeholder for _TrainingExecutor, which serves the implementation of tf.estimator.train_and_evaluate.

PiperOrigin-RevId: 168240151

---
Commit 10ba148f7 authored by Peter Hawkins<phawkins@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Switch control_flow_ops library to use Resource variants of Stack operators, instead of deprecated Ref variants.

PiperOrigin-RevId: 168234822

---
Commit ca43fe82b authored by Ali Yahya<alive@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
TFE: Improves the interfaces of tape.watch_variable() and implicit_grad().

tape.watch_variable() replaces tape.watch() and now is called on ResourceVariable objects instead of their underlying handles.

implicit_grad() now returns a list of (gradient, variable) pairs to be consistent with tf.Optimizer's interface.

PiperOrigin-RevId: 168232055

---
Commit b72862dfc authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
internal change

PiperOrigin-RevId: 168225993

---
Commit da3280f4d authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Re-enable tsan for sdca_estimator_test.

PiperOrigin-RevId: 168186374

---
Commit c936c1155 authored by Yifei Feng<yifeif@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Fix pip tests for contrib/gan.
- Add *_impl.py so tests can still access removed symbols.
- Add /python directory layer to make *_impy.py and __init__.py not in the same dir.

PiperOrigin-RevId: 168161722

---
Commit ce9a2b00f authored by Toby Boyd<tobyboyd@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Performance guide update

PiperOrigin-RevId: 168159289

---
Commit 3bce4f9a0 authored by Shanqing Cai<cais@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
TFE: expose tfe.num_gpus()

PiperOrigin-RevId: 168154345

---
Commit 67a7cbc28 authored by Jianwei Xie<xiejw@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Changed the default eval throttle secs from 2 min to 10 mins.

PiperOrigin-RevId: 168120323

---
Commit 92bed178f authored by Eugene Brevdo<ebrevdo@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Reduce cmake log mess.

* Echo off for the .bat scripts.
* TF cmake: disable warnings in some of the patched projects (gif,jpeg,lmdb).

PiperOrigin-RevId: 168119914

---
Commit 702d59582 authored by joshkyh<joshkyh@users.noreply.github.com>
Committed by Yifei Feng<fengyifei2026@gmail.com>:
Corrected hyperlink for audio training tutorial (#12923)

---
Commit 877c9deca authored by Frank Chen<frankchn@gmail.com>
Committed by Yifei Feng<fengyifei2026@gmail.com>:
Reverse change eb75ded6 so that internal tests will pass. (#12933)

As support for int64 global steps is not ready in TPUs, I am reversing this change so that our internal performance and regression tests will pass.
---
Commit 665966438 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Re-enable grpc_session_test.

PiperOrigin-RevId: 168078694

---
Commit 405def792 authored by Chris Leary<leary@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
[XLA] Switch CallInliner to use CallGraph::VisitNodes.

PiperOrigin-RevId: 168078645

---
Commit aba3466f1 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Exposes Head and factory methods in tf.contrib.estimator.

PiperOrigin-RevId: 168071246

---
Commit b76565b39 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Some profiler fixes and cleanup.

PiperOrigin-RevId: 168069346

---
Commit 32ffc5a81 authored by Jonas<sauercrowd@users.noreply.github.com>
Committed by Yifei Feng<fengyifei2026@gmail.com>:
Just a dot in order to be consistent (#12919)

added a dot to the `7` to make clear it's a float (like every other number)
---
Commit 0753b0c79 authored by Alexandre Passos<apassos@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Scope the scalar cache in the context.

PiperOrigin-RevId: 168065417

---
Commit 48deb206b authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Migrate TFGAN features to third_party.

PiperOrigin-RevId: 168060880

---
Commit d2ae1311f authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Fixing an issue in the BUILD file of the LSH ops.

PiperOrigin-RevId: 168056645

---
Commit 2f440eda4 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Expose NumpyReader for reading timeseries data.

PiperOrigin-RevId: 168055838

---
Commit be1916ce7 authored by Daniel Grazian<dgr@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Added functionality to allow `SqlDataset` to interpret a database column as various numeric types, including several integer types and `dtypes.float64`.

PiperOrigin-RevId: 168055827

---
Commit fa2000a0b authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Supporting nightly windows pip packages.

PiperOrigin-RevId: 168054959

---
Commit a263ea626 authored by Asim Shankar<ashankar@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
eager: Treat eager tensors as constants during graph construction.

Unless capturing is explicitly enabled.

PiperOrigin-RevId: 168052675

---
Commit 6e402d0d2 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Make TODO a bit more specific.

PiperOrigin-RevId: 168051381

---
Commit c779384bc authored by Daniel Grazian<dgr@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Added code example to the doc string for `SqlDataset`.

PiperOrigin-RevId: 168049037

---
Commit ff6dd474a authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Use self._in_graph_mode consistently in ResourceVariable
instead of sometimes getting it from the context.

Also: fix formatting of a comment and use a more precise test to detect
if initial_value is set.
PiperOrigin-RevId: 168047258

---
Commit f331f528b authored by Alexandre Passos<apassos@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Removes "fast paths" which are not fast in eager mode.

PiperOrigin-RevId: 168046278

---
Commit 86f1713e5 authored by Jianwei Xie<xiejw@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Introduces TrainSpec and EvalSpec.

PiperOrigin-RevId: 168040435

---
Commit c8b9e92f0 authored by Asim Shankar<ashankar@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
eager: Move "register_function" to context.py

This will allow function registration from other
modules without having to import "function.py".
(And besides, the function really does belong on the context).

PiperOrigin-RevId: 168040411

---
Commit 74137f994 authored by Shanqing Cai<cais@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Fix signed int overflow issue in tensor_id.cc

When a node name has a long numeric suffix, e.g.,
"foo/y_0/gradient_debug_09684b60f2184c67b744721915034528" (as has happened with tfdbg GradientsDebugger),

the parsing algorithm in ParseTensorName() may experience signed int overflow. Replacing the types with "unsigned int" resolves the issue.

PiperOrigin-RevId: 168039195

---
Commit 450c3b562 authored by Rohan Jain<rohanj@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Using rendezvous manager to pass args / rets between devices during function remote execution. This enables CPU->GPU remote device executions now.

PiperOrigin-RevId: 168038285

---
Commit 82cc6529f authored by Jianwei Xie<xiejw@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Fixes the wording about StopIteration.

PiperOrigin-RevId: 168034451

---
Commit fb5588002 authored by Gunhan Gulsoy<gunan@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
Add a statement on install/index.md on what os are supported.

PiperOrigin-RevId: 168032996

---
Commit f83f6b9ef authored by Chris Leary<leary@google.com>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
[XLA] Handle higher-order HLOs (e.g. While) in CallInliner and test.

PiperOrigin-RevId: 168029345

---
Commit 8988ae365 authored by A. Unique TensorFlower<gardener@tensorflow.org>
Committed by TensorFlower Gardener<gardener@tensorflow.org>:
BEGIN_PUBLIC
Automated g4 rollback of changelist 167916124

PiperOrigin-RevId: 168916710
This commit is contained in:
Patrick Nguyen 2017-09-15 18:14:40 -07:00 committed by TensorFlower Gardener
parent c48d6964bc
commit a373b1f742
104 changed files with 3005 additions and 632 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
README.md
View File

@ -45,11 +45,11 @@ GPU packages on all platforms and Windows CPU-only packages will arrive soon!
**Individual whl files**
* Linux CPU-only: [Python 2](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-cpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON2,label=cpu-slave/lastSuccessfulBuild/artifact/pip_test/whl/tensorflow-1.3.0-cp27-none-linux_x86_64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-cpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON2,label=cpu-slave)) / [Python 3.4](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-cpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON3,label=cpu-slave/lastSuccessfulBuild/artifact/pip_test/whl/tensorflow-1.3.0-cp34-cp34m-linux_x86_64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-cpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON3,label=cpu-slave/)) / [Python 3.5](https://ci.tensorflow.org/view/Nightly/job/nightly-python35-linux-cpu/lastSuccessfulBuild/artifact/pip_test/whl/tensorflow-1.3.0-cp35-cp35m-linux_x86_64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-python35-linux-cpu/))
* Linux GPU: [Python 2](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-linux-gpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON2,label=gpu-linux/lastSuccessfulBuild/artifact/pip_test/whl/tensorflow_gpu-1.3.0-cp27-none-linux_x86_64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-linux-gpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON2,label=gpu-linux/)) / [Python 3.4](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-linux-gpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON3,label=gpu-linux/lastSuccessfulBuild/artifact/pip_test/whl/tensorflow_gpu-1.3.0-cp34-cp34m-linux_x86_64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-linux-gpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON3,label=gpu-linux/)) / [Python 3.5](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-linux-gpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON3.5,label=gpu-linux/lastSuccessfulBuild/artifact/pip_test/whl/tensorflow_gpu-1.3.0-cp35-cp35m-linux_x86_64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-linux-gpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON3.5,label=gpu-linux/))
* Mac CPU-only: [Python 2](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-cpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON2,label=mac-slave/lastSuccessfulBuild/artifact/pip_test/whl/tensorflow-1.3.0-py2-none-any.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-cpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON2,label=mac-slave/)) / [Python 3](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-cpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON3,label=mac-slave/lastSuccessfulBuild/artifact/pip_test/whl/tensorflow-1.3.0-py3-none-any.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-cpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON3,label=mac-slave/))
* Windows CPU-only: [Python 3.5 64-bit](https://ci.tensorflow.org/view/Nightly/job/nightly-win/M=windows,PY=35/lastSuccessfulBuild/artifact/cmake_build/tf_python/dist/tensorflow-1.3.0-cp35-cp35m-win_amd64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-win/M=windows,PY=35/)) / [Python 3.6 64-bit](https://ci.tensorflow.org/view/Nightly/job/nightly-win/M=windows,PY=36/lastSuccessfulBuild/artifact/cmake_build/tf_python/dist/tensorflow-1.3.0-cp36-cp36m-win_amd64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-win/M=windows,PY=36/))
* Windows GPU: [Python 3.5 64-bit](https://ci.tensorflow.org/view/Nightly/job/nightly-win/M=windows-gpu,PY=35/lastSuccessfulBuild/artifact/cmake_build/tf_python/dist/tensorflow_gpu-1.3.0-cp35-cp35m-win_amd64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-win/M=windows-gpu,PY=35/)) / [Python 3.6 64-bit](https://ci.tensorflow.org/view/Nightly/job/nightly-win/M=windows-gpu,PY=36/lastSuccessfulBuild/artifact/cmake_build/tf_python/dist/tensorflow_gpu-1.3.0-cp36-cp36m-win_amd64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-win/M=windows-gpu,PY=36/))
* Linux CPU-only: [Python 2](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-cpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON2,label=cpu-slave/lastSuccessfulBuild/artifact/pip_test/whl/tensorflow-1.4.0dev-cp27-none-linux_x86_64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-cpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON2,label=cpu-slave)) / [Python 3.4](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-cpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON3,label=cpu-slave/lastSuccessfulBuild/artifact/pip_test/whl/tensorflow-1.4.0dev-cp34-cp34m-linux_x86_64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-cpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON3,label=cpu-slave/)) / [Python 3.5](https://ci.tensorflow.org/view/Nightly/job/nightly-python35-linux-cpu/lastSuccessfulBuild/artifact/pip_test/whl/tensorflow-1.4.0dev-cp35-cp35m-linux_x86_64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-python35-linux-cpu/))
* Linux GPU: [Python 2](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-linux-gpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON2,label=gpu-linux/lastSuccessfulBuild/artifact/pip_test/whl/tensorflow_gpu-1.4.0dev-cp27-none-linux_x86_64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-linux-gpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON2,label=gpu-linux/)) / [Python 3.4](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-linux-gpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON3,label=gpu-linux/lastSuccessfulBuild/artifact/pip_test/whl/tensorflow_gpu-1.4.0dev-cp34-cp34m-linux_x86_64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-linux-gpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON3,label=gpu-linux/)) / [Python 3.5](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-linux-gpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON3.5,label=gpu-linux/lastSuccessfulBuild/artifact/pip_test/whl/tensorflow_gpu-1.4.0dev-cp35-cp35m-linux_x86_64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-linux-gpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON3.5,label=gpu-linux/))
* Mac CPU-only: [Python 2](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-cpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON2,label=mac-slave/lastSuccessfulBuild/artifact/pip_test/whl/tensorflow-1.4.0dev-py2-none-any.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-cpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON2,label=mac-slave/)) / [Python 3](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-cpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON3,label=mac-slave/lastSuccessfulBuild/artifact/pip_test/whl/tensorflow-1.4.0dev-py3-none-any.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-matrix-cpu/TF_BUILD_IS_OPT=OPT,TF_BUILD_IS_PIP=PIP,TF_BUILD_PYTHON_VERSION=PYTHON3,label=mac-slave/))
* Windows CPU-only: [Python 3.5 64-bit](https://ci.tensorflow.org/view/Nightly/job/nightly-win/M=windows,PY=35/lastSuccessfulBuild/artifact/cmake_build/tf_python/dist/tensorflow-1.4.0dev-cp35-cp35m-win_amd64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-win/M=windows,PY=35/)) / [Python 3.6 64-bit](https://ci.tensorflow.org/view/Nightly/job/nightly-win/M=windows,PY=36/lastSuccessfulBuild/artifact/cmake_build/tf_python/dist/tensorflow-1.4.0dev-cp36-cp36m-win_amd64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-win/M=windows,PY=36/))
* Windows GPU: [Python 3.5 64-bit](https://ci.tensorflow.org/view/Nightly/job/nightly-win/M=windows-gpu,PY=35/lastSuccessfulBuild/artifact/cmake_build/tf_python/dist/tensorflow_gpu-1.4.0dev-cp35-cp35m-win_amd64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-win/M=windows-gpu,PY=35/)) / [Python 3.6 64-bit](https://ci.tensorflow.org/view/Nightly/job/nightly-win/M=windows-gpu,PY=36/lastSuccessfulBuild/artifact/cmake_build/tf_python/dist/tensorflow_gpu-1.4.0dev-cp36-cp36m-win_amd64.whl) ([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-win/M=windows-gpu,PY=36/))
* Android: [demo APK](https://ci.tensorflow.org/view/Nightly/job/nightly-android/lastSuccessfulBuild/artifact/out/tensorflow_demo.apk), [native libs](http://ci.tensorflow.org/view/Nightly/job/nightly-android/lastSuccessfulBuild/artifact/out/native/)
([build history](https://ci.tensorflow.org/view/Nightly/job/nightly-android/))

1
tensorflow/BUILD
View File

@ -472,6 +472,7 @@ cc_binary(
"//tensorflow:darwin": [
"-Wl,-exported_symbols_list", # This line must be directly followed by the exported_symbols.lds file
"//tensorflow/c:exported_symbols.lds",
"-Wl,-install_name,@rpath/libtensorflow.so",
],
"//tensorflow:windows": [],
"//tensorflow:windows_msvc": [],

4
tensorflow/c/version_script.lds
View File

@ -1,8 +1,8 @@
VERS_1.0 {
# Export symbols in c_api.h.
global:
TF_*;
TFE_*;
*TF_*;
*TFE_*;
# Hide everything else.
local:

1
tensorflow/cc/BUILD
View File

@ -336,6 +336,7 @@ cc_library(
":cc_ops",
":cc_ops_internal",
":grad_op_registry",
":gradients",
],
alwayslink = 1,
)

12
tensorflow/cc/gradients/math_grad.cc
View File

@ -696,6 +696,18 @@ Status MeanGrad(const Scope& scope, const Operation& op,
}
REGISTER_GRADIENT_OP("Mean", MeanGrad);
Status LgammaGrad(const Scope& scope, const Operation& op,
const std::vector<Output>& grad_inputs,
std::vector<Output>* grad_outputs) {
auto grad = grad_inputs[0];
Scope grad_scope = scope.WithControlDependencies(grad);
auto x = ConjugateHelper(grad_scope, op.input(0));
auto dx = Mul(scope, grad, Digamma(scope, x));
grad_outputs->push_back(dx);
return scope.status();
}
REGISTER_GRADIENT_OP("Lgamma", LgammaGrad);
Status MinOrMaxGrad(const Scope& scope, const Operation& op,
const std::vector<Output>& grad_inputs,
std::vector<Output>* grad_outputs) {

12
tensorflow/cc/gradients/math_grad_test.cc
View File

@ -821,5 +821,17 @@ TEST_F(NaryGradTest, Minimum) {
RunTest(x, x_init_value, y, shape);
}
TEST_F(NaryGradTest, Lgamma) {
TensorShape shape({3, 2});
auto x = Placeholder(scope_, DT_FLOAT, Placeholder::Shape(shape));
auto y = Lgamma(scope_, x);
// Select values to avoid instability when computing finite differences.
// Ref: https://en.wikipedia.org/wiki/File:Gamma_plot.svg
Tensor x_init_value =
test::AsTensor<float>({-3.5f, -2.5f, -1.5f, 1.0f, 2.0f, 3.5f}, {3, 2});
RunTest(x, x_init_value, y, shape);
// TODO(suharshs): add test case for complex values
}
} // namespace
} // namespace tensorflow

82
tensorflow/cc/gradients/nn_grad.cc
View File

@ -18,6 +18,7 @@ limitations under the License.
#include "tensorflow/cc/ops/standard_ops.h"
#include "tensorflow/cc/framework/grad_op_registry.h"
#include "tensorflow/cc/framework/gradients.h"
namespace tensorflow {
namespace ops {
@ -118,6 +119,87 @@ Status BiasAddGradHelper(const Scope& scope, const Operation& op,
}
REGISTER_GRADIENT_OP("BiasAdd", BiasAddGradHelper);
Status Conv2DGrad(const Scope& scope, const Operation& op,
const std::vector<Output>& grad_inputs,
std::vector<Output>* grad_outputs) {
string data_format;
string padding;
std::vector<int32> strides;
bool use_cudnn_on_gpu;
auto attrs = op.output(0).node()->attrs();
TF_RETURN_IF_ERROR(GetNodeAttr(attrs, "data_format", &data_format));
TF_RETURN_IF_ERROR(GetNodeAttr(attrs, "padding", &padding));
TF_RETURN_IF_ERROR(GetNodeAttr(attrs, "strides", &strides));
TF_RETURN_IF_ERROR(GetNodeAttr(attrs, "use_cudnn_on_gpu",
&use_cudnn_on_gpu));
Conv2DBackpropInput::Attrs input_attrs;
input_attrs.DataFormat(data_format);
input_attrs.UseCudnnOnGpu(use_cudnn_on_gpu);
auto dx_1 = Conv2DBackpropInput(scope, Shape(scope, op.input(0)),
op.input(1), grad_inputs[0],
strides, padding, input_attrs);
grad_outputs->push_back(dx_1);
Conv2DBackpropFilter::Attrs filter_attrs;
filter_attrs.DataFormat(data_format);
filter_attrs.UseCudnnOnGpu(use_cudnn_on_gpu);
auto dx_2 = Conv2DBackpropFilter(scope, op.input(0),
Shape(scope, op.input(1)), grad_inputs[0],
strides, padding, filter_attrs);
grad_outputs->push_back(dx_2);
return scope.status();
}
REGISTER_GRADIENT_OP("Conv2D", Conv2DGrad);
Status MaxPoolGradHelper(const Scope& scope, const Operation& op,
const std::vector<Output>& grad_inputs,
std::vector<Output>* grad_outputs) {
string data_format;
string padding;
std::vector<int32> strides;
std::vector<int32> ksize;
auto attrs = op.output(0).node()->attrs();
TF_RETURN_IF_ERROR(GetNodeAttr(attrs, "data_format", &data_format));
TF_RETURN_IF_ERROR(GetNodeAttr(attrs, "ksize", &ksize));
TF_RETURN_IF_ERROR(GetNodeAttr(attrs, "padding", &padding));
TF_RETURN_IF_ERROR(GetNodeAttr(attrs, "strides", &strides));
internal::MaxPoolGrad::Attrs grad_attrs;
grad_attrs.DataFormat(data_format);
auto dx = internal::MaxPoolGrad(scope, op.input(0),
op.output(0),
grad_inputs[0],
ksize, strides,
padding, grad_attrs);
grad_outputs->push_back(dx);
return scope.status();
}
REGISTER_GRADIENT_OP("MaxPool", MaxPoolGradHelper);
Status MaxPoolGradV2Helper(const Scope& scope, const Operation& op,
const std::vector<Output>& grad_inputs,
std::vector<Output>* grad_outputs) {
string data_format;
string padding;
auto attrs = op.output(0).node()->attrs();
TF_RETURN_IF_ERROR(GetNodeAttr(attrs, "data_format", &data_format));
TF_RETURN_IF_ERROR(GetNodeAttr(attrs, "padding", &padding));
MaxPoolGradV2::Attrs grad_attrs;
grad_attrs.DataFormat(data_format);
auto dx = MaxPoolGradV2(scope, op.input(0),
op.output(0),
grad_inputs[0],
op.input(1),
op.input(2),
padding,
grad_attrs);
grad_outputs->push_back(dx);
grad_outputs->push_back(NoGradient());
grad_outputs->push_back(NoGradient());
return scope.status();
}
REGISTER_GRADIENT_OP("MaxPoolV2", MaxPoolGradV2Helper);
} // anonymous namespace
} // namespace ops
} // namespace tensorflow

27
tensorflow/cc/gradients/nn_grad_test.cc
View File

@ -138,5 +138,32 @@ TEST_F(NNGradTest, BiasAddGradHelper) {
RunTest({x, bias}, {shape, bias_shape}, {y}, {shape});
}
TEST_F(NNGradTest, Conv2DGrad) {
TensorShape shape({1, 2, 2, 1});
auto x = Placeholder(scope_, DT_FLOAT, Placeholder::Shape(shape));
Tensor filter = test::AsTensor<float>({0.5f}, {1, 1, 1, 1});
const std::vector<int> strides{1, 1, 1, 1};
auto y = Conv2D(scope_, x, filter, strides, "SAME");
RunTest(x, shape, y, shape);
}
TEST_F(NNGradTest, MaxPoolGradHelper) {
TensorShape shape({1, 2, 2, 1});
auto x = Placeholder(scope_, DT_FLOAT, Placeholder::Shape(shape));
const std::vector<int> ksize{1, 2, 2, 1};
const std::vector<int> strides{1, 1, 1, 1};
auto y = MaxPool(scope_, x, ksize, strides, "SAME");
RunTest(x, shape, y, shape);
}
TEST_F(NNGradTest, MaxPoolGradV2Helper) {
TensorShape shape({1, 2, 2, 1});
auto x = Placeholder(scope_, DT_FLOAT, Placeholder::Shape(shape));
Tensor ksize = test::AsTensor<int>({1, 2, 2, 1}, {4});
Tensor strides = test::AsTensor<int>({1, 1, 1, 1}, {4});
auto y = MaxPoolV2(scope_, x, ksize, strides, "SAME");
RunTest(x, shape, y, shape);
}
} // namespace
} // namespace tensorflow

2
tensorflow/compiler/xla/service/gpu/ir_emitter_unnested.cc
View File

@ -999,7 +999,7 @@ Status IrEmitterUnnested::EmitRowReduction(
// for (shuffle_distance = 16; shuffle_distance > 0; shuffle_distance /= 2)
// partial_result = Reducer(
// partial_result,
// __shfl_down(partial_result, shuffle_distance));
// __shfl_down_sync(CUDA_WARP_ALL, partial_result, shuffle_distance));
// if (lane_id == 0)
// AtomicReducer(&output[y], partial_result);
// }

24
tensorflow/compiler/xla/service/gpu/llvm_gpu_backend/gpu_backend_lib.cc
View File

@ -71,7 +71,18 @@ const int kDefaultInlineThreshold = 1100;
// Gets the libdevice filename for a particular compute capability. When
// presented with a GPU we don't recognize, we just return the libdevice from
// compute_20.
static string GetLibdeviceFilename(std::pair<int, int> compute_capability) {
static string GetLibdeviceFilename(const string& libdevice_dir_path,
std::pair<int, int> compute_capability) {
// Since CUDA 9.0, all GPU versions are included in a single file
const char* unified_libdevice_filename = "libdevice.10.bc";
std::vector<string> unified_libdevice_files;
const tensorflow::Status status =
tensorflow::Env::Default()->GetMatchingPaths(
tensorflow::io::JoinPath(libdevice_dir_path, unified_libdevice_filename),
&unified_libdevice_files);
if (status.ok() && unified_libdevice_files.size() == 1) {
return unified_libdevice_filename;
}
// There are only four libdevice files: compute_{20,30,35,50}. Each GPU
// version gets mapped to one of these. Note in particular that sm_60 and
// sm_61 map to libdevice.compute_30.
@ -101,7 +112,7 @@ static string GetLibdeviceFilename(std::pair<int, int> compute_capability) {
}
// Gets the GPU name as it's known to LLVM for a given compute capability. If
// we see an unrecognized compute capability, we return "sm_20".
// we see an unrecognized compute capability, we return "sm_30".
static string GetSmName(std::pair<int, int> compute_capability) {
static auto* m = new std::map<std::pair<int, int>, int>({{{2, 0}, 20},
{{2, 1}, 21},
@ -114,8 +125,10 @@ static string GetSmName(std::pair<int, int> compute_capability) {
{{5, 3}, 53},
{{6, 0}, 60},
{{6, 1}, 61},
{{6, 2}, 62}});
int sm_version = 20;
{{6, 2}, 62},
// TODO: Change this to 70 once LLVM NVPTX supports it
{{7, 0}, 60}});
int sm_version = 30;
auto it = m->find(compute_capability);
if (it != m->end()) {
sm_version = it->second;
@ -306,7 +319,8 @@ tensorflow::Status LinkLibdeviceIfNecessary(
llvm::Linker linker(*module);
string libdevice_path = tensorflow::io::JoinPath(
libdevice_dir_path, GetLibdeviceFilename(compute_capability));
libdevice_dir_path, GetLibdeviceFilename(libdevice_dir_path,
compute_capability));
TF_RETURN_IF_ERROR(tensorflow::Env::Default()->FileExists(libdevice_path));
VLOG(1) << "Linking with libdevice from: " << libdevice_path;
std::unique_ptr<llvm::Module> libdevice_module =

9
tensorflow/compiler/xla/service/hlo_verifier.cc
View File

@ -531,6 +531,15 @@ StatusOr<bool> HloVerifier::Run(HloModule* module) {
<< " computation: " << computation.get();
}
}
if (instruction->opcode() == HloOpcode::kBroadcast) {
// If you see this failure then someone has confused the difference
// between the HLO broadcast op, and the UserComputation broadcast
// op. See https://groups.google.com/forum/#!topic/xla-dev/9LqijHmTt_I
// or ComputationLowerer::Visit()
TF_RET_CHECK(instruction->dimensions().size() ==
ShapeUtil::Rank(instruction->operand(0)->shape()))
<< "Broadcast HLO has invalid number of dimensions.";
}
auto previous = instructions.find(instruction->name());
TF_RET_CHECK(previous == instructions.end())

81
tensorflow/contrib/android/java/org/tensorflow/contrib/android/TensorFlowInferenceInterface.java
View File

@ -20,6 +20,7 @@ import android.os.Build.VERSION;
import android.os.Trace;
import android.text.TextUtils;
import android.util.Log;
import java.io.ByteArrayOutputStream;
import java.io.FileInputStream;
import java.io.IOException;
import java.io.InputStream;
@ -78,10 +79,35 @@ public class TensorFlowInferenceInterface {
throw new RuntimeException("Failed to load model from '" + model + "'", e);
}
}
try {
loadGraph(is, g);
if (VERSION.SDK_INT >= 18) {
Trace.beginSection("initializeTensorFlow");
Trace.beginSection("readGraphDef");
}
// TODO(ashankar): Can we somehow mmap the contents instead of copying them?
byte[] graphDef = new byte[is.available()];
final int numBytesRead = is.read(graphDef);
if (numBytesRead != graphDef.length) {
throw new IOException(
"read error: read only "
+ numBytesRead
+ " of the graph, expected to read "
+ graphDef.length);
}
if (VERSION.SDK_INT >= 18) {
Trace.endSection(); // readGraphDef.
}
loadGraph(graphDef, g);
is.close();
Log.i(TAG, "Successfully loaded model from '" + model + "'");
if (VERSION.SDK_INT >= 18) {
Trace.endSection(); // initializeTensorFlow.
}
} catch (IOException e) {
throw new RuntimeException("Failed to load model from '" + model + "'", e);
}
@ -105,8 +131,30 @@ public class TensorFlowInferenceInterface {
this.runner = sess.runner();
try {
loadGraph(is, g);
if (VERSION.SDK_INT >= 18) {
Trace.beginSection("initializeTensorFlow");
Trace.beginSection("readGraphDef");
}
int baosInitSize = is.available() > 16384 ? is.available() : 16384;
ByteArrayOutputStream baos = new ByteArrayOutputStream(baosInitSize);
int numBytesRead;
byte[] buf = new byte[16384];
while ((numBytesRead = is.read(buf, 0, buf.length)) != -1) {
baos.write(buf, 0, numBytesRead);
}
byte[] graphDef = baos.toByteArray();
if (VERSION.SDK_INT >= 18) {
Trace.endSection(); // readGraphDef.
}
loadGraph(graphDef, g);
Log.i(TAG, "Successfully loaded model from the input stream");
if (VERSION.SDK_INT >= 18) {
Trace.endSection(); // initializeTensorFlow.
}
} catch (IOException e) {
throw new RuntimeException("Failed to load model from the input stream", e);
}
@ -269,8 +317,8 @@ public class TensorFlowInferenceInterface {
/**
* Copy a byte sequence into the input Tensor with name {@link inputName} as a string-valued
* scalar tensor. In the TensorFlow type system, a "string" is an arbitrary sequence of
* bytes, not a Java {@code String} (which is a sequence of characters).
* scalar tensor. In the TensorFlow type system, a "string" is an arbitrary sequence of bytes, not
* a Java {@code String} (which is a sequence of characters).
*/
public void feedString(String inputName, byte[] src) {
addFeed(inputName, Tensor.create(src));
@ -278,9 +326,8 @@ public class TensorFlowInferenceInterface {
/**
* Copy an array of byte sequences into the input Tensor with name {@link inputName} as a
* string-valued one-dimensional tensor (vector). In the TensorFlow type system, a "string"
* is an arbitrary sequence of bytes, not a Java {@code String} (which is a sequence of
* characters).
* string-valued one-dimensional tensor (vector). In the TensorFlow type system, a "string" is an
* arbitrary sequence of bytes, not a Java {@code String} (which is a sequence of characters).
*/
public void feedString(String inputName, byte[][] src) {
addFeed(inputName, Tensor.create(src));
@ -458,27 +505,10 @@ public class TensorFlowInferenceInterface {
}
}
private void loadGraph(InputStream is, Graph g) throws IOException {
private void loadGraph(byte[] graphDef, Graph g) throws IOException {
final long startMs = System.currentTimeMillis();
if (VERSION.SDK_INT >= 18) {
Trace.beginSection("loadGraph");
Trace.beginSection("readGraphDef");
}
// TODO(ashankar): Can we somehow mmap the contents instead of copying them?
byte[] graphDef = new byte[is.available()];
final int numBytesRead = is.read(graphDef);
if (numBytesRead != graphDef.length) {
throw new IOException(
"read error: read only "
+ numBytesRead
+ " of the graph, expected to read "
+ graphDef.length);
}
if (VERSION.SDK_INT >= 18) {
Trace.endSection(); // readGraphDef.
Trace.beginSection("importGraphDef");
}
@ -490,7 +520,6 @@ public class TensorFlowInferenceInterface {
if (VERSION.SDK_INT >= 18) {
Trace.endSection(); // importGraphDef.
Trace.endSection(); // loadGraph.
}
final long endMs = System.currentTimeMillis();

2
tensorflow/contrib/estimator/python/estimator/extenders.py
View File

@ -54,7 +54,7 @@ def add_metrics(estimator, metric_fn):
```
Args:
estimator: A ${tf.estimator.Esitmator} object.
estimator: A ${tf.estimator.Estimator} object.
metric_fn: A function which should obey the following signature:
- Args: can only have following four arguments in any order:
* predictions: Predictions `Tensor` or dict of `Tensor` created by given

1
tensorflow/contrib/layers/__init__.py
View File

@ -51,6 +51,7 @@ See the @{$python/contrib.layers} guide.
@@unit_norm
@@bow_encoder
@@embed_sequence
@@maxout
@@apply_regularization
@@l1_l2_regularizer

5
tensorflow/contrib/layers/python/layers/feature_column.py
View File

@ -939,6 +939,11 @@ class _OneHotColumn(
weighted_column = sparse_ops.sparse_merge(sp_ids=sparse_id_column,
sp_values=weight_tensor,
vocab_size=self.length)
# Remove (?, -1) index
weighted_column = sparse_ops.sparse_slice(
weighted_column,
[0, 0],
weighted_column.dense_shape)
return sparse_ops.sparse_tensor_to_dense(weighted_column)
dense_id_tensor = sparse_ops.sparse_tensor_to_dense(sparse_id_column,

30
tensorflow/contrib/layers/python/layers/feature_column_test.py
View File

@ -31,6 +31,7 @@ from tensorflow.python.feature_column import feature_column as fc_core
from tensorflow.python.framework import constant_op
from tensorflow.python.framework import dtypes
from tensorflow.python.framework import sparse_tensor as sparse_tensor_lib
from tensorflow.python.ops import lookup_ops
from tensorflow.python.ops import parsing_ops
from tensorflow.python.ops import state_ops
from tensorflow.python.ops import variable_scope
@ -319,6 +320,35 @@ class FeatureColumnTest(test.TestCase):
self.assertEqual(one_hot.sparse_id_column.name, "ids_weighted_by_weights")
self.assertEqual(one_hot.length, 3)
def testMissingValueInOneHotColumnForWeightedSparseColumn(self):
# Github issue 12583
ids = fc.sparse_column_with_keys("ids", ["marlo", "omar", "stringer"])
weighted_ids = fc.weighted_sparse_column(ids, "weights")
one_hot = fc.one_hot_column(weighted_ids)
features = {
'ids': constant_op.constant([['marlo', 'unknown', 'omar']]),
'weights': constant_op.constant([[2., 4., 6.]])
}
one_hot_tensor = feature_column_ops.input_from_feature_columns(
features, [one_hot])
with self.test_session() as sess:
sess.run(variables.global_variables_initializer())
sess.run(lookup_ops.tables_initializer())
self.assertAllEqual([[2., 6., 0.]], one_hot_tensor.eval())
def testMissingValueInOneHotColumnForSparseColumnWithKeys(self):
ids = fc.sparse_column_with_keys("ids", ["marlo", "omar", "stringer"])
one_hot = fc.one_hot_column(ids)
features = {
'ids': constant_op.constant([['marlo', 'unknown', 'omar']])
}
one_hot_tensor = feature_column_ops.input_from_feature_columns(
features, [one_hot])
with self.test_session() as sess:
sess.run(variables.global_variables_initializer())
sess.run(lookup_ops.tables_initializer())
self.assertAllEqual([[1., 1., 0.]], one_hot_tensor.eval())
def testOneHotColumnDeepCopy(self):
a = fc.sparse_column_with_keys("a", ["a", "b", "c", "d"])
column = fc.one_hot_column(a)

7
tensorflow/contrib/layers/python/layers/layers.py
View File

@ -50,6 +50,7 @@ from tensorflow.python.ops import standard_ops
from tensorflow.python.ops import variable_scope
from tensorflow.python.ops import variables as tf_variables
from tensorflow.python.training import moving_averages
from tensorflow.python.layers.maxout import maxout
# TODO(b/28426988): Replace legacy_* fns migrated from slim.
# TODO(b/28426988): Remove legacy_* when all uses have migrated to new API.
@ -92,7 +93,8 @@ __all__ = ['avg_pool2d',
'unit_norm',
'legacy_fully_connected',
'legacy_linear',
'legacy_relu']
'legacy_relu',
'maxout']
DATA_FORMAT_NCHW = 'NCHW'
DATA_FORMAT_NHWC = 'NHWC'
@ -811,7 +813,8 @@ def batch_norm(inputs,
if data_format == DATA_FORMAT_NCHW:
mean = array_ops.reshape(mean, params_shape_broadcast)
variance = array_ops.reshape(variance, params_shape_broadcast)
beta = array_ops.reshape(beta, params_shape_broadcast)
if beta is not None:
beta = array_ops.reshape(beta, params_shape_broadcast)
if gamma is not None:
gamma = array_ops.reshape(gamma, params_shape_broadcast)

7
tensorflow/contrib/layers/python/layers/layers_test.py
View File

@ -2636,6 +2636,13 @@ class BatchNormTest(test.TestCase):
data_format='NCHW', shape=shape, is_training=True)
self.assertAllClose(nhwc, nchw, atol=1e-4, rtol=1e-4)
def testBatchNormBeta(self):
# Test case for 11673
with self.test_session() as sess:
a = array_ops.placeholder(dtypes.float32, shape=(10, 10, 10, 10))
b = _layers.batch_norm(a, center=False, data_format='NCHW',
zero_debias_moving_mean=True)
sess.run(variables_lib.global_variables_initializer())
class LayerNormTest(test.TestCase):

22
tensorflow/contrib/learn/python/learn/datasets/mnist.py
View File

@ -30,7 +30,7 @@ from tensorflow.python.framework import random_seed
from tensorflow.python.platform import gfile
# CVDF mirror of http://yann.lecun.com/exdb/mnist/
SOURCE_URL = 'https://storage.googleapis.com/cvdf-datasets/mnist/'
DEFAULT_SOURCE_URL = 'https://storage.googleapis.com/cvdf-datasets/mnist/'
def _read32(bytestream):
@ -215,7 +215,8 @@ def read_data_sets(train_dir,
dtype=dtypes.float32,
reshape=True,
validation_size=5000,
seed=None):
seed=None,
source_url=DEFAULT_SOURCE_URL):
if fake_data:
def fake():
@ -227,28 +228,31 @@ def read_data_sets(train_dir,
test = fake()
return base.Datasets(train=train, validation=validation, test=test)
if not source_url: # empty string check
source_url = DEFAULT_SOURCE_URL
TRAIN_IMAGES = 'train-images-idx3-ubyte.gz'
TRAIN_LABELS = 'train-labels-idx1-ubyte.gz'
TEST_IMAGES = 't10k-images-idx3-ubyte.gz'
TEST_LABELS = 't10k-labels-idx1-ubyte.gz'
local_file = base.maybe_download(TRAIN_IMAGES, train_dir,
SOURCE_URL + TRAIN_IMAGES)
source_url + TRAIN_IMAGES)
with gfile.Open(local_file, 'rb') as f:
train_images = extract_images(f)
local_file = base.maybe_download(TRAIN_LABELS, train_dir,
SOURCE_URL + TRAIN_LABELS)
source_url + TRAIN_LABELS)
with gfile.Open(local_file, 'rb') as f:
train_labels = extract_labels(f, one_hot=one_hot)
local_file = base.maybe_download(TEST_IMAGES, train_dir,
SOURCE_URL + TEST_IMAGES)
source_url + TEST_IMAGES)
with gfile.Open(local_file, 'rb') as f:
test_images = extract_images(f)
local_file = base.maybe_download(TEST_LABELS, train_dir,
SOURCE_URL + TEST_LABELS)
source_url + TEST_LABELS)
with gfile.Open(local_file, 'rb') as f:
test_labels = extract_labels(f, one_hot=one_hot)
@ -262,13 +266,13 @@ def read_data_sets(train_dir,
train_images = train_images[validation_size:]
train_labels = train_labels[validation_size:]
options = dict(dtype=dtype, reshape=reshape, seed=seed)
train = DataSet(train_images, train_labels, **options)
validation = DataSet(validation_images, validation_labels, **options)
test = DataSet(test_images, test_labels, **options)
return base.Datasets(train=train, validation=validation, test=test)

7
tensorflow/contrib/makefile/compile_nsync.sh
View File

@ -214,12 +214,12 @@ for arch in $archs; do
armeabi-v7a) toolchain="arm-linux-androideabi-4.9"
sysroot_arch="arm"
bin_prefix="arm-linux-androideabi"
march_option="-march=armv7-a"
march_option="-march=armv7-a -mfloat-abi=softfp -mfpu=neon"
;;
armeabi-v7a-hard) toolchain="arm-linux-androideabi-4.9"
sysroot_arch="arm"
bin_prefix="arm-linux-androideabi"
march_option="-march=armv7-a"
march_option="-march=armv7-a -mfpu=neon"
;;
mips) toolchain="mipsel-linux-android-4.9"
sysroot_arch="mips"
@ -265,8 +265,7 @@ for arch in $archs; do
-I$(NDK_ROOT)/sources/cxx-stl/gnu-libstdc++/4.9/libs/'"$arch"'/include \
-I../../platform/c++11 -I../../platform/gcc \
-I../../platform/posix -pthread
PLATFORM_CFLAGS=-std=c++11 -Wno-narrowing '"$march_option"' \
-mfloat-abi=softfp -mfpu=neon -fPIE
PLATFORM_CFLAGS=-std=c++11 -Wno-narrowing '"$march_option"' -fPIE
PLATFORM_LDFLAGS=-pthread
MKDEP=${CC} -M -std=c++11
PLATFORM_C=../../platform/c++11/src/nsync_semaphore_mutex.cc \

4
tensorflow/contrib/makefile/sub_makefiles/android/Makefile.in
View File

@ -52,7 +52,9 @@ $(INFERENCE_SO_PATH): $(LIB_OBJS) $(INFERENCE_OBJS)
@mkdir -p $(dir $@)
$(CXX) $(CXXFLAGS) $(INCLUDES) \
-o $@ $(INFERENCE_OBJS) $(LIB_OBJS) \
$(LIBFLAGS) $(LDFLAGS) -shared $(LIBS)
$(LIBFLAGS) $(LDFLAGS) \
-shared -Wl,-soname,$(INFERENCE_SO_NAME) \
$(LIBS)
$(INFERENCE_SO_NAME): $(INFERENCE_SO_PATH)

9
tensorflow/contrib/metrics/python/ops/metric_ops.py
View File

@ -34,6 +34,7 @@ from tensorflow.python.ops import metrics_impl
from tensorflow.python.ops import nn
from tensorflow.python.ops import state_ops
from tensorflow.python.ops import variable_scope
from tensorflow.python.ops import weights_broadcast_ops
from tensorflow.python.util.deprecation import deprecated
@ -651,7 +652,7 @@ def _streaming_confusion_matrix_at_thresholds(
label_is_neg = math_ops.logical_not(label_is_pos)
if weights is not None:
broadcast_weights = _broadcast_weights(
broadcast_weights = weights_broadcast_ops.broadcast_weights(
math_ops.to_float(weights), predictions)
weights_tiled = array_ops.tile(array_ops.reshape(
broadcast_weights, [1, -1]), [num_thresholds, 1])
@ -955,7 +956,7 @@ def streaming_specificity_at_sensitivity(
def streaming_sensitivity_at_specificity(
predictions, labels, specificity, weights=None, num_thresholds=200,
metrics_collections=None, updates_collections=None, name=None):
"""Computes the specificity at a given sensitivity.
"""Computes the sensitivity at a given specificity.
The `streaming_sensitivity_at_specificity` function creates four local
variables, `true_positives`, `true_negatives`, `false_positives` and
@ -1924,7 +1925,7 @@ def streaming_covariance(predictions,
weighted_predictions = predictions
weighted_labels = labels
else:
weights = _broadcast_weights(weights, labels)
weights = weights_broadcast_ops.broadcast_weights(weights, labels)
batch_count = math_ops.reduce_sum(weights) # n_B in eqn
weighted_predictions = math_ops.multiply(predictions, weights)
weighted_labels = math_ops.multiply(labels, weights)
@ -2051,7 +2052,7 @@ def streaming_pearson_correlation(predictions,
# Broadcast weights here to avoid duplicate broadcasting in each call to
# `streaming_covariance`.
if weights is not None:
weights = _broadcast_weights(weights, labels)
weights = weights_broadcast_ops.broadcast_weights(weights, labels)
cov, update_cov = streaming_covariance(
predictions, labels, weights=weights, name='covariance')
var_predictions, update_var_predictions = streaming_covariance(

5
tensorflow/contrib/slim/python/slim/data/tfexample_decoder_test.py
View File

@ -228,7 +228,10 @@ class TFExampleDecoderTest(test.TestCase):
image_shape = (2, 3, 3)
unused_image, serialized_example = self.GenerateImage(
image_format='jpeg', image_shape=image_shape)
with self.assertRaises(TypeError):
# decode_raw support uint16 now so ValueError will be thrown instead.
with self.assertRaisesRegexp(
ValueError,
'true_fn and false_fn must have the same type: uint16, uint8'):
unused_decoded_image = self.RunDecodeExample(
serialized_example,
tfexample_decoder.Image(dtype=dtypes.uint16),

2
tensorflow/contrib/timeseries/README.md
View File

@ -2,7 +2,7 @@
TensorFlow Time Series (TFTS) is a collection of ready-to-use classic models
(state space, autoregressive), and flexible infrastructure for building
high-performance time series models whatever the architecture. It includes tools
high-performance time series models with custom architectures. It includes tools
for chunking and batching a series, and for saving model state across chunks,
making use of parallel computation even when training sequential models on long
series (using truncated backpropagation).

254
tensorflow/contrib/verbs/rdma.cc
View File

@ -165,9 +165,10 @@ void RdmaAdapter::Process_CQ() {
RdmaBuffer* ab = rc->tx_ack_buffer_;
ab->SendNextItem();
// find buffer
RdmaBuffer* tb = rc->FindBuffer(rm.name_);
RdmaTensorBuffer* tb =
reinterpret_cast<RdmaTensorBuffer*>(rc->FindBuffer(rm.name_));
tb->SetBufferStatus(remote, idle);
worker_env_->compute_pool->Schedule([tb]() { tb->SendNextItem(); });
worker_env_->compute_pool->Schedule([tb]() { tb->ReSendNextItem(); });
} else if (rm.type_ == RDMA_MESSAGE_BUFFER_REQUEST) {
// remote host requests to create a tensor buffer;
// send ack to release remote tx message buffer
@ -198,7 +199,8 @@ void RdmaAdapter::Process_CQ() {
RdmaBuffer* ab = rc->tx_ack_buffer_;
ab->SendNextItem();
// find buffer
RdmaBuffer* tb = rc->FindBuffer(rm.name_);
RdmaTensorBuffer* tb =
reinterpret_cast<RdmaTensorBuffer*>(rc->FindBuffer(rm.name_));
CHECK(rm.buffer_size_ == tb->size_)
<< "rm.buffer_size = " << rm.buffer_size_
<< "tb->size_ = " << tb->size_ << "rm.name_ = " << rm.name_;
@ -208,7 +210,7 @@ void RdmaAdapter::Process_CQ() {
tb->SetRemoteMR(rmr, true);
tb->SetBufferStatus(local, idle);
tb->SetBufferStatus(remote, idle);
worker_env_->compute_pool->Schedule([tb]() { tb->SendNextItem(); });
worker_env_->compute_pool->Schedule([tb]() { tb->ReSendNextItem(); });
} else if (rm.type_ == RDMA_MESSAGE_TENSOR_WRITE) {
// tensor RDMA write completed
worker_env_->compute_pool->Schedule([rm, rc]() {
@ -624,6 +626,12 @@ RdmaMessageBuffer::RdmaMessageBuffer(RdmaChannel* channel, string name)
RdmaTensorBuffer::RdmaTensorBuffer(RdmaChannel* channel, string name)
: RdmaBuffer(channel, name) {}
RdmaTensorBuffer::~RdmaTensorBuffer() {
for (Itable it = retable.begin(); it != retable.end(); ++it) {
delete (it->second);
}
}
// Send the next ack from the buffer's job queue.
void RdmaAckBuffer::SendNextItem() {
uint32_t imm_data = LookupBufferIndex("rx_ack_buffer");
@ -655,6 +663,99 @@ void RdmaMessageBuffer::SendNextItem() {
}
}
Rendezvous::DoneCallback RdmaTensorBuffer::getRecvTensorCallback(
const string& key_with_step_id, const string& key, int64 step_id,
const Rendezvous::ParsedKey& parsed) {
Rendezvous::DoneCallback cb = [this, key_with_step_id, key, step_id, parsed](
const Status& status, const Rendezvous::Args& send_args,
const Rendezvous::Args& recv_args, const Tensor& in, bool is_dead) {
CHECK(status.ok()) << "RecvLocalAsync was not ok, key" << key_with_step_id
<< " error message: " << status.error_message();
size_t buffer_size = RdmaMessage::kMessageTotalBytes;
size_t tensor_bytes = 0;
// Figures out which device the tensor is hosted on.
Device* src_dev = nullptr;
Status s = channel_->adapter_->worker_env_->device_mgr->LookupDevice(
parsed.src_device, &src_dev);
CHECK(s.ok()) << "src device not found";
// Does the device have the right incarnation number we expect?
CHECK(src_dev->attributes().incarnation() == parsed.src_incarnation)
<< "RecvTensor expects a different device incarnation: "
<< parsed.src_incarnation << " vs. "
<< src_dev->attributes().incarnation()
<< ". Your worker job was probably restarted. Check your "
<< "worker job for the reason why it was restarted.";
Device* dst_dev = nullptr;
// destination is on CPU.
s = channel_->adapter_->worker_env_->device_mgr->LookupDevice("CPU:0",
&dst_dev);
CHECK(s.ok()) << "dst device not found";
AllocatorAttributes dst_alloc_attr;
dst_alloc_attr.set_on_host(true);
bool can_memcpy = DataTypeCanUseMemcpy(in.dtype());
// string tensor needs to be serialized
Tensor copy;
TensorProto proto;
if (src_dev->tensorflow_gpu_device_info() &&
(!send_args.alloc_attrs.on_host())) {
CHECK(send_args.device_context)
<< "send dev name: " << src_dev->name()
<< " gpu_info: " << src_dev->tensorflow_gpu_device_info();
if (can_memcpy) {
AllocatorAttributes host_alloc_attrs;
host_alloc_attrs.set_gpu_compatible(true);
host_alloc_attrs.set_on_host(true);
Allocator* alloc = ProcessState::singleton()->GetCUDAHostAllocator(0);
copy = Tensor(alloc, in.dtype(), in.shape());
tensor_bytes = in.TotalBytes();
buffer_size += tensor_bytes;
GPUUtil::CopyGPUTensorToCPU(
src_dev, send_args.device_context, &in, &copy,
[this, copy, tensor_bytes, buffer_size, key, in, step_id,
key_with_step_id, is_dead, send_args, recv_args](const Status& s) {
CHECK(s.ok()) << "copy tensor from gpu sync";
StringPiece copy_buf;
copy_buf = copy.tensor_data();
PostCopyOperations(true, buffer_size, tensor_bytes, key, in,
step_id, is_dead, key_with_step_id, &copy,
NULL, &copy_buf, send_args, recv_args);
});
} else {
// "val" is on a GPU. No longer uses GPUUtil to fill the proto, use
// aync instead
GPUUtil::SetProtoFromGPU(
in, src_dev, send_args.device_context, &proto, is_dead,
[this, proto, buffer_size, key, in, step_id, key_with_step_id,
is_dead, send_args, recv_args](const Status& s) mutable {
CHECK(s.ok()) << "copy proto from gpu sync";
auto tensor_bytes = proto.ByteSize();
buffer_size += tensor_bytes;
PostCopyOperations(false, buffer_size, tensor_bytes, key, in,
step_id, is_dead, key_with_step_id, NULL,
&proto, NULL, send_args, recv_args);
});
}
} else {
// tensor is in CPU memory.
StringPiece copy_buf;
if (can_memcpy) {
copy_buf = in.tensor_data();
tensor_bytes = in.TotalBytes();
} else {
in.AsProtoTensorContent(&proto);
tensor_bytes = proto.ByteSize();
}
buffer_size += tensor_bytes;
PostCopyOperations(can_memcpy, buffer_size, tensor_bytes, key, in,
step_id, is_dead, key_with_step_id, &copy, &proto,
&copy_buf, send_args, recv_args);
}
};
return cb;
}
// Send the next tensor from the buffer's job queue.
void RdmaTensorBuffer::SendNextItem() {
// get the key
@ -666,6 +767,7 @@ void RdmaTensorBuffer::SendNextItem() {
queue_.pop();
}
}
// send the tensor if a key is acquired.
if (key_with_step_id != "") {
VLOG(2) << "try to send tensor: " << key_with_step_id;
@ -675,107 +777,54 @@ void RdmaTensorBuffer::SendNextItem() {
CHECK(key.compare(name_) == 0);
Rendezvous::ParsedKey parsed;
Rendezvous::ParseKey(key, &parsed);
Rendezvous::DoneCallback cb = [this, key_with_step_id, key, step_id,
parsed](const Status& status,
const Rendezvous::Args& send_args,
const Rendezvous::Args& recv_args,
const Tensor& in, bool is_dead) {
CHECK(status.ok()) << "RecvLocalAsync was not ok, key" << key_with_step_id
<< " error message: " << status.error_message();
size_t buffer_size = RdmaMessage::kMessageTotalBytes;
size_t tensor_bytes = 0;
// Figures out which device the tensor is hosted on.
Device* src_dev = nullptr;
Status s = channel_->adapter_->worker_env_->device_mgr->LookupDevice(
parsed.src_device, &src_dev);
CHECK(s.ok()) << "src device not found";
// Does the device have the right incarnation number we expect?
CHECK(src_dev->attributes().incarnation() == parsed.src_incarnation)
<< "RecvTensor expects a different device incarnation: "
<< parsed.src_incarnation << " vs. "
<< src_dev->attributes().incarnation()
<< ". Your worker job was probably restarted. Check your "
<< "worker job for the reason why it was restarted.";
Device* dst_dev = nullptr;
// destination is on CPU.
s = channel_->adapter_->worker_env_->device_mgr->LookupDevice("CPU:0",
&dst_dev);
CHECK(s.ok()) << "dst device not found";
AllocatorAttributes dst_alloc_attr;
dst_alloc_attr.set_on_host(true);
bool can_memcpy = DataTypeCanUseMemcpy(in.dtype());
// string tensor needs to be serialized
Tensor copy;
TensorProto proto;
if (src_dev->tensorflow_gpu_device_info() &&
(!send_args.alloc_attrs.on_host())) {
CHECK(send_args.device_context)
<< "send dev name: " << src_dev->name()
<< " gpu_info: " << src_dev->tensorflow_gpu_device_info();
if (can_memcpy) {
AllocatorAttributes host_alloc_attrs;
host_alloc_attrs.set_gpu_compatible(true);
host_alloc_attrs.set_on_host(true);
Allocator* alloc = ProcessState::singleton()->GetCUDAHostAllocator(0);
copy = Tensor(alloc, in.dtype(), in.shape());
tensor_bytes = in.TotalBytes();
buffer_size += tensor_bytes;
GPUUtil::CopyGPUTensorToCPU(
src_dev, send_args.device_context, &in, &copy,
[this, copy, tensor_bytes, buffer_size, key, in, step_id,
key_with_step_id, is_dead](const Status& s) {
CHECK(s.ok()) << "copy tensor from gpu sync";
StringPiece copy_buf;
copy_buf = copy.tensor_data();
PostCopyOperations(true, buffer_size, tensor_bytes, key, in,
step_id, is_dead, key_with_step_id, &copy,
NULL, &copy_buf);
});
} else {
// "val" is on a GPU. No longer uses GPUUtil to fill the proto, use
// aync instead
GPUUtil::SetProtoFromGPU(
in, src_dev, send_args.device_context, &proto, is_dead,
[this, proto, buffer_size, key, in, step_id, key_with_step_id,
is_dead](const Status& s) mutable {
CHECK(s.ok()) << "copy proto from gpu sync";
auto tensor_bytes = proto.ByteSize();
buffer_size += tensor_bytes;
PostCopyOperations(false, buffer_size, tensor_bytes, key, in,
step_id, is_dead, key_with_step_id, NULL,
&proto, NULL);
});
}
} else {
// tensor is in CPU memory.
StringPiece copy_buf;
if (can_memcpy) {
copy_buf = in.tensor_data();
tensor_bytes = in.TotalBytes();
} else {
in.AsProtoTensorContent(&proto);
tensor_bytes = proto.ByteSize();
}
buffer_size += tensor_bytes;
PostCopyOperations(can_memcpy, buffer_size, tensor_bytes, key, in,
step_id, is_dead, key_with_step_id, &copy, &proto,
&copy_buf);
}
// maybe some margin for string tensor?
};
Rendezvous::DoneCallback cb =
getRecvTensorCallback(key_with_step_id, key, step_id, parsed);
channel_->adapter_->worker_env_->rendezvous_mgr->RecvLocalAsync(step_id,
parsed, cb);
}
}
void RdmaTensorBuffer::ReSendNextItem() {
// get the key
string key_with_step_id = "";
{
mutex_lock lock{mu_};
if (!requeue.empty()) {
key_with_step_id = requeue.front();
requeue.pop();
}
}
// send the tensor if a key is acquired.
if (key_with_step_id != "") {
VLOG(2) << "try to send tensor: " << key_with_step_id;
string key;
int64 step_id;
VerbsUtil::GetKeyAndStepId(key_with_step_id, key, step_id);
CHECK(key.compare(name_) == 0);
Rendezvous::ParsedKey parsed;
Rendezvous::ParseKey(key, &parsed);
Rendezvous::DoneCallback cb =
getRecvTensorCallback(key_with_step_id, key, step_id, parsed);
ReItem* item;
{
mutex_lock lock{mu_};
Itable it = retable.find(key_with_step_id);
CHECK(it != retable.end()) << "Could not find dup-recv context";
item = it->second;
retable.erase(it);
}
cb(Status::OK(), item->send_args, item->recv_args, item->in, item->is_dead);
delete (item);
}
}
void RdmaTensorBuffer::PostCopyOperations(
bool can_memcpy, size_t buffer_size, size_t tensor_bytes, const string& key,
const Tensor& in, int64 step_id, bool is_dead,
const string& key_with_step_id, const Tensor* copy,
const TensorProto* proto, const StringPiece* copy_buf) {
const TensorProto* proto, const StringPiece* copy_buf,
const Rendezvous::Args& send_args, const Rendezvous::Args& recv_args) {
// prepare message
RdmaMessage rm;
rm.name_size_ = key.size();
@ -793,9 +842,12 @@ void RdmaTensorBuffer::PostCopyOperations(
VLOG(2) << "Extend RDMA buffer from " << size_ << " to " << buffer_size;
}
CreateCPUBuffer(buffer_size, false);
// Need to be received again, put into the re-recv queue and the table
requeue.push(key_with_step_id);
ReItem* item = new ReItem(send_args, recv_args, in, is_dead);
retable.insert(std::pair<string, ReItem*>(key_with_step_id, item));
mu_.unlock();
// put back the key since it is not sent;
EnqueueItem(key_with_step_id);
// no longer used: put back the key since it is not sent;
// ask the remote to create the same buffer
rm.type_ = RDMA_MESSAGE_BUFFER_REQUEST;
rm.remote_addr_ = reinterpret_cast<uint64_t>(buffer_);
@ -841,9 +893,11 @@ void RdmaTensorBuffer::PostCopyOperations(
}
Write(imm_data, buffer_size);
} else {
// Need to be received again, put into the re-recv queue and the table
requeue.push(key_with_step_id);
ReItem* item = new ReItem(send_args, recv_args, in, is_dead);
retable.insert(std::pair<string, ReItem*>(key_with_step_id, item));
mu_.unlock();
// put back the key since it is not sent;
EnqueueItem(key_with_step_id);
}
}

50
tensorflow/contrib/verbs/rdma.h
View File

@ -28,6 +28,7 @@ limitations under the License.
#include <vector>
#include "tensorflow/core/distributed_runtime/worker_env.h"
#include "tensorflow/core/framework/rendezvous.h"
#include "tensorflow/core/framework/tensor.h"
#include "tensorflow/core/framework/tensor_shape.h"
#include "tensorflow/core/framework/types.h"
@ -224,14 +225,57 @@ class RdmaMessageBuffer : public RdmaBuffer {
class RdmaTensorBuffer : public RdmaBuffer {
public:
explicit RdmaTensorBuffer(RdmaChannel* channel, string name);
virtual ~RdmaTensorBuffer() override {}
virtual ~RdmaTensorBuffer() override;
void SendNextItem() override;
void PostCopyOperations(bool can_memcpy, size_t buffer_size,
size_t tensor_bytes, const string& key,
const Tensor& in, int64 step_id, bool is_dead,
const string& key_with_step_id, const Tensor* copy,
const TensorProto* proto,
const StringPiece* copy_buf);
const TensorProto* proto, const StringPiece* copy_buf,
const Rendezvous::Args& send_args,
const Rendezvous::Args& recv_args);
void ReSendNextItem();
private:
Rendezvous::DoneCallback getRecvTensorCallback(
const string& key_with_step_id, const string& key, int64 step_id,
const Rendezvous::ParsedKey& parsed);
struct ReItem {
Rendezvous::Args send_args;
Rendezvous::Args recv_args;
Tensor in;
bool is_dead;
ReItem(const Rendezvous::Args& send_args_,
const Rendezvous::Args& recv_args_, const Tensor& in_, bool is_dead_)
: send_args(send_args_),
recv_args(recv_args_),
in(in_),
is_dead(is_dead_) {
if (send_args.device_context) {
send_args.device_context->Ref();
}
if (recv_args.device_context) {
recv_args.device_context->Ref();
}
}
~ReItem() {
if (send_args.device_context) {
send_args.device_context->Unref();
}
if (recv_args.device_context) {
recv_args.device_context->Unref();
}
}
};
typedef std::map<string, ReItem*> Table;
typedef Table::iterator Itable;
std::queue<string> requeue GUARDED_BY(mu_);
Table retable GUARDED_BY(mu_);
};
struct RdmaMessage {

6
tensorflow/core/BUILD
View File

@ -790,13 +790,16 @@ cc_library(
]) + if_mkl([
"//tensorflow/core/kernels:mkl_concat_op",
"//tensorflow/core/kernels:mkl_conv_op",
"//tensorflow/core/kernels:mkl_cwise_ops_common",
"//tensorflow/core/kernels:mkl_fused_batch_norm_op",
"//tensorflow/core/kernels:mkl_identity_op",
"//tensorflow/core/kernels:mkl_input_conversion_op",
"//tensorflow/core/kernels:mkl_lrn_op",
"//tensorflow/core/kernels:mkl_pooling_ops",
"//tensorflow/core/kernels:mkl_relu_op",
"//tensorflow/core/kernels:mkl_reshape_op",
"//tensorflow/core/kernels:mkl_tfconv_op",
"//tensorflow/core/kernels:mkl_aggregate_ops",
]),
)
@ -2481,10 +2484,13 @@ tf_cc_test_mkl(
"//tensorflow/cc:cc_ops",
"//tensorflow/cc:scope",
"//tensorflow/cc:sendrecv_ops",
"//tensorflow/core/kernels:mkl_aggregate_ops",
"//tensorflow/core/kernels:mkl_concat_op",
"//tensorflow/core/kernels:mkl_conv_op",
"//tensorflow/core/kernels:mkl_cwise_ops_common",
"//tensorflow/core/kernels:mkl_fused_batch_norm_op",
"//tensorflow/core/kernels:mkl_identity_op",
"//tensorflow/core/kernels:mkl_input_conversion_op",
"//tensorflow/core/kernels:mkl_lrn_op",
"//tensorflow/core/kernels:mkl_pooling_ops",
"//tensorflow/core/kernels:mkl_relu_op",

4
tensorflow/core/common_runtime/mkl_cpu_allocator.h
View File

@ -75,12 +75,12 @@ class MklCPUAllocator : public Allocator {
// Hooks provided by this allocator for memory allocation routines from MKL
static inline void* MallocHook(size_t size) {
VLOG(2) << "MklCPUAllocator: In MallocHook";
VLOG(3) << "MklCPUAllocator: In MallocHook";
return cpu_allocator()->AllocateRaw(kAlignment, size);
}
static inline void FreeHook(void* ptr) {
VLOG(2) << "MklCPUAllocator: In FreeHook";
VLOG(3) << "MklCPUAllocator: In FreeHook";
cpu_allocator()->DeallocateRaw(ptr);
}

5
tensorflow/core/distributed_runtime/rpc/grpc_worker_cache.cc
View File

@ -69,9 +69,8 @@ class GrpcWorkerCache : public WorkerCachePartial {
} else {
SharedGrpcChannelPtr channel = channel_cache_->FindWorkerChannel(target);
if (!channel) return nullptr;
WorkerInterface* ret = NewGrpcRemoteWorker(&live_rpc_counter_, channel,
&completion_queue_, &logger_);
return ret;
return NewGrpcRemoteWorker(&live_rpc_counter_, channel,
&completion_queue_, &logger_);
}
}

18
tensorflow/core/distributed_runtime/worker_cache_partial.cc
View File

@ -29,7 +29,7 @@ namespace tensorflow {
bool WorkerCachePartial::GetDeviceLocalityNonBlocking(
const string& device_name, DeviceLocality* locality) {
mutex_lock lock(mu_); // could use reader lock
const auto& iter = device_status_cache_.find(device_name);
auto iter = device_status_cache_.find(device_name);
if (iter != device_status_cache_.end()) {
*locality = iter->second.locality();
return true;
@ -44,16 +44,8 @@ void WorkerCachePartial::GetDeviceLocalityAsync(const string& device_name,
// If cache entry was empty, make one try to fill it by RPC.
SchedClosure([this, &device_name, locality, done]() {
Status s = RefreshDeviceStatus(device_name);
if (s.ok()) {
if (!GetDeviceLocalityNonBlocking(device_name, locality)) {
mutex_lock lock(mu_);
const auto& iter = device_status_cache_.find(device_name);
if (iter == device_status_cache_.end()) {
s = errors::Unavailable("No known remote device: ", device_name);
} else {
s = errors::Internal("Failed to find locality for ", device_name);
}
}
if (s.ok() && !GetDeviceLocalityNonBlocking(device_name, locality)) {
s = errors::Unavailable("No known remote device: ", device_name);
}
done(s);
});
@ -70,7 +62,9 @@ Status WorkerCachePartial::RefreshDeviceStatus(const string& device_name) {
s = errors::InvalidArgument("Bad device name to RefreshDeviceStatus: ",
device_name);
}
auto deleter = [this, task](WorkerInterface* wi) { ReleaseWorker(task, wi); };
auto deleter = [this, &task](WorkerInterface* wi) {
ReleaseWorker(task, wi);
};
std::unique_ptr<WorkerInterface, decltype(deleter)> rwi(CreateWorker(task),
deleter);
if (s.ok() && !rwi.get()) {

167
tensorflow/core/graph/mkl_layout_pass.cc
View File

@ -256,6 +256,7 @@ class MklLayoutRewritePass : public GraphOptimizationPass {
public:
MklLayoutRewritePass() {
// NOTE: names are alphabetically sorted.
csinfo_.addn = "AddN";
csinfo_.avg_pool = "AvgPool";
csinfo_.avg_pool_grad = "AvgPoolGrad";
csinfo_.bias_add = "BiasAdd";
@ -279,17 +280,31 @@ class MklLayoutRewritePass : public GraphOptimizationPass {
csinfo_.mkl_conv2d_with_bias = "_MklConv2DWithBias";
csinfo_.mkl_conv2d_with_bias_backprop_bias =
"_MklConv2DWithBiasBackpropBias";
csinfo_.relu = "Relu";
csinfo_.relu_grad = "ReluGrad";
csinfo_.reshape = "Reshape";
csinfo_.split = "Split";
csinfo_.relu = "Relu";
csinfo_.relu_grad = "ReluGrad";
csinfo_.reshape = "Reshape";
csinfo_.split = "Split";
// Element-wise ops. Ensure you also add any new ops to IsOpElementWise
// in the MklUtil.h (IsMklElementWiseOp method) to ensure that the
// MklInputConversion op is added before it.
csinfo_.add = "Add";
csinfo_.maximum = "Maximum";
csinfo_.mul = "Mul";
csinfo_.squared_difference = "SquaredDifference";
csinfo_.sub = "Sub";
// End - element-wise ops. See note above.
// NOTE: names are alphabetically sorted.
rinfo_.push_back({csinfo_.addn, mkl_op_registry::GetMklOpName(csinfo_.addn), CopyAttrsAddN,
AddNRewrite, nullptr});
rinfo_.push_back({csinfo_.add,
mkl_op_registry::GetMklOpName(csinfo_.add),
CopyAttrsDataType, AlwaysRewrite, nullptr});
rinfo_.push_back({csinfo_.avg_pool,
GetMklOpName(csinfo_.avg_pool),
mkl_op_registry::GetMklOpName(csinfo_.avg_pool),
CopyAttrsPooling, AlwaysRewrite, nullptr});
rinfo_.push_back({csinfo_.avg_pool_grad,
GetMklOpName(csinfo_.avg_pool_grad),
mkl_op_registry::GetMklOpName(csinfo_.avg_pool_grad),
CopyAttrsPooling, AlwaysRewrite, nullptr});
// BiasAddGrad gets written into Conv2DWithBiasBackpropBias depending
// on if context contains Conv2D.
@ -303,50 +318,62 @@ class MklLayoutRewritePass : public GraphOptimizationPass {
CopyAttrsBiasAddGrad, ContextMatchRewrite,
&biasaddgrad_matmul_context_});
rinfo_.push_back({csinfo_.concat,
GetMklOpName(csinfo_.concat),
mkl_op_registry::GetMklOpName(csinfo_.concat),
CopyAttrsConcat, AlwaysRewrite, nullptr});
rinfo_.push_back({csinfo_.concatv2,
GetMklOpName(csinfo_.concatv2),
mkl_op_registry::GetMklOpName(csinfo_.concatv2),
CopyAttrsConcatV2, AlwaysRewrite, nullptr});
rinfo_.push_back({csinfo_.conv2d,
GetMklOpName(csinfo_.conv2d),
mkl_op_registry::GetMklOpName(csinfo_.conv2d),
CopyAttrsConv2D, AlwaysRewrite, nullptr});
rinfo_.push_back({csinfo_.conv2d_grad_filter,
GetMklOpName(csinfo_.conv2d_grad_filter),
mkl_op_registry::GetMklOpName(csinfo_.conv2d_grad_filter),
CopyAttrsConv2D, AlwaysRewrite, nullptr});
rinfo_.push_back({csinfo_.conv2d_grad_input,
GetMklOpName(csinfo_.conv2d_grad_input),
mkl_op_registry::GetMklOpName(csinfo_.conv2d_grad_input),
CopyAttrsConv2D, AlwaysRewrite, nullptr});
rinfo_.push_back({csinfo_.fused_batch_norm,
GetMklOpName(csinfo_.fused_batch_norm),
mkl_op_registry::GetMklOpName(csinfo_.fused_batch_norm),
CopyAttrsFusedBatchNorm, AlwaysRewrite, nullptr});
rinfo_.push_back({csinfo_.fused_batch_norm_grad,
GetMklOpName(csinfo_.fused_batch_norm_grad),
mkl_op_registry::GetMklOpName(csinfo_.fused_batch_norm_grad),
CopyAttrsFusedBatchNorm, AlwaysRewrite, nullptr});
rinfo_.push_back({csinfo_.identity,
GetMklOpName(csinfo_.identity),
mkl_op_registry::GetMklOpName(csinfo_.identity),
CopyAttrsIdentity, AlwaysRewrite, nullptr});
rinfo_.push_back({csinfo_.lrn,
GetMklOpName(csinfo_.lrn),
mkl_op_registry::GetMklOpName(csinfo_.lrn),
CopyAttrsLRN, AlwaysRewrite, nullptr});
rinfo_.push_back({csinfo_.lrn_grad,
GetMklOpName(csinfo_.lrn_grad),
mkl_op_registry::GetMklOpName(csinfo_.lrn_grad),
CopyAttrsLRN, AlwaysRewrite, nullptr});
rinfo_.push_back({csinfo_.max_pool,
GetMklOpName(csinfo_.max_pool),
mkl_op_registry::GetMklOpName(csinfo_.max_pool),
CopyAttrsPooling, NonDepthBatchWisePoolRewrite, nullptr});
rinfo_.push_back({csinfo_.max_pool_grad,
GetMklOpName(csinfo_.max_pool_grad),
mkl_op_registry::GetMklOpName(csinfo_.max_pool_grad),
CopyAttrsPooling, AlwaysRewrite, nullptr});
rinfo_.push_back({csinfo_.maximum,
mkl_op_registry::GetMklOpName(csinfo_.maximum),
CopyAttrsDataType, AlwaysRewrite, nullptr});
rinfo_.push_back({csinfo_.mul,
mkl_op_registry::GetMklOpName(csinfo_.mul),
CopyAttrsDataType, AlwaysRewrite, nullptr});
rinfo_.push_back({csinfo_.relu,
GetMklOpName(csinfo_.relu),
CopyAttrsRelu, AlwaysRewrite, nullptr});
mkl_op_registry::GetMklOpName(csinfo_.relu),
CopyAttrsDataType, AlwaysRewrite, nullptr});
rinfo_.push_back({csinfo_.relu_grad,
GetMklOpName(csinfo_.relu_grad),
CopyAttrsRelu, AlwaysRewrite, nullptr});
mkl_op_registry::GetMklOpName(csinfo_.relu_grad),
CopyAttrsDataType, AlwaysRewrite, nullptr});
rinfo_.push_back({csinfo_.reshape,
GetMklOpName(csinfo_.reshape),
mkl_op_registry::GetMklOpName(csinfo_.reshape),
CopyAttrsReshape, AlwaysRewrite, nullptr});
rinfo_.push_back({csinfo_.squared_difference,
mkl_op_registry::GetMklOpName(csinfo_.squared_difference),
CopyAttrsDataType, AlwaysRewrite, nullptr});
rinfo_.push_back({csinfo_.sub,
mkl_op_registry::GetMklOpName(csinfo_.sub),
CopyAttrsDataType, AlwaysRewrite, nullptr});
// Add info about which ops to add workspace edge to and the slots.
wsinfo_.push_back({csinfo_.lrn, csinfo_.lrn_grad, 0, 2, 1, 3});
@ -429,6 +456,8 @@ class MklLayoutRewritePass : public GraphOptimizationPass {
/// Structure to store all constant strings
/// NOTE: names are alphabetically sorted.
typedef struct {
string addn;
string add;
string avg_pool;
string avg_pool_grad;
string bias_add;
@ -446,15 +475,19 @@ class MklLayoutRewritePass : public GraphOptimizationPass {
string matmul;
string max_pool;
string max_pool_grad;
string maximum;
string mkl_conv2d;
string mkl_conv2d_grad_input;
string mkl_conv2d_grad_filter;
string mkl_conv2d_with_bias;
string mkl_conv2d_with_bias_backprop_bias;
string mul;
string relu;
string relu_grad;
string reshape;
string split;
string squared_difference;
string sub;
} ConstStringsInfo;
private:
@ -502,15 +535,6 @@ class MklLayoutRewritePass : public GraphOptimizationPass {
return N;
}
// Get the name of Mkl op from original TensorFlow op
// We prefix 'Mkl' to the original op to get Mkl op.
// TODO(nhasabni) We should move this to mkl_util.h.
inline string GetMklOpName(const string& name) const {
// Prefix that we add to Tensorflow op name to construct Mkl op name.
const char* const kMklOpPrefix = "_Mkl";
return string(kMklOpPrefix) + name;
}
// Can op represented by node 'n' run on DEVICE_CPU?
// Op can run on CPU with MKL if the runtime assigned device or the
// user requested device contains device CPU, or both are empty.
@ -604,6 +628,19 @@ class MklLayoutRewritePass : public GraphOptimizationPass {
return false;
}
static bool AddNRewrite(const Node* n, const ContextInfo* c) {
CHECK_NOTNULL(n);
int num;
CHECK_EQ(GetNodeAttr(n->def(), "N", &num).ok(), true);
// Condition that specifies non-batch-wise and non-depth-wise pooling.
if (num == 2) {
return true;
}
return false;
}
// Is BiasAddGrad node in 'n' is associated with Conv2DWithBias node
// specified in contextinfo 'ci'. Function updates fwd_node to point
// to Conv2DWithBias node if 'n' is associated with Conv2DWithBias.
@ -907,15 +944,16 @@ class MklLayoutRewritePass : public GraphOptimizationPass {
// We need operator-specific function to copy attributes because the framework
// does not provide any generic function for it.
// NOTE: names are alphabetically sorted.
static void CopyAttrsAddN(const Node* orig_node, NodeBuilder* nb);
static void CopyAttrsBiasAddGrad(const Node* orig_node, NodeBuilder* nb);
static void CopyAttrsConcat(const Node* orig_node, NodeBuilder* nb);
static void CopyAttrsConcatV2(const Node* orig_node, NodeBuilder* nb);
static void CopyAttrsConv2D(const Node* orig_node, NodeBuilder* nb);
static void CopyAttrsDataType(const Node* orig_node, NodeBuilder* nb);
static void CopyAttrsFusedBatchNorm(const Node* orig_node, NodeBuilder* nb);
static void CopyAttrsIdentity(const Node* orig_node, NodeBuilder* nb);
static void CopyAttrsLRN(const Node* orig_node, NodeBuilder* nb);
static void CopyAttrsPooling(const Node* orig_node, NodeBuilder* nb);
static void CopyAttrsRelu(const Node* orig_node, NodeBuilder* nb);
static void CopyAttrsReshape(const Node* orig_node, NodeBuilder* nb);
static void CopyAttrsSplit(const Node* orig_node, NodeBuilder* nb);
@ -1334,7 +1372,7 @@ void MklLayoutRewritePass::AddWorkSpaceEdgeIfNeeded(
TF_CHECK_OK(GetNodeAttr(orig_node->def(), "T", &T));
for (auto ws : wsinfo_) {
if (orig_node->type_string() == ws.fwd_op &&
mkl_op_registry::IsMklOp(GetMklOpName(orig_node->type_string()), T)) {
mkl_op_registry::IsMklOp(mkl_op_registry::GetMklOpName(orig_node->type_string()), T)) {
// If this op is a fwd op, then we need to check if there is an
// edge from this node's fwd_slot to bwdop's bwd_slot. If there is
// an edge, then we just add an attribute on this node for setting
@ -1360,7 +1398,7 @@ void MklLayoutRewritePass::AddWorkSpaceEdgeIfNeeded(
nb->Attr("workspace_enabled", false);
}
} else if (orig_node->type_string() == ws.bwd_op &&
mkl_op_registry::IsMklOp(GetMklOpName(orig_node->type_string()),
mkl_op_registry::IsMklOp(mkl_op_registry::GetMklOpName(orig_node->type_string()),
T)) {
// If this op is a bwd op, then we need to add workspace edge and
// it's Mkl tensor edge between its corresponding fwd op and this
@ -1376,7 +1414,7 @@ void MklLayoutRewritePass::AddWorkSpaceEdgeIfNeeded(
if (e->src_output() == ws.fwd_slot &&
// We would have rewritten the forward op, so we need to use
// GetMklOpName call to get its Mkl name.
e->src()->type_string() == GetMklOpName(ws.fwd_op) &&
e->src()->type_string() == mkl_op_registry::GetMklOpName(ws.fwd_op) &&
e->dst_input() == ws.bwd_slot) {
nb->Attr("workspace_enabled", true);
CHECK_NOTNULL(ws_tensors);
@ -1455,6 +1493,20 @@ void MklLayoutRewritePass::CopyAttrsConv2D(const Node* orig_node,
nb->Attr("use_cudnn_on_gpu", use_cudnn_on_gpu);
}
void MklLayoutRewritePass::CopyAttrsAddN(const Node* orig_node,
NodeBuilder* nb) {
DataType T;
int N;
// Get all attributes from old node.
TF_CHECK_OK(GetNodeAttr(orig_node->def(), "T", &T));
TF_CHECK_OK(GetNodeAttr(orig_node->def(), "N", &N));
// Add attributes to new node.
nb->Attr("T", T);
nb->Attr("N", N);
}
void MklLayoutRewritePass::CopyAttrsBiasAddGrad(const Node* orig_node,
NodeBuilder* nb) {
DataType T;
@ -1527,8 +1579,8 @@ void MklLayoutRewritePass::CopyAttrsPooling(const Node* orig_node,
nb->Attr("data_format", data_format);
}
void MklLayoutRewritePass::CopyAttrsRelu(const Node* orig_node,
NodeBuilder* nb) {
void MklLayoutRewritePass::CopyAttrsDataType(const Node* orig_node,
NodeBuilder* nb) {
DataType T;
// Get all attributes from old node.
@ -1894,7 +1946,15 @@ Status MklLayoutRewritePass::RewriteNode(std::unique_ptr<Graph>* g,
}
// Get all inputs.
const int num_inputs = orig_node->in_edges().size();
int num_inputs = orig_node->in_edges().size();
// Drop count for control edges from inputs
for (const Edge* e : orig_node->in_edges()) {
if (e->IsControlEdge()) {
num_inputs--;
}
}
gtl::InlinedVector<Node*, 4> control_edges;
gtl::InlinedVector<std::pair<Node*, int>, 4> inputs(num_inputs);
FillInputs(orig_node, &control_edges, &inputs);
@ -2008,7 +2068,34 @@ MklLayoutRewritePass::CheckForNodeRewrite(const Node* n) const {
// BiasAddGrad is not an Mkl layer, so we make an exception for it.
if (n->type_string() != csinfo_.bias_add_grad) {
if (!mkl_op_registry::IsMklOp(GetMklOpName(n->type_string()), T)) {
if (!mkl_op_registry::IsMklOp(mkl_op_registry::GetMklOpName(n->type_string()), T)) {
return nullptr;
}
}
// For elementwise node, we reuse the Eigen implementation and pass the MKL
// metadata tensor through so we can avoid conversions. However, if all
// incoming edges are in TF format, we don't need all this overhead, so
// replace the elementwise node only if at least one of its parents is a MKL
// node.
//
// TODO(vrane): Add implementation for element-wise ops that doesn't reuse
// eigen code to reduce cross-library dependency.
if (mkl_op_registry::IsMklElementWiseOp(
mkl_op_registry::GetMklOpName(n->type_string()), T)) {
bool incoming_mkl_edge = false;
for (auto parent : n->in_edges()) {
if (mkl_op_registry::IsMklOp(
mkl_op_registry::GetMklOpName(parent->src()->type_string()), T)) {
incoming_mkl_edge = true;
break;
} else {
VLOG(1) << "Non-MKL parent is: " << parent->src()->type_string();
}
}
if (incoming_mkl_edge == false) {
VLOG(1) << "Skipping replacement of elementwise node which has no MKL "
"parents.";
return nullptr;
}
}

288
tensorflow/core/graph/mkl_layout_pass_test.cc
View File

@ -133,19 +133,19 @@ TEST_F(MklLayoutPassTest, Basic) {
InitGraph(
"node { name: 'A' op: 'Input'}"
"node { name: 'B' op: 'Input'}"
"node { name: 'C' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'C' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }"
"node { name: 'D' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'D' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Mul);D(Mul)|"
"A(Input);B(Input);C(Zeta);D(Zeta)|"
"A->C;A->D;B->C:1;B->D:1");
}
// Test set 1: Conv2D + AddBias
// C=_MklConv2D(A,M,B,N); E=BiasAdd(C,D); Z=Sub(E,Y) (for interleaved ordering)
// C=_MklConv2D(A,B,M,N); E=BiasAdd(C,D); Z=Sub(E,Y) (for contiguous ordering)
// C=_MklConv2D(A,M,B,N); E=BiasAdd(C,D); Z=Zeta(E,Y) (for interleaved ordering)
// C=_MklConv2D(A,B,M,N); E=BiasAdd(C,D); Z=Zeta(E,Y) (for contiguous ordering)
TEST_F(MklLayoutPassTest, NodeMerge_Conv2DWithBias_Positive) {
CHECK_EQ(kTensorOrdering, MklTfTensorOrdering::TENSORS_CONTIGUOUS);
InitGraph(
@ -166,18 +166,18 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DWithBias_Positive) {
" attr { key: 'data_format' value { s: 'NCHW' } }"
" input: ['C', 'D'] }"
"node { name: 'Y' op: 'Input'}"
"node { name: 'Z' op: 'Sub'"
"node { name: 'Z' op: 'Zeta'"
" attr {key: 'T' value { type: DT_FLOAT } }"
" input: ['E', 'Y']}");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);D(Input);DMT/_0(Const);E(_MklConv2DWithBias);"
"M(_MklInput);N(_MklInput);Y(Input);Z(Sub)|A->E;"
"M(_MklInput);N(_MklInput);Y(Input);Z(Zeta)|A->E;"
"A:control->DMT/_0:control;B->E:1;D->E:2;DMT/_0->E:5;E->Z;M->E:3;"
"N->E:4;Y->Z:1");
}
// C=_MklConv2D(A,M:1,B,N:1); E=BiasAdd(C,D); Z=Sub(E,Y) (for interleaved)
// C=_MklConv2D(A,B,M:1,N:1); E=BiasAdd(C,D); Z=Sub(E,Y) (for contiguous)
// C=_MklConv2D(A,M:1,B,N:1); E=BiasAdd(C,D); Z=Zeta(E,Y) (for interleaved)
// C=_MklConv2D(A,B,M:1,N:1); E=BiasAdd(C,D); Z=Zeta(E,Y) (for contiguous)
// Test for correct output slots selected
TEST_F(MklLayoutPassTest, NodeMerge_Conv2DWithBias_Positive1) {
CHECK_EQ(kTensorOrdering, MklTfTensorOrdering::TENSORS_CONTIGUOUS);
@ -199,17 +199,17 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DWithBias_Positive1) {
" attr { key: 'data_format' value { s: 'NCHW' } }"
" input: ['C', 'D'] }"
"node { name: 'Y' op: 'Input'}"
"node { name: 'Z' op: 'Sub'"
"node { name: 'Z' op: 'Zeta'"
" attr {key: 'T' value { type: DT_FLOAT } }"
" input: ['E', 'Y']}");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);D(Input);DMT/_0(Const);E(_MklConv2DWithBias);"
"M(_MklInput2);N(_MklInput2);Y(Input);Z(Sub)|A->E;"
"M(_MklInput2);N(_MklInput2);Y(Input);Z(Zeta)|A->E;"
"A:control->DMT/_0:control;B->E:1;D->E:2;DMT/_0->E:5;E->Z;"
"M:1->E:3;N:1->E:4;Y->Z:1");
}
// C=Conv2D(A,B); E=BiasAdd(C,D); Z=Sub(E,Y);
// C=Conv2D(A,B); E=BiasAdd(C,D); Z=Zeta(E,Y);
// This is a case of node rewrite followed by node merge.
// We will first rewrite Conv2D to _MklConv2D, and then merge _MklConv2D
// with BiasAdd to produce _MklConv2DWithBias.
@ -231,12 +231,12 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DWithBias_Positive2) {
" attr { key: 'data_format' value { s: 'NCHW' } }"
" input: ['C', 'D'] }"
"node { name: 'Y' op: 'Input'}"
"node { name: 'Z' op: 'Sub'"
"node { name: 'Z' op: 'Zeta'"
" attr {key: 'T' value { type: DT_FLOAT } }"
" input: ['E', 'Y']}");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);D(Input);DMT/_0(Const);DMT/_1(Const);"
"DMT/_2(Const);E(_MklConv2DWithBias);Y(Input);Z(Sub)|"
"DMT/_2(Const);E(_MklConv2DWithBias);Y(Input);Z(Zeta)|"
"A->E;A:control->DMT/_0:control;A:control->DMT/_1:control;"
"A:control->DMT/_2:control;B->E:1;D->E:2;DMT/_0->E:3;DMT/_1->E:4;"
"DMT/_2->E:5;E->Z;Y->Z:1");
@ -286,7 +286,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DWithBias_Negative_Dataflow1) {
"M(_MklInput);N(_MklInput)|A->C;B->C:1;D->F;E->F:1;M->C:2;N->C:3");
}
// _MklConv2D has two outgoing edges: BiasAdd and some other dummy node (Add).
// _MklConv2D has two outgoing edges: BiasAdd and some other dummy node (Zeta).
// Merge should not be done in such case.
TEST_F(MklLayoutPassTest, NodeMerge_Conv2DWithBias_Negative_Dataflow2) {
InitGraph(
@ -308,12 +308,12 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DWithBias_Negative_Dataflow2) {
" attr { key: 'data_format' value { s: 'NCHW' } }"
" input: ['D', 'E'] }" // Conv2D has two outputs.
// No merge should happen.
"node { name: 'G' op: 'Add'"
"node { name: 'G' op: 'Zeta'"
" attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['C', 'E'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(_MklConv2D);D(Input);E(Input);F(BiasAdd);"
"G(Add);M(_MklInput);N(_MklInput)|A->C;B->C:1;C->G;D->F;"
"G(Zeta);M(_MklInput);N(_MklInput)|A->C;B->C:1;C->G;D->F;"
"E->F:1;E->G:1;M->C:2;N->C:3");
}
@ -362,7 +362,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_Positive) {
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" attr { key: 'padding' value { s: 'SAME' } }"
" input: ['A', 'B', 'C', 'M', 'N', 'O']}"
"node { name: 'E' op: 'Sub'"
"node { name: 'E' op: 'Zeta'"
" attr {key: 'T' value { type: DT_FLOAT } }"
" input: ['D', 'A']}"
"node { name: 'F' op: 'Int32Input'}"
@ -387,7 +387,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_Positive) {
" input: ['E'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Input);D(_MklConv2DWithBias);DMT/_0(Const);"
"E(Sub);F(Int32Input);G(_MklConv2DBackpropFilter);H(Int32Input);"
"E(Zeta);F(Int32Input);G(_MklConv2DBackpropFilter);H(Int32Input);"
"I(_MklConv2DBackpropInput);J(_MklConv2DWithBiasBackpropBias);"
"M(_MklInput);N(_MklInput);O(_MklInput)|A->D;A->E:1;A->G;B->D:1;"
"B->I:1;C->D:2;D->E;DMT/_0->J:1;E->G:2;E->I:2;E->J;"
@ -413,7 +413,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_Negative1) {
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" attr { key: 'padding' value { s: 'SAME' } }"
" input: ['A', 'B', 'C', 'M', 'N', 'O']}"
"node { name: 'E' op: 'Sub'"
"node { name: 'E' op: 'Zeta'"
" attr {key: 'T' value { type: DT_FLOAT } }"
" input: ['D', 'A']}"
"node { name: 'F' op: 'Int32Input'}"
@ -438,7 +438,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_Negative1) {
" input: ['E'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Input);D(_MklConv2DWithBias);"
"E(Sub);F(Int32Input);G(_MklConv2DBackpropFilter);H(Int32Input);"
"E(Zeta);F(Int32Input);G(_MklConv2DBackpropFilter);H(Int32Input);"
"I(_MklConv2DBackpropInput);J(BiasAddGrad);"
"M(_MklInput);N(_MklInput);O(_MklInput)|A->D;A->E:1;A->G:2;B->D:1;"
"B->I:1;C->D:2;D->E;E->G;E->I:2;E->J;F->G:1;H->I;M->D:3;M->G:3;"
@ -463,7 +463,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_Negative2) {
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" attr { key: 'padding' value { s: 'SAME' } }"
" input: ['B', 'A', 'C', 'M', 'N', 'O']}"
"node { name: 'E' op: 'Sub'"
"node { name: 'E' op: 'Zeta'"
" attr {key: 'T' value { type: DT_FLOAT } }"
" input: ['D', 'A']}"
"node { name: 'F' op: 'Int32Input'}"
@ -488,7 +488,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_Negative2) {
" input: ['E'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Input);D(_MklConv2DWithBias);"
"E(Sub);F(Int32Input);G(_MklConv2DBackpropFilter);H(Int32Input);"
"E(Zeta);F(Int32Input);G(_MklConv2DBackpropFilter);H(Int32Input);"
"I(_MklConv2DBackpropInput);J(BiasAddGrad);"
"M(_MklInput);N(_MklInput);O(_MklInput)|A->D:1;A->E:1;A->G;B->D;"
"B->I:1;C->D:2;D->E;E->G:2;E->I:2;E->J;F->G:1;H->I;M->D:3;M->G:3;"
@ -512,7 +512,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_BpropFilter_Positive) {
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" attr { key: 'padding' value { s: 'SAME' } }"
" input: ['A', 'B', 'C', 'M', 'N', 'O']}"
"node { name: 'E' op: 'Sub'"
"node { name: 'E' op: 'Zeta'"
" attr {key: 'T' value { type: DT_FLOAT } }"
" input: ['D', 'A']}"
"node { name: 'F' op: 'Int32Input'}"
@ -529,7 +529,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_BpropFilter_Positive) {
" input: ['E'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Input);D(_MklConv2DWithBias);DMT/_0(Const);"
"E(Sub);F(Int32Input);G(_MklConv2DBackpropFilter);"
"E(Zeta);F(Int32Input);G(_MklConv2DBackpropFilter);"
"H(_MklConv2DWithBiasBackpropBias);M(_MklInput);N(_MklInput);"
"O(_MklInput)|A->D;A->E:1;A->G;B->D:1;C->D:2;D->E;DMT/_0->H:1;"
"E->G:2;E->H;E:control->DMT/_0:control;F->G:1;M->D:3;M->G:3;"
@ -553,7 +553,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_BpropFilter_Negative1) {
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" attr { key: 'padding' value { s: 'SAME' } }"
" input: ['A', 'B', 'C', 'M', 'N', 'O']}"
"node { name: 'E' op: 'Sub'"
"node { name: 'E' op: 'Zeta'"
" attr {key: 'T' value { type: DT_FLOAT } }"
" input: ['D', 'A']}"
"node { name: 'F' op: 'Int32Input'}"
@ -570,7 +570,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_BpropFilter_Negative1) {
" input: ['E'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Input);D(_MklConv2DWithBias);"
"E(Sub);F(Int32Input);G(_MklConv2DBackpropFilter);H(BiasAddGrad);"
"E(Zeta);F(Int32Input);G(_MklConv2DBackpropFilter);H(BiasAddGrad);"
"M(_MklInput);N(_MklInput);O(_MklInput)|A->D;A->E:1;A->G:2;B->D:1;"
"C->D:2;D->E;E->G;E->H;F->G:1;M->D:3;M->G:3;N->D:4;N->G:4;O->D:5;"
"O->G:5");
@ -593,7 +593,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_BpropFilter_Negative2) {
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" attr { key: 'padding' value { s: 'SAME' } }"
" input: ['B', 'A', 'C', 'M', 'N', 'O']}"
"node { name: 'E' op: 'Sub'"
"node { name: 'E' op: 'Zeta'"
" attr {key: 'T' value { type: DT_FLOAT } }"
" input: ['D', 'A']}"
"node { name: 'F' op: 'Int32Input'}"
@ -610,7 +610,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_BpropFilter_Negative2) {
" input: ['E'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Input);D(_MklConv2DWithBias);"
"E(Sub);F(Int32Input);G(_MklConv2DBackpropFilter);H(BiasAddGrad);"
"E(Zeta);F(Int32Input);G(_MklConv2DBackpropFilter);H(BiasAddGrad);"
"M(_MklInput);N(_MklInput);O(_MklInput)|A->D:1;A->E:1;A->G;B->D;"
"C->D:2;D->E;E->G:2;E->H;F->G:1;M->D:3;M->G:3;N->D:4;N->G:4;O->D:5;"
"O->G:5");
@ -618,8 +618,8 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_BpropFilter_Negative2) {
// No _MklConv2DWithBias in context, but _MklConv2D in context.
// No rewrite for BiasAddGrad should happen.
// C=_MklConv2D(A,M,B,N); D=Sub(C,A); E=BiasAddGrad(D) (for interleaved)
// C=_MklConv2D(A,B,M,N); D=Sub(C,A); E=BiasAddGrad(D) (for contiguous)
// C=_MklConv2D(A,M,B,N); D=Zeta(C,A); E=BiasAddGrad(D) (for interleaved)
// C=_MklConv2D(A,B,M,N); D=Zeta(C,A); E=BiasAddGrad(D) (for contiguous)
TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_Neg_NoMklConv2DWithBias) {
InitGraph(
"node { name: 'A' op: 'Input'}"
@ -633,7 +633,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_Neg_NoMklConv2DWithBias) {
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" attr { key: 'padding' value { s: 'SAME' } }"
" input: ['A', 'B', 'M', 'N']}"
"node { name: 'D' op: 'Sub'"
"node { name: 'D' op: 'Zeta'"
" attr {key: 'T' value { type: DT_FLOAT } }"
" input: ['C', 'A']}"
"node { name: 'E' op: 'BiasAddGrad'"
@ -641,21 +641,21 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_Neg_NoMklConv2DWithBias) {
" attr { key: 'data_format' value { s: 'NCHW' } }"
" input: ['D'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(_MklConv2D);D(Sub);E(BiasAddGrad);"
"A(Input);B(Input);C(_MklConv2D);D(Zeta);E(BiasAddGrad);"
"M(_MklInput);N(_MklInput)|A->C;A->D:1;B->C:1;C->D;D->E;"
"M->C:2;N->C:3");
}
// No Conv2D in the context for BiasAddGrad. No rewrite should happen.
// C=Add(A,B); D=Sub(C,A); E=BiasAddGrad(D)
// C=Polygamma(A,B); D=Zeta(C,A); E=BiasAddGrad(D)
TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_Negative_NoConv2D) {
InitGraph(
"node { name: 'A' op: 'Input'}"
"node { name: 'B' op: 'Input'}"
"node { name: 'C' op: 'Add'"
"node { name: 'C' op: 'Polygamma'"
" attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B']}"
"node { name: 'D' op: 'Sub'"
"node { name: 'D' op: 'Zeta'"
" attr {key: 'T' value { type: DT_FLOAT } }"
" input: ['C', 'A']}"
"node { name: 'E' op: 'BiasAddGrad'"
@ -663,13 +663,13 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_Negative_NoConv2D) {
" attr { key: 'data_format' value { s: 'NCHW' } }"
" input: ['D'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Add);D(Sub);E(BiasAddGrad)|"
"A(Input);B(Input);C(Polygamma);D(Zeta);E(BiasAddGrad)|"
"A->C;A->D:1;B->C:1;C->D;D->E");
}
// No Conv2D in the context for BiasAddGrad, but MatMul in context.
// Rewrite should happen, but name of BiasAddGrad does not change.
// C=MatMul(A,B); D=Sub(C,A); E=BiasAddGrad(D)
// C=MatMul(A,B); D=Zeta(C,A); E=BiasAddGrad(D)
TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_Negative_NoConv2D_MatMul) {
InitGraph(
"node { name: 'A' op: 'Input'}"
@ -679,7 +679,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_Negative_NoConv2D_MatMul) {
" attr { key: 'transpose_a' value { b: false } }"
" attr { key: 'transpose_b' value { b: false } }"
" input: ['A', 'B']}"
"node { name: 'D' op: 'Sub'"
"node { name: 'D' op: 'Zeta'"
" attr {key: 'T' value { type: DT_FLOAT } }"
" input: ['C', 'A']}"
"node { name: 'E' op: 'BiasAddGrad'"
@ -687,12 +687,12 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_Negative_NoConv2D_MatMul) {
" attr { key: 'data_format' value { s: 'NCHW' } }"
" input: ['D'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(MatMul);D(Sub);E(BiasAddGrad)|"
"A(Input);B(Input);C(MatMul);D(Zeta);E(BiasAddGrad)|"
"A->C;A->D:1;B->C:1;C->D;D->E");
}
// Test set 3: MatMul..BiasAddGrad -> BiasAddGrad rewrite tests
// C=MatMul(A,B); D=Sub(C,A); E=BiasAddGrad(D)
// C=MatMul(A,B); D=Zeta(C,A); E=BiasAddGrad(D)
TEST_F(MklLayoutPassTest, NodeMerge_MatMulBiasAddGrad_Positive) {
InitGraph(
"node { name: 'A' op: 'Input'}"
@ -702,7 +702,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_MatMulBiasAddGrad_Positive) {
" attr { key: 'transpose_a' value { b: false } }"
" attr { key: 'transpose_b' value { b: false } }"
" input: ['A', 'B']}"
"node { name: 'D' op: 'Sub'"
"node { name: 'D' op: 'Zeta'"
" attr {key: 'T' value { type: DT_FLOAT } }"
" input: ['C', 'A']}"
"node { name: 'E' op: 'BiasAddGrad'"
@ -710,20 +710,20 @@ TEST_F(MklLayoutPassTest, NodeMerge_MatMulBiasAddGrad_Positive) {
" attr { key: 'data_format' value { s: 'NCHW' } }"
" input: ['D'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(MatMul);D(Sub);E(BiasAddGrad)|"
"A(Input);B(Input);C(MatMul);D(Zeta);E(BiasAddGrad)|"
"A->C;A->D:1;B->C:1;C->D;D->E");
}
// No MatMul in the context for BiasAddGrad. No rewrite should happen.
// C=Add(A,B); D=Sub(C,A); E=BiasAddGrad(D)
// C=Polygamma(A,B); D=Zeta(C,A); E=BiasAddGrad(D)
TEST_F(MklLayoutPassTest, NodeMerge_MatMulBiasAddGrad_Negative_NoMatMul) {
InitGraph(
"node { name: 'A' op: 'Input'}"
"node { name: 'B' op: 'Input'}"
"node { name: 'C' op: 'Add'"
"node { name: 'C' op: 'Polygamma'"
" attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B']}"
"node { name: 'D' op: 'Sub'"
"node { name: 'D' op: 'Zeta'"
" attr {key: 'T' value { type: DT_FLOAT } }"
" input: ['C', 'A']}"
"node { name: 'E' op: 'BiasAddGrad'"
@ -731,7 +731,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_MatMulBiasAddGrad_Negative_NoMatMul) {
" attr { key: 'data_format' value { s: 'NCHW' } }"
" input: ['D'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Add);D(Sub);E(BiasAddGrad)|"
"A(Input);B(Input);C(Polygamma);D(Zeta);E(BiasAddGrad)|"
"A->C;A->D:1;B->C:1;C->D;D->E");
}
@ -752,10 +752,10 @@ TEST_F(MklLayoutPassTest, NodeRewrite_Conv2D_Basic) {
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" attr { key: 'padding' value { s: 'SAME' } }"
" input: ['A', 'B']}"
"node { name: 'D' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'D' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['B', 'C'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(_MklConv2D);D(Mul);DMT/_0(Const);"
"A(Input);B(Input);C(_MklConv2D);D(Zeta);DMT/_0(Const);"
"DMT/_1(Const)|A->C;A:control->DMT/_0:control;"
"A:control->DMT/_1:control;B->C:1;B->D;C->D:1;DMT/_0->C:2;"
"DMT/_1->C:3");
@ -781,11 +781,11 @@ TEST_F(MklLayoutPassTest, NodeRewrite_Conv2D_Positive1) {
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" attr { key: 'padding' value { s: 'SAME' } }"
" input: ['A', 'C']}"
"node { name: 'E' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'E' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['C', 'D'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(_MklConv2D);D(_MklConv2D);DMT/_0(Const);"
"DMT/_1(Const);DMT/_2(Const);E(Mul)|A->C;A->D;"
"DMT/_1(Const);DMT/_2(Const);E(Zeta)|A->C;A->D;"
"A:control->DMT/_0:control;A:control->DMT/_1:control;"
"A:control->DMT/_2:control;B->C:1;C->D:1;C->E;"
"C:2->D:3;D->E:1;DMT/_0->C:2;DMT/_1->C:3;DMT/_2->D:2");
@ -803,10 +803,10 @@ TEST_F(MklLayoutPassTest, NodeRewrite_Conv2D_Negative_UnsupportedType) {
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" attr { key: 'padding' value { s: 'SAME' } }"
" input: ['A', 'B']}"
"node { name: 'D' op: 'Mul' attr { key: 'T' value { type: DT_HALF } }"
"node { name: 'D' op: 'Zeta' attr { key: 'T' value { type: DT_HALF } }"
" input: ['B', 'C'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(HalfInput);B(HalfInput);C(Conv2D);D(Mul)|"
"A(HalfInput);B(HalfInput);C(Conv2D);D(Zeta)|"
"A->C;B->C:1;B->D;C->D:1");
}
@ -822,11 +822,11 @@ TEST_F(MklLayoutPassTest, NodeRewrite_Conv2DGradFilter_Positive) {
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" attr { key: 'padding' value { s: 'SAME' } }"
" input: ['A', 'B', 'C']}"
"node { name: 'E' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'E' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'D'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Int32Input);C(Input);D(_MklConv2DBackpropFilter);"
"DMT/_0(Const);DMT/_1(Const);DMT/_2(Const);E(Mul)|"
"DMT/_0(Const);DMT/_1(Const);DMT/_2(Const);E(Zeta)|"
"A->D;A->E;A:control->DMT/_0:control;A:control->DMT/_1:control;"
"A:control->DMT/_2:control;B->D:1;C->D:2;D->E:1;DMT/_0->D:3;"
"DMT/_1->D:4;DMT/_2->D:5");
@ -844,11 +844,11 @@ TEST_F(MklLayoutPassTest, NodeRewrite_Conv2DGradInput_Positive) {
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" attr { key: 'padding' value { s: 'SAME' } }"
" input: ['B', 'A', 'C']}"
"node { name: 'E' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'E' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'D'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Int32Input);C(Input);D(_MklConv2DBackpropInput);"
"DMT/_0(Const);DMT/_1(Const);DMT/_2(Const);E(Mul)|"
"DMT/_0(Const);DMT/_1(Const);DMT/_2(Const);E(Zeta)|"
"A->D:1;A->E;B->D;B:control->DMT/_0:control;"
"B:control->DMT/_1:control;B:control->DMT/_2:control;C->D:2;"
"D->E:1;DMT/_0->D:3;DMT/_1->D:4;DMT/_2->D:5");
@ -869,11 +869,11 @@ TEST_F(MklLayoutPassTest, NodeRewrite_Concat_Basic) {
" attr { key: 'T' value { type: DT_FLOAT } }"
" attr { key: 'N' value { i: 2 } }"
" input: ['A', 'B:0', 'B:1']}"
"node { name: 'E' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'E' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['C', 'D'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Const);B(InputList);C(Input);D(_MklConcat);DMT/_0(Const);"
"DMT/_1(Const);DMT/_2(Const);E(Mul)|A->D;A:control->DMT/_0:control;"
"DMT/_1(Const);DMT/_2(Const);E(Zeta)|A->D;A:control->DMT/_0:control;"
"A:control->DMT/_1:control;A:control->DMT/_2:control;B->D:1;"
"B:1->D:2;C->E;D->E:1;DMT/_0->D:3;DMT/_1->D:4;DMT/_2->D:5");
}
@ -908,12 +908,12 @@ TEST_F(MklLayoutPassTest, NodeRewrite_Concat_Input_Mkl) {
" attr { key: 'T' value { type: DT_FLOAT } }"
" attr { key: 'N' value { i: 2 } }"
" input: ['G', 'E', 'F']}"
"node { name: 'I' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'I' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'H'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Input);D(Input);DMT/_0(Const);DMT/_1(Const);"
"DMT/_2(Const);DMT/_3(Const);DMT/_4(Const);E(_MklConv2D);"
"F(_MklConv2D);G(Const);H(_MklConcat);I(Mul)|A->E;A->I;"
"F(_MklConv2D);G(Const);H(_MklConcat);I(Zeta)|A->E;A->I;"
"A:control->DMT/_2:control;A:control->DMT/_3:control;"
"B->E:1;C->F;C:control->DMT/_0:control;C:control->DMT/_1:control;"
"D->F:1;DMT/_0->F:2;DMT/_1->F:3;DMT/_2->E:2;DMT/_3->E:3;"
@ -935,7 +935,7 @@ TEST_F(MklLayoutPassTest, NodeRewrite_Concat_Input_MixedMkl) {
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" attr { key: 'padding' value { s: 'SAME' } }"
" input: ['A', 'B']}"
"node { name: 'F' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'F' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['C', 'D']}"
"node { name: 'G' op: 'Const' "
" attr { key: 'dtype' value { type: DT_INT32 } }"
@ -946,12 +946,12 @@ TEST_F(MklLayoutPassTest, NodeRewrite_Concat_Input_MixedMkl) {
" attr { key: 'T' value { type: DT_FLOAT } }"
" attr { key: 'N' value { i: 2 } }"
" input: ['G', 'E', 'F']}"
"node { name: 'I' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'I' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'H'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Input);D(Input);DMT/_0(Const);DMT/_1(Const);"
"DMT/_2(Const);DMT/_3(Const);E(_MklConv2D);F(Mul);G(Const);"
"H(_MklConcat);I(Mul)|A->E;A->I;A:control->DMT/_0:control;"
"DMT/_2(Const);DMT/_3(Const);E(_MklConv2D);F(Zeta);G(Const);"
"H(_MklConcat);I(Zeta)|A->E;A->I;A:control->DMT/_0:control;"
"A:control->DMT/_1:control;B->E:1;C->F;D->F:1;DMT/_0->E:2;"
"DMT/_1->E:3;DMT/_2->H:3;DMT/_3->H:5;E->H:1;E:2->H:4;F->H:2;"
"G->H;G:control->DMT/_2:control;G:control->DMT/_3:control;H->I:1");
@ -973,11 +973,11 @@ TEST_F(MklLayoutPassTest, NodeRewrite_ConcatV2_Basic) {
" attr { key: 'Tidx' value { type: DT_INT32 } }"
" attr { key: 'N' value { i: 2 } }"
" input: ['B:0', 'B:1', 'A']}"
"node { name: 'E' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'E' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['C', 'D'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Const);B(InputList);C(Input);D(_MklConcatV2);DMT/_0(Const);"
"DMT/_1(Const);DMT/_2(Const);E(Mul)|A->D:2;B->D;B:1->D:1;"
"DMT/_1(Const);DMT/_2(Const);E(Zeta)|A->D:2;B->D;B:1->D:1;"
"B:control->DMT/_0:control;B:control->DMT/_1:control;"
"B:control->DMT/_2:control;C->E;D->E:1;DMT/_0->D:3;"
"DMT/_1->D:4;DMT/_2->D:5");
@ -1014,12 +1014,12 @@ TEST_F(MklLayoutPassTest, NodeRewrite_ConcatV2_Input_Mkl) {
" attr { key: 'Tidx' value { type: DT_INT32 } }"
" attr { key: 'N' value { i: 2 } }"
" input: ['E', 'F', 'G']}"
"node { name: 'I' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'I' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'H'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Input);D(Input);DMT/_0(Const);DMT/_1(Const);"
"DMT/_2(Const);DMT/_3(Const);DMT/_4(Const);E(_MklConv2D);"
"F(_MklConv2D);G(Const);H(_MklConcatV2);I(Mul)|A->E;A->I;"
"F(_MklConv2D);G(Const);H(_MklConcatV2);I(Zeta)|A->E;A->I;"
"A:control->DMT/_2:control;A:control->DMT/_3:control;B->E:1;C->F;"
"C:control->DMT/_0:control;C:control->DMT/_1:control;"
"D->F:1;DMT/_0->F:2;DMT/_1->F:3;DMT/_2->E:2;DMT/_3->E:3;"
@ -1041,7 +1041,7 @@ TEST_F(MklLayoutPassTest, NodeRewrite_ConcatV2_Input_MixedMkl) {
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" attr { key: 'padding' value { s: 'SAME' } }"
" input: ['A', 'B']}"
"node { name: 'F' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'F' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['C', 'D']}"
"node { name: 'G' op: 'Const' "
" attr { key: 'dtype' value { type: DT_INT32 } }"
@ -1053,12 +1053,12 @@ TEST_F(MklLayoutPassTest, NodeRewrite_ConcatV2_Input_MixedMkl) {
" attr { key: 'Tidx' value { type: DT_INT32 } }"
" attr { key: 'N' value { i: 2 } }"
" input: ['E', 'F', 'G']}"
"node { name: 'I' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'I' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'H'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Input);D(Input);DMT/_0(Const);DMT/_1(Const);"
"DMT/_2(Const);DMT/_3(Const);E(_MklConv2D);F(Mul);G(Const);"
"H(_MklConcatV2);I(Mul)|A->E;A->I;A:control->DMT/_0:control;"
"DMT/_2(Const);DMT/_3(Const);E(_MklConv2D);F(Zeta);G(Const);"
"H(_MklConcatV2);I(Zeta)|A->E;A->I;A:control->DMT/_0:control;"
"A:control->DMT/_1:control;B->E:1;C->F;D->F:1;DMT/_0->E:2;"
"DMT/_1->E:3;DMT/_2->H:4;DMT/_3->H:5;E->H;E:2->H:3;"
"E:control->DMT/_2:control;E:control->DMT/_3:control;F->H:1;"
@ -1071,10 +1071,10 @@ TEST_F(MklLayoutPassTest, NodeRewrite_Relu_Positive) {
"node { name: 'B' op: 'Relu'"
" attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A'] }"
"node { name: 'C' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'C' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(_MklRelu);C(Mul);DMT/_0(Const)|A->B;A->C;"
"A(Input);B(_MklRelu);C(Zeta);DMT/_0(Const)|A->B;A->C;"
"A:control->DMT/_0:control;B->C:1;DMT/_0->B:1");
}
@ -1085,10 +1085,10 @@ TEST_F(MklLayoutPassTest, NodeRewrite_ReluGrad_Positive) {
"node { name: 'C' op: 'ReluGrad'"
" attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }"
"node { name: 'D' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'D' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'C'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(_MklReluGrad);D(Mul);DMT/_0(Const);"
"A(Input);B(Input);C(_MklReluGrad);D(Zeta);DMT/_0(Const);"
"DMT/_1(Const)|A->C;A->D;A:control->DMT/_0:control;"
"A:control->DMT/_1:control;B->C:1;C->D:1;DMT/_0->C:2;DMT/_1->C:3");
}
@ -1102,10 +1102,10 @@ TEST_F(MklLayoutPassTest, NodeRewrite_ReluReluGrad_Positive) {
"node { name: 'C' op: 'ReluGrad'"
" attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }"
"node { name: 'D' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'D' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'C'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(_MklRelu);C(_MklReluGrad);D(Mul);DMT/_0(Const);"
"A(Input);B(_MklRelu);C(_MklReluGrad);D(Zeta);DMT/_0(Const);"
"DMT/_1(Const)|A->B;A->C;A->D;A:control->DMT/_0:control;"
"A:control->DMT/_1:control;B->C:1;B:1->C:3;C->D:1;DMT/_0->B:1;"
"DMT/_1->C:2");
@ -1121,10 +1121,10 @@ TEST_F(MklLayoutPassTest, NodeRewrite_AvgPool_Positive) {
" attr { key: 'padding' value { s: 'VALID' } }"
" attr { key: 'strides' value { list: {i: 1, i:1, i:2, i:2} } }"
" input: ['A'] }"
"node { name: 'C' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'C' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(_MklAvgPool);C(Mul);DMT/_0(Const)|A->B;A->C;"
"A(Input);B(_MklAvgPool);C(Zeta);DMT/_0(Const)|A->B;A->C;"
"A:control->DMT/_0:control;B->C:1;DMT/_0->B:1");
}
@ -1139,10 +1139,10 @@ TEST_F(MklLayoutPassTest, NodeRewrite_AvgPoolGrad_Positive) {
" attr { key: 'padding' value { s: 'VALID' } }"
" attr { key: 'strides' value { list: {i: 1, i:1, i:2, i:2} } }"
" input: ['A', 'B'] }"
"node { name: 'D' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'D' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['B', 'C'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Int32Input);B(Input);C(_MklAvgPoolGrad);D(Mul);DMT/_0(Const);"
"A(Int32Input);B(Input);C(_MklAvgPoolGrad);D(Zeta);DMT/_0(Const);"
"DMT/_1(Const)|A->C;A:control->DMT/_0:control;"
"A:control->DMT/_1:control;B->C:1;B->D;C->D:1;DMT/_0->C:2;"
"DMT/_1->C:3");
@ -1166,10 +1166,10 @@ TEST_F(MklLayoutPassTest, NodeRewrite_AvgPoolAvgPoolGrad_Positive) {
" attr { key: 'padding' value { s: 'VALID' } }"
" attr { key: 'strides' value { list: {i: 1, i:1, i:2, i:2} } }"
" input: ['I', 'B'] }"
"node { name: 'D' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'D' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'C'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(_MklAvgPool);C(_MklAvgPoolGrad);D(Mul);DMT/_0(Const);"
"A(Input);B(_MklAvgPool);C(_MklAvgPoolGrad);D(Zeta);DMT/_0(Const);"
"DMT/_1(Const);I(Int32Input)|A->B;A->D;A:control->DMT/_0:control;"
"B->C:1;B:1->C:3;C->D:1;DMT/_0->B:1;DMT/_1->C:2;I->C;"
"I:control->DMT/_1:control");
@ -1188,12 +1188,12 @@ TEST_F(MklLayoutPassTest, NodeRewrite_FusedBatchNormGrad_Positive) {
" attr { key: 'epsilon' value { f: 0.0001 } }"
" attr { key: 'is_training' value { b: true } }"
" input: ['A', 'B', 'C', 'D', 'E'] }"
"node { name: 'G' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'G' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'F'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Input);D(Input);DMT/_0(Const);DMT/_1(Const);"
"DMT/_2(Const);DMT/_3(Const);DMT/_4(Const);E(Input);"
"F(_MklFusedBatchNormGrad);G(Mul)|A->F;A->G;"
"F(_MklFusedBatchNormGrad);G(Zeta)|A->F;A->G;"
"A:control->DMT/_0:control;A:control->DMT/_1:control;"
"A:control->DMT/_2:control;A:control->DMT/_3:control;"
"A:control->DMT/_4:control;B->F:1;C->F:2;D->F:3;"
@ -1214,12 +1214,12 @@ TEST_F(MklLayoutPassTest, NodeRewrite_FusedBatchNorm_Positive) {
" attr { key: 'epsilon' value { f: 0.0001 } }"
" attr { key: 'is_training' value { b: true } }"
" input: ['A', 'B', 'C', 'D', 'E'] }"
"node { name: 'G' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'G' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'F'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Input);D(Input);DMT/_0(Const);DMT/_1(Const);"
"DMT/_2(Const);DMT/_3(Const);DMT/_4(Const);E(Input);"
"F(_MklFusedBatchNorm);G(Mul)|A->F;A->G;"
"F(_MklFusedBatchNorm);G(Zeta)|A->F;A->G;"
"A:control->DMT/_0:control;A:control->DMT/_1:control;"
"A:control->DMT/_2:control;A:control->DMT/_3:control;"
"A:control->DMT/_4:control;B->F:1;C->F:2;D->F:3;"
@ -1268,12 +1268,12 @@ TEST_F(MklLayoutPassTest, MaxPoolLRN_Positive) {
" attr { key: 'depth_radius' value { i: 2 } }"
" input: ['E', 'F', 'B'] }"
"node { name: 'H' op: 'Input'}"
"node { name: 'I' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'I' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['H', 'G'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(_MklLRN);C(_MklMaxPool);D(Input);DMT/_0(Const);DMT/_1(Const);"
"DMT/_2(Const);E(_MklMaxPoolGrad);F(Input);G(_MklLRNGrad);H(Input);"
"I(Mul)|A->B;A:control->DMT/_0:control;B->C;B->E;B->G:2;B:1->G:3;"
"I(Zeta)|A->B;A:control->DMT/_0:control;B->C;B->E;B->G:2;B:1->G:3;"
"B:2->C:1;B:2->E:4;B:2->G:6;B:3->G:7;B:control->DMT/_1:control;C->E:1;"
"C:1->E:3;C:2->E:5;C:3->E:7;D->E:2;DMT/_0->B:1;DMT/_1->E:6;DMT/_2->G:5;"
"E->G;E:1->G:4;E:control->DMT/_2:control;F->G:1;G->I:1;H->I");
@ -1301,11 +1301,11 @@ TEST_F(MklLayoutPassTest, LRN_Positive) {
" attr { key: 'data_format' value { s: 'NCHW' } }"
" attr { key: 'depth_radius' value { i: 2 } }"
" input: ['C', 'D', 'B'] }"
"node { name: 'F' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'F' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['C', 'E'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(_MklLRN);C(Input);D(Input);DMT/_0(Const);DMT/_1(Const);"
"DMT/_2(Const);E(_MklLRNGrad);F(Mul)|"
"DMT/_2(Const);E(_MklLRNGrad);F(Zeta)|"
"A->B;A:control->DMT/_0:control;B->E:2;B:1->E:3;B:2->E:6;B:3->E:7;"
"C->E;C->F;C:control->DMT/_1:control;C:control->DMT/_2:control;"
"D->E:1;DMT/_0->B:1;DMT/_1->E:4;DMT/_2->E:5;E->F:1");
@ -1323,10 +1323,10 @@ TEST_F(MklLayoutPassTest, LRN_Negative1) {
" attr { key: 'data_format' value { s: 'NCHW' } }"
" attr { key: 'depth_radius' value { i: 2 } }"
" input: ['A'] }"
"node { name: 'C' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'C' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(_MklLRN);C(Mul);DMT/_0(Const)|"
"A(Input);B(_MklLRN);C(Zeta);DMT/_0(Const)|"
"A->B;A->C;A:control->DMT/_0:control;B->C:1;DMT/_0->B:1");
}
@ -1344,11 +1344,11 @@ TEST_F(MklLayoutPassTest, LRN_Negative2) {
" attr { key: 'data_format' value { s: 'NCHW' } }"
" attr { key: 'depth_radius' value { i: 2 } }"
" input: ['A', 'B', 'C'] }"
"node { name: 'E' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'E' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'D'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Input);D(_MklLRNGrad);DMT/_0(Const);"
"DMT/_1(Const);DMT/_2(Const);DMT/_3(Const);DMT/_4(Const);E(Mul)|"
"DMT/_1(Const);DMT/_2(Const);DMT/_3(Const);DMT/_4(Const);E(Zeta)|"
"A->D;A->E;A:control->DMT/_0:control;A:control->DMT/_1:control;"
"A:control->DMT/_2:control;A:control->DMT/_3:control;"
"A:control->DMT/_4:control;B->D:1;C->D:2;D->E:1;DMT/_0->D:3;"
@ -1386,12 +1386,12 @@ TEST_F(MklLayoutPassTest, LRN_Negative3) {
" attr { key: 'data_format' value { s: 'NCHW' } }"
" attr { key: 'depth_radius' value { i: 2 } }"
" input: ['C', 'B', 'D'] }"
"node { name: 'G' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'G' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['E', 'F'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(_MklLRN);C(Input);D(Input);DMT/_0(Const);DMT/_1(Const);"
"DMT/_2(Const);DMT/_3(Const);DMT/_4(Const);DMT/_5(Const);"
"DMT/_6(Const);E(_MklLRNGrad);F(_MklLRNGrad);G(Mul)|A->B;"
"DMT/_6(Const);E(_MklLRNGrad);F(_MklLRNGrad);G(Zeta)|A->B;"
"A:control->DMT/_0:control;B->E:2;"
"B->F:1;B:1->E:3;B:2->E:6;B:2->F:5;B:3->E:7;C->E;C->F;"
"C:control->DMT/_1:control;C:control->DMT/_2:control;"
@ -1421,11 +1421,11 @@ TEST_F(MklLayoutPassTest, NodeWorkspace_MaxPool_Positive) {
" attr { key: 'padding' value { s: 'VALID' } }"
" attr { key: 'strides' value { list: {i: 1, i:1, i:2, i:2} } }"
" input: ['C', 'B', 'D'] }"
"node { name: 'F' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'F' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['C', 'E'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(_MklMaxPool);C(Input);D(Input);DMT/_0(Const);"
"DMT/_1(Const);DMT/_2(Const);E(_MklMaxPoolGrad);F(Mul)|"
"DMT/_1(Const);DMT/_2(Const);E(_MklMaxPoolGrad);F(Zeta)|"
"A->B;A:control->DMT/_0:control;B->E:1;B:1->E:3;B:2->E:5;B:3->E:7;"
"C->E;C->F;C:control->DMT/_1:control;C:control->DMT/_2:control;"
"D->E:2;DMT/_0->B:1;DMT/_1->E:4;DMT/_2->E:6;E->F:1");
@ -1444,10 +1444,10 @@ TEST_F(MklLayoutPassTest, NodeWorkspace_MaxPool_Negative1) {
" attr { key: 'padding' value { s: 'VALID' } }"
" attr { key: 'strides' value { list: {i: 1, i:1, i:2, i:2} } }"
" input: ['A'] }"
"node { name: 'C' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'C' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(_MklMaxPool);C(Mul);DMT/_0(Const)|"
"A(Input);B(_MklMaxPool);C(Zeta);DMT/_0(Const)|"
"A->B;A->C;A:control->DMT/_0:control;B->C:1;DMT/_0->B:1");
}
@ -1466,11 +1466,11 @@ TEST_F(MklLayoutPassTest, NodeWorkspace_MaxPool_Negative2) {
" attr { key: 'padding' value { s: 'VALID' } }"
" attr { key: 'strides' value { list: {i: 1, i:1, i:2, i:2} } }"
" input: ['A', 'B', 'C'] }"
"node { name: 'E' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'E' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'D'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Input);D(_MklMaxPoolGrad);DMT/_0(Const);"
"DMT/_1(Const);DMT/_2(Const);DMT/_3(Const);DMT/_4(Const);E(Mul)|"
"DMT/_1(Const);DMT/_2(Const);DMT/_3(Const);DMT/_4(Const);E(Zeta)|"
"A->D;A->E;A:control->DMT/_0:control;A:control->DMT/_1:control;"
"A:control->DMT/_2:control;A:control->DMT/_3:control;"
"A:control->DMT/_4:control;B->D:1;C->D:2;D->E:1;DMT/_0->D:3;"
@ -1489,10 +1489,10 @@ TEST_F(MklLayoutPassTest, NodeWorkspace_MaxPool_Negative3) {
" attr { key: 'padding' value { s: 'VALID' } }"
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" input: ['A'] }"
"node { name: 'C' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'C' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(MaxPool);C(Mul)|A->B;A->C;B->C:1");
"A(Input);B(MaxPool);C(Zeta)|A->B;A->C;B->C:1");
}
// Test MaxPool handling for batch-wise pooling (NCHW)
@ -1507,10 +1507,10 @@ TEST_F(MklLayoutPassTest, NodeWorkspace_MaxPool_Negative4) {
" attr { key: 'padding' value { s: 'VALID' } }"
" attr { key: 'strides' value { list: {i: 2, i:1, i:1, i:1} } }"
" input: ['A'] }"
"node { name: 'C' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'C' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(MaxPool);C(Mul)|A->B;A->C;B->C:1");
"A(Input);B(MaxPool);C(Zeta)|A->B;A->C;B->C:1");
}
// Test MaxPool handling for depth-wise pooling (NHWC)
@ -1525,10 +1525,10 @@ TEST_F(MklLayoutPassTest, NodeWorkspace_MaxPool_Negative5) {
" attr { key: 'padding' value { s: 'VALID' } }"
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" input: ['A'] }"
"node { name: 'C' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'C' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(MaxPool);C(Mul)|A->B;A->C;B->C:1");
"A(Input);B(MaxPool);C(Zeta)|A->B;A->C;B->C:1");
}
// Test MaxPool handling for depth-wise pooling (NCHW)
@ -1543,10 +1543,10 @@ TEST_F(MklLayoutPassTest, NodeWorkspace_MaxPool_Negative6) {
" attr { key: 'padding' value { s: 'VALID' } }"
" attr { key: 'strides' value { list: {i: 1, i:2, i:1, i:1} } }"
" input: ['A'] }"
"node { name: 'C' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'C' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(MaxPool);C(Mul)|A->B;A->C;B->C:1");
"A(Input);B(MaxPool);C(Zeta)|A->B;A->C;B->C:1");
}
// Test MaxPool handling for batch-wise pooling (NHWC)
@ -1561,10 +1561,10 @@ TEST_F(MklLayoutPassTest, NodeWorkspace_MaxPool_Negative7) {
" attr { key: 'padding' value { s: 'VALID' } }"
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" input: ['A'] }"
"node { name: 'C' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'C' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(MaxPool);C(Mul)|A->B;A->C;B->C:1");
"A(Input);B(MaxPool);C(Zeta)|A->B;A->C;B->C:1");
}
// Test MaxPool handling for batch-wise pooling (NHWC)
@ -1579,10 +1579,10 @@ TEST_F(MklLayoutPassTest, NodeWorkspace_MaxPool_Negative8) {
" attr { key: 'padding' value { s: 'VALID' } }"
" attr { key: 'strides' value { list: {i: 2, i:1, i:1, i:1} } }"
" input: ['A'] }"
"node { name: 'C' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'C' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(MaxPool);C(Mul)|A->B;A->C;B->C:1");
"A(Input);B(MaxPool);C(Zeta)|A->B;A->C;B->C:1");
}
// Test MaxPool handling for depth-wise pooling (NHWC)
@ -1597,10 +1597,10 @@ TEST_F(MklLayoutPassTest, NodeWorkspace_MaxPool_Negative9) {
" attr { key: 'padding' value { s: 'VALID' } }"
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" input: ['A'] }"
"node { name: 'C' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'C' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(MaxPool);C(Mul)|A->B;A->C;B->C:1");
"A(Input);B(MaxPool);C(Zeta)|A->B;A->C;B->C:1");
}
// Test MaxPool handling for depth-wise pooling (NHWC)
@ -1615,10 +1615,10 @@ TEST_F(MklLayoutPassTest, NodeWorkspace_MaxPool_Negative10) {
" attr { key: 'padding' value { s: 'VALID' } }"
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:2} } }"
" input: ['A'] }"
"node { name: 'C' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'C' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }");
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(MaxPool);C(Mul)|A->B;A->C;B->C:1");
"A(Input);B(MaxPool);C(Zeta)|A->B;A->C;B->C:1");
}
/////////////////////////////////////////////////////////////////////
@ -1636,10 +1636,10 @@ TEST_F(MklLayoutPassTest, NodeRewrite_Conv2D_DeviceTest) {
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" attr { key: 'padding' value { s: 'SAME' } }"
" input: ['A', 'B']}"
"node { name: 'D' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'D' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['B', 'C'] }", kGPUDevice);
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Conv2D);D(Mul)|A->C;B->C:1;B->D;C->D:1");
"A(Input);B(Input);C(Conv2D);D(Zeta)|A->C;B->C:1;B->D;C->D:1");
}
TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_DeviceTest) {
@ -1657,7 +1657,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_DeviceTest) {
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" attr { key: 'padding' value { s: 'SAME' } }"
" input: ['A', 'B', 'C', 'M', 'N', 'O']}"
"node { name: 'E' op: 'Sub'"
"node { name: 'E' op: 'Zeta'"
" attr {key: 'T' value { type: DT_FLOAT } }"
" input: ['D', 'A']}"
"node { name: 'F' op: 'BiasAddGrad'"
@ -1666,7 +1666,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DBackprop_DeviceTest) {
" input: ['E'] }", kGPUDevice);
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Input);D(_MklConv2DWithBias);"
"E(Sub);F(BiasAddGrad);M(_MklInput);N(_MklInput);"
"E(Zeta);F(BiasAddGrad);M(_MklInput);N(_MklInput);"
"O(_MklInput)|A->D;A->E:1;B->D:1;C->D:2;D->E;E->F;"
"M->D:3;N->D:4;O->D:5");
}
@ -1683,10 +1683,10 @@ TEST_F(MklLayoutPassTest, NodeRewrite_Conv2DGradFilter_DeviceTest) {
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" attr { key: 'padding' value { s: 'SAME' } }"
" input: ['A', 'B', 'C']}"
"node { name: 'E' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'E' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'D'] }", kGPUDevice);
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Int32Input);C(Input);D(Conv2DBackpropFilter);E(Mul)|"
"A(Input);B(Int32Input);C(Input);D(Conv2DBackpropFilter);E(Zeta)|"
"A->D;A->E;B->D:1;C->D:2;D->E:1");
}
@ -1696,10 +1696,10 @@ TEST_F(MklLayoutPassTest, NodeRewrite_Relu_DeviceTest) {
"node { name: 'B' op: 'Relu'"
" attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A'] }"
"node { name: 'C' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'C' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }", kGPUDevice);
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Relu);C(Mul)|A->B;A->C;B->C:1");
"A(Input);B(Relu);C(Zeta)|A->B;A->C;B->C:1");
}
TEST_F(MklLayoutPassTest, NodeRewrite_ReluGrad_DeviceTest) {
@ -1709,10 +1709,10 @@ TEST_F(MklLayoutPassTest, NodeRewrite_ReluGrad_DeviceTest) {
"node { name: 'C' op: 'ReluGrad'"
" attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }"
"node { name: 'D' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'D' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'C'] }", kGPUDevice);
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(ReluGrad);D(Mul)|A->C;A->D;B->C:1;C->D:1");
"A(Input);B(Input);C(ReluGrad);D(Zeta)|A->C;A->D;B->C:1;C->D:1");
}
TEST_F(MklLayoutPassTest, NodeRewrite_MaxPool_DeviceTest) {
@ -1725,10 +1725,10 @@ TEST_F(MklLayoutPassTest, NodeRewrite_MaxPool_DeviceTest) {
" attr { key: 'padding' value { s: 'VALID' } }"
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" input: ['A'] }"
"node { name: 'C' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'C' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }", kGPUDevice);
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(MaxPool);C(Mul)|A->B;A->C;B->C:1");
"A(Input);B(MaxPool);C(Zeta)|A->B;A->C;B->C:1");
}
TEST_F(MklLayoutPassTest, NodeRewrite_AvgPool_DeviceTest) {
@ -1741,10 +1741,10 @@ TEST_F(MklLayoutPassTest, NodeRewrite_AvgPool_DeviceTest) {
" attr { key: 'padding' value { s: 'VALID' } }"
" attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }"
" input: ['A'] }"
"node { name: 'C' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'C' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'B'] }", kGPUDevice);
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(AvgPool);C(Mul)|A->B;A->C;B->C:1");
"A(Input);B(AvgPool);C(Zeta)|A->B;A->C;B->C:1");
}
// Concat Op test: Concat with no Mkl layer feeding it
@ -1762,10 +1762,10 @@ TEST_F(MklLayoutPassTest, NodeRewrite_Concat_DeviceTest) {
" attr { key: 'T' value { type: DT_FLOAT } }"
" attr { key: 'N' value { i: 2 } }"
" input: ['A', 'B:0', 'B:1']}"
"node { name: 'E' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'E' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['C', 'D'] }", kGPUDevice);
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Const);B(InputList);C(Input);D(Concat);E(Mul)|A->D;"
"A(Const);B(InputList);C(Input);D(Concat);E(Zeta)|A->D;"
"B->D:1;B:1->D:2;C->E;D->E:1");
}
@ -1784,10 +1784,10 @@ TEST_F(MklLayoutPassTest, NodeRewrite_ConcatV2_DeviceTest) {
" attr { key: 'Tidx' value { type: DT_INT32 } }"
" attr { key: 'N' value { i: 2 } }"
" input: ['B:0', 'B:1', 'A']}"
"node { name: 'E' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'E' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['C', 'D'] }", kGPUDevice);
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Const);B(InputList);C(Input);D(ConcatV2);E(Mul)|"
"A(Const);B(InputList);C(Input);D(ConcatV2);E(Zeta)|"
"A->D:2;B->D;B:1->D:1;C->E;D->E:1");
}
@ -1804,11 +1804,11 @@ TEST_F(MklLayoutPassTest, NodeRewrite_FusedBatchNorm_DeviceTest) {
" attr { key: 'epsilon' value { f: 0.0001 } }"
" attr { key: 'is_training' value { b: true } }"
" input: ['A', 'B', 'C', 'D', 'E'] }"
"node { name: 'G' op: 'Mul' attr { key: 'T' value { type: DT_FLOAT } }"
"node { name: 'G' op: 'Zeta' attr { key: 'T' value { type: DT_FLOAT } }"
" input: ['A', 'F'] }", kGPUDevice);
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(Input);D(Input);E(Input);"
"F(FusedBatchNorm);G(Mul)|A->F;A->G;B->F:1;C->F:2;D->F:3;"
"F(FusedBatchNorm);G(Zeta)|A->F;A->G;B->F:1;C->F:2;D->F:3;"
"E->F:4;F->G:1");
}
@ -1832,12 +1832,12 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DWithBias_DeviceTest) {
" attr { key: 'data_format' value { s: 'NCHW' } }"
" input: ['C', 'D'] }"
"node { name: 'Y' op: 'Input'}"
"node { name: 'Z' op: 'Sub'"
"node { name: 'Z' op: 'Zeta'"
" attr {key: 'T' value { type: DT_FLOAT } }"
" input: ['E', 'Y']}", kGPUDevice);
EXPECT_EQ(DoMklLayoutOptimizationPass(),
"A(Input);B(Input);C(_MklConv2D);D(Input);E(BiasAdd);"
"M(_MklInput);N(_MklInput);Y(Input);Z(Sub)|A->C;"
"M(_MklInput);N(_MklInput);Y(Input);Z(Zeta)|A->C;"
"B->C:1;C->E;D->E:1;E->Z;M->C:2;N->C:3;Y->Z:1");
}
@ -1853,7 +1853,7 @@ static void BM_MklLayoutRewritePass(int iters, int op_nodes) {
random::SimplePhilox rnd(&philox);
for (int op = 0; op < op_nodes; op++) {
s += strings::Printf(
"node { name: 'op%04d' op: 'Mul' attr { key: 'T' value { "
"node { name: 'op%04d' op: 'Zeta' attr { key: 'T' value { "
"type: DT_FLOAT } } input: ['in%04d', 'in%04d' ] }",
op, rnd.Uniform(10), rnd.Uniform(10));
}

159
tensorflow/core/graph/mkl_tfconversion_pass.cc
View File

@ -64,6 +64,15 @@ namespace tensorflow {
// in the Mkl format. Non-compliant ops accept inputs and outputs in the
// TensorFlow format.
//
// ADDENDUM: For element-wise ops, we may or may not need a conversion to
// take place before we hit the op. For this, we add a new op before each
// element-wise MKL op to deal with the inputs, called _MklInputConversion.
// This pass has been enhanced to add this capability.
//
// The _MklInputConversion op will check the inputs to the elementwise op and
// make sure that either both are in MKL format or both are in TF format,
// depending on their initial state and whether broadcast is needed or not.
class MklToTfConversionPass : public GraphOptimizationPass {
public:
MklToTfConversionPass() {}
@ -87,6 +96,16 @@ class MklToTfConversionPass : public GraphOptimizationPass {
return mkl_op_registry::IsMklOp(op_name, T);
}
// Is the input Op supported by Mkl-specific layout AND
// is it element-wise?
//
// @input op_name string of the op
// @input T Datatype to use for checking input op
// @return true if op is Mkl supported; false, otherwise.
inline bool IsMklElementWiseOp(const string& op_name, DataType T) const {
return mkl_op_registry::IsMklElementWiseOp(op_name, T);
}
// Insert layout conversion node on the edge pointed by 'e' from graph 'g'.
//
// Edge will be deleted once a call to this function is successful.
@ -96,6 +115,17 @@ class MklToTfConversionPass : public GraphOptimizationPass {
// @return Success:OK() if insertion is successful, otherwise returns
// appropriate error status code.
Status InsertConversionNodeOnEdge(std::unique_ptr<Graph>* g, Edge*);
// For element-wise ops, we need to sanitize the inputs. For this, we add a
// new node at the input of the replacement element-wise node that checks
// the inputs and converts one/both of them as required. See the op code
// comments for details.
//
// Insert input conversion node as parent of 'n' from graph 'g'.
//
// @return Success:OK() if insertion is successful, otherwise returns
// appropriate error status code.
Status InsertInputConversionNode(std::unique_ptr<Graph>* g, Node*);
};
// We register MklToTf insertion for phase 2 in post-partition grouping
@ -171,6 +201,92 @@ Status MklToTfConversionPass::InsertConversionNodeOnEdge(
return Status::OK();
}
Status MklToTfConversionPass::InsertInputConversionNode(
std::unique_ptr<Graph>* g, Node* n) {
CHECK_NOTNULL(n);
// Get the input nodes and edges
std::vector<const Edge*> edges;
TF_CHECK_OK(n->input_edges(&edges));
if (edges.size() != 4) {
return Status(error::Code::INVALID_ARGUMENT,
"MKL Binary Element-wise op should have exactly 2 data"
" inputs and 2 metadata inputs");
}
// Sanity check: ensure that both inputs are of the expected type, and the
// same type as input type
CHECK_EQ(BaseType(edges[0]->src()->output_type(edges[0]->src_output())),
BaseType(edges[1]->src()->output_type(edges[1]->src_output())));
CHECK_EQ(BaseType(edges[0]->src()->output_type(edges[0]->src_output())),
BaseType(n->input_type(0)));
// Check ordering of edges
for (uint i = 0; i < 4; i++) {
CHECK_EQ((edges[i]->dst_input() == i), true);
}
// Build the conversion node and specify src as input.
Node* conversion_node = nullptr;
TF_CHECK_OK(
NodeBuilder((*g)->NewName("MklInputConversion"), "_MklInputConversion")
.Input(edges[0]->src(), edges[0]->src_output())
.Input(edges[1]->src(), edges[1]->src_output())
.Input(edges[2]->src(), edges[2]->src_output())
.Input(edges[3]->src(), edges[3]->src_output())
.Device(n->def().device())
.Attr("T", n->input_type(0))
.Finalize(&**g, &conversion_node));
CHECK_NOTNULL(conversion_node);
// Change the destination of any control edges to the InputConversion node
if (edges.size() != n->in_edges().size()) {
std::vector<const Edge*> edges_to_remove;
for (const Edge* e : n->in_edges()) {
if (e->IsControlEdge()) {
CHECK_NOTNULL((*g)->AddControlEdge(e->src(), conversion_node));
edges_to_remove.push_back(e);
}
}
for (const Edge* e : edges_to_remove) {
(*g)->RemoveEdge(e);
}
}
string data_format;
if (GetNodeAttr(edges[0]->src()->def(), "data_format", &data_format) ==
Status::OK()) {
conversion_node->AddAttr("data_format", data_format);
}
// Get assigned device from destination node and apply it to conversion node.
// We want conversion node to be on the same device as the destination node.
conversion_node->set_assigned_device_name(n->assigned_device_name());
// Set the Mkl op label for this op.
conversion_node->AddAttr("_kernel", mkl_op_registry::kMklOpLabel);
// Now that we have added edges from src->conversion_node, let's add edge from
// output of conversion_node to the element-wise node.
CHECK_NOTNULL((*g)->AddEdge(conversion_node, 0, n, edges[0]->dst_input()));
CHECK_NOTNULL((*g)->AddEdge(conversion_node, 1, n, edges[1]->dst_input()));
CHECK_NOTNULL((*g)->AddEdge(conversion_node, 2, n, edges[2]->dst_input()));
CHECK_NOTNULL((*g)->AddEdge(conversion_node, 3, n, edges[3]->dst_input()));
VLOG(1) << "MklToTfConversionPass - InputConversion: Inserting input "
<< "conversion node on: " << n->type_string() << " successful.";
// Remove src->dst edge now.
(*g)->RemoveEdge(edges[0]);
(*g)->RemoveEdge(edges[1]);
(*g)->RemoveEdge(edges[2]);
(*g)->RemoveEdge(edges[3]);
return Status::OK();
}
bool MklToTfConversionPass::RunPass(std::unique_ptr<Graph>* g) {
bool result = false;
@ -239,6 +355,49 @@ bool MklToTfConversionPass::RunPass(std::unique_ptr<Graph>* g) {
DumpGraph("After MklToTfConversionPass", &**g);
//---------------------------------------------------------------------------
// Check all nodes and add an input-conversion-node if the node is an mkl
// element-wise node.
VLOG(1) << "Before running MklToTfConversionPass - InputConversion";
std::vector<Node*> candidate_nodes;
std::vector<Node*> order;
GetReversePostOrder(**g, &order); // This will give us topological sort.
for (Node* n : order) {
// If node is not an op or it does not have a datatype, then skip.
DataType datatype;
if (!n->IsOp() || (GetNodeAttr(n->def(), "T", &datatype) != Status::OK())) {
continue;
}
if (IsMklElementWiseOp(n->type_string(), datatype)) {
// If the input node is an input-conversion op, skip
Node* input_node = nullptr;
TF_CHECK_OK(n->input_node(0, &input_node));
DataType input_datatype;
if ((GetNodeAttr(n->def(), "T", &input_datatype) == Status::OK()) &&
(input_node->type_string().compare("_MklInputConversion") == 0)) {
continue;
}
VLOG(1) << "MklToTfConversionPass: InputConversion: Scheduled node "
<< n->name() << " for inserting input conversion node";
candidate_nodes.push_back(const_cast<Node*>(n));
}
}
// Process all candidate edges and insert conversion nodes on them.
for (Node* n : candidate_nodes) {
// Even if we insert conversion node on a single node, we
// need to return true.
if (InsertInputConversionNode(g, n) == Status::OK()) {
VLOG(1) << "MklToTfConversionPass: Inserted conversion "
<< "on node " << n->name();
result = true;
}
}
DumpGraph("After MklToTfConversionPass - InputConversion", &**g);
// We need to return true even if we insert one conversion node
// anywhere in the graph.
return result;

45
tensorflow/core/kernels/BUILD
View File

@ -2340,7 +2340,10 @@ tf_kernel_library(
tf_kernel_library(
name = "svd_op",
prefix = "svd_op",
deps = LINALG_DEPS,
deps = LINALG_DEPS + if_cuda([
":cuda_solvers",
":transpose_functor",
]),
)
cc_library(
@ -2938,7 +2941,7 @@ tf_kernel_library(
"//tensorflow/core:framework",
"//tensorflow/core:lib",
"//tensorflow/core:nn_ops_op_lib",
],
] + if_cuda(["@cub_archive//:cub"]),
)
tf_kernel_library(
@ -5501,6 +5504,22 @@ tf_mkl_kernel_library(
],
)
tf_mkl_kernel_library(
name = "mkl_input_conversion_op",
hdrs = ["mkl_tfconv_op.h"],
prefix = "mkl_input_conversion",
deps = [
":bounds_check",
":ops_util",
"//tensorflow/core:core_cpu",
"//tensorflow/core:framework",
"//tensorflow/core:lib",
"//tensorflow/core:lib_internal",
"//tensorflow/core:nn_ops_op_lib",
"//third_party/mkl:intel_binary_blob",
],
)
tf_mkl_kernel_library(
name = "mkl_pooling_ops",
srcs = [
@ -5543,6 +5562,14 @@ tf_mkl_kernel_library(
],
)
tf_mkl_kernel_library(
name = "mkl_aggregate_ops",
prefix = "mkl_aggregate_ops",
deps = MATH_DEPS + [
"//third_party/mkl:intel_binary_blob",
],
)
tf_mkl_kernel_library(
name = "mkl_concat_op",
prefix = "mkl_concat_op",
@ -5575,6 +5602,20 @@ tf_mkl_kernel_library(
],
)
tf_mkl_kernel_library(
name = "mkl_cwise_ops_common",
hdrs = [
"cwise_ops.h",
"cwise_ops_common.h",
"cwise_ops_gradients.h",
],
prefix = "mkl_cwise_ops_common",
deps = NN_DEPS + [
"cwise_op",
"//third_party/mkl:intel_binary_blob",
],
)
cc_library(
name = "dataset",
srcs = ["dataset.cc"],

23
tensorflow/core/kernels/bias_op_gpu.cu.cc
View File

@ -173,15 +173,20 @@ __global__ void BiasGradNCHW_SharedAtomics(const T* output_backprop,
// Accumulate the results in the shared memory into the first element.
// No syncthreads is needed since this is only in the same warp.
int32 thread_index = threadIdx.x;
if (thread_index < 16) s_data[thread_index] += s_data[thread_index + 16];
if (thread_index < 8) s_data[thread_index] += s_data[thread_index + 8];
if (thread_index < 4) s_data[thread_index] += s_data[thread_index + 4];
if (thread_index < 2) s_data[thread_index] += s_data[thread_index + 2];
if (thread_index < 1) s_data[thread_index] += s_data[thread_index + 1];
// The first thread writes out the accumulated result to the global location.
if (thread_index == 0) {
CudaAtomicAdd(bias_backprop + bias_index, T(s_data[0]));
if (thread_index < 16) {
s_data[thread_index] += s_data[thread_index + 16];
__syncwarp(0xFFFF);
if (thread_index < 8) s_data[thread_index] += s_data[thread_index + 8];
__syncwarp(0xFF);
if (thread_index < 4) s_data[thread_index] += s_data[thread_index + 4];
__syncwarp(0xF);
if (thread_index < 2) s_data[thread_index] += s_data[thread_index + 2];
__syncwarp(0x3);
if (thread_index == 0) {
T val = T(s_data[0] + s_data[1]);
// The first thread writes out the accumulated result to global location.
CudaAtomicAdd(bias_backprop + bias_index, val);
}
}
}

38
tensorflow/core/kernels/cuda_solvers.cc
View File

@ -174,7 +174,7 @@ Status CudaSolver::CopyLapackInfoToHostAsync(
}
info_checker_callback(status, host_lapack_infos);
};
auto cb =
std::bind(wrapped_info_checker_callback, context_,
std::move(info_checker_callback), std::move(host_lapack_infos));
@ -188,6 +188,7 @@ Status CudaSolver::CopyLapackInfoToHostAsync(
// numeric types.
#define TF_CALL_LAPACK_TYPES(m) \
m(float, S) m(double, D) m(std::complex<float>, C) m(std::complex<double>, Z)
#define TF_CALL_LAPACK_TYPES_NO_COMPLEX(m) m(float, S) m(double, D)
// Macros to construct cusolverDn method names.
#define DN_SOLVER_FN(method, lapack_prefix) cusolverDn##lapack_prefix##method
@ -327,6 +328,41 @@ static inline Status GetrsImpl(SolverFnT solver, OpKernelContext* context,
TF_CALL_LAPACK_TYPES(GETRS_INSTANCE);
template <typename Scalar, typename BufSizeFnT, typename SolverFnT>
static inline Status GesvdImpl(BufSizeFnT bufsize, SolverFnT solver,
OpKernelContext* context,
cusolverDnHandle_t cusolver_dn_handle,
signed char jobu, signed char jobvt, int m,
int n, Scalar* A, int lda, Scalar* S, Scalar* U,
int ldu, Scalar* VT, int ldvt,
int* dev_lapack_info) {
/* Get amount of workspace memory required. */
int lwork;
TF_RETURN_IF_CUSOLVER_ERROR(bufsize(cusolver_dn_handle, m, n, &lwork));
/* Allocate device memory for workspace. */
ScratchSpace<Scalar> dev_workspace(context, lwork, /* on_host */ false);
/* Launch the solver kernel. */
TF_RETURN_IF_CUSOLVER_ERROR(solver(
cusolver_dn_handle, jobu, jobvt, m, n, CUDAComplex(A), lda, S,
CUDAComplex(U), ldu, CUDAComplex(VT), ldvt,
CUDAComplex(dev_workspace.mutable_data()), lwork, NULL, dev_lapack_info));
return Status::OK();
}
#define GESVD_INSTANCE(Scalar, lapack_prefix) \
template <> \
Status CudaSolver::Gesvd<Scalar>( \
signed char jobu, signed char jobvt, int m, int n, Scalar* dev_A, \
int lda, Scalar* dev_S, Scalar* dev_U, int ldu, Scalar* dev_VT, \
int ldvt, int* dev_lapack_info) const { \
return GesvdImpl(DN_BUFSIZE_FN(gesvd, lapack_prefix), \
DN_SOLVER_FN(gesvd, lapack_prefix), context_, \
cusolver_dn_handle_, jobu, jobvt, m, n, dev_A, lda, \
dev_S, dev_U, ldu, dev_VT, ldvt, dev_lapack_info); \
}
TF_CALL_LAPACK_TYPES_NO_COMPLEX(GESVD_INSTANCE);
//=============================================================================
// Wrappers of cuBlas computational methods begin here.
//

16
tensorflow/core/kernels/cuda_solvers.h
View File

@ -258,13 +258,23 @@ class CudaSolver {
Status Syevd(cusolverEigMode_t jobz, cublasFillMode_t uplo, int n, Scalar*
dev_A, int lda, Scalar* dev_W, int* dev_lapack_info) const;
*/
// Singular value decomposition.
// See: http://docs.nvidia.com/cuda/cusolver/#cuds-lt-t-gt-gesvd
template <typename Scalar>
Status Gesvd(signed char jobu, signed char jobvt, int m, int n, Scalar* dev_A,
int lda, Scalar* dev_S, Scalar* dev_U, int ldu, Scalar* dev_VT,
int ldvt, int* dev_lapack_info);
*/
int lda, Scalar* dev_S, Scalar* dev_U, int ldu, Scalar* dev_VT,
int ldvt, int* dev_lapack_info) const;
/*
// Batched linear solver using LU factorization from getrfBatched.
// See:
http://docs.nvidia.com/cuda/cublas/index.html#cublas-lt-t-gt-getrsbatched
template <typename Scalar>
Status GetrsBatched(cublasOperation_t trans, int n, int nrhs,
const Scalar* dev_Aarray[], int lda, const int* devIpiv,
Scalar* dev_Barray[], int ldb, int* info, int batch_size)
const;
*/
private:
OpKernelContext* context_; // not owned.

4
tensorflow/core/kernels/cwise_ops.h
View File

@ -139,7 +139,7 @@ struct scalar_left : private Binary {
typedef Tout result_type;
const Tin* left;
EIGEN_DEVICE_FUNC inline scalar_left(const scalar_left& other) = default;
inline scalar_left(const scalar_left& other) = default;
template <typename... Args>
EIGEN_DEVICE_FUNC inline explicit scalar_left(const Tin* c, Args... args)
@ -169,7 +169,7 @@ struct scalar_right : private Binary {
typedef Tout result_type;
const Tin* right;
EIGEN_DEVICE_FUNC inline scalar_right(const scalar_right& other) = default;
inline scalar_right(const scalar_right& other) = default;
template <typename... Args>
EIGEN_DEVICE_FUNC inline explicit scalar_right(const Tin* c, Args... args)

2
tensorflow/core/kernels/cwise_ops_common.cc
View File

@ -20,7 +20,9 @@ namespace tensorflow {
BinaryOpShared::BinaryOpShared(OpKernelConstruction* ctx, DataType out,
DataType in)
: OpKernel(ctx) {
#ifndef INTEL_MKL
OP_REQUIRES_OK(ctx, ctx->MatchSignature({in, in}, {out}));
#endif
}
void BinaryOpShared::SetUnimplementedError(OpKernelContext* ctx) {

1
tensorflow/core/kernels/decode_raw_op.cc
View File

@ -105,6 +105,7 @@ REGISTER(Eigen::half);
REGISTER(float);
REGISTER(double);
REGISTER(int32);
REGISTER(uint16);
REGISTER(uint8);
REGISTER(int16);
REGISTER(int8);

30
tensorflow/core/kernels/depthwise_conv_op_gpu.cu.cc
View File

@ -22,6 +22,7 @@ limitations under the License.
#include "tensorflow/core/platform/types.h"
#include "tensorflow/core/util/cuda_kernel_helper.h"
#include "tensorflow/core/util/tensor_format.h"
#include "external/cub_archive/cub/util_ptx.cuh"
#if !defined(_MSC_VER)
#define UNROLL _Pragma("unroll")
@ -1015,6 +1016,21 @@ __global__ void __launch_bounds__(640, 2)
}
}
// Device function to compute sub-warp sum reduction for a power-of-two group of
// neighboring threads.
template<int kWidth, typename T>
__device__ __forceinline__ T WarpSumReduce(T val) {
// support only power-of-two widths.
assert(__popc(kWidth) == 1);
int sub_warp = cub::LaneId() / kWidth;
int zeros = sub_warp * kWidth;
unsigned mask = ((1UL << kWidth) - 1) << zeros;
for (int delta = kWidth / 2; delta > 0; delta /= 2) {
val += CudaShuffleXor(mask, val, delta);
}
return val;
}
// CUDA kernel to compute the depthwise convolution backward w.r.t. filter in
// NHWC format, tailored for small images up to 32x32. Stride and depth
// multiplier must be 1. Padding must be 'SAME'. Only use this kernel if
@ -1127,6 +1143,7 @@ __launch_bounds__(1024, 2) void DepthwiseConv2dBackpropFilterGPUKernelNHWCSmall(
// Note: the condition to reach this is uniform across the entire block.
__syncthreads();
unsigned active_threads = CudaBallot(CUDA_WARP_ALL, depth_in_range);
if (depth_in_range) {
const T* const out_ptr = inout_offset + output;
@ -1140,7 +1157,7 @@ __launch_bounds__(1024, 2) void DepthwiseConv2dBackpropFilterGPUKernelNHWCSmall(
T val = out1 * tile_ptr[0] + out2 * tile_ptr[tile_offset];
// Warp-accumulate pixels of the same depth and write to accumulator.
for (int delta = 16; delta >= kBlockSlices; delta /= 2) {
val += CudaShuffleDown(val, delta);
val += CudaShuffleDown(active_threads, val, delta);
}
if (!(thread_idx & 32 - kBlockSlices) /* lane_idx < kBlockSlices */) {
*accum_ptr = val;
@ -1164,9 +1181,7 @@ __launch_bounds__(1024, 2) void DepthwiseConv2dBackpropFilterGPUKernelNHWCSmall(
if (filter_depth < in_depth) {
T val = accum_data[i];
// Warp-accumulate the pixels of the same depth from the accumulator.
for (int delta = kAccumPixels / 2; delta > 0; delta /= 2) {
val += CudaShuffleDown(val, delta);
}
val = WarpSumReduce<kAccumPixels>(val);
if (!(thread_idx & kAccumPixels - 1)) {
CudaAtomicAdd(filter_offset + filter, val);
}
@ -1382,6 +1397,7 @@ __launch_bounds__(1024, 2) void DepthwiseConv2dBackpropFilterGPUKernelNCHWSmall(
// Note: the condition to reach this is uniform across the entire block.
__syncthreads();
unsigned active_threads = CudaBallot(CUDA_WARP_ALL, slice_in_range);
if (slice_in_range) {
const T* const out_ptr = inout_offset + output;
@ -1395,7 +1411,7 @@ __launch_bounds__(1024, 2) void DepthwiseConv2dBackpropFilterGPUKernelNCHWSmall(
T val = out1 * tile_ptr[0] + out2 * tile_ptr[tile_offset];
// Warp-accumulate pixels of the same depth and write to accumulator.
for (int delta = 16 / kBlockSlices; delta > 0; delta /= 2) {
val += CudaShuffleDown(val, delta);
val += CudaShuffleDown(active_threads, val, delta);
}
if (!(thread_idx & 32 / kBlockSlices - 1)) {
*accum_ptr = val;
@ -1419,9 +1435,7 @@ __launch_bounds__(1024, 2) void DepthwiseConv2dBackpropFilterGPUKernelNCHWSmall(
if (filter_depth < in_depth) {
T val = accum_data[i];
// Warp-accumulate pixels of the same depth from the accumulator.
for (int delta = kAccumPixels / 2; delta > 0; delta /= 2) {
val += CudaShuffleDown(val, delta);
}
val = WarpSumReduce<kAccumPixels>(val);
if (!(thread_idx & kAccumPixels - 1)) {
CudaAtomicAdd(filter_offset + filter, val);
}

2
tensorflow/core/kernels/fill_functor.cc
View File

@ -20,6 +20,7 @@ limitations under the License.
#include "third_party/eigen3/unsupported/Eigen/CXX11/Tensor"
#include "tensorflow/core/framework/tensor_types.h"
#include "tensorflow/core/framework/types.h"
#include "tensorflow/core/framework/variant_encode_decode.h"
namespace tensorflow {
namespace functor {
@ -50,6 +51,7 @@ DEFINE_SETZERO_CPU(int32);
DEFINE_SETZERO_CPU(int64);
DEFINE_SETZERO_CPU(complex64);
DEFINE_SETZERO_CPU(complex128);
DEFINE_SETZERO_CPU(Variant);
#undef DEFINE_SETZERO_CPU
#ifdef TENSORFLOW_USE_SYCL

273
tensorflow/core/kernels/mkl_aggregate_ops.cc Normal file
View File

@ -0,0 +1,273 @@
/* 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.
==============================================================================*/
// See docs in ../ops/math_ops.cc.
#ifdef INTEL_MKL
#define EIGEN_USE_THREADS
#include <numeric>
#include "tensorflow/core/framework/numeric_op.h"
#include "tensorflow/core/framework/register_types.h"
#include "tensorflow/core/lib/gtl/inlined_vector.h"
#include "tensorflow/core/platform/logging.h"
#include "mkl_dnn.h"
#include "mkl_dnn_types.h"
#include "tensorflow/core/util/mkl_util.h"
namespace tensorflow {
typedef Eigen::ThreadPoolDevice CPUDevice;
template <typename Device, typename T>
class MklAddNOp : public OpKernel {
public:
explicit MklAddNOp(OpKernelConstruction* context) : OpKernel(context) {}
void Compute(OpKernelContext* ctx) override {
const int num = ctx->num_inputs();
OP_REQUIRES(ctx, num / 2 == 2,
errors::InvalidArgument("Only additions of two arguments "
"supported by MKL. Num inputs: ",
num));
MklAddNOpContext mkl_context;
const Tensor& input0 = MklGetInput(ctx, 0);
GetMklShape(ctx, 0, &(mkl_context.input1_shape));
bool input1_in_mkl_format = mkl_context.input1_shape.IsMklTensor();
const Tensor& input1 = MklGetInput(ctx, 1);
GetMklShape(ctx, 1, &(mkl_context.input2_shape));
bool input2_in_mkl_format = mkl_context.input2_shape.IsMklTensor();
mkl_context.in_dims = input1_in_mkl_format
? mkl_context.input1_shape.GetDimension()
: input0.dims();
mkl_context.in_dims = input2_in_mkl_format
? mkl_context.input2_shape.GetDimension()
: input1.dims();
// Generate size, stride for input if input is in MKL format.
ExtractMklOpParams(&mkl_context.in1_sizes,
&mkl_context.in1_strides, input0, &mkl_context.input1_shape);
ExtractMklOpParams(&mkl_context.in2_sizes,
&mkl_context.in2_strides, input1, &mkl_context.input2_shape);
std::vector<float> coeff(2, 1.0);
mkl_context.MklCreateInputLayouts(ctx);
CHECK_EQ(dnnSumCreate_F32(&mkl_context.Eltwise, mkl_context.attributes, 2,
mkl_context.lt_input1, &coeff[0]),
E_SUCCESS);
Tensor mkl_tmp_input1_buf_tensor, mkl_tmp_input2_buf_tensor;
mkl_context.MklPrepareAddNInputs(ctx, &mkl_tmp_input1_buf_tensor,
&mkl_tmp_input2_buf_tensor);
Tensor* output = nullptr;
if (input1_in_mkl_format || input2_in_mkl_format) {
TensorShape tf_shape;
mkl_context.output_shape.SetMklTensor(true);
mkl_context.output_shape.SetMklLayout(mkl_context.Eltwise, dnnResourceDst);
mkl_context.output_shape.SetTfLayout(
mkl_context.in_dims, mkl_context.in1_sizes, mkl_context.in1_strides);
if (input1_in_mkl_format == true) {
mkl_context.output_shape.SetTfDimOrder(mkl_context.in_dims,
mkl_context.input1_shape.GetTfToMklDimMap());
} else {
mkl_context.output_shape.SetTfDimOrder(mkl_context.in_dims,
mkl_context.input2_shape.GetTfToMklDimMap());
}
tf_shape.AddDim(dnnLayoutGetMemorySize_F32(static_cast<dnnLayout_t>(
mkl_context.output_shape.GetMklLayout())) /
sizeof(T));
AllocateOutputSetMklShape(ctx, 0, &output, tf_shape,
mkl_context.output_shape);
} else {
const TensorShape& o_shape = input1.shape();
mkl_context.output_shape.SetMklTensor(false);
AllocateOutputSetMklShape(ctx, 0, &output, o_shape,
mkl_context.output_shape);
}
mkl_context.Eltwise_res[dnnResourceDst] =
static_cast<void*>(output->flat<T>().data());
// Execute convolution
CHECK_EQ(dnnExecute_F32(mkl_context.Eltwise, mkl_context.Eltwise_res),
E_SUCCESS);
mkl_context.MklCleanup();
}
void ExtractMklOpParams(size_t** out_sizes, size_t** out_strides,
const Tensor& input, const MklShape* input_shape) {
bool input_in_mkl_format = input_shape->IsMklTensor();
int in_dims = input_in_mkl_format
? input_shape->GetDimension()
: input.dims();
size_t* in_sizes = new size_t[in_dims];
size_t* in_strides = new size_t[in_dims];
if (input_in_mkl_format) {
for (int i = 0; i < in_dims; i++) {
in_sizes[i] = input_shape->GetSizes()[i];
in_strides[i] = input_shape->GetStrides()[i];
}
} else {
for (int i = 0; i < in_dims; i++) {
in_sizes[i] =
input.dim_size((in_dims - 1) - i);
}
in_strides[0] = 1;
for (int i = 1; i < in_dims; i++) {
in_strides[i] =
in_strides[i - 1] * in_sizes[i - 1];
}
}
*out_sizes = in_sizes;
*out_strides = in_strides;
}
private:
typedef struct {
int in_dims;
size_t* in1_sizes;
size_t* in1_strides;
size_t* in2_sizes;
size_t* in2_strides;
dnnPrimitive_t Eltwise = nullptr;
dnnPrimitiveAttributes_t attributes = nullptr;
void* Eltwise_res[dnnResourceNumber];
dnnLayout_t lt_input1 = nullptr, lt_input2 = nullptr;
MklShape input1_shape, input2_shape, output_shape;
void MklCreateInputLayouts(OpKernelContext* context) {
bool input1_in_mkl_format = input1_shape.IsMklTensor();
if (!input1_in_mkl_format) {
CHECK_EQ(
dnnLayoutCreate_F32(&lt_input1, in_dims, in1_sizes, in1_strides),
E_SUCCESS);
} else {
lt_input1 = static_cast<dnnLayout_t>(input1_shape.GetCurLayout());
}
bool input2_in_mkl_format = input2_shape.IsMklTensor();
if (!input2_in_mkl_format) {
CHECK_EQ(
dnnLayoutCreate_F32(&lt_input2, in_dims, in2_sizes, in2_strides),
E_SUCCESS);
} else {
lt_input2 = static_cast<dnnLayout_t>(input2_shape.GetCurLayout());
}
}
void MklPrepareAddNInputs(OpKernelContext* context,
Tensor* mkl_tmp_input1_buf_tensor,
Tensor* mkl_tmp_input2_buf_tensor) {
bool mkl_convert_input1, mkl_convert_input2;
dnnPrimitive_t mkl_prim_convert_input1 = nullptr,
mkl_prim_convert_input2 = nullptr;
dnnLayout_t mkl_lt_internal_input1 = nullptr,
mkl_lt_internal_input2 = nullptr;
void *mkl_buf_convert_input1 = nullptr, *mkl_buf_convert_input2 = nullptr;
dnnResourceType_t dnnResourceMultipleSrc2 =
(dnnResourceType_t)(dnnResourceMultipleSrc + 1);
// Compare with internal layouts and convert if needed
const Tensor& input1 = MklGetInput(context, 0);
void* mkl_buf_input1 =
const_cast<void*>(static_cast<const void*>(input1.flat<T>().data()));
CHECK_EQ(dnnLayoutCreateFromPrimitive_F32(
&mkl_lt_internal_input1, Eltwise, dnnResourceMultipleSrc),
E_SUCCESS);
mkl_convert_input1 =
!dnnLayoutCompare_F32(mkl_lt_internal_input1, lt_input1);
if (mkl_convert_input1) {
CHECK_EQ(dnnConversionCreate_F32(&mkl_prim_convert_input1, lt_input1,
mkl_lt_internal_input1),
E_SUCCESS);
AllocTmpBuffer(context, mkl_tmp_input1_buf_tensor,
mkl_lt_internal_input1, &mkl_buf_convert_input1);
CHECK_EQ(
dnnConversionExecute_F32(mkl_prim_convert_input1, mkl_buf_input1,
mkl_buf_convert_input1),
E_SUCCESS);
dnnDelete_F32(mkl_prim_convert_input1);
}
dnnLayoutDelete_F32(mkl_lt_internal_input1);
Eltwise_res[dnnResourceMultipleSrc] =
(mkl_convert_input1) ? mkl_buf_convert_input1 : mkl_buf_input1;
const Tensor& input2 = MklGetInput(context, 1);
void* mkl_buf_input2 =
const_cast<void*>(static_cast<const void*>(input2.flat<T>().data()));
CHECK_EQ(dnnLayoutCreateFromPrimitive_F32(
&mkl_lt_internal_input2, Eltwise, dnnResourceMultipleSrc2),
E_SUCCESS);
mkl_convert_input2 =
!dnnLayoutCompare_F32(mkl_lt_internal_input2, lt_input2);
if (mkl_convert_input2) {
CHECK_EQ(dnnConversionCreate_F32(&mkl_prim_convert_input2, lt_input2,
mkl_lt_internal_input2),
E_SUCCESS);
AllocTmpBuffer(context, mkl_tmp_input2_buf_tensor,
mkl_lt_internal_input2, &mkl_buf_convert_input2);
CHECK_EQ(
dnnConversionExecute_F32(mkl_prim_convert_input2, mkl_buf_input2,
mkl_buf_convert_input2),
E_SUCCESS);
dnnDelete_F32(mkl_prim_convert_input2);
}
dnnLayoutDelete_F32(mkl_lt_internal_input2);
Eltwise_res[dnnResourceMultipleSrc2] =
(mkl_convert_input2) ? mkl_buf_convert_input2 : mkl_buf_input2;
}
void MklCleanup() {
bool input1_in_mkl_format = input1_shape.IsMklTensor();
bool input2_in_mkl_format = input2_shape.IsMklTensor();
dnnDelete_F32(Eltwise);
if (!input1_in_mkl_format) {
dnnLayoutDelete_F32(lt_input1);
delete [] in1_sizes;
delete [] in1_strides;
}
if (!input2_in_mkl_format) {
dnnLayoutDelete_F32(lt_input2);
delete [] in2_sizes;
delete [] in2_strides;
}
}
} MklAddNOpContext;
};
#define REGISTER_MKL_CPU(T) \
REGISTER_KERNEL_BUILDER(Name("_MklAddN") \
.Device(DEVICE_CPU) \
.TypeConstraint<T>("T") \
.Label(mkl_op_registry::kMklOpLabel), \
MklAddNOp<CPUDevice, T>);
TF_CALL_float(REGISTER_MKL_CPU);
#undef REGISTER_MKL_CPU
} // namespace tensorflow
#endif // INTEL_MKL

2
tensorflow/core/kernels/mkl_conv_ops.cc
View File

@ -406,8 +406,10 @@ class MklConv2DOp : public OpKernel {
CHECK_EQ(dnnConversionCreate_F32(&mkl_prim_convert_filter, lt_filter,
mkl_lt_internal_filter),
E_SUCCESS);
mkl_buf_convert_filter = const_cast<void*>(
static_cast<const void*>(output_filter->flat<T>().data()));
CHECK_EQ(
dnnConversionExecute_F32(mkl_prim_convert_filter, mkl_buf_filter,
mkl_buf_convert_filter),

88
tensorflow/core/kernels/mkl_cwise_ops_common.cc Normal file
View File

@ -0,0 +1,88 @@
/* 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.
==============================================================================*/
#ifdef INTEL_MKL
// See docs in ../ops/math_ops.cc.
#define EIGEN_USE_THREADS
#include <iostream>
#include <vector>
#include "tensorflow/core/kernels/cwise_ops_common.h"
#include "tensorflow/core/util/mkl_util.h"
namespace tensorflow {
typedef Eigen::ThreadPoolDevice CPUDevice;
template <typename Device, typename Functor>
class MklBinaryOp : public BinaryOp<Device, Functor> {
public:
explicit MklBinaryOp(OpKernelConstruction* context)
: BinaryOp<Device, Functor>(context) {}
void Compute(OpKernelContext* context) override {
auto in0 = context->input(0);
auto in1 = context->input(1);
VLOG(1) << "Shapes (start mklbinaryop compute): "
<< in0.shape().DebugString() << " _and_ "
<< in1.shape().DebugString();
// Call the TensorFlow BinaryOp Compute method
BinaryOp<Device, Functor>::Compute(context);
auto out = context->mutable_output(0);
VLOG(1) << "Shapes (output): " << out->shape().DebugString();
// Pass input shape through to ouput shape
ForwardMklMetaDataInToOut(context, 0, 0);
out = context->mutable_output(0);
VLOG(1) << "Shapes (output): " << out->shape().DebugString();
}
};
//---------- Registration macros for various element-wise ops -----------
// We will need to redefine "REGISTER" to include the mkl_op_registry flag
#pragma push_macro("REGISTER")
#undef REGISTER
#define REGISTER(OP, D, N, F, T) \
REGISTER_KERNEL_BUILDER(Name(N) \
.Device(DEVICE_##D) \
.TypeConstraint<T>("T") \
.Label(mkl_op_registry::kMklOpLabel), \
OP<D##Device, F<T>>);
REGISTER5(MklBinaryOp, CPU, "_MklAdd", functor::add, float, Eigen::half, double,
int32, int64);
REGISTER7(MklBinaryOp, CPU, "_MklSub", functor::sub, float, Eigen::half, double,
int32, int64, complex64, complex128);
REGISTER5(MklBinaryOp, CPU, "_MklMul", functor::mul, float, Eigen::half, double,
uint8, int32);
REGISTER5(MklBinaryOp, CPU, "_MklMaximum", functor::maximum, float, Eigen::half,
double, int32, int64);
REGISTER5(MklBinaryOp, CPU, "_MklSquaredDifference",
functor::squared_difference, float, Eigen::half, double, int32,
int64);
#undef REGISTER
#pragma pop_macro("REGISTER")
//-----------------------------------------------------------------------
} // end namespace tensorflow
#endif // INTEL_MKL

4
tensorflow/core/kernels/mkl_identity_op.cc
View File

@ -41,9 +41,9 @@ class MklIdentityOp : public OpKernel {
bool input_in_mkl_format = mkl_shape_input.IsMklTensor();
if (input_in_mkl_format) {
ForwarMklTensorInToOut(context, 0, 0);
ForwardMklTensorInToOut(context, 0, 0);
} else {
FowardTfTensorInToOut(context, 0, 0);
ForwardTfTensorInToOut(context, 0, 0);
}
}

259
tensorflow/core/kernels/mkl_input_conversion_op.cc Normal file
View File

@ -0,0 +1,259 @@
/* 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.
==============================================================================*/
#ifdef INTEL_MKL
#include <algorithm>
#include <vector>
#include "tensorflow/core/framework/numeric_op.h"
#include "tensorflow/core/framework/op.h"
#include "tensorflow/core/framework/op_kernel.h"
#include "tensorflow/core/framework/register_types.h"
#include "tensorflow/core/framework/tensor.h"
#include "tensorflow/core/framework/tensor_shape.h"
#include "tensorflow/core/kernels/ops_util.h"
#include "tensorflow/core/platform/cpu_info.h"
#include "tensorflow/core/platform/macros.h"
#include "tensorflow/core/util/tensor_format.h"
#include "tensorflow/core/kernels/mkl_tfconv_op.h"
#include "tensorflow/core/util/mkl_util.h"
namespace tensorflow {
typedef Eigen::ThreadPoolDevice CPUDevice;
///////////////////////////////////////////////////////////
// Op kernel
// Checks and ensures that the 2 inputs are compatible for mkl binary ops.
// Here's the basic logic:
//
// if both inputs are in TF format:
// pass the inputs through to the output
// else if both inputs are in mkl format:
// if both have the same shape:
// pass the inputs through to the output
// else:
// convert both to TF
// else if one is TF and one is MKL:
// if broadcast is needed:
// convert the MKL format input to TF format
// else:
// convert the TF format input to MKL format
///////////////////////////////////////////////////////////
template <typename Device, typename T>
class MklInputConversionOp : public OpKernel {
public:
explicit MklInputConversionOp(OpKernelConstruction* context)
: OpKernel(context) {
OP_REQUIRES_OK(context, context->GetAttr("data_format", &data_format_str));
OP_REQUIRES_OK(context, context->GetAttr("T", &op_data_type));
has_avx512f_ = port::TestCPUFeature(port::CPUFeature::AVX512F);
}
private:
void Compute(OpKernelContext* context) override {
// Check if input tensors are in MKL format.
const Tensor& input_tensor_0 = MklGetInput(context, 0);
MklShape input_shape_0;
GetMklShape(context, 0, &input_shape_0);
const Tensor& input_tensor_1 = MklGetInput(context, 1);
MklShape input_shape_1;
GetMklShape(context, 1, &input_shape_1);
bool tf_shapes_are_same = MklCompareShapes(&context->input(0).shape(),
&context->input(1).shape());
VLOG(1) << "MklInputConversionOp: Input shapes are "
<< (tf_shapes_are_same ? "*same*" : "*different*") << ": "
<< context->input(0).shape().DebugString() << " and "
<< context->input(1).shape().DebugString();
// - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -
// if both inputs are in TF format, just copy input tensors to output.
if (!input_shape_0.IsMklTensor() && !input_shape_1.IsMklTensor()) {
VLOG(1) << "MklInputConversionOp: No conversion needed, "
<< "copying TF inputs to output";
ForwardTfTensorInToOut(context, 0, 0);
ForwardTfTensorInToOut(context, 1, 1);
return;
}
// - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -
// If both inputs are in MKL format
if (input_shape_0.IsMklTensor() && input_shape_1.IsMklTensor()) {
// If both have the same shape, pass them through
if (tf_shapes_are_same) {
VLOG(1) << "MklInputConversionOp: No conversion needed, "
<< "copying MKL inputs with identical shapes to output";
ForwardMklTensorInToOut(context, 0, 0);
ForwardMklTensorInToOut(context, 1, 1);
return;
}
// Sanity check
bool mkl_shapes_are_same =
MklCompareShapes(&input_shape_0, &input_shape_1);
if (mkl_shapes_are_same) {
CHECK(false) << "MklInputConversionOp: Unexpected: TF shapes are "
"different but MKL shapes are same";
}
// Both have different shapes, so broadcast will be necessary.
// Convert to TF and pass both tensors through (we can't do broadcast
// with MKL tensors)
VLOG(1) << "MklInputConversionOp: Broadcast needed, "
<< "converted MKL inputs to TF format";
MklToTfOp<Device, T>::ConvertMklToTf(this, context, data_format_str,
op_data_type, has_avx512f_, 0);
MklToTfOp<Device, T>::ConvertMklToTf(this, context, data_format_str,
op_data_type, has_avx512f_, 1);
SetDummyMklShapeOutput(context, 0);
SetDummyMklShapeOutput(context, 1);
return;
}
// - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -
// One input is MKL and one is TF. If no broadcast is needed, convert
// the TF tensor to MKL, otherwise convert the MKL tensor to TF format
VLOG(1) << "MklInputConversionOp: Inputs in different formats (MKL/TF)";
const Tensor* mkl_tensor;
const MklShape* mkl_shape;
const Tensor* tf_tensor;
MklShape* tf_mkl_shape;
uint mkl_tensor_index;
uint tf_tensor_index;
if (input_shape_0.IsMklTensor() && !input_shape_1.IsMklTensor()) {
mkl_tensor = &input_tensor_0;
mkl_shape = &input_shape_0;
mkl_tensor_index = 0;
tf_tensor = &input_tensor_1;
tf_mkl_shape = &input_shape_1;
tf_tensor_index = 1;
} else if (!input_shape_0.IsMklTensor() && input_shape_1.IsMklTensor()) {
mkl_tensor = &input_tensor_1;
mkl_shape = &input_shape_1;
mkl_tensor_index = 1;
tf_tensor = &input_tensor_0;
tf_mkl_shape = &input_shape_0;
tf_tensor_index = 0;
} else {
CHECK(false) << "MklInputConversionOp: Unexpected combination of input "
"shapes for MKL "
<< "element-wise op";
}
// Broadcast is needed if the shapes are not the same
bool broadcast_needed;
size_t in0_size = 1;
for (size_t i = 0; i < mkl_shape->GetDimension(); ++i)
in0_size *= mkl_shape->tf_dim_size(i);
size_t in1_size = 1;
for (size_t i = 0; i < tf_tensor->shape().dims(); ++i)
in1_size *= tf_tensor->shape().dim_size(i);
broadcast_needed = (in0_size != in1_size);
if (!broadcast_needed) {
// Both shapes are same, convert the TF input to MKL
VLOG(1) << "MklInputConversionOp: No broadcast needed.";
VLOG(1) << "MklInputConversionOp: Converting input " << tf_tensor_index
<< " to MKL format";
// Create MklShape
Tensor* tensor_out;
MklShape mkl_output_mkl_shape;
mkl_output_mkl_shape.SetMklTensor(true);
mkl_output_mkl_shape.SetTfLayout(mkl_shape->GetDimension(),
mkl_shape->GetSizes(),
mkl_shape->GetStrides());
mkl_output_mkl_shape.SetTfDimOrder(mkl_shape->GetDimension());
// ** Temporarily borrow the layout from the MKL input **
mkl_output_mkl_shape.SetMklLayout(mkl_shape->GetCurLayout());
// Create output tensor
AllocateOutputSetMklShape(context, tf_tensor_index, &tensor_out,
mkl_tensor->shape(), mkl_output_mkl_shape);
// Since the shapes are the same, use information from the other tensor
tf_mkl_shape->SetTfLayout(mkl_shape->GetDimension(),
mkl_shape->GetSizes(), mkl_shape->GetStrides());
// Convert the data format
tf_mkl_shape->GetConvertedFlatData(
mkl_shape->GetCurLayout(),
const_cast<T*>(tf_tensor->flat<T>().data()),
const_cast<T*>(tensor_out->flat<T>().data()));
// ** Release the borrowed layout to avoid double deletion
// in the destructor call **
mkl_output_mkl_shape.SetMklLayout(nullptr);
// -- The tensor in MKL format passes through --
ForwardMklTensorInToOut(context, mkl_tensor_index, mkl_tensor_index);
} else {
// Broadcast is needed, so convert the MKL input to TF
VLOG(1) << "MklInputConversionOp: Broadcast needed.";
VLOG(1) << "MklInputConversionOp: Converting input " << mkl_tensor_index
<< " to TF format";
MklToTfOp<Device, T>::ConvertMklToTf(this, context, data_format_str,
op_data_type, has_avx512f_,
mkl_tensor_index);
SetDummyMklShapeOutput(context, mkl_tensor_index);
// The tensor in TF format passes through
ForwardTfTensorInToOut(context, tf_tensor_index, tf_tensor_index);
}
VLOG(1) << "MklInputConversionOp: Shapes (output): "
<< context->mutable_output(0)->shape().DebugString() << " and "
<< context->mutable_output(1)->shape().DebugString();
VLOG(1) << "MklInputConversion completed successfully.";
}
private:
/// Data format of the operation
string data_format_str;
/// Data type of the operation
DataType op_data_type;
/// CPUIDInfo
bool has_avx512f_ = false;
};
///////////////////////////////////////////////////////////
// Register kernel
///////////////////////////////////////////////////////////
#define REGISTER_CPU(T) \
REGISTER_KERNEL_BUILDER(Name("_MklInputConversion") \
.Device(DEVICE_CPU) \
.TypeConstraint<T>("T") \
.Label(mkl_op_registry::kMklOpLabel), \
MklInputConversionOp<CPUDevice, T>);
TF_CALL_NUMBER_TYPES(REGISTER_CPU);
#undef REGISTER_CPU
} // namespace tensorflow
#endif // INTEL_MKL

136
tensorflow/core/kernels/mkl_tfconv_op.h Normal file
View File

@ -0,0 +1,136 @@
/* 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.
==============================================================================*/
#ifdef INTEL_MKL
#ifndef TENSORFLOW_CORE_KERNELS_MKL_TFCONV_OP_H_
#define TENSORFLOW_CORE_KERNELS_MKL_TFCONV_OP_H_
#include <algorithm>
#include <vector>
#include "tensorflow/core/framework/numeric_op.h"
#include "tensorflow/core/framework/op.h"
#include "tensorflow/core/framework/op_kernel.h"
#include "tensorflow/core/framework/register_types.h"
#include "tensorflow/core/framework/tensor.h"
#include "tensorflow/core/framework/tensor_shape.h"
#include "tensorflow/core/kernels/ops_util.h"
#include "tensorflow/core/platform/cpu_info.h"
#include "tensorflow/core/platform/macros.h"
#include "tensorflow/core/util/tensor_format.h"
#include "mkl_dnn.h"
#include "mkl_dnn_types.h"
#include "tensorflow/core/util/mkl_util.h"
namespace tensorflow {
typedef Eigen::ThreadPoolDevice CPUDevice;
///////////////////////////////////////////////////////////
// Op kernel
///////////////////////////////////////////////////////////
template <typename Device, typename T>
class MklToTfOp : public OpKernel {
public:
explicit MklToTfOp(OpKernelConstruction* context) : OpKernel(context) {
OP_REQUIRES_OK(context, context->GetAttr("data_format", &data_format_str));
OP_REQUIRES_OK(context, context->GetAttr("T", &op_data_type));
has_avx512f_ = port::TestCPUFeature(port::CPUFeature::AVX512F);
}
void Compute(OpKernelContext* context) override {
ConvertMklToTf(this, context, data_format_str, op_data_type, has_avx512f_,
0);
VLOG(1) << "MKLToTFConversion complete successfully.";
}
static void ConvertMklToTf(OpKernel* op_kernel, OpKernelContext* context,
string data_format_str, DataType op_data_type,
bool has_avx512f, uint input_number) {
// Check that input tensor is in MKL format.
const Tensor& input_tensor = MklGetInput(context, input_number);
MklShape input_shape;
GetMklShape(context, input_number, &input_shape);
// if input is already in Tf format, then just copy input tensor to output.
if (!input_shape.IsMklTensor()) {
context->set_output(input_number, input_tensor);
VLOG(1) << "MKLToTFConversion: No conversion needed, "
<< "copying input to output";
return;
}
// Check that input data type is same as operator data type and that it is
// same as output data type.
DataType input_data_type = op_kernel->input_type(input_number);
DataType output_data_type = op_kernel->output_type(input_number);
CHECK_EQ(op_data_type, input_data_type);
CHECK_EQ(op_data_type, output_data_type);
TensorShape output_shape;
size_t ndims = input_shape.GetDimension();
size_t* in_sizes = new size_t[ndims];
for (size_t i = 0; i < ndims; i++) {
// Outermost to innermost dimension
output_shape.AddDim(input_shape.GetSizes()[input_shape.tf_dim_idx(i)]);
in_sizes[i] = input_shape.GetSizes()[i];
}
// Allocate output tensor.
Tensor* output_tensor = NULL;
OP_REQUIRES_OK(context,
context->allocate_output(input_number, output_shape, &output_tensor));
dnnLayout_t output_layout =
static_cast<dnnLayout_t>(input_shape.GetTfLayout());
// Execute DNNConversion.
void* input_buffer =
static_cast<void*>(const_cast<T*>(input_tensor.flat<T>().data()));
delete[] in_sizes;
void* output_buffer =
static_cast<void*>(const_cast<T*>(output_tensor->flat<T>().data()));
input_shape.GetConvertedFlatData(output_layout, input_buffer,
output_buffer);
VLOG(1) << "MKLToTFConversion complete successfully.";
}
private:
/// Data format of the operation
string data_format_str;
/// Data type of the operation
DataType op_data_type;
/// CPUIDInfo
bool has_avx512f_ = false;
};
///////////////////////////////////////////////////////////
// Register kernel
///////////////////////////////////////////////////////////
#define REGISTER_CPU(T) \
REGISTER_KERNEL_BUILDER(Name("_MklToTf") \
.Device(DEVICE_CPU) \
.TypeConstraint<T>("T") \
.Label(mkl_op_registry::kMklOpLabel), \
MklToTfOp<CPUDevice, T>);
TF_CALL_NUMBER_TYPES(REGISTER_CPU);
#undef REGISTER_CPU
} // namespace tensorflow
#endif // TENSORFLOW_CORE_KERNELS_MKL_TFCONV_OP_H_
#endif // INTEL_MKL

413
tensorflow/core/kernels/svd_op_gpu.cu.cc Normal file
View File

@ -0,0 +1,413 @@
/* 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.
==============================================================================*/
// See docs in ../ops/linalg_ops.cc.
// TODO(shamanDevel): Enable complex inputs. This will require a specialization
// of Gesvd for complex inputs as well as a new kernel
// definition to output the singular values as reals
// instead of complex values. The current CPU implementation
// outputs the singular values as complex values and then
// casts them to reals in the python wrapper.
#if GOOGLE_CUDA
#define EIGEN_USE_GPU
#include <algorithm>
#include <vector>
#include "tensorflow/core/framework/kernel_def_builder.h"
#include "tensorflow/core/framework/op_kernel.h"
#include "tensorflow/core/framework/register_types.h"
#include "tensorflow/core/framework/tensor_shape.h"
#include "tensorflow/core/framework/types.h"
#include "tensorflow/core/kernels/cuda_solvers.h"
#include "tensorflow/core/kernels/linalg_ops_common.h"
#include "tensorflow/core/kernels/transpose_functor.h"
#include "tensorflow/core/lib/core/errors.h"
#include "tensorflow/core/platform/logging.h"
#include "tensorflow/core/platform/stream_executor.h"
#include "tensorflow/core/platform/types.h"
#include "tensorflow/core/util/cuda_kernel_helper.h"
#include "third_party/eigen3/unsupported/Eigen/CXX11/Tensor"
namespace tensorflow {
static const char kErrMsg[] =
"Singular Value Decomposition was not successful. The input might not be "
"valid.";
typedef Eigen::GpuDevice GPUDevice;
namespace {
// This kernel computes the reduction
// V' = sum_i (M_i * U_i,1 * S_i).
// The result is stored in V[batch] and has the same sign as the
// real value of V (which should be computed)
template <class Scalar>
__global__ void ComputeValueOfVKernel(Cuda2DLaunchConfig config, int64 m,
int64 ldu, const Scalar* M,
const Scalar* U, const Scalar* S,
Scalar* V) {
CUDA_AXIS_KERNEL_LOOP(batch, config.virtual_thread_count, x) {
CUDA_AXIS_KERNEL_LOOP(i, config.virtual_thread_count, y) {
Scalar v = M[i + m * batch] * U[ldu * (i + m * batch)] * S[batch];
CudaAtomicAdd(V + batch, v);
}
}
}
// Extracts the sign of V
// V[i] = V[i]>=0 ? 1 : 0
template <class Scalar>
__global__ void ExtractSignOfVKernel(CudaLaunchConfig config, Scalar* V) {
CUDA_1D_KERNEL_LOOP(i, config.virtual_thread_count) {
V[i] = V[i] >= 0 ? Scalar(1) : Scalar(-1);
}
}
}
// Scalar: The input scalar type (can be complex)
template <class Scalar>
class SvdOpGpu : public AsyncOpKernel {
public:
using RealScalar = typename Eigen::NumTraits<Scalar>::Real;
explicit SvdOpGpu(OpKernelConstruction* context) : AsyncOpKernel(context) {
OP_REQUIRES_OK(context, context->GetAttr("compute_uv", &compute_uv_));
OP_REQUIRES_OK(context, context->GetAttr("full_matrices", &full_matrices_));
}
void RunSVD(OpKernelContext* context, DoneCallback done, int64 m, int64 n,
int64 p, int64 batch_size, Scalar* input_ptr,
RealScalar* outputS_ptr, Scalar* outputU_ptr,
Scalar* outputVT_ptr, int* dev_info_ptr, CudaSolver& solver) {
// Save the input matrix
// Needed for the n=1 fix, see below, since SVD destroys the input
Tensor input_copy;
if (compute_uv_ && n == 1) {
OP_REQUIRES_OK_ASYNC(
context,
context->allocate_temp(DataTypeToEnum<Scalar>::v(),
TensorShape({batch_size, m}), &input_copy),
done);
const GPUDevice& d = context->eigen_device<GPUDevice>();
d.memcpy(input_copy.flat<Scalar>().data(), input_ptr,
batch_size * m * sizeof(Scalar));
}
for (int64 batch = 0; batch < batch_size; ++batch) {
Scalar* input = input_ptr + batch * m * n;
RealScalar* outputS = outputS_ptr + batch * p;
Scalar* outputU = NULL;
Scalar* outputVT = NULL;
char jobu = 'N';
char jobvt = 'N';
if (compute_uv_) {
if (full_matrices_) {
outputU = outputU_ptr + batch * m * m;
outputVT = outputVT_ptr + batch * n * n;
jobu = 'A';
jobvt = 'A';
} else {
outputU = outputU_ptr + batch * m * p;
outputVT = outputVT_ptr + batch * n * p;
jobu = 'S';
jobvt = 'S';
}
}
OP_REQUIRES_OK_ASYNC(
context, solver.Gesvd(jobu, jobvt, m, n, input, m, outputS, outputU,
m, outputVT, n, dev_info_ptr + batch),
done);
}
// This is a bug in cuSolver:
// If n is one, then outputVT only contains zeros instead of ones.
// Hence, I need to fill outputVT manually
// The question is: +1 or -1?
// -> Compute U*S and compare sign against M
// But because S is zero except for the first entry, the multiplication
// simplifies a lot.
// However, what happens if M contains zeros? At these indices, it is
// impossible to determine the value of V.
// -> Compute V for all rows in M to cope for zeros.
// 1. V' = sum_i (M_i * U_i,1 * S_i)
// 2. V = {1, V'>=0, -1, V'<0}
// TODO: what is with complex values?
if (compute_uv_ && n == 1) {
// 1. compute the (batched) sum
const GPUDevice& d = context->eigen_device<GPUDevice>();
d.memset(outputVT_ptr, 0, batch_size * sizeof(Scalar));
Cuda2DLaunchConfig cfg2D = GetCuda2DLaunchConfig(batch_size, m, d);
ComputeValueOfVKernel<<<cfg2D.block_count, cfg2D.thread_per_block, 0,
d.stream()>>>(
cfg2D, m, full_matrices_ ? m : p, input_copy.flat<Scalar>().data(),
outputU_ptr, outputS_ptr, outputVT_ptr);
// 2. clamp V to -1 or +1
CudaLaunchConfig cfg1D = GetCudaLaunchConfig(batch_size, d);
ExtractSignOfVKernel<<<cfg1D.block_count, cfg1D.thread_per_block, 0,
d.stream()>>>(cfg1D, outputVT_ptr);
}
}
void CheckResult(OpKernelContext* context, DoneCallback done,
const std::vector<DeviceLapackInfo>& dev_info,
CudaSolver& solver, Tensor& catch1, Tensor& catch2) {
auto info_checker = [context, dev_info, done, catch1, catch2](
const Status& status, const std::vector<HostLapackInfo>& /* unused */) {
Status full_status = status;
if (!full_status.ok()) {
full_status.Update(errors::InvalidArgument(kErrMsg));
}
OP_REQUIRES_OK_ASYNC(context, full_status, done);
done();
};
OP_REQUIRES_OK_ASYNC(context, solver.CopyLapackInfoToHostAsync(
dev_info, std::move(info_checker)),
done);
}
// The SVD if m >= n
// TODO: can the two cases (MgeqN and MlessN) be simplified,
// common boilerplate be reduced, or even combined in one method?
void PerformSVD_MgeqN(OpKernelContext* context, DoneCallback done, int64 m,
int64 n, int64 p, const gtl::ArraySlice<int32>& perm,
const Tensor& M, Tensor* S, Tensor* U, Tensor* V) {
TensorShape shapeRaw = M.shape();
shapeRaw.RemoveDim(shapeRaw.dims() - 1);
shapeRaw.RemoveDim(shapeRaw.dims() - 1);
// Transpose M, because cuSolver expects it to be column-major
TensorShape input_shape = shapeRaw;
input_shape.AddDim(n);
input_shape.AddDim(m);
Tensor input_copy;
OP_REQUIRES_OK_ASYNC(
context, context->allocate_temp(M.dtype(), input_shape, &input_copy),
done);
auto device = context->eigen_device<GPUDevice>();
OP_REQUIRES_OK_ASYNC(context, DoTranspose(device, M, perm, &input_copy),
done);
// I need to transpose U at the end
// Not V, because cuSolver work column-major
Tensor u_copy;
if (compute_uv_) {
TensorShape u_shape;
if (full_matrices_) {
u_shape = U->shape();
} else {
u_shape = shapeRaw;
u_shape.AddDim(p);
u_shape.AddDim(m);
}
OP_REQUIRES_OK_ASYNC(
context, context->allocate_temp(U->dtype(), u_shape, &u_copy), done);
}
// get the pointers to the data
Scalar* input_ptr;
RealScalar* outputS_ptr;
Scalar* outputU_ptr = NULL;
Scalar* outputV_ptr = NULL;
auto input_reshaped = input_copy.template flat_inner_dims<Scalar, 3>();
input_ptr = input_reshaped.data();
outputS_ptr = S->template flat_inner_dims<RealScalar, 2>().data();
if (compute_uv_) {
outputU_ptr = u_copy.template flat_inner_dims<Scalar, 3>().data();
outputV_ptr = V->template flat_inner_dims<Scalar, 3>().data();
}
// call the SVD
const int64 batch_size = input_reshaped.dimension(0);
std::vector<DeviceLapackInfo> dev_info;
dev_info.emplace_back(context, batch_size, "gesvd");
CudaSolver solver(context);
RunSVD(context, done, m, n, p, batch_size, input_ptr, outputS_ptr,
outputU_ptr, outputV_ptr, dev_info.back().mutable_data(), solver);
// Transpose U
if (compute_uv_) {
OP_REQUIRES_OK_ASYNC(context, DoTranspose(device, u_copy, perm, U), done);
}
// now check if the SVD operation succeeded or not
CheckResult(context, done, dev_info, solver, input_copy, u_copy);
}
// The SVD if m < n
void PerformSVD_MlessN(OpKernelContext* context, DoneCallback done, int64 m,
int64 n, int64 p, const gtl::ArraySlice<int32>& perm,
const Tensor& M, Tensor* S, Tensor* U, Tensor* V) {
// Perform the SVD on M'
// Reuse the input buffer or make a copy for the SVD depending on whether
// this op owns the
// input buffer exclusively. This is needed because the SVD modifies the
// input
Tensor input_copy;
OP_REQUIRES_OK_ASYNC(context, context->forward_input_or_allocate_temp(
{0}, DataTypeToEnum<Scalar>::value,
M.shape(), &input_copy),
done);
if (!M.SharesBufferWith(input_copy)) {
const GPUDevice& d = context->eigen_device<GPUDevice>();
d.memcpy(input_copy.flat<Scalar>().data(), M.flat<Scalar>().data(),
M.NumElements() * sizeof(Scalar));
}
// I need to transpose V at the end
Tensor v_copy;
if (compute_uv_) {
TensorShape v_shape;
if (full_matrices_) {
v_shape = V->shape();
} else {
TensorShape shapeRaw = M.shape();
shapeRaw.RemoveDim(shapeRaw.dims() - 1);
shapeRaw.RemoveDim(shapeRaw.dims() - 1);
v_shape = shapeRaw;
v_shape.AddDim(p);
v_shape.AddDim(n);
}
OP_REQUIRES_OK_ASYNC(
context, context->allocate_temp(V->dtype(), v_shape, &v_copy), done);
}
// get the pointers to the data
Scalar* input_ptr;
RealScalar* outputS_ptr;
Scalar* outputU_ptr = NULL;
Scalar* outputV_ptr = NULL;
auto input_reshaped = input_copy.template flat_inner_dims<Scalar, 3>();
input_ptr = input_reshaped.data();
outputS_ptr = S->template flat_inner_dims<RealScalar, 2>().data();
if (compute_uv_) {
// Note that U and V are flipped
outputU_ptr = v_copy.template flat_inner_dims<Scalar, 3>().data();
outputV_ptr = U->template flat_inner_dims<Scalar, 3>().data();
}
// call the SVD
const int64 batch_size = input_reshaped.dimension(0);
std::vector<DeviceLapackInfo> dev_info;
dev_info.emplace_back(context, batch_size, "gesvd");
CudaSolver solver(context);
// Note that m and n are flipped
RunSVD(context, done, n, m, p, batch_size, input_ptr, outputS_ptr,
outputU_ptr, outputV_ptr, dev_info.back().mutable_data(), solver);
// Transpose V
if (compute_uv_) {
auto device = context->eigen_device<GPUDevice>();
OP_REQUIRES_OK_ASYNC(context, DoTranspose(device, v_copy, perm, V), done);
}
// now check if the SVD operation succeeded or not
CheckResult(context, done, dev_info, solver, input_copy, v_copy);
}
void ComputeAsync(OpKernelContext* context, DoneCallback done) final {
const Tensor& input = context->input(0);
const int ndims = input.dims();
const int64 m = input.dim_size(ndims - 2);
const int64 n = input.dim_size(ndims - 1);
const int64 p = std::min(m, n);
// Validate inputs.
OP_REQUIRES_ASYNC(
context, ndims >= 2,
errors::InvalidArgument("Input must have rank >= 2, got ", ndims),
done);
// output tensors.
Tensor* outputU = NULL;
Tensor* outputS = NULL;
Tensor* outputV = NULL;
// compute shapes
TensorShape shapeRaw = input.shape();
shapeRaw.RemoveDim(shapeRaw.dims() - 1);
shapeRaw.RemoveDim(shapeRaw.dims() - 1);
TensorShape shapeS = shapeRaw;
TensorShape shapeU = shapeRaw;
TensorShape shapeV = shapeRaw;
shapeS.AddDim(p);
if (compute_uv_) {
if (full_matrices_) {
shapeU.AddDim(m);
shapeU.AddDim(m);
shapeV.AddDim(n);
shapeV.AddDim(n);
} else {
shapeU.AddDim(m);
shapeU.AddDim(p);
shapeV.AddDim(n);
shapeV.AddDim(p);
}
} else {
shapeU = TensorShape({0});
shapeV = TensorShape({0});
}
// allocate output
OP_REQUIRES_OK_ASYNC(context, context->allocate_output(0, shapeS, &outputS),
done);
OP_REQUIRES_OK_ASYNC(context, context->allocate_output(1, shapeU, &outputU),
done);
OP_REQUIRES_OK_ASYNC(context, context->allocate_output(2, shapeV, &outputV),
done);
if (n == 0 || m == 0) {
// If X is an empty matrix (0 rows, 0 col), X * X' == X.
// Therefore, we return X.
done();
return;
}
// Prepare permutation
std::vector<int32> perm;
for (size_t i = 0; i < ndims - 2; ++i) perm.push_back(i);
perm.push_back(ndims - 1); // transpose last two dimensions
perm.push_back(ndims - 2);
gtl::ArraySlice<int32> permAS(perm);
// call implementations
if (m >= n) {
PerformSVD_MgeqN(context, done, m, n, p, permAS, input, outputS, outputU,
outputV);
} else {
PerformSVD_MlessN(context, done, m, n, p, permAS, input, outputS, outputU,
outputV);
}
}
private:
bool compute_uv_;
bool full_matrices_;
};
// TODO: add support for complex types
REGISTER_LINALG_OP_GPU("Svd", (SvdOpGpu<float>), float);
REGISTER_LINALG_OP_GPU("Svd", (SvdOpGpu<double>), double);
REGISTER_LINALG_OP_GPU("BatchSvd", (SvdOpGpu<float>), float);
REGISTER_LINALG_OP_GPU("BatchSvd", (SvdOpGpu<double>), double);
} // namespace tensorflow
#endif // GOOGLE_CUDA

2
tensorflow/core/kernels/tensor_array_ops.cc
View File

@ -1069,7 +1069,7 @@ class TensorArrayUnpackOrScatterOp : public OpKernel {
} else {
OP_REQUIRES(
ctx, max_index < array_size,
errors::InvalidArgument("Max scatter index must be <= array size (",
errors::InvalidArgument("Max scatter index must be < array size (",
max_index, " vs. ", array_size, ")"));
}
element_shape.RemoveDim(0);

103
tensorflow/core/ops/math_ops.cc
View File

@ -498,6 +498,24 @@ Returns x + y element-wise.
[here](http://docs.scipy.org/doc/numpy/user/basics.broadcasting.html)
)doc");
REGISTER_OP("_MklAdd")
.Input("x: T")
.Input("y: T")
.Input("mkl_x: uint8")
.Input("mkl_y: uint8")
.Output("z: T")
.Output("mkl_z: uint8")
.Attr(
"T: {half, float, double, uint8, int8, int16, int32, int64, complex64, "
"complex128, string}")
.SetShapeFn(shape_inference::BroadcastBinaryOpShapeFn)
.Doc(R"doc(
Returns x + y element-wise.
*NOTE*: `Add` supports broadcasting. `AddN` does not. More about broadcasting
[here](http://docs.scipy.org/doc/numpy/user/basics.broadcasting.html)
)doc");
REGISTER_OP("Sub")
.BINARY_MORE()
.SetShapeFn(shape_inference::BroadcastBinaryOpShapeFn)
@ -508,6 +526,19 @@ Returns x - y element-wise.
[here](http://docs.scipy.org/doc/numpy/user/basics.broadcasting.html)
)doc");
REGISTER_OP("_MklSub")
.BINARY_FEWER()
.Input("mkl_x: uint8")
.Input("mkl_y: uint8")
.Output("mkl_z: uint8")
.SetShapeFn(shape_inference::BroadcastBinaryOpShapeFn)
.Doc(R"doc(
Returns x - y element-wise.
*NOTE*: `Sub` supports broadcasting. More about broadcasting
[here](http://docs.scipy.org/doc/numpy/user/basics.broadcasting.html)
)doc");
REGISTER_OP("Mul")
.BINARY_MORE()
.SetIsCommutative()
@ -519,6 +550,20 @@ Returns x * y element-wise.
[here](http://docs.scipy.org/doc/numpy/user/basics.broadcasting.html)
)doc");
REGISTER_OP("_MklMul")
.BINARY_MORE()
.Input("mkl_x: uint8")
.Input("mkl_y: uint8")
.Output("mkl_z: uint8")
.SetIsCommutative()
.SetShapeFn(shape_inference::BroadcastBinaryOpShapeFn)
.Doc(R"doc(
Returns x * y element-wise.
*NOTE*: `Mul` supports broadcasting. More about broadcasting
[here](http://docs.scipy.org/doc/numpy/user/basics.broadcasting.html)
)doc");
REGISTER_OP("Div")
.BINARY_MORE()
.SetShapeFn(shape_inference::BroadcastBinaryOpShapeFn)
@ -577,6 +622,20 @@ Returns (x - y)(x - y) element-wise.
[here](http://docs.scipy.org/doc/numpy/user/basics.broadcasting.html)
)doc");
REGISTER_OP("_MklSquaredDifference")
.BINARY_FEWER()
.Input("mkl_x: uint8")
.Input("mkl_y: uint8")
.Output("mkl_z: uint8")
.SetIsCommutative()
.SetShapeFn(shape_inference::BroadcastBinaryOpShapeFn)
.Doc(R"doc(
Returns (x - y)(x - y) element-wise.
*NOTE*: `SquaredDifference` supports broadcasting. More about broadcasting
[here](http://docs.scipy.org/doc/numpy/user/basics.broadcasting.html)
)doc");
#undef BINARY_FEWER
#undef BINARY_MORE
@ -594,6 +653,23 @@ Returns the max of x and y (i.e. x > y ? x : y) element-wise.
[here](http://docs.scipy.org/doc/numpy/user/basics.broadcasting.html)
)doc");
REGISTER_OP("_MklMaximum")
.Input("x: T")
.Input("y: T")
.Input("mkl_x: uint8")
.Input("mkl_y: uint8")
.Output("z: T")
.Output("mkl_z: uint8")
.Attr("T: {half, float, double, int32, int64}")
.SetIsCommutative()
.SetShapeFn(shape_inference::BroadcastBinaryOpShapeFn)
.Doc(R"doc(
Returns the max of x and y (i.e. x > y ? x : y) element-wise.
*NOTE*: `Maximum` supports broadcasting. More about broadcasting
[here](http://docs.scipy.org/doc/numpy/user/basics.broadcasting.html)
)doc");
REGISTER_OP("Minimum")
.Input("x: T")
.Input("y: T")
@ -2604,4 +2680,31 @@ Equivalent to np.digitize.
@end_compatibility
)doc");
#ifdef INTEL_MKL
REGISTER_OP("_MklAddN")
.Input("inputs: N * T")
.Input("mkl_input: N * uint8")
.Output("sum: T")
.Output("mkl_sum: uint8")
.Attr("N: int >= 1")
.Attr("T: numbertype")
.SetIsCommutative()
.SetIsAggregate()
.SetShapeFn([](InferenceContext* c) {
ShapeHandle cur = c->input(c->num_inputs() - 1);
for (int i = c->num_inputs() - 2; i >= 0; --i) {
TF_RETURN_WITH_CONTEXT_IF_ERROR(c->Merge(c->input(i), cur, &cur),
"From merging shape ", i,
" with other shapes.");
}
c->set_output(0, cur);
return Status::OK();
})
.Doc(R"doc(
Add two input tensors element wise using mkl kernel sum.
inputs: Must all be the same size and shape.
)doc");
#endif // INTEL_MKL
} // namespace tensorflow

23
tensorflow/core/ops/nn_ops.cc
View File

@ -3238,6 +3238,29 @@ REGISTER_OP("_MklToTf")
.Doc(R"doc(
MKL operator to convert a tensor from MKL layout to TensorFlow layout.
NOTE Do not invoke this operator directly in Python. Graph rewrite pass is
expected to invoke these operators.
)doc");
REGISTER_OP("_MklInputConversion")
.Input("input_0: T")
.Input("input_1: T")
.Input("mkl_input_0: uint8")
.Input("mkl_input_1: uint8")
.Output("output_0: T")
.Output("output_1: T")
.Output("mkl_output_0: uint8")
.Output("mkl_output_1: uint8")
// All datatypes supported by element-wise ops
.Attr(
"T: {half, float, double, uint8, int8, uint16, int16, int32, int64, "
"complex64, complex128}")
.Attr(GetConvnetDataFormatAttrString())
.Doc(R"doc(
MKL operator to process the inputs to an elementwise MKL op. Both inputs
need to be either in TF or in MKL format. This op is added before every
element-wise MKL op.
NOTE Do not invoke this operator directly in Python. Graph rewrite pass is
expected to invoke these operators.
)doc");

19
tensorflow/core/ops/ops.pbtxt
View File

@ -15865,6 +15865,25 @@ op {
}
summary: "Transforms a serialized tensorflow.TensorProto proto into a Tensor."
}
op {
name: "SerializeTensor"
input_arg {
name: "tensor"
description: "A Tensor of type `T`."
type: "T"
}
output_arg {
name: "serialized"
description: "A serialized TensorProto proto of the input tensor."
type_attr: DT_STRING
}
attr {
name: "T"
type: "type"
description: "The type of the input tensor."
}
summary: "Transforms a Tensor into a serialized TensorProto proto."
}
op {
name: "Placeholder"
output_arg {

2
tensorflow/core/ops/parsing_ops.cc
View File

@ -26,7 +26,7 @@ using shape_inference::ShapeHandle;
REGISTER_OP("DecodeRaw")
.Input("bytes: string")
.Output("output: out_type")
.Attr("out_type: {half,float,double,int32,uint8,int16,int8,int64}")
.Attr("out_type: {half,float,double,int32,uint16,uint8,int16,int8,int64}")
.Attr("little_endian: bool = true")
.SetShapeFn([](InferenceContext* c) {
// Note: last dimension is data dependent.

2
tensorflow/core/ops/string_ops.cc
View File

@ -381,7 +381,7 @@ input = b'thirteen'
position = [1, 5, 7]
length = [3, 2, 1]
output = [b'hir', b'ee', b'n"]
output = [b'hir', b'ee', b'n']
```
input: Tensor of strings

2
tensorflow/core/platform/cuda_libdevice_path_test.cc
View File

@ -27,7 +27,7 @@ TEST(CudaLibdevicePathTest, LibdevicePath) {
VLOG(2) << "Libdevice root = " << LibdeviceRoot();
std::vector<string> libdevice_files;
TF_EXPECT_OK(Env::Default()->GetMatchingPaths(
io::JoinPath(LibdeviceRoot(), "libdevice.compute_*.bc"),
io::JoinPath(LibdeviceRoot(), "libdevice.*.bc"),
&libdevice_files));
EXPECT_LT(0, libdevice_files.size());
}

4
tensorflow/core/public/version.h
View File

@ -19,12 +19,12 @@ limitations under the License.
// TensorFlow uses semantic versioning, see http://semver.org/.
#define TF_MAJOR_VERSION 1
#define TF_MINOR_VERSION 3
#define TF_MINOR_VERSION 4
#define TF_PATCH_VERSION 0
// TF_VERSION_SUFFIX is non-empty for pre-releases (e.g. "-alpha", "-alpha.1",
// "-beta", "-rc", "-rc.1")
#define TF_VERSION_SUFFIX ""
#define TF_VERSION_SUFFIX "-dev"
#define TF_STR_HELPER(x) #x
#define TF_STR(x) TF_STR_HELPER(x)

76
tensorflow/core/util/cuda_kernel_helper.h
View File

@ -25,6 +25,29 @@ limitations under the License.
#include "tensorflow/core/platform/logging.h"
#include "tensorflow/core/platform/stream_executor.h"
#include "tensorflow/core/platform/types.h"
#include "cuda/include/cuda.h"
// Mask for all 32 threads in a warp.
#define CUDA_WARP_ALL 0xFFFFFFFF
#if defined(CUDA_VERSION) && CUDA_VERSION < 9000
// CUDA 9.0 introduces a new, light-weight barrier synchronization primitive
// that operates at the warp-scope. This is required to ensure visibility of
// reads/writes among threads that can make indepenent progress on Volta.
// For previous CUDA versions these synchronizations not necessary, and we
// define an empty function as a convenience for backward compatibility.
__device__ inline void __syncwarp(unsigned mask=CUDA_WARP_ALL) {}
// CUDA 9.0 deprecates the warp-intrinsic functions (shfl, ballot, etc.) in
// favor of synchronizing versions. These ensure that all warp lanes specified
// in mask execute the intrinsic in convergence. Here we provide legacy mappings
// to the less-verbose routines provided in previous versions of CUDA.
#define __ballot_sync(mask, predicate) __ballot(predicate)
#define __shfl_sync(mask, val, srcLane, width) __shfl(val, srcLane, width)
#define __shfl_down_sync(mask, val, delta, width) __shfl_down(val, delta, width)
#define __shfl_up_sync(mask, val, delta, width) __shfl_up(val, delta, width)
#define __shfl_xor_sync(mask, val, laneMask, width) __shfl_xor(val, laneMask, width)
#endif
// Usage of GetCudaLaunchConfig, GetCuda2DLaunchConfig, and
// GetCuda3DLaunchConfig:
@ -613,82 +636,95 @@ EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T tf_max(const T& x, const T& y) {
return x < y ? y : x;
}
__device__ EIGEN_ALWAYS_INLINE unsigned CudaBallot(unsigned mask,
int predicate) {
return __ballot_sync(mask, predicate);
}
template <typename T>
__device__ EIGEN_ALWAYS_INLINE T CudaShuffle(T value, int srcLane,
__device__ EIGEN_ALWAYS_INLINE T CudaShuffle(unsigned mask, T value,
int srcLane,
int width = warpSize) {
return __shfl(value, srcLane, width);
return __shfl_sync(mask, value, srcLane, width);
}
// Variant of the (undocumented) version from the CUDA SDK, but using unsigned
// instead of float for lo and hi (which is incorrect with ftz, for example).
// A bug has been filed with NVIDIA and will be fixed in the next CUDA release.
// TODO(csigg): remove when the bug is fixed in the next CUDA release.
__device__ EIGEN_ALWAYS_INLINE double CudaShuffle(double value, int srcLane,
__device__ EIGEN_ALWAYS_INLINE double CudaShuffle(unsigned mask,
double value, int srcLane,
int width = warpSize) {
unsigned lo, hi;
asm volatile("mov.b64 {%0,%1}, %2;" : "=r"(lo), "=r"(hi) : "d"(value));
hi = __shfl(hi, srcLane, width);
lo = __shfl(lo, srcLane, width);
hi = __shfl_sync(mask, hi, srcLane, width);
lo = __shfl_sync(mask, lo, srcLane, width);
asm volatile("mov.b64 %0, {%1,%2};" : "=d"(value) : "r"(lo), "r"(hi));
return value;
}
template <typename T>
__device__ EIGEN_ALWAYS_INLINE T CudaShuffleUp(T value, int delta,
__device__ EIGEN_ALWAYS_INLINE T CudaShuffleUp(unsigned mask,
T value, int delta,
int width = warpSize) {
return __shfl_up(value, delta, width);
return __shfl_up_sync(mask, value, delta, width);
}
// Variant of the (undocumented) version from the CUDA SDK, but using unsigned
// instead of float for lo and hi (which is incorrect with ftz, for example).
// A bug has been filed with NVIDIA and will be fixed in the next CUDA release.
// TODO(csigg): remove when the bug is fixed in the next CUDA release.
__device__ EIGEN_ALWAYS_INLINE double CudaShuffleUp(double value, int delta,
__device__ EIGEN_ALWAYS_INLINE double CudaShuffleUp(unsigned mask,
double value, int delta,
int width = warpSize) {
unsigned lo, hi;
asm volatile("mov.b64 {%0,%1}, %2;" : "=r"(lo), "=r"(hi) : "d"(value));
hi = __shfl_up(hi, delta, width);
lo = __shfl_up(lo, delta, width);
hi = __shfl_up_sync(mask, hi, delta, width);
lo = __shfl_up_sync(mask, lo, delta, width);
asm volatile("mov.b64 %0, {%1,%2};" : "=d"(value) : "r"(lo), "r"(hi));
return value;
}
template <typename T>
__device__ EIGEN_ALWAYS_INLINE T CudaShuffleDown(T value, int delta,
__device__ EIGEN_ALWAYS_INLINE T CudaShuffleDown(unsigned mask,
T value, int delta,
int width = warpSize) {
return __shfl_down(value, delta, width);
return __shfl_down_sync(mask, value, delta, width);
}
// Variant of the (undocumented) version from the CUDA SDK, but using unsigned
// instead of float for lo and hi (which is incorrect with ftz, for example).
// A bug has been filed with NVIDIA and will be fixed in the next CUDA release.
// TODO(csigg): remove when the bug is fixed in the next CUDA release.
__device__ EIGEN_ALWAYS_INLINE double CudaShuffleDown(double value, int delta,
__device__ EIGEN_ALWAYS_INLINE double CudaShuffleDown(unsigned mask,
double value, int delta,
int width = warpSize) {
unsigned lo, hi;
asm volatile("mov.b64 {%0,%1}, %2;" : "=r"(lo), "=r"(hi) : "d"(value));
hi = __shfl_down(hi, delta, width);
lo = __shfl_down(lo, delta, width);
hi = __shfl_down_sync(mask, hi, delta, width);
lo = __shfl_down_sync(mask, lo, delta, width);
asm volatile("mov.b64 %0, {%1,%2};" : "=d"(value) : "r"(lo), "r"(hi));
return value;
}
template <typename T>
__device__ EIGEN_ALWAYS_INLINE T CudaShuffleXor(T value, int laneMask,
__device__ EIGEN_ALWAYS_INLINE T CudaShuffleXor(unsigned mask,
T value, int laneMask,
int width = warpSize) {
return __shfl_xor(value, laneMask, width);
return __shfl_xor_sync(mask, value, laneMask, width);
}
// Variant of the (undocumented) version from the CUDA SDK, but using unsigned
// instead of float for lo and hi (which is incorrect with ftz, for example).
// A bug has been filed with NVIDIA and will be fixed in the next CUDA release.
// TODO(csigg): remove when the bug is fixed in the next CUDA release.
__device__ EIGEN_ALWAYS_INLINE double CudaShuffleXor(double value, int laneMask,
__device__ EIGEN_ALWAYS_INLINE double CudaShuffleXor(unsigned mask,
double value, int laneMask,
int width = warpSize) {
unsigned lo, hi;
asm volatile("mov.b64 {%0,%1}, %2;" : "=r"(lo), "=r"(hi) : "d"(value));
hi = __shfl_xor(hi, laneMask, width);
lo = __shfl_xor(lo, laneMask, width);
hi = __shfl_xor_sync(mask, hi, laneMask, width);
lo = __shfl_xor_sync(mask, lo, laneMask, width);
asm volatile("mov.b64 %0, {%1,%2};" : "=d"(value) : "r"(lo), "r"(hi));
return value;
}

131
tensorflow/core/util/mkl_util.h
View File

@ -65,6 +65,8 @@ class MklShape {
void SetDimensions(const size_t dimension) { dimension_ = dimension; }
void SetMklLayout(dnnLayout_t mklLayout) { mklLayout_ = mklLayout; }
void SetMklLayout(const void* primitive, size_t resourceType) {
CHECK_EQ(
dnnLayoutCreateFromPrimitive_F32(&mklLayout_, (dnnPrimitive_t)primitive,
@ -135,6 +137,7 @@ class MklShape {
size_t GetDimension() const { return dimension_; }
const size_t* GetSizes() const { return sizes_; }
int64 dim_size(int index) const { return sizes_[index]; }
int64 tf_dim_size(int index) const { return sizes_[tf_to_mkl_dim_map_[index]]; }
const size_t* GetStrides() const { return strides_; }
const size_t* GetTfToMklDimMap() const { return tf_to_mkl_dim_map_; }
size_t tf_dim_idx(int index) const { return tf_to_mkl_dim_map_[index]; }
@ -581,7 +584,7 @@ inline void CopyTfTensorInToOutWithShape(OpKernelContext* context,
context->set_output(idx_data_out, output);
}
inline void FowardTfTensorInToOut(OpKernelContext* context,
inline void ForwardTfTensorInToOut(OpKernelContext* context,
int idx_in, int idx_out) {
int num_inputs = context->num_inputs();
int num_outputs = context->num_outputs();
@ -598,7 +601,7 @@ inline void FowardTfTensorInToOut(OpKernelContext* context,
}
}
inline void ForwarMklTensorInToOut(OpKernelContext* context,
inline void ForwardMklTensorInToOut(OpKernelContext* context,
int idx_in, int idx_out) {
int num_inputs = context->num_inputs();
int num_outputs = context->num_outputs();
@ -616,6 +619,98 @@ inline void ForwarMklTensorInToOut(OpKernelContext* context,
}
}
// Forward the MKL shape ONLY (used in elementwise and other ops where
// we call the eigen implementation and MKL shape is not used)
inline void ForwardMklMetaDataInToOut(OpKernelContext* context,
uint idx_data_in, uint idx_data_out) {
uint idx_meta_in = GetTensorMetaDataIndex(idx_data_in, context->num_inputs());
uint idx_meta_out =
GetTensorMetaDataIndex(idx_data_out, context->num_outputs());
if (IsRefType(context->input_dtype(idx_data_in))) {
context->forward_ref_input_to_ref_output(idx_meta_in, idx_meta_out);
} else {
context->set_output(idx_meta_out, context->input(idx_meta_in));
}
}
// Set a dummy MKL shape (called when the output is in TF format)
inline void SetDummyMklShapeOutput(OpKernelContext* context,
uint idx_data_out) {
MklShape mkl_shape_output;
mkl_shape_output.SetMklTensor(false);
AllocateOutputSetMklShape(context, idx_data_out, mkl_shape_output);
}
// Checks if the TF shape for both MKL tensors is the same or not
// Returns: true if both TF shapes are the same, false otherwise
inline bool MklCompareShapes(const MklShape* input_shape_0,
const MklShape* input_shape_1) {
// Check for number of dimensions
if (input_shape_0->GetDimension() != input_shape_1->GetDimension()) {
return false;
}
// Check size of each dimension
size_t ndims = input_shape_0->GetDimension();
for (size_t i = 0; i < ndims; i++) {
if (input_shape_0->dim_size(i) != input_shape_1->dim_size(i)) {
return false;
}
}
return true;
}
// Checks if the TF shape for both tensors is the same or not
// Returns: true if TF shapes for both are the same, false otherwise
inline bool MklCompareShapes(const MklShape* input_shape_0,
const TensorShape* input_shape_1) {
// Check for number of dimensions
if (input_shape_0->GetDimension() != input_shape_1->dims()) {
return false;
}
// Check size of each dimension
size_t ndims = input_shape_0->GetDimension();
for (size_t i = 0; i < ndims; i++) {
if (input_shape_0->tf_dim_size(i) != input_shape_1->dim_size(i)) {
return false;
}
}
return true;
}
// Checks if the TF shape for both tensors is the same or not
// Returns: true if TF shapes for both are the same, false otherwise
inline bool MklCompareShapes(const TensorShape* input_shape_0,
const MklShape* input_shape_1) {
return MklCompareShapes(input_shape_1, input_shape_0);
}
// Checks if the TF shape for both tensors is the same or not
// Returns: true if TF shapes for both are the same, false otherwise
inline bool MklCompareShapes(const TensorShape* input_shape_0,
const TensorShape* input_shape_1) {
// Check for number of dimensions
if (input_shape_0->dims() != input_shape_1->dims()) {
return false;
}
// Check size of each dimension
size_t ndims = input_shape_0->dims();
for (size_t i = 0; i < ndims; i++) {
if (input_shape_0->dim_size(i) != input_shape_1->dim_size(i)) {
return false;
}
}
return true;
}
// TODO(intel_tf): Remove this routine when faster MKL layout conversion is
// out.
inline void MklNHWCToNCHW(const Tensor& input, Tensor** output) {
const float* buf_in = input.flat<float>().data();
float* buf_out = (*output)->flat<float>().data();
@ -652,11 +747,19 @@ namespace mkl_op_registry {
static const char* kMklOpLabel = "MklOp";
static const char* kMklOpLabelPattern = "label='MklOp'";
// Get the name of Mkl op from original TensorFlow op
// We prefix 'Mkl' to the original op to get Mkl op.
inline string GetMklOpName(const string& name) {
// Prefix that we add to Tensorflow op name to construct Mkl op name.
const char* const kMklOpPrefix = "_Mkl";
return string(kMklOpPrefix) + name;
}
// Check whether opname with type T is registered as MKL-compliant.
//
// @input: name of the op
// @input: T datatype to be used for checking op
// @return: true if opname is registered as Mkl op
// @return: true if opname is registered as Mkl op; false otherwise
static inline bool IsMklOp(const std::string& op_name, DataType T) {
string kernel = KernelsRegisteredForOp(op_name);
bool result =
@ -667,6 +770,28 @@ static inline bool IsMklOp(const std::string& op_name, DataType T) {
return result;
}
// Check whether opname with type T is registered as MKL-compliant and
// is element-wise.
//
// @input: name of the op
// @input: T datatype to be used for checking op
// @return: true if opname is registered as element-wise Mkl op; false otherwise
static inline bool IsMklElementWiseOp(const std::string& op_name, DataType T) {
if (!IsMklOp(op_name, T)) {
return false;
}
bool result = (0 == op_name.compare(GetMklOpName("Add")) ||
0 == op_name.compare(GetMklOpName("Sub")) ||
0 == op_name.compare(GetMklOpName("Mul")) ||
0 == op_name.compare(GetMklOpName("Maximum")) ||
0 == op_name.compare(GetMklOpName("SquaredDifference")));
VLOG(1) << "mkl_op_registry::" << op_name
<< " is elementwise MKL op: " << result;
return result;
}
} // namespace mkl_op_registry
} // namespace tensorflow

2
tensorflow/docs_src/about/bib.md
View File

@ -37,7 +37,7 @@ system, we suggest you cite this whitepaper.
<pre>
@misc{tensorflow2015-whitepaper,
title={ {TensorFlow}: Large-Scale Machine Learning on Heterogeneous Systems},
url={http://tensorflow.org/},
url={https://www.tensorflow.org/},
note={Software available from tensorflow.org},
author={
Mart\'{\i}n~Abadi and

2
tensorflow/docs_src/extend/estimators.md
View File

@ -15,7 +15,7 @@ as regressors and classifiers:
Construct a neural network regression model.
* @{tf.estimator.DNNLinearCombinedClassifier}:
Construct a neural network and linear combined classification model.
* @{tf.estimator.DNNRegressor}:
* @{tf.estimator.DNNLinearCombinedRegressor}:
Construct a neural network and linear combined regression model.
But what if none of `tf.estimator`'s predefined model types meets your needs?

2
tensorflow/docs_src/get_started/get_started.md
View File

@ -447,7 +447,7 @@ estimator = tf.estimator.Estimator(model_fn=model_fn)
x_train = np.array([1., 2., 3., 4.])
y_train = np.array([0., -1., -2., -3.])
x_eval = np.array([2., 5., 8., 1.])
y_eval = np.array([-1.01, -4.1, -7, 0.])
y_eval = np.array([-1.01, -4.1, -7., 0.])
input_fn = tf.estimator.inputs.numpy_input_fn(
{"x": x_train}, y_train, batch_size=4, num_epochs=None, shuffle=True)
train_input_fn = tf.estimator.inputs.numpy_input_fn(

2
tensorflow/docs_src/install/install_c.md
View File

@ -35,7 +35,7 @@ enable TensorFlow for C:
OS="linux" # Change to "darwin" for Mac OS
TARGET_DIRECTORY="/usr/local"
curl -L \
"https://storage.googleapis.com/tensorflow/libtensorflow/libtensorflow-${TF_TYPE}-${OS}-x86_64-1.3.0.tar.gz" |
"https://storage.googleapis.com/tensorflow/libtensorflow/libtensorflow-${TF_TYPE}-${OS}-x86_64-1.4.0-dev.tar.gz" |
sudo tar -C $TARGET_DIRECTORY -xz
The `tar` command extracts the TensorFlow C library into the `lib`

2
tensorflow/docs_src/install/install_go.md
View File

@ -35,7 +35,7 @@ steps to install this library and enable TensorFlow for Go:
TF_TYPE="cpu" # Change to "gpu" for GPU support
TARGET_DIRECTORY='/usr/local'
curl -L \
"https://storage.googleapis.com/tensorflow/libtensorflow/libtensorflow-${TF_TYPE}-$(go env GOOS)-x86_64-1.3.0.tar.gz" |
"https://storage.googleapis.com/tensorflow/libtensorflow/libtensorflow-${TF_TYPE}-$(go env GOOS)-x86_64-1.4.0-dev.tar.gz" |
sudo tar -C $TARGET_DIRECTORY -xz
The `tar` command extracts the TensorFlow C library into the `lib`

18
tensorflow/docs_src/install/install_java.md
View File

@ -34,7 +34,7 @@ following to the project's `pom.xml` to use the TensorFlow Java APIs:
<dependency>
<groupId>org.tensorflow</groupId>
<artifactId>tensorflow</artifactId>
<version>1.3.0</version>
<version>1.4.0-dev</version>
</dependency>
```
@ -63,7 +63,7 @@ As an example, these steps will create a Maven project that uses TensorFlow:
<dependency>
<groupId>org.tensorflow</groupId>
<artifactId>tensorflow</artifactId>
<version>1.3.0</version>
<version>1.4.0-dev</version>
</dependency>
</dependencies>
</project>
@ -122,7 +122,7 @@ refer to the simpler instructions above instead.
Take the following steps to install TensorFlow for Java on Linux or Mac OS:
1. Download
[libtensorflow.jar](https://storage.googleapis.com/tensorflow/libtensorflow/libtensorflow-1.3.0.jar),
[libtensorflow.jar](https://storage.googleapis.com/tensorflow/libtensorflow/libtensorflow-1.4.0-dev.jar),
which is the TensorFlow Java Archive (JAR).
2. Decide whether you will run TensorFlow for Java on CPU(s) only or with
@ -141,7 +141,7 @@ Take the following steps to install TensorFlow for Java on Linux or Mac OS:
OS=$(uname -s | tr '[:upper:]' '[:lower:]')
mkdir -p ./jni
curl -L \
"https://storage.googleapis.com/tensorflow/libtensorflow/libtensorflow_jni-${TF_TYPE}-${OS}-x86_64-1.3.0.tar.gz" |
"https://storage.googleapis.com/tensorflow/libtensorflow/libtensorflow_jni-${TF_TYPE}-${OS}-x86_64-1.4.0-dev.tar.gz" |
tar -xz -C ./jni
### Install on Windows
@ -149,10 +149,10 @@ Take the following steps to install TensorFlow for Java on Linux or Mac OS:
Take the following steps to install TensorFlow for Java on Windows:
1. Download
[libtensorflow.jar](https://storage.googleapis.com/tensorflow/libtensorflow/libtensorflow-1.3.0.jar),
[libtensorflow.jar](https://storage.googleapis.com/tensorflow/libtensorflow/libtensorflow-1.4.0-dev.jar),
which is the TensorFlow Java Archive (JAR).
2. Download the following Java Native Interface (JNI) file appropriate for
[TensorFlow for Java on Windows](https://storage.googleapis.com/tensorflow/libtensorflow/libtensorflow_jni-cpu-windows-x86_64-1.3.0.zip).
[TensorFlow for Java on Windows](https://storage.googleapis.com/tensorflow/libtensorflow/libtensorflow_jni-cpu-windows-x86_64-1.4.0-dev.zip).
3. Extract this .zip file.
@ -200,7 +200,7 @@ must be part of your `classpath`. For example, you can include the
downloaded `.jar` in your `classpath` by using the `-cp` compilation flag
as follows:
<pre><b>javac -cp libtensorflow-1.3.0.jar HelloTF.java</b></pre>
<pre><b>javac -cp libtensorflow-1.4.0-dev.jar HelloTF.java</b></pre>
### Running
@ -214,11 +214,11 @@ two files are available to the JVM:
For example, the following command line executes the `HelloTF` program on Linux
and Mac OS X:
<pre><b>java -cp libtensorflow-1.3.0.jar:. -Djava.library.path=./jni HelloTF</b></pre>
<pre><b>java -cp libtensorflow-1.4.0-dev.jar:. -Djava.library.path=./jni HelloTF</b></pre>
And the following command line executes the `HelloTF` program on Windows:
<pre><b>java -cp libtensorflow-1.3.0.jar;. -Djava.library.path=jni HelloTF</b></pre>
<pre><b>java -cp libtensorflow-1.4.0-dev.jar;. -Djava.library.path=jni HelloTF</b></pre>
If the program prints <tt>Hello from <i>version</i></tt>, you've successfully
installed TensorFlow for Java and are ready to use the API. If the program

22
tensorflow/docs_src/install/install_linux.md
View File

@ -172,7 +172,7 @@ Take the following steps to install TensorFlow with Virtualenv:
virtualenv environment:
<pre>(tensorflow)$ <b>pip3 install --upgrade \
https://storage.googleapis.com/tensorflow/linux/cpu/tensorflow-1.3.0-cp34-cp34m-linux_x86_64.whl</b></pre>
https://storage.googleapis.com/tensorflow/linux/cpu/tensorflow-1.4.0dev-cp34-cp34m-linux_x86_64.whl</b></pre>
If you encounter installation problems, see
[Common Installation Problems](#common_installation_problems).
@ -277,7 +277,7 @@ take the following steps:
<pre>
$ <b>sudo pip3 install --upgrade \
https://storage.googleapis.com/tensorflow/linux/cpu/tensorflow-1.3.0-cp34-cp34m-linux_x86_64.whl</b>
https://storage.googleapis.com/tensorflow/linux/cpu/tensorflow-1.4.0dev-cp34-cp34m-linux_x86_64.whl</b>
</pre>
If this step fails, see
@ -464,7 +464,7 @@ Take the following steps to install TensorFlow in an Anaconda environment:
<pre>
(tensorflow)$ <b>pip install --ignore-installed --upgrade \
https://storage.googleapis.com/tensorflow/linux/cpu/tensorflow-1.3.0-cp34-cp34m-linux_x86_64.whl</b></pre>
https://storage.googleapis.com/tensorflow/linux/cpu/tensorflow-1.4.0dev-cp34-cp34m-linux_x86_64.whl</b></pre>
<a name="ValidateYourInstallation"></a>
@ -632,14 +632,14 @@ This section documents the relevant values for Linux installations.
CPU only:
<pre>
https://storage.googleapis.com/tensorflow/linux/cpu/tensorflow-1.3.0-cp27-none-linux_x86_64.whl
https://storage.googleapis.com/tensorflow/linux/cpu/tensorflow-1.4.0dev-cp27-none-linux_x86_64.whl
</pre>
GPU support:
<pre>
https://storage.googleapis.com/tensorflow/linux/gpu/tensorflow_gpu-1.3.0-cp27-none-linux_x86_64.whl
https://storage.googleapis.com/tensorflow/linux/gpu/tensorflow_gpu-1.4.0dev-cp27-none-linux_x86_64.whl
</pre>
Note that GPU support requires the NVIDIA hardware and software described in
@ -651,14 +651,14 @@ Note that GPU support requires the NVIDIA hardware and software described in
CPU only:
<pre>
https://storage.googleapis.com/tensorflow/linux/cpu/tensorflow-1.3.0-cp34-cp34m-linux_x86_64.whl
https://storage.googleapis.com/tensorflow/linux/cpu/tensorflow-1.4.0dev-cp34-cp34m-linux_x86_64.whl
</pre>
GPU support:
<pre>
https://storage.googleapis.com/tensorflow/linux/gpu/tensorflow_gpu-1.3.0-cp34-cp34m-linux_x86_64.whl
https://storage.googleapis.com/tensorflow/linux/gpu/tensorflow_gpu-1.4.0dev-cp34-cp34m-linux_x86_64.whl
</pre>
Note that GPU support requires the NVIDIA hardware and software described in
@ -670,14 +670,14 @@ Note that GPU support requires the NVIDIA hardware and software described in
CPU only:
<pre>
https://storage.googleapis.com/tensorflow/linux/cpu/tensorflow-1.3.0-cp35-cp35m-linux_x86_64.whl
https://storage.googleapis.com/tensorflow/linux/cpu/tensorflow-1.4.0dev-cp35-cp35m-linux_x86_64.whl
</pre>
GPU support:
<pre>
https://storage.googleapis.com/tensorflow/linux/gpu/tensorflow_gpu-1.3.0-cp35-cp35m-linux_x86_64.whl
https://storage.googleapis.com/tensorflow/linux/gpu/tensorflow_gpu-1.4.0dev-cp35-cp35m-linux_x86_64.whl
</pre>
@ -689,14 +689,14 @@ Note that GPU support requires the NVIDIA hardware and software described in
CPU only:
<pre>
https://storage.googleapis.com/tensorflow/linux/cpu/tensorflow-1.3.0-cp36-cp36m-linux_x86_64.whl
https://storage.googleapis.com/tensorflow/linux/cpu/tensorflow-1.4.0dev-cp36-cp36m-linux_x86_64.whl
</pre>
GPU support:
<pre>
https://storage.googleapis.com/tensorflow/linux/gpu/tensorflow_gpu-1.3.0-cp36-cp36m-linux_x86_64.whl
https://storage.googleapis.com/tensorflow/linux/gpu/tensorflow_gpu-1.4.0dev-cp36-cp36m-linux_x86_64.whl
</pre>

10
tensorflow/docs_src/install/install_mac.md
View File

@ -109,7 +109,7 @@ Take the following steps to install TensorFlow with Virtualenv:
TensorFlow in the active Virtualenv is as follows:
<pre> $ <b>pip3 install --upgrade \
https://storage.googleapis.com/tensorflow/mac/cpu/tensorflow-1.3.0-py2-none-any.whl</b></pre>
https://storage.googleapis.com/tensorflow/mac/cpu/tensorflow-1.4.0dev-py2-none-any.whl</b></pre>
If you encounter installation problems, see
[Common Installation Problems](#common-installation-problems).
@ -230,7 +230,7 @@ take the following steps:
issue the following command:
<pre> $ <b>sudo pip3 install --upgrade \
https://storage.googleapis.com/tensorflow/mac/cpu/tensorflow-1.3.0-py2-none-any.whl</b> </pre>
https://storage.googleapis.com/tensorflow/mac/cpu/tensorflow-1.4.0dev-py2-none-any.whl</b> </pre>
If the preceding command fails, see
[installation problems](#common-installation-problems).
@ -339,7 +339,7 @@ Take the following steps to install TensorFlow in an Anaconda environment:
TensorFlow for Python 2.7:
<pre> (tensorflow)$ <b>pip install --ignore-installed --upgrade \
https://storage.googleapis.com/tensorflow/mac/cpu/tensorflow-1.3.0-py2-none-any.whl</b></pre>
https://storage.googleapis.com/tensorflow/mac/cpu/tensorflow-1.4.0dev-py2-none-any.whl</b></pre>
<a name="ValidateYourInstallation"></a>
@ -512,7 +512,7 @@ This section documents the relevant values for Mac OS installations.
<pre>
https://storage.googleapis.com/tensorflow/mac/cpu/tensorflow-1.3.0-py2-none-any.whl
https://storage.googleapis.com/tensorflow/mac/cpu/tensorflow-1.4.0dev-py2-none-any.whl
</pre>
@ -520,7 +520,7 @@ https://storage.googleapis.com/tensorflow/mac/cpu/tensorflow-1.3.0-py2-none-any.
<pre>
https://storage.googleapis.com/tensorflow/mac/cpu/tensorflow-1.3.0-py3-none-any.whl
https://storage.googleapis.com/tensorflow/mac/cpu/tensorflow-1.4.0dev-py3-none-any.whl
</pre>

4
tensorflow/docs_src/install/install_sources.md
View File

@ -342,10 +342,10 @@ Invoke `pip install` to install that pip package.
The filename of the `.whl` file depends on your platform.
For example, the following command will install the pip package
for TensorFlow 1.3.0 on Linux:
for TensorFlow 1.4.0dev on Linux:
<pre>
$ <b>sudo pip install /tmp/tensorflow_pkg/tensorflow-1.3.0-py2-none-any.whl</b>
$ <b>sudo pip install /tmp/tensorflow_pkg/tensorflow-1.4.0dev-py2-none-any.whl</b>
</pre>
## Validate your installation

3
tensorflow/docs_src/install/install_windows.md
View File

@ -153,6 +153,9 @@ TensorFlow}.
If the system outputs an error message instead of a greeting, see [Common
installation problems](#common_installation_problems).
There is also a helpful [script](https://gist.github.com/mrry/ee5dbcfdd045fa48a27d56664411d41c)
for Windows TensorFlow installation issues.
## Common installation problems
We are relying on Stack Overflow to document TensorFlow installation problems

2
tensorflow/docs_src/programmers_guide/graphs.md
View File

@ -319,7 +319,7 @@ described below.
* **`target`.** If this argument is left empty (the default), the session will
only use devices in the local machine. However, you may also specify a
`grpc://` URL to specify the address of a TensorFlow server, which gives the
session access to all devices on machines that that server controls. See
session access to all devices on machines that this server controls. See
@{tf.train.Server} for details of how to create a TensorFlow
server. For example, in the common **between-graph replication**
configuration, the @{tf.Session} connects to a @{tf.train.Server} in the same

2
tensorflow/docs_src/programmers_guide/tensors.md
View File

@ -147,7 +147,7 @@ Passing a single number, however, returns a subvector of a matrix, as follows:
```python
my_row_vetor = my_matrix[2]
my_row_vector = my_matrix[2]
my_column_vector = my_matrix[:, 3]
```

2
tensorflow/docs_src/tutorials/layers.md
View File

@ -270,7 +270,7 @@ The `padding` argument specifies one of two enumerated values
(case-insensitive): `valid` (default value) or `same`. To specify that the
output tensor should have the same width and height values as the input tensor,
we set `padding=same` here, which instructs TensorFlow to add 0 values to the
edges of the output tensor to preserve width and height of 28. (Without padding,
edges of the input tensor to preserve width and height of 28. (Without padding,
a 5x5 convolution over a 28x28 tensor will produce a 24x24 tensor, as there are
24x24 locations to extract a 5x5 tile from a 28x28 grid.)

2
tensorflow/examples/android/README.md
View File

@ -37,7 +37,7 @@ on API >= 14 devices.
4. [TF
Speech](https://github.com/tensorflow/tensorflow/blob/master/tensorflow/examples/android/src/org/tensorflow/demo/SpeechActivity.java):
Runs a simple speech recognition model built by the [audio training
tutorial](https://www.tensorflow.org/tutorials/image_retraining). Listens
tutorial](https://www.tensorflow.org/versions/master/tutorials/audio_recognition). Listens
for a small set of words, and highlights them in the UI when they are
recognized.

2
tensorflow/examples/android/src/org/tensorflow/demo/CameraActivity.java
View File

@ -110,7 +110,7 @@ public abstract class CameraActivity extends Activity implements OnImageAvailabl
rgbBytes = new int[previewWidth * previewHeight];
onPreviewSizeChosen(new Size(previewSize.width, previewSize.height), 90);
}
ImageUtils.convertYUV420SPToARGB8888(bytes, rgbBytes, previewWidth, previewHeight, false);
ImageUtils.convertYUV420SPToARGB8888(bytes, previewWidth, previewHeight, rgbBytes);
} catch (final Exception e) {
LOGGER.e(e, "Exception!");
return;

126
tensorflow/examples/android/src/org/tensorflow/demo/env/ImageUtils.java vendored
View File

@ -27,7 +27,7 @@ import java.io.FileOutputStream;
public class ImageUtils {
@SuppressWarnings("unused")
private static final Logger LOGGER = new Logger();
static {
try {
System.loadLibrary("tensorflow_demo");
@ -98,6 +98,66 @@ public class ImageUtils {
// Always prefer the native implementation if available.
private static boolean useNativeConversion = true;
public static void convertYUV420SPToARGB8888(
byte[] input,
int width,
int height,
int[] output) {
if (useNativeConversion) {
try {
ImageUtils.convertYUV420SPToARGB8888(input, output, width, height, false);
return;
} catch (UnsatisfiedLinkError e) {
LOGGER.w(
"Native YUV420SP -> RGB implementation not found, falling back to Java implementation");
useNativeConversion = false;
}
}
// Java implementation of YUV420SP to ARGB8888 converting
final int frameSize = width * height;
for (int j = 0, yp = 0; j < height; j++) {
int uvp = frameSize + (j >> 1) * width;
int u = 0;
int v = 0;
for (int i = 0; i < width; i++, yp++) {
int y = 0xff & input[yp];
if ((i & 1) == 0) {
v = 0xff & input[uvp++];
u = 0xff & input[uvp++];
}
output[yp] = YUV2RGB(y, u, v);
}
}
}
private static int YUV2RGB(int y, int u, int v) {
// Adjust and check YUV values
y = (y - 16) < 0 ? 0 : (y - 16);
u -= 128;
v -= 128;
// This is the floating point equivalent. We do the conversion in integer
// because some Android devices do not have floating point in hardware.
// nR = (int)(1.164 * nY + 2.018 * nU);
// nG = (int)(1.164 * nY - 0.813 * nV - 0.391 * nU);
// nB = (int)(1.164 * nY + 1.596 * nV);
int y1192 = 1192 * y;
int r = (y1192 + 1634 * v);
int g = (y1192 - 833 * v - 400 * u);
int b = (y1192 + 2066 * u);
// Clipping RGB values to be inside boundaries [ 0 , kMaxChannelValue ]
r = r > kMaxChannelValue ? kMaxChannelValue : (r < 0 ? 0 : r);
g = g > kMaxChannelValue ? kMaxChannelValue : (g < 0 ? 0 : g);
b = b > kMaxChannelValue ? kMaxChannelValue : (b < 0 ? 0 : b);
return 0xff000000 | ((r << 6) & 0xff0000) | ((g >> 2) & 0xff00) | ((b >> 10) & 0xff);
}
public static void convertYUV420ToARGB8888(
byte[] yData,
byte[] uData,
@ -114,56 +174,28 @@ public class ImageUtils {
yData, uData, vData, out, width, height, yRowStride, uvRowStride, uvPixelStride, false);
return;
} catch (UnsatisfiedLinkError e) {
LOGGER.w("Native YUV -> RGB implementation not found, falling back to Java implementation");
LOGGER.w(
"Native YUV420 -> RGB implementation not found, falling back to Java implementation");
useNativeConversion = false;
}
}
int i = 0;
for (int y = 0; y < height; y++) {
int pY = yRowStride * y;
int uv_row_start = uvRowStride * (y >> 1);
int pUV = uv_row_start;
int pV = uv_row_start;
int yp = 0;
for (int j = 0; j < height; j++) {
int pY = yRowStride * j;
int pUV = uvRowStride * (j >> 1);
for (int x = 0; x < width; x++) {
int uv_offset = pUV + (x >> 1) * uvPixelStride;
out[i++] =
YUV2RGB(
convertByteToInt(yData, pY + x),
convertByteToInt(uData, uv_offset),
convertByteToInt(vData, uv_offset));
for (int i = 0; i < width; i++) {
int uv_offset = pUV + (i >> 1) * uvPixelStride;
out[yp++] = YUV2RGB(
0xff & yData[pY + i],
0xff & uData[uv_offset],
0xff & vData[uv_offset]);
}
}
}
private static int convertByteToInt(byte[] arr, int pos) {
return arr[pos] & 0xFF;
}
private static int YUV2RGB(int nY, int nU, int nV) {
nY -= 16;
nU -= 128;
nV -= 128;
if (nY < 0) nY = 0;
// This is the floating point equivalent. We do the conversion in integer
// because some Android devices do not have floating point in hardware.
// nR = (int)(1.164 * nY + 2.018 * nU);
// nG = (int)(1.164 * nY - 0.813 * nV - 0.391 * nU);
// nB = (int)(1.164 * nY + 1.596 * nV);
final int foo = 1192 * nY;
int nR = foo + 1634 * nV;
int nG = foo - 833 * nV - 400 * nU;
int nB = foo + 2066 * nU;
nR = Math.min(kMaxChannelValue, Math.max(0, nR));
nG = Math.min(kMaxChannelValue, Math.max(0, nG));
nB = Math.min(kMaxChannelValue, Math.max(0, nB));
return 0xff000000 | ((nR << 6) & 0x00ff0000) | ((nG >> 2) & 0x0000FF00) | ((nB >> 10) & 0xff);
}
/**
* Converts YUV420 semi-planar data to ARGB 8888 data using the supplied width and height. The
@ -176,7 +208,7 @@ public class ImageUtils {
* @param height The height of the input image.
* @param halfSize If true, downsample to 50% in each dimension, otherwise not.
*/
public static native void convertYUV420SPToARGB8888(
private static native void convertYUV420SPToARGB8888(
byte[] input, int[] output, int width, int height, boolean halfSize);
/**
@ -193,7 +225,7 @@ public class ImageUtils {
* @param halfSize If true, downsample to 50% in each dimension, otherwise not.
* @param output A pre-allocated array for the ARGB 8:8:8:8 output data.
*/
public static native void convertYUV420ToARGB8888(
private static native void convertYUV420ToARGB8888(
byte[] y,
byte[] u,
byte[] v,
@ -215,7 +247,7 @@ public class ImageUtils {
* @param width The width of the input image.
* @param height The height of the input image.
*/
public static native void convertYUV420SPToRGB565(
private static native void convertYUV420SPToRGB565(
byte[] input, byte[] output, int width, int height);
/**
@ -228,7 +260,7 @@ public class ImageUtils {
* @param width The width of the input image.
* @param height The height of the input image.
*/
public static native void convertARGB8888ToYUV420SP(
private static native void convertARGB8888ToYUV420SP(
int[] input, byte[] output, int width, int height);
/**
@ -241,7 +273,7 @@ public class ImageUtils {
* @param width The width of the input image.
* @param height The height of the input image.
*/
public static native void convertRGB565ToYUV420SP(
private static native void convertRGB565ToYUV420SP(
byte[] input, byte[] output, int width, int height);
/**

2
tensorflow/examples/ios/README.md
View File

@ -30,7 +30,7 @@ cp ~/graphs/inception5h/* tensorflow/examples/ios/simple/data/
long time since it is big (~450MB). For example, if you want to run the
simple example, then:
```bash
cd tensorflow/ios/simple
cd tensorflow/examples/ios/simple
pod install
open tf_simple_example.xcworkspace # obs, not the .xcodeproj directory
```

7
tensorflow/examples/speech_commands/train.py
View File

@ -15,9 +15,10 @@
r"""Simple speech recognition to spot a limited number of keywords.
This is a self-contained example script that will train a very basic audio
recognition model in TensorFlow. It can download the necessary training data,
and runs with reasonable defaults to train within a few hours even only using a
CPU. For more information see http://tensorflow.org/tutorials/audio_recognition.
recognition model in TensorFlow. It downloads the necessary training data and
runs with reasonable defaults to train within a few hours even only using a CPU.
For more information, please see
https://www.tensorflow.org/tutorials/audio_recognition.
It is intended as an introduction to using neural networks for audio
recognition, and is not a full speech recognition system. For more advanced

3
tensorflow/java/src/main/java/org/tensorflow/Tensor.java
View File

@ -515,8 +515,6 @@ public final class Tensor implements AutoCloseable {
private static int elemByteSize(DataType dataType) {
switch (dataType) {
case UINT8:
return 1;
case FLOAT:
case INT32:
return 4;
@ -524,6 +522,7 @@ public final class Tensor implements AutoCloseable {
case INT64:
return 8;
case BOOL:
case UINT8:
return 1;
case STRING:
throw new IllegalArgumentException("STRING tensors do not have a fixed element size");

6
tensorflow/java/src/main/native/tensor_jni.cc
View File

@ -41,8 +41,11 @@ size_t elemByteSize(TF_DataType dtype) {
// have the same byte sizes. Validate that:
switch (dtype) {
case TF_BOOL:
case TF_UINT8:
static_assert(sizeof(jboolean) == 1,
"Java boolean not compatible with TF_BOOL");
static_assert(sizeof(jbyte) == 1,
"Java byte not compatible with TF_UINT8");
return 1;
case TF_FLOAT:
case TF_INT32:
@ -90,6 +93,7 @@ void writeScalar(JNIEnv* env, jobject src, TF_DataType dtype, void* dst,
CASE(TF_DOUBLE, jdouble, "doubleValue", "()D", Double);
CASE(TF_INT32, jint, "intValue", "()I", Int);
CASE(TF_INT64, jlong, "longValue", "()J", Long);
CASE(TF_UINT8, jbyte, "byteValue", "()B", Byte);
#undef CASE
case TF_BOOL: {
jclass clazz = env->FindClass("java/lang/Boolean");
@ -134,6 +138,7 @@ size_t write1DArray(JNIEnv* env, jarray array, TF_DataType dtype, void* dst,
CASE(TF_INT32, jint, Int);
CASE(TF_INT64, jlong, Long);
CASE(TF_BOOL, jboolean, Boolean);
CASE(TF_UINT8, jbyte, Byte);
#undef CASE
default:
throwException(env, kIllegalStateException, "invalid DataType(%d)",
@ -168,6 +173,7 @@ size_t read1DArray(JNIEnv* env, TF_DataType dtype, const void* src,
CASE(TF_INT32, jint, Int);
CASE(TF_INT64, jlong, Long);
CASE(TF_BOOL, jboolean, Boolean);
CASE(TF_UINT8, jbyte, Byte);
#undef CASE
default:
throwException(env, kIllegalStateException, "invalid DataType(%d)",

17
tensorflow/python/BUILD
View File

@ -3785,6 +3785,7 @@ py_library(
"layers/convolutional.py",
"layers/core.py",
"layers/layers.py",
"layers/maxout.py",
"layers/normalization.py",
"layers/pooling.py",
],
@ -3865,6 +3866,22 @@ py_test(
],
)
py_test(
name = "layers_maxout_test",
size = "small",
srcs = ["layers/maxout_test.py"],
main = "layers/maxout_test.py",
srcs_version = "PY2AND3",
deps = [
":client_testlib",
":framework_for_generated_wrappers",
":layers",
":math_ops",
":nn_ops",
":random_ops",
],
)
py_test(
name = "layers_utils_test",
size = "small",

16
tensorflow/python/kernel_tests/decode_raw_op_test.py
View File

@ -93,6 +93,22 @@ class DecodeRawOpTest(test.TestCase):
result = decode.eval(feed_dict={in_bytes: [""]})
self.assertEqual(len(result), 1)
def testToUInt16(self):
with self.test_session():
in_bytes = array_ops.placeholder(dtypes.string, shape=[None])
decode = parsing_ops.decode_raw(in_bytes, out_type=dtypes.uint16)
self.assertEqual([None, None], decode.get_shape().as_list())
# Use FF/EE/DD/CC so that decoded value is higher than 32768 for uint16
result = decode.eval(feed_dict={in_bytes: [b"\xFF\xEE\xDD\xCC"]})
self.assertAllEqual(
[[0xFF + 0xEE * 256, 0xDD + 0xCC * 256]], result)
with self.assertRaisesOpError(
"Input to DecodeRaw has length 3 that is not a multiple of 2, the "
"size of uint16"):
decode.eval(feed_dict={in_bytes: ["123", "456"]})
if __name__ == "__main__":
test.main()

4
tensorflow/python/kernel_tests/metrics_test.py
View File

@ -3538,7 +3538,7 @@ class MeanPerClassAccuracyTest(test.TestCase):
weights_queue = data_flow_ops.FIFOQueue(
6, dtypes=dtypes_lib.float32, shapes=(1, 1))
_enqueue_vector(sess, weights_queue, [1.0])
_enqueue_vector(sess, weights_queue, [1.0])
_enqueue_vector(sess, weights_queue, [0.5])
_enqueue_vector(sess, weights_queue, [1.0])
_enqueue_vector(sess, weights_queue, [0.0])
_enqueue_vector(sess, weights_queue, [1.0])
@ -3551,7 +3551,7 @@ class MeanPerClassAccuracyTest(test.TestCase):
variables.local_variables_initializer().run()
for _ in range(6):
sess.run(update_op)
desired_output = np.mean([2.0 / 2.0, 1.0 / 2.0])
desired_output = np.mean([2.0 / 2.0, 0.5 / 1.5])
self.assertAlmostEqual(desired_output, mean_accuracy.eval())
def testMultipleUpdatesWithMissingClass(self):

1
tensorflow/python/kernel_tests/segment_reduction_ops_test.py
View File

@ -645,7 +645,6 @@ class SparseSegmentReductionOpTest(SparseSegmentReductionHelper):
with self.assertRaisesOpError(r"Segment id 0 out of range \[0, 0\)"):
s.eval()
class SegmentReductionOpBenchmark(test.Benchmark):
outer_dim_options = [2**x for x in range(9, 14, 2)]
ratio_options = [2**x for x in range(1, 6, 2)]

47
tensorflow/python/kernel_tests/svd_op_test.py
View File

@ -41,23 +41,19 @@ class SvdOpTest(test.TestCase):
linalg_ops.svd(vector)
def _GetSvdOpTest(dtype_, shape_, use_static_shape_):
def _GetSvdOpTest(dtype_, shape_, use_static_shape_, use_gpu_):
is_complex = dtype_ in (np.complex64, np.complex128)
is_single = dtype_ in (np.float32, np.complex64)
# The gpu version returns results that are much less precise
precision_factor = 100 if use_gpu_ else 1
tol = precision_factor * (3e-4 if is_single else 1e-12)
def CompareSingularValues(self, x, y):
if is_single:
tol = 5e-5
else:
tol = 1e-14
self.assertAllClose(x, y, atol=(x[0] + y[0]) * tol)
def CompareSingularVectors(self, x, y, rank):
if is_single:
atol = 5e-4
else:
atol = 5e-14
# We only compare the first 'rank' singular vectors since the
# remainder form an arbitrary orthonormal basis for the
# (row- or column-) null space, whose exact value depends on
@ -72,13 +68,9 @@ def _GetSvdOpTest(dtype_, shape_, use_static_shape_):
sum_of_ratios = np.sum(np.divide(y, x), -2, keepdims=True)
phases = np.divide(sum_of_ratios, np.abs(sum_of_ratios))
x *= phases
self.assertAllClose(x, y, atol=atol)
self.assertAllClose(x, y, atol=2 * tol)
def CheckApproximation(self, a, u, s, v, full_matrices):
if is_single:
tol = 1e-5
else:
tol = 1e-14
# Tests that a ~= u*diag(s)*transpose(v).
batch_shape = a.shape[:-2]
m = a.shape[-2]
@ -99,10 +91,6 @@ def _GetSvdOpTest(dtype_, shape_, use_static_shape_):
# Tests that x[...,:,:]^H * x[...,:,:] is close to the identity.
xx = math_ops.matmul(x, x, adjoint_a=True)
identity = array_ops.matrix_band_part(array_ops.ones_like(xx), 0, 0)
if is_single:
tol = 1e-5
else:
tol = 1e-14
self.assertAllClose(identity.eval(), xx.eval(), atol=tol)
def Test(self):
@ -116,7 +104,7 @@ def _GetSvdOpTest(dtype_, shape_, use_static_shape_):
for compute_uv in False, True:
for full_matrices in False, True:
with self.test_session() as sess:
with self.test_session(use_gpu = use_gpu_) as sess:
if use_static_shape_:
x_tf = constant_op.constant(x_np)
else:
@ -167,14 +155,15 @@ def _GetSvdOpTest(dtype_, shape_, use_static_shape_):
if __name__ == "__main__":
for dtype in np.float32, np.float64, np.complex64, np.complex128:
for rows in 1, 2, 5, 10, 32, 100:
for cols in 1, 2, 5, 10, 32, 100:
for batch_dims in [(), (3,)] + [(3, 2)] * (max(rows, cols) < 10):
shape = batch_dims + (rows, cols)
for use_static_shape in True, False:
name = "%s_%s_%s" % (dtype.__name__, "_".join(map(str, shape)),
use_static_shape)
setattr(SvdOpTest, "testSvd_" + name,
_GetSvdOpTest(dtype, shape, use_static_shape))
for use_gpu in False, True:
for dtype in np.float32, np.float64, np.complex64, np.complex128:
for rows in 1, 2, 5, 10, 32, 100:
for cols in 1, 2, 5, 10, 32, 100:
for batch_dims in [(), (3,)] + [(3, 2)] * (max(rows, cols) < 10):
shape = batch_dims + (rows, cols)
for use_static_shape in True, False:
name = "%s_%s_%s_%s" % (dtype.__name__, "_".join(map(str, shape)),
use_static_shape, use_gpu)
setattr(SvdOpTest, "testSvd_" + name,
_GetSvdOpTest(dtype, shape, use_static_shape, use_gpu))
test.main()

10
tensorflow/python/layers/base.py
View File

@ -25,24 +25,20 @@ from __future__ import print_function
import collections
import copy
import functools
import re
import weakref
from six.moves import xrange # pylint: disable=redefined-builtin
import numpy as np
import six
from tensorflow.python.eager import context
from tensorflow.python.estimator import util as estimator_util
from tensorflow.python.framework import ops
from tensorflow.python.framework import dtypes
from tensorflow.python.framework import ops
from tensorflow.python.framework import tensor_shape
from tensorflow.python.ops import array_ops
from tensorflow.python.ops import variables as tf_variables
from tensorflow.python.ops import variable_scope as vs
from tensorflow.python.util import nest
from tensorflow.python.ops import variables as tf_variables
from tensorflow.python.platform import tf_logging as logging
from tensorflow.python.util import nest
class Layer(object):

16
tensorflow/python/layers/convolutional.py
View File

@ -20,23 +20,13 @@ from __future__ import absolute_import
from __future__ import division
from __future__ import print_function
import six
from six.moves import xrange # pylint: disable=redefined-builtin
import numpy as np
from tensorflow.python.eager import context
from tensorflow.python.framework import ops
from tensorflow.python.ops import array_ops
from tensorflow.python.ops import control_flow_ops
from tensorflow.python.ops import nn
from tensorflow.python.ops import math_ops
from tensorflow.python.ops import init_ops
from tensorflow.python.ops import standard_ops
from tensorflow.python.ops import variable_scope as vs
from tensorflow.python.framework import tensor_shape
from tensorflow.python.layers import base
from tensorflow.python.layers import utils
from tensorflow.python import framework
from tensorflow.python.ops import array_ops
from tensorflow.python.ops import init_ops
from tensorflow.python.ops import nn
class _Conv(base.Layer):

7
tensorflow/python/layers/core.py
View File

@ -22,6 +22,7 @@ from __future__ import absolute_import
from __future__ import division
from __future__ import print_function
import six
from six.moves import xrange # pylint: disable=redefined-builtin
import numpy as np
@ -29,15 +30,13 @@ import numpy as np
from tensorflow.python.eager import context
from tensorflow.python.framework import ops
from tensorflow.python.framework import tensor_shape
from tensorflow.python.layers import base
from tensorflow.python.layers import utils
from tensorflow.python.ops import array_ops
from tensorflow.python.ops import init_ops
from tensorflow.python.ops import math_ops
from tensorflow.python.ops import nn
from tensorflow.python.ops import standard_ops
from tensorflow.python.ops import variable_scope as vs
from tensorflow.python.layers import base
from tensorflow.python.layers import utils
class Dense(base.Layer):

108
tensorflow/python/layers/maxout.py Normal file
View File

@ -0,0 +1,108 @@
# 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.
# =============================================================================
# pylint: disable=unused-import,g-bad-import-order
"""Contains the maxout layer
"""
from __future__ import absolute_import
from __future__ import division
from __future__ import print_function
from tensorflow.python.framework import ops
from tensorflow.python.ops import math_ops
from tensorflow.python.ops import gen_array_ops
from tensorflow.python.layers import base
def maxout(inputs, num_units, axis=-1, name=None):
"""Adds a maxout op from https://arxiv.org/abs/1302.4389
"Maxout Networks" Ian J. Goodfellow, David Warde-Farley, Mehdi Mirza, Aaron Courville,
Yoshua Bengio
Usually the operation is performed in the filter/channel dimension. This can also be
used after fully-connected layers to reduce number of features.
Arguments:
inputs: Tensor input
num_units: Specifies how many features will remain after maxout in the `axis` dimension
(usually channel). This must be multiple of number of `axis`.
axis: The dimension where max pooling will be performed. Default is the
last dimension.
name: Optional scope for name_scope.
Returns:
A `Tensor` representing the results of the pooling operation.
Raises:
ValueError: if num_units is not multiple of number of features.
"""
return MaxOut(num_units=num_units, axis=axis, name=name)(inputs)
class MaxOut(base.Layer):
"""Adds a maxout op from https://arxiv.org/abs/1302.4389
"Maxout Networks" Ian J. Goodfellow, David Warde-Farley, Mehdi Mirza, Aaron Courville, Yoshua
Bengio
Usually the operation is performed in the filter/channel dimension. This can also be
used after fully-connected layers to reduce number of features.
Arguments:
inputs: Tensor input
num_units: Specifies how many features will remain after maxout in the `axis` dimension
(usually channel).
This must be multiple of number of `axis`.
axis: The dimension where max pooling will be performed. Default is the
last dimension.
name: Optional scope for name_scope.
Returns:
A `Tensor` representing the results of the pooling operation.
Raises:
ValueError: if num_units is not multiple of number of features.
"""
def __init__(self,
num_units,
axis=-1,
name=None,
**kwargs):
super(MaxOut, self).__init__(
name=name, trainable=False, **kwargs)
self.axis = axis
self.num_units = num_units
def call(self, inputs):
inputs = ops.convert_to_tensor(inputs)
shape = inputs.get_shape().as_list()
num_channels = shape[self.axis]
if num_channels % self.num_units:
raise ValueError('number of features({}) is not '
'a multiple of num_units({})'
.format(num_channels, self.num_units))
shape[self.axis] = -1
shape += [num_channels // self.num_units]
# Dealing with batches with arbitrary sizes
for i in range(len(shape)):
if shape[i] is None:
shape[i] = gen_array_ops.shape(inputs)[i]
outputs = math_ops.reduce_max(gen_array_ops.reshape(inputs, shape), -1, keep_dims=False)
return outputs

61
tensorflow/python/layers/maxout_test.py Normal file
View File

@ -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.
# =============================================================================
# pylint: disable=unused-import,g-bad-import-order
from __future__ import absolute_import
from __future__ import division
from __future__ import print_function
from tensorflow.python.layers import maxout
from tensorflow.python.layers import convolutional as conv_layers
from tensorflow.python.layers import core as core_layers
from tensorflow.python.ops import random_ops
from tensorflow.python.platform import test
import numpy as np
"""
Contains the maxout layer tests
"""
class MaxOutTest(test.TestCase):
def test_simple(self):
inputs = random_ops.random_uniform((64, 10, 36), seed=1)
graph = maxout.maxout(inputs, num_units=3)
self.assertEqual(graph.get_shape().as_list(), [64, 10, 3])
def test_fully_connected(self):
inputs = random_ops.random_uniform((64, 50), seed=1)
graph = core_layers.dense(inputs, 50)
graph = maxout.maxout(graph, num_units=10)
self.assertEqual(graph.get_shape().as_list(), [64, 10])
def test_nchw(self):
inputs = random_ops.random_uniform((10, 100, 100, 3), seed=1)
graph = conv_layers.conv2d(inputs, 10, 3, padding="SAME")
graph = maxout.maxout(graph, num_units=1)
self.assertEqual(graph.get_shape().as_list(), [10, 100, 100, 1])
def test_invalid_shape(self):
inputs = random_ops.random_uniform((10, 100, 100, 3), seed=1)
graph = conv_layers.conv2d(inputs, 3, 10, strides=(1, 1))
with self.assertRaisesRegexp(ValueError, 'number of features'):
graph = maxout.maxout(graph, num_units=2)
if __name__ == '__main__':
test.main()

12
tensorflow/python/layers/normalization.py
View File

@ -26,24 +26,18 @@ import numpy as np
from tensorflow.python.eager import context
from tensorflow.python.framework import constant_op
from tensorflow.python.framework import dtypes
from tensorflow.python.framework import tensor_shape
from tensorflow.python.framework import ops
from tensorflow.python.framework import tensor_shape
from tensorflow.python.layers import base
from tensorflow.python.layers import utils
from tensorflow.python.ops import array_ops
from tensorflow.python.ops import nn
from tensorflow.python.ops import gen_resource_variable_ops
from tensorflow.python.ops import resource_variable_ops
from tensorflow.python.ops import math_ops
from tensorflow.python.ops import init_ops
from tensorflow.python.ops import standard_ops
from tensorflow.python.ops import state_ops
from tensorflow.python.ops import variable_scope as vs
from tensorflow.python.training import moving_averages
from tensorflow.python.framework import tensor_util
from tensorflow.python.ops import variables
from tensorflow.python.layers import base
from tensorflow.python.layers import utils
class BatchNormalization(base.Layer):

14
tensorflow/python/layers/pooling.py
View File

@ -20,21 +20,11 @@ from __future__ import absolute_import
from __future__ import division
from __future__ import print_function
import six
from six.moves import xrange # pylint: disable=redefined-builtin
import numpy as np
from tensorflow.python.framework import ops
from tensorflow.python.ops import array_ops
from tensorflow.python.ops import control_flow_ops
from tensorflow.python.ops import nn
from tensorflow.python.ops import init_ops
from tensorflow.python.ops import standard_ops
from tensorflow.python.ops import variable_scope as vs
from tensorflow.python.framework import tensor_shape
from tensorflow.python.layers import base
from tensorflow.python.layers import utils
from tensorflow.python import framework
from tensorflow.python.ops import array_ops
from tensorflow.python.ops import nn
class _Pooling1D(base.Layer):

5
tensorflow/python/layers/utils.py
View File

@ -20,13 +20,8 @@ from __future__ import absolute_import
from __future__ import division
from __future__ import print_function
import six
from six.moves import xrange # pylint: disable=redefined-builtin
import numpy as np
from tensorflow.python.ops import variables
from tensorflow.python.ops import control_flow_ops
from tensorflow.python.ops import math_ops
from tensorflow.python.framework import ops
from tensorflow.python.framework import tensor_util

5
tensorflow/python/ops/metrics_impl.py
View File

@ -259,11 +259,10 @@ def _streaming_confusion_matrix(labels, predictions, num_classes, weights=None):
update_op: An operation that increments the confusion matrix.
"""
# Local variable to accumulate the predictions in the confusion matrix.
cm_dtype = dtypes.int64 if weights is not None else dtypes.float64
total_cm = _create_local(
'total_confusion_matrix',
shape=[num_classes, num_classes],
dtype=cm_dtype)
dtype=dtypes.float64)
# Cast the type to int64 required by confusion_matrix_ops.
predictions = math_ops.to_int64(predictions)
@ -282,7 +281,7 @@ def _streaming_confusion_matrix(labels, predictions, num_classes, weights=None):
# Accumulate the prediction to current confusion matrix.
current_cm = confusion_matrix.confusion_matrix(
labels, predictions, num_classes, weights=weights, dtype=cm_dtype)
labels, predictions, num_classes, weights=weights, dtype=dtypes.float64)
update_op = state_ops.assign_add(total_cm, current_cm)
return total_cm, update_op

12
tensorflow/python/training/moving_averages.py
View File

@ -278,14 +278,12 @@ class ExponentialMovingAverage(object):
# Create an ExponentialMovingAverage object
ema = tf.train.ExponentialMovingAverage(decay=0.9999)
# Create the shadow variables, and add ops to maintain moving averages
# of var0 and var1.
maintain_averages_op = ema.apply([var0, var1])
# Create an op that will update the moving averages after each training
# step. This is what we will use in place of the usual training op.
with tf.control_dependencies([opt_op]):
training_op = tf.group(maintain_averages_op)
# Create the shadow variables, and add ops to maintain moving averages
# of var0 and var1. This also creates an op that will update the moving
# averages after each training step. This is what we will use in place
# of the usual training op.
training_op = ema.apply([var0, var1])
...train the model by running training_op...
```

1
tensorflow/tensorflow.bzl
View File

@ -980,6 +980,7 @@ check_deps = rule(
def tf_custom_op_library(name, srcs=[], gpu_srcs=[], deps=[]):
cuda_deps = [
clean_dep("//tensorflow/core:stream_executor_headers_lib"),
"@local_config_cuda//cuda:cuda_headers",
"@local_config_cuda//cuda:cudart_static",
]
deps = deps + tf_custom_op_library_additional_deps()

4
tensorflow/tf_exported_symbols.lds
View File

@ -1,6 +1,6 @@
*tensorflow*
*perftools*gputools*
*tf_*
TF_*
TFE_*
*TF_*
*TFE_*
*nsync_*

4
tensorflow/tf_version_script.lds
View File

@ -2,8 +2,8 @@ tensorflow {
global:
*tensorflow*;
*perftools*gputools*;
TF_*;
TFE_*;
*TF_*;
*TFE_*;
*nsync_*;
local:
*;

3
tensorflow/tools/ci_build/windows/cpu/pip/build_tf_windows.sh
View File

@ -60,8 +60,11 @@ reinstall_tensorflow_pip ${PIP_NAME}
# Define no_tensorflow_py_deps=true so that every py_test has no deps anymore,
# which will result testing system installed tensorflow
# TODO(pcloudy): Remove TF_SAVER_LENIENT_NAMES after
# https://github.com/tensorflow/tensorflow/issues/12844 is fixed.
bazel test -c opt $BUILD_OPTS -k --test_output=errors \
--define=no_tensorflow_py_deps=true --test_lang_filters=py \
--test_tag_filters=-no_pip,-no_windows \
--build_tag_filters=-no_pip,-no_windows --build_tests_only \
--test_env=TF_SAVER_LENIENT_NAMES=True \
//${PY_TEST_DIR}/tensorflow/python/...

3
tensorflow/tools/ci_build/windows/gpu/pip/build_tf_windows.sh
View File

@ -61,8 +61,11 @@ reinstall_tensorflow_pip ${PIP_NAME}
# Define no_tensorflow_py_deps=true so that every py_test has no deps anymore,
# which will result testing system installed tensorflow
# GPU tests are very flaky when running concurrently, so set local_test_jobs=1
# TODO(pcloudy): Remove TF_SAVER_LENIENT_NAMES after
# https://github.com/tensorflow/tensorflow/issues/12844 is fixed.
bazel test -c opt $BUILD_OPTS -k --test_output=errors \
--define=no_tensorflow_py_deps=true --test_lang_filters=py \
--test_tag_filters=-no_pip,-no_windows,-no_windows_gpu,-no_gpu,-no_pip_gpu \
--build_tag_filters=-no_pip,-no_windows,-no_windows_gpu,-no_gpu,-no_pip_gpu \
--test_env=TF_SAVER_LENIENT_NAMES=True \
--local_test_jobs=1 --build_tests_only //${PY_TEST_DIR}/tensorflow/python/...

2
tensorflow/tools/docker/Dockerfile.devel
View File

@ -72,7 +72,7 @@ RUN mkdir /bazel && \
RUN git clone https://github.com/tensorflow/tensorflow.git && \
cd tensorflow && \
git checkout r1.3
git checkout r1.4
WORKDIR /tensorflow
# TODO(craigcitro): Don't install the pip package, since it makes it

Some files were not shown because too many files have changed in this diff Show More