Commit Graph

11986 Commits

Author SHA1 Message Date
Smit Hinsu
bf61dd6420 Disable MLIR bridge for NMS image ops test
MLIR bridge doesn't support tf.NonMaxSuppressionV4 legalization that is
conditionally generated by non_max_suppression_padded function.

PiperOrigin-RevId: 320197235
Change-Id: If7242133254680b366771ced50de074ed6180563
2020-07-08 11:02:29 -07:00
Ruoxin Sang
02ad000479 Support dynamic outputs for XLA on demand ops.
PiperOrigin-RevId: 317902879
Change-Id: I6b6dfa54855d5996ac15d4b5c48a5db5dc230025
2020-07-01 11:11:47 -07:00
Rahul Joshi
e60cf08994 [NFC] Re-instate use of FuncOp::isPublic()
PiperOrigin-RevId: 317774876
Change-Id: I6a832236377d403b1ebd24ecbc26025c37dc1c13
2020-06-22 18:24:39 -07:00
Berkin Ilbeyi
2f9984ff83 Allocate input/output buffers at the correct memory space.
PiperOrigin-RevId: 317761146
Change-Id: Ia27b6765a547451575240e76ea419f705e29bd32
2020-06-22 17:31:14 -07:00
Feng Liu
557916fa3f Fuse the back to back tfr.cast with unranked input to tf.EnsureShape
After this step, all the tfr ops should be raised to the tf ops

PiperOrigin-RevId: 317756602
Change-Id: I15731f231caf6c47eac1719fc1c2ed3d01cff515
2020-06-22 17:19:46 -07:00
Tim Shen
3980537192 [MLIR] Remove TupleSelectOp from LHLO.
LHLO uses output-parameter, but TupleSelectOp outputs into a tuple on the device. The current type constraints are wrong, and there is not enough expressiveness in LHLO to define a device-memory representation that can be passed in between kernel launches.

PiperOrigin-RevId: 317749542
Change-Id: I6d3350cbf9decf006f239a7208f6bbef0175ac61
2020-06-22 16:08:30 -07:00
Andy Ly
dc756d14ca Sync TensorFlow MLIR ODS with TensorFlow op registry.
tf.BesselI0e and tf.BesselI1e are moved to tf_ops.td as they are now python generated ops, so they do not have summaries and descriptions in the op registry. int8 and int16 support has been added to tf.Acos, tf.Atan, tf.Inv, tf.Reciprocal, tf.Round, and tf.Tan.

PiperOrigin-RevId: 317748438
Change-Id: Icb86560d5118c38c69819fd02daa3a3841e113b2
2020-06-22 16:00:20 -07:00
Yanhua Sun
b2f0928940 Add DeviceIndex xla op.
DeviceIndex op: given a list of device names, this operation returns the index of the device this op runs.  In the case of XLA, we are not executing on any device, we return the length of the list.

PiperOrigin-RevId: 317740778
Change-Id: I0679aa0adc5508b83502eee0d2044584577ed5b4
2020-06-22 15:06:38 -07:00
Feng Liu
31acd85aee Add quantization support for tfl.transpose_conv
PiperOrigin-RevId: 317725618
Change-Id: If98427ba8b732edc7e1bd30f9c0f11d565d08a84
2020-06-22 13:46:24 -07:00
George Karpenkov
9ca89a201c [TF2XLA] [NFC] Refactor datastructures for resource variables to not require snapshotting for compilation
Previously, `BuildXlaCompilerArguments` required taking a snapshot of all
resource variables in order to start compiling.
In this CL it can operate with a span of pointers to resource variables instead
(we CHECK at runtime that the lock is held).
That refactoring allows to launch the XLA compilation without creating an extra
reference to the underlying Tensor of the passed resource variables.

PiperOrigin-RevId: 317706126
Change-Id: I37a97601a08f165b23b4745e1b032bf91c21c313
2020-06-22 12:55:49 -07:00
Davide Libenzi
44067f0783 Make XRT CPU/GPU use MaybeOwning buffer interface, so the new copy protection CL won't break aliasing.
PiperOrigin-RevId: 317700747
Change-Id: Ie7b5bb1989cd4359b30ad86a450de5bff0962c31
2020-06-22 12:32:43 -07:00
Christian Sigg
8785b4f5b8 Make gpu_ftz_test ready for CUDA 11.
CUDA 11 changed the libdevice implementation of expf from previously two to one ex2.approx.ftz. Change the CHECK directives to handle both cases.

PiperOrigin-RevId: 317698875
Change-Id: Idd7b28c77427f299b80fa1b7f4b9be8c7881f963
2020-06-22 12:24:29 -07:00
Rahul Joshi
251923169d [NFC] Eliminate use of .getBlocks() when not needed
Also use llvm::hasSingleElement() instead of .size() == 1

PiperOrigin-RevId: 317675565
Change-Id: I4f0e8892957c2b20e115584fe7424da68d53b67a
2020-06-22 10:34:01 -07:00
Feng Liu
22f6939be9 Apply default quantization parameter before quantization pass
By this order, the default quantization parameter is only applied on the
activations and the weight quantization parameter will use the parameters from
the weight content.

PiperOrigin-RevId: 317672365
Change-Id: Ib7b02ae19105db124721242ea51a5ccc1d5aa68e
2020-06-22 10:05:11 -07:00
Alex Zinenko
bb73be3e36 [mlir][xla] LHLO-to-Affine: simplify DotOp conversion with nested builders
MLIR recently introduced a new idiom for constructing loop nests. Use it to
make the legalization of LHLO to affine loops more concise and readable.

PiperOrigin-RevId: 317649515
Change-Id: Idfab27b4655d6df90d940fb7b064ea9941d8a700
2020-06-22 07:28:39 -07:00
Stephan Herhut
12f5cd7dde Propagate noalias and alignment properties of TensorFlow ABI into kernels.
PiperOrigin-RevId: 317646198
Change-Id: I20a649c3127586106c250e213e6b6700fb302495
2020-06-22 07:13:58 -07:00
Tamara Norman
0868ca7bb2 Allow a shape to be passed to CopyToHostAsync
PiperOrigin-RevId: 317611333
Change-Id: I4526f9dbd1b223eb23fe928326afca0eb133c2f5
2020-06-22 01:55:19 -07:00
Andy Ly
6d2ce43b03 Remove empty description fields in tf_generated_ops.td and tf_ops.td. (NFC)
PiperOrigin-RevId: 317466704
Change-Id: Ic0bbec0c7013c2f2238a4f4e5763632c846f1337
2020-06-20 08:40:03 -07:00
Sanjoy Das
7f2bfd5709 Have the HloEvaluator print out intermediate values on VLOG(100)
This is useful for tracing through very small HLO programs.

PiperOrigin-RevId: 317432566
Change-Id: I1b70ed0f81ed57915ecccaf8cb85fefc857e2151
2020-06-19 22:14:31 -07:00
David Majnemer
728a4a4405 Unbreak the MSVC build
Don't use inline variables, our MSVC builds do not like them.

Instead, simulate them using inline functions + function-scope static variables.

PiperOrigin-RevId: 317428053
Change-Id: Icf8838c159ab9f132ad32360633046f4c2224a79
2020-06-19 21:15:11 -07:00
Smit Hinsu
b1fdd334e7 Rewrite masks while applying ellipsis mask to tf.StridedSlice op
begin_mask and end_mask needs to be adjusted while simplifying StridedSlice op by removing ellipsis_mask. new_axis_mask is already removed before this and we don't yet support shrink_axis_mask.

PiperOrigin-RevId: 317419960
Change-Id: Ie4a5f404f95f5b909065311a54cbbed64d4ccf4b
2020-06-19 19:28:31 -07:00
Andy Ly
f3cb68f0ac Remove assertions checking device map size in TPUVariableRuntimeReformattingPass.
It is possible for more than one replicated device to be populated (e.g. replicated host).

PiperOrigin-RevId: 317419924
Change-Id: I1d8fa95654f324557d78adf4570adc6c3cfdabb4
2020-06-19 19:22:40 -07:00
George Karpenkov
aa7ff6aa28 [TF2XLA] Set up aliasing for resource variables even when not returning a tuple
PiperOrigin-RevId: 317414582
Change-Id: I45cd1f314331cb86a0257e7b7cf9d0639be84e99
2020-06-19 18:20:41 -07:00
Sanjoy Das
f840a62268 Rollback "[TF:TRT] Cosmetic fix."
PiperOrigin-RevId: 317407274
Change-Id: I73cd486acf9091e6678e553ab9b0545288f73324
2020-06-19 17:15:15 -07:00
George Karpenkov
6116b7f911 [XLA] [client] Implement a RunAsync overload which does not need a vector of shapes
PiperOrigin-RevId: 317406952
Change-Id: I69d8cc8a68ffdfbf70e2969f5df5e6adba7d2e1d
2020-06-19 17:09:13 -07:00
George Karpenkov
62683d061c [XLA] Rollback of rollback of "Implement LocalClient::Run which supports buffer donation"
PiperOrigin-RevId: 317400695
Change-Id: I56f1f8df347d5a3b2bad9526c7315c63ad6ddadb
2020-06-19 16:33:47 -07:00
Kuangyuan Chen
e3809400b2 Import initialization graph in SignatureDef SavedModels as an MLIR function in
TF saved model dialect.
-Mark the init function referenced by SessionInitializerOp as an exported
function with the reserved name "__tf_saved_model_session_initializer".
-Remove variable init logic in init function if it is already imported as a global tensor.
-Add a canonicalizer to SessionInitializerOp to remove empty init functions.
-Return error on SessionInitializerOp in IREE compiler pipeline.

PiperOrigin-RevId: 317395165
Change-Id: Idb7e54cac08add9fc8f2ccfbaf341135fdf59e3b
2020-06-19 15:54:34 -07:00
Jacques Pienaar
90443c2b17 Enable test for generating tf_ops
PiperOrigin-RevId: 317389351
Change-Id: I54f6ac53a974cf603ed5fe1a30b6fbb464c80d28
2020-06-19 15:19:10 -07:00
Bixia Zheng
f129485019 [TF:TRT] Cosmetic fix.
Rewrite two lines of #if into #if GOOGLE_CUDA && GOOGLE_TENSORRT.

PiperOrigin-RevId: 317386436
Change-Id: Icc8ae27a17900b6f0a198d32c6d73345084eab50
2020-06-19 15:04:44 -07:00
Sanjoy Das
94bf57d06c Do not try to compile trivially dead branches in the Case tf2xla lowering
This is important for the upcoming DeviceIndex op which can be used to select
one of many implementations depending on the device, and some of them may not be
compilable by tf2xla.

PiperOrigin-RevId: 317376420
Change-Id: I6428df6f4da238e5d2bc3618d51c579e34454945
2020-06-19 14:18:26 -07:00
Kuangyuan Chen
2ec0214b48 Internal rollback for saved model importer.
PiperOrigin-RevId: 317354577
Change-Id: I6fdc05921b2f1d85b34b1751ed00f65525d45ad3
2020-06-19 12:15:52 -07:00
Yunxing Dai
a8456eae42 Check the value shape of set dimension size in shape inference.
PiperOrigin-RevId: 317351738
Change-Id: Ia185e7745753711ca9ebf657522ef7422c9696ca
2020-06-19 12:07:45 -07:00
A. Unique TensorFlower
2417a15bf1 Clean up post fusion using DCE.
PiperOrigin-RevId: 317349120
Change-Id: I479d9967323d86e924315d2b1302bafd01ed4151
2020-06-19 11:52:26 -07:00
A. Unique TensorFlower
9aea03e98d Add StatefulRandom op to TF MLIR.
PiperOrigin-RevId: 317347849
Change-Id: I19b7b1b3157d065cc556cce7aba72b79ed9b776f
2020-06-19 11:42:56 -07:00
Bixia Zheng
9152edc1f0 [TF:TRT] Add flag TF_TRT_ALLOW_ENGINE_NATIVE_SEGMENT_EXECUTION.
The default value of the flag is True. When the flag value is false, the
bridge will report an error when the native segment of a TRTEngineOp is
executed.

Add test cases.

PiperOrigin-RevId: 317340632
Change-Id: Iacded09b38e63442bbd93076a079d385fb8a77e6
2020-06-19 11:09:33 -07:00
Rahul Joshi
1332b9365e [NFC] Adopt FuncOp::isPublic() and friends
PiperOrigin-RevId: 317335755
Change-Id: I6f1b25798150bb3a255d4572831d3be4747d28c7
2020-06-19 10:52:28 -07:00
Thomas Joerg
ef1cabc7a8 [XLA:GPU] Split reduce ops with large but non-consecutive reduction dimensions.
PiperOrigin-RevId: 317330616
Change-Id: Icdcf320b233479c2f74c5b40ee4c8d9a73a6088a
2020-06-19 10:35:29 -07:00
Bixia Zheng
57f9d638c0 [TF:TRT] Add a prefix to the warning messages from TF-TRT.
Add LOG_WARNING_WITH_PREFIX to common/utils.h. Replace the use of LOG(WARNING)
with this new macro.

PiperOrigin-RevId: 317330336
Change-Id: Ife0aa0347dd72f6eb0f8805af4d46a7d4cb099ea
2020-06-19 10:31:29 -07:00
Tamara Norman
07c54454ee Add an option such that the cached host_value can be discarded
PiperOrigin-RevId: 317315157
Change-Id: I9d7145390a526003069321c7e04794e139a53c09
2020-06-19 08:56:31 -07:00
T.J. Alumbaugh
9e7d5ef6f2 Full int8 quantization BatchMatMul
PiperOrigin-RevId: 317304259
Change-Id: Icf96d9d129db30b965e36f5c8befd27762b173b2
2020-06-19 07:33:08 -07:00
Hanhan Wang
0c7e61d660 Remove the canonicalize pattern for folding a pad op into the following conv op.
Basically rolledback for cl/305641881, the pattern could hurt performance
because the operation can't be fully tiled in Linalg transformation. In this
context, not everyone wants this pattern, so remove it from canonicalize
patterns.

PiperOrigin-RevId: 317302072
Change-Id: I19aa64e14eecccfd738ad3f775f3670974bc68f9
2020-06-19 07:14:26 -07:00
Alexander Belyaev
b58e600045 [XLA][MLIR] Enable xla_hlo.ReshapeOp -> xla_lhlo.ReshapeOp conversion.
PiperOrigin-RevId: 317284676
Change-Id: Ia845183efcfabe77f6eb66d8c56dcbfc82653982
2020-06-19 04:30:25 -07:00
Smit Hinsu
c662daf489 Override CustomCall in MlirHloBuilder
Also, enable mlir bridge for image ops compilers test. ResizeBilinear op
lowering usese CustomCall in case of TPU lowerings.

PiperOrigin-RevId: 317272443
Change-Id: I134c828cdc76552a0cbfdeb7c65532aa986314e2
2020-06-19 02:26:50 -07:00
A. Unique TensorFlower
9d1ec55aed Integrate LLVM at https://github.com/llvm/llvm-project/commit/7f0d7f326316
PiperOrigin-RevId: 317270655
Change-Id: Ic80ab697da45212c8d58bcda989e5ee0a330b565
2020-06-19 02:19:17 -07:00
George Karpenkov
9f20b156bc [XLA:GPU] [NFC] Clarify the precondition for the fast reduction emitter
PiperOrigin-RevId: 317266013
Change-Id: I384acac279f0db53f195d5b43318c38c87a1739c
2020-06-19 01:17:59 -07:00
A. Unique TensorFlower
397494a231 Hoisting unconditional converts from conditional branch computations.
PiperOrigin-RevId: 317239618
Change-Id: If3b16ff4f2bbcf38ee1ca51f5e8b187c58ab8e91
2020-06-18 20:58:13 -07:00
Gaurav Jain
85ad8031f6 Expand dtype support for Neg
PiperOrigin-RevId: 317237033
Change-Id: I59c5e45d469f7bf704976b66bc122aaac3982b5e
2020-06-18 20:31:51 -07:00
David Majnemer
4a14e778d6 [XLA] Introduce ManifestCheckingTest
PiperOrigin-RevId: 317229603
Change-Id: Ibcc9ea3895d520024f5d80d52330aeb3b970585d
2020-06-18 19:20:02 -07:00
A. Unique TensorFlower
723751b20e Add set_outfeed_config in XLA HloInstruction.
PiperOrigin-RevId: 317222410
Change-Id: I5de8a5067f1002a9d656d4e26d145ffe3fe372ed
2020-06-18 18:19:13 -07:00
Lucy Fox
83b09270dc Update FusedKernelMatcher pass to use upstream util to get stripped op name.
Added this util upstream in D81435, so now using that instead and deleting the unneeded code here.

PiperOrigin-RevId: 317209256
Change-Id: Id2d8a1fca34ca85e59a05a85bf7f6f59b425c7c1
2020-06-18 16:57:18 -07:00