1 /* Copyright 2017 The TensorFlow Authors. All Rights Reserved. 2 3 Licensed under the Apache License, Version 2.0 (the "License"); 4 you may not use this file except in compliance with the License. 5 You may obtain a copy of the License at 6 7 http://www.apache.org/licenses/LICENSE-2.0 8 9 Unless required by applicable law or agreed to in writing, software 10 distributed under the License is distributed on an "AS IS" BASIS, 11 WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. 12 See the License for the specific language governing permissions and 13 limitations under the License. 14 ==============================================================================*/ 15 16 // Utilities for working with XLA layout and shapes. 17 18 #ifndef TENSORFLOW_COMPILER_MLIR_XLA_LAYOUT_UTIL_H_ 19 #define TENSORFLOW_COMPILER_MLIR_XLA_LAYOUT_UTIL_H_ 20 21 #include <functional> 22 #include <vector> 23 24 #include "tensorflow/compiler/xla/client/xla_builder.h" 25 #include "tensorflow/compiler/xla/service/hlo_sharding.h" 26 #include "tensorflow/compiler/xla/shape.h" 27 #include "tensorflow/compiler/xla/statusor.h" 28 #include "tensorflow/compiler/xla/xla_data.pb.h" 29 30 namespace mlir { 31 32 // XLA Layout preferences. Currently, when it comes to TPU, there are two 33 // primary layout choices for any XLA argumetns (parameter or resource): (1) 34 // CompactChunkPadded and (2) Linear. CompactChunkPadded is the native TPU 35 // layout while Linear is native host (CPU) layout. 36 // This enum allows the caller of XLA to progogate layout preference to the XLA 37 // compiler. 38 // kNoPreference: the generic layout where the XLA compiler has the freedom 39 // to assign any layout. 40 // kTpuPreferCompactChunkPaddedLayout: use native TPU layout on TPU. 41 // kTpuPreferLinearLayout: use native CPU layout on TPU. The compiler may 42 // insert transformation TPU kernels. 43 // As the layout of any argument will change from a native host layout to a 44 // native TPU layout either on host or on device, XLA compiler and TPU runtime 45 // must be in coordination to transform the parameters in a consistent way. 46 enum class XlaLayoutPreference { 47 kNoPreference = 0, 48 kTpuPreferCompactChunkPaddedLayout = 1, 49 kTpuPreferLinearLayout = 2 50 }; 51 52 // The following defines the layout preference of an xla tensor. 53 // The return value of LayoutPreferenceFn can be used in 54 // ShapeRepresentationFn. 55 typedef std::function<xla::StatusOr<XlaLayoutPreference>( 56 const xla::Shape& shape)> 57 LayoutPreferenceFn; 58 59 typedef std::function<xla::StatusOr<xla::Shape>( 60 const xla::Shape& shape, bool fast_mem, 61 XlaLayoutPreference layout_preference)> 62 ShapeRepresentationFn; 63 64 // Return a LayoutPreferenceFn that always uses kNoPreference layout. 65 LayoutPreferenceFn UseNoPreferenceLayoutFn(); 66 67 // Rewrites the layout of xla_shape if there is tiled sharding. 68 xla::Status RewriteLayoutWithShardedShape( 69 const std::optional<xla::HloSharding>& sharding, bool use_fast_memory, 70 const LayoutPreferenceFn& layout_preference_fn, 71 const ShapeRepresentationFn& shape_representation_fn, 72 xla::Shape* xla_shape); 73 74 // Adds reshapes to fix the layout of an output, if a shape_representation_fn or 75 // sharding is present. 76 xla::StatusOr<xla::XlaOp> ReshapeWithCorrectRepresentationAndSharding( 77 xla::XlaBuilder* builder, xla::XlaOp original, xla::Shape original_shape, 78 const LayoutPreferenceFn& layout_preference_fn, 79 const ShapeRepresentationFn& shape_representation_fn, 80 std::optional<xla::OpSharding> sharding, bool fast_mem); 81 82 } // namespace mlir 83 84 #endif // TENSORFLOW_COMPILER_TF2XLA_SHAPE_UTIL_H_ 85