1 /* Copyright 2019 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 #include "tensorflow/compiler/xla/service/optimize_input_output_buffer_alias.h" 16 17 #include <queue> 18 #include <string> 19 20 #include "absl/container/flat_hash_map.h" 21 #include "absl/memory/memory.h" 22 #include "absl/strings/str_format.h" 23 #include "tensorflow/compiler/xla/shape_util.h" 24 #include "tensorflow/compiler/xla/status_macros.h" 25 #include "tensorflow/compiler/xla/util.h" 26 #include "tensorflow/core/lib/core/errors.h" 27 #include "tensorflow/core/platform/types.h" 28 29 namespace xla { 30 namespace { 31 32 // Returns true if the given shape is a non-nested tuple. 33 bool IsNonNestedTuple(const Shape& shape) { 34 return shape.IsTuple() && !ShapeUtil::IsNestedTuple(shape); 35 } 36 37 } // namespace 38 39 StatusOr<bool> OptimizeInputOutputBufferAlias::Build( 40 const Shape& input_shape, const Shape& output_shape, 41 HloInputOutputAliasConfig* alias_config) { 42 bool changed = false; 43 TF_RET_CHECK(LayoutUtil::HasLayout(input_shape)); 44 TF_RET_CHECK(LayoutUtil::HasLayout(output_shape)); 45 VLOG(1) << "input_shape:" << input_shape.ToString(); 46 VLOG(1) << "output_shape:" << output_shape.ToString(); 47 48 // For all buffers defined by the parameter, build a map from the byte 49 // size to the list of the buffers of that size. 50 absl::flat_hash_map<int64, std::queue<ShapeIndex>> size_to_input_index; 51 ShapeUtil::ForEachSubshape( 52 input_shape, [&](const Shape& subshape, const ShapeIndex& index) { 53 if (subshape.IsTuple()) { 54 return; 55 } 56 int64 bytes = size_func_(subshape); 57 size_to_input_index[bytes].push(index); 58 }); 59 60 // For each result buffer shape index, take the first unused parameter 61 // buffer that matches the size. 62 TF_RETURN_IF_ERROR(ShapeUtil::ForEachSubshapeWithStatus( 63 output_shape, [&](const Shape& subshape, const ShapeIndex& index) { 64 if (subshape.IsTuple()) { 65 return Status::OK(); 66 } 67 int64 bytes = size_func_(subshape); 68 69 auto it = size_to_input_index.find(bytes); 70 if (it != size_to_input_index.end() && !it->second.empty()) { 71 changed = true; 72 const ShapeIndex& input_index = it->second.front(); 73 const ShapeIndex& output_index = index; 74 if (!alias_config->ParameterHasAlias(0, input_index) && 75 !alias_config->OutputHasAlias(output_index)) { 76 TF_RETURN_IF_ERROR(alias_config->SetUpAlias( 77 output_index, 0, input_index, 78 HloInputOutputAliasConfig::AliasKind::kSystemAlias)); 79 } 80 VLOG(3) << "Set up alias from with param index " 81 << it->second.front().ToString() << ", shape size " << bytes 82 << " and result subshape " 83 << ShapeUtil::HumanStringWithLayout(subshape) << " at index " 84 << index.ToString(); 85 it->second.pop(); 86 } 87 return Status::OK(); 88 })); 89 return changed; 90 } 91 92 StatusOr<bool> OptimizeInputOutputBufferAlias::Run(HloModule* module) { 93 // User buffer alias only work for modules with 1 parameter. 94 if (module->entry_computation()->num_parameters() != 1) { 95 return false; 96 } 97 98 HloInputOutputAliasConfig* alias_config = 99 &module->input_output_alias_config(); 100 101 return Build(module->entry_computation()->parameter_instruction(0)->shape(), 102 module->entry_computation()->root_instruction()->shape(), 103 alias_config); 104 } 105 106 } // namespace xla 107
